Skip to content

Commit 8e49017

Browse files
apstenku123apstenku123
committed
docs(triton-sm121f-tma): GB10 C-tile TMA EXECUTE attempt -- NO-GO (UTMALDG illegal on sm_121a, sm_121f 48KB-shared cap)
EXECUTED evidence on gb10 (sm_121, CUDA 13.2): the grounded C-tile TMA lowers to real UTMALDG.2D SASS (tma_load=2) but does NOT execute. sm_121f (ArchFix family target) fails ptxas with 0x18020=98KB shared > 0xc000=48KB cap; sm_121a compiles but the executed UTMALDG faults cudaErrorIllegalInstruction on first launch (same copy_sm90.h:96 class). Diagnosis (sm_121f makes TMA execute) FALSIFIED. No-regression under sm_121f: path_c F0 chunk_precompute mma.sync GEMM PASS (6.10ms, 10.4x, parity cb 1.22e-04 ss 3.74e-06); SIMT vadd PASS; committed dstates SIMT route parity 4.88e-04 OPT ~1102ms unchanged. ArchFix 48KB static-shared cap is a real global regression for >48KB kernels (sm_121a permitted full Blackwell shared). coalescing-pays-off: UNANSWERED -- no executing TMA exists on GB10.
1 parent 784a3c3 commit 8e49017

1 file changed

Lines changed: 91 additions & 0 deletions

File tree

‎docs/TRITON-SM121F-TMA.md‎

Lines changed: 91 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,91 @@
1+
# sm_121f ArchFix + C-tile TMA EXECUTE attempt on GB10 (Blackwell sm_121)
2+
3+
Date: 2026-06-08. Host: NVIDIA GB10 (CC 12.1, aarch64-linux), CUDA 13.2, NVRTC/nvcc 13.2.
4+
tilelang HEAD: `a80e4fa8` (merge/upstream-codegen-reorg). All GPU exec on gb10 under mutex.
5+
6+
## Goal
7+
After the ArchFix (CC 12.x -> sm_121f FAMILY arch, commit a80e4fa8), re-enable the
8+
iter-6 routed C-tile TMA load and measure whether the coalesced TMA EXECUTES and
9+
drops the dstates kernel §P1 ms vs the prior routed best (~1102 ms).
10+
11+
## VERDICT: NO-GO. The C-tile TMA still does NOT execute on GB10.
12+
13+
The diagnosis ("sm_121 fully supports cp.async.bulk.tensor TMA; sm_121f is the fix")
14+
is FALSIFIED on this hardware by executed evidence:
15+
16+
1. The grounded C-tile TMA lowers to REAL `tl::tma_load` / `UTMALDG.2D` SASS
17+
(tma_load_count=2; SASS shows `UTMALDG.2D`, `UTMACCTL.PF`, `USETMAXREG`).
18+
2. Under **sm_121f** (the ArchFix family target) the TMA kernel **fails to COMPILE**:
19+
`ptxas error: uses too much shared data (0x18020=98KB, 0xc000=48KB max)`.
20+
The `f` family target caps STATIC shared at 48KB; the double-buffered TMA tile
21+
plan needs ~98KB. It never reaches execution.
22+
3. Under **sm_121a** (arch-specific; the documented alternative) the SAME kernel
23+
COMPILES (sm_121a permits the full Blackwell shared), but at runtime the
24+
executed `UTMALDG.2D` **FAULTS** with `cudaErrorIllegalInstruction`
25+
("an illegal instruction was encountered") on the FIRST launch -- the same
26+
copy_sm90.h:96 class fault as iter-6. compute-sanitizer could not even flush a
27+
report: the illegal instruction hard-aborts the CUDA context/process.
28+
(OFF baseline parity passed 4.882812e-04 immediately before the OPT fault.)
29+
30+
So: sm_121f = won't compile (48KB cap); sm_121a = compiles but UTMALDG faults.
31+
Neither makes the C-tile TMA execute. The coalescing-pays-off question is therefore
32+
UNANSWERABLE from an executing TMA -- there is no executing TMA on GB10.
33+
34+
## The instruction
35+
`copy_sm90.h:96` emits (CUDA >= 12.8 CTA variant):
36+
`cp.async.bulk.tensor.2d.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint`
37+
-> compiles to `UTMALDG.2D ... desc[...]`. The descriptor is a by-value
38+
`__grid_constant__ const CUtensorMap arg1_desc`. The executed UTMALDG traps illegal
39+
on GB10 regardless of f/a suffix.
40+
41+
## ArchFix regression (sm_121f 48KB static-shared cap)
42+
The ArchFix is NOT free: the family `f` target caps static shared at 48KB whereas
43+
the prior `a` suffix permitted the full Blackwell shared. Any tilelang kernel
44+
needing >48KB static shared that compiled under the pre-ArchFix sm_121a now FAILS
45+
to compile under sm_121f. Observed directly: the routed dstates kernel's 80KB-shared
46+
variant (`0x14000`) ptxas-fails under sm_121f / sm_120f / plain sm_121 but compiles
47+
clean under sm_121a. This is a real global regression introduced by a80e4fa8.
48+
49+
## No-regression checks under sm_121f (clean committed HEAD)
50+
- path_c F0 `chunk_precompute_fwd_cuda_prim` (the tensor-core GEMM): **PASS**.
51+
COMPILE ok; codegen `tl::mma_sync x4` + ldmatrix x4 (real mma.sync); per-call
52+
median 6.10 ms vs serial 63.76 ms = 10.4x; parity cb=1.22e-04 ss=3.74e-06
53+
dac=4.885e-04 (gate 5e-4) PASS. The production path_c GEMM executes correctly
54+
under sm_121f -- mma.sync HMMA is fine.
55+
- Trivial SIMT vadd kernel: PASS (maxdiff 0) under sm_121f.
56+
- Committed dstates SIMT route (HEAD default, no TMA): parity OFF/OPT 4.882812e-04
57+
PASS; OPT median ~1101-1102 ms (== prior routed best), OFF ~1478 ms. Unchanged.
58+
- Standalone fp16 GEMM (128x128 tile, HMMA.16816.F32, NOT path_c): FAULTS illegal
59+
instruction at the HMMA (tvm_kernels.cu:45) under BOTH sm_121f AND sm_121a AND
60+
with disable_tma. This is config-specific (the path_c F0 mma.sync path works);
61+
reproduces at CLEAN committed HEAD -> a pre-existing GB10/HMMA-tile condition,
62+
NOT caused by the ArchFix. Flagged honestly; does not affect path_c.
63+
64+
## §P1 ms verdict
65+
- routed EXEC §P1 ms now: ~1101-1102 ms (committed SIMT route, no TMA). Same as 1102.
66+
- coalescing-pays-off: UNANSWERED. The TMA never executes, so no executed
67+
coalesced-TMA ms exists to compare. ms_delta vs 1102 = 0 (no change; TMA dead).
68+
- gap to native 1.12 ms: still ~1.0x10^3x off; unchanged. The routed best stays
69+
banked-addressing-fold 1102 ms (native ~1.12 ms).
70+
71+
## Reproduce (gb10, under mutex)
72+
Shim on PYTHONPATH so PtrAnalysis subprocess seeds ptr-state and grounding fires:
73+
```
74+
PYTHONPATH=poc/triton_frontend/_cxx/build:. \
75+
TL_ARCH_SUFFIX=a TL_TMA_ROUTE=1 \
76+
python poc/triton_frontend/_test_harness/tridao_parity/measure_dstates_tma_exec.py
77+
# -> tma_load_count=2; OFF parity PASS; OPT launch -> illegal instruction
78+
```
79+
(TL_ARCH_SUFFIX / TL_TMA_ROUTE are experiment-only env gates added during this
80+
investigation; the committed HEAD source is unmodified -- the TMA route is NOT the
81+
committed default. The committed default is the SIMT route at ~1102 ms.)
82+
83+
## Recommendation
84+
1. REVERT or re-scope ArchFix a80e4fa8: the sm_121f family target's 48KB static-shared
85+
cap is a real regression for >48KB kernels. If TMA is wanted, sm_121a (arch-specific)
86+
is required for the shared capacity -- but the UTMALDG still faults, so TMA is not
87+
usable on this GB10 regardless. Keep CC 12.x on sm_121a (the pre-ArchFix behavior)
88+
to preserve full shared for the SIMT/GEMM kernels that DO work.
89+
2. The C-tile TMA is a dead end on this GB10 silicon/driver (UTMALDG illegal).
90+
Stay on the committed SIMT route (~1102 ms). The next real lever is 128-bit
91+
vectorization of the SIMT loads (the swizzled-shared-dest float4 blocker), NOT TMA.

0 commit comments

Comments
 (0)