Warps and warp scheduling: Exercises — Hinglish deep dive

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

Yeh page ek self-test ladder hai. Har rung pehle se mushkil hai. Solution ko cover karo, problem try karo, phir reveal karo. Yahan jo bhi number compute karo ge, woh machine-checked hain.

Shuru karne se pehle, yeh rahe sirf woh words jo tumhe chahiye — sab Warps and warp scheduling se liye gaye hain:

Do arithmetic tools jinpar hum rely karte hain:

Level 1 — Recognition

Exercise 1.1 (L1)

Ek block blockDim = (32, 1, 1) ke saath launch hota hai. Kitne warps use honge, aur kya woh full hain?

Recall Solution

Hum kya karte hain: warp-count formula apply karo. Threads $= 32 \times 1 \times 1 = 32$ . $warps = ⌈ \frac{32}{32} ⌉ = 1$ Yeh kaisa dikhta hai: ek full "32 seats ki row", har seat occupied. Answer: 1 warp, fully occupied (32/32 active).

Exercise 1.2 (L1)

Ek block blockDim = (48, 1, 1) ke saath launch hota hai. Kitne warps? Kitne thread slots waste hote hain?

Recall Solution

Threads $= 48$ . $warps = ⌈ \frac{48}{48 \to 48/32 = 1.5} ⌉ = ⌈ 1.5 ⌉ = 2$ Do warps $2 \times 32 = 64$ slots dete hain lekin sirf 48 real threads hain. Wasted slots $= 64 - 48 = 16$ . Answer: 2 warps, 16 wasted slots (warp 1 runs 32/32, warp... ruko — recount karo): warp 0 = threads 0–31 (full), warp 1 = threads 32–47 (16 active, 16 idle). Toh sara waste warp 1 mein hai.

Exercise 1.3 (L1)

True ya false: ek SM apne execution units par individual threads schedule karta hai.

Recall Solution

False. Scheduling ka quantum warp hai. Ek warp ke saare 32 threads lockstep mein ek saath issue hote hain; scheduler kabhi ek akela thread nahi pick karta. (Dekhein SIMT-vs-SIMD — kyun SIMT per-thread state rakhta hai lekin phir bhi per-warp issue karta hai.)

Level 2 — Application

Exercise 2.1 (L2)

blockDim = (16, 4, 1). (a) Kitne warps? (b) Thread $(t_{x}, t_{y}, t_{z}) = (3, 2, 0)$ kis warp mein jaata hai?

Recall Solution

(a) Threads $= 16 \times 4 \times 1 = 64 \Rightarrow ⌈ 64/32 ⌉ = 2$ warps, dono full. (b) Linear index $D_{x} = 16, D_{y} = 4$ ke saath: $i = 0 \cdot (16 \cdot 4) + 2 \cdot 16 + 3 = 32 + 3 = 35$ Warp $= ⌊ 35/32 ⌋ = 1$ . Answer: 2 warps; thread (3,2,0) warp 1 mein hai.

Exercise 2.2 (L2)

Ek SM zyada se zyada 64 warps hold kar sakta hai. Tum abhi 48 active warps chala rahe ho. Occupancy kitni hai (percentage mein)?

Recall Solution

$Occupancy = \frac{48}{64} = 0.75 = 75%$ Answer: 75% occupancy. Dekhein Occupancy-vs-Performance — 75% aksar kaafi hoti hai; 100% chase karna ulta pad sakta hai agar usse register spills force hon.

Exercise 2.3 (L2)

Global memory latency $L = 400$ cycles hai. Har warp memory operations ke beech $W = 8$ independent instructions karta hai. Latency fully hide karne ke liye roughly kitne warps resident hone chahiye (single scheduler)?

Recall Solution

Yeh formula kyun: jab ek warp $L$ cycles wait karta hai, doosre warps ko scheduler ko feed karte rehna chahiye, har cycle ek instruction. Har warp $W$ useful instructions offer karta hai phir woh bhi stall ho jaata hai. $N_{warps} = ⌈ \frac{L}{W} ⌉ = ⌈ \frac{400}{8} ⌉ = 50$ Answer: 50 warps 400-cycle latency hide karne ke liye.

Level 3 — Analysis

Exercise 3.1 (L3)

blockDim = (17, 8, 1) (136 threads). Kitne warps, aur last warp mein exactly kitne active threads hain?

Recall Solution

Threads $= 17 \times 8 = 136 \Rightarrow ⌈ 136/32 ⌉ = ⌈ 4.25 ⌉ = 5$ warps. Pehle 4 warps indices 0–127 use karte hain (128 threads, sab full). Last warp = indices 128–135. Last warp mein active $= 136 - 128 = 8$ . Answer: 5 warps; last warp sirf 8/32 threads run karta hai (24 idle lanes → us warp ka 75% waste). Launch ko (16,8,1) = 128 threads = exactly 4 full warps kar ke fix karo.

Exercise 3.2 (L3)

Ek warp if (threadIdx.x % 2 == 0) A(); else B(); hit karta hai. Branch A 30 cycles leta hai, B 20 cycles. Divergent execution time kya hai, aur sirf A ke no-divergence cost se slowdown kitna hai?

Recall Solution

Kyunki odd/even lanes split ho jaate hain, warp mein dono branches present hain → woh masks ke saath sequentially chalte hain: $T_{diverged} = T_{A} + T_{B} = 30 + 20 = 50 cycles$ No-divergence baseline (sab lanes A lein) $= 30$ cycles. $Slowdown = \frac{50}{30} \approx 1.67 \times$ Answer: 50 cycles, ≈1.67× slower. Dekhein Branch-Divergence-Patterns: % 2 predicate sabse bura type hai — yeh guarantee karta hai ki har warp diverge kare.

Exercise 3.3 (L3)

3.2 jaisa hi code lekin predicate hai threadIdx.x < 32 ek 32-thread warp ke andar jahan har thread ka threadIdx.x 0–31 hai. Kya yeh warp diverge karta hai? Uska time kya hai?

Recall Solution

32 mein se har ek thread < 32 satisfy karta hai, toh sab branch A lete hain, koi B nahi leta. Koi divergence nahi → $T = T_{A} = 30$ cycles. Answer: Koi divergence nahi; 30 cycles. Lesson: divergence depend karta hai ki ek warp ke lanes agree karte hain ya nahi, na ki source mein branch exist karne par.

Level 4 — Synthesis

Exercise 4.1 (L4)

Ek SM mein 4 warp schedulers hain, har ek 1 instruction/cycle issue karta hai ( $T_{issue} = 1$ ). Average stall latency $L = 300$ cycles. Chaalon scheduler sab busy rakhne ke liye kitne warps resident hone chahiye, yeh maan ke har warp stall se pehle 1 instruction deta hai (worst case)?

Recall Solution

Har scheduler ko $L$ cycles ke liye har cycle ek fresh ready warp chahiye: $warps \geq N \times \frac{L}{T _{issue}} = 4 \times \frac{300}{1} = 1200$ Answer: 1200 warps jitna independent work. Kyunki ek real SM ~64 warps tak limited hai, yeh 1-instruction warps se impossible hai — tumhe har warp ke andar instruction-level parallelism bhi chahiye (stall se pehle har warp ke andar zyada independent instructions).

Exercise 4.2 (L4)

4.1 continue karo, lekin ab har warp stall se pehle $W = 20$ independent instructions karta hai. Kitne warps (instructions nahi) chahiye? Kya yeh 64-warp cap ke saath achievable hai?

Recall Solution

$warps = \frac{N \cdot L / T _{issue}}{W} = \frac{1200}{20} = 60$ $60 \leq 64$ , toh haan, achievable hai. Answer: 60 warps, 64-warp limit ke andar achievable (occupancy $= 60/64 \approx 93.8%$ ).

Exercise 4.3 (L4)

Tumhare paas ek kernel hai jisko 32 registers/thread chahiye. SM ke paas 65536 total registers hain aur maximum 2048 resident threads. Kaunsa limit — registers ya thread cap — occupancy decide karta hai, aur max resident warp count kya hai?

Recall Solution

Register-limited thread count: $65536/32 = 2048$ threads. Hardware thread cap: 2048 threads. Dono exactly tie karte hain → limit dono mein se koi bhi hai (dono 2048 dete hain). Warps $= 2048/32 = 64$ → full occupancy. Answer: ==Dono limits 2048 threads = 64 warps = 100% occupancy dete hain==. Agar registers 40/thread ho jaayein, toh $65536/40 = 1638$ threads $\Rightarrow ⌊ 1638/32 ⌋ = 51$ warps → register-limited. (Dekhein Register-Pressure.)

Level 5 — Mastery

Exercise 5.1 (L5)

Ek grid ek $100 \times 100$ image process karta hai blockDim = (16, 16, 1) ke saath. (a) Threads per block? (b) Warps per block? (c) Wasted slots per block? (d) Agar grid $⌈ 100/16 ⌉^{2}$ blocks ka hai, total blocks kitne?

Recall Solution

(a) $16 \times 16 = 256$ threads. (b) $⌈ 256/32 ⌉ = 8$ warps, sab full (256 evenly divide hota hai 32 se). (c) $8 \times 32 - 256 = 0$ wasted slots — perfect, kyunki 16×16 ek multiple hai 32 ka. (d) Blocks per dimension $= ⌈ 100/16 ⌉ = ⌈ 6.25 ⌉ = 7$ . Total $= 7 \times 7 = 49$ blocks. Answer: 256 threads, 8 full warps, 0 waste, 49 blocks. (Grid 100×100 se bahar edge threads waste karta hai, lekin har block internally warp-clean hai — bounds if se handle hota hai, jo 5.2 mein dekhein.)

Exercise 5.2 (L5)

5.1 ke edge blocks if (x < 100 && y < 100) se guard karte hain. Ek warp jo boundary straddle karta hai — maano 20 lanes image ke andar aur 12 bahar — uska divergence cost kya hai agar "inside" work $T = 40$ cycles ka hai aur "outside" lanes kuch nahi karte (0 cycles)?

Recall Solution

Do groups: 20 lanes body run karte hain ( $T = 40$ ), 12 lanes empty else lete hain (0). Woh phir bhi serialize karte hain, lekin empty branch ki cost 0 hai: $T_{diverged} = 40 + 0 = 40 cycles$ Baseline (koi divergence nahi, sab 32 body karte hain) bhi 40 hota. Slowdown $= 1 \times$ — divergence yahan free hai kyunki disabled lanes ka koi kaam tha hi nahi. Answer: 40 cycles, koi real penalty nahi. Masked-off lanes single 40-cycle body pass ke dauran "free mein" idle rehte hain. Divergence sirf tab cost karta hai jab dono paths mein real kaam ho.

Exercise 5.3 (L5, Synthesis + Analysis)

Kernel A: har warp 50% diverge karta hai (warp ka aadha $T_{A} = 100$ costs, aadha $T_{B} = 100$ costs). Kernel B: koi divergence nahi, lekin sirf 32 resident warps hain Kernel A ke 64 ke mukable. Memory latency $L = 400$ , warps ko $W = 25$ ke saath latency hide karne ke liye $\geq ⌈ L / W ⌉$ chahiye. Kaun sa kernel latency zyada aasaani se hide karta hai, aur Kernel A ka per-warp divergence slowdown kya hai?

Recall Solution

Latency-hiding need: $⌈ 400/25 ⌉ = 16$ warps.

Kernel A: 64 warps $\geq 16$ ✓ latency hide karta hai.
Kernel B: 32 warps $\geq 16$ ✓ yeh bhi latency hide karta hai. Dono bar clear karte hain, toh latency yahan deciding factor nahi hai. Kernel A divergence slowdown per warp: $(T_{A} + T_{B}) / T_{A} = (100 + 100) /100 = 2 \times$ . Answer: Dono latency hide karte hain; lekin Kernel A har warp par 2× divergence tax bharta hai. Jab latency dono taraf cover ho, Kernel B jeet ta hai — uske warps divergent region mein 2× faster hain. Lesson: ek baar latency hide ho jaaye, toh divergence fix karna zyada warps pile karne se better hota hai.