Skip to content

Commit ad8f28b

Browse files
apstenku123apstenku123
committed
docs(triton-route): MEASURED §P1 parity + exec-ms + prologue breakdown for routed Tri-Dao bwd
- dstates: PASS multi-K-trip (cs!=BK) small 3.05e-05 AND §P1 4.88e-04 (full 29.36M nz) - other 6/7: exact defects (3 segfault, dcb %1044 emit-raise, dx arg4_numel_243 de-mono gap, ddAcs no-compare) - routed dstates §P1 = 1473ms/kernel (nsys+events+wall agree) vs native 1.14ms = ~1292x - strip-and-time: ~100% is the un-fused pointer-arith prologue, GEMM ~0ms - GO/NO-GO: NO-GO whole bwd (1/7 correct, perf non-viable); dstates BLOCKER-1+2 RESOLVED
1 parent f4b77a1 commit ad8f28b

1 file changed

Lines changed: 140 additions & 0 deletions

File tree

‎docs/TRITON-ROUTE-P1.md‎

Lines changed: 140 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,140 @@
1+
# TRITON-ROUTE-P1 — Routed Tri-Dao bwd: parity + exec-ms + prologue breakdown
2+
3+
MEASURED from committed HEAD on gb10 (NVIDIA GB10 sm_121, aarch64-linux).
4+
- gb10 tilelang HEAD: `db5fe57c117460ce84d72cd7b07f153f51c4ef0a` (de-monomorph: symbolic flat-arg buffer extents)
5+
- local mirror HEAD: `48b9236a86ef0ebd0b22ec3625b374b7d2fbf874` (same logical fix, separate repo)
6+
- Working tree clean (only `3rdparty/tvm` submodule pointer = built state). No uncommitted source.
7+
- All 7 PrimFunc JSONs regenerated FRESH from HEAD via `emit_pf_json.py` before measuring.
8+
9+
RULE#1: real-strided / §P1 numbers ONLY. No degenerate (single-K-trip / cs==BK) or
10+
tiny-config number reported as a real-strided/§P1 result. Routed exec ms reported ONLY
11+
for the parity-PASS kernel. Every failing kernel reports its EXACT defect; no fabricated PASS.
12+
13+
---
14+
15+
## 1. Per-kernel parity — REAL strided multi-K-trip (cs=64 != BLOCK_K=32, 2 K-trips/chunk)
16+
17+
### 1a. SMALL real-strided config (b1 nh8 hd64 ds64 nc4 cs64 ng2 s256, grid (1,4,8))
18+
19+
| Kernel | Status @ HEAD | MAXDIFF | ALLCLOSE 1e-3 | Notes |
20+
|--------------------------------|------------------------|--------------|---------------|-------|
21+
| `_chunk_scan_bwd_dstates` | **PASS** | 3.051758e-05 | True | nz 131072/131072 full; cs=64,BK=32 = 2 K-trips |
22+
| `_chunk_scan_bwd_dc` | SEGFAULT (routed launch) | n/a | — | dumps core at kernel launch |
23+
| `_chunk_scan_bwd_dcb` | NO-ROUTE / compile fail | n/a | — | emit RAISES `PtrState references unresolved SSA value '%1044'`; compiled variant fails nvcc (`mma.h` error limit) |
24+
| `_chunk_scan_bwd_dx` | NO-LAUNCH (de-mono incomplete) | n/a | — | `tvm InternalError: arg4_numel_243 used but not passed as API argument` |
25+
| `_chunk_state_bwd_db` | SEGFAULT (routed launch) | n/a | — | dumps core at kernel launch |
26+
| `_chunk_state_bwd_dx` | SEGFAULT (routed launch) | n/a | — | dumps core at kernel launch |
27+
| `_chunk_state_bwd_ddAcs_stable`| NO-COMPARE | n/a | — | atomic/return-only output; harness pin_kernel also KeyError BLOCK_SIZE_M |
28+
29+
**SMALL real-strided result: 1/7 numerically correct (dstates).**
30+
31+
### 1b. PRODUCTION §P1 config (b1 nh112 hd64 ds64 nc64 cs64 ng8 s4096, grid (1,64,112), numel 29,360,128)
32+
33+
Only `_chunk_scan_bwd_dstates` routes+launches; only it has a §P1 parity number:
34+
35+
| Kernel | Status | MAXDIFF | ALLCLOSE 1e-3 | nonzero |
36+
|---------------------------|--------|--------------|---------------|----------------|
37+
| `_chunk_scan_bwd_dstates` | **PASS** | 4.882812e-04 | True | 29360128/29360128 (native 29360127) |
38+
39+
The SAME compiled kernel that passed at small (1,4,8) launches+passes at §P1 (1,64,112) with no
40+
segfault — the de-monomorph (symbolic flat-arg extents) fix is CONFIRMED for dstates: numel grew
41+
from ~131K to 29.36M with the same baked PrimFunc.
42+
43+
**§P1 prod result: 1/7 numerically correct (dstates).**
44+
45+
---
46+
47+
## 2. Routed EXEC ms/kernel @ §P1 (parity-PASS kernel only)
48+
49+
Authoritative on-device kernel time, 3 independent methods, all agree:
50+
51+
| Method | dstates §P1 routed ms/kernel |
52+
|-------------------------------------|------------------------------|
53+
| nsys `cuda_gpu_kern_sum` (1 launch) | **1473.43 ms** (avg of 4 kern instances 1465–1478ms) |
54+
| CUDA events (N=100) | 1470.52 ms |
55+
| Python wall (N=50) | 1476.32 / 1472.43 ms |
56+
57+
Device time ≈ wall time → the cost is **genuine on-GPU kernel runtime**, NOT host/launch overhead.
58+
59+
**vs native: native `_chunk_scan_bwd_dstates` @ §P1 = 1.14 ms/kernel.**
60+
Routed = ~1473 ms → **≈1292x SLOWER than native.**
61+
62+
---
63+
64+
## 3. Prologue vs GEMM breakdown — strip-and-time (MEASURED)
65+
66+
The routed kernel body is: (a) a ~133-`for`-loop **prologue** that materializes the Triton
67+
pointer arithmetic (`tl.arange`, broadcasts, int32-overflow guards) as serial per-thread
68+
element-wise loops over `bool[4096]` / `int[4096]` arrays, followed by (b) a tiny **GEMM**
69+
(2× `mma.sync` in a 4-iter K-loop with `ldmatrix`), then (c) carry/epilogue.
70+
71+
Strip-and-time (inject `return;` immediately before the GEMM block, recompile through the
72+
identical TileLang adapter path, same §P1 args, CUDA-event timed):
73+
74+
| Variant | DEVICE ms/kernel @ §P1 |
75+
|----------------------------------|------------------------|
76+
| FULL kernel | 1473.29 ms |
77+
| PROLOGUE-ONLY (return pre-GEMM) | 1474.99 ms |
78+
| ⇒ GEMM + epilogue contribution | ~0 ms (within noise) |
79+
80+
**~100% of the 1473ms is the prologue. The actual GEMM is negligible.**
81+
82+
### Is this inherent to the per-block model?
83+
YES — as currently lowered. The routed PrimFunc emits each Triton pointer-offset computation
84+
as a literal element-wise loop (e.g. `tile_binop[i] = (val <= INT_MAX) & (INT_MIN <= val)`)
85+
over 4096-element arrays, run **serially by every one of the 128 threads in every one of the
86+
7168 blocks**, instead of folding the offsets into the load addressing the way native Triton
87+
does. Roughly ~20 such `[4096]` loops × 128 threads × 7168 blocks of redundant scalar work
88+
dwarf the 2-instruction GEMM. This is a structural property of the current whole-block
89+
walker lowering (pointer math not fused into addressing, not thread-parallelized, not
90+
vectorized) — not a launch artifact and not fixable by tuning; it needs the frontend to
91+
fold pointer arithmetic into load/store addressing.
92+
93+
---
94+
95+
## 4. GO / NO-GO
96+
97+
**NO-GO for production routing of the Tri-Dao bwd as a whole.**
98+
99+
Honest scorecard, measured from committed HEAD:
100+
- Correctness: **1/7** kernels numerically correct at real-strided multi-K-trip (small AND §P1):
101+
only `_chunk_scan_bwd_dstates`. The other 6 fail at HEAD: 3 SEGFAULT at launch (dc, db, state_dx),
102+
1 NO-ROUTE emit-raise (dcb %1044), 1 NO-LAUNCH de-mono-incomplete (dx arg4_numel_243),
103+
1 NO-COMPARE atomic-output (ddAcs).
104+
- Runnability: dstates IS de-monomorphized — same baked kernel runs small→§P1 with no segfault. GOOD.
105+
- Performance: even the one correct kernel is **~1292x slower than native** (1473ms vs 1.14ms),
106+
and the strip-and-time proves ~100% of that is the un-fused pointer-arithmetic prologue.
107+
108+
### What is solid / reproducible
109+
- `_chunk_scan_bwd_dstates` multi-K-trip (cs!=BK) correctness: bit-close at small (3.05e-05)
110+
AND §P1 (4.88e-04), full output blocks written. The original BLOCKER-1 (multi-K-trip
111+
K-accumulation) and BLOCKER-2 (monomorphization segfault) are RESOLVED for dstates.
112+
113+
### Exact residuals (RAISED, not papered over)
114+
1. **6/7 kernels do not route/run at HEAD** — each with the exact defect above. The §P1 GO
115+
requires fixing the dc/db/state_dx launch segfaults, the dx de-mono gap (`arg4_numel_243`
116+
not threaded as an API arg — same class of fix as the dstates extent symbolization, but
117+
for an index/numel symbol), and the dcb emit defect (`%1044` unresolved PtrState).
118+
2. **Performance is non-viable** even for the passing kernel: the prologue must fold Triton
119+
pointer arithmetic into addressing (eliminate the serial `[4096]` bool/int offset loops)
120+
before any routed kernel can compete with native.
121+
122+
---
123+
124+
## 5. Reproduce (from committed HEAD on gb10)
125+
126+
```
127+
ssh gb10
128+
cd /home/dave/source/tilelang && source /home/dave/cppmega-venv/bin/activate
129+
# 1. regen PrimFunc JSONs from HEAD
130+
for k in _chunk_scan_bwd_dstates _chunk_scan_bwd_dc _chunk_scan_bwd_dcb _chunk_scan_bwd_dx \
131+
_chunk_state_bwd_db _chunk_state_bwd_dx _chunk_state_bwd_ddAcs_stable; do
132+
python poc/triton_frontend/_test_harness/tridao_parity/emit_pf_json.py $k; done
133+
# 2. small real-strided parity (per kernel, isolated so one segfault doesn't abort the rest)
134+
python poc/triton_frontend/_test_harness/tridao_parity/parity_all7.py small _chunk_scan_bwd_dstates
135+
# 3. §P1 production dstates parity + timing
136+
python poc/triton_frontend/_test_harness/tridao_parity/parity_prod_dstates.py
137+
# 4. prologue strip-and-time
138+
python /tmp/strip_time_final.py full
139+
python /tmp/strip_time_final.py prologue
140+
```

0 commit comments

Comments
 (0)