Why cutlass don't have nn layout gemm for int8 tensorcore?

[LeiWang1999]

only have TN layout I think

cutlass_tensorop_i8816gemm_s8_256x128_64x2_tn_align16
cutlass_tensorop_i8816gemm_s8_128x256_64x2_tn_align16
cutlass_tensorop_i8816gemm_s8_128x128_64x2_tn_align16
cutlass_tensorop_i8816gemm_s8_64x256_64x2_tn_align16
cutlass_tensorop_i8816gemm_s8_256x64_64x2_tn_align16
cutlass_tensorop_i8816gemm_s8_64x128_64x2_tn_align16
cutlass_tensorop_i8816gemm_s8_128x64_64x2_tn_align16
cutlass_tensorop_i8816gemm_s8_64x64_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_256x128_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_128x256_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_128x128_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_64x256_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_256x64_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_64x128_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_128x64_64x2_tn_align16
cutlass_tensorop_s8_i8816gemm_s8_64x64_64x2_tn_align16
cutlass_tensorop_i16832gemm_s8_256x128_64x3_tn_align16
cutlass_tensorop_i16832gemm_s8_128x256_64x3_tn_align16
cutlass_tensorop_i16832gemm_s8_256x64_64x4_tn_align16
cutlass_tensorop_i16832gemm_s8_64x256_64x4_tn_align16
cutlass_tensorop_i16832gemm_s8_128x128_64x5_tn_align16
cutlass_tensorop_i16832gemm_s8_128x64_64x6_tn_align16
cutlass_tensorop_i16832gemm_s8_64x128_64x6_tn_align16
cutlass_tensorop_i16832gemm_s8_64x64_64x10_tn_align16
cutlass_tensorop_i16832gemm_s8_128x64_128x3_tn_align16
cutlass_tensorop_i16832gemm_s8_64x128_128x3_tn_align16
cutlass_tensorop_i16832gemm_s8_64x64_128x5_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_256x128_64x3_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_128x256_64x3_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_256x64_64x4_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_64x256_64x4_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_128x128_64x5_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_128x64_64x6_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_64x128_64x6_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_64x64_64x10_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_128x64_128x3_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_64x128_128x3_tn_align16
cutlass_tensorop_s8_i16832gemm_s8_64x64_128x5_tn_align16

[hwu36]

tensor core instruction is TN layout. ldmatrix can only transpose 16bit data. that's why we do not support it.

if you need to run NN layout, you can use the non tensor core version.

[LeiWang1999]

Hello there, tensor core instruction is TN layout, is that when we said TN layout, we mean T is matrix B, N is matrix A? because as the docs described, tensor core instruction is row.col layout , which means A is row major and B is col major/

[jackkosaian]

Hi, @LeiWang1999. As discussed here, TN means that operand A is in mode "T" (transposed; row major) and operand B is in mode "N" (non-transposed; column major). This is also the requirement listed by cuBLASLt (see here, starting with "To use IMMA kernels").