tvm_tirx.backend.cuda.op.tcgen05_encode_instr_descriptor

tvm_tirx.backend.cuda.op.tcgen05_encode_instr_descriptor#

tvm_tirx.backend.cuda.op.tcgen05_encode_instr_descriptor(desc, *, d_dtype, a_dtype, b_dtype, M, N, K, trans_a, trans_b, n_cta_groups=1, neg_a=False, neg_b=False, sat_d=False, is_sparse=False)#

TVM intrinsic to create instruction descriptor for tcgen05 MMA without block scaling

Parameters:

  • desc (Expr) – The pointer to the instruction descriptor.

  • d_dtype (str) – The datatype of resultant matrix D.

  • a_dtype (str) – The datatype of multiplicand matrix A.

  • b_dtype (str) – The datatype of multiplicand matrix B.

  • M (int) – The size of non-reduction dimension of Matrix A.

  • N (int) – The size of non-reduction dimension of Matrix B.

  • K (int) – The size of reduction dimension of Matrix A/B.

  • trans_a (bool) – Whether the multiplicand matrix A is transposed. True for M/N major, False for K major.

  • trans_b (bool) – Whether the multiplicand matrix B is transposed. True for M/N major, False for K major.

  • n_cta_groups (int) – The number of CTA groups involved in the MMA operation.

  • neg_a (bool) – Whether to negate the multiplicand matrix A.

  • neg_b (bool) – Whether to negate the multiplicand matrix B.

  • sat_d (bool) – Whether to saturate the resultant matrix D.

  • is_sparse (bool) – Whether the MMA operation is sparse.