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

__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<<<num_blocks, 256>>>(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

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

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.