Skip to content

Commit c5c5891

Browse files
apstenku123apstenku123
committed
docs(triton-ctile-tma): iteration-6 EXECUTEMEASURE — C-tile TMA launches (mod64 fixed, real linux shim) but UTMALDG faults at exec; NO-GO with exact remaining reason
1 parent bbb6aa7 commit c5c5891

1 file changed

Lines changed: 121 additions & 0 deletions

File tree

‎docs/TRITON-CTILE-TMA.md‎

Lines changed: 121 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,121 @@
1+
# Triton C-tile executed TMA — iteration 6 (EXECUTEMEASURE)
2+
3+
Date: 2026-06-07. Host: gb10 (NVIDIA GB10, sm_121 / cap 12.1, aarch64-linux, CUDA 13.2).
4+
Repro source: `/home/dave/source/tilelang` @ tilelang HEAD `a82f63ea` + tvm submodule
5+
commit `2a4f9754` (codegen-llvm alloca alignment fix landed this iteration).
6+
7+
## Executive result — NO-GO (C-tile TMA launches but FAULTS at execution)
8+
9+
The C-tile TMA now **launches** (the mod64 tensorMap-create abort is fixed) and the
10+
UTMALDG instruction is **present and reached at runtime**, but it **faults with
11+
`CUDA_ERROR_ILLEGAL_INSTRUCTION` inside `tl::tma_load` (copy_sm90.h:96)**. The kernel
12+
does not run to completion, so NO valid §P1 execution-ms delta can be measured. The key
13+
question ("does one real coalesced TMA load measurably drop §P1 ms?") is **not yet
14+
answerable** — the load does not execute correctly.
15+
16+
## What this iteration discovered and fixed (executed evidence)
17+
18+
### 1. gb10 was missing the linux PtrAnalysis C++ shim (root cause of "UTMALDG=0")
19+
The prior REBUILDGROUND result reporting "tma_load=2 ... LDG drop to 38" was NOT
20+
reproducible on gb10. On gb10 the only built shim artifacts were **Mach-O darwin** `.so`
21+
files (rsync'd from the Mac). With no loadable `_triton_frontend_cxx` on linux, the
22+
frontend fell back to the **MVP scalar path**: `ctx.ptr_states` empty ->
23+
`_emit_ptrstate_tile_load_copynode` **never fires** -> the C-tile grounding code at
24+
`memory.py:1391-1404` is **dead** -> generated CUDA had **zero** tma/cp.async/mbarrier,
25+
UTMALDG=0, LDG=128.
26+
27+
Fix: built the **real linux shim** (not the stub) on gb10 against the existing Triton
28+
aarch64 install:
29+
- `TRITON_INSTALL_DIR=/home/dave/source/triton-aarch64-install` (has `libTritonIR.a` etc.)
30+
- MLIR/LLVM from `/home/dave/.triton/llvm/llvm-ac5dc54d-almalinux-arm64`
31+
- pybind11 from the venv
32+
- output: `_triton_frontend_cxx.cpython-313-aarch64-linux-gnu.so` (build dir
33+
`poc/triton_frontend/_cxx/build_linux`, also copied into `_cxx/build/` for
34+
auto-discovery; stale darwin `.so` removed).
35+
36+
With the real shim, the **codegen** SASS shows the grounding firing:
37+
| build | UTMALDG | LDG | spill(STL+LDL) | HMMA |
38+
|-------|---------|-----|----------------|------|
39+
| OFF (prologue_opt=False) | 0 | 66 | 658 | 32 |
40+
| OPT (prologue_opt=True) | **10** | 96 | 272 | 32 |
41+
42+
So the C-tile grounding DOES emit real TMA in codegen (UTMALDG 0 -> 10). Note: LDG did
43+
**not** drop to 38 (it rose 66 -> 96); the prior "38" figure is not reproducible here.
44+
45+
### 2. tensorMap 64-byte alignment was NOT actually fixed for this kernel (now fixed)
46+
The aligned(64) fix at `codegen_c_host.cc:438` only covers the **C-host** backend. This
47+
kernel is lowered via the **LLVM backend** (`codegen_cpu.cc`), where the TMA descriptor is
48+
materialized through a `tvm_ffi_any` stack alloca emitted as `align 8`:
49+
```
50+
%stack_ffi_any539 = alloca [35 x %0], align 8 <- BEFORE
51+
```
52+
At align 8 the descriptor address is mod64 != 0 and the runtime validator FATALs:
53+
`tensorMap address must be 64-byte aligned, but got ... mod64=16` (also seen mod64=48).
54+
55+
Fix (committed, tvm `2a4f9754`): in `codegen_cpu.cc` `CreateIntrinsic` stack-alloca, pin
56+
the `tvm_ffi_any` branch to `setAlignment(Align(64))` — mirroring the existing
57+
`tensormap` branch. After rebuilding `libtvm_compiler.so` (ninja, no GPU needed):
58+
```
59+
%stack_ffi_any539 = alloca [35 x %0], align 64 <- AFTER
60+
```
61+
Result: the mod64 abort is **gone**; the kernel **launches**.
62+
63+
### 3. Remaining blocker — UTMALDG faults at execution (NO-GO)
64+
With alignment fixed, the §P1 OPT kernel launches (grid=(1,64,112) block=(256,1,1)) and
65+
then crashes. `compute-sanitizer --tool memcheck` pinpoints it exactly:
66+
```
67+
Illegal instruction
68+
at void tl::tma_load<...>(const CUtensorMap_st&, ...)+0x18f0 in copy_sm90.h:96
69+
by thread (128,0,0) in block (0,0,0)
70+
Device Frame: _chunk_scan_bwd_dstates_kernel_kernel+0x1050 in tvm_kernels.cu:171
71+
```
72+
i.e. the `cp.async.bulk.tensor.2d` (UTMALDG) instruction itself faults. The descriptor is
73+
now statically valid (passes `copy.cc:988 is_one(global_stride[0])` via the grounding)
74+
and 64-byte aligned, but the descriptor encoding is **invalid at execution** on sm_121 for
75+
the real C-tensor layout (grounded innermost ds-stride=1, with the outer k-axis stride and
76+
global dims coming from the symbolic runtime args). The TMA load is genuinely reached and
77+
genuinely wrong — it is not codegen-only and not a silent fallback.
78+
79+
Note: a **plain TileLang matmul** (TMA and even non-TMA) also raised
80+
`CUDA_ERROR_ILLEGAL_INSTRUCTION` on this same gb10 config in spot checks, suggesting a
81+
broader sm_121/CUDA-13.2 codegen-arch interaction may compound the descriptor issue. The
82+
OFF dstates path (HMMA-only, no TMA) runs cleanly and passes parity, so the dstates
83+
illegal-instruction is specifically tied to the TMA/async pipelined path (sanitizer frame =
84+
`tl::tma_load`), not the GEMM.
85+
86+
## Parity (unaffected paths still correct)
87+
- §P1 dstates OFF path: **MAXDIFF = 4.882812e-04, allclose@1e-3 PASS** (unchanged; the
88+
alignment fix does not touch the OFF/scalar path).
89+
- §P1 dstates OPT path: cannot be parity-checked — kernel crashes before producing output.
90+
- Native baseline (mamba_ssm triton): **~1.13 ms** median (min 1.130), N=60 — matches the
91+
prior 1.12 ms reference.
92+
93+
## GO / NO-GO
94+
**NO-GO for "one real executed C-tile TMA load with measured ms delta."** Two real
95+
blockers were cleared this iteration (missing linux shim; LLVM-backend descriptor
96+
alignment), advancing the state from "UTMALDG=0, grounding dead" to "UTMALDG launches".
97+
The exact remaining reason (RULE#1) is: **the grounded C-tile TMA descriptor faults with
98+
illegal instruction inside `tl::tma_load` (copy_sm90.h:96) on sm_121** — the descriptor
99+
encoding / sm90 bulk-tensor path is not valid for this de-monomorphized layout on this GPU.
100+
Resolving it requires a correct sm_121 TMA descriptor (correct global dims/box/strides for
101+
the partially-grounded C layout, and likely an sm_100/sm_120-class bulk-tensor path rather
102+
than the sm90 template) — a deeper fix than this iteration's grounding+alignment scope.
103+
104+
## Reproduce from HEAD
105+
```
106+
# build linux shim (no GPU):
107+
cmake -S poc/triton_frontend/_cxx -B poc/triton_frontend/_cxx/build_linux -GNinja \
108+
-DMLIR_DIR=/home/dave/.triton/llvm/llvm-ac5dc54d-almalinux-arm64/lib/cmake/mlir \
109+
-DLLVM_DIR=/home/dave/.triton/llvm/llvm-ac5dc54d-almalinux-arm64/lib/cmake/llvm \
110+
-DTRITON_INSTALL_DIR=/home/dave/source/triton-aarch64-install \
111+
-Dpybind11_DIR=<venv>/lib/python3.13/site-packages/pybind11/share/cmake/pybind11 \
112+
-DPython3_EXECUTABLE=<venv>/bin/python
113+
ninja -C poc/triton_frontend/_cxx/build_linux
114+
cp poc/triton_frontend/_cxx/build_linux/_triton_frontend_cxx.*.so poc/triton_frontend/_cxx/build/
115+
# rebuild tvm with the alignment fix (no GPU):
116+
ninja -C build lib/libtvm_compiler.so
117+
# codegen (UTMALDG 0->10):
118+
PYTHONPATH=.../build:. python poc/triton_frontend/_test_harness/tridao_parity/sass_dstates.py
119+
# exec (GPU): launches, then illegal instruction in tl::tma_load:
120+
CUDA_LAUNCH_BLOCKING=1 .../compute-sanitizer python /tmp/run_opt_only.py
121+
```

0 commit comments

Comments
 (0)