Noesis
6.2.11 · D1 · HinglishGPU Architecture

Foundations — Warp divergence penalties

2,375 words11 min read↑ Read in English

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

Is page pe assume kiya gaya hai ki aapne parent note ke words mein se koi bhi pehle nahi dekha. Hum har ek word build karenge — thread, warp, SIMT, mask, cycle, branch — ek picture se, ek aisi order mein jahan har idea sirf pehle se bane ideas use karta hai. Jab aap finish karein, parent note padne jaiye aur har symbol ka matlab pehle se clear hoga.

1. Thread — ek worker jo ek recipe follow kar raha hai

Yeh topic isko kyun zaroori samajhta hai? Kyunki GPU ki puri trick hai in hazaron workers ko ek saath chalana. Bahut saare ke baare mein baat karne se pehle, hume ek ke baare mein crystal clear hona chahiye: ek thread exactly ek worker hai, ek recipe hai, ek waqt mein ek step hai.

Figure mein har worker ke paas same recipe card hai lekin alag idx hai — ek personal ID number. Wahi ID hai jo worker #3 ko worker #17 se alag behave karati hai. idx ko yaad rakho; yeh almost sabhi divergence ka source hai.

Recall Woh kaunsi single value hai jo ek thread ko same code run karne wale uske neighbor se alag behave karati hai?

Uska thread index idx ::: per-thread ID number, jaise threadIdx.x.

2. idx — personal ID number

32 workers ki ek row imagine karo, jo 0 se 31 tak numbered jerseys pehne hue khade hain. Jab recipe kehti hai if (idx < 16), workers 0–15 jawab dete hain "haan, main hoon" aur workers 16–31 jawab dete hain "nahi". Recipe identical hai; sirf jersey number alag hai. Yeh single fact baad mein pure warp ko split kara dega.

3. Warp — 32 workers lockstep mein march karte hue

Yeh woh picture hai jo baaki sab kuch sensible banati hai:

Dekho ek conductor ka baton (instruction pointer, burnt orange mein) jo ek step ko point kar raha hai. Saare 32 workers zaroor us step par hone chahiye. Har worker ke paas ek baton nahi hai — 32 ki puri team ke liye ek baton hai. Divergence kyun hurt karti hai uski poori wajah yahi hai: ek baton share karne wali team do workers ko do alag steps par simultaneously nahi rakh sakti.

(AMD same idea ko 64 threads ke saath banata hai aur ise wavefront kehta hai. Same picture, wider team.)

4. SIMT — "ek instruction, bahut saare workers ko batai gayi"

Exactly yeh tool kyun aur "har thread ko independently run karo" nahi? Kyunki independent threads ko apna khud ka baton, decoder, aur clock chahiye hoga — mehenga. SIMT is sawaal ka jawab deta hai "main 32 workers ko ek ko command karne ki cost mein kaise command karun?" Jawab: unhe ek shared command do aur har ek ko alag data par apply karne do.

Pakad, jiske baare mein poora topic hai: ek shared command ka matlab har worker ko same step obey karne mein capable hona chahiye. Agar step 5 hai "left mudo" lekin worker #17 ki recipe kehti hai "right mudo," shared baton atki hui hai — aur Section 6 dikhata hai hardware iske baare mein kya karta hai.

Recall SIMT 32 threads ko kya share karne deta hai, hardware cost bachata hai?

Ek instruction decoder / ek instruction pointer (baton) ::: saare 32 ko broadcast, har ek apne khud ke data[idx] par apply karta hai.

Related deeper notes: 6.2.8-SIMT-execution-model aur 6.2.9-Warp-scheduling.

5. Branch — recipe mein fork

32 marching workers ko ek fork tak pahunchte hue imagine karo. Ek sign ek sawaal poochta hai — "kya tera idx 16 se kam hai?" Har worker apne jersey number se jawab deta hai, isliye sab ek hi taraf point nahi karte.

  • Agar har worker same jawab deta hai (sab "haan" ya sab "nahi"), team saath milkar ek path par mud jaati hai. Koi problem nahi — shared baton abhi bhi kaam karta hai, sab same step par hain. Yeh ek uniform branch hai.
  • Agar workers alag-alag jawab dete hain, team split ho jaati hai. Yeh divergence hai, aur shared baton ek saath do taraf point nahi kar sakta.

Yahi impossibility poora crisis hai. Hume ek aur tool chahiye dekhne ke liye ki hardware isse kaise escape karta hai.

6. Mask — on/off switches ki ek row

Yeh hardware ka clever escape hai. Woh baton ko do taraf point nahi kar sakta, isliye woh iske bajaye yeh karta hai:

  1. Baton ko path X pe point karo. Mask ko sirf un threads ke liye on karo jinhoune X choose kiya; baaki sabko off karo. Path X run karo. Off threads idle rehte hain — present, powered, lekin unke results discard kar diye jaate hain.
  2. Baton ko path Y pe point karo. Mask flip karo: Y-threads ke liye on, X-threads ke liye off. Path Y run karo.
  3. Dono paths done → saare switches waapis on karo → reconverge aur saath milkar march karo.

Figure mein switches ki do rows padho. Notice karo ki path X aur path Y ek ke baad ek run hote hain, kabhi saath nahi. Yahi serialization hai: do paths jo ho sakti thi simultaneous wo ab time mein stack ho gayi hain. Har row mein greyed-out (masked-off) workers wasted capacity hain — divergence ki poori cost unhi grey squares mein rehti hai.

Recall Jab threads diverge karte hain, kya hardware dono paths ek saath run karta hai?

Nahi — woh unhe ek ke baad ek run karta hai, mask use karke yeh switch karte hue ki kaun se threads har pass par active hain ::: masked-off threads idle rehte hain aur ek clock tick burn karte hain.

7. Cycle — hardware clock ki ek tick

Seconds ki jagah cycles mein kyun measure karte hain? Kyunki seconds specific chip ki speed par depend karte hain, lekin wasted-to-useful cycles ka ratio har chip par same hota hai. Jab parent note likhta hai cycles aur cycles, toh yeh inhi ticks ko count kar raha hai. Do paths jo overlap kar sakti thi ab ticks cost karti hain ki jagah.

Aapke paas ab wo har symbol hai jo parent note use karta hai. idx, warp, SIMT, branch, mask, cycle, sum-vs-max — har ek ek picture hai. Parent ke formulas sirf inhi pictures par arithmetic hain.

8. Foundations topic ko kaise feed karte hain

Thread = one worker, one recipe

idx = personal ID per thread

Warp = 32 threads share one baton

SIMT = one instruction to many threads

Branch = fork in the recipe

Divergence = threads answer the fork differently

Mask = on off switches per thread

Cycle = one clock tick of cost

Warp Divergence Penalties 6.2.11

Upar se neeche padho: is map par kuch bhi kisi aisi box ko use nahi karta jo uske upar pehle se na bani ho. Do roots (ek single thread, aur uska idx) warp mein grow karte hain, jise banane layak hone ke liye SIMT chahiye, jo branches ko dangerous banata hai, jo masks force karta hai, jiski cost hum cycles mein count karte hain — aur wahi cost topic hai.

9. Aage kahan jaana hai

  • Lockstep rule depth mein: 6.2.8-SIMT-execution-model
  • Warps ko har cycle run karne ke liye kaise pick kiya jaata hai: 6.2.9-Warp-scheduling
  • Doosra tarika jisme memory access order aapko hurt karta hai: 6.2.10-Memory-coalescing
  • CPUs ek related fork problem ko kaise alag deal karte hain: 8.4.2-Branch-prediction
  • Algorithms design karna taki warps diverge na karein: 9.3.5-Parallel-algorithmsdesign
  • Penalty hide karne ke liye enough warps alive rakhna: 6.2.12-Occupancy-optimization
  • Real GPUs par practical fixes: 10.1.3-CUDA-optimization-patterns

Equipment checklist

Right side cover karo; kya aap reveal karne se pehle answer de sakte ho?

  • Ek thread hai ::: instructions ka ek stream — ek single worker jo ek recipe follow karta hai, ek waqt mein ek step.
  • idx hai ::: unique whole-number ID (0,1,2,…) jo otherwise-identical threads ko alag data touch karata hai.
  • Ek warp hai ::: 32 threads ki ek fixed team jo same instant par same instruction execute karne ke liye force ki jaati hai, ek instruction pointer share karti hai.
  • "Puri team ke liye ek baton" picture represent karta hai ::: single shared instruction pointer — poora warp hamesha same recipe step par hota hai.
  • SIMT stands for aur matlab hai ::: Single Instruction, Multiple Thread — ek broadcast command jo bahut saare threads ke dwara apne data par apply ki jaati hai.
  • Ek branch hai ::: recipe mein ek fork (if/else, while) jahan threads alag-alag taraf ja sakti hain.
  • Warp divergence hai ::: ek warp ke threads ka ek branch ka alag-alag jawab dena, isliye shared baton do taraf point nahi kar sakta.
  • Ek execution mask hai ::: 32 on/off switches ki ek row jo decide karti hai ki kaun se threads current instruction par act karte hain; off threads idle rehte hain.
  • Reconvergence hai ::: saare divergent paths finish hone ke baad har mask switch ko waapis on karna, taki warp phir saath milkar march kare.
  • Ek cycle hai ::: ek clock tick, woh unit jisme hum divergence cost count karte hain.
  • Diverged cost versus ideal cost hai ::: saare path lengths ka sum () versus maximum single path ().