persistent threads gpu
**Persistent Threads** is a **GPU programming pattern where a fixed number of threads remain alive for the entire program duration** — repeatedly fetching work from a shared queue rather than being launched and terminated for each work item, reducing kernel launch overhead and enabling dynamic load balancing.
**Traditional GPU Programming**
- For each work batch: Launch kernel → threads process items → kernel exits.
- Kernel launch overhead: ~5–15 μs per launch.
- Problem: Variable-size work items → some thread blocks finish early, GPU underutilized.
**Persistent Thread Pattern**
```cuda
__global__ void persistent_kernel(WorkQueue* queue) {
// Launch exactly: num_SMs * warps_per_SM threads
while (true) {
WorkItem item;
if (!queue->try_pop(&item)) break; // Atomic dequeue
process(item); // Variable-cost work
}
}
// Launch once, process all work
persistent_kernel<<>>(queue);
```
**Work Queue Implementation**
- Global atomic counter: `atomicAdd(&head, 1)` to claim work items.
- Lock-free circular buffer: Multiple producers + multiple consumers.
- Warps fetch work from queue independently — natural load balancing.
**Benefits**
- **Zero launch overhead**: Single kernel launch for all work.
- **Dynamic load balancing**: Fast warps process more items automatically.
- **Producer-consumer**: CPU or other kernels enqueue work while persistent kernel runs.
- **Variable workload**: Handles irregular work (e.g., sparse BFS, ray tracing).
**Challenges**
- **Deadlock risk**: If queue empty and threads waiting — need termination condition.
- **Synchronization**: Work queue access must be atomic — contention at high work rates.
- **Occupancy constraint**: Must launch exactly the right number of threads to maximize occupancy without over-subscribing SMs.
**Use Cases**
- **Ray tracing**: Each ray has variable path length — persistent warps fetch ray tasks.
- **BFS / graph algorithms**: Frontier work queue — variable per-vertex work.
- **Stream processing**: Continuous stream of incoming work items.
Persistent threads are **a powerful pattern for irregular, dynamic GPU workloads** — they trade the simplicity of fixed-size kernel launches for the flexibility needed by graph algorithms, simulation systems, and real-time streaming applications where work size and arrival time cannot be predicted at launch time.