[QST] FP4 Tensor Core mma.sync Instruction Unsupported on SM_101 Architecture

[KarlDe1]

The current kernel uses the following FP4 Tensor Core instruction:
mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3

This instruction is only supported on SM_120+ (Hopper/Blackwell architectures). Our target platform is SM_101, which does not have native FP4 Tensor Core support, so the instruction cannot execute.

Is there a recommended method to emulate this FP4 Tensor Core MMA on SM_101 while maintaining equivalent numerical results?

[depaulmillz]

SM101 supports the tcgen05 instructions. These can be used instead to do MXFP4 or NVFP4 compute.

[KarlDe1]

SM101 supports the tcgen05 instructions. These can be used instead to do MXFP4 or NVFP4 compute.

@depaulmillz Thanks for your reply,

Yes, I know that sm101 supports tcgen05 instructions, but according to the PTX 8.7 documentation corresponding to sm101, it appears that tcgen05 does not support the m16n8k64 tile shape?

[depaulmillz]

That is correct tcgen05 does not support 16x8x64 tile shape. The smallest tile shape that is supported is a 128x8x64. Are you specifically looking for a small tile along just N or this exact tile size?

[KarlDe1]

That is correct tcgen05 does not support 16x8x64 tile shape. The smallest tile shape that is supported is a 128x8x64. Are you specifically looking for a small tile along just N or this exact tile size?

Yes, I’m specifically looking for this exact tile size 16x8x64.

[depaulmillz]

I'm not sure what approach will give the best perf (the best approach will likely vary based on the problem shape), but I would suspect using 128x8x64 tcgen05 instructions and padding a 16x8x64 tile in TMEM would be a good start. This will require a different approach to pipelining where A, SFA, and SFB need to be staged in TMEM through TMEM store instructions rather than copy.