mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-22 06:50:25 +08:00
Chennes
e4fba8678a
Some ldmatrix lane maps step beyond the final SM75 MMA micro-tile, and column overflow must carry into the next row. Layout remapping previously linearized the source indices and optional offset before converting them back to multidimensional coordinates, which normalized these addresses as a side effect. Since #2836 preserves the multidimensional indices, the out-of-range coordinates reach the swizzle layout unchanged and the native SM75 MMA GEMM paths fault on illegal shared-memory accesses. Wrap each complete row-major source offset within its operand region in ldmatrix_a and ldmatrix_b before rebuilding the source coordinates. Leading dimensions remain untouched, so the wrap cannot cross into another stage of a multi-version buffer. When the coordinates are provably in range, the simplifier removes the wrap. SM70 and SM80 generated sources remain byte-identical, while the SM75 addresses match the pre-#2836 address mapping. Add multi-stage SM75 GEMM coverage and a lowered-offset bounds regression that cross-compiles explicitly for sm_75 and therefore runs on any CUDA runner.