This document is relevant for: Inf2, Trn1, Trn1n

nki.isa.tensor_tensor#

nki.isa.tensor_tensor(data1, data2, op, dtype=None, mask=None, **kwargs)[source]#

Perform an element-wise operation of input two tiles using Vector Engine. The two tiles must have the same partition axis size and the same number of elements per partition. All input and output tiles can be in either SBUF or PSUM.

The element-wise operator is specified using the op field and can be any binary operator supported by NKI (see Supported Math Operators for details). For bitvec operators, the input/output data types must be integer types and Vector Engine treats all input elements as bit patterns without any data type casting. For arithmetic operators, there is no restriction on the input/output data types, but the engine automatically casts input data types to float32 and performs the element-wise operation in float32 math. The float32 results are cast to the target data type specified in the dtype field before written into the output tile. If the dtype field is not specified, it is default to be the same as the data type of data1 or data2, whichever has the highest precision.

Note, if you need broadcasting capability in the free dimension for either input tile, you should consider using nki.isa.tensor_scalar API instead, which has better performance than nki.isa.tensor_tensor in general.

Estimated instruction cost:

Given N is the number of elements per partition in data1/data2,

  • If one input tile is in PSUM and the other is in SBUF, N Vector Engine cycles

  • If all of the below conditions are met, also N Vector Engine cycles: - both input tiles are in SBUF, - input/output data types are all bfloat16, - the operator is add, multiply or subtract, - Input tensor data is contiguous along the free dimension (that is, stride in each partition is 1 element)

  • Otherwise, 2N Vector Engine cycles

Parameters:

  • data1 – lhs input operand of the element-wise operation

  • data2 – rhs input operand of the element-wise operation

  • op – a binary math operator (see Supported Math Operators for supported operators)

  • mask – (optional) a compile-time constant predicate that controls whether/how this instruction is executed (see NKI API Masking for details)

  • dtype – (optional) data type to cast the output type to (see Supported Data Types for more information); if not specified, it will default to be the same as the data type of the input tiles, or whichever input type has the highest precision (see NKI Type Promotion for more information);

Returns:

an output tile of the element-wise operation

Example:

import neuronxcc.nki.isa as nisa
import neuronxcc.nki.language as nl
import numpy as np
...
  ##################################################################
  # Example 1: add two tiles, a and b, of the same
  # shape (128, 512) element-wise and get
  # the addition result in tile c
  ##################################################################
  i_p = nl.arange(128)[:, None]
  i_f = nl.arange(512)[None, :]
  c = nisa.tensor_tensor(a[i_p, i_f], b[i_p, i_f], np.add)

This document is relevant for: Inf2, Trn1, Trn1n