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.