CUDA Graphs — Capture and Replay
CUDA Graphs — Capture and Replay
Thesis
CUDA Graphs record GPU work once and replay it with less CPU launch overhead. Useful when shapes are static. No GPU timings on this host.
Mental model
capture → instantiate → launch repeatedly
Worked examples
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
my_kernel<<<grid, block, 0, stream>>>(...);
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&exec, graph, NULL, NULL, 0);
cudaGraphLaunch(exec, stream);
PyTorch: torch.cuda.make_graphed_callables when shapes are fixed.
Trap
Capturing across host syncs or with dynamic shapes. Graphs fail or help little.
Boring rule
Use graphs only after a profiler shows launch overhead and shapes are stable.
Try this
- Nsight Systems graphed vs ungated loops on a GPU box.
- Read CUDA Programming Guide — Graphs.
- Check current vLLM/TensorRT graph toggles.
Sources
- CUDA C++ Programming Guide — Graphs