Warp divergence penalties: Exercises — Hinglish deep dive

6.2.11 · D4 · Hardware › GPU Architecture › Warp divergence penalties

Yeh page Warp divergence penalties (6.2.11) ka self-test ladder hai. Upar se neeche karo: har level demand badhata hai — divergence spot karne se lekar, uska cost compute karne tak, nested/loop structure analyse karne tak, fixes design karne tak, aur trade-offs master karne tak.

Har solution ko cover karo aur pehle khud try karo. Apna answer pakka karne ke baad hi reveal karo.

Shuru karne se pehle, do symbols jo hum har jagah use karenge, zero se define kiye gaye hain:

Hum average kyun nahi karte, add kyun karte hain: SIMT mein ek instruction decoder hai jo 32 lanes ko feed karta hai. Woh ek baar mein sirf ek path ki instructions ko point kar sakta hai. Doosre path ke threads abhi bhi "wahan" hain — woh bas masked hain: unke outputs discard ho jaate hain, lekin cycles phir bhi tik karte rehte hain. Toh ek path jo sirf ek thread leta hai, uska full length ka cost lagta hai. Yeh dhyan rakhna — yahi aadhe exercises ka trap hai.

Neeche di gayi picture har exercise ke liye mental model hai is page par. Shuru karne se pehle isse study karo: top bar ideal duniya dikhata hai jahan do paths (A = 12 cycles, B = 20 cycles) parallel mein overlap hote hain, toh warp $max (12, 20) = 20$ par finish karta hai. Bottom bar diverged duniya dikhata hai jahan wahi do paths back-to-back run hote hain, toh warp $12 + 20 = 32$ par finish karta hai. Dayi taraf ka amber gap pure divergence waste hai — bilkul wahi quantity jo tum neeche ke problems mein $T_{div} - T_{ideal}$ ke roop mein compute karte ho.

Level 1 — Recognition

Recall

Solution L1.1 Koi divergence nahi. Har thread index 0–31 condition < 32 satisfy karta hai, toh saare 32 threads condition ko true evaluate karte hain aur path A saath mein lete hain. else path B koi nahi leta, toh woh kabhi issue nahi hota. Kyunki saare threads ek path share karte hain, serialize karne ki koi zaroorat nahi. Divergence ke liye ek hi warp ke andar asahmat hona zaroori hai.

Worked example

L1.2 — Divergent line dhundho Neeche konsi line intra-warp divergence cause karti hai? (warp = threads 0–31)

int idx = threadIdx.x;          // (a)
float x = data[idx];            // (b)
if (idx % 2 == 0) y = x*x;      // (c)
data[idx] = y;                  // (d)

Recall

Solution L1.2 Line (c). Predicate idx % 2 == 0 even lanes (0,2,…,30) ke liye true aur odd lanes (1,3,…,31) ke liye false hai. Usi warp mein 32 threads split ho jaate hain — 16 true, 16 false — toh hardware ko body ko ek baar even-mask ke saath aur ek baar (yahan empty, implicit else) odd-mask ke saath issue karna padta hai. Lines (a), (b), (d) mein koi branch nahi hai: har thread apne address par same op karta hai, jo uniform control flow hai (dekho coalescing, divergence se alag concern).

Level 2 — Application

Recall

Solution L2.1 Dono paths kam se kam ek thread le raha hai, toh dono count karte hain. $T_{ideal} = max (15, 25) = 25 cycles$ $T_{div} = 15 + 25 = 40 cycles$ $Efficiency = \frac{25}{40} = 0.625 = 62.5%$ Thread count (8 vs 24) arithmetic mein kabhi nahi aaya — sirf kya ek path liya jaata hai aur woh kitna lamba hai.

Worked example

L2.2 — Four-way switch Ek warp 4-way if/else if chain run karta hai. Path costs aur kitne threads har ek lete hain:

path	cycles	threads
0	20	8
1	30	10
2	25	8
3	5	6

$T_{div}$ , $T_{ideal}$ , efficiency dhundho.

Recall

Solution L2.2 Chaaon paths mein se har ek ka kam se kam ek taker hai, toh chaaon serialize karte hain: $T_{div} = 20 + 30 + 25 + 5 = 80 cycles$ $T_{ideal} = max (20, 30, 25, 5) = 30 cycles$ $Efficiency = \frac{30}{80} = 0.375 = 37.5%$ Yeh parent note ke Example 2 se match karta hai — data-dependent 4-way branching brutal hai.

Recall

Solution L2.3 Do paths liye ja rahe hain (lane 0 → path1; lanes 1–31 → path2), toh dono poore count karte hain: $T_{div} = 100 + 50 = 150 cycles, T_{ideal} = max (100, 50) = 100$ $Efficiency = \frac{100}{150} = 0.6 \overline{6} = 66.7%$ "Weighted" guess $50 + 100/32 \approx 53.1$ cycles galat hai kyunki 100-cycle path 32 threads mein share nahi hota — sirf ek lane ke liye bhi uska poora length cost lagta hai.

Level 3 — Analysis

Worked example

L3.1 — Nested branches, worst case

if (A) { if (B) p1; else p2; }
else   { if (C) p3; else p4; }

Path costs: p1=10, p2=14, p3=8, p4=12. Ek warp mein predicates is tarah split hote hain ki chaaon leaves mein se har ek ko kam se kam ek thread milta hai. $T_{div}$ aur efficiency do. Phir best-case $T_{div}$ do agar warp aise arrange ho ki sirf p1 aur p3 kabhi reach ho.

Recall

Solution L3.1 Worst case — 4 distinct paths reached: $T_{div} = 10 + 14 + 8 + 12 = 44 cycles, T_{ideal} = max (10, 14, 8, 12) = 14 cycles$ $Efficiency = \frac{14}{44} \approx 0.318 = 31.8%$ Best arranged case — sirf p1 aur p3 liye gaye (e.g. B true jab bhi A, C true jab bhi not-A): $T_{div} = 10 + 8 = 18 cycles, T_{ideal} = max (10, 8) = 10 cycles$ $Efficiency = \frac{10}{18} \approx 0.556 = 55.6%$ Nesting depth ne cost set nahi ki — distinct leaves jo actually reach hue unki number ne ki. Tree figure dekho.

Recall

Solution L3.2 Loop back-edge ek branch hai: har iteration par, woh lanes jinki i unki counts[idx] tak pahunch gayi hai drop out ho jaati hain (masked), lekin warp ko iterations run karte rehna padta hai jab tak max-count lane (20) finish nahi ho jaati. Toh warp $max = 20$ iterations worth of issue slots saari 32 lanes ke liye execute karta hai, jinmein se sirf $\sum = 320$ useful hain. $Utilisation = \frac{320}{32 \times 20} = \frac{320}{640} = 0.5 = 50%$ Short lanes ke exit hone ke baad har iteration slot ka aadha masked lanes par waste ho jaata hai.

Recall

Solution L3.3 $Eff_{X} = \frac{m a x ( 40 , 40 , 40 )}{120} = \frac{40}{120} = 33.3%$ $Eff_{Y} = \frac{m a x ( 10 , 10 , 10 , 10 , 10 )}{120} = \frac{10}{120} = 8.3%$ Warp Y zyada bura hai. Same total serialized time, lekin Y ka ideal (best single path) bahut chhota hai, toh Y ka zyada time pure divergence waste hai. Lesson: zyada, chhote, equally-populated paths divergence-inefficiency ka sabse kharab shape hain — ideal shrink hota hai jabki sum wahi rehta hai.

Level 4 — Synthesis

Recall

Solution L4.1 idx % 2 har lane mein alternate karta hai, split guarantee karta hai. Ise warp-granular predicate se replace karo:

int warp_id = idx / 32;
if (warp_id % 2 == 0) heavy(); else light();

Ab ek warp ke saare 32 lanes same warp_id share karte hain, toh ek warp poori tarah heavy ya poori tarah light hota hai — ek ek path each. Ek all-heavy warp ke liye: sirf ek path liya jaata hai jo 60 cycles ka hai, $T_{div} = 60 = T_{ideal}$ : $Efficiency = \frac{60}{60} = 100%$ Total work unchanged hai, lekin divergence penalty eliminate ho gayi. Yeh parent note ki alignment strategy hai, aur yeh directly occupancy planning mein jaata hai.

Recall

Solution L4.2 Pehle, rule kahan se aata hai? Do-path conditional run karne ke do tarike hain:

Branch (serialize): $T_{A} + T_{B}$ pay karo — har path ek baar issue hota hai, back-to-back.
Predicate (no branch): har lane dono paths run karta hai aur unwanted result mask karta hai, toh woh work ka $T_{A} + T_{B}$ bhi pay karta hai — lekin woh branch instruction, mask-stack push/pop, aur reconvergence delete kar deta hai, jo roughly ek "extra max-path" worth overhead cost karte hain. Toh predication ka effective charge roughly $2 max (T_{A}, T_{B})$ hota hai: tum bade path ka kaam do baar karte ho (ek baar real, ek baar masked) bina branch machinery ke. Predication ka charge branching ke charge se neeche set karo: $branch charge T_{A} + T_{B} > predication charge 2 max (T_{A}, T_{B}) \Rightarrow predicate is cheaper.$ Yeh inequality exactly crossover hai: yeh sirf tab hold karta hai jab dono paths short aur balanced hain (toh unka sum do baar bade wale se zyada hota hai), aur tab fail hota hai jab ek path doosre ko dwarf karta hai (tab $2 max$ sum se aage nikal jaata hai aur real branch jeet jaata hai).

Ab apply karo $T_{A} = 4$ , $T_{B} = 3$ ke saath: $T_{A} + T_{B} = 7, 2 max (T_{A}, T_{B}) = 2 \times 4 = 8.$ Kyunki $7 < 8$ hai, serialized-branch sum predication charge se pehle hi kam hai — yeh branch itna short aur balanced hai ki real branch raw-cost mein equal-ya-better choice hai, phir bhi practice mein compilers yahan predicate karte hain kyunki yeh divergence stall aur reconvergence ko remove karta hai essentially same cycle count par. Bahut lopsided long branches ke liye ( $T_{A} + T_{B} > 2 max$ ) tum real branch prefer karoge taaki badi path un lanes par na pay ho jo use nahi chahti. CPUs par related idea: branch prediction, jo GPUs mein nahi hota.

Recall

Solution L4.3 Sorted, har warp uniform hai → single path → $T_{div} = T_{ideal}$ , efficiency $= 100%$ har warp ke liye. Chaar warp-flavours 20, 30, 25, 5 cost karte hain respectively; average per-warp time hai $\frac{20 + 30 + 25 + 5}{4} = \frac{80}{4} = 20 cycles/warp.$ Compare karo: mixed warp ne wahi chaar tarah ke kaam ke liye 80 cycles spend kiye. Sorting ne yahan $4 \times$ speed-up dilaya purely divergence ko kill karke — classic parallel-algorithm preprocessing trick, jo CUDA optimization patterns ka bhi ek staple hai.

Level 5 — Mastery

Recall

Solution L5.1 Har warp ko two-number method se price karo. Ek warp jo do liye gaye paths mein split hota hai uski cost sum hoti hai; ek uniform warp ki cost uska single path hoti hai.

Warp 0 (uniform, ek 30-cycle path): $T_{div} = 30$ .
Warps 1, 2, 3 (har ek 30 aur 18 mein split hota hai, dono liye jaate hain): $T_{div} = 30 + 18 = 48$ har ek.

Saare chaar warps mein issued cycles sum karo: $Total issued = 30 + 3 \times 48 = 30 + 144 = 174 cycles.$ All-uniform ideal har warp ko sirf uska sabse lamba single path ek baar charge karta hai ( $max (30, 18) = 30$ ), chaar warps mein: $T_{ideal}^{block} = 4 \times 30 = 120 cycles.$ $Block efficiency = \frac{120}{174} \approx 0.690 = 69.0%.$ Ek clean warp teen dirty walon ko rescue nahi kar sakta — divergence per-warp hai aur penalties simply block mein add hoti hain.

Recall

Solution L5.2 Option A — har warp split karta hai, dono paths liye jaate hain: per-warp $= 40 + 10 = 50$ ; chaar warps $= 4 \times 50 = 200$ cycles. Option B — 2 warps uniform-heavy ( $40$ each) + 2 warps uniform-light ( $10$ each): $2 \times 40 + 2 \times 10 = 80 + 20 = 100 cycles.$ $Speed-up = \frac{200}{100} = 2.0 \times .$ Sorting ne issued cycles aadhe kar diye kyunki har warp internally uniform ban gaya. Yeh upar sab ka synthesis hai: divergence recognize karo (L1), price karo (L2), split analyse karo (L3), fix chuno (L4), aur scale par win prove karo (L5).

Recall

Solution L5.3 Dono ke liye divergent region: $25 + 15 = 40$ cycles.

Warp P: total $40$ cycles.
Warp Q: paths reconverge hone ke baad, saare 32 lanes phir se active hain aur ek instruction stream share karte hain, toh 10-cycle tail ek baar issue hota hai: $40 + 10 = 50$ cycles. Tail uniform control flow hai — koi split nahi, koi serialization nahi — toh woh ek baar add hota hai, per path sum nahi hota. Reconvergence exactly woh point hai jahan divergence stack pop hota hai aur warp full-width SIMT par wapas jaata hai.

Recall

Self-test: ek formula jo saath le jaao Tum kisi bhi divergent region ko kaise price karte ho? ::: $T_{div} = \sum (un paths ki lengths jinke \geq 1 taker hain)$ ; $T_{ideal} = max (woh lengths)$ ; efficiency $= T_{ideal} / T_{div}$ . Populations timing affect nahi karti; pehle/baad ka uniform code ek baar run hota hai.