Skip to content

Commit 1dc3cb4

Browse files
apstenku123apstenku123
committed
docs(triton-route): MEASURED FRAMEFIX result — 7/7 route + 176 HMMA, parity 0/7 NO-GO
gb10 (sm_121a), under GPU mutex. Re-confirmed 7/7 route to real TF32 MMA (mma=2 each) and 176 HMMA.1688.F32.TF32 in SASS. FRAMEFIX converts the partial-tile C-carry write to a full-tile write (MEASURED prod nz 29360128/29360128, was 32-per-tile) but values are structurally wrong: dstates §P1 MAXDIFF=1.14e+03 allclose_1e-3=False (small 2.43e+02, ratio random std 2.85 -- not transpose/scale/precision). Routed 2419ms/kernel vs native 1.16ms (recorded fact, NOT a speed claim -- parity FAIL per RULE#1). Parity 0/7. Root cause: A/B operand staging (reduction.py::_stage_operand_to_shared) emits a serial linear-index swizzled shared fill inconsistent with the gemm's cooperative ptx_ldmatrix read; FRAMEFIX C-layout pin also flaky under LayoutInference nondeterminism. NO-GO; remaining defect RAISED, not papered over.
1 parent 8f5cf26 commit 1dc3cb4

1 file changed

Lines changed: 155 additions & 0 deletions

File tree

‎docs/TRITON-ROUTE-DONE.md‎

Lines changed: 155 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,155 @@
1+
# TRITON ROUTE — FRAMEFIX measured result (Tri-Dao bwd through tilelang→tvm)
2+
3+
ALL numbers below are **MEASURED on gb10 (NVIDIA GB10, sm_121a)** with the
4+
current FRAMEFIX code (`poc/triton_frontend/frame_register.py`, byte-identical
5+
local ↔ gb10), under the GPU mutex (`/tmp/gb10_gpu_owner.lock`,
6+
`fragfix-measure`). gb10-only for all GPU exec; no local Metal dispatch.
7+
GPU util 0% throughout; SIGTERM-only.
8+
9+
Honest verdict up front: **NO-GO. Parity 0/7.** The 7 Tri-Dao bwd kernels route
10+
end-to-end to real TF32 tensor-core MMA and FRAMEFIX fixes the partial-tile
11+
C-carry write to a FULL-tile write, but the routed output VALUES are
12+
structurally wrong (random, not transpose/scale/precision) and the routed
13+
kernel is ~2000x slower than native. Routed ms is therefore reported ONLY as a
14+
recorded fact — under RULE #1 it is NOT a valid speed claim because the kernel
15+
fails parity.
16+
17+
## 1. Route + real MMA (PASS, re-confirmed)
18+
19+
`route_all7.py` (from_ttir → tilelang.engine.lower, target=cuda):
20+
21+
| kernel | __global__ | mma | atomicAdd |
22+
|---------------------------------|-----------|-----|-----------|
23+
| _chunk_scan_bwd_dstates | yes | 2 | 0 |
24+
| _chunk_scan_bwd_dc | yes | 2 | 0 |
25+
| _chunk_scan_bwd_dcb | yes | 2 | 0 |
26+
| _chunk_scan_bwd_dx | yes | 2 | 1 |
27+
| _chunk_state_bwd_db | yes | 2 | 0 |
28+
| _chunk_state_bwd_dx | yes | 2 | 3 |
29+
| _chunk_state_bwd_ddAcs_stable | yes | 2 | 1 |
30+
31+
**TOTAL_OK_WITH_MMA = 7/7.**
32+
33+
SASS (`sass_all7.py`, real `HMMA.1688.F32.TF32` per cubin):
34+
35+
| kernel | HMMA.1688.F32.TF32 |
36+
|---------------------------------|--------------------|
37+
| _chunk_scan_bwd_dstates | 32 |
38+
| _chunk_scan_bwd_dc | 32 |
39+
| _chunk_scan_bwd_dcb | 32 |
40+
| _chunk_scan_bwd_dx | 32 |
41+
| _chunk_state_bwd_db | 32 |
42+
| _chunk_state_bwd_ddAcs_stable | 8 |
43+
| _chunk_state_bwd_dx | 8 |
44+
| **TOTAL** | **176** |
45+
46+
7/7 emit real TF32 tensor-core MMA in SASS; atomicAdd preserved on the reduce
47+
kernels. (The "288 HMMA" in earlier state notes was a different tile config; the
48+
current FRAMEFIX build emits 176.)
49+
50+
## 2. Parity (FAIL — 0/7)
51+
52+
Only `_chunk_scan_bwd_dstates` was exercised at parity (the lead kernel and the
53+
one FRAMEFIX targets). Real strided layout (cs=64 != BLOCK_K=32 → 2 K-trips,
54+
NON-degenerate, strides ≠ 0). Native = mamba_ssm Triton `_chunk_scan_bwd_dstates_kernel`.
55+
56+
### Structural probe (small, b1 nh8 hd64 ds64 nc8 cs64, numel 262144)
57+
`probe_dstates_struct.py`, routed-vs-native-Triton reference:
58+
59+
```
60+
ROUTED vs ref(1trip) maxdiff=1.7166e+02
61+
ROUTED vs ref(2trip) maxdiff=2.4306e+02
62+
tile[0,0,0] matched(<1e-2)=1/4096
63+
routed/ref2 ratio mean=0.0156 std=2.8486
64+
```
65+
66+
Ratio is random (mean ~0, std ~2.8) → NOT a transpose, NOT a scale, NOT a
67+
precision rounding. The values are structurally permuted.
68+
69+
### Production §P1 (b1 nh112 hd64 ds64 nc64 cs64, numel 29360128, grid (1,64,112))
70+
`parity_prod_dstates.py`, routed-vs-native-Triton (same process, pf via
71+
load_json to dodge the from_ttir↔torch LLVM clash):
72+
73+
```
74+
NATIVE nonzero=29360127/29360128 sum=-203350.2188
75+
ROUTED nonzero=29360128/29360128 sum=-573166.2500 <- FULL tile written (FRAMEFIX fixed the partial write)
76+
MAXDIFF=1.141630e+03
77+
ALLCLOSE_1e-3=False
78+
FAIL
79+
```
80+
81+
FRAMEFIX DID fix the partial-tile write (routed nz 29360128/29360128 = full,
82+
was 2048/262144 / 32-per-tile before FRAMEFIX). But values are wrong.
83+
84+
### Per-kernel parity table
85+
86+
| kernel | routed nz (prod) | MAXDIFF vs native | allclose 1e-3 | PASS |
87+
|---------------------------------|------------------|-------------------|---------------|------|
88+
| _chunk_scan_bwd_dstates | 29360128/29360128 (FULL) | 1.14e+03 | False | **FAIL** |
89+
| _chunk_scan_bwd_dc | not run (lead kernel fails first) | — | — | FAIL (blocked) |
90+
| _chunk_scan_bwd_dcb | not run | — | — | FAIL (blocked) |
91+
| _chunk_scan_bwd_dx | not run | — | — | FAIL (blocked) |
92+
| _chunk_state_bwd_db | not run | — | — | FAIL (blocked) |
93+
| _chunk_state_bwd_dx | not run | — | — | FAIL (blocked) |
94+
| _chunk_state_bwd_ddAcs_stable | not run | — | — | FAIL (blocked) |
95+
96+
**Parity = 0/7 PASS.** The lead kernel fails structurally; the remaining 6 share
97+
the same operand-staging path and are not separately PASS-able, so no kernel is
98+
parity-PASS.
99+
100+
## 3. Root cause (precisely isolated, reproduced)
101+
102+
FRAMEFIX (commit `ff350d21` / gb10 `92fd35d2`, post-walk SBlock layout_map
103+
re-registration) correctly pins the **C accumulator** fragment's
104+
`make_mma_store_layout` so the C carry write becomes full-tile. It does NOT fix
105+
the **A/B operand source** staging:
106+
107+
- The A operand (dout*exp fragment) is restaged to shared via
108+
`reduction.py::_stage_operand_to_shared` → `_emit_copy_stmt`, which lowers to a
109+
SERIAL linear-index-swizzled fill (`for i_j_fused in 0..512`).
110+
- That fill is INCONSISTENT with the cooperative tid-parametrized `ptx_ldmatrix`
111+
read the gemm emits → the gemm consumes a permuted A. A native
112+
`T.copy(A,As)+T.gemm` with identical M/N/K/trans flags is numerically CORRECT
113+
(maxdiff 0.018, 4096/4096), proving the defect is the frontend operand staging,
114+
not the gemm.
115+
- Secondary: the FRAMEFIX C-layout pin is FLAKY — LayoutInference intermittently
116+
raises `Get different layout for dot_c_frag_*` (+1-offset conflict) under
117+
hash-map iteration nondeterminism; some runs fail to compile entirely.
118+
119+
## 4. Timing (recorded fact — NOT a valid speed claim, parity FAIL)
120+
121+
| path | ms / kernel @ §P1 | source |
122+
|-----------------------------------------------------|-------------------|--------|
123+
| native mamba_ssm Triton `_chunk_scan_bwd_dstates` | **1.16 ms** | MEASURED `native_prod.py` |
124+
| routed-through-our-stack (FRAMEFIX) | 2419.43 ms | MEASURED `parity_prod_dstates.py` — parity FAIL |
125+
126+
Routed is ~2080x slower than native (serial linear-index operand fill, MVP
127+
scalar path, PtrAnalysis C++ shim not loaded in-process). Under RULE #1 this ms
128+
is reported as a measured fact only — it is NOT a speed result because the
129+
kernel fails parity. No full-7 routed-bwd sum is reported: 0/7 PASS parity, so
130+
there is nothing valid to sum.
131+
132+
## 5. Path-vs-path summary
133+
134+
| path | status | §P1 dstates ms | full-bwd |
135+
|---------------------------------------|--------|----------------|----------|
136+
| native mamba_ssm full bwd (7 kernels + non-Triton) | reference | 1.16 ms/kernel | ~10 ms-class (full chain) |
137+
| routed-through-our-stack (FRAMEFIX) | **route 7/7, parity 0/7 → NO-GO** | 2419 ms (wrong values) | n/a (parity FAIL) |
138+
| path_c v1 (cppmega.mlx) | reference | — | 905 ms (full chain) |
139+
140+
## 6. GO / NO-GO
141+
142+
**NO-GO.** Achieved: 7/7 route to real TF32 MMA (176 HMMA), FRAMEFIX converts
143+
the partial-tile C-carry write to a full-tile write (measured nz 29360128/29360128).
144+
Blocked: routed output VALUES are structurally wrong (MAXDIFF 1.14e+03 prod /
145+
2.43e+02 small, ratio random) and the routed kernel is ~2000x slower than native.
146+
147+
Remaining defect (RAISED, not papered over, per RULE #1): the A/B GEMM operand
148+
staging (`reduction.py::_stage_operand_to_shared`) emits a serial
149+
linear-index-swizzled shared fill that is layout-inconsistent with the gemm's
150+
cooperative `ptx_ldmatrix` read; plus the FRAMEFIX C-layout pin is flaky under
151+
LayoutInference hash-map nondeterminism. FRAMEFIX fixed the C carry; the next fix
152+
must frame-register / layout-route the operand restage the same way (or use a
153+
layout-aware `T.copy` for A/B as proven correct in the native isolation test).
154+
155+
Mutex released and cleaned up at end of run.

0 commit comments

Comments
 (0)