Warp divergence penalties

Overview

When threads in a warp take different execution paths through conditional code, the GPU must serialize those paths, destroying the parallel efficiency that makes GPUs powerful. This is warp divergence, and it can reduce effective throughput to 1/32nd of peak on NVIDIA GPUs.

GPUs use SIMT (Single Instruction Multiple Thread): one instruction decoder broadcasts to32 execution units. When threads diverge, the hardware must execute each branch serially, masking off inactive threads.

The Mechanism

How the Hardware Handles Divergence

Step-by-step execution model:

1. All threads reach a branch together (same instruction pointer)
2. Each thread evaluates condition → produces true/false
3. Hardware creates execution mask: bitmask marking which threads go way
4. Execute first path with mask A active (other threads idle, burn power)
5. Push second path onto divergence stack
6. Execute second path with mask B active (first group now idle)
7. Reconverge when both paths meet at a common point

Where $T_i$ is the execution time of path $i$.

Why this formula? Without divergence, all32 threads execute in parallel: $T_{\text{ideal}} = \max(T_1, T_2, \ldots, T_k)$. With divergence, we serialize: each path runs sequentially. The speedup loss is:

Derivation for worst-case (each thread different):

If every thread takes a unique path through a conditional with $n$ instructions:

• Ideal: 1 warp executes $n$ instructions in time $n \cdot t_{\text{cycle}}$
• Diverged: 32 sequential executions, each $n$ cycles → $32n \cdot t_{\text{cycle}}$

Throughput reduction: $32\times$ slower

__global__ void divergent_kernel(int *data, int n) {
    int idx = threadIdx.x;
    if (idx < 16) {
        data[idx] = data[idx] * 2;      // Path A: 10 instructions
    } else {
        data[idx] = data[idx] + 100;    // Path B: 8 instructions  
    }
}

What happens inside one warp (threads 0-31):

1. All 32 threads evaluate idx < 16
2. Threads 0-15 → true, threads 16-31 → false
3. Execute path A (multiply): threads 0-15 active, 16-31 masked off (idle) → 10 instruction cycles
4. Execute path B (add): threads 16-31 active, 0-15 masked off → 8 instruction cycles
5. Reconverge at kernel exit

Total cost: 10 + 8 = 18 instruction cycles

Without divergence (if all took same path): max(10, 8) = 10 cycles

Efficiency: 10/18 = 55.6%

Why this step? The hardware can't skip instructions for inactive threads in SIMT—it must decode and issue them, just with masked outputs. This is the serialization penalty.

__global__ void data_dependent(int *types, float *data) {
    int idx = threadIdx.x;
    int type = types[idx];  // Each thread loads different value
    
    if (type == 0) {
        data[idx] = sqrt(data[idx]);        // 20 cycles
    } else if (type == 1) {
        data[idx] = exp(data[idx]);         // 30 cycles
    } else if (type == 2) {
        data[idx] = log(data[idx]);         // 25 cycles
    } else {
        data[idx] = data[idx] * data[idx];  // 5 cycles
    }
}

Scenario: In one warp, 8 threads have type=0, 10 have type=1, 8 have type=2, 6 have type=3.

Execution:

1. Path type=0: sqrt, 20 cycles,8 threads active
2. Path type=1: exp, 30 cycles, 10 threads active
3. Path type=2: log, 25 cycles, 8 threads active
4. Path type=3: multiply, 5 cycles, 6 threads active

Total: 20 + 30 + 25 + 5 = 80 cycles

Ideal parallel: max(20,30,25,5) = 30 cycles

Efficiency: 30/80 = 37.5%

Why this matters: Data-dependent branching where each thread might take a different path is catastrophic for GPU efficiency. This is common in ray tracing (different materials), particle systems (different particle types), or sparse matrix operations.

__global__ void loop_divergent(int *counts, float *data) {
    int idx = threadIdx.x;
    for (int i = 0; i < counts[idx]; i++) {  // Each thread different count
        data[idx] += i * 0.5f;
    }
}

If counts[] = [5, 20, 3, 15 ...] across a warp:

• Thread with count=3 finishes first, goes idle
• Thread with count=20 keeps entire warp busy until it finishes
• Effective iterations: $\max(\text{counts}) = 20$
• Utilization: $\frac{\sum \text{counts}}{32 \cdot \max(\text{counts})}$

Why? The loop exit is a branch. On each iteration, some threads exit (condition false) while others continue. Hardware must keep executing iterations until the last thread finishes, with a shrinking active mask.

Common Patterns and Mitigation

Why it feels right: You think, "Only2 out of 32 threads take the else branch, so 93% are doing the same thing—shouldn't that be okay?"

Why it's wrong: The hardware executes both branches fully for the entire warp, regardless of how many threads take each path. Even if only 1 thread takes a different path, you pay the full cost of both paths.

Example:

if (idx == 0) {
    // 100 instructions - only thread 0
} else {
    // 50 instructions - threads 1-31
}

Cost: 100 + 50 = 150 cycles, not 50 + (100/32) = 53 cycles.

The fix: Understand that divergence is per-warp, and cost is sum of all paths, not weighted by thread count. Restructure to avoid branching altogether, or ensure divergence happens at warp boundaries.

Why it feels right: You think nested conditionals create more divergence points.

Why it's nuanced: Nested ifs within the same divergent path don't add extra serialization—they're part of one path's instruction count. But nested ifs that create orthogonal divergence patterns multiply the number of execution paths.

Example:

if (A) {
    if (B) { path1; } else { path2; }
} else {
    if (C) { path3; } else { path4; }
}

Worst case: 4 distinct paths if distribution is bad. Cost: sum of all path lengths.

The fix: Think in terms of "distinct execution paths" across the warp, not syntactic nesting depth.

Mitigation Strategies

1. Align divergence with warp boundaries

// Bad: divergence within warp
if (idx % 2 == 0) { ... }
 
// Good: divergence at warp granularity
int warp_id = idx / 32;
if (warp_id % 2 == 0) { ... }  // Entire warps take same path

Why this works: If all threads in a warp take the same path, there's no divergence. Restructure data/computation so similar work groups by32s.

2. Predication for short branches

Modern GPUs support predicated execution for small conditionals:

// Compiler may convert to predication (no branch):
result = (condition) ? a : b;  // Both computed, one selected per thread

When this helps: If both paths are<~5 instructions, predication avoids divergence overhead. Both sides execute, but results are masked. Only beneficial if $T_A + T_B$ (serialized) $> 2 \cdot \max(T_A, T_B)$ (predicated).

3. Sort/bucket data

Preprocess data so similar items are adjacent, aligned to warp size:

// Before kernel: sort particles by type in groups of 32
// Then: each warp processes uniform type, zero divergence

4. Dynamic paralelism / warp-level intrinsics

Use __ballotsync, __shfl_sync to detect and handle divergence explicitly:

unsigned mask = __ballot_sync(0xFFFFFFFF, condition);
if (mask == 0xFFFFFFFF) {
    // All threads true: fast path
} else if (mask == 0) {
    // All threads false: fast path
} else {
    // Mixed: handle carefully
}

Performance Characteristics

For a warp split into $k$ groups executing paths of lengths $I_1, I_2, \ldots, I_k$:

Binary split (50/50) with equal path lengths:50% waste.

Worst case (32-way split, equal lengths): $\frac{32I - I}{32I} = 96.875\%$ waste.

Real-world measurements:

• Well-structured code (warp-aligned divergence): <5% loss
• Typical ray tracing (material branching): 30-50% loss
• Pathological data-dependent (sparse matrices): 70-90% loss

You tell all the boys to work on problem A while the girls sit and wait. Then you tell all the girls to work on problem B while the boys sit and wait. Even though the kids could have worked in parallel, you had to serialize it because you can only give one instruction at a time.

That's warp divergence. The GPU can only say one thing to all 32 threads, so when they need to do different things, it has to go one at a time through each group. The more different groups, the more waiting everyone does.

The fix? Organize your class so all the boys are in one room and all the girls in another (warp-aligned divergence), or give worksheets that don't have "if you're a boy/girl" splits at all (branchless code).

Connections

• 6.2.8-SIMT-execution-model - Fundamental execution model that causes divergence
• 6.2.9-Warp-scheduling - How scheduler handles divergent warps
• 6.2.10-Memory-coalescing - Memory efficiency also requires uniform access patterns
• 8.4.2-Branch-prediction - Why CPU branch prediction doesn't apply to GPUs
• 9.3.5-Parallel-algorithmsdesign - Designing algorithms to avoid divergence
• 6.2.12-Occupancy-optimization - Divergence reduces effective occupancy
• 10.1.3-CUDA-optimization-patterns - Practical patterns for divergence-free code

#flashcards/hardware

What is a warp in GPU architecture? :: A group of 32 threads (NVIDIA) or 64 threads (AMD wavefront) that execute the same instruction simultaneously in lockstep. They share one instruction pointer and are the fundamental unit of GPU execution.

Concept Map

Hinglish (regional understanding)