SIMT, Warps, and Branch Divergence
REFERENCE CONTEXT · CUDA 12.2; RTX 2070 SUPER; Turing sm_75
Core Concept
CUDA threads execute in 32-lane warps under the SIMT model. When lanes in one warp follow different control-flow paths, the hardware generally executes each path while masking inactive lanes.
Why It Affects Performance
Divergence lowers the number of useful lanes per issued instruction. The goal is not to remove every branch, but to keep lanes in a warp on similar paths when the branch bodies are substantial.
Example or Pseudocode
if ((threadIdx.x & 1) == 0) out[i] = heavy_a(x[i]);
else out[i] = heavy_b(x[i]);Even and odd lanes take different expensive paths, so both paths may execute. Grouping inputs or splitting kernels can help, but introduces reordering and launch costs.
Profiling and Performance Judgment
Inspect Branch Efficiency, WarpStateStats, and source/SASS control flow. Short branches may be predicated, so source code alone does not prove divergence cost.
Five-Minute Exercise
Compare branching on threadIdx.x % 2 with branching on blockIdx.x % 2. Which is more likely to create intra-warp divergence?
Show answer
The thread-index branch splits lanes inside each warp. A block-index branch usually sends every warp in a block down one path, although blocks can still have uneven workloads.