persistent kernel gpu
**Persistent GPU Kernels** is the **programming technique where a single GPU kernel runs continuously for the lifetime of the application (or a large phase of it), consuming work items from a global queue rather than launching a new kernel for each batch of work — eliminating the 5-20 μs kernel launch overhead per invocation and enabling GPU-side scheduling, dynamic work generation, and fine-grained producer-consumer patterns that the traditional launch-per-batch model cannot efficiently support**.
**The Kernel Launch Overhead Problem**
Each GPU kernel launch involves: CPU-side API call, command buffer insertion, GPU command processor dispatch, and resource allocation. Total overhead: 5-20 μs per launch. For workloads with small kernels (50 μs of compute): launch overhead is 10-30% of total time. For iterative algorithms with 1000+ launches: cumulative overhead of 5-20 ms becomes significant.
**Persistent Kernel Architecture**
```
__global__ void persistent_kernel(WorkQueue* queue) {
while (true) {
WorkItem item = queue->dequeue(); // atomic pop
if (item.is_terminate()) return;
process(item); // actual computation
// Optionally: enqueue new work items
}
}
```
Key design elements:
- **Global Work Queue**: Lock-free MPMC (multi-producer multi-consumer) queue in GPU global memory. atomicAdd-based or ring buffer with atomic head/tail pointers.
- **Grid-Level Persistence**: Launch enough thread blocks to fill the GPU (100-200% occupancy). Blocks never exit — they loop, dequeueing and processing work items indefinitely.
- **Dynamic Load Balancing**: Every thread block pulls work from the same queue — naturally load-balanced. No block-level partitioning needed.
- **Termination**: CPU inserts a poison pill / terminate signal. All blocks detect and exit gracefully.
**Use Cases**
- **Graph Algorithms**: BFS, SSSP, PageRank where each iteration generates a variable frontier. Persistent kernel avoids relaunching for each level — 2-5× speedup on small graphs.
- **Ray Tracing**: Persistent wavefront scheduler — each warp processes one ray, pulling new rays from the queue when the current ray terminates.
- **Simulation**: Agent-based models, particle systems where work per step varies. Persistent kernel adapts to variable workload without CPU intervention.
- **Server-Side Inference**: GPU acts as a persistent service processing inference requests from a queue. No per-request kernel launch overhead.
**Challenges**
- **Deadlock Risk**: If the grid cannot fit all required blocks simultaneously, some blocks wait for resources held by sleeping blocks — deadlock. Solution: limit block count to guaranteed concurrent capacity.
- **Starvation**: If the queue is empty, persistent blocks spin-wait — wasting GPU resources. Solution: yield (cooperative groups) or backoff.
- **CUDA Graphs Alternative**: For fixed computation patterns, CUDA Graphs provide launch overhead reduction without persistent kernel complexity. Persistent kernels are better for dynamic/unpredictable workloads.
Persistent GPU Kernels is **the programming pattern that transforms the GPU from a batch processor to a continuous computing engine** — enabling dynamic, data-driven workloads that cannot be efficiently decomposed into fixed-size kernel launches.