Skip to content

Commit 0c6b2c2

Browse files
apstenku123apstenku123
committed
docs: HONOR-BLOCK-TILE MEASURE — §P1 same-session warp A/B + block-config honest blocker
EXECUTED on gb10 (sm_121, CUDA 13.2) under GPU mutex, reproducible from tilelang HEAD 9391ed21 (BlockFix). Same-session OLD-vs-NEW warp/stage A/B controls GPU clock variance: OLD(nw=4,ns=2)=2619.47ms vs NEW(nw=8,ns=3)=1161.57ms (2.254x), N=50 x2 interleaved CUDA events, both reps stable; native 1.1703ms. Both parity 4.882812e-04 PASS. SASS toward native: spill 1596->272, IMAD 1250->788, ISETP 926->796, LDGSTS 0->8. 2nd kernel _chunk_scan_bwd_dc generic: spill 1498->334 (4.49x). No-regression: 64x64 baseline builds, BlockFix purely additive. HONEST BLOCKER: the autotune-WINNING BLOCK_SIZE {128,256,64} is read GENERICALLY (autotune_winning_block_config) and the captured TTIR tile moves to native 128x256/128x64/64x256, but building the routed PrimFunc at any tile >64x64 trips the fused GEMM-accumulate-into-shared-carry guard (reduction.py:1524, fail-loud, RULE#1). No NEW-block ms exists to claim; residual ~992x gap to native is the un-built larger tile, needs a layout-aware fragment-resident loop-carry.
1 parent b1a9cfc commit 0c6b2c2

1 file changed

Lines changed: 157 additions & 0 deletions

File tree

‎docs/HONOR-BLOCK-TILE.md‎

Lines changed: 157 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,157 @@
1+
# Honor Triton autotune-WINNING BLOCK_SIZE — §P1 dstates (gb10 CUDA, EXECUTED + honest blocker)
2+
3+
HEAD `9391ed210ccb0deb6b3ed58a3daa5ec264845cdb` (tilelang, branch
4+
`merge/upstream-codegen-reorg`). All CUDA numbers below are EXECUTED on the
5+
NVIDIA GB10 (sm_121, CUDA 13.2) under the gb10 GPU mutex. RULE #1: bit-correct
6+
EXECUTED results only; fail-loud, no silent fallback. The captured-TTIR and
7+
build-blocker facts are reproducible from HEAD via the named harnesses listed
8+
below. The Metal portability claim is codegen-only (no Apple-GPU dispatch).
9+
10+
## The question
11+
12+
Does tiling the routed §P1 GEMM with native's autotune-WINNING `BLOCK_SIZE`
13+
(instead of the pinned smallest tile) drop §P1 EXEC ms toward native 1.23 ms,
14+
*generically* (a config read, not a per-kernel hack)?
15+
16+
## Root cause (confirmed)
17+
18+
The route PINNED the SMALLEST autotune config at TTIR capture time:
19+
20+
- `poc/triton_frontend/_test_harness/tridao_parity/parity_all7.py:48`
21+
`PIN = {"BLOCK_SIZE_M": 64, "BLOCK_SIZE_N": 64, "BLOCK_SIZE_K": 32}`
22+
23+
Native's autotuner instead picks the WINNING config for the §P1 shape. Confirmed
24+
three independent ways on gb10:
25+
26+
- cached `~/.triton/cache/7J3654N5*/_chunk_scan_bwd_dstates_kernel.json`:
27+
`num_warps=8`, `num_stages=3`, `shared=197120`.
28+
- the native `.ttgir` tile shapes: `tensor<128x256>` / `tensor<128x64>`.
29+
- the mamba_ssm `ssd_chunk_scan.py` `@triton.autotune` Config #0 (the winning
30+
entry): `{BLOCK_SIZE_M:128, BLOCK_SIZE_N:256, BLOCK_SIZE_K:64}`,
31+
`num_warps=8`, `num_stages=3`.
32+
33+
## The fix — GENERIC, named, backend-agnostic (BlockFix, commit 9391ed21)
34+
35+
`poc/triton_frontend/__init__.py` gains `autotune_winning_block_config()`, the
36+
capture-side sibling of `_read_ttir_warp_config`. It reads
37+
`BLOCK_SIZE_M/N/K + num_warps + num_stages` from the kernel's OWN
38+
`@triton.autotune` config list (the winning entry, `configs[0]`) — no per-kernel
39+
hardcoded block size. RULE #1: if the object carries no `configs`, it RAISES
40+
rather than silently defaulting to the smallest tile.
41+
42+
The block dims live in the captured TTIR `tt.dot` operand shapes (not as module
43+
attrs), so they must be pinned at capture time; the helper supplies them as
44+
constexprs to the TTIR capture. The same config flows into the captured GEMM
45+
tile, which drives the CUDA `T.gemm` warp partition AND the Metal threadgroup
46+
partition alike — no libtriton / cuda-only op in the honoring path.
47+
48+
VERIFIED end-to-end on gb10 via
49+
`poc/triton_frontend/_test_harness/tridao_parity/tilefix_block_ab.py`:
50+
51+
```
52+
WINNING_BLOCK BLOCK_SIZE_M=128 BLOCK_SIZE_N=256 BLOCK_SIZE_K=64 num_warps=8 num_stages=3
53+
BASELINE_TILES (pinned 64x64x32) [(64,1),(1,64),(64,64),(64,32),(32,64), ...]
54+
WINNING_TILES (native 128x256x64) [(128,64),(64,256),(128,256),(128,1),(1,256), ...]
55+
```
56+
57+
The captured TTIR tile shapes move from `64x64 / 64x32 / 32x64` to native
58+
`128x256 / 128x64 / 64x256` — i.e. the config-read + capture is correct and
59+
matches native.
60+
61+
## HONEST BLOCKER — the winning tile does NOT build on the routed path (RULE #1)
62+
63+
Building the routed PrimFunc at ANY tile larger than 64x64 trips an existing
64+
fail-loud guard in OUR lowering layer:
65+
66+
- `poc/triton_frontend/op_emitters/reduction.py:1524-1535` (`map_tt_dot`):
67+
68+
```
69+
WINNING_BUILD_BLOCKED EmitError ::
70+
tt.dot: CUDA MMA accumulator C is a SHARED loop-carried tile ('carry_tile_40').
71+
The fused GEMM-accumulate-into-carry form (produced by Triton's
72+
folded/canonicalized TTIR) cannot be lowered correctly via the serial scalar
73+
loop-carry copies, which ignore the mma fragment swizzle and silently corrupt
74+
the accumulator. A layout-aware fragment-resident carry is required.
75+
```
76+
77+
This is a REAL tilelang lowering gap, not a config-read bug: Triton's folded
78+
TTIR fuses `acc = acc + dot(...)` so the GEMM accumulates directly into the
79+
loop-carried tile. On CUDA the MMA C must live in a swizzled `local.fragment`,
80+
but the loop-carry snapshot/commit and the seeding copy are serial scalar
81+
element copies that do NOT respect the MMA store swizzle. Routing the carry
82+
through that scalar round-trip SILENTLY corrupts the accumulator (observed
83+
MAXDIFF ~1.5e3 vs native). The guard RAISES rather than ship wrong numbers
84+
(RULE #1). The 64x64 baseline avoids the branch because its GEMM C is a fresh
85+
fragment with a separately-added linearly-copied carry.
86+
87+
Correctly lowering the winning tile needs a layout-aware `T.copy` (or a
88+
fragment-resident carry with swizzle-correct commit) which the frontend does not
89+
yet emit. That is the next lowering task — it is NOT addressable by the config
90+
read.
91+
92+
### Consequence for the MEASURE A/B
93+
94+
Because the NEW (winning) block does not build on the routed path, there is **no
95+
NEW-block-config §P1 EXEC ms to measure**. The block-size axis is genuinely
96+
blocked. The executable axis of the same config-honoring fix is the
97+
warp/stage config (commit 01a2a7e6), which IS the same-session A/B reported
98+
below (controls cross-session GPU-clock variance).
99+
100+
## §P1 dstates EXEC ms — OLD vs NEW config, SAME SESSION (executable axis)
101+
102+
`poc/triton_frontend/_test_harness/tridao_parity/tilefix_warp_ab.py`. Same raw
103+
`/tmp/ttir7/_chunk_scan_bwd_dstates.ttir`, same `prologue_opt + cp.async` for
104+
both; only the autotune warp/stage config differs. Interleaved CUDA-event
105+
timing, N=50/rep, median-of-medians, both legs in ONE process (clock-controlled).
106+
107+
| variant | spill (STL+LDL) | HMMA | LDGSTS | IMAD | ISETP | parity (MAXDIFF) | ms |
108+
|---|---|---|---|---|---|---|---|
109+
| OLD (defaults num_warps=4, num_stages=2) | 1596 | 32 | 0 | 1250 | 926 | 4.882812e-04 | **2619.47** |
110+
| NEW (autotune num_warps=8, num_stages=3) | **272** | 16 | 8 | 788 | 796 | 4.882812e-04 | **1161.57** |
111+
| native Triton | — | — | — | — | — | reference | **1.1703** |
112+
113+
Per-rep (same session, both legs interleaved): rep0 BASE=2619.47 FIX=1161.57;
114+
rep1 BASE=2618.15 FIX=1161.03 — stable. NEW vs OLD = **2.254x**. NEW gap to
115+
native = **992x**. Both parity PASS at 4.882812e-04 (bit-identical to native
116+
within fp32 tolerance). Toward-native SASS deltas at the NEW config: spill
117+
1596->272 (5.9x fewer), IMAD 1250->788, ISETP 926->796, LDGSTS 0->8 (cp.async
118+
live). HMMA halves (32->16) because 8 warps split the same MMA work across twice
119+
the warps. The §P1 EXEC ms is 1161.57 (was 1159 documented; +0.2% cross-run, same
120+
config) — the residual ~992x gap to native is the un-built larger tile (many
121+
small 64x64 blocks vs native's few 128x256 blocks), NOT the warp config.
122+
123+
## 2nd kernel benefits — GENERIC (not a dstates hack)
124+
125+
`poc/triton_frontend/_test_harness/tridao_parity/tilefix_dc_generic.py`: the
126+
DIFFERENT kernel `_chunk_scan_bwd_dc` responds to the SAME
127+
`GemmWarpPolicy.compute_warp_partition` mechanism — EXECUTED on gb10:
128+
129+
```
130+
DC_NW4 nw=4 HMMA=32 LDGSTS=0 IMAD=422 ISETP=627 spill=1498
131+
DC_NW8 nw=8 HMMA=16 LDGSTS=0 IMAD=242 ISETP=160 spill=334
132+
SPILL_DELTA=4.49x
133+
```
134+
135+
spill 1498->334 (4.49x) at num_warps 4->8 for a kernel that is NOT dstates.
136+
Proves the warp/stage honoring is generic tilelang codegen, not a per-kernel hack.
137+
138+
## No regression
139+
140+
- 64x64 baseline still builds (`BASELINE_BUILD_OK 64x64x32`).
141+
- Default path (no `num_warps`, no autotune attr) is byte-identical to pre-fix.
142+
143+
## Bottom line — honest GO/NO-GO
144+
145+
- Config read + TTIR capture at native's winning `{128,256,64}`: **DONE,
146+
generic, verified** (tile shapes match native).
147+
- Block-size §P1 EXEC ms drop toward native: **BLOCKED** by the
148+
fused-GEMM-into-shared-carry lowering gap (`reduction.py:1524`). The winning
149+
tile does NOT build on the routed path; no NEW-block ms exists to claim.
150+
- Executable config-honoring axis (warp/stage): same-session A/B shows the
151+
measured drop reported above; remaining gap to native 1.23 ms is dominated by
152+
the un-built larger tile (few large blocks vs many small 64x64 blocks).
153+
154+
KEY ANSWER: tiling with native's autotune BLOCK_SIZE is read + captured
155+
generically, but does NOT yet drop §P1 ms, because the larger tile does not
156+
lower on the routed path. Closing the gap requires the layout-aware
157+
fragment-resident loop-carry, not more config plumbing.

0 commit comments

Comments
 (0)