CUDA Graphs is a CUDA runtime feature that records a sequence of GPU kernel launches and memory operations into a reusable graph object, then replays the entire sequence with a single API call. The key insight is that in traditional eager execution each kernel launch requires a round-trip from CPU to GPU driver — typically 5-10 microseconds of overhead — and that overhead compounds across thousands of small kernels per training step or inference pass. CUDA Graphs amortize all of that to a single launch call regardless of how many nodes the graph contains.

<svg viewBox="0 0 760 470" xmlns="http://www.w3.org/2000/svg" font-family="-apple-system,Segoe UI,Roboto,sans-serif">
  <rect width="760" height="470" fill="#0d1117"/>

  <!-- Title -->
  <text x="380" y="22" fill="#e6edf3" font-size="13" font-weight="700" text-anchor="middle">CUDA Graphs</text>
  <text x="380" y="36" fill="#8b98a5" font-size="10" text-anchor="middle">record kernel launch sequences once, replay with a single API call — eliminating per-kernel CPU launch overhead</text>

  <!-- ── PANEL 1: Eager vs Graph timeline ── -->
  <rect x="8" y="44" width="236" height="210" rx="6" fill="#0e1320" stroke="#334155" stroke-width="1"/>
  <text x="126" y="58" fill="#f87171" font-size="10" font-weight="700" text-anchor="middle">Eager vs Graph Launch</text>

  <!-- Eager mode -->
  <text x="18" y="74" fill="#94a3b8" font-size="9" font-weight="600">Eager (standard):</text>
  <!-- timeline bar -->
  <rect x="18" y="78" width="216" height="6" rx="2" fill="#1e293b"/>
  <!-- K1 -->
  <rect x="18" y="76" width="32" height="10" rx="2" fill="#1d4ed8" stroke="#60a5fa" stroke-width="0.8"/>
  <text x="34" y="84" fill="#bfdbfe" font-size="7.5" text-anchor="middle">K1</text>
  <!-- overhead 1 -->
  <rect x="50" y="76" width="18" height="10" rx="1" fill="#7f1d1d" stroke="#f87171" stroke-width="0.6"/>
  <text x="59" y="84" fill="#fca5a5" font-size="6.5" text-anchor="middle">OH</text>
  <!-- K2 -->
  <rect x="68" y="76" width="32" height="10" rx="2" fill="#1d4ed8" stroke="#60a5fa" stroke-width="0.8"/>
  <text x="84" y="84" fill="#bfdbfe" font-size="7.5" text-anchor="middle">K2</text>
  <!-- overhead 2 -->
  <rect x="100" y="76" width="18" height="10" rx="1" fill="#7f1d1d" stroke="#f87171" stroke-width="0.6"/>
  <text x="109" y="84" fill="#fca5a5" font-size="6.5" text-anchor="middle">OH</text>
  <!-- K3 -->
  <rect x="118" y="76" width="32" height="10" rx="2" fill="#1d4ed8" stroke="#60a5fa" stroke-width="0.8"/>
  <text x="134" y="84" fill="#bfdbfe" font-size="7.5" text-anchor="middle">K3</text>
  <!-- overhead 3 -->
  <rect x="150" y="76" width="18" height="10" rx="1" fill="#7f1d1d" stroke="#f87171" stroke-width="0.6"/>
  <text x="159" y="84" fill="#fca5a5" font-size="6.5" text-anchor="middle">OH</text>
  <!-- K4 -->
  <rect x="168" y="76" width="32" height="10" rx="2" fill="#1d4ed8" stroke="#60a5fa" stroke-width="0.8"/>
  <text x="184" y="84" fill="#bfdbfe" font-size="7.5" text-anchor="middle">K4</text>
  <!-- overhead legend -->
  <rect x="206" y="76" width="18" height="10" rx="1" fill="#7f1d1d" stroke="#f87171" stroke-width="0.6"/>
  <text x="215" y="84" fill="#fca5a5" font-size="6.5" text-anchor="middle">OH</text>

  <text x="18" y="98" fill="#f87171" font-size="8">5-10 µs CPU overhead per kernel launch</text>

  <!-- Graph mode -->
  <text x="18" y="116" fill="#94a3b8" font-size="9" font-weight="600">Graph (after capture):</text>
  <!-- single launch bar -->
  <rect x="18" y="120" width="8" height="10" rx="1" fill="#166534" stroke="#4ade80" stroke-width="0.8"/>
  <text x="22" y="128" fill="#86efac" font-size="5.5" text-anchor="middle">L</text>
  <!-- fused kernel block -->
  <rect x="28" y="120" width="150" height="10" rx="2" fill="#14532d" stroke="#22c55e" stroke-width="1"/>
  <text x="103" y="128" fill="#86efac" font-size="7.5" text-anchor="middle">K1 + K2 + K3 + K4 (replayed on GPU)</text>

  <text x="18" y="140" fill="#4ade80" font-size="8">Single cudaGraphLaunch call — &lt;1 µs CPU overhead</text>

  <!-- speedup callout -->
  <rect x="18" y="148" width="216" height="40" rx="4" fill="#0a1f10" stroke="#166534" stroke-width="0.8"/>
  <text x="126" y="162" fill="#4ade80" font-size="10" font-weight="700" text-anchor="middle">5-10x launch overhead reduction</text>
  <text x="126" y="174" fill="#86efac" font-size="8.5" text-anchor="middle">+20-40% end-to-end throughput for</text>
  <text x="126" y="184" fill="#86efac" font-size="8.5" text-anchor="middle">small-kernel workloads (inference, RL steps)</text>

  <!-- constraint box -->
  <rect x="18" y="196" width="216" height="28" rx="4" fill="#1a1206" stroke="#d97706" stroke-width="0.8"/>
  <text x="126" y="208" fill="#fcd34d" font-size="8.5" font-weight="600" text-anchor="middle">Constraint: static kernel args + shapes</text>
  <text x="126" y="218" fill="#94a3b8" font-size="8" text-anchor="middle">shape change = recapture (or cudaGraphExecUpdate)</text>

  <!-- ── PANEL 2: Stream capture lifecycle ── -->
  <rect x="252" y="44" width="252" height="210" rx="6" fill="#0e1320" stroke="#334155" stroke-width="1"/>
  <text x="378" y="58" fill="#93c5fd" font-size="10" font-weight="700" text-anchor="middle">Stream Capture Lifecycle</text>

  <!-- Step boxes -->
  <!-- Step 1: Begin capture -->
  <rect x="270" y="66" width="216" height="36" rx="5" fill="#172554" stroke="#2563eb" stroke-width="1"/>
  <text x="290" y="80" fill="#93c5fd" font-size="9" font-weight="700">① cudaStreamBeginCapture(stream)</text>
  <text x="290" y="92" fill="#94a3b8" font-size="8">GPU ops on stream are recorded, not executed</text>

  <!-- arrow down -->
  <line x1="378" y1="103" x2="378" y2="113" stroke="#475569" stroke-width="1.5"/>
  <polygon points="374,113 382,113 378,119" fill="#475569"/>

  <!-- Step 2: Record ops -->
  <rect x="270" y="120" width="216" height="36" rx="5" fill="#1a2e05" stroke="#16a34a" stroke-width="1"/>
  <text x="290" y="134" fill="#6ee7b7" font-size="9" font-weight="700">② kernelA&lt;&lt;&lt;...&gt;&gt;&gt;(); kernelB&lt;&lt;&lt;...&gt;&gt;&gt;();</text>
  <text x="290" y="146" fill="#94a3b8" font-size="8">memcpy, sync, any CUDA op — all recorded as nodes</text>

  <!-- arrow down -->
  <line x1="378" y1="157" x2="378" y2="167" stroke="#475569" stroke-width="1.5"/>
  <polygon points="374,167 382,167 378,173" fill="#475569"/>

  <!-- Step 3: End capture + instantiate -->
  <rect x="270" y="174" width="216" height="36" rx="5" fill="#1c1633" stroke="#7c3aed" stroke-width="1"/>
  <text x="290" y="188" fill="#c4b5fd" font-size="9" font-weight="700">③ cudaStreamEndCapture → cudaGraph_t</text>
  <text x="290" y="200" fill="#94a3b8" font-size="8">cudaGraphInstantiate → cudaGraphExec_t</text>

  <!-- arrow down -->
  <line x1="378" y1="211" x2="378" y2="221" stroke="#475569" stroke-width="1.5"/>
  <polygon points="374,221 382,221 378,227" fill="#475569"/>

  <!-- Step 4: Launch -->
  <rect x="270" y="228" width="216" height="22" rx="5" fill="#0a1f10" stroke="#166534" stroke-width="1"/>
  <text x="378" y="243" fill="#4ade80" font-size="9" font-weight="700" text-anchor="middle">④ cudaGraphLaunch(exec, stream)  — repeat N×</text>

  <!-- replay arc -->
  <path d="M 270 239 Q 255 239 255 200 Q 255 185 270 182" stroke="#22c55e" stroke-width="1.2" stroke-dasharray="4,3" fill="none"/>
  <polygon points="266,179 270,185 274,179" fill="#22c55e"/>
  <text x="249" y="213" fill="#4ade80" font-size="7.5" text-anchor="middle" transform="rotate(-90,249,213)">replay</text>

  <!-- ── PANEL 3: Dependency DAG ── -->
  <rect x="512" y="44" width="240" height="210" rx="6" fill="#0e1320" stroke="#334155" stroke-width="1"/>
  <text x="632" y="58" fill="#fcd34d" font-size="10" font-weight="700" text-anchor="middle">Graph as Dependency DAG</text>

  <!-- DAG nodes -->
  <!-- Root: memcpy H2D -->
  <rect x="570" y="66" width="90" height="22" rx="4" fill="#1c1107" stroke="#d97706" stroke-width="1"/>
  <text x="615" y="81" fill="#fcd34d" font-size="8.5" text-anchor="middle">memcpyH2D (weights)</text>

  <!-- arrows to 2 parallel kernels -->
  <line x1="600" y1="89" x2="566" y2="108" stroke="#475569" stroke-width="1"/>
  <polygon points="562,106 568,112 570,104" fill="#475569"/>
  <line x1="630" y1="89" x2="664" y2="108" stroke="#475569" stroke-width="1"/>
  <polygon points="668,106 662,112 660,104" fill="#475569"/>

  <!-- parallel kernel A -->
  <rect x="528" y="110" width="72" height="22" rx="4" fill="#14253d" stroke="#2563eb" stroke-width="1"/>
  <text x="564" y="125" fill="#93c5fd" font-size="8.5" text-anchor="middle">QK matmul</text>

  <!-- parallel kernel B -->
  <rect x="628" y="110" width="72" height="22" rx="4" fill="#14253d" stroke="#2563eb" stroke-width="1"/>
  <text x="664" y="125" fill="#93c5fd" font-size="8.5" text-anchor="middle">V matmul</text>

  <!-- arrows to softmax -->
  <line x1="564" y1="133" x2="600" y2="152" stroke="#475569" stroke-width="1"/>
  <polygon points="596,150 602,156 604,148" fill="#475569"/>
  <line x1="664" y1="133" x2="628" y2="152" stroke="#475569" stroke-width="1"/>
  <polygon points="632,150 626,156 624,148" fill="#475569"/>

  <!-- softmax -->
  <rect x="570" y="154" width="90" height="22" rx="4" fill="#1a1206" stroke="#d97706" stroke-width="1"/>
  <text x="615" y="169" fill="#fcd34d" font-size="8.5" text-anchor="middle">softmax + scale</text>

  <!-- arrow to output matmul -->
  <line x1="615" y1="177" x2="615" y2="196" stroke="#475569" stroke-width="1"/>
  <polygon points="611,196 619,196 615,202" fill="#475569"/>

  <!-- output proj -->
  <rect x="570" y="200" width="90" height="22" rx="4" fill="#0a1f10" stroke="#16a34a" stroke-width="1"/>
  <text x="615" y="215" fill="#6ee7b7" font-size="8.5" text-anchor="middle">output proj + memcpyD2H</text>

  <!-- labels -->
  <text x="524" y="100" fill="#64748b" font-size="7.5">parallel</text>
  <text x="524" y="109" fill="#64748b" font-size="7.5">streams</text>
  <text x="530" y="148" fill="#64748b" font-size="7.5">join</text>

  <!-- ── 3 Cards ── -->
  <!-- Card 1: API surface -->
  <rect x="8" y="262" width="236" height="100" rx="6" fill="#0e1320" stroke="#334155" stroke-width="1"/>
  <text x="126" y="276" fill="#93c5fd" font-size="9.5" font-weight="700" text-anchor="middle">Key API Calls</text>
  <text x="18" y="289" fill="#fcd34d" font-size="8.5" font-weight="600">Capture:</text>
  <text x="18" y="300" fill="#e5e7eb" font-size="8">cudaStreamBeginCapture / EndCapture</text>
  <text x="18" y="311" fill="#fcd34d" font-size="8.5" font-weight="600">Build graph:</text>
  <text x="18" y="322" fill="#e5e7eb" font-size="8">cudaGraphInstantiate → cudaGraphExec_t</text>
  <text x="18" y="333" fill="#fcd34d" font-size="8.5" font-weight="600">Launch:</text>
  <text x="18" y="344" fill="#e5e7eb" font-size="8">cudaGraphLaunch(exec, stream)</text>
  <text x="18" y="355" fill="#6b7280" font-size="8">Update without recapture: cudaGraphExecUpdate()</text>

  <!-- Card 2: torch.cuda.graph -->
  <rect x="252" y="262" width="252" height="100" rx="6" fill="#0e1320" stroke="#334155" stroke-width="1"/>
  <text x="378" y="276" fill="#6ee7b7" font-size="9.5" font-weight="700" text-anchor="middle">PyTorch / Framework Integration</text>
  <text x="262" y="289" fill="#e5e7eb" font-size="8.5">torch.cuda.graph() context manager wraps CUDA Graph API</text>
  <text x="262" y="302" fill="#e5e7eb" font-size="8">g = torch.cuda.CUDAGraph()</text>
  <text x="262" y="313" fill="#e5e7eb" font-size="8">with torch.cuda.graph(g): model(x_static)</text>
  <text x="262" y="324" fill="#e5e7eb" font-size="8">g.replay()  # zero Python overhead per step</text>
  <text x="262" y="337" fill="#fcd34d" font-size="8.5">Requires static input tensors — use copy_ to update values</text>
  <text x="262" y="349" fill="#6b7280" font-size="8">TensorRT, JAX XLA, and MXNet also use graph-level capture</text>

  <!-- Card 3: When to use -->
  <rect x="512" y="262" width="240" height="100" rx="6" fill="#0e1320" stroke="#334155" stroke-width="1"/>
  <text x="632" y="276" fill="#fcd34d" font-size="9.5" font-weight="700" text-anchor="middle">When to Use (and Not)</text>
  <text x="522" y="289" fill="#4ade80" font-size="8.5" font-weight="600">Use when:</text>
  <text x="522" y="300" fill="#e5e7eb" font-size="8">• Inference: fixed batch size + sequence length</text>
  <text x="522" y="311" fill="#e5e7eb" font-size="8">• Training: same graph per step (steady-state loop)</text>
  <text x="522" y="322" fill="#e5e7eb" font-size="8">• RL env steps: identical observation + action shapes</text>
  <text x="522" y="333" fill="#f87171" font-size="8.5" font-weight="600">Avoid when:</text>
  <text x="522" y="344" fill="#e5e7eb" font-size="8">• Dynamic shapes (variable seq len, conditional branches)</text>
  <text x="522" y="355" fill="#e5e7eb" font-size="8">• Kernels that call CPU callbacks or host functions</text>

  <!-- ── Bottom stat bar ── -->
  <rect x="8" y="368" width="744" height="28" rx="5" fill="#0f1922" stroke="#1e3a5f" stroke-width="0.8"/>
  <text x="28" y="384" fill="#4ade80" font-size="9" font-weight="600">5-10 µs → &lt;1 µs</text>
  <text x="28" y="394" fill="#6b7280" font-size="8">launch overhead per kernel</text>
  <line x1="130" y1="370" x2="130" y2="394" stroke="#1e3a5f" stroke-width="0.7"/>
  <text x="145" y="384" fill="#93c5fd" font-size="9" font-weight="600">+20-40% throughput</text>
  <text x="145" y="394" fill="#6b7280" font-size="8">small-kernel workloads</text>
  <line x1="270" y1="370" x2="270" y2="394" stroke="#1e3a5f" stroke-width="0.7"/>
  <text x="285" y="384" fill="#fcd34d" font-size="9" font-weight="600">3 API calls</text>
  <text x="285" y="394" fill="#6b7280" font-size="8">BeginCapture / Instantiate / Launch</text>
  <line x1="415" y1="370" x2="415" y2="394" stroke="#1e3a5f" stroke-width="0.7"/>
  <text x="430" y="384" fill="#f87171" font-size="9" font-weight="600">Static shapes required</text>
  <text x="430" y="394" fill="#6b7280" font-size="8">recapture on shape change</text>
  <line x1="570" y1="370" x2="570" y2="394" stroke="#1e3a5f" stroke-width="0.7"/>
  <text x="585" y="384" fill="#c4b5fd" font-size="9" font-weight="600">cudaGraphExecUpdate</text>
  <text x="585" y="394" fill="#6b7280" font-size="8">tweak args without full recapture</text>

  <!-- Footer -->
  <text x="380" y="420" fill="#4b5563" font-size="8" text-anchor="middle">CUDA Graphs shift kernel scheduling work from runtime CPU to compile-time — every replayed launch is a lookup not a round-trip through the driver</text>
  <text x="380" y="432" fill="#4b5563" font-size="8" text-anchor="middle">Used in NVIDIA TensorRT, torch.compile, vLLM continuous batching, and training frameworks to close the gap between theoretical and realized FLOPS</text>
</svg>

The capture-instantiate-replay lifecycle has three phases. First, the application calls cudaStreamBeginCapture on a CUDA stream, after which every GPU operation submitted to that stream is recorded as a node in an internal graph rather than executed immediately. After all operations are recorded, cudaStreamEndCapture returns a cudaGraph_t handle representing the DAG of nodes and their dependencies. That graph is compiled into an executable form with cudaGraphInstantiate, producing a cudaGraphExec_t object. From that point on, cudaGraphLaunch submits the entire sequence to the GPU in one call.

The dependency graph captures parallelism automatically. When kernels are submitted to different CUDA streams during capture, the graph records them as parallel nodes. Memory copies, kernel launches, event synchronizations, and host function calls all become nodes with typed edges encoding their dependencies. This means the GPU scheduler sees the full DAG at launch time and can overlap independent nodes across compute and memory copy engines simultaneously.

PyTorch exposes CUDA Graphs through torch.cuda.graph(). The pattern requires static input tensors: allocate placeholder tensors, run a warmup pass to trigger CuBLAS and cuDNN workspace allocation, then capture with g = torch.cuda.CUDAGraph() and the context manager torch.cuda.graph(g). Subsequent forward passes call g.replay() with fresh data copied into the static tensors using tensor.copy_(). This eliminates all Python interpreter and CUDA driver overhead from the steady-state loop, which is why vLLM, TensorRT-LLM, and FlashInfer use graph capture for their inference kernels.

Static shapes are the fundamental constraint. The captured graph hardcodes kernel grid dimensions, tensor addresses, and argument values at capture time. A change in batch size or sequence length invalidates the graph and requires recapture. cudaGraphExecUpdate provides a lighter path — it replaces node parameters without recompiling the graph topology — but it cannot change kernel grid dimensions. Production inference servers typically maintain a pool of graphs captured at common batch sizes (1, 2, 4, 8, 16, ...) and select the smallest covering graph for each incoming request.

Conditional graphs and device graphs (CUDA 12.4+) extend the model to dynamic control flow. Device graphs can be launched from within a running kernel, enabling recursive or iterative patterns that previously required CPU synchronization. Conditional nodes (if/while) allow the graph to branch based on GPU-side predicates without returning to the CPU.

When CUDA Graphs are not appropriate: workloads with variable-length inputs that change every step, operations that must trigger host-side callbacks, workflows with significant Python-side control flow between GPU calls that cannot be captured, and any case where the graph topology itself must change dynamically. For those cases, torch.compile with inductor can still extract fusion and overlap benefits without requiring strict static shapes.