simt execution model divergence

**SIMT Execution and Warp Divergence** characterizes **the single-instruction-multiple-thread execution model where all threads in a warp must execute same instruction, forcing serialized computation of divergent control flow and enabling fine-grained synchronization via warp voting functions.** **SIMT Execution Model Fundamentals** - **Warp Definition**: 32 threads executing in lockstep (Ampere, Hopper). All threads execute same instruction simultaneously (same program counter). - **Program Counter Synchronicity**: All threads in warp share PC. Branches create divergence; some threads take branch, others don't. - **Instruction Level Parallelism (ILP)**: Warp issues 1-4 instructions per cycle (depending on available execution units, latency). Dual-issue allows concurrent FP32 + memory operations. - **SIMT vs SIMD**: SIMT scalar (each thread has scalar registers), SIMD vector (threads share vector registers). SIMT simpler programming model. **Warp Divergence at Branch Points** - **Branch Condition**: if (thread_id < 16) {...}. Some threads take branch, others skip. - **Divergence Impact**: Warp serializes: execute if-branch code with active threads masking (inactive threads stall). Then execute else-branch for alternate threads. - **Serial Execution**: Both branches executed sequentially (not parallel). Effective throughput halved if 50/50 branch distribution (worst case). - **Convergence Stack**: Hardware maintains predication masks tracking which threads active. Stack-based mechanism (IPDOM tree) manages nesting. **Predicated Execution** - **Predicate Register**: Boolean flag per thread (32-bit register with predicate bits). Instruction conditional on predicate (@p0 instruction executes if p0 true for thread). - **Predication Implementation**: All instructions in branch executed, but predicate masks results. Inactive threads produce side effects (state unchanged). - **Branch Elimination**: Small if-else blocks predicated (no explicit branch). Reduces branch misprediction penalty, enables better ILP. - **Predicate Overhead**: Extra instruction (set predicate), + masked instruction execution (no branch, but no result storage). Faster than explicit branch if block small (<4 instructions). **Branch Reconvergence via IPDOM Stack** - **Instruction Level Dominance (IPDOM)**: Reverse dominance in CFG (control flow graph). IPDOM identifies post-dominating blocks (executed after all branches reconverge). - **Reconvergence Point**: IPDOM target = block where all branches from divergence point rejoin. All threads active again. - **Stack Mechanism**: Upon branch, hardware pushes divergence info (predicate masks, target) on stack. Upon reaching reconvergence, pops stack. - **Nesting Complexity**: Nested divergence (if within if) creates stack depth > 1. Deep nesting (>8 levels) possible but rare. **Warp Voting Functions** - **__ballot_sync(mask, predicate)**: Ballot across warp. Returns 32-bit integer with bit i set if thread i's predicate true. Mask specifies participating threads. - **__any_sync(mask, predicate)**: Reduction AND. Returns 1 if any thread's predicate true, else 0 across masked warp. - **__all_sync(mask, predicate)**: Reduction AND. Returns 1 if all threads' predicate true, else 0. - **Use Cases**: ballot() for warp-level histogram; any() for early exit (any thread found solution); all() for synchronization (all threads ready). **Avoiding Divergence via Data-Dependent Branching Analysis** - **Divergence Detection**: Profiler reports "warp stall due to branch" metric. Indicates branch frequency and impact. - **Data-Dependent Patterns**: Analysis of branch conditions determines if thread divergence likely. Example: if (array[tid] > threshold) may have high divergence if array values random. - **Sorting Trick**: For highly-divergent conditionals, sort data by condition value. Clusters threads with same condition together (better branch prediction, less divergence). - **Early Exit**: Loop termination conditions checked via ballot(). Mask inactive threads (data processed), continue active threads. Reduces warp idleness. **Structured vs Unstructured Control Flow** - **Structured Flow**: Single entry/exit loops, if-else blocks. Compiler easily determines reconvergence points. Simple hardware handling. - **Unstructured Flow**: Multiple exits (break, return), goto statements. Complicates reconvergence analysis. Modern GPUs handle but with overhead. - **Best Practice**: Favor structured loops/conditionals. Avoid deep nesting. Minimize branches in hot kernels. **Performance Implications** - **Branch Prediction**: Modern GPUs (Hopper) have branch predictors similar to CPUs. Predicted branches have <5 cycle penalty (vs ~15 cycles misprediction). - **Occupancy Trade-off**: Loop divergence (some threads exit early) may limit occupancy (warps with all threads done freed). Improved throughput overall. - **Warp Efficiency Metric**: Percentage of threads executing useful work. Divergence reduces warp efficiency (inactive threads masked). Target >80% warp efficiency.