Two answers to the same idle GPU

Small kernels have two costs that have nothing to do with arithmetic: the CPU pays launch setup for every kernel it submits, and the GPU drains its pipeline between one kernel’s tail and the next kernel’s start. As kernels got faster (single-digit microseconds on modern parts), both overheads stayed constant — until they were the runtime.

CUDA graphs and kernel fusion are the two standard answers, and they answer different questions:

• CUDA graphs are a work-submission mechanism. They move the CPU cost to a one-time instantiation and let the GPU’s own scheduler walk the dependency DAG. The kernels inside are untouched — same cubins, same SASS, same grids.
• Kernel fusion is a compiler transformation. It rewrites the kernels themselves into one kernel, so intermediate values never leave the chip and the pipeline never drains between stages.

This page is about the first mechanism in depth — how NVIDIA actually implements graphs — because understanding exactly what the graph scheduler does (and deliberately does not do) is what defines the gap a SASS-level fusion pass works in.

What a CUDA graph is

A cudaGraph_t is a DAG. An operation is a node; a dependency is an edge. The documented node types: kernel launches, CPU function calls, memcpy, memset, empty nodes, event record/wait, external semaphore signal/wait, conditional nodes, memory alloc/free nodes, and child graphs (nested graphs as single nodes).

Two properties of edges matter for everything below:

• An edge constrains order, not placement. Once a node’s dependencies complete, “scheduling is left up to the CUDA system” — independent branches may run concurrently, on whatever SMs and copy engines are free.
• A default edge (edge type 0) is a full dependency with memory-synchronizing behavior: the downstream node observes complete execution of the upstream node, with its memory writes visible. Between two kernel nodes, that is a grid-wide barrier plus a memory flush — the same semantics as two kernels on one stream.

CUDA 12.3 added edge data (ports + dependency type), whose one non-default use so far is Programmatic Dependent Launch — more on that below, since it is the only place the “full barrier between kernels” rule bends.

Building the DAG: explicit API vs stream capture

The explicit graph API is what it looks like: cudaGraphCreate, then cudaGraphAddNode / cudaGraphAddKernelNode / cudaGraphAddMemcpyNode with dependency arrays. Tedious, but exact.

Stream capture is how real code gets graphed — including library internals the caller never sees. cudaStreamBeginCapture(stream) puts the stream in capture mode; from then on, work “launched” into it does not execute — it is appended to an internal capture graph. cudaStreamEndCapture returns the graph:

cudaStreamBeginCapture(stream);
kernel_A<<<..., stream>>>(...);   // recorded, not run
libraryCall(stream);              // records whatever it launches
kernel_C<<<..., stream>>>(...);
cudaStreamEndCapture(stream, &graph);

Cross-stream fork/join survives capture through events. Recording an event in a capturing stream produces a captured event — a handle to the set of nodes recorded so far; a cudaStreamWaitEvent on it pulls that stream into the same capture graph and adds the dependency edges. The rule that keeps this sane: every stream forked out of a capture must be joined back to the origin stream before cudaStreamEndCapture, or the whole capture fails.

Whichever path built it, the cudaGraph_t at this point is a template — a portable description. Nothing has been validated against the hardware, and nothing can run yet.

Paying the setup cost once

cudaGraphInstantiate(&graphExec, graph, ...) turns the template into an executable graph (cudaGraphExec_t). The programming guide describes the contract in one sentence: instantiation “takes a snapshot of the graph template, validates it, and performs much of the setup and initialization of work with the aim of minimizing what needs to be done at launch.”

Concretely, this is where everything a per-kernel stream launch would redo every time gets done once:

• Validation — topology checks, node-type restrictions, context/device legality.
• Work-description construction — the launch descriptors for every node (kernel image reference, grid/block dims, parameter buffer contents) are pre-built instead of assembled per launch.
• Scheduling precomputation — the topology is analyzed so the device can walk it without host involvement; since CUDA 12.0, node creation order is used as a scheduling heuristic (earlier-created nodes are preferred), which is why node order in well-behaved capture code tends to match intuition.
• Snapshot semantics — the exec graph is frozen as of instantiation. Later edits to the template do nothing to it.

Instantiation is the dominant one-time cost of adopting graphs. NVIDIA’s own microbenchmarks (CUDA 11.8 vs 12.6):

The right column is the point of the whole mechanism: on Ampere and later, re-launching an already-uploaded graph is constant-time on the CPU — about 2.5 µs whether the graph holds ten kernels or a thousand (pre-Ampere stays O(n)). Compare with 2 µs + 200 ns per node under CUDA 11.8, and with the per-kernel cost of stream launches the graph replaced.

Upload is the device-side half: cudaGraphUpload (or the first launch, implicitly) stages the instantiated work descriptions onto the GPU. First launch is still O(n) on the CPU for exactly this reason; the upload flag and explicit call exist so that cost can be moved off the critical path.

What the device does with the DAG

cudaGraphLaunch(graphExec, stream) submits the whole graph as one item of stream work. From there, the host is out of the loop: the GPU front-end walks the pre-built topology, and — in NVIDIA’s words — dependencies are “enforced with hardware accelerations, instead of relying solely on CUDA streams and events, when possible.” Independent branches issue as soon as their inputs are ready; dependent nodes wait on device-side completion signaling rather than host round-trips.

What does not change is what a kernel node is: an ordinary grid launch of an ordinary cubin. Each node carries its own launch configuration, its own parameter buffer, its own register allocation and shared-memory plan, baked into its SASS long before the graph existed. And each default edge between two kernel nodes carries full stream-order semantics: the upstream grid runs to completion, its writes are flushed and visible, and only then does the downstream grid begin.

Later toolkit releases stretched the execution model further, always within the same “whole DAG, launched once” frame:

• Conditional nodes (12.3+) — IF / WHILE / SWITCH nodes whose body is a child graph, with the condition value evaluated on the device (cudaGraphSetConditional from kernel code). Loops and branches without leaving the graph.
• Device graph launch — a kernel launching an instantiated, uploaded graph from the device, enabling device-side schedulers that pick which graph runs next with no host involvement.

Living with a frozen graph

The snapshot that makes launch cheap makes change awkward. CUDA’s answer is a tiered update story:

• Whole-graph update — cudaGraphExecUpdate applies parameters from a topologically identical template. The pairing of old to new nodes is deterministic only under strict ordering rules (same API call order per stream, same dependency-array order, consistent sink-node order), so the supported pattern is: re-run the exact same capture code, update, and if the update is rejected, destroy and re-instantiate.
• Individual node update — cudaGraphExecKernelNodeSetParams and friends patch one node’s parameters (kernel args, copy addresses) in place, skipping topology checks. The cheap path when only a few pointers change — which in practice is the common case: same network, new tensor addresses.
• Enable/disable — cudaGraphNodeSetEnabled turns kernel/memcpy/memset nodes into empty nodes and back, so one superset graph can be specialized per launch.

What you can never do is change the shape: topology, node types, kernel functions, grid geometry — any of those moves and you re-instantiate. A graph is a replay device for a fixed computation with retargetable pointers.

What graphs deliberately don't do

Everything above is launch mechanics. Nothing in the architecture touches the kernels themselves — and that is a design decision, not an oversight. The graph’s unit of reasoning is the node; a kernel node is an opaque reference to a cubin image. Concretely, no matter how the graph is scheduled:

• Intermediate data still round-trips through global memory. A dependent edge means writes flushed and visible; the downstream kernel re-reads them from L2/DRAM. There is no mechanism for one node to hand a value to the next in registers or shared memory — the nodes never share an SM-resident state.
• The pipeline still drains at every dependent edge. Except for PDL’s prologue overlap, node N+1’s blocks start after node N’s blocks finish. Ten fused-able stages in a graph still pay nine grid-wide barriers and nine memory round-trips; they just stop paying the CPU to schedule them.
• Each node keeps its own resource plan. Register counts, shared memory, occupancy — all per-kernel, all frozen in SASS at ptxas time. Nothing reallocates registers across what used to be a kernel boundary, because as far as the runtime is concerned there still is one.
• Shapes are frozen. Grid geometry is baked at capture; retuning it means re-capture and re-instantiation.

Graphs answer “the CPU can’t feed the GPU fast enough.” Fusion answers “the GPU keeps stopping between stages and spilling intermediates to DRAM.” A real workload usually wants both — and the fusion half is exactly where a SASS-level toolchain operates, because the fusion decision has to happen to the machine code the graph nodes point at.

How we fuse kernels

Every technique below was first proven by hand, at Triton source level, in a series of megakernel campaigns on Jetson AGX Orin (sm_87) — whole decode tokens of Qwen3.5-4B, Nemotron 3 Nano 4B, and Gemma 4 fused into persistent kernels or staged CUDA graphs, dropped into stock llama.cpp and vLLM for 1.26–1.29× end-to-end. sass2mlir’s job is to turn each of them from a hand-rolled, per-model source hack into a compiler pass over post-ptxas SASS — applicable to any shipped cubin, with no source and no Python, and composable with CUDA graphs: fuse first, then capture the fused result.

The unit of fusion is always the same move: delete a grid barrier and move the consumer’s math into the producer’s tile loop — into its epilogue, its prologue, or through a counter-based mini-join. What varies is the legality argument and the cost model, and that is exactly what a compiler formalizes.

Fusions that delete stages

Epilogue fusion — the core technique

The producer kernel’s accumulator never touches memory: the consumer’s elementwise math (activation, norm, residual add) runs in the same registers, in the same instruction stream, right after the final multiply-add. The campaign version fused a 4-tap depthwise conv + SiLU into a GEMV epilogue; the norm variant factors rstd out as a scalar applied once per output element.

Legality is an ownership argument: the epilogue is race-free iff each output element is owned by exactly one tile — which is what made conv+SiLU safe against an in-place conv-window shift, and what makes the general case checkable.

How sass2mlir enables it: in the lifted SSA this is mechanical. The pass takes two adjacent kernels (adjacency comes from sass2mlir intercept’s run.csv or a captured graph), inlines the consumer’s body after the producer’s stores, and replaces each consumer LDG of the intermediate with the SSA value feeding the producer’s STG — def-use does the wiring. The ownership legality is proved, not assumed: uniformity analysis plus address analysis show the store and load addresses coincide per-thread and are disjoint across tiles. And the campaigns supply the cost model: the prologue-side SwiGLU fusion measured net-negative because the producer→consumer thread mapping wasn’t injective, so every consumer thread re-loaded and re-computed — the pass fuses epilogue-side on 1:1 ownership and refuses otherwise.

Counter-based mini-joins — N producers, one consumer, no barrier

When N tiles produce the slices one consumer stage needs as a whole row, you don’t need a grid barrier: each tile ATOM.ADDs a per-group counter as it stores (the arrive is the release), and the tile that observes the count reach N-1 — the last arriver — loads the full row and runs the consumer math inline. Used in the campaigns to fold a gated RMS-norm × SiLU into a DeltaNet epilogue (8 tiles per head, self-resetting counters) and to fold a grouped norm into the last-arriving head of each group. Each use deletes an entire stage and its barrier.

How sass2mlir enables it: the pattern is recognizable at SASS level — producer stores → full barrier → consumer whole-row loads — and the rewrite is a counter allocation plus a conditional consumer block predicated on the arrive’s return value. It’s a structural pass (like warp-specialization, it rewrites control flow rather than peepholing inside it), and the release/acquire ordering is small enough to pin down completely in a corpus-level test rather than trust by inspection.

Segment merging — back-to-back GEMVs as one tile walk

Several small GEMVs in a row (qkv projections, ffn gate+up) each pay the persistent-loop straggler imbalance and each begins every tile iteration with a dependent map load (plane, row from a table). The fusion concatenates the weight planes into one tile space walked as a single strided loop, and decodes (plane, row) arithmetically from the tile index — threshold compares, bitwise-identical, free. Measured: the map-load form runs at 30 vs 62 GB/s and 115 vs 38 regs; the merge itself cut the monolith 240→168 regs and straggle 9.4%→3.3%.

How sass2mlir enables it: two passes. mapload-to-arithmetic finds the dependent LDG/LDC at a loop head feeding an address computation — a three-node def-use chain in the IR — and replaces it with IMAD/ISETP/LOP3 threshold decode when the table is piecewise-constant. The kernel-pair merge is the same inlining machinery as epilogue fusion, minus the data dependence.

Grid-barrier synthesis — fusion’s enabling infrastructure

Every stage-deleting fusion above crosses a grid-wide dependency, so the fused kernel needs barriers the source compiler never emitted. The campaign protocol, hardened over three models: one ATOM.E.ADD.STRONG.GPU acq_rel arrive per CTA, an RMW spin on the same counter, an unconditional trailing acquire by everyone including the last arriver, then BAR.SYNC — with slots never reused within a token and self-cleaned at the end (~1.1 µs per barrier; 257 per token after the barrier diet, down from 321).

How sass2mlir enables it: grid-barrier-synth inserts exactly this sequence at fusion boundaries — and, critically, enforces the two legality rules the campaigns learned from miscompiles: never substitute a volatile-load spin for the RMW spin (that variant silently broke at 321-barrier scale), and every cross-CTA buffer read needs an L2-coherent or atomic load, because L1 is not coherent across SMs. A SASS-level tool audits what actually shipped — the campaigns hit a Triton stack that silently dropped .cv cache modifiers, so the source’s intent was not the binary’s behavior. The protocol itself is a small concurrent algorithm — small enough to reason about completely, and exercised at 321-barrier scale in regression.

Fusions that fix the inner loop

Dequant strength reduction — I2F → PRMT

Q4/int8 weight decode spends its ALU in integer→float conversion, and I2F issues at 16/SM/clk — an eighth of the FFMA rate. The replacement is bit-exact and all full-rate ops: XOR-bias the bytes (0x80808080), PRMT each into the mantissa of 2²³, and an FADD against the magic constant recovers the signed value. Verified exact over all 256 int8 values; worth 1.19→1.30× end-to-end on a 4B model.

How sass2mlir enables it: a peephole pass over I2F-on-extracted-quant chains — trivially recognizable in the dialect. And “exact for all 256 values” is a finite claim: the pass is validated by exhaustively checking every possible input — the strongest correctness story a compiler transform can have, no theorem prover required.

Tiles-in-flight scheduling — the u4 GEMV form

Keep four weight tiles’ loads in flight before the first reduction, values held in bf16 until the multiply. Measured 139→171 GB/s — and bitwise identical, because the accumulation order never changes: it’s purely an instruction-scheduling transformation that buys memory-level parallelism.

How sass2mlir enables it: this is the concrete target pattern for the instruction-scheduling pass on the roadmap — and because the transform is bitwise-identical by construction, the existing crosscheck / byte-diff machinery is the oracle. No new verification infrastructure needed.

Load widening and cache policy

Scalarized LDG.E.U8 streams (the signature of struct-field-decode weight layouts) widen to LDG.E.64/.128 when alignment is provable — worth up to 3.7× on an affected plane in the campaigns (47→175 GB/s). Single-use weight streams want the evict-first policy bit so they don’t thrash L2.

How sass2mlir enables it: widening is gated on an alignment proof from const-bank offsets and strides — the campaigns’ counter-case (the same hint on fp16 scale offsets miscompiled at rel 1.81) is exactly the premise the pass must verify. Cache/eviction policy is a control-field attribute edit on the IR, the same class of move as the reuse/yield edits — except sass2mlir writes the SASS field directly, so it bypasses the toolchain bugs the campaigns hit (evict_first + .cg crashing ptxas; .cv silently dropped) and only has to encode the hardware’s legality, not the compiler’s.

Occupancy-driven register allocation

The campaigns’ central law: spills are nearly free; occupancy is not. The winning configurations were always found by capping registers to reach the next co-residency rung — SMs × floor(65536 / (threads × regs)) — e.g. 96 regs/58 spills buying a 5th CTA per SM (70 CTAs on 14 SMs) and the best token time, while every “fewer CTAs, spill-free” variant lost.

How sass2mlir enables it: the register allocator’s free mode gets the co-residency ladder as its cost model — pick the cap that lands on the next rung, let it spill. One hard-won legality constraint comes along: an aggressive cap that spills ~160 regs/thread through local memory was observed to break grid-barrier visibility (run-to-run drift, root-caused to the spill regime), so spill placement around spin/barrier regions is a legality question, not a cost question.

PDL: overlap without fusion (sm_90+)

For the kernel boundaries that shouldn’t be fused — register-incompatible stages, or kernels you don’t control — Hopper’s programmatic dependent launch is the graph-edge relaxation described above, driven by hand: the producer signals GRIDDEPCONTROL.LAUNCH_DEPENDENTS in its epilogue, the consumer launches early, runs its independent prologue (weight prefetch doesn’t touch the producer’s output), and executes GRIDDEPCONTROL.WAIT only before its first dependent load.

How sass2mlir enables it: the handshake is two instructions inserted at SASS level — pdl-synth finds the consumer’s first load that depends on the producer’s output (def-use through the intermediate buffer), puts the WAIT there, and puts the signal after the producer’s last store to that buffer. The dependent-edge bubble shrinks to a prologue’s worth of overlap while both kernels keep their own register budgets. (sm_90 and later only — on sm_87 the equivalent answer is the fusions above.)

Selective, not monolithic

The campaigns measured both directions, so the design rule is empirical, not aesthetic: a fully monolithic megakernel loses whenever its stages have heterogeneous register demands — Qwen3.5’s monolith pays its fattest stage’s 252-reg allocation everywhere, lands at 2 CTAs/SM and 57 GB/s, and loses to ~155 lean kernels in a CUDA graph (24.8 vs 30.4 tok/s) — while Nemotron’s monolith wins (+28%) because its register budget survived the merge. And SASS can’t repeal this: registers-per-thread is fixed per launch, so a fused kernel pays max-over-stages no matter how good the allocator is.

So sass2mlir’s fusion is selective by construction:

1. Keep the workload as lean staged kernels — the shape that already wins.
2. Apply epilogue fusions, mini-joins, and segment merges where the legality proofs hold and the fused kernel stays on its occupancy rung.
3. Overlap the remaining boundaries with PDL on Hopper-class parts.
4. Hand the shrunken sequence back to a CUDA graph — fuse first, capture after.

Delivery reuses what already exists: sass2mlir intercept captures the kernel sequence and its launch geometry, the passes fire through sass2mlir apply, and the patched cubin drops into the original slot. The campaigns’ interposer shows the extreme end of that road — a captured 554-node vLLM decode graph substituted at capture time for a one-node graph — but the everyday case is quieter: the same graph, fewer nodes, each one the same kernel it was, minus the stages that never needed to exist.

References

• CUDA Programming Guide — CUDA Graphs: node/edge semantics, stream capture rules, instantiation contract, update rules, conditional nodes, graph memory nodes.
• NVIDIA Developer Blog — Constant Time Launch for Straight-Line CUDA Graphs: the 11.8 → 12.6 instantiation and launch-overhead measurements quoted above, and the node-creation-order scheduling heuristic.
• NVIDIA Developer Blog — Enabling Dynamic Control Flow in CUDA Graphs with Device Graph Launch: device-side graph instantiation flags, upload, and launch.