Skip to content

Commit 294b515

Browse files
apstenku123apstenku123
committed
docs(nvptx-ab): NVPTX backend feature-gap A/B — NO-GO, source-codegen wins
Per-kernel A/B of TVM LLVM NVPTX (target.build.nvptx, codegen_nvptx.cc) vs default source-codegen (codegen_cuda.cc -> NVRTC). MEASURED gb10 sm_121: plain elementwise compiles on NVPTX (real PTX provenance); every T.gemm / tensor-core kernel FAILS (unknown intrinsic tl.tl_shuffle_elect; LLVM bad-signature SIGABRT on F0). codegen_nvptx CreateIntrinsic has zero mma/ldmatrix/cp.async/mbarrier coverage. Routed-dstates prologue confirmed backend-orthogonal (IR-level serial loop). NVPTX not default for any path.
1 parent ad8f28b commit 294b515

1 file changed

Lines changed: 106 additions & 0 deletions

File tree

‎docs/NVPTX-BACKEND-AB.md‎

Lines changed: 106 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,106 @@
1+
# NVPTX backend vs source-codegen+NVRTC — A/B for all our kernels
2+
3+
**Date:** 2026-06-07 · **HW:** NVIDIA GB10 (sm_121, aarch64-linux) · **TVM:** db5fe57 · tilelang db5fe57c
4+
**Mutex:** acquired/released cleanly on gb10 (`nvptxab`). GPU-adjacent work ran gb10-only; no local Metal dispatch.
5+
**Provenance commit:** `6f014dfb` (probes) — this doc + analysis.
6+
7+
## TL;DR / NET VERDICT — NO-GO for NVPTX as a default on any path
8+
9+
The TVM LLVM NVPTX backend (`src/target/cuda/llvm/codegen_nvptx.cc`, `target.build.nvptx`)
10+
is **feature-incomplete for every tensor-core kernel we ship.** It compiles only
11+
plain elementwise TIR. The instant a kernel uses `T.gemm` / pipelined shared-mem
12+
copies — i.e. ALL of our real GEMM/SSD/fp8 prims — NVPTX cannot lower the TileLang
13+
intrinsics and either throws `unknown intrinsic` or hard-`SIGABRT`s in LLVM.
14+
15+
**Source-codegen (`codegen_cuda.cc` -> NVRTC -> ptxas) remains the only working path
16+
for 4 of the 5 kernels.** There is no EXEC-ms A/B to report for the tensor-core
17+
kernels because NVPTX never produces a runnable module for them. That feature gap
18+
is the finding — it is reported honestly, not hidden.
19+
20+
## HOW to switch the backend (surveyed)
21+
22+
- Default path: `tilelang.jit(...)` -> lowers TIR -> `target.build.tilelang_cuda`
23+
-> `codegen_cuda.cc` emits **CUDA C++ text** -> NVRTC -> PTX -> cubin.
24+
- Backend kinds live behind FFI globals: `tvm.ffi.get_global_func("target.build.tilelang_cuda")`
25+
(default text codegen) vs `tvm.ffi.get_global_func("target.build.nvptx")` (LLVM NVPTX,
26+
`codegen_nvptx.cc`). To force NVPTX you build a `Target({"kind":"nvptx","mcpu":"sm_121"},
27+
host=Target("llvm"))` and feed the **same lowered device module** to `target.build.nvptx`.
28+
- `tilelang/utils/target.py` `SUPPORTED_TARGETS` does NOT list `nvptx` — there is no
29+
first-class tilelang flag to select it; you go through the raw TVM FFI build func
30+
(what the probes do). This alone means NVPTX is off the supported-frontend path today.
31+
32+
## codegen_nvptx.cc intrinsic coverage (read at db5fe57)
33+
34+
`CodeGenNVPTX::CreateIntrinsic` (lines 260-293) handles ONLY:
35+
`tvm_warp_shuffle{,_up,_down}` (32-bit), `tvm_warp_activemask`, `atomic_add` (32-bit),
36+
plus thread/block-index sreg reads. **Everything else falls through to
37+
`CodeGenLLVM::CreateIntrinsic`**, which `TVM_FFI_THROW`s on any unknown op
38+
(codegen_llvm.cc:1483).
39+
40+
There is **NO** handler for any of the intrinsics our GEMM kernels emit, measured
41+
in the device TIR of a plain fp16 matmul:
42+
`ptx_ldmatrix`, `ptx_mma`, `ptx_arrive_barrier`, `ptx_arrive_barrier_expect_tx`,
43+
`ptx_fence_barrier_init`, `ptx_init_barrier_thread_count`, and TileLang's own
44+
`tl.tl_shuffle_elect`. No `tvm_mma_sync`, no `cp.async`, no `ldmatrix`. The
45+
tensor-core / async-copy / mbarrier surface is entirely absent.
46+
47+
## Per-kernel A/B table
48+
49+
| # | Kernel | source-codegen | NVPTX | Winner | Notes |
50+
|---|--------|----------------|-------|--------|-------|
51+
| 5 | plain elementwise add (1M f32, `nvptx_elem_probe.py`) | OK | **OK** | tie (both compile) | NVPTX emits real PTX: `// Generated by LLVM NVPTX Back-End`, `.version 8.8 .target sm_121`, `.visible .func add2_kernel`. Provenance MEASURED on gb10. No tensor cores -> no gap. No EXEC-ms taken (trivial bandwidth kernel, not representative). |
52+
| 5b | plain fp16 `T.gemm` matmul 1024³ (`nvptx_ab_probe.py`) | OK via `tilelang.jit` (src 5651 chars, `__global__`+`tl::`+`ldmatrix`) | **FAIL** | source-codegen | NVPTX: `InternalError: unknown intrinsic tl.tl_shuffle_elect` in `CodeGenNVPTX::AddFunction` -> codegen_llvm.cc:1483. TIR carries ptx_ldmatrix×2, ptx_mma×2, 4× ptx_*barrier. MEASURED gb10. |
53+
| 2 | path_c F0 `chunk_precompute_fwd_cuda_prim` (the 14.5× TC GEMM, 24.57->1.13ms) | OK (src backend, prior MEASURED 14.5×) | **FAIL (HARD ABORT)** | source-codegen | NVPTX `SIGABRT` at LLVM `Instructions.cpp:757` `"Calling a function with bad signature!"` — uncatchable. Default src path unaffected; 1.13ms stands on source backend. No NVPTX ms exists. |
54+
| 1 | ROUTED Triton `_chunk_scan_bwd_dstates` (§P1 shim, 1479ms serial prologue) | OK (src backend, only working path) | **FAIL** | source-codegen | Same T.gemm/tensor-core intrinsic wall as #5b/#2 (identical `tl.tl_shuffle_elect` + ptx_ldmatrix/ptx_mma/ptx_*barrier set). NVPTX never compiles it -> no EXEC-ms either backend. See prologue note below. |
55+
| 3 | path_c B2 `chunk_scan_combine_bwd_cuda_prim` (~905ms) / §27 gemm | OK (src backend) | **FAIL (by construction)** | source-codegen | T.gemm prim; emits the same unhandled intrinsic set proven by #5b/#2. Not separately run (identical wall, conserve gb10 time). EXTRAPOLATION from the shared root cause, labeled as such. |
56+
| 4 | path_c fp8 native MMA prim (`fp8_matmul_path_c.py`) | OK (src backend) | **FAIL (by construction)** | source-codegen | fp8 MMA needs `ptx_mma` + ldmatrix + barriers — none in codegen_nvptx. EXTRAPOLATION from shared root cause. |
57+
58+
MEASURED: #5, #5b, #2 (run on gb10 2026-06-07). EXTRAPOLATION (shared root cause, not
59+
separately executed): #1 exec, #3, #4 — flagged honestly.
60+
61+
## WHY (LLVM sched/regalloc vs NVRTC->ptxas)
62+
63+
The intended question — does LLVM's instruction-scheduler / register-allocator /
64+
vectorizer beat NVRTC->ptxas — **cannot be answered for any tensor-core kernel**,
65+
because LLVM NVPTX never gets a valid module to schedule: lowering aborts at the
66+
intrinsic-translation stage, before any scheduling/regalloc runs. The cause is NOT
67+
codegen quality; it is a missing intrinsic-lowering surface in `codegen_nvptx.cc`
68+
(no mma/ldmatrix/cp.async/mbarrier translation). For the one kernel where both
69+
backends DO compile (plain elementwise), the kernel is pure memory traffic with no
70+
arithmetic to schedule, so it is not representative and we took no ms. Net: on our
71+
workload the comparison is moot — source-codegen wins by being the only path that
72+
emits a tensor-core kernel at all.
73+
74+
## Routed-dstates PROLOGUE — backend-orthogonal (confirmed by construction)
75+
76+
The 1479ms `_chunk_scan_bwd_dstates` cost is dominated by a **serial prologue that
77+
exists in the TIR before codegen** (an IR-level serial-loop artifact). Two facts make
78+
the orthogonality airtight: (1) the text-codegen backend choice happens strictly
79+
AFTER the TIR is fixed, so swapping `tilelang_cuda` for `nvptx` cannot alter the
80+
serial-loop structure; (2) NVPTX never compiles this kernel anyway (same tensor-core
81+
intrinsic wall), so it cannot have improved it. The real fix is the frontend
82+
prologue-opt (GEMMify / auto-fuse the serial reduction), NOT a backend swap.
83+
84+
## Recommendation
85+
86+
**NVPTX default: NO** — for any path. It is not feature-complete for our kernels
87+
(zero tensor-core/async-copy/mbarrier coverage), it is not even exposed as a
88+
supported tilelang target, and on the one elementwise case it merely ties. Making it
89+
default would break every GEMM/SSD/fp8 prim we ship (RULE #1: that is a silent-breakage
90+
trap, rejected). Keep source-codegen (`codegen_cuda.cc` -> NVRTC) as the default.
91+
92+
If NVPTX is ever wanted, it is a large engineering project: add `CreateIntrinsic`
93+
lowering for `ptx_mma`/`ptx_ldmatrix`/`cp.async`/the four mbarrier intrinsics/
94+
`tl.tl_shuffle_elect`, fix the bad-signature abort, and only THEN can an honest
95+
sched/regalloc A/B even begin. Out of scope here; flagged for the user to extend.
96+
97+
## Reproduce
98+
99+
```
100+
# gb10 only (GPU mutex), tilelang+tvm SOURCE at /home/dave/source/tilelang
101+
ssh gb10 'cd /home/dave/source/tilelang && source /home/dave/cppmega-venv/bin/activate
102+
python /tmp/nvptx_elem_probe.py # -> NVPTX OK (real PTX)
103+
python /tmp/nvptx_ab_probe.py # -> NVPTX FAIL: unknown intrinsic tl.tl_shuffle_elect
104+
python /tmp/nvptx_ops_probe.py # -> lists the 6 unhandled ptx_* intrinsics; default jit OK'
105+
# probes tracked at tilelang poc/nvptx_ab/ (commit 6f014dfb)
106+
```

0 commit comments

Comments
 (0)