mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Diao Han
78a4ff6587
* [BugFix][Hopper] Select the widest legal WGMMA N instead of gcd block_N extents that are legal single WGMMA widths but not powers of two (96, 112, 160, 176) were decomposed into n32 / n16 instructions, because the instruction width came from gcd(warp_col_tiles, 256). wgmma.h already instantiates every multiple of 8 up to 256, so only the width selection was the constraint. Select the widest legal width that tiles warp_col_tiles exactly, keeping gcd as the fallback. Widening is restricted to multiples of 16: the N % 16 == 8 class already miscomputes (#2593), and the integer instruction tables only instantiate multiples of 16 (plus 8 and 24). The result is never narrower than before, so no existing shape regresses. The sparse generator carried the same line and now shares the helper. * [Testing] Cover T.gemm_sp at the widened WGMMA N values The width selector is shared with the sparse generator, so assert the emitted wgmma_sp instruction and the sparse numeric result for block_N in {96, 112, 160} as well. Verified as a real regression: with the selector reverted the sparse path emits n32 / n16 / n32 exactly like the dense one.