gpu warp divergence
**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.