mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-11-01 08:23:24 +08:00
The example and the maint warp-specialized benchmark each carried a private macro with the blockscaled_chunk_kmajor tile addressing spelled out by hand. Move that knowledge into the library next to T.mma_gemm_blockscaled: - T.ue4m3_scale_tile_source_coords: the single source of the packed scale tile addressing (tiles are contiguous in this layout); words_per_kblock stays a parameter so future contracts (block_K != 256, MX scale packings) reuse it. - T.copy_ue4m3_scale_tile: cooperative whole-block staging macro with trace-time dtype/shape validation, used by the example. The scale operand stays visible in user kernels; only the layout math is hidden. The maint kernel keeps its producer-warp-group staging schedule and reuses the coords helper. Validated: nvf4 language + example CLI tests (91 passed incl. a new validation test), example and WS benchmark --verify pass. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>