feat(sm120): add tileN = 8,16 for the dense F8F6F4 GEMM

[sunnypatell]

sm120_rr_smem_copy_selector_B has taken a TileShapeN parameter since #3292, and uses it to pick between the 1-, 2- and 4-matrix ldmatrix atoms. The block-scaled builder passes the real tile width. The dense F8F6F4 builder never does, so it always resolves to the default of 32 and takes the 4-matrix atom no matter how narrow the tile is.

The result is that any tile narrower than 32 fails to build, and it fails with a message that never names the tile as the cause:

cute/atom/copy_atom.hpp(206): error: static assertion failed with
  "TiledCopy uses too few vals for selected CopyAtom"

This forwards size<1>(TileShape_MNK{}) exactly the way the block-scaled builder already does. 128x16 and 128x8 go from not building to building and computing correctly. It's the same narrow-tile capability #3176 and #3292 added on the block-scaled side, which the dense path never picked up.

why this cannot change an existing kernel

Every branch of the selector is TileShapeN < 16 → x1, < 32 → x2, else x4, and the default is 32. Any tile with N >= 32 therefore resolves to the identical atom it resolves to today. The only configurations whose behaviour changes are ones that currently fail to compile, so there is no kernel in the library or the profiler that this can move.

verified

RTX 5070 Ti (sm_120a), CUDA 13.0.3, driver 610.74, reference checks via TestSmall:

tile	before	after
128x64x64, 128x32x64	builds, passes	builds, passes
128x16x64, 128x8x64	fails to compile	builds, passes

Beyond what's in the test file, I ran a wider matrix while narrowing this down: e4m3 and e5m2, cooperative and pingpong, N = 64/32/16/8, all passing. Pingpong at N=8 is covered in the committed tests specifically because that's the shape that produced reference failures on the block-scaled path during #3292 review.

Non-regression: full SM120 tensorop suite green, 35/35 across tensorop_f32_sm120 (15), tensorop_f16_sm120 (14) and sm120_grouped_gemm_device_tensorop (6).

Tests go into the existing sm120_gemm_f8_f8_f32_tensor_op.cu alongside the 128x64x64 case, matching the one-test-per-tile-shape layout the block-scaled narrow-N tests already use, so there's no new file and no CMakeLists change.

deliberately not included

The array, sparse and blockwise SM120 builders have the same omission. I left them alone because each needs its own verification matrix and I've only proven the dense F8F6F4 path on hardware I have. Happy to extend in a follow-up if you want them, or to fold them in here if you'd rather it land as one change.

Two things I checked and they are pre-existing, not affected by this: sub-8-bit operands (e2m1) don't build on this builder at N=64 or N=32 either, so that's a separate constraint; and cutlass_library emits 128x16 and 128x8 only for the block-scaled generators, so nothing currently shipped was broken by the missing forward.

@depaulmillz you're the closest thing to an owner I can identify from blame here, since #3176 is where small tile N started and this is the sibling builder that didn't get it. Point me at someone else if that's not right.

[Copilot]

Pull request overview

This PR fixes SM120 dense (non-blockscaled) F8F6F4 GEMM mainloop construction for narrow-N tiles by forwarding the actual tile-N (size<1>(TileShape_MNK{})) into sm120_rr_smem_copy_selector_B, allowing correct selection of x1/x2/x4 ldmatrix atoms for N=8/16.

Changes:

• Forward TileShapeN to detail::sm120_rr_smem_copy_selector_B in the SM120 dense MMA builder so narrow-N tiles can compile.
• Add unit tests covering tile shapes 128x32x64, 128x16x64, and 128x8x64 (including a pingpong schedule case) for the dense F8 path.

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated no comments.

File	Description
include/cutlass/gemm/collective/builders/sm120_mma_builder.inl	Passes tile-N into the shared-memory copy atom selector for operand B to enable narrow-N tile support.
test/unit/gemm/device/sm120_tensorop_gemm/sm120_gemm_f8_f8_f32_tensor_op.cu	Adds regression tests ensuring 128x16 and 128x8 dense F8 tiles compile and produce correct results (plus coverage for pingpong schedule).

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.