Indicates whether the types \(L\) (the left matrix), \(R\) (the right matrix) and \(A\) (the accumulator matrix) can participate in a matrix multiply accumulate operation.

The types \(L\), \(R\) and \(A\) can participate in a matrix multiply accumulate if

1. The element types of \(L\) and \(R\) are integral scalars of the same bitwidth or they are floating point scalars of the same conversion rank

2. The element types of \(L\), \(R\) and \(A\) are selected from a row in the following table:

   Element Type of \(L\) and \(R\)	Element type of \(A\)
   \(8\) bit integral types	\(32\) bit signed integral types
   __nv_fp8_e4m3 __nv_fp8_e5m2 __half	__half float
   __nv_bfloat16 __nv_tf32 float	float
   double	double

3. All of \(L\), \(R\) and \(A\) have the same rank and the rank is either \(2\) or \(3\).

4. In the case where the rank is \(2\), the shapes of \(L\), \(R\) and \(A\) are taken from the table for some integers \(N\), \(K\) and \(M\):

   Shape of \(L\)	Shape of \(R\)	Shape of \(A\)
   \(N \times K\)	\(K \times M\)	\(N \times M\)

5. In the case where the rank is \(3\), the shapes of \(L\), \(R\) and \(A\) are taken from the table for some integers \(A\), \(B\), \(C\), \(N\), \(K\) and \(M\):

   Shape of \(L\)	Shape of \(R\)	Shape of \(A\)
   \(A \times N \times K\)	\(B \times K \times M\)	\(C \times N \times M\)

   Additionally \(A\) and \(B\) satisfy the following predicates:

   \[\begin{split}\begin{cases} A = C & \text{or} \\ A = 1 & \end{cases} \quad \quad \begin{cases} B = C & \text{or} \\ B = 1 & \end{cases}\end{split}\]

Indicates whether the types \(L\) and \(R\) may participate in a matrix multiply operation.

The types \(L\) and \(R\) may participate in a matrix multiply operation if:

1. The element types of \(L\) and \(R\) are integral scalars whose bitwidth is \(8\) or they are floating point scalars of the same conversion rank.

2. \(L\) and \(R\) have the same rank and that rank is either \(2\) or \(3\).

3. In the case where the rank is \(2\), the shapes of \(L\) and \(R\) are taken from the table below for some integers \(N\), \(K\) and \(M\):

   Shape of \(L\)	Shape of \(R\)
   \(N \times K\)	\(K \times M\)

   The hypothetical ct::extents specialization of dimensions \(N \times M\) must be tile compatible.

4. In the case where the rank is \(3\), the shapes of \(L\) and \(R\) are taken from the table below for some integers \(A\), \(B\), \(N\), \(K\) and \(M\):

   Shape of \(L\)	Shape of \(R\)
   \(A \times N \times K\)	\(B \times K \times M\)

   Additionally, either \(A = B\) or at least one of them is equal to \(1\). The hypothetical ct::extents specialization whose dimensions are \(\operatorname{max}(A, B) \times N \times M\) must be tile compatible.

Yields the result type of performing matrix multiplication on operands of types \(L\) and \(R\)

The result type is a specialization of ct::tile. The element type of the result is determined by the element types of \(L\) and \(R\) according to the following table:

The shape of the result depends on the rank of \(L\) and \(R\):

1. If \(L\) and \(R\) have rank is \(2\) with shapes \(N \times K\) and \(K \times M\) respectively, the result shape is \(N \times M\).

2. If \(L\) and \(R\) have rank \(3\) with shapes \(A \times N \times K\) and \(B \times K \times M\), the result shape is \(C \times N \times M\) where \(C\) is the larger of \(A\) and \(B\).

Performs matrix multiply on lhs and rhs and adds the result to acc. For rank \(3\) arguments, the matrix multiply accumulate is performed for each set of matrices along the first dimension of the operands.

If the operands have rank \(3\), the lhs undergoes broadcast conversion to the shape \(A_0 \times L_1 \times L_2\) and rhs undergoes broadcast conversion to the shape \(A_0 \times R_1 \times R_2\).

Let \(a\), \(b\) and \(c\) denote the converted operands lhs, rhs, and acc respectively.

For rank \(2\) arguments, the result is an approximation of the matrix \(r\) determined by

For rank \(3\) arguments, the result is an approximation of the matrix \(r\) determined by

namespace ct = ::cuda::tiles;
using i32x2x4 = ct::tile<int, ct::shape<2, 4>>;
using i32x4x2 = ct::tile<int, ct::shape<4, 2>>;
using i32x2x2 = ct::tile<int, ct::shape<2, 2>>;

auto lhs = ct::element_cast<float>(ct::iota<i32x2x4>());
auto rhs = ct::element_cast<float>(ct::iota<i32x4x2>());
auto acc = ct::element_cast<float>(ct::iota<i32x2x2>());

auto r = ct::mma(lhs, rhs, acc);

namespace ct = ::cuda::tiles;
using i32x1x2x4 = ct::tile<int, ct::shape<1, 2, 4>>;
using i32x1x4x2 = ct::tile<int, ct::shape<1, 4, 2>>;
using i32x1x2x2 = ct::tile<int, ct::shape<1, 2, 2>>;

auto lhs = ct::element_cast<float>(
             ct::cat<0>(ct::iota<i32x1x2x4>(), ct::iota<i32x1x2x4>()));
auto rhs = ct::element_cast<float>(
             ct::cat<0>(ct::iota<i32x1x4x2>(), -ct::iota<i32x1x4x2>()));
auto acc = ct::element_cast<float>(
             ct::cat<0>(ct::iota<i32x1x2x2>(), -ct::iota<i32x1x2x2>()));

auto r = ct::mma(lhs, rhs, acc);

Performs matrix multiply on lhs and rhs. For rank \(3\) arguments, the matrix multiply is performed for each set of matrices along the first dimension of the operands.

If the operands have rank \(3\), the lhs undergoes broadcast conversion to the shape \(\operatorname{max}(L_0, R_0) \times L_1 \times L_2\) and rhs undergoes broadcast conversion to the shape \(\operatorname{max}(L_0, R_0) \times R_1 \times R_2\).

Let \(a\) and \(b\) denote the converted operands lhs and rhs respectively.

For rank \(2\) arguments, the result is an approximation of the matrix \(r\) determined by

For rank \(3\) arguments, the result is an approximation of the matrix \(r\) determined by

namespace ct = ::cuda::tiles;
using i32x2x4 = ct::tile<int, ct::shape<2, 4>>;
using i32x4x2 = ct::tile<int, ct::shape<4, 2>>;

auto lhs = ct::element_cast<float>(ct::iota<i32x2x4>());
auto rhs = ct::element_cast<float>(ct::iota<i32x4x2>());

auto r = ct::matmul(lhs, rhs);

namespace ct = ::cuda::tiles;
using i32x1x2x4 = ct::tile<int, ct::shape<1, 2, 4>>;
using i32x1x4x2 = ct::tile<int, ct::shape<1, 4, 2>>;

auto lhs = ct::element_cast<float>(ct::cat<0>(ct::iota<i32x1x2x4>(), ct::iota<i32x1x2x4>()));
auto rhs = ct::element_cast<float>(ct::cat<0>(ct::iota<i32x1x4x2>(), -ct::iota<i32x1x4x2>()));

auto r = ct::matmul(lhs, rhs);

Performs block scaled matrix multiply and accumulate on the operands lhs, rhs, and acc using the scale factors specified in lhsScale and rhsScale. For rank \(3\) operands, the matrix multiply is batched along the first dimension and all inputs are broadcasted to match the batch length of the accumulator.

The left and right matrices are subdivided into vectors. Each vector is associated with a separate scaling factor which is used to scale the values in the vector before performing the matrix multiply. The number of elements which correspond to a single scaling factor is called the vector length.

For rank \(2\) arguments, the result is an approximation of the matrix \(r\) determined by

where \(V = \frac{L_1}{LS_1}\) indicates the vector length.

For rank \(3\) operands, the inputs undergo broadcast conversion as follows:

• lhs is broadcasted to shape \(A_0 \times L_1 \times L_2\)

• lhsScale is broadcasted to shape \(A_0 \times LS_1 \times LS_2\)

• rhs is broadcasted to shape \(A_0 \times R_1 \times R_2\)

• rhsScale is broadcasted to shape \(A_0 \times RS_1 \times RS_2\)

The result is an approximation of the matrix \(r\) determined by the following equation on the converted operands:

where \(V = \frac{L_2}{LS_2}\) indicates the vector length.

The atomic constraint validates the following:

Shape Constraints

All arguments have the same rank and the rank is \(2\) or \(3\).

For rank \(2\) operands, the shape of each operand has the following format:

for some integers \(N\), \(M\), \(S\), and \(V\). The value of \(V\) (the vector size) must be consistent with the type constraint table below.

For rank \(3\) operands, the shape of each operand has the following format:

For some integers \(N\), \(M\), \(S\), \(V\), \(B_0\), \(B_1\), \(B_2\), \(B_3\), and \(B\). The value of \(V\) (the vector size) must be consistent with the type constraint table below. For each \(B_i\), either \(B_i = B\) or \(B_i = 1\).

Type Constraints

The element types of \(L\) and \(R\) must be the same. The element types of \(LS\) and \(RS\) must be the same. The element type of \(A\) must be float.

The element types and vector size, \(V\), must be selected from one of the rows in the following table:

namespace ct = ::cuda::tiles;
using e5m2 = __nv_fp8_e5m2;
using e8m0 = __nv_fp8_e8m0;

e5m2 lhsData[2][64] = {};
lhsData[0][0]  = e5m2(2.0f);   lhsData[0][1]  = e5m2(3.0f);
lhsData[1][0]  = e5m2(4.0f);   lhsData[1][1]  = e5m2(6.0f);
lhsData[0][32] = e5m2(0.5f);   lhsData[0][33] = e5m2(2.0f);
lhsData[1][32] = e5m2(-1.0f);  lhsData[1][33] = e5m2(4.0f);

e8m0 lhsScaleData[2][2] = {
    {e8m0(1.0f), e8m0(2.0f)},
    {e8m0(4.0f), e8m0(8.0f)}
};

e5m2 rhsData[64][2] = {};
rhsData[0][0]  = e5m2(-3.0f);  rhsData[0][1]  = e5m2(2.0f);
rhsData[1][0]  = e5m2(0.5f);   rhsData[1][1]  = e5m2(1.5f);
rhsData[32][0] = e5m2(-2.0f);  rhsData[32][1] = e5m2(2.0f);
rhsData[33][0] = e5m2(2.0f);   rhsData[33][1] = e5m2(0.5f);

e8m0 rhsScaleData[2][2] = {
    {e8m0(16.0f), e8m0(1.0f)},
    {e8m0(0.5f),  e8m0(4.0f)}
};

float accData[2][2] = {
  {3.0f, -9.0f},
  {5.0f, 7.0f}
};

using i32x2x64 = ct::tile<int, ct::shape<2, 64>>;
using i32x2x2 = ct::tile<int, ct::shape<2, 2>>;
using i32x64x2 = ct::tile<int, ct::shape<64, 2>>;

auto lhs = ct::load(&lhsData[0][0] + ct::iota<i32x2x64>());
auto lhsScale = ct::load(&lhsScaleData[0][0] + ct::iota<i32x2x2>());
auto rhs = ct::load(&rhsData[0][0] + ct::iota<i32x64x2>());
auto rhsScale = ct::load(&rhsScaleData[0][0] + ct::iota<i32x2x2>());
auto acc = ct::load(&accData[0][0] + ct::iota<i32x2x2>());

auto r = ct::scaled_mma(lhs, lhsScale, rhs, rhsScale, acc);