|
| 1 | +# TRITON-ROUTE-FAST — MEASURE phase (gb10 sm_121, mutex held) |
| 2 | + |
| 3 | +Status: **NO-GO** on the goal (correct + native-fast Tri-Dao bwd through our stack). |
| 4 | +All numbers below are MEASURED live on gb10 (aarch64-linux, sm_121) at HEAD |
| 5 | +`eade89be` (StructuredGemm: cooperative gemm operand staging + deterministic C |
| 6 | +layout, BUG1/BUG2). RULE#1 honored: routed ms reported, parity is fabrication-free, |
| 7 | +the exact remaining defect is named. |
| 8 | + |
| 9 | +## Headline |
| 10 | + |
| 11 | +| metric | MEASURED | target | |
| 12 | +|---|---|---| |
| 13 | +| parity PASS (prod §P1) | **0/7** | 7/7 @ 1e-3 | |
| 14 | +| parity PASS (small, real 2-K-trip) | **0/7** | 7/7 @ 1e-3 | |
| 15 | +| routed ms/kernel (dstates, prod §P1) | **2547.6 ms** | ~1.16 ms | |
| 16 | +| native ms/kernel (dstates, prod §P1) | **1.14228 ms** | (baseline) | |
| 17 | +| routed vs native | **~2230x SLOWER** | ~1x | |
| 18 | +| 7/7 route + MMA in CUDA | **YES** (mma/ldmatrix present all 7) | yes | |
| 19 | + |
| 20 | +The serial-staging bottleneck the StructuredGemm fix targeted is NOT the live |
| 21 | +bottleneck in this environment: routed ms is ~2547 ms (vs the prior 2419/2535 ms), |
| 22 | +i.e. essentially UNCHANGED. The dominant cost is a SEPARATE upstream defect (below). |
| 23 | + |
| 24 | +## Root cause (verified at source, not inferred) |
| 25 | + |
| 26 | +`poc.triton_frontend.ptr_analysis.shim_available() == False` on gb10. The C++ |
| 27 | +`_triton_frontend_cxx` PtrAnalysis shim is UNBUILT for aarch64-linux (only macOS |
| 28 | +`-darwin.so` artifacts exist; final link of vendored `libTriton*.a` macOS archives |
| 29 | +fails on aarch64). The frontend emits a RuntimeWarning and falls back to the MVP |
| 30 | +scalar path, which lowers multi-element tile operand LOADS as per-element |
| 31 | +masked-GATHER (`carry_index_*`) sequences. That gather PERMUTES the `[hd,BK]` |
| 32 | +operand → structural permutation (MAXDIFF 2.6–2.8e2 small / 1.0–1.17e3 prod — a |
| 33 | +permutation, not a transpose/scale/precision miss), AND the per-element gather + |
| 34 | +serial epilogue are the ~2547 ms cost. |
| 35 | + |
| 36 | +The BUG1 cooperative `T.copy` operand staging (eade89be) IS present in the routed |
| 37 | +CUDA, but the operand still ARRIVES permuted from the scalar gather upstream of the |
| 38 | +staging copy, so parity stays broken and the gather cost dominates. The real fix is |
| 39 | +to BUILD the C++ shim on gb10 so multi-element tile loads use the vendored |
| 40 | +triton-shared PtrAnalysis (no gather) — a separate, pre-existing infra defect. |
| 41 | + |
| 42 | +## Per-kernel parity (small cfg: b1 nh8 hd64 ds64 nc4 cs64 ng2 s256, REAL 2-K-trip, cs64≠BK32, strides≠0) |
| 43 | + |
| 44 | +| kernel | status | MAXDIFF | note | |
| 45 | +|---|---|---|---| |
| 46 | +| `_chunk_scan_bwd_dstates` | FAIL | 2.78e2 | correctly marshalled (== proven prod harness map); full tile written, values permuted | |
| 47 | +| `_chunk_scan_bwd_dc` | SEGV | — | generic positional marshaller OOB: walker reorders params (carry_index gather buffers FIRST), positional packing → wrong buffers | |
| 48 | +| `_chunk_scan_bwd_dcb` | SEGV | — | same marshaller defect | |
| 49 | +| `_chunk_scan_bwd_dx` | SEGV | — | same marshaller defect | |
| 50 | +| `_chunk_state_bwd_db` | SEGV | — | same marshaller defect | |
| 51 | +| `_chunk_state_bwd_dx` | SEGV | — | same marshaller defect | |
| 52 | +| `_chunk_state_bwd_ddAcs_stable` | N/A | — | atomic ddA accumulator output (no directly-comparable buffer) | |
| 53 | + |
| 54 | +prod §P1 (S4096 c64 g8 H112 P64 N64): `_chunk_scan_bwd_dstates` FAIL MAXDIFF=1.165e3, |
| 55 | +full tile written (routed nz 29360128/29360128 == native nz), values permuted. |
| 56 | + |
| 57 | +The 6 non-dstates kernels share the same underlying defect; the SEGV is a harness |
| 58 | +marshalling limitation (param reorder), not a new fault — each needs per-kernel |
| 59 | +hand-mapping like dstates' prod harness, after which they'd hit the SAME gather |
| 60 | +permutation. |
| 61 | + |
| 62 | +## ms table (parity-PASS kernels only — there are NONE; dstates ms shown for the record) |
| 63 | + |
| 64 | +| path | ms/kernel @ prod §P1 | source | |
| 65 | +|---|---|---| |
| 66 | +| native triton `_chunk_scan_bwd_dstates` | 1.14228 | MEASURED gb10 | |
| 67 | +| routed (our stack) `_chunk_scan_bwd_dstates` | 2547.6 | MEASURED gb10 (parity FAIL — informational only) | |
| 68 | +| path_c v1 (full chain) | 905 (prior) | EXTRAPOLATION / prior run | |
| 69 | + |
| 70 | +RULE#1: routed ms is informational here because NO kernel passes parity. There is no |
| 71 | +honest "routed-our-stack fast" number to report as a result. |
| 72 | + |
| 73 | +## NO PERF REGRESSION on other frontend GEMM paths (user CRITICAL ask) |
| 74 | + |
| 75 | +The StructuredGemm change (eade89be) touches exactly 3 python files |
| 76 | +(`op_emitters/reduction.py`, `frame_register.py`, `__init__.py`), inside |
| 77 | +`map_tt_dot`'s operand-staging closure. |
| 78 | + |
| 79 | +- Live CUDA-ms before/after for `fla_dot_exp2` / `matmul` / `dot_reduce_atomic` / |
| 80 | + `dot_reduce_atomic_trans_b` is **NOT obtainable on gb10**: the full TTIR→PrimFunc |
| 81 | + lowering for ANY `tt.dot` kernel is BLOCKED here because `_lower_ttir` needs the |
| 82 | + same UNBUILT C++ shim (`round_trip_through_cxx_shim`; in-process MLIR has no `tt` |
| 83 | + dialect → "Dialect `tt' not found"). The block is IDENTICAL before and after. |
| 84 | +- Measurable surface = frontend unit suite. BEFORE/AFTER by file-swap (the 3 files |
| 85 | + reverted to `88106531` vs HEAD `eade89be`), identical command: |
| 86 | + `pytest test_dot_reduce_atomic test_op_emitters_reduction test_op_mapping |
| 87 | + test_op_emitters_memory test_reducer_corpus test_builder_context`: |
| 88 | + |
| 89 | + | build | result | |
| 90 | + |---|---| |
| 91 | + | AFTER (eade89be) | **121 passed, 10 skipped, 3 failed** | |
| 92 | + | BEFORE (88106531) | **121 passed, 10 skipped, 3 failed** | |
| 93 | + |
| 94 | + IDENTICAL — same 3 failures both ways (pre-existing: 2× ranked-offset store bounds, |
| 95 | + 1× tma-fallback-when-triton-available). ZERO correctness regression from the change. |
| 96 | +- 7/7 Tri-Dao bwd kernels still compile to CUDA with MMA/ldmatrix tensor-core |
| 97 | + instructions (route intact). |
| 98 | + |
| 99 | +Honest verdict: NO correctness regression on other frontend paths (unit suite |
| 100 | +byte-identical pass/fail before/after). Live CUDA-ms before/after for the named GEMM |
| 101 | +kernels could not be measured because their lowering is shim-gated in this env |
| 102 | +(blocked identically pre/post change) — reported as a measurement gap, not a PASS. |
| 103 | + |
| 104 | +## GO/NO-GO |
| 105 | + |
| 106 | +**NO-GO.** Routed Tri-Dao bwd is neither correct (0/7 parity) nor fast (2547 ms vs |
| 107 | +1.14 ms native). The blocker is the UNBUILT aarch64 C++ PtrAnalysis shim forcing the |
| 108 | +scalar masked-gather path; the StructuredGemm operand-staging fix is in place and |
| 109 | +correct in isolation but is downstream of the permuting gather, so it cannot rescue |
| 110 | +parity or perf until the shim is built on gb10. |
0 commit comments