CUDA Graphs: What They Are and How LLM Inference Uses Them (2026)

What Is a CUDA Graph?

A CUDA Graph is a recorded sequence of GPU operations (kernels, memory copies, and the dependencies between them) that the CPU submits to the GPU as one unit. You capture the work once, instantiate it, then replay it as many times as you need. Each replay skips the per-kernel launch cost the CPU would otherwise pay. NVIDIA's own V100 example shows the effect: a kernel that runs for 2.9µs cost 9.6µs per step with a synchronize after every launch, 3.8µs with overlapped launches, and 3.4µs as a captured graph. LLM inference is the clearest modern case. Decode produces one token per sequence per step and runs hundreds of small kernels, so launch overhead becomes a real share of step time. Serving engines such as vLLM capture graphs for a fixed set of batch sizes at startup and replay the matching one on every step.

Launch Overhead: The Numbers

NVIDIA's introductory CUDA Graphs post (CUDA 10.1, Tesla V100) runs one short kernel 20 times per step, then repeats the step. The kernel itself takes 2.9µs. What changes between the three versions is only how the CPU submits it.

The graph wins because instantiation does the setup work once (validating the launch parameters, building the dependency structure) and each cudaGraphLaunch submits the whole batch of work in one call. The longer the kernels run, the smaller this saving is as a share of total time. A 2ms matrix multiply gains almost nothing. A decode step made of hundreds of 10µs kernels gains a lot.

How CUDA Graph Capture Works

The common path is stream capture. Start capture on a stream, issue your kernels and copies as usual, end capture, and CUDA hands you a graph of what was issued instead of running it.

cudaGraph_t graph;
cudaGraphExec_t instance;

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
for (int k = 0; k < NKERNEL; ++k) {
  shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d);
}
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);

for (int step = 0; step < NSTEP; ++step) {
  cudaGraphLaunch(instance, stream);
  cudaStreamSynchronize(stream);
}

Three rules follow from what a graph is. The kernels, their launch configurations, and the memory addresses they touch are frozen at capture, so a replay reads and writes the same buffers every time; you copy new inputs into those buffers before each launch. Anything the CPU decides at runtime (a Python branch, a host-side allocation, a device-to-host copy that the CPU waits on) cannot happen inside the captured region. And a new input shape means a different graph. PyTorch exposes the same mechanism through torch.cuda.CUDAGraph and the torch.cuda.graph context manager, with static input and output tensors standing in for the frozen buffers.

Why LLM Decode Needs CUDA Graphs

Prefill processes the whole prompt in one forward pass with large matrix multiplies, so its kernels are long and launch cost disappears into them. Decode is the opposite. Each step generates one token per sequence, every layer runs its normalization, projections, attention, and MLP kernels on a handful of rows, and a 60-layer model can issue several hundred kernels per token. At batch size 1 to 32 many of those kernels finish in microseconds, and eager execution leaves the GPU idle while the CPU launches the next one.

Capturing the decode step removes that gap. Engines capture one graph per batch size in a configured list, pad each live batch up to the nearest captured size, and replay it. This is also why decode-heavy features interact with graphs: speculative decoding verifies 1 + k tokens per sequence per step, so the engine needs graphs for that query length too, and KV cache layouts must keep stable addresses for the attention kernels inside the graph.

CUDA Graphs in vLLM: cudagraph_mode

vLLM v1 exposes one setting, CompilationConfig.cudagraph_mode, and treats graph capture separately from torch.compile. It calls a batch of pure decode (query length 1, or 1 + the number of speculative tokens) a uniform decode batch, and anything with prefill a non-uniform batch.

Backend support decides which mode you can use. The vLLM design doc notes that only FlashAttention 3 supported unified full capture at the time of writing, while FlashInfer, FlashMLA, and Mamba support graphs for pure decode batches only. For engine-level differences, see SGLang vs vLLM and nano-vLLM, which implements graph capture in a few hundred lines.

CUDA Graph Limitations

• Static shapes. One graph per shape. Engines capture a list of batch sizes and pad, which spends compute on padding rows.
• Startup cost. Capturing dozens of batch sizes adds time and GPU memory before the server takes traffic.
• Frozen addresses. Inputs must be copied into the captured buffers, and allocations inside the region need the capture-aware allocator.
• No host logic inside. Data-dependent control flow and CPU synchronization break capture, which is why vLLM keeps some ops eager in PIECEWISE mode.

FAQ