Part IV · the second capstone and the craft
Debugging warp-specialized kernels & how the compiler lowers them
Lessons 08–12 built kernels with a dozen barriers, warpgroups on opposite phases, scalar mailboxes, and fp32/fp16 aliasing in one TMEM allocation. Here is the brutal fact that makes this paradigm different from everything in gpu_kernels: no compiler pass verifies your phases or your byte counts. One wrong phase, one missing fence, one mis-scoped collective — and the kernel hangs or silently corrupts, and the compiler happily emits it. That is why this paradigm needs a debugging discipline, and why it helps to know how the kernel actually lowers to PTX. This lesson is the craft.
arrive/wait phase, every expect_tx byte count, every fence.proxy_async, every collective scope — is a correctness obligation the toolchain does not check for us. __syncthreads() at least had one well-known footgun (divergent control flow); a Blackwell async kernel has a dozen, and the failures are non-local: a wrong arrival count in the prologue hangs the kernel on the second K-tile, a missing fence corrupts a tile three iterations later. The forced move is not "write better the first time." It is to adopt a repeatable debugging discipline — a worksheet and a methodology — and to be able to read the generated CUDA, because the generated code is the ground truth the silent bugs hide in.
The worksheet — the core mental tool
Before you change a single line, fill in four columns for the kernel. This is the model you validate the generated code against; without it you are pattern-matching symptoms blind. The worksheet forces you to make every implicit assumption explicit — and async bugs are almost always an implicit assumption that turned out false.
| Column | What you write down |
|---|---|
| ROLES | The exact threads / warps / warpgroups / CTAs that issue each async op. "The producer warp issues the TMA" is not enough — write down whether it is warp_id == 0 or wg_id == 0, and which single lane calls elect.sync. Most scope bugs are visible right here. |
| STORAGE | Where each live tile sits at each step: GMEM / SMEM / TMEM / registers. The A-tile is in SMEM stage 0; the accumulator is in TMEM columns 0–127; P aliases S's TMEM bytes after S is consumed. If you can't name the location, you can't reason about reuse. |
| HANDOFF | For every transfer: producer, consumer, the signal object (which named barrier), the arrival count, the phase each side expects, and the fence or drain required. This is the column where deadlocks and corruption live. One row per transfer; if a transfer has no row, that's your bug. |
| LIFETIME | The earliest point each storage slot may be reused. The SMEM stage-0 buffer is free only after the MMA that reads it signals empty; TMEM may be deallocated only after the epilogue's last tcgen05.ld; a TMA store's source SMEM is reusable only after wait_group(0). Premature reuse is the silent-corruption engine. |
Then validate the generated CUDA against the worksheet, checking exactly these:
- Role guards match. The generated guard (
(warp_id_in_cta >> 2) == 0etc.) selects the threads the ROLES column says it should. - Barrier inits precede the guarded branches. Every
mbarrier_initruns before any branch that waits on it — and runs on a thread that actually reaches it. - Collectives aren't narrowed. No
cta_sync()/ cluster-wide op sits inside a lane or warp guard that excludes threads the collective needs. - Arrive/wait phases match. The phase the producer flips to is the phase the consumer is waiting to differ from, every round.
- Lifetimes are respected. TMA drains, TMEM deallocation, and SMEM reuse happen only when the LIFETIME column permits — never earlier.
The 8-step methodology
The same loop every time. It is deliberately boring; the discipline is the value. Note steps 1–3 come before you touch the kernel, and step 8 (measure) comes dead last.
- Reproduce at the MINIMUM failing shape. Shrink batch, sequence, K-tiles until the failure still reproduces. A bug that hangs on tile 2 is far easier to read than one that corrupts at sequence 32768.
- Verify compilation succeeds with the right API and target. A kernel built for the wrong compute capability silently falls off the TMA /
tcgen05path — and then "wrong results" is really "wrong build." - Save & inspect the generated CUDA. Dump it to a file. This is the ground truth; your mental model is a hypothesis about it.
- Build the worksheet (roles / storage / handoff / lifetime) from your intent.
- Validate code vs worksheet using the five checks above.
- Classify the failure into one of the four classes below — the class tells you which tool to reach for.
- Change ONE handoff element at a time. One arrival count, one phase, one fence. Changing two means you can't attribute the result.
- Re-test correctness BEFORE measuring performance. A fast wrong answer is worthless; verify against the reference first, profile second.
Widget · symptom → diagnosis triage
The fastest way into the four classes is to start from what you actually observe. Pick the symptom you're seeing; the triage tells you the likely failure class, the first thing to check, and the right tool to reach for. Everything here is driven by the tables in the rest of the lesson — this is those tables, made queryable.
The four failure classes
Every async-kernel bug lands in one of four buckets. The class is not cosmetic: it tells you whether to reach for synccheck or racecheck, whether you have to restart the process, and whether the bug is in your phases or your byte counts.
Class 1 · Deadlock — symptom: hangs, then an unspecified launch failure
Every deadlock is the same shape: a wait that never returns because the thing it waits on never happens. The art is finding which arrival never came.
| Cause | Why it hangs | Fix |
|---|---|---|
| Arrival-count mismatch | init(128) but a guard lets only one thread arrive → 127 arrivals are forever missing → the wait never returns. | Set the init count to the actual number of arrivers, not the warpgroup size you assumed. |
| Barrier init nested in a warpgroup branch | if wg_id == 1: guards the init, so it never runs on CTA thread 0; the consumer waits on an uninitialized barrier. | Move all barrier inits to the top level, executed by CTA thread 0 only, before any role branch. |
cta_sync() inside a warpgroup branch | cta_sync needs all CTA threads to arrive; threads in the other warpgroup branch never reach it. | Keep cta_sync at the top level, outside any divergent role branch (use warpgroup_sync for within-role sync). |
| Scheduler state skipped | Some threads bypass next_tile() and never advance the work iterator → they wait for a tile that the others already consumed. | Ensure every participating thread runs the scheduler step; don't guard it behind a role mask. |
| Phase desync (off-by-one tile count) | TMA iterates K_TILES but MMA iterates K_TILES − 1 → their phases drift → deadlock on the second tile. | Both producer and consumer must iterate exactly K_TILES. The "second tile" signature is diagnostic. |
| Initial-phase error | Producer and consumer both start on the same phase → the consumer's first wait sees "already flipped" or both block immediately. | Opposite initial phases: producer starts phase = 1 (passes immediately), consumer starts phase = 0 (blocks until the first arrival). Lesson 10's ring. |
Class 2 · Crash / poisoned context
torch.randn on the GPU — because the context is dead, not because your next fix is wrong. Restart the Python process before testing the next fix. Engineers lose hours "fixing" a bug that was already fixed two attempts ago, fooled by errors that are just the corpse of the first crash. When you hit class 2, kill the interpreter.
| Trigger | Why it crashes | Fix |
|---|---|---|
pool.alloc after pool.commit | The SMEM/TMEM pool is finalized at commit; allocating after it corrupts the pool's bookkeeping. | Do all allocs before commit. Treat commit as sealing the allocation arena. |
tcgen05.alloc / dealloc under a lane guard | TMEM alloc/dealloc need full-warp participation; issuing them from one lane (lane_id == 0) is illegal and tears down the context. | Issue tcgen05.alloc/dealloc from the whole warp, not behind an elect.sync / lane mask. |
Missing cta_sync() before tcgen05.dealloc | TMEM is freed while the writeback path is still reading the accumulator out of it → illegal access on freed TMEM. | Fence with cta_sync() after the last epilogue read, before dealloc (this is a LIFETIME-column violation). |
| Out-of-range GMEM / SMEM access | A scheduler index or a shape that isn't tile-aligned addresses past the buffer. | Check the scheduler's tile indices and that the problem shape is aligned to the tile shape; pad if needed. |
Class 3 · Wrong results — the symptom shape classifies the cause
This is the most useful diagnostic table in the lesson. You do not guess at wrong-results bugs — you look at the shape of the error tensor, and the pattern names the cause.
| Symptom shape | What it means |
|---|---|
| ROW-STRIPE corruption (whole rows wrong) | Phase mismatch, tile-index error, or role-ownership confusion — a consumer read the wrong round's tile, or two roles disagree on who owns which rows. |
NaN | Descriptor / operand setup error, or uninitialized accumulators — the math ran on garbage. |
| FINITE but patterned | A stale or partial tile, or an undrained store — the data is real, just from the wrong moment in time. |
And the specific fixes, each one a concrete handoff bug:
| Bug | Mechanism | Fix |
|---|---|---|
tcgen05.commit issued outside elect.sync | All 32 lanes issue the commit → 31 of them produce empty commit groups that signal the barrier prematurely → TMA overwrites the SMEM tile while the MMA is still reading it → corruption. | Issue tcgen05.commit only from the single elected lane (inside elect.sync). |
Missing fence.proxy_async("shared::cta") before a TMA store | The TMA (async proxy) engine does not see the SMEM writes the threads just made → it stores stale data. | Insert the proxy fence so thread-domain SMEM writes are visible to the async-copy proxy before the store is issued. |
Missing cp.async.bulk.commit_group + wait_group(0) after a store | The destination SMEM is reused before the store has drained → the next iteration overwrites bytes still in flight. | Commit and wait on the store group before reusing the staging SMEM (a LIFETIME violation on the store path). |
| Persistent-kernel intermittent failures at small sizes | Phase-reset races between consecutive tiles — one tile's phase reset overlaps the next tile's first wait. | Serialize the per-tile phase reset; the small-size / intermittent signature points straight at a reset race. |
Class 4 · Correct but slow
The kernel is right but far off the roofline. Nine times out of ten, a primitive you thought you used never made it into the generated code — which is exactly why this lesson's second half is about lowering. Check the generated CUDA first; the missing instruction is the bug.
| Symptom in the generated code / profile | Cause | Where to look |
|---|---|---|
No cp.async.bulk.tensor in the generated CUDA | TMA lowering failed — the copy fell back to a thread copy. | The dispatch flag (dispatch="tma"), the target compute capability, and the operand layout the descriptor requires. |
No tcgen05 path | Tensor-core instructions didn't generate; the MMA fell back to a slower path. | The MMA dispatch (dispatch="tcgen05") and operand/accumulator layout. |
| TMA and MMA not overlapping (Nsight Systems timeline alternates instead of overlapping) | Shallow pipeline or phase serialization — the producer and consumer are taking turns, not running together. | Pipeline depth (too few stages), and phases that accidentally serialize the ring. |
| Good at small shapes, poor at large | Register spill / occupancy / staging-buffer pressure as the tile grows. | Register usage (setmaxnreg budget), staging buffer count, and the EPI_N chunking that keeps the epilogue off the spill path. |
The toolchain — which tool for which class
The real tools, mapped to the classes they catch. The point of classifying first (step 6) is that the class chooses the tool; you don't run all of them blindly.
| Tool | Catches | Maps to |
|---|---|---|
compute-sanitizer --tool synccheck | Divergent or invalid barrier usage. | Class 1, causes #2/#3 (init in a branch, cta_sync in a branch). |
compute-sanitizer --tool racecheck | SMEM read-before-write and buffer-reuse hazards. | Class 3, the stale-tile bugs (premature reuse, missing drain). |
compute-sanitizer --tool memcheck | Out-of-bounds / misaligned accesses → illegal access. | Class 2, the crashes (OOB GMEM/SMEM, freed-TMEM read). |
compute-sanitizer --tool initcheck | Uninitialized reads. | Class 3, the NaN path (uninitialized accumulators / operands). |
| Nsight Systems timeline | Copy-engine vs SM activity — they should overlap, not alternate. | Class 4, the no-overlap symptom. |
| Nsight Compute Warp State / stall reasons | stall_long_scoreboard = memory wait · stall_barrier = barrier wait · stall_not_selected = occupancy. The SASS / instruction-mix view confirms tcgen05 / cp.async.bulk are present. | Class 4 (which stall dominates), and confirming class-4 missing-instruction. |
cuda-gdb · info cuda threads | Which warps are parked on a barrier. | Class 1 — see exactly which role is stuck and on which barrier. |
cuda-gdbinfo cuda threads: it shows every warp's PC, so you can see which warps are parked on the wait and which are off doing something else. If the producer warps are all sitting at a wait and the consumer never reached its arrive, you've localized the missing arrival — now go to the worksheet's HANDOFF row for that barrier and check the count and phase.
Inspecting the generated CUDA
The generated code is where the silent bugs live, so reading it is a core skill, not a last resort. In the toolchain you dump the device source from the compiled module:
// dump the generated CUDA for the device kernel
ex.mod.imports[0].inspect_source("cuda")Then search for a handful of landmarks. Each one confirms (or refutes) a specific worksheet assumption:
| Search for | It confirms |
|---|---|
if (threadIdx.x < 1) | The CTA-thread-0 guard — usually the barrier-init block. Verify your mbarrier_init calls live here, before any role branch. |
mbarrier_init | The barriers exist and are initialized; cross-check the counts against the HANDOFF column. |
tcgen05 | The tensor-core path generated (its absence is a class-4 bug). |
cp.async.bulk.tensor | TMA lowering succeeded (its absence is a class-4 bug). |
cta_sync(); | The collectives are where you expect — and crucially, not nested inside a role guard. |
The guard-expression mapping — a debugging tool
When you wrote the kernel you thought in roles: "warpgroup 0," "warp 0," "lane 0." The compiler lowers those to arithmetic on threadIdx.x. To validate the generated code, you have to read the arithmetic back into roles. Memorize this table — it turns an opaque if into a role.
| Your intent | Generated guard |
|---|---|
wg_id == 0 (warpgroup 0) | (warp_id_in_cta >> 2) == 0 |
warp_id == 0 | (warp_id_in_cta & 3) == 0 |
lane_id == 0 | (threadIdx.x % 32) == 0 |
cta_sync — or a TMEM alloc, or a barrier init — ends up inside one of these masks, you've found the bug. A cta_sync() under (warp_id_in_cta >> 2) == 0 will hang forever (class 1): three of the four warpgroups never reach it. A tcgen05.alloc under (threadIdx.x % 32) == 0 will poison the context (class 2): TMEM alloc needs the full warp. Reading the guard back into a role is how you catch a mis-scoped collective before you've burned an afternoon on the symptom.
Widget · spot the bug — barrier-init placement
The single most common class-1 deadlock: a barrier init that the compiler hoists into a role branch (or that you wrote there), so CTA thread 0 — the only thread that should run it — never does. Toggle where the mbarrier_init lives and watch the consumer either proceed or hang. The generated guard is shown so you can read the scope back the way the previous section teaches.
How the compiler lowers it
Now the second half, and the reason class-4 bugs even exist. You've been told to "read the generated CUDA." But what produced it? Knowing the lowering pipeline tells you where a missing tcgen05 or cp.async.bulk went — and it makes the central claim of this lesson concrete: there is no pass that checks your phases.
The ladder, from the tile-level IR you author down to two outputs:
copy, gemm, reduction, with TileLayout-typed accesses and scope ids like T.cta_id / T.thread_id. This is the level you actually wrote.BindTarget — attach the hardware target (compute capability, etc.)LowerTIRx (below). After it, the module is "plain TIR."SplitHostDevice — separate host orchestration from device kerneltcgen05 / cp.async.bulk instructions (or, if lowering failed, without them).The pass that matters: LowerTIRx
LowerTIRx is where your tile primitives become hardware. It is itself two sub-passes, and class-4 bugs are born in the first one:
copy, gemm, reduction — with the dispatcher-selected backend body. This is exactly where dispatch="tma" picks the TMA-vs-cp.async copy backend, and dispatch="tcgen05" picks the MMA backend. If the dispatch flag, target, or operand layout is wrong, the dispatcher picks a fallback — and this is where class-4's "no TMA / no tcgen05 in the output" originates.TileLayout-typed accesses into concrete addresses via addr = data + elem_offset + layout.apply(coord), and lowers scope ids — T.cta_id, T.thread_id — to hardware thread axes via launch_thread. This is where "lane 0" becomes (threadIdx.x % 32) == 0: the guard-expression mapping from earlier is this pass's output.After LowerTIRx the module is plain TIR — the tile abstraction is gone, replaced by concrete addresses, thread-axis guards, and backend instruction bodies.
The only verifier in the pipeline
VerifyMemory, and all it gates is host-vs-device pointer dereference (you can't deref a host pointer on the device, or vice versa). That's it. Nothing verifies barrier phase correctness. Nothing verifies that an expect_tx byte count matches the bytes a TMA actually moves. Nothing checks that a collective isn't trapped in a role mask. The compiler will lower a kernel with a phase desync or a wrong byte count into perfectly valid PTX and emit it without a word of complaint. That absence is the entire reason the worksheet and the 8-step discipline exist. The correctness of every handoff is on you, because no pass is checking.
Why it's built this way — progressive lowering
This is MLIR-style progressive lowering: each stage rewrites the program from one dialect to a lower one — tile primitives → plain TIR → CUDA — preserving the math while moving to a representation closer to the hardware. (The companion ai_compilers track is the whole story of how this lowering works as a discipline.)
The contrast that explains the burden: a compiler like Triton runs a pipeline that does auto-insert cp.async plus the commit/wait pairing for you — its passes own the pipelining, so its passes own that correctness. Here, you author the barriers, the phases, and the byte counts explicitly in the tile IR. The compiler faithfully lowers exactly what you wrote. The flip side of the control that let you hand-build FA4's twelve-barrier choreography is that you own the correctness of every one of those barriers — there is no pass to fall back on.
synccheck for deadlocks, racecheck for stale tiles, memcheck for crashes, initcheck for NaN, Nsight for no-overlap, cuda-gdb for parked warps. Read the symptom shape (row-stripe = phase, NaN = uninit, finite-patterned = stale). Read the generated guards back into roles. And when the context is poisoned, restart Python before you trust the next result. The compiler lowered your tile primitives through LowerTIRx — but it never checked them.
ai_compilers track (the compiler's view of everything we author by hand here). For the profiling tools — Nsight Systems timelines, Nsight Compute stall reasons, reading a roofline — see gpu_kernels/20 · the profiling toolkit.