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**.