persistent kernel
**Persistent Kernels** are the **GPU programming technique where a kernel is launched once and runs indefinitely, continuously polling for new work from a shared queue rather than being launched and terminated for each task** — eliminating the repeated kernel launch overhead by keeping GPU threads alive and ready, achieving sub-microsecond task dispatch latency compared to the 3-10 µs of standard kernel launches, critical for workloads with many small tasks like graph processing, dynamic neural network execution, and real-time systems.
**Standard vs. Persistent Kernel Model**
```
Standard model:
CPU: launch(task1) → wait → launch(task2) → wait → launch(task3)
GPU: [idle][task1][idle][task2][idle][task3]
Overhead: 3-10 µs per launch
Persistent model:
CPU: push(task1) → push(task2) → push(task3) (to GPU-visible queue)
GPU: [persistent kernel: poll → task1 → poll → task2 → poll → task3]
Overhead: ~100ns per task (queue polling)
```
**Implementation Pattern**
```cuda
__global__ void persistent_kernel(TaskQueue *queue, Result *results) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
while (true) {
// Poll for work
Task task;
if (tid == 0) {
// Block leader atomically dequeues task
task = atomicDequeue(queue);
if (task.type == TERMINATE) break;
}
// Broadcast task to all threads in block
task = __shfl_sync(0xFFFFFFFF, task, 0);
// Process task cooperatively
process(task, results, tid);
// Signal completion
if (tid == 0) atomicIncrement(&task.done_flag);
}
}
// Launch once, runs forever
persistent_kernel<<>>(d_queue, d_results);
// CPU feeds work by writing to queue
submit_task(h_queue, new_task); // GPU picks up in ~100ns
```
**Benefits and Costs**
| Aspect | Standard Kernels | Persistent Kernels |
|--------|-----------------|-------------------|
| Launch overhead | 3-10 µs per kernel | ~0 (launched once) |
| Task dispatch | µs level | ~100 ns |
| GPU utilization | Variable (idle between launches) | Continuous |
| Dynamic work | New kernel per shape/size | Same kernel handles all |
| Resource occupancy | Released between launches | Held permanently |
| Programming complexity | Simple | High |
**Challenges**
- **Occupancy starvation**: Persistent kernel occupies SMs → other kernels can't run.
- **Deadlock risk**: All blocks waiting on queue → no blocks available for dependent tasks.
- **Power consumption**: Polling threads consume power even when idle.
- **Debugging**: Long-running kernels are harder to debug and profile.
**Use Cases**
| Application | Why Persistent | Benefit |
|------------|---------------|--------|
| Graph processing | Irregular, many small tasks | 5-10× throughput |
| Dynamic neural networks | Variable computation per sample | Sub-ms dispatch |
| Real-time inference | Latency-critical, steady stream | Minimal tail latency |
| Task graph execution | Fine-grained dependencies | Avoid launch per task |
| Ray tracing | Dynamic workload distribution | Better load balancing |
**Modern Alternatives**
- **CUDA Graphs**: Pre-record kernel sequence → replay as batch (less flexible but simpler).
- **CUDA Dynamic Parallelism**: Kernels launch other kernels (limited to 24 levels).
- **Cooperative Groups + Grid Sync**: All blocks coordinate without persistent model.
Persistent kernels are **the advanced GPU programming technique for absolute minimum dispatch latency** — while CUDA Graphs handle the common case of repeated fixed sequences, persistent kernels provide the ultimate flexibility for workloads where task structure is dynamic and unpredictable, enabling GPU programming models that more closely resemble event-driven systems than the traditional bulk-synchronous launch-and-wait paradigm.