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

__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<<<num_SMs, WARPS_PER_SM * 32>>>(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.

Persistent Threads GPU

Take persistent threads gpu further