GPU Warp Divergence is the performance penalty that occurs when threads within the same warp (typically 32 threads executing in lockstep) take different paths at a branch instruction — forcing the GPU to serialize the divergent paths by executing each branch sequentially and masking inactive threads, wasting execution slots and reducing the effective parallelism that is the GPU's fundamental performance advantage.

How SIMT Execution Works

• GPU executes threads in groups called warps (NVIDIA, 32 threads) or wavefronts (AMD, 32/64 threads).
• All threads in a warp execute the SAME instruction at the SAME time (Single Instruction, Multiple Threads).
• No divergence: All 32 threads active → 100% utilization.
• With divergence: Only a subset active per branch → utilization drops.

Divergence Example

if (threadIdx.x < 16) {
    // Path A — threads 0-15 execute, 16-31 idle
    a[threadIdx.x] = compute_A();
} else {
    // Path B — threads 16-31 execute, 0-15 idle
    a[threadIdx.x] = compute_B();
}
// Both paths reconverge here → all 32 threads active again

• Without divergence: 1 pass. With divergence: 2 passes → 50% efficiency.

Cost of Divergence

Sources of Divergence

• Data-dependent branches: if (data[tid] > threshold) — diverges if data varies within warp.
• Thread ID branches: if (tid % 4 == 0) — predictable divergence pattern.
• Loop iteration counts: while (data[tid]) — threads exit loop at different times.
• Switch statements: Multiple paths from single branch → multi-way divergence.

Minimizing Divergence

1. Reorganize data: Sort/partition data so threads in same warp take same path.

• Compact: Move "yes" elements together, "no" elements together → separate warps.

2. Predication over branching: For short branches, compute both paths and select result.

• result = (condition) ? path_A : path_B; — no divergence, both computed.

3. Warp-level primitives: __ballot_sync(), __shfl_sync() — collective operations avoid branches. 4. Algorithm redesign: Replace branching with arithmetic (branchless min/max, bitwise selection).

Reconvergence

• After divergent section, threads must reconverge to resume lockstep execution.
• Stack-based reconvergence (traditional): Hardware push/pop divergence stack.
• Independent Thread Scheduling (Volta+): Each thread has own PC → more flexible but reconvergence still matters for performance.

GPU warp divergence is the single most common source of GPU underutilization — understanding and minimizing divergence through data reorganization, predication, and algorithm design is essential for writing high-performance GPU kernels that achieve the theoretical throughput of the hardware.