Warps and warp scheduling: Worked examples — Hinglish deep dive

6.2.5 · D3 · Hardware › GPU Architecture › Warps and warp scheduling

Ye parent topic ka "roll up your sleeves" wala page hai. Hum yahan theory dobara nahi padhayenge — instead hum har tarah ki situation ko dhundhte hain jo parent note mein ho sakti hai, aur har ek ko number tak le jaate hain. Agar koi term unfamiliar lage, parent note mein uski neenv rakhi gayi hai; yahan hum usse use karte hain.

Kuch bhi start karne se pehle, teen numbers ka ek plain-language recap — kyunki neeche ka har example inhi se bana hai:

Recall Teen quantities jo har example mein kaam aati hain

Warp size ::: 32 threads (NVIDIA hardware par fixed). Ek warp 32 threads ka ek bundle hota hai jo ek instruction fetch share karte hain.
Linear thread index ::: ek 3D (x,y,z) thread coordinate ko ek single counting number mein flatten karo, x sabse fast chalega.
Warp ID ::: woh linear index 32 se divide karke, neeche ki taraf round karo. Toh threads 0–31 → warp 0, 32–63 → warp 1, aur aise hi aage.

Scenario matrix

Is table ko ek checklist ki tarah socho. Har cell ek class of problem hai. Neeche, har worked example par un cell(s) ki stamp hai jo woh clear karta hai. Aakhir mein, koi bhi cell chhuti nahi — exam mein koi aisa warp scenario nahi milega jise humne skip kiya ho.

#	Case class	Isme kya tricky hai	Covered by
A	Clean 1D block, multiple of 32	"No waste" baseline	Ex 1
B	Block not a multiple of 32	Tail mein partial (wasted) warp	Ex 2
C	2D block, row-major flatten	Warps rows ko kaat ke jaate hain	Ex 3
C2	True 3D block, non-unit depth	$t_{z}$ term actually fire karta hai	Ex 3b
D	Degenerate: block 32 se chhota	Ek warp, mostly idle lanes	Ex 4
E	Latency hiding — enough warps?	$⌈ L / W ⌉$ warps chahiye	Ex 5
F	Limiting case: loads ke beech zero compute	Koi hiding possible nahi — memory bound	Ex 6
G	Warp divergence, dono branches li gayi	Serialized paths, slowdown factor	Ex 7
H	Divergence edge: sabhi threads same branch	Zero penalty (real divergence nahi)	Ex 8
I	Real-world word problem	Waste bachane ke liye block dims chunna	Ex 9
J	Exam twist: occupancy vs. warps-needed	Hardware cap ideal number ko beat karta hai	Ex 10

Group 1 — Threads kahan jaate hain? (cells A, B, C, C2, D)

Charon ke peeche ek hi formula:

Recall Total threads vs. total warps — rounding rule precisely samjho

Maano $P = D_{x} D_{y} D_{z}$ total thread count hai. Do baatein kabhi confuse mat karna:

Sabse bada warp ID hai $⌊(P - 1) /32 ⌋$ (aakhri linear index ka floor, jo ki $P - 1$ hai kyunki indices 0 se count hote hain). Warp IDs chalte hain $0, 1, \dots, ⌊(P - 1) /32 ⌋$ .
Warps ki sankhya isliye hai $⌊(P - 1) /32 ⌋ + 1$ , aur yeh $⌈ P /32 ⌉$ (total ka ceiling) ke barabar hai.

Ye dono kyun agree karte hain? Likho $P = 32 q + r$ , jahan $q$ 32 ke complete groups ki whole number hai aur $r$ remainder hai ( $0 \leq r \leq 31$ ). Do cases:

Agar $r = 0$ : block 32 ka exact multiple hai. Aakhri index hai $P - 1 = 32 q - 1$ , toh $⌊(P - 1) /32 ⌋ = q - 1$ , jisse milte hain $q - 1 + 1 = q$ warps. Aur $⌈ P /32 ⌉ = ⌈ q ⌉ = q$ . Same.
Agar $r > 0$ : aakhri index hai $P - 1 = 32 q + (r - 1)$ jahan $0 \leq r - 1 \leq 30$ , toh $⌊(P - 1) /32 ⌋ = q$ , jisse milte hain $q + 1$ warps. Aur $⌈ P /32 ⌉ = q + 1$ (fraction ko upar round karo). Same.

Toh dono cases mein dono sides identical value par aate hain — yahi wajah hai ki "aakhri index ko floor karo, phir ek add karo" (jo top warp ka ID batata hai) "total ko ceil karo" (jo batata hai kitne warps hain) ke barabar hota hai.

Ek-line takeaway ::: $⌊(P - 1) /32 ⌋ + 1 = ⌈ P /32 ⌉$ har positive integer $P$ ke liye, kyunki $P = 32 q + r$ split karne par dono sides $q$ (agar $r = 0$ ) ya $q + 1$ (agar $r > 0$ ) ke barabar hoti hain.

Example 1 — Cell A: clean 1D block

Worked example Baseline:

blockDim = (256, 1, 1) Kitne warps hain, aur kya sab full hain?

Forecast: padhne se pehle warp count guess karo. (Hint: divide karo.)

Flatten karo. $D_{y} = D_{z} = 1$ ke saath, formula $linear = t_{x}$ tak collapse ho jaata hai. Threads bas $0, 1, \dots, 255$ hain. Yeh step kyun? 1D block ko koi cross-row bookkeeping ki zarurat nahi — linear index hi x-index hai.
32 se divide karo. $256/32 = 8$ exactly. Yeh step kyun? Warp count hai $⌈ 256/32 ⌉ = 8$ ; equivalently top warp ID hai $⌊ 255/32 ⌋ = 7$ , aur $7 + 1 = 8$ warps.
Fullness check karo. Har warp exactly 32 consecutive indices cover karta hai bina kisi leftover ke, kyunki 256, 32 ka multiple hai. Yeh step kyun? Waste tab hi aata hai jab count 32 se divisible nahi hoti.

Answer: 8 warps, sab 32/32 active, 0 wasted lanes.

Verify: $8 \times 32 = 256$ = block size. Koi thread unaccounted nahi, koi lane idle nahi. ✓

Example 2 — Cell B: 32 ka multiple nahi

Worked example Tail warp:

blockDim = (100, 1, 1) Kitne warps hain, aur kitne lanes waste hote hain?

Forecast: kya 3 warps hain ya 4? Aur waste kya hai?

Do tarike, ek answer. Total threads $P = 100$ . Warp count $= ⌈ 100/32 ⌉ = ⌈ 3.125 ⌉ = 4$ . Cross-check aakhri index se: $⌊ 99/32 ⌋ = 3$ top warp ID hai, toh warps hain 0,1,2,3 → $3 + 1 = 4$ warps. Yeh step kyun? Upar wala recall box prove karta hai ki ceil-the-total aur floor-the-last-index-plus-one dono ek hi 4 dete hain. Dono raaste round up karte hain, kyunki ek bhi leftover thread ko poora warp chahiye.
Tail dhundho. Warps 0,1,2 mein indices 0–95 hain (woh hai $3 \times 32 = 96$ threads, sab full). Warp 3 mein indices 96–99 hain → sirf 4 active threads. Yeh step kyun? Threads warps mein order mein bhrte hain; aakhri full group ke baad jo bacha woh partial warp banata hai.
Waste count karo. Warp 3 mein $32 - 4 = 28$ idle lanes hain. Yeh step kyun? SM warp 3 ke liye ek full warp ka scheduling slot reserve karta hai chahe 28 lanes kuch nahi karte — yahi cost hai 32 ka multiple na hone ki.

Answer: 4 warps; aakhri wala 4/32 full; 28 wasted lanes.

Verify: $3 \times 32 + 4 = 100$ ✓. Tail warp ka waste fraction $= 28/32 = 87.5%$ .

Example 3 — Cell C: 2D block, warps rows ko kaat ke jaate hain

Worked example Row-major reality:

blockDim = (48, 4, 1) Warp 1 mein kaun se threads aate hain?

Forecast: guess karo ki warp 1 ek hi row mein rehta hai ya nahi.

Total & flatten rule. $48 \times 4 = 192$ threads. $D_{z} = 1$ ke saath $t_{z}$ term zero hai, toh linear index $= t_{y} \cdot 48 + t_{x}$ . Yeh step kyun? Row $t_{y} = 0$ linear 0–47 occupy karta hai, row $t_{y} = 1$ 48–95, aur aise aage. Rows 48 wide hain, 32 nahi — toh warp boundaries row boundaries se align nahi hongi.
Warp 1 = linear 32–63. Linear 32–47 row 0 se aate hain ( $t_{x} = 32..47$ ). Linear 48–63 row 1 se aate hain ( $t_{x} = 0..15$ ). Yeh step kyun? Warp 1 row 0 aur row 1 ke beech ki seam par hai — yeh classic "warp cuts across rows" situation hai.
Warps count karo. $192/32 = 6$ exactly → 6 full warps, koi waste nahi (192, 32 ka multiple hai chahe rows nahi hain). Yeh step kyun? Waste total thread count vs. 32 par depend karta hai, row width par nahi.

Figure walkthrough: har colored cell ek thread hai, jaise block actually hai — 48 columns (x) across, 4 rows (y) down, x fastest chalta hai. Cells ko warp ID ke hisaab se tint kiya gaya hai, toh warp 0 ek color, warp 1 doosra, etc. White mein outlined single warp, warp 1 hai: dhyan do ki woh white group top row ke beech mein shuru hota hai (x=32–47, yellow arrow), phir neeche jump karta hai aur doosri row ke start par continue karta hai (x=0–15). Yeh visual "wrap-around" poora point hai — kyunki ek row 48 wide hai lekin warp sirf 32, warp boundaries rows ke through slice karte hain unse align hone ki jagah.

Answer: warp 1 = row 0 ka tail (x=32–47) plus row 1 ka head (x=0–15); kul 6 warps, 0 waste.

Verify: $6 \times 32 = 192 = 48 \times 4$ ✓. Warp 1 spans linear 32–63 = row 0 se 16 threads + row 1 se 16 = 32 ✓.

Example 3b — Cell C2: ek sachi 3D block (depth actually matters)

Worked example

$t_{z}$ term fire karta hai: blockDim = (8, 4, 3) Linear index 70 wala thread kis slice aur row mein hai, aur kitne warps hain?

Forecast: guess karo ki index 70 pehle, doosre, ya teesre z-slice mein hai.

Total threads. $P = D_{x} D_{y} D_{z} = 8 \times 4 \times 3 = 96$ . Yeh hamaara pehla block hai jahan $D_{z} \neq = 1$ , toh $t_{z} (D_{x} D_{y})$ term finally kaam karta hai. Yeh step kyun? Ek z-slice mein $D_{x} D_{y} = 8 \times 4 = 32$ threads hote hain. Toh har naya z-slice 32 ka ek poora plane add karta hai — teen slices stack hain.
Linear 70 locate karo. Ek slice 32 threads ki hai, toh slice $t_{z} = ⌊ 70/32 ⌋ = 2$ (teesri slice, kyunki z 0 se count hota hai). Us slice ke andar offset hai $70 - 2 \times 32 = 6$ . Phir row $t_{y} = ⌊ 6/8 ⌋ = 0$ aur column $t_{x} = 6 - 0 \times 8 = 6$ . Yeh step kyun? Hum flatten formula ko invert kar rahe hain: sabse bada place-value ( $t_{z}$ ) pehle peel off karo, phir $t_{y}$ , phir $t_{x}$ — exactly mixed-radix number se digits padhne jaisa.
Warps count karo. $⌈ 96/32 ⌉ = 3$ warps, aur yahan beautifully har z-slice exactly ek warp hai (kyunki ek slice = 32 threads). Yeh step kyun? Jab $D_{x} D_{y}$ khud 32 ka multiple ho, warp boundaries cleanly slice boundaries par pad jaate hain — koi cross-slice straddling nahi.

Answer: linear 70 → thread coordinate $(t_{x}, t_{y}, t_{z}) = (6, 0, 2)$ , teesre z-slice mein; block mein 3 warps hain, ek per slice.

Verify: index rebuild karo: $t_{z} (D_{x} D_{y}) + t_{y} D_{x} + t_{x} = 2 \times 32 + 0 \times 8 + 6 = 70$ ✓. Index 70 ka warp $= ⌊ 70/32 ⌋ = 2$ , slice $t_{z} = 2$ se match karta hai ✓.

Example 4 — Cell D (degenerate): block ek warp se chhota

Group 2 — Kya scheduler busy rehega? (cells E, F, J)

Parent ka latency-hiding count yaad karo. Plain words mein: jab ek warp $L$ cycles memory ke liye wait karta hai, scheduler doosre warps chalata hai. Agar ek warp memory requests ke beech $W$ useful instructions karta hai, toh itne warps chahiye ki "baaki sab ka kaam" wait ko fill kare.

Example 5 — Cell E: latency hide karne ke liye kaafi warps

Example 6 — Cell F (limiting case): loads ke beech zero compute

Worked example Unhidable stall:

$L = 400$ , $W = 1$ Ek kernel jo back-to-back sirf loads stream karta hai. Ise hide karne ke liye kitne warps?

Forecast: kya yeh possible bhi hai?

Plug in karo. $N = ⌈ 400/1 ⌉ = 400$ warps. Yeh step kyun? Memory ops ke beech sirf 1 instruction ke saath, har warp 400-cycle wait ka sirf 1 cycle cover karta hai — toh 400 warps chahiye honge.
Reality check. Max hai 64 warps/SM ⇒ 400 ≫ 64 ⇒ hide nahi ho sakta. Kernel memory-bound hai. Yeh step kyun? Yeh formula ki degenerate limit hai: jaise $W \to 1$ , $N \to L$ , jo hardware ceiling ko cross kar jaata hai. Koi occupancy ek aisi kernel nahi bacha sakta jis mein overlap karne ke liye koi arithmetic hi nahi.

Answer: 400 warps chahiye; impossible → memory-bound, throughput bandwidth se limited hai, scheduling se nahi.

Verify: $⌈ 400/1 ⌉ = 400 > 64$ ✓. Dekho Occupancy-vs-Performance: ek point ke baad, zyada occupancy help karna band kar deti hai — yahi wajah hai.

Example 10 — Cell J (exam twist): ideal warps vs. hardware cap

Pehle, woh scheduler-level formula jo yeh example use karta hai. Parent note count per scheduler chalata hai, toh use karne se pehle uske symbols name karna zaroori hai:

Formula Scheduler-level warp demand

$Warps needed \geq N_{sched} \times \frac{L}{T _{issue}} / I_{warp}$ jahan (sab cycle-counts clock cycles mein hain, sab counts dimensionless whole numbers hain):

$N_{sched}$ = number of warp schedulers per SM (independent hardware units jo har cycle ek warp choose karte hain issue karne ke liye; modern SMs mein 4 hote hain). Unit: schedulers (ek plain count).
$L$ = memory stall latency. Unit: cycles.
$T_{issue}$ = cycles between successive issues from one scheduler ( $T_{issue} = 1$ matlab ek scheduler har cycle ek instruction fire kar sakta hai). Unit: cycles per instruction.
$I_{warp}$ = independent instructions jo ek warp khud stall hone se pehle issue kar sakta hai. Unit: instructions per warp.

Units ke saath padhne par: $N_{sched} \cdot L / T_{issue}$ ki units hain $schedulers \times cycles / (cycles/instruction) = instruction-slots$ — instruction slots ki total sankhya jo ek stall ke dauran filled rehni chahiye. $I_{warp}$ (instructions per warp) se divide karne par "instructions" cancel ho jaate hain, bacha warps.

Worked example Trap:

$N_{sched} = 4$ , $L = 400$ cyc, $T_{issue} = 1$ cyc/inst, $I_{warp} = 20$ inst/warp; hardware max = 64 warps/SM Kya SM fully fed ho sakta hai?

Forecast: kya 64 warps kaafi hain? (Exam chahta hai ki tum haan kaho.)

Instruction budget. $N_{sched} \cdot L / T_{issue} = 4 \times 400/1 = 1600$ instruction slots per SM "in flight" rehne chahiye. Yeh step kyun? Char schedulers har cycle fire karte hain 400 cycles ke liye = $4 \times 400 = 1600$ slots fill karne hain stall rehne tak.
Warps mein convert karo. $I_{warp} = 20$ se divide karo: $1600/20 = 80$ warps. Yeh step kyun? Warps woh unit hai jise tum actually schedule kar sakte ho; yeh instruction budget ko warp count mein convert karta hai upar dikhaye unit cancellation ke zariye.
Cap par lag gaye. Hardware max 64 warps/SM allow karta hai, lekin chahiye 80. Kyunki $80 > 64$ , warp-level parallelism akele 16 warps se short padta hai. Yeh step kyun? Yahi twist hai: ideal answer (80) SM physically hold kar sakne se (64) zyada hai. Hardware cap, formula nahi, binding limit hai — toh tumhe instruction-level parallelism bhi exploit karna hoga har warp ke andar (ek warp ke andar independent instructions back-to-back issue karna) deficit pura karne ke liye. Zyada ILP registers cost karti hai — dekho Register-Pressure.

Answer: ideally 80 warps chahiye, lekin cap hai 64 → 16 short; gap ILP se close karo, sirf occupancy se nahi.

Verify: $4 \cdot 400/1 = 1600$ ; $1600/20 = 80$ ; $80 > 64$ (deficit $= 16$ ) ✓.

Group 3 — Divergence cost (cells G, H)

Jab ek warp ke threads mein disagree ho ki kaun sa branch lena hai, warp dono branches ek ke baad ek chalata hai, har baar un lanes ko mask kar ke jo participate nahi karni chahiye. Pattern-level tricks ke liye dekho Branch-Divergence-Patterns.

Example 7 — Cell G: dono branches populated

Worked example Split warp: 20 lanes

if leti hain ( $T_{if} = 30$ cyc), 12 else ( $T_{else} = 30$ cyc) Ek undivided warp ke comparison mein slowdown kya hai?

Forecast: kya 20-vs-12 split answer change karta hai? Haan ya nahi guess karo.

Divergent time. Dono branches chalte hain: $T_{diverged} = 30 + 30 = 60$ cycles. Yeh step kyun? Warp if chalata hai 12 lanes masked off ke saath, phir else 20 lanes masked off ke saath — sequentially. Lane count per branch time nahi change karta; masked lanes abhi bhi same cycles tick karte hain.
Baseline. Koi divergence nahi hoti toh ek branch cost hoti: $30$ cycles. Yeh step kyun? Agar sab 32 lanes agree kar len, sirf ek path chalta hai.
Slowdown. $60/30 = 2 \times$ . Yeh step kyun? Yeh classic worst case hai: equal-cost branches ⇒ exactly double time, split ratio chahe kuch bhi ho.

Figure walkthrough: top blue bar baseline hai — agar sab 32 lanes agree kar len, sirf if block chalta hai, 30 cycles, khatam. Neeche ki row dikhati hai ki divergence actually kya karta hai: pehle green if bar chalta hai (20 lanes live, 12 masked off aur idle), aur tabhi tak lal else bar chalta hai jab pehla khatam ho (12 lanes live, 20 masked off). Kyunki dono side-by-side ki jagah end-to-end chalte hain, yellow double-arrow poore 60 cycles span karta hai — visibly blue bar se double. Har colored bar mein masked lanes "rent pay kar rahe hain" (cycles occupy kar rahe hain) bina koi useful kaam kiye, yahi exactly divergence penalty hai.

Answer: 60 cycles, 2× slowdown — aur 31-vs-1 split mein bhi 2× hi hota.

Verify: $30 + 30 = 60$ ; $60/30 = 2.0$ ✓. Figure jo key insight dikhata hai: 12 idle if-lanes aur 20 idle else-lanes cycles waste karte hain lekin bachate koi nahi.

Example 8 — Cell H (edge case): sab threads agree karte hain

Worked example Real divergence nahi: sab 32 lanes

if leti hain ( $T_{if} = 30$ ), 0 else Penalty kya hai?

Forecast: kya else block kuch cost karta hai agar koi enter hi nahi karta?

Populated branches. $n_{else} = 0$ , toh else path skip ho jaata hai — else ke liye warp mask all-zero hai aur hardware use jump kar jaata hai. Yeh step kyun? Divergence sirf un branches ko serialize karta hai jinhe kam se kam ek lane leta hai. Ek empty branch free hai.
Time. $T = T_{if} = 30$ cycles. Slowdown $= 30/30 = 1 \times$ . Yeh step kyun? Yeh divergence formula ki boundary hai: jaise hi $n_{else} \to 0$ , $T_{diverged} \to T_{if}$ aur penalty gayab ho jaati hai.

Answer: 30 cycles, 1× (koi penalty nahi) — fully convergent warp.

Verify: $n_{else} = 0$ ke saath, $T = 30 + 0 \cdot 30 = 30$ ; $30/30 = 1.0$ ✓. Moral: branch decisions ko warp boundaries se align karna (sab 32 agree karte hain) penalty khatam kar deta hai.

Group 4 — Word problem (cell I)

Example 9 — Cell I: image kernel ke liye block dims chunna

Worked example Real world: ek

$1920 \times 1080$ image process karo, ek thread per pixel Tumhe 2D block shape chunni hai. Ek "balanced-looking" (30, 30) ko ek "aligned" (32, 8) se compare karo. Kaun sa kam lanes waste karta hai, aur har block mein kitne warps hain?

Forecast: square block dekhne mein efficient lagta hai — kya woh actually better hai? Compute karne se pehle guess karo.

Block (30,30). Threads $= 30 \times 30 = 900$ . Warp count $= ⌈ 900/32 ⌉ = ⌈ 28.125 ⌉ = 29$ . Aakhri (partial) warp mein $900 - 28 \times 32 = 900 - 896 = 4$ threads hain → us tail warp mein $32 - 4 = 28$ wasted lanes. Yeh step kyun? 900, 32 ka multiple nahi hai, toh hum 2D block ke andar Cell B mein aa jaate hain — ek partial tail warp appear hota hai, bilkul Example 2 ki tarah.
Block (32,8). Threads $= 32 \times 8 = 256$ . Warp count $= 256/32 = 8$ exactly, 0 wasted lanes. Yeh step kyun? blockDim.x ko warp width (32) ke barabar banana clean warp boundaries guarantee karta hai — row width equals warp width, toh warps rows ko straddle bhi nahi karte (Example 3 se unlike).
Decide karo. (32,8) kuch waste nahi karta aur warps ko rows se aligned rakhta hai, jo coalesced memory access mein bhi help karta hai (dekho GPU-Memory-Hierarchy). (30,30) 28 lanes waste karta hai aur rows misalign karta hai. Yeh step kyun? Kam wasted lanes + aligned memory ⇒ higher effective throughput. "Balanced-looking" square block woh trap hai jo question ne tumhare liye set kiya tha.

Answer: (32,8) chunno — 8 warps, 0 wasted lanes — (30,30) ki jagah, jisme 29 warps aur 28 wasted lanes hain.

Verify: $⌈ 900/32 ⌉ = 29$ aur $900 - 896 = 4$ active (28 idle) ✓; $256/32 = 8$ exact, 0 idle ✓. (Image bahut saare aisi blocks se tile hoti hai; yahan hum per-block efficiency compare kar rahe hain, jo determine karta hai ki har block lanes waste karta hai ya nahi.)

Recall Quick self-test

Ek (72,1,1) block: kitne warps aur kitne wasted lanes? ::: $⌈ 72/32 ⌉ = 3$ warps; tail warp mein $72 - 64 = 8$ active hain → 24 wasted lanes. Ek (8,4,3) block: kitne warps? ::: $⌈ 96/32 ⌉ = 3$ warps (ek per z-slice, kyunki har slice exactly 32 threads hai). $L = 300$ , $W = 15$ : hide karne ke liye kitne warps chahiye? ::: $⌈ 300/15 ⌉ = 20$ warps. 16 lanes if leti hain (cost 40), 16 else (cost 40): slowdown? ::: $(40 + 40) /40 = 2 \times$ .