CUDA Graph

CUDA Graphs are a mechanism to capture a sequence of GPU operations as a graph and replay them with minimal CPU overhead — eliminating the per-kernel launch overhead that limits performance for workloads with many small GPU operations.

The Problem CUDA Graphs Solve

Each CUDA kernel launch: ~5–15 μs CPU overhead for driver processing.
DNN inference: 100–1000 small kernels per inference step.
At 100 kernels × 10 μs = 1ms overhead per inference → unacceptable for latency-sensitive applications.

CUDA Graph Concepts

Graph Node: A GPU operation (kernel, memcpy, memset, event).
Edge: Dependency between nodes.
Graph Instantiation: Compile graph to an executable graph.
Graph Launch: Execute instantiated graph — single CPU call for all operations.

Creating a CUDA Graph

cudaGraph_t graph;
cudaGraphExec_t instance;

// Method 1: Stream Capture
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel_A<<<grid, block, 0, stream>>>();
cudaMemcpyAsync(dst, src, size, kind, stream);
kernel_B<<<grid, block, 0, stream>>>();
cudaStreamEndCapture(stream, &graph);

// Instantiate and launch repeatedly
cudaGraphInstantiate(&instance, graph, 0);
for (int iter = 0; iter < N; iter++)
    cudaGraphLaunch(instance, stream);

Performance Benefits

Eliminates per-kernel CPU launch overhead.
Pre-optimizes execution order and dependency resolution.
Memory transfer scheduling: Optimal DMA sequencing determined once.
Typical speedup: 5–30% for inference workloads (more for smaller models).

Operator Fusion

Separate kernels: Multiple global memory round-trips (read-compute-write-read-compute-write).
Fused kernel: Single kernel processes multiple operators → data stays in registers/shared memory.
Manual fusion: Combine elementwise ops into single kernel.
Compiler fusion: XLA (TensorFlow), TorchInductor (PyTorch), TVM automate fusion.
FlashAttention: Fuses QK^T matmul + softmax + V matmul → 4x memory bandwidth reduction.

CUDA Graphs and operator fusion are the key to closing the gap between raw GPU compute and actual inference throughput — at batch size 1, these optimizations are often the difference between 1ms and 5ms latency, directly determining real-time applicability of AI applications.

CUDA Graph

Take cuda graph further