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.

Go deeper with CFSGPT

Get AI-powered deep-dives, save terms, and run advanced simulations — free account.

Create Free Account