GPU Kernel Launch Overhead
GPU Kernel Launch Overhead is the fixed latency cost (typically 3-10 microseconds) incurred each time the CPU dispatches a computation kernel to the GPU — which becomes a significant performance bottleneck when an application launches thousands of small kernels per second, as the launch overhead can dominate actual computation time, motivating kernel fusion, CUDA Graphs, and persistent kernel techniques to amortize or eliminate this per-launch cost.
Kernel Launch Pipeline
1. CPU prepares kernel arguments and grid configuration.
2. CPU writes launch command to GPU command buffer (driver overhead).
3. Command is submitted to GPU command processor.
4. GPU command processor decodes and schedules work.
5. GPU SMs begin executing threads.
- Steps 1-4: ~3-10 µs of overhead before any GPU thread runs.
- For large kernels (1ms+ runtime): 3-10 µs overhead is negligible.
- For tiny kernels (1-10 µs runtime): Overhead is 50-90% of total time!
Launch Overhead Breakdown
| Component | Typical Latency | Notes |
|---|---|---|
| Driver API call | 1-3 µs | CPU-side driver processing |
| Command buffer write | 0.5-1 µs | PCIe MMIO or host memory |
| GPU command processing | 1-3 µs | Decode, resource allocation |
| SM scheduling | 0.5-2 µs | Warp creation, register allocation |
| Total | 3-10 µs | Per kernel launch |
Impact on ML Workloads
- PyTorch eager mode: Each operation (add, matmul, relu) → separate kernel launch.
- A single transformer layer: ~20-50 kernel launches.
- 32-layer model forward pass: ~600-1600 kernel launches.
- At 5 µs each: 3-8 ms of pure launch overhead → significant for inference.
Mitigation Strategies
| Strategy | How | Overhead Reduction |
|---|---|---|
| Kernel fusion | Combine multiple ops into one kernel | Eliminate intermediate launches |
| CUDA Graphs | Record sequence → replay as single dispatch | Amortize to ~1 µs total |
| Persistent kernels | Kernel stays running, polls for new work | Near-zero per-task overhead |
| torch.compile | Fuse operations at graph level | 50-80% fewer launches |
| TensorRT/TVM | Aggressive pre-compilation fusion | Minimal launches |
CUDA Graphs
// Record sequence of kernels
cudaGraph_t graph;
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel_a<<<grid, block, 0, stream>>>(...);
kernel_b<<<grid, block, 0, stream>>>(...);
kernel_c<<<grid, block, 0, stream>>>(...);
cudaStreamEndCapture(stream, &graph);
// Create executable graph (one-time cost)
cudaGraphExec_t instance;
cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
// Replay entire sequence with single launch (repeated)
cudaGraphLaunch(instance, stream); // ~1µs for entire sequence- CUDA Graphs reduce per-launch overhead by 50-90% for repeated kernel sequences.
- Perfect for: Inference (same operations repeated), training loops with fixed structure.
Kernel Fusion in Practice
# Unfused (3 kernel launches):
y = torch.relu(x @ W + b) # matmul, add, relu = 3 kernels
# Fused (1 kernel launch via torch.compile):
@torch.compile
def fused_linear_relu(x, W, b):
return torch.relu(x @ W + b) # Compiled to single fused kernelGPU kernel launch overhead is the hidden performance tax that makes naive GPU programming inefficient — while individual launches are microseconds, the cumulative cost across thousands of small operations makes kernel fusion and CUDA Graphs essential optimizations for any GPU application that needs to maximize throughput, particularly in ML inference where latency budgets are tight and every microsecond of overhead directly impacts response time.