GPU Warp Divergence is the performance penalty that occurs when threads within the same warp (NVIDIA, 32 threads) or wavefront (AMD, 64 threads) take different execution paths at a branch instruction — forcing the SIMT processor to serialize the divergent paths by executing each branch sequentially while masking inactive threads, potentially halving or worse the effective throughput of divergent code sections.

How SIMT Execution Creates Divergence

GPU hardware executes one instruction across all threads in a warp simultaneously. When a conditional branch is encountered:

• If ALL threads take the same path: no penalty, full throughput.
• If SOME threads take the if-path and others the else-path: the hardware first executes the if-path with else-threads masked (inactive), then executes the else-path with if-threads masked. Both paths execute sequentially — the cost is the SUM of both paths, not the MAX.

Divergence Impact

// High divergence — every other thread takes a different path
if (threadIdx.x % 2 == 0) {
    path_A();  // 16 threads active, 16 masked
} else {
    path_B();  // 16 threads active, 16 masked
}
// Effective utilization: 50% (both paths execute sequentially)

// No divergence — all threads in a warp take the same path
if (threadIdx.x / 32 == some_condition) {
    path_A();  // entire warp goes one way
} else {
    path_B();  // different warp goes other way
}
// Effective utilization: 100%

Mitigation Strategies

• Data Reorganization: Sort or bin data so that threads within a warp process similar work (e.g., particles of the same type, pixels in the same region). Coherent data produces coherent branches.
• Thread Reassignment: Instead of assigning thread-to-data statically, use a work queue where each warp pulls homogeneous work items.
• Predication: For short divergent code (a few instructions), compilers replace branches with predicated execution — both paths compute, and a select instruction picks the correct result. Eliminates the branch entirely at the cost of executing redundant instructions.
• Warp Specialization: Assign different warps to different code paths rather than letting a single warp encounter the branch. More warps but each runs at full efficiency.

Nested Divergence

Nested branches compound the problem: a two-level nested if-else can reduce utilization to 25% (4 serial paths with 8 active threads each in a 32-thread warp). Deeply branching code (recursive tree traversal, interpreters) causes severe divergence and should be restructured or moved to the CPU.

Measurement

NVIDIA Nsight Compute reports "warp execution efficiency" — the ratio of active threads to total threads across all executed instructions. Values below 80% indicate significant divergence worth optimizing.

GPU Warp Divergence is the fundamental tension between the GPU's SIMT execution model and data-dependent control flow — the performance cliff that programmers must understand and design around to achieve the throughput that makes GPU computing worthwhile.