qutaoandClaude Fable 5 790a3f0ed1 Support M or N tail tiles in the SM120 NVFP4 example ...

Production shapes rarely keep M (batch x seq) a multiple of 128. Lift
the divisibility requirements to the validated envelope:

- scale packers zero-pad rows to full 128-row tiles (the same ceil
  convention CUTLASS/cuBLAS/CuTeDSL use for the blocked SF tensor), so
  the packed source always has ceil(rows/128)*128 rows;
- the example accepts an arbitrary M or an N tail: TMA loads zero-fill
  out-of-bounds rows and the C store is predicated. N keeps the
  16-byte-aligned-rows rule (N % 8 == 0), the same contiguous-dim
  alignment CuTeDSL validates;
- simultaneous M and N tails are rejected with a clear error: the
  fragment-to-global copy lowering emits the column guard bound from
  the wrong dimension in that case (upstream bug, tracked separately).

Tests: tail-shape runtime correctness (257x384, 130x128, 128x136 with
exact binary-scale references), rejection paths, and packer padding
round-trip.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

2026-07-10 15:07:03 +08:00

..

test_tilelang_quantize_nvfp4_scale_layout.py

Support M or N tail tiles in the SM120 NVFP4 example

2026-07-10 15:07:03 +08:00