GPU Glossary
/
Host Software·2 min read

CUDA Graph

A graph of kernel launches and other GPU work that the host submits to the device in a single step.

A CUDA Graph is a graph of kernel launches and other work that the host submits to the device all at once, instead of one launch at a time.

Why use them#

CUDA Graphs mainly cut overhead: the host's cost of identifying, configuring and submitting many kernels in a short time. Each launch takes on the order of microseconds. When hundreds of kernels must launch within milliseconds, that overhead becomes very noticeable. This is common in low-latency LLM inference.

Key idea Launching a whole graph at once replaces hundreds of microsecond-scale launch costs with one submission.

Creating a graph#

Graphs are most often created with the stream capture API in the CUDA Runtime. It records every operation issued on a single CUDA stream so it can be replayed later:

// capture
cudaStreamBeginCapture(stream);
kernelGemm<<<{32, 20},64,19200,stream>>>(a, b, c);
kernelEpilogue<<<{256,2},{8,32},0,stream>>>(c, c);
cudaStreamEndCapture(stream, &graph);

// launch
cudaGraphInstantiate(&graphExec, graph, flags);
cudaGraphLaunch(graphExec, stream);

NVIDIA documents the Runtime interface to CUDA Graphs here.

PyTorch wraps this API, for example in the torch.cuda.graph context manager. That is how CUDA Graphs are usually captured for neural network training and inference.

Example#

Below is a CUDA Graph captured on a B200 GPU running a torch.Linear layer:

┌─────────────────────────────────────────────────────────────────────────┐
│  ┌───────────────────────────────────────────────────────────────────┐  │
│  │                          NODE 0: KERNEL                           │  │
│  ├───────────────────────────────────────────────────────────────────┤  │
│  │  ID:         0 (topoId: 1)                                        │  │
│  │  Kernel:     cutlass3x_sm100_simt_sgemm_f32_f32_f32_f32_f32_      │  │
│  │              64x32x16_1x1x1_3_tnn_align1_bias_f32_relu            │  │
│  │              <<<{32,20},64,19200>>>                               │  │
│  │  Node handle: 0x0000564604539520                                  │  │
│  │  Func handle: 0x0000564603AFCC00                                  │  │
│  └───────────────────────────────────────────────────────────────────┘  │
│                              │                                          │
│                              │                                          │
│                              ▼                                          │
│  ┌───────────────────────────────────────────────────────────────────┐  │
│  │                          NODE 1: KERNEL                           │  │
│  ├───────────────────────────────────────────────────────────────────┤  │
│  │  ID:         1 (topoId: 0)                                        │  │
│  │  Kernel:     _ZN8cublasLt8epilogue4impl12globalKernelILi8E...     │  │
│  │              <<<{256,2},{8,32},0>>>                               │  │
│  │  Node handle: 0x0000564604539C88                                  │  │
│  │  Func handle: 0x00005646044770F0                                  │  │
│  └───────────────────────────────────────────────────────────────────┘  │
└─────────────────────────────────────────────────────────────────────────┘

Limitation: not portable#

Kernels in a graph are identified by pointers, such as 0x564603AFCC00, and so are inputs and outputs. These and other references to device resources mean a CUDA Graph can't be serialized or moved elsewhere, short of fully checkpointing and restoring host and device memory.