CUDA Graphs

Capture a sequence of kernels/copies once and replay it with one launch, reducing CPU overhead (valuable for many small kernels, e.g., deep learning inference).

Capture a sequence of kernels/copies once and replay it with one launch, reducing CPU overhead (valuable for many small kernels, e.g., deep learning inference).

25.1 Stream capture

cudaGraph_t graph;
cudaGraphExec_t graphExec;

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
for (int k = 0; k < 20; ++k) {
    kernelA<<<grid, block, 0, stream>>>(d_x);
    kernelB<<<grid, block, 0, stream>>>(d_x);
}
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&graphExec, graph, 0);

for (int iter = 0; iter < 1000; ++iter)
    cudaGraphLaunch(graphExec, stream);          // replay the whole sequence
cudaStreamSynchronize(stream);
cudaGraphExecDestroy(graphExec); cudaGraphDestroy(graph);

25.2 Explicit graph construction & updates

cudaGraphCreate, cudaGraphAddKernelNode, cudaGraphAddMemcpyNode, dependencies between nodes; cudaGraphExecKernelNodeSetParams / cudaGraphExecUpdate to change parameters without re-instantiating; conditional nodes (if/while) and device-side graph launch in CUDA 12.x. Frameworks (PyTorch torch.cuda.graphs) expose graphs for inference speedups.

Why graphs are faster

Every kernel launch costs some CPU time to prepare and submit, typically a few microseconds. For large kernels this does not matter, but a model that runs hundreds of small kernels per step can spend more time launching than computing, leaving the GPU idle between kernels. A graph records the whole sequence and its dependencies once; replaying it submits all of it with a single call, and the driver can schedule the work more efficiently because it sees everything in advance.

The trade-off is flexibility: a graph replays exactly what was captured, with the same kernels, launch sizes and pointers. It suits work that repeats in the same shape, such as training or inference steps with fixed batch sizes.

Common mistakes

Expecting a replay to see new inputs placed in a different buffer: the graph still uses the captured addresses. Copy new data into the same buffers, or update the nodes with cudaGraphExecUpdate or cudaGraphExecKernelNodeSetParams.

Calling synchronous or forbidden APIs during capture, such as cudaMalloc, cudaMemcpy (not async) or cudaDeviceSynchronize: capture is invalidated and fails. Allocate before capturing, or use cudaMallocAsync on the captured stream.

Assuming CPU code inside the capture loop runs on each replay: only GPU work submitted to the stream is recorded.

Re-instantiating the graph every iteration, which throws away the savings. Instantiate once and replay many times.