Warp divergence penalties — Hinglish notes

6.2.11 · Hardware › GPU Architecture

Overview

Jab ek warp ke threads conditional code mein alag-alag execution paths lete hain, toh GPU ko un paths ko serialize karna padta hai, jisse woh parallel efficiency destroy ho jaati hai jo GPUs ko powerful banati hai. Isse warp divergence kehte hain, aur yeh NVIDIA GPUs par effective throughput ko peak ka 1/32 tak reduce kar sakta hai.

GPUs SIMT (Single Instruction Multiple Thread) use karte hain: ek instruction decoder 32 execution units ko broadcast karta hai. Jab threads diverge karte hain, hardware ko har branch ko serially execute karna padta hai, inactive threads ko mask off karke.

The Mechanism

Hardware Divergence Ko Kaise Handle Karta Hai

Step-by-step execution model:

Saare threads ek saath branch par pahunchte hain (same instruction pointer)
Har thread condition evaluate karta hai → true/false produce karta hai
Hardware execution mask create karta hai: bitmask jo mark karta hai ki kaun se threads kaun si taraf jaate hain
Pehla path execute karo mask A active ke saath (baaki threads idle, power burn karte hain)
Doosre path ko divergence stack par push karo
Doosra path execute karo mask B active ke saath (pehla group ab idle)
Reconverge jab dono paths ek common point par milte hain

T_{diverged} = i = 1 \sum k T_{i}

Jahan $T_{i}$ path $i$ ki execution time hai.

Yeh formula kyun? Divergence ke bina, saare 32 threads parallel execute karte hain: $T_{ideal} = max (T_{1}, T_{2}, \dots, T_{k})$ . Divergence ke saath, hum serialize karte hain: har path sequentially run karta hai. Speedup loss yeh hai:

Efficiency = \frac{max ( T _{1} , T _{2} , \dots , T _{k} )}{\sum _{i = 1}^{k} T _{i}}

Worst-case ke liye derivation (har thread alag):

Agar har thread ek unique path leta hai ek conditional mein jisme $n$ instructions hain:

Ideal: 1 warp $n$ instructions execute karta hai time $n \cdot t_{cycle}$ mein
Diverged: 32 sequential executions, har ek $n$ cycles → $32 n \cdot t_{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  
    }
}

Ek warp ke andar kya hota hai (threads 0-31):

Saare 32 threads idx < 16 evaluate karte hain
Threads 0-15 → true, threads 16-31 → false
Path A execute karo (multiply): threads 0-15 active, 16-31 masked off (idle) → 10 instruction cycles
Path B execute karo (add): threads 16-31 active, 0-15 masked off → 8 instruction cycles
Kernel exit par Reconverge

Total cost: 10 + 8 = 18 instruction cycles

Divergence ke bina (agar sab same path lete): max(10, 8) = 10 cycles

Efficiency: 10/18 = 55.6%

Yeh step kyun? SIMT mein hardware inactive threads ke liye instructions skip nahi kar sakta—use unhe decode aur issue karna padta hai, bas masked outputs ke saath. Yahi serialization penalty hai.

__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: Ek warp mein, 8 threads ka type=0 hai, 10 ka type=1, 8 ka type=2, 6 ka type=3.

Execution:

Path type=0: sqrt, 20 cycles, 8 threads active
Path type=1: exp, 30 cycles, 10 threads active
Path type=2: log, 25 cycles, 8 threads active
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%

Yeh kyun matter karta hai: Data-dependent branching jahan har thread alag path le sakta hai, GPU efficiency ke liye catastrophic hai. Yeh ray tracing (alag materials), particle systems (alag particle types), ya sparse matrix operations mein common hai.

__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;
    }
}

Agar counts[] = [5, 20, 3, 15 ...] ek warp mein:

count=3 wala thread pehle khatam hota hai, idle ho jaata hai
count=20 wala thread poore warp ko busy rakhta hai jab tak woh khatam na ho
Effective iterations: $max (counts) = 20$
Utilization: $\frac{\sum counts}{32 \cdot m a x ( counts )}$

Kyun? Loop exit ek branch hai. Har iteration par, kuch threads exit karte hain (condition false) jabki baaki continue karte hain. Hardware ko iterations execute karte rehna padta hai jab tak aakhri thread khatam na ho, shrinking active mask ke saath.

Common Patterns aur Mitigation

Kyun sahi lagta hai: Tum sochte ho, "32 mein se sirf 2 threads else branch lete hain, toh 93% same kaam kar rahe hain—kya yeh okay nahi hona chahiye?"

Kyun galat hai: Hardware dono branches ko poori tarah se poore warp ke liye execute karta hai, chahe kitne bhi threads koi bhi path lein. Agar sirf 1 thread alag path leta hai, tab bhi tum dono paths ki poori cost pay karte ho.

Example:

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

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

Fix: Samjho ki divergence per-warp hoti hai, aur cost sabhi paths ka sum hoti hai, thread count se weighted nahi. Bilkul branching avoid karne ke liye restructure karo, ya ensure karo ki divergence warp boundaries par ho.

Kyun sahi lagta hai: Tum sochte ho nested conditionals zyada divergence points create karte hain.

Kyun nuanced hai: Ek same divergent path ke andar nested ifs extra serialization add nahi karte—woh ek path ke instruction count ka hissa hain. Lekin nested ifs jo orthogonal divergence patterns create karte hain, execution paths ki sankhya multiply kar dete hain.

Example:

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

Worst case: 4 distinct paths agar distribution buri ho. Cost: sabhi path lengths ka sum.

Fix: "Warp mein distinct execution paths" ke terms mein socho, na ki syntactic nesting depth mein.

Mitigation Strategies

Divergence ko warp boundaries ke saath align karo

// 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

Yeh kyun kaam karta hai: Agar ek warp ke saare threads same path lete hain, toh koi divergence nahi hai. Data/computation ko restructure karo taaki similar work 32s mein group ho.

Chote branches ke liye Predication

Modern GPUs chote conditionals ke liye predicated execution support karte hain:

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

Yeh kab help karta hai: Agar dono paths <~5 instructions hain, toh predication divergence overhead avoid karta hai. Dono sides execute hoti hain, lekin results masked hote hain. Tabhi beneficial hai jab $T_{A} + T_{B}$ (serialized) $> 2 \cdot max (T_{A}, T_{B})$ (predicated).

Data ko Sort/bucket karo

Data ko preprocess karo taaki similar items adjacent hon, warp size ke aligned:

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

Dynamic parallelism / warp-level intrinsics

Divergence ko explicitly detect aur handle karne ke liye __ballot_sync, __shfl_sync use karo:

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

Ek warp ke liye jo $k$ groups mein split hai aur paths of lengths $I_{1}, I_{2}, \dots, I_{k}$ execute kar raha hai:

Instructions executed = i = 1 \sum k I_{i}

Instructions useful = i max I_{i}

Wasted work ratio = \frac{\sum I _{i} - max I _{i}}{\sum I _{i}}

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

Worst case (32-way split, equal lengths): $\frac{32 I - I}{32 I} = 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

Recall

Ek 12-saal ke bachche ko explain karo Socho tum teacher ho, aur tumhari class mein 32 bachche ek worksheet kar rahe hain. Tum ek waqt mein sirf EK instruction de sakte ho, aur sabko sunna padta hai. Worksheet kehti hai "Agar tum ladke ho, problem A solve karo. Agar tum ladki ho, problem B solve karo."

Tum saare ladkon ko problem A par kaam karne kehte ho jabki ladkiyaan baith kar wait karti hain. Phir tum saari ladkiyon ko problem B par kaam karne kehte ho jabki ladke baith kar wait karte hain. Bhale hi bachche parallel mein kaam kar sakte the, tumhe use serialize karna pada kyunki tum ek waqt mein sirf ek instruction de sakte ho.

Yahi warp divergence hai. GPU saare 32 threads ko sirf ek cheez keh sakta hai, isliye jab unhe alag-alag kaam karna hota hai, use har group ke through ek-ek karke jaana padta hai. Jitne zyada alag groups, utna zyada sabka wait karna.

Fix? Apni class ko organize karo taaki saare ladke ek room mein hon aur saari ladkiyaan doosre mein (warp-aligned divergence), ya aise worksheets do jo "agar ladka/ladki ho" splits na rakhen (branchless code).

Connections

6.2.8-SIMT-execution-model - Fundamental execution model jo divergence cause karta hai
6.2.9-Warp-scheduling - Scheduler divergent warps ko kaise handle karta hai
6.2.10-Memory-coalescing - Memory efficiency bhi uniform access patterns maangti hai
8.4.2-Branch-prediction - Kyun CPU branch prediction GPUs par apply nahi hoti
9.3.5-Parallel-algorithmsdesign - Divergence avoid karne ke liye algorithms design karna
6.2.12-Occupancy-optimization - Divergence effective occupancy reduce karta hai
10.1.3-CUDA-optimization-patterns - Divergence-free code ke liye practical patterns

#flashcards/hardware

What is a warp in GPU architecture? :: 32 threads (NVIDIA) ya 64 threads (AMD wavefront) ka ek group jo lockstep mein same instruction simultaneously execute karta hai. Woh ek instruction pointer share karte hain aur GPU execution ka fundamental unit hain.

What is warp divergence?

Jab ek single warp ke threads conditional code (if/else, loops) mein alag-alag execution paths lete hain, GPU ko har path ki execution serialize karni padti hai, inactive threads masked off hote hain lekin phir bhi time consume karte hain.

How does the GPU handle warp divergence mechanically?

Hardware har branch path ke liye execution masks create karta hai, pehle path ko kuch threads active ke saath execute karta hai (baaki idle), remaining paths ko divergence stack par push karta hai, har path ko serially execute karta hai, phir jab saare paths complete hote hain toh reconverge karta hai.

What is the performance cost formula for divergence with k paths?

T_diverged = sabhi k paths ke T_i ka sum, versus T_ideal = parallel mein max(T_i). Efficiency = max(T_i) / sum(T_i). Worst case 32 distinct paths ke saath: 96.875% wasted work.

Why doesn't "only 1 thread diverges" mean low cost?

Kyunki hardware ko dono branch paths poore warp duration ke liye poori tarah execute karne padte hain, chahe kitne bhi threads koi bhi path lein. Cost path lengths ka sum hai, thread count se weighted nahi.

What is the key strategy to minimize divergence?

Divergence ko warp boundaries (32-thread groups) ke saath align karo. Ensure karo ki ek warp ke saare threads same path lein, data/work ko organize karke taaki similar items 32s ke multiples mein grouped hon.

When does predicated execution help with divergence?

Chote branches ke liye (<~5 instructions per path) jahan dono sides compute karna aur per-thread results select karna branches serialize karne se sasta ho. Compiler maskable instructions mein convert karta hai bina actual branch ke. Tabhi beneficial hai jab

T_{A} + T_{B}

(serialized)

> 2 \cdot max (T_{A}, T_{B})

(predicated).

What is the worst-case divergence scenario?

Data-dependent branching jahan ek warp ka har thread bilkul alag execution path leta hai (jaise ray tracing mein 32 alag material types, random loop counts). Isse saare 32 paths serialize hote hain: 32x slowdown.

How do loops cause divergence?

Loop exit conditions branches hain. Har iteration par, kuch threads exit kar sakte hain (condition false) jabki baaki continue karte hain, divergence create karte hue. Warp ko iterations execute karte rehna padta hai jab tak aakhri thread khatam na ho, shrinking active masks ke saath.

What is __ballot_sync used for in divergence handling?

Ek bitmask return karta hai jo dikhata hai ki warp mein kaun se threads condition pass kiye. Uniform (sab true/false) versus divergent cases detect karne ki allow karta hai, fast-path optimizations enable karta hai jab saare threads agree karte hain.