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.

Persistent Kernel GPU

Take persistent kernel gpu further