Skip to content

Commit f4b77a1

Browse files
apstenku123apstenku123
committed
docs(triton-route): MEASURED-from-HEAD PROVE+MEASURE — honest NO-GO
Routed Tri-Dao bwd, all numbers re-measured on gb10 from committed tilelang HEAD aada7de1 (=gb10 6a5500d7), pf JSONs regenerated before every run. - Route: 4/7 fully lower to CUDA+mma (dcb EmitError %1044; db/dx unbound arg4). - Parity real-strided 1e-3 all-elements: 0/7. dstates FAILs MAXDIFF 2.41e2 at cs=64!=BK (2 K-trips), writes only half the blocks. The prior bit-exact PASS was the DEGENERATE single-K-trip tiny config (cs==BK==32) only. - §P1: routed dstates is MONOMORPHIZED (fixed extent 1048576, baked tiny grid) -> SEGFAULTS at §P1 (1,64,112). No routed §P1 EXEC ms exists. Native §P1=1.17ms. - Routed dstates EXEC (tiny, the only runnable PASS): 4.46 ms/kernel vs native 0.017ms — residual = whole-block ~168-loop prologue, not operand staging. - USER ASK delivered (no regression): matmul SEGFAULT->1.343 ms (MAXDIFF 2.60e-2, TF32 trunc); fla_dot_exp2 SEGFAULT->0.0322 ms (MAXDIFF 1.26e-5, ALLCLOSE True). Honest GO/NO-GO: NO-GO for the headline goal; USER-ASK sub-results are GO.
1 parent 5d141ef commit f4b77a1

1 file changed

Lines changed: 178 additions & 0 deletions

File tree

‎docs/TRITON-ROUTE-RUN.md‎

Lines changed: 178 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,178 @@
1+
# TRITON-ROUTE-RUN — routed Tri-Dao bwd: PROVE + MEASURE (MEASURED from HEAD)
2+
3+
All numbers below are **MEASURED on gb10 (NVIDIA GB10 sm_121, aarch64-linux)** and
4+
**reproducible from the committed HEAD**:
5+
6+
- tilelang frontend HEAD: `aada7de1` (local) = `6a5500d7` (gb10 `/home/dave/source/tilelang`)
7+
— *"fix(triton-frontend): sink SBlock alloc_buffers INSIDE kernel launch"*
8+
- gb10 working tree clean at that commit; pf JSONs regenerated from HEAD before every run.
9+
10+
RULE #1: no fabricated PASS. Every claim here was re-run from the committed HEAD after
11+
regenerating the routed PrimFunc JSON via `emit_pf_json.py`. Where the route fails or the
12+
launch segfaults, that is reported as the literal verdict — not papered over.
13+
14+
---
15+
16+
## 1. Routing coverage (honest, from HEAD)
17+
18+
Two distinct route gates:
19+
20+
- **emit / PrimFunc level** (`from_ttir` → `tvm.ir.save_json`): 6/7 produce a PrimFunc.
21+
`_chunk_scan_bwd_dcb` RAISES `EmitError: PtrState references unresolved SSA value '%1044'`
22+
— it does NOT route.
23+
- **full lowering to CUDA with real `mma`** (`tilelang.engine.lower`, `route_all7.py`): **4/7**.
24+
25+
```
26+
ROUTE _chunk_scan_bwd_dstates global=True mma=2 atomicAdd=0
27+
ROUTE _chunk_scan_bwd_dc global=True mma=2 atomicAdd=0
28+
ROUTEFAIL _chunk_scan_bwd_dcb EmitError: PtrState references unresolved SSA value '%1044'
29+
ROUTE _chunk_scan_bwd_dx global=True mma=2 atomicAdd=1
30+
ROUTEFAIL _chunk_state_bwd_db InternalError: variables (arg4,) used but not passed as API args
31+
ROUTEFAIL _chunk_state_bwd_dx InternalError: variables (arg4,) used but not passed as API args
32+
ROUTE _chunk_state_bwd_ddAcs_stable global=True mma=2 atomicAdd=1
33+
TOTAL_OK_WITH_MMA=4/7
34+
```
35+
36+
The prior "7/7 route to TVM PrimFunc w/ real T.gemm" was measured at the *emit* level;
37+
at *full lowering* only 4/7 survive (`db`/`dx` fail with an unbound `arg4`).
38+
39+
---
40+
41+
## 2. Parity (routed vs native triton, torch.allclose 1e-3, all elements)
42+
43+
### 2a. dstates — DEGENERATE single-K-trip (cs == BK == 32) — PASS
44+
45+
`parity_tiny.py` (cfg b1 nh2 hd64 ds64 nc2 **cs32** ng2 s256, grid (1,2,2), numel 16384):
46+
47+
```
48+
NATIVE nz=16384/16384 sum=-1240.0947
49+
ROUTED nz=16384/16384 sum=-1240.0947
50+
MAXDIFF=0.000000e+00 ALLCLOSE_1e-3=True -> PASS
51+
```
52+
53+
This is a bit-exact PASS, **but only because cs==BK gives a SINGLE K-trip** (no
54+
K-accumulation across trips). It does not exercise the strided multi-K-trip path.
55+
56+
### 2b. dstates — REAL STRIDED multi-K-trip (cs=64 != BK=32, 2 K-trips) — FAIL
57+
58+
`parity_all7.py small _chunk_scan_bwd_dstates` (cfg b1 nh8 hd64 ds64 nc4 **cs64** ng2 s256,
59+
grid (1,4,8)):
60+
61+
```
62+
NATIVE nz=131072/131072 sum=-18940.2539
63+
ROUTED nz=65536/131072 sum=-9057.1641 <-- only HALF the blocks written
64+
MAXDIFF=2.409756e+02 ALLCLOSE_1e-3=False -> FAIL
65+
```
66+
67+
**The original DEFECT 2 (MAXDIFF ~2.41e2) is NOT fixed for the real strided / multi-K-trip
68+
case.** It only "passes" in the degenerate single-K-trip tiny config. The routed kernel writes
69+
exactly half the output blocks (65536/131072) under the 2-K-trip config — consistent with a
70+
K-accumulation / second-K-trip defect, not operand staging.
71+
72+
### 2c. Other 6 kernels (small, real strided) — none confirmed PASS
73+
74+
| kernel | small-mode verdict (from HEAD) |
75+
|---|---|
76+
| `_chunk_scan_bwd_dstates` | FAIL — MAXDIFF 2.41e2, half blocks |
77+
| `_chunk_scan_bwd_dc` | SEGFAULT at routed launch |
78+
| `_chunk_scan_bwd_dcb` | does not route (EmitError %1044) |
79+
| `_chunk_scan_bwd_dx` | SEGFAULT at routed launch |
80+
| `_chunk_state_bwd_db` | InternalError on lower (arg4 unbound) |
81+
| `_chunk_state_bwd_dx` | InternalError on lower (arg4 unbound) |
82+
| `_chunk_state_bwd_ddAcs_stable` | NO_CAPTURE — atomic/return output; harness grid KeyError(BLOCK_SIZE_M) |
83+
84+
**Parity PASS count (real strided, all elements, 1e-3): 0 / 7.**
85+
The single bit-exact PASS (2a) is the degenerate single-K-trip config only.
86+
87+
### 2d. dstates — PRODUCTION §P1 — SEGFAULT (monomorphized kernel)
88+
89+
The routed dstates PrimFunc declares **every buffer at a FIXED flat extent 1048576** and bakes
90+
its grid/loop bounds to the tiny dims — it is **monomorphized**, not parameterized. §P1 needs
91+
grid (1,64,112) and a 29,360,128-element dout. Feeding §P1-sized tensors segfaults at the
92+
routed launch (both `parity_prod_dstates.py` and `parity_all7.py prod` core-dump). Native §P1
93+
runs fine. **There is no routed §P1 to compare.**
94+
95+
---
96+
97+
## 3. EXEC ms (real kernel runtime — LAUNCH + CUDA-event timed)
98+
99+
### 3a. Native §P1 reference (re-measured)
100+
101+
```
102+
NATIVE_P1_GPU_EVENT_MS = 1.17 ms/kernel grid=(1,64,112) numel=29,360,128
103+
```
104+
(matches the 1.14 ms reference.)
105+
106+
### 3b. Routed dstates EXEC ms — ONLY at the tiny config it can run
107+
108+
`parity_tiny` is a single-K-trip bit-exact PASS, so per RULE #1 (exec ms only for parity-PASS),
109+
this is the only routed dstates EXEC number that is allowed:
110+
111+
```
112+
ROUTED dstates GPU-event = 4.46 ms/kernel (tiny grid (1,2,2), numel 16384)
113+
NATIVE dstates same tiny config = 0.017 ms/kernel
114+
```
115+
116+
**Residual bottleneck:** the routed kernel takes 4.46 ms to fill 4 blocks / 16384 elements —
117+
~260x the native time for the SAME work — because it runs a fixed whole-block prologue
118+
(the ~168-loop monomorphized prologue) independent of useful work. This is the dominant
119+
residual, NOT operand staging.
120+
121+
There is **no routed §P1 EXEC ms** (§2d: it segfaults), so the routed-vs-native-1.14ms
122+
comparison at §P1 cannot be made from HEAD.
123+
124+
### 3c. USER ASK — fla_dot_exp2 + matmul GPU-EXEC before/after (REAL kernel runtime)
125+
126+
The generic arg-marshaller segfaulted both when the runtime buffer was real-sized (smaller than
127+
the monomorphized declared 1048576 extent). Fix = pack each buffer to the declared extent
128+
(`fit()` in the exec harness). Both go from core-dump to running:
129+
130+
| kernel | BEFORE | AFTER (ms/kernel) | MAXDIFF | note |
131+
|---|---|---|---|---|
132+
| matmul (64x64x64) | SEGFAULT (core dump) | **1.343** | 2.60e-2 | TF32-MMA truncation (ALLCLOSE False @1e-3; precision, not a value bug) |
133+
| fla_dot_exp2 (16x16x16) | SEGFAULT (core dump) | **0.0322** | 1.26e-5 | ALLCLOSE_1e-3 = True |
134+
135+
No regression: both were a hard crash before, both run after. Same root cause as the dstates
136+
prod segfault (undersized buffer vs monomorphized declared extent).
137+
138+
---
139+
140+
## 4. GO / NO-GO
141+
142+
**NO-GO** for the headline goal ("routed Tri-Dao bwd RUNNABLE + NUMERICALLY CORRECT 1e-3,
143+
measured by EXEC ms vs native 1.14ms").
144+
145+
- Runnable: only the degenerate single-K-trip tiny config; §P1 segfaults (monomorphized).
146+
- Numerically correct 1e-3 all-elements real strided: **0/7** (dstates FAILs 2.41e2 at 2 K-trips;
147+
half blocks written). The earlier "FIXED / MAXDIFF 0" was the single-K-trip config only.
148+
- EXEC ms vs 1.14ms: cannot be produced — routed §P1 does not launch.
149+
150+
**Delivered (GO) sub-results, reproducible from HEAD:**
151+
- USER ASK: matmul SEGFAULT→1.343 ms; fla_dot_exp2 SEGFAULT→0.0322 ms (no regression).
152+
- dstates single-K-trip bit-exact PASS (MAXDIFF 0).
153+
- Honest residual root-cause: monomorphized fixed-extent PrimFunc + whole-block prologue
154+
(≈168 loops), and an unfixed 2nd-K-trip K-accumulation defect (half blocks at cs=64).
155+
156+
**Remaining defects to fix (the real blockers):**
157+
1. **Monomorphization** — routed PrimFunc bakes extent 1048576 + tiny grid; cannot run §P1.
158+
Must parameterize buffer extents/grid by the runtime dims.
159+
2. **Multi-K-trip K-accumulation** — at cs!=BK only half the output blocks are written and
160+
MAXDIFF=2.41e2. The single-K-trip PASS masks this.
161+
3. **Route gaps** — `dcb` EmitError %1044; `db`/`dx` unbound `arg4` at lowering.
162+
163+
---
164+
165+
## 5. Reproduce (gb10, GPU mutex held)
166+
167+
```bash
168+
ssh gb10 'cd /home/dave/source/tilelang && source /home/dave/cppmega-venv/bin/activate
169+
# regenerate routed PrimFunc JSONs from HEAD
170+
for n in _chunk_scan_bwd_dstates _chunk_scan_bwd_dc _chunk_scan_bwd_dx \
171+
_chunk_state_bwd_db _chunk_state_bwd_dx _chunk_state_bwd_ddAcs_stable; do
172+
python poc/triton_frontend/_test_harness/tridao_parity/emit_pf_json.py $n; done
173+
python poc/triton_frontend/_test_harness/tridao_parity/route_all7.py # 4/7 mma
174+
python poc/triton_frontend/_test_harness/tridao_parity/parity_tiny.py # single-K PASS
175+
python poc/triton_frontend/_test_harness/tridao_parity/parity_all7.py small _chunk_scan_bwd_dstates # real-strided FAIL 2.41e2
176+
python /tmp/exec_matmul2.py # matmul AFTER 1.343 ms
177+
python /tmp/exec_fla3.py # fla AFTER 0.0322 ms'
178+
```

0 commit comments

Comments
 (0)