Profile a single-request decode on an H100 and you will find the GPU idle a measurable fraction of the time — not for lack of work, but waiting on the CPU to tell it what to do next. A decode step is hundreds of tiny kernel launches, each costing ~5µs of host-side dispatch before a single flop executes. This essay is the full anatomy of that overhead wall: the launch arithmetic, the CUDA-graph escape hatch with its bucketing tax, the compiler route through torch.compile, the radical fix of persistent kernels — and the Python scheduler sitting underneath all of it.
01 Anatomy of a wasted microsecond
What a kernel launch actually costs, and why the arithmetic gets worse every GPU generation.
Every cudaLaunchKernel call is a host-side transaction: the CPU packs arguments, crosses into the driver, the driver validates parameters and enqueues a command on a stream, and control returns. The GPU may not start the kernel for several more microseconds — it has to drain the queue, resolve stream dependencies, and find an SM. Measured end to end, one launch costs 3–10µs on a quiet system (PyTorch's own CUDA-graphs guide quotes 5–15µs; eBPF-based profilers like launchlate measure 5–20µs in ideal conditions, ballooning to hundreds of microseconds under contention). The number that matters for serving is the steady-state one: budget ~5µs per launch and you will not be far wrong.
Now count the launches. A transformer decode step is not one kernel; it is a long sequence of small ones. Per layer, a fused serving stack like vLLM issues roughly: one RMSNorm, one fused QKV GEMM, one attention kernel (paged FlashAttention), one output projection, one more RMSNorm, one fused MLP (gate+up), one down-projection, and — under tensor parallelism — one or two all-reduces. Call it 8–12 dispatches per layer before fusion, 2–4 after the fusions that every serious stack applies. For an 80-layer 70B-class model that is 160–320 launches per token. The brief version of the math:
Nlaunch × L ≈ 200 × 5µs = 1.0 ms of pure dispatch overhead per token
One millisecond per token is a hard ceiling of 1,000 tokens/second before the GPU does any useful work — and that ceiling is set entirely by the host. On top of the driver cost sits Python: the GIL, argument packing in PyTorch's dispatcher, and the serving framework's per-iteration Python — scheduler decisions, sampling-parameter prep, detokenizer calls. In vLLM that host-side Python is typically another few hundred microseconds per step, roughly constant regardless of model size.
Here is the part most treatments get backwards: the overhead wall does not bite hardest on the biggest models. Write the per-token time budget explicitly:
Ttoken = Tgpu + Nlaunch·L + Tpython
For a 70B model in FP16 at batch 1 on an H100, Tgpu is dominated by streaming 140 GB of weights through 3.35 TB/s of HBM3: ~42 ms per token. The 1 ms of launch overhead is 2% — noise against the memory wall. But take an 8B model: 16 GB of weights is ~4.8 ms of memory time, launches are ~80 × 5µs = 0.4 ms, Python is ~0.3 ms. Overhead is suddenly ~13% of every token. And the trend is hostile: each GPU generation raises memory bandwidth (shrinking Tgpu) while launch latency stays stubbornly ~5µs. On an 8 TB/s-class GPU the same 8B decode drops toward 2 ms of GPU time against ~0.7 ms of overhead — more than a quarter of the step spent waiting on the host. Faster GPUs make the CPU bottleneck worse, not better. That is the overhead wall, and it is why small-model, low-batch, interactive serving — the exact shape of chat products — is where this fight happens.
Batching divides the wall. The overhead terms N·L and Tpython are per step, not per token: at batch 64, that 1 ms of dispatch becomes ~15µs per token, and the wall all but disappears into the memory traffic of 64 concurrent KV-cache reads. This is why the optimization ordering matters and why the wall is specifically a low-batch phenomenon. Continuous batching — Orca's innovation, now universal — keeps the batch full, which is the cheapest overhead optimization ever shipped: it does not reduce N or L, it divides them by the batch size. Everything in this essay is about what to do when batching has done all it can — the interactive tail, where batch 1–8 is the product requirement, not a tuning failure.
Open nsys on a decode step and the wall is visible directly. The CPU "Runtime API" row shows a dense palisade of cudaLaunchKernel calls; the GPU row shows kernels with dark gaps between them. Each gap is the interval between one kernel's completion and the next kernel's start — time the SMs sit idle while the host dispatches. When the gaps are as wide as the kernels, you are CPU-bound and no amount of kernel optimization will help. The diagram below is that timeline, drawn to the same arithmetic: ten kernels, 116µs of execution, 54µs of gaps.
Both timelines show the same ten kernels on the same 4 px/µs scale. In (a) the host dispatches each kernel individually and the GPU idles ~6µs between completions; in (b) one graph launch replays the recorded sequence with the kernels packed back-to-back. The takeaway is the ratio: 54µs of gaps against 116µs of work means nearly a third of the step is host overhead — and that ratio is independent of kernel speed. Faster kernels make it worse.
How do you confirm which overhead you have? In nsys, enable the CUDA API trace: the Runtime API row shows cudaLaunchKernel calls on the CPU timeline, the GPU row shows kernel executions. Line up each kernel's start with the launch that enqueued it, and the delta decomposes into three parts: Python think time between launches (visible as gaps on the CPU row itself), driver and queueing delay, and stream-dependency stalls. Dense CPU row plus gappy GPU row means launch overhead; a gappy CPU row means Python. For production attribution, eBPF tools like launchlate put a uprobe on cudaLaunchKernel and a kprobe on kernel entry and histogram the delta — under contention it routinely measures 30–80µs per launch, several times the quiet-system number, which is why lab benchmarks understate the wall.
One more multiplier is arriving via speculative decoding. A draft-then-verify step runs k draft-model forwards plus one target-model verify per accepted token — each phase its own launch sequence. EAGLE-style drafting roughly doubles to triples N per output token while the GPU work per token stays flat, so the host tax per accepted token grows. Every technique that adds steps per token — speculative decoding, multi-token prediction, agentic tool loops — pays the overhead wall once per step.
The overhead wall is a ratio, not a constant: (N·L) / Tgpu. Everything that shrinks Tgpu — more bandwidth, smaller models, shorter contexts — raises the ratio. B200-class memory bandwidth will make batch-1 serving of 8B models roughly one-third host overhead unless the launches go away.
02 CUDA graphs: capture once, replay forever
The standard escape hatch — and the bucketing tax it charges.
CUDA graphs attack the N·L term directly. Instead of dispatching hundreds of kernels per step, you run the step once under capture: the driver records every kernel, memcpy, and dependency into a graph object without changing what executes. Then you instantiate it and replay the whole sequence with a single cudaGraphLaunch. The graph captures the structure — which kernels, in which order, reading and writing which addresses — not the data. One launch replaces hundreds, and the GPU-side scheduler walks the recorded dependency list with no host involvement. This is the single highest-leverage overhead optimization in serving, which is why every major stack does it.
The catch is that a captured graph is rigid. During capture you may not allocate memory, synchronize with the host, or change the kernel sequence — and every tensor address is baked in. A decode step with batch size 7 and a decode step with batch size 8 are different graphs. vLLM's answer, and the industry-standard one, is bucketing: pre-capture graphs for a fixed set of batch-size "buckets" and pad every real batch up to the nearest bucket. vLLM's default capture set is 1, 2, 4, then multiples of 8 from 8 to 248, then multiples of 16 from 256 up to max_cudagraph_capture_size (512 tokens on Hopper, 1024 on datacenter Blackwell) — 51 graphs in the default configuration. A batch that arrives larger than the biggest bucket skips graph mode entirely and falls back to eager launches.
Bucketing trades three things against each other. First, padding waste: a batch of 9 runs in the 16-bucket and 44% of its lanes are dummy rows doing useless work. Averaged over a uniform batch-size distribution from 1 to 512, the vLLM default wastes 3.3% of compute on padding; a coarse power-of-two scheme (10 buckets) wastes 24.5%. The waste concentrates exactly where latency matters most — small batches. Second, capture memory: NVIDIA reserves ~64 KB of device memory per kernel launch recorded in a graph (pre-CUDA 12.4 / driver 550+; the figure is documented in PyTorch's CUDAGraph Trees guide). At ~150 launches per decode graph that is ~9.4 MB per bucket, or ~0.5 GB for the default 51. That sounds modest until speculative decoding enters: wider draft/verify graphs multiply it, and a recent vLLM regression that expanded default capture sizes burned 5.12 GiB on graph memory and starved the KV cache to 0.94 GiB — the engine could barely start (vLLM PR #54418). Third, recapture latency: the first time a bucket is hit, the step runs under capture — a one-time stall on the order of a second — and any control-flow change (a new bucket, a different attention mask shape) invalidates the recording.
The replay path itself has a subtle mechanism worth knowing. vLLM does not rewrite graph node parameters per step; it keeps static padded buffers — input tensors allocated once at bucket size — and before each replay it copies the live batch's data into them with a few small host-to-device copies. The graph then reads the static addresses it was captured with. The per-step host work collapses to: copy inputs in (a few µs), one cudaGraphLaunch, copy the sampled token out. Everything between the copies is device-side.
The rigidity has a precise API shape, worth stating exactly. During stream capture you may not: allocate or free device memory (the graph records against a dedicated pool, which is why vLLM pre-allocates its static buffers); synchronize with the host (no .item(), no .cpu(), no blocking copies — one sync aborts the capture); or branch on values computed inside the captured region. Multi-stream capture is legal but every cross-stream dependency must be recorded as an event the graph can replay; serving stacks capture on a single stream and keep it simple.
There is a middle setting between "recapture everything" and "static buffers plus copies": cudaGraphExecUpdate, which rewrites node parameters — a memcpy node's source address, a kernel's launch arguments — in an already-instantiated graph without re-recording. It is the API embodiment of "structure static, data dynamic." vLLM's decode path mostly does not need it, because the static-buffer-plus-copy trick covers the common case with less machinery; update is the right tool when addresses change but topology does not, for example KV-cache page tables that grow across steps.
Three bucketing schemes evaluated on batch sizes 1–512 with a uniform distribution: coarse powers of two (10 buckets), the vLLM default (51), and a fine step-4 scheme (128). Padding waste collapses fast and then barely moves — going from 51 to 128 buckets buys 2 points of waste for 2.5× the capture memory. The operating region sits where the waste curve has flattened but memory is still modest; vLLM's default lands in it, which is why 51 buckets survived years of tuning. The memory curve is the 64 KB-per-launch figure times ~150 launches per decode graph; wider speculative-decode graphs shift the whole orange curve up, which is exactly the failure mode that burned 5 GiB in vLLM PR #54418.
Two practical consequences follow. First, graphs are a decode optimization: prefill shapes vary with prompt length, so mainstream stacks do not graph prefill — the capture cost and bucket explosion are not worth it. Decode is one token per sequence per step, a fixed, highly repetitive structure; it is the 95%-of-steps case and graphs erase nearly all of its dispatch overhead. Second, graphs compose badly with dynamism anywhere in the recorded region: a new sampler, a changed mask, an extra all-reduce from a different parallelism layout — each is a new graph. The bucket count is not just a memory knob; it is a configuration-space knob, and every new axis multiplies it.
03 torch.compile and cudagraph trees
The compiler route: fewer launches by construction, graphs without hand bucketing.
If CUDA graphs reduce the cost per launch, torch.compile reduces the number of launches. Inductor traces the model's Python, fuses what it can — RMSNorm, residual adds, and quantization scales collapse into single Triton kernels; pointwise epilogues fold into their producers — and emits one kernel where eager PyTorch issued five. A decode step that launched 200 kernels might launch 120 after compilation. Since overhead is N·L, cutting N by 40% cuts the host tax by 40% before any graph is captured. vLLM has moved this way explicitly: recent versions compile the model with torch.compile and then capture CUDA graphs over the compiled region, stacking both reductions.
The interesting part is how Inductor handles the dynamism that hand-written graph code manages with buckets. CUDAGraph Trees — enabled by torch.compile(mode="reduce-overhead") — re-record a graph for every unique input shape the compiled function sees, organizing the recordings in a tree keyed by shape. A decode loop that visits batch sizes 7, 13, and 31 ends up with three recordings, captured lazily on first encounter, replayed thereafter. No hand-maintained bucket list; the shapes you actually serve are the shapes you capture. The price is the same 64 KB-per-launch memory per recording, which is why the docs warn that rapidly changing shapes make re-recording unprofitable and suggest padding to a few fixed shapes — buckets by another name — or setting torch._inductor.config.triton.cudagraph_skip_dynamic_graphs=True to graph only the static-shape regions.
Two more compiler mechanisms matter here. Triton is the reason "fewer, bigger kernels" is now the default shape of optimized code: writing a fused kernel used to mean hand-authoring CUDA, a weeks-long project; Triton (and embedded DSLs like ThunderKittens, which packages the load-compute-store-finish pattern for rapid kernel iteration) made fusion a days-long project, so stacks fuse aggressively and the launch count falls by construction. Autotuning is the other half: max-autotune benchmarks GEMM tile configurations per shape instead of trusting heuristics, which matters because a decode GEMM at batch 1 is a skinny, memory-bound problem where the wrong tile config wastes the little parallelism available.
The tradeoff is paid upfront: compile time. A full inductor compile of a 70B-class model takes minutes on first run (tens of seconds for 8B-class; both are estimates that vary wildly with max-autotune and machine), and it is paid again on every code change during development. Serving stacks amortize it by compiling once at engine startup — vLLM's compilation_config exists precisely to make that one-time cost configurable — and caching the artifacts. The subtler cost is graph breaks: any Python-side data-dependent control flow in the traced region (an if on a tensor value, a Python loop over a dynamic range) splits the compiled graph, and each fragment is launched separately — reintroducing the overhead the compiler was hired to remove. The discipline this imposes is real: the hot loop must be traceable, branch-free Python, or the compiler cannot help you.
A concrete graph break, since this is where compiler projects die: sampling. The model forward is traceable, but if next_token == eos: stop branches on a tensor value — a break at the exact boundary between "model" and "everything else." Every production stack resolves it the same way: compile and capture the model forward, leave sampling, detokenization, and scheduling outside the compiled region. The break is not eliminated; it is placed where it costs one launch instead of two hundred.
Two operational notes. First, cache the artifacts: Inductor writes compiled kernels to TORCHINDUCTOR_CACHE_DIR — point it at persistent storage or every container restart repays the minutes-long compile. Second, max-autotune is a CI-time activity, not a deploy-time one: benchmark tile configs once per (model, GPU) pair, ship the chosen configs, and compile in default or reduce-overhead mode at startup. Autotuning during a rollout is how you get paged.
Graphs and compilers are complements, not substitutes. torch.compile shrinks N (fewer, bigger kernels); CUDA graphs shrink L (one launch instead of hundreds). vLLM's stack does both — compile first, capture the compiled region — because the product of the two reductions is what finally makes the host overhead negligible.
04 Persistent kernels: never leave the GPU
The radical fix — one launch for the whole session, with the loop on the device.
Graphs reduce a step to one launch. Persistent kernels ask why there is a per-step launch at all. The idea: launch one cooperative kernel at session start that never exits. Inside it, a device-side loop pulls work from a queue in global memory — an atomic work counter, or a mailbox the host updates through mapped memory — runs the decode step, grid-synchronizes, and repeats. The host's per-token job collapses to a single 4-byte write ("here is the next batch descriptor") plus one tiny device-to-host copy of the sampled token. No per-step dispatch, no Python in the loop, no graph recapture on shape change — the kernel reads the batch size from the work descriptor and branches on-device.
The mechanism has three load-bearing pieces. First, cooperative launch (cudaLaunchCooperativeKernel): the entire grid must be resident simultaneously, because phases are separated by grid.sync() — a device-wide barrier. That constrains the grid to what fits on the SMs at once and makes occupancy a correctness issue, not just a performance one. Second, device-side scheduling: CTAs claim tiles of work by atomically incrementing a global counter, the same pattern high-performance GEMMs use internally, extended to the whole model. Third, host signaling without relaunch: the host writes the next work descriptor into host-pinned, device-mapped memory; the persistent grid polls it (or blocks on a device-side semaphore). The round trip the host used to pay per token becomes a cache-coherent write.
This is not hypothetical. The XL-persistent-kernel essay on this site walks through a speculative-decode design where the draft/verify loop, acceptance-rate feedback, and block-size adaptation all run inside one persistent mega-kernel — the host never participates in the accept/reject decision. The general principle transfers: any control decision the host was making per step (which bucket, which graph, whether to continue) can move into the kernel if the state it needs lives in device memory.
Why, then, is everything not a persistent kernel? Three hard problems. Collectives. NCCL is a host-launched API; there are no device-side NCCL collectives, so a tensor-parallel all-reduce inside a persistent kernel needs custom device-side communication or a split design where comms happen in separate host-launched kernels — which reintroduces the launches you were eliminating. Multi-GPU persistent serving is an open research problem, not an engineering task. Debugging. A hung persistent kernel looks exactly like a GPU hang: cuda-gdb on a never-exiting grid, timeouts instead of error codes, and printf-style debugging through device memory. The edit-debug cycle that takes seconds for host code takes tens of minutes. Portability. Grid-sync across 132 SMs on an H100 has different occupancy math than 84 SMs on an A100 or a Blackwell GB200; the "one kernel" has to be re-tuned per architecture, and a grid sized for one GPU may not even launch on another.
The left panel is the conventional decode step: six host-side phases per token, each crossing the CPU–GPU boundary at least once — schedule, launch, synchronize/sample, copy the token back, detokenize, enqueue the next step. The right panel is the persistent alternative: one cooperative launch for the entire session, with the decode loop (work-pull, grid-wide barrier, decode, token write) running on-device and the host reduced to a 4-byte "next batch" write per step. The takeaway is architectural: every per-step host decision is a candidate for device-side migration, but collectives, debuggability, and per-architecture occupancy tuning are the taxes.
The honest placement of persistent kernels on the maturity curve: single-GPU, single-model research systems can use them today; production multi-GPU serving cannot, because the NCCL problem has no clean answer yet. They are the direction the field is walking, not the tool you deploy this quarter. The tool you deploy this quarter is graphs plus compilation — which is why the ordering of this essay runs from the standard fix to the radical one.
05 The scheduler is Python
Below the launches, the control plane itself runs in an interpreter.
Everything above assumed the host at least dispatches quickly. It often does not, because the thing deciding what to dispatch is Python. In vLLM, every decode step runs scheduler.schedule() — scanning the waiting and running queues, applying the scheduling policy (priority, preemption, chunked-prefill budgeting), consulting the block manager for KV-cache space, and assembling the batch descriptor — all in interpreted Python, plus the sampling-parameter preparation and per-request detokenization that follow the step. This is the Tpython term in the budget, and unlike launch overhead it scales with the number of requests: queue scans are O(running + waiting), sampling prep is O(batch), detokenization is O(batch). At batch 128 the scheduler's Python can cost milliseconds — more than the launches it is scheduling.
The direction of travel is to push the control plane down the stack. vLLM v1's EngineCore architecture was the structural precondition: it split the engine into a dedicated core where the scheduler and the model executor communicate over queues, rather than one interleaved Python loop doing both. That separation is what makes a compiled or native scheduler possible — the scheduling step becomes a well-typed function over batch descriptors instead of a tangle of interpreter state. The end state various stacks are converging on: the per-step scheduling decision compiled (torch.compile over the scheduler's tensor logic) or rewritten in C++/Rust, with Python demoted to policy configuration — "use this eviction policy" — rather than per-step execution.
Do not forget the unglamorous tail of the host path: tokenizers and detokenizers. Detokenization runs per token per request on the critical path between the sampled token id and the bytes the user sees; a Python detokenizer doing incremental decode for a hundred concurrent requests is real, measurable latency that no CUDA graph touches. The stacks that take overhead seriously have moved detokenization to native code or batched it off the step path. It is the same lesson as the launches: profile the host, not just the device.
There is also a scheduling trick that requires no scheduler rewrite at all: overlap. The scheduler's work for step t+1 — queue scans, block allocation, batch assembly — does not depend on step t's GPU results (the only results it needs are the finished token ids, a few bytes). So run the scheduler concurrently with GPU execution: while the device replays the graph for step t, the host prepares step t+1. Done right, Tpython hides inside Tgpu and the host tax drops out of the critical path. The constraint is the dependency edge — sampling for step t+1 needs step t's logits — so the overlap covers scheduling and input preparation, not the whole pipeline. It is pipelining, applied to the control plane.
Kernel dispatch — ~5µs per launch, N launches per step. Fixed by graphs (one launch) or compilation (fewer launches).
Scheduler Python — queue scans, block-manager bookkeeping, sampling prep. Grows with batch; fixed by compiling or nativizing the scheduler.
Detokenizer — per-token-per-request Python on the user-visible path. Fixed by native/batched detokenization.
Graph management — bucket selection, recapture stalls, static-buffer copies. Grows with configuration space; fixed by cudagraph trees or persistent kernels.
Overhead optimization has a ceiling, and the ceiling has a name: the memory wall. Graphs, compilers, and native schedulers buy back the gaps between kernels — they cannot make the weights smaller or the bandwidth wider. Do this work after batching is right (batching amortizes everything) and before exotic kernel surgery (which optimizes the part that is already a minority of the step).
06 Open questions
What a device-driven serving stack would actually require.
Can the whole serving loop go device-side? The persistent-kernel work shows the decode loop can. The harder parts are admission (new requests arriving mid-session need their KV pages allocated — the block manager lives on the host), preemption (evicting a sequence means rewriting page tables the kernel is actively reading), and multi-tenancy (fair sharing of SMs between models). Each is a host-side control decision that would need a device-side equivalent. Nobody has demonstrated all three in one system.
Do graphs compose with disaggregated serving? In split prefill/decode architectures (Splitwise, DistServe, Mooncake-style), the KV cache moves between machines mid-request — an NCCL or RDMA transfer whose shape varies per request. CUDA graphs can capture NCCL work (PyTorch's CUDAGraph Trees do per-device recording with NCCL support), but the transfer sizes are request-dependent, which fights the static-shape requirement exactly where the system's value proposition lives. Whether the transfer control plane can be captured, or must stay host-driven, is unresolved.
Does Python ever leave the hot loop? The economic argument says yes: at some batch size the interpreter tax exceeds the launch tax it was hired to schedule around, and the industry keeps rediscovering this (TensorRT-LLM's C++ executor never had the problem; it also never had vLLM's iteration velocity). The likely equilibrium is a split — Python for policy, compiled code for mechanism — because the scheduler's decisions change weekly while the dispatch never does. Watch for the scheduler compile boundary: whichever stack first ships a compiled per-step scheduler with vLLM-class policy richness wins the next round of this fight.
The meta-question: what does the overhead ratio look like on the next two GPU generations? If launch latency stays ~5µs while bandwidth doubles again, the (N·L)/Tgpu ratio for interactive serving keeps climbing, and techniques that look exotic today — persistent kernels, device-side scheduling — become the obviously correct engineering. The overhead wall is not a fixed obstacle; it is a moving one, and it moves toward the host.