mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Chennes
b7489e2608
The SM75 GEMM emitter derives MMA shapes from input dtype bits alone, so it already emits the (16,8,8) shape for fp16->fp16 accumulation (identical to the registered fp16->fp32 case) and the (8,8,16) shape for uint8 inputs (identical to signed int8). But the C++ dispatcher table only covered the three configs tested by #2198 (fp16->fp32, int8, int4), so requesting fp16->fp16-accum or uint8 GEMM on Turing failed at nvcc with an opaque `tl::mma_sync: unsupported configuration` static_assert. Add the two missing tl::detail SM75 atoms and register their dispatchers, following the precedent set by #2198's hand-written int4/uint4 structs (CuTe ships only fp16->fp32 and signed-int8 SM75 atoms). The uint8 atom uses the non-saturating PTX form to match the existing signed-int8 path. Verified on a TITAN RTX (SM75): emitted PTX is mma.sync.aligned.m16n8k8.row.col.f16.f16.f16.f16 and mma.sync.aligned.m8n8k16.row.col.s32.u8.u8.s32, and end-to-end GEMM numerics match torch.