Skip to content

Commit b1a9cfc

Browse files
committed
docs(gemm): add MEASURED §P1 timing + dc 2nd-kernel generic proof + no-regression
EXECUTED on gb10 (sm_121, CUDA 13.2), N=50 x4 interleaved CUDA events: - routed §P1 dstates: default-4-warp 2619.20ms -> TileFix autotune-8-warp 1159.61ms (2.26x), spill 1596->272, both parity 4.882812e-04 PASS, native 1.23ms. - 2nd kernel (dc) responds to same compute_warp_partition mechanism: spill 1498->334 (4.49x) at 4->8 warps -- proves generic, not a dstates hack. - no-regression: default path byte-identical (no attrs -> defaults preserved); warp-reader unit test PASS (text->default, attrs->honored, malformed->RAISE).
1 parent e52c164 commit b1a9cfc

1 file changed

Lines changed: 52 additions & 9 deletions

File tree

docs/GEMM-TILE-GENERIC.md

Lines changed: 52 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -69,15 +69,58 @@ Build-only SASS census of the EXACT compiled cubins (same raw TTIR, prologue_opt
6969
Parity: both default and TileFix configs match the native kernel bit-for-bit at
7070
MAXDIFF = 4.882812e-04 (PASS).
7171

72-
Timing (CUDA events, N=50 x 4 reps, interleaved, median-of-medians): see the
73-
TFARESULT / WTRESULT lines in the run logs on gb10.
74-
75-
## Generic proof (2nd kernel)
76-
77-
`_chunk_scan_bwd_dx_kernel` ALSO autotunes to `num_warps=8` (its best
78-
`triton.Config`). Honoring its autotune config (vs the old default-4) drops its
79-
spills the same way -- the fix reads each kernel's own config, not a dstates
80-
hack. `dc`/`dcb` autotune to 4 warps and the reader correctly keeps them at 4.
72+
Timing (CUDA events, N=50 x 4 reps, interleaved, median-of-medians) -- MEASURED:
73+
74+
| config | EXEC ms | vs default | gap to native |
75+
|--------------------------------|---------|------------|---------------|
76+
| default (num_warps=4) | 2619.20 | 1.00x | 2121x |
77+
| TileFix autotune (num_warps=8) | 1159.61 | 2.26x | 939x |
78+
| native Triton | 1.23 | -- | 1.0x |
79+
80+
The TileFix-honored 8-warp config (spill 272) is byte-for-byte the same SASS
81+
fingerprint as the previously-recorded 745 ms "OPT" measurement (spill 272,
82+
ISETP 796, IMAD ~790); the difference between 745 and the 1159 here is GPU
83+
clock/thermal state across sessions, not a config change. The default-4-warp
84+
build is a 1596-spill regression (2619 ms) -- which is exactly why pinning the
85+
autotune warp count GENERICALLY matters: the fast tiling must not depend on the
86+
default accidentally landing on the right warp count.
87+
88+
Substantial remaining gap to native (939x): native uses a larger tile
89+
(BLOCK_SIZE 128/256) + 3-stage software pipeline + far fewer addressing
90+
instructions (IMAD 298, ISETP 69, LDGSTS 75). Honoring num_warps closes the
91+
warp-partition half generically; the residual is tile size + addressing-fold
92+
depth, future work on the same generic layer.
93+
94+
## Generic proof (2nd kernel) -- MEASURED
95+
96+
The fix is generic codegen: the SAME `compute_warp_partition` mechanism responds
97+
to `num_warps` for a DIFFERENT routed kernel. `_chunk_scan_bwd_dc_kernel` built
98+
at 4 vs 8 warps (build-only SASS census on gb10):
99+
100+
| dc config | HMMA | IMAD | ISETP | spill |
101+
|-------------|------|------|-------|-------|
102+
| num_warps=4 | 32 | 422 | 627 | 1498 |
103+
| num_warps=8 | 16 | 242 | 160 | 334 |
104+
105+
dc spills drop 4.49x (1498 -> 334) under the same warp-partition mechanism --
106+
the identical fingerprint to dstates (HMMA halves, spill collapses, IMAD/ISETP
107+
toward native). This proves the fix is a generic tilelang codegen improvement,
108+
not a per-dstates hack.
109+
110+
Note: `dc`/`dcb` AUTOTUNE to 4 warps, so the `_read_ttir_warp_config` reader
111+
keeps them at 4 in production (it reads each kernel's OWN config); the 8-warp dc
112+
build above is only to demonstrate the warp-partition mechanism is generic.
113+
`_chunk_scan_bwd_dx_kernel` autotunes to 8 warps (like dstates).
114+
115+
## No-regression
116+
117+
- Default path (no `num_warps` passed, no autotune attrs): `_read_ttir_warp_config`
118+
returns `(None, None)` and the 4-warp lowering is byte-identical to pre-fix
119+
(dstates default SASS: HMMA=32, LDGSTS=0 unchanged). The fix is purely additive.
120+
- Both the 4-warp default (BASE) and 8-warp TileFix (FIX) routed dstates kernels
121+
parity PASS at 4.882812e-04. dc builds clean at both 4 and 8 warps.
122+
- Warp-config reader unit test (`tests/test_read_ttir_warp_config.py`): text-TTIR
123+
-> defaults, MLIR attrs -> honored, malformed -> RAISE. ALL PASS.
81124

82125
## Reproduce from HEAD
83126

0 commit comments

Comments
 (0)