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**
```cuda
cudaGraph_t graph;
cudaGraphExec_t instance;
// Method 1: Stream Capture
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel_A<<>>();
cudaMemcpyAsync(dst, src, size, kind, stream);
kernel_B<<>>();
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.