GPU Warp Divergence
GPU Warp Divergence is the performance degradation that occurs when threads within a warp take different execution paths at a branch — forcing the hardware to serialize both paths by masking inactive threads, effectively halving or worse the warp's throughput for each divergent branch.
Divergence Mechanics:
- SIMT Execution Model: all 32 threads in a warp execute the same instruction simultaneously; when a conditional branch evaluates differently across threads, the warp must execute both taken and not-taken paths sequentially
- Active Mask: hardware maintains a bitmask indicating which threads are active for the current instruction; inactive threads execute the instruction but their results are discarded (no register writeback, no memory store)
- Reconvergence Point: after both paths complete, the warp reconverges and resumes full-width execution; the compiler inserts synchronization stack entries to track reconvergence points
- Nested Divergence: divergence within an already-divergent path creates further serialization; worst case is 32 unique paths executed sequentially — reducing warp throughput to 1/32
Common Divergence Patterns:
- Thread-ID Conditional: if(threadIdx.x < N) creates divergence within warps where some threads satisfy the condition and others don't; only the boundary warp(s) actually diverge — warps entirely within or outside the range execute without penalty
- Data-Dependent Branching: if(data[tid] > threshold) evaluates differently based on input data; highly irregular data causes severe divergence; sorted or clustered data reduces divergence within warps
- Loop Divergence: while(data[tid]) where each thread iterates a different number of times; the warp continues until the last thread finishes — threads that exit early waste cycles waiting
- Switch Statements: multi-way branches where different threads take different cases; N unique paths selected requires N serial executions of the warp
Mitigation Strategies:
- Data Reorganization: sorting data so adjacent threads process similar values reduces data-dependent divergence; worth the sorting overhead for kernels with many divergent branches
- Predication: the compiler converts short branches (few instructions) into predicated execution — both paths execute but results are conditionally committed; eliminates branch divergence overhead for branches shorter than the predication threshold (~7 instructions on modern architectures)
- Warp-Level Voting: __any_sync/__all_sync allow warps to collectively evaluate conditions before branching — if all threads agree, no divergence occurs; the fast path avoids the branch entirely
- Thread Coarsening: assigning multiple work items per thread and processing them in a loop can convert inter-thread divergence into intra-thread sequential execution — trades parallelism for reduced divergence
- Algorithm Redesign: replacing conditional logic with arithmetic (branchless code) eliminates divergence entirely; example: min/max using conditional assignment instead of if-else branches
Measurement and Analysis:
- Branch Efficiency Metric: Nsight Compute reports branch efficiency as (executed_instructions / (executed_instructions + replay_instructions)) — values below 90% indicate significant divergence
- Active Thread Occupancy: profilers show average active threads per warp per instruction — ideal is 32; divergent code shows averages below the warp width
- Instruction Replay: divergent warps replay instructions for each path; profiled as instruction replay overhead — high replay ratios indicate divergence as the primary performance bottleneck
GPU warp divergence is a fundamental SIMT execution constraint that requires parallel programmers to think in terms of warp-uniform control flow — in well-optimized GPU code, divergent branches are either eliminated through branchless techniques, minimized through data reorganization, or confined to boundary warps where their impact is negligible.