6.2.5 · D1Warps and warp scheduling
Ek GPU threads ko ek-ek karke nahi, balki 32 ke fixed bundles mein chalata hai, aur bundle ka har thread exactly usi instruction ko usi moment par run karne ke liye majboor hota hai. Is topic mein jo bhi hai — warps, scheduling, divergence, occupancy — sab usi ek design choice ka consequence hai: threads ko bundle karo, control share karo, waiting ko chhupaao.
Yeh page kuch bhi assume nahi karta. Warps and warp scheduling ko theek se padhne se pehle, hum har word, symbol, aur picture build karenge jis par woh tika hua hai. Upar se neeche padho; har block agla earn karta hai.
Warps se pehle, hume woh atom chahiye jisse woh bante hain.
Ek thread ek worker hai jo aapke program ki ek copy run karta hai , data ke apne ek chote se piece par kaam karta hua. Agar aap ek million numbers ke do arrays add karte ho, toh aap ek million threads launch karte ho aur thread number i compute karta hai C[i] = A[i] + B[i].
Picture: ek chota sa worker ek index number i pakde hua, memory ke ek slot ke upar khada hai.
Threads kabhi akele nahi rehte. Woh containers ke ek nested set mein aate hain, jise hum aage draw karte hain.
Grid ek launch ke liye threads ki poori fauj hai.Block grid ke andar ek team hai — usi block ke threads fast shared memory ke zariye baat kar sakte hain.Warp sabse chota squad hai jo hardware actually move karta hai: 32 threads ek saath glued.
Yeh nesting bilkul Thread-Blocks-and-Grids ka subject hai; yahan hume sirf boxes dekhne ki zarurat hai.
==threadIdx== woh coordinate label hai jo har thread carry karta hai taaki usse pata chale ki data ka kaun sa piece mera hai . Iske teen parts hain jo likhe jaate hain threadIdx.x, threadIdx.y, threadIdx.z.
Picture: figure s01 mein, ek block ka har cell ek thread hai; uska (x, y, z) usi block ke andar sirf uska row/column/layer address hai.
Intuition Teen numbers kyun aur ek kyun nahi?Real data aksar 2D (ek image) ya 3D (ek volume) hota hai. Ek thread ko "main column x, row y par hoon" kehne ki suvidha dena thread → pixel mapping ko natural banata hai. Lekin neecha ka hardware workers ki ek seedhi line hai, toh hume un coordinates ko ek single number mein flatten karna hoga. Woh flattening agla symbol hai.
==blockDim== batata hai ki ek block ke har axis par kitne threads hain : blockDim.x, blockDim.y, blockDim.z. Agar blockDim = (16, 8, 1) hai toh block 16 wide, 8 tall, 1 deep hai → 16 × 8 = 128
Picture: figure s01 mein cells ki grid ki width, height, aur depth.
Hardware saare threads ko single-file line mein lagata hai. Us line mein thread ki jagah dhundhne ke liye hum block ko row by row padhte hain (pehle x fastest, phir y, phir z). Ise row-major order kehte hain.
Worked example Formula ko17 × 8 (16, 7, 0) lo, matlab last column, last row, blockDim = (17, 8, 1).
lin = 0 × ( 17 × 8 ) + 7 × 17 + 16 = 0 + 119 + 16 = 135
Picture kyun row-major matter karta hai yeh dikhati hai: yeh decide karta hai ki kaun se threads ek hi warp mein jaate hain . Ek single warp 32 consecutive line positions ka ek slice hota hai — toh ek warp do rows ki boundary cross kar sakta hai.
Ab hum warp ko arithmetically define kar sakte hain.
Definition Warp (arithmetic form)Ek warp 32 consecutive linear indices ka ek run hai. Ek thread jis warp mein belong karta hai woh hai:
warpID = ⌊ 32 lin ⌋
Do naye symbols aaye. Inse milte hain:
⌊ x ⌋ ⌊ x ⌋ nearest whole number par round down karo . ⌊ 135/32 ⌋ = ⌊ 4.21 ⌋ = 4
Picture: number line par kharo, phir nearest integer tick par baayi taraf jao.
⌈ x ⌉ ⌈ x ⌉ nearest whole number par round up karo . ⌈ 136/32 ⌉ = ⌈ 4.25 ⌉ = 5 5 warps chahiye.
Picture: number line par nearest integer tick ki taraf daayein jao.
Intuition Yahan ceiling kyun chahiyeHardware sirf poore warps de sakta hai — aap "4.25 warps" nahi maang sakte. Toh warps ki count hamesha round up hoti hai: koi bhi bacha hua thread phir bhi ek poora warp cost karta hai. Yehi waste ka source hai, aur isliye parent note aapse baar baar kehta hai ki block sizes 32 ka multiple rakho.
Common mistake Floor vs ceiling mix-upFloor use karo jab poochho "kaun sa warp ek given thread mein hai?" (index → bucket).
Ceiling use karo jab poochho "in threads ko kitne warps chahiye?" (count → containers).
Lockstep matlab warp ka har thread ek hi clock tick par exactly wohi instruction execute karta hai . "Similar kaam" nahi — identical instruction, bas alag-alag data par.
Picture: 32 rower ek hi drumbeat par kheench rahe hain; ek instruction fetch saari 32 oars chalata hai.
SIMT = Single Instruction, Multiple Thread . Ek instruction stream, 32 threads us par sawaar, har ek ke apne data aur apne registers ke saath. Ise plain SIMD se compare karein SIMT-vs-SIMD mein.
Lekin branches hote hain — kabhi kabhi kuch threads ko ek instruction skip karni chahiye . Hardware ise ek mask se handle karta hai.
Active mask ek 32-bit on/off switch hai, har thread ke liye ek bit. 1 matlab "yeh thread instruction karta hai"; 0 matlab "yeh ek baar bahar baitho, kuch produce mat karo."
Picture: warp ke upar 32 light-switches ki ek strip; fetch ki gayi instruction chalti hai, lekin sirf lit threads act karte hain.
Intuition Mask divergence ko expensive kyun banata haiEk warp ek saath do alag instructions nahi chala sakta. Agar 20 threads if branch chahte hain aur 12 else chahte hain, toh warp if code chalata hai 12 switches off ke saath, phir else code chalata hai 20 switches off ke saath — dono paths ek ke baad ek hote hain. Woh serial replay hi warp divergence hai.
Ek cycle GPU ke clock ki ek tick hai — hardware jitna chota time measure karta hai. "400 cycles" ka matlab sirf "400 ticks ka wait" hai.
L L cycles ki woh number hai jo aapko result ka wait karna padta hai — khaskar door global memory se padhna, jo slow hai (saikdon cycles). GPU-Memory-Hierarchy mein dekho ki woh doori kahaan se aati hai.
Picture: ek warp "memory se fetch" ka ek arrow chalata hai; jawab L
Ek warp stalled hai jab woh aage nahi badh sakta — usually us latency L __syncthreads() par wait kar raha hota hai.
Intuition Kaafi warps rakhne ka poora pointEk stalled warp ek waste engine hai — jab tak koi doosra ready warp jump in karke khali ticks fill na kar sake. Woh swap latency hiding hai: jab warp 0 L
Occupancy = SM kitna warps se bhara hai , fraction ke roop mein:
Occupancy = Max Warps per SM Active Warps per SM necessary hai lekin sufficient nahi — Occupancy-vs-Performance mein explore kiya gaya hai.
Do symbols ise support karte hain:
SM (Streaming Multiprocessor): GPU par ek physical engine jo kaafi warps rakhta hai aur schedulers contain karta hai. Ise ek factory floor samjho; poore GPU mein kaafi hote hain.Warp scheduler: ek SM ke andar traffic controller jo har cycle ek ready warp pick karke issue karta hai.
Occupancy ko free mein 1 tak push kyun nahi kiya ja sakta: har warp registers (Register-Pressure ) aur shared memory claim karta hai. Zyada warps → har warp ke liye kam. Woh trade-off in foundations se real tuning story ka bridge hai.
Thread one worker one index
threadIdx and blockDim coordinates
Linear index row-major flatten
Floor and ceiling rounding
Warp 32 consecutive threads
Active mask on-off switches
Warp divergence serial replay
Latency hiding swap warps
Warps and warp scheduling
Self-test: right side cover karo aur aage padhne se pehle har ek ka jawab zor se bolo.
Thread kya hota hai, ek sentence mein? Ek worker jo apne data slice par program ki ek copy run karta hai.
threadIdx aur blockDim kya describe karte hain?Thread ka coordinate label, aur block ke har axis par kitne threads hain.
Thread (3,2,0) ko (16,4,1) block mein flatten karo. 0 + 2 × 16 + 3 = 35
⌊ x ⌋ ⌈ x ⌉ Floor round down karta hai; ceiling round up karta hai.
Kaun sa thread ka warp ID deta hai, aur kaun sa warp count? Floor warp ID deta hai; ceiling warps ki number deta hai.
136 threads ke block ke liye kitne warps? ⌈ 136/32 ⌉ = 5
"Lockstep / SIMT" ka matlab kya hai? Saare 32 threads same clock tick par same instruction chalate hain, har ek ke apne data ke saath.
Active mask kya hai? Ek 32-bit on/off switch jo decide karta hai ki kaun se threads current instruction par act karte hain.
Divergence slow kyun hai? Dono branches ek ke baad ek opposite masks ke saath chalte hain, kyunki ek warp ek time par sirf ek instruction chalata hai.
Stall kya hai, aur ise kaise chhupaate hain? Ek warp jo aage nahi badh sakta; ek doosra ready warp free mein swap in kiya jaata hai khali cycles fill karne ke liye.
Latency L W ⌈ L / W ⌉
Occupancy define karo. Active warps per SM divided by max warps per SM.