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.