CUDA cores and execution model: Exercises — Hinglish deep dive

6.2.3 · D4 · Hardware › GPU Architecture › CUDA cores and execution model

Parent note: CUDA cores and execution model. Prerequisites jo tumhare paas open ho sakti hain: 6.2.01-GPU-vs-CPU-architecture, 6.2.02-GPU-memory-hierarchy, 6.1.03-SIMD-vs-SIMT, 6.2.04-occupancy-and-performance.

Shuru karne se pehle, vocabulary ke teen chhote pieces jo poora page inhi par tikaa hai, bilkul zero se banaye gaye hain:

Definition CUDA index words (

blockIdx.x, blockDim.x, threadIdx.x) Jab tum GPU par kaam launch karte ho toh threads ko equal-sized blocks (platoons) mein organize karte ho, aur saare blocks milke ek grid (poori army) banate hain. CUDA har running thread ko teen built-in numbers deta hai taaki woh pata kar sake "main kaun hoon?":

threadIdx.x = apne block ke andar meri seat number (0, 1, 2, … blockDim.x − 1 tak).
blockDim.x = har block mein kitne threads hain (platoon size — har block ke liye same).
blockIdx.x = main kis block mein hoon (kaun sa platoon, 0, 1, 2, … numbered).

.x ka matlab sirf "x-direction" hai — CUDA tumhe blocks/threads ko 3 dimensions tak count karne deta hai, lekin is page ke har problem mein 1D hai, isliye sirf .x chahiye. Agar poora build-up chahiye, parent note dekho: CUDA cores and execution model.

Level 1 — Recognition

Exercise 1.1. Ek warp mein kitne threads hote hain, aur unke liye "in lockstep" ka kya matlab hai?

Recall Solution 1.1

Ek warp mein 32 threads hote hain. "In lockstep" ka matlab hai ki saare 32 same cycle par same instruction execute karte hain — jaise 32 dancers jo kabhi bhi alag move nahi kar sakte, hamesha ek hi move karte hain. Woh sirf us data mein alag hote hain jo har ek apne registers mein rakhta hai, kabhi us instruction mein nahin jo ek given tick par run karte hain.

Exercise 1.2. Sach ya jhooth: ek CUDA core ek poora processor hai jaise CPU core, jiske paas apna branch predictor aur instruction decoder hota hai.

Recall Solution 1.2

Jhooth. Ek CUDA core sirf ek ALU (arithmetic logic unit) hai — ek chhota calculator jo ek cycle mein ek FP32 ya integer operation karta hai. Iske paas koi private decoder, branch predictor, ya independent instruction stream nahin hoti. Woh apne warp ke baaki 31 cores ke saath saara control logic share karta hai. Yahi sharing ki wajah se GPUs per-ALU saste hain lekin branchy code mein bure hain.

Exercise 1.3. Architecture hierarchy ke char levels ke naam batao, sabse chhote se sabse bade tak.

Recall Solution 1.3

CUDA core (ek scalar ALU, 1 op/cycle)
Warp (32 threads in lockstep)
Streaming Multiprocessor (SM) (kai cores + schedulers + registers + shared memory rakhta hai)
GPU chip (kai SMs rakhta hai)

Level 2 — Application

Exercise 2.1. Ek grid blockDim.x = 256 ke saath launch hua hai. Block blockIdx.x = 7 mein position threadIdx.x = 40 wale thread ka global thread ID compute karo.

Recall Solution 2.1

Indexing formula use karo: $Global thread ID = blockIdx.x \times blockDim.x + threadIdx.x .$ Har piece ka matlab: blockIdx.x hai kaun sa block (kaun sa platoon), blockDim.x hai threads per block (platoon size), aur threadIdx.x hai tumhara seat apne block ke andar. Block index ko block size se multiply karne par saare pehle platoons skip ho jaate hain; apna seat add karne par tum exactly land karte ho. $= 7 \times 256 + 40 = 1792 + 40 = 1832 .$

Exercise 2.2. Tumhare paas N = 1{,}000{,}000 elements hain aur tumne threadsPerBlock = 256 choose kiya. Standard ceiling launch formula use karke, kitne blocks launch karoge?

Recall Solution 2.2

Formula upar round karta hai (yahi floor bracket $⌊ \cdot ⌋$ hai jo ek number ko next whole value se aage nudge karne par apply hota hai) taaki koi element uncovered na rahe: $numBlocks = ⌊ \frac{N + threadsPerBlock - 1}{threadsPerBlock} ⌋ .$ + threadsPerBlock - 1 kyun? Plain integer division neeche round karta hai, jisse aakhri partly-full block drop ho jaata. Pehle ek-se-kam-ek-block add karna kisi bhi remainder ke hote hue round-up force karta hai — ye exactly disguise mein ceiling hai. $= ⌊ \frac{1 , 000 , 000 + 255}{256} ⌋ = ⌊ \frac{1 , 000 , 255}{256} ⌋ = 3907.$ Yeh $3907 \times 256 = 1, 000, 192$ threads launch karta hai — extra $192$ threads kernel mein if (i < N) guard se band ho jaate hain.

Exercise 2.3. Ek block mein 256 threads hain. Kitne warps hain? Agar block mein 250 threads hain toh?

Recall Solution 2.3

Warps 32 ke chunks mein bante hain, isliye thread count ko ceiling bracket se upar round karte hain: $warps = ⌈ \frac{threads}{32} ⌉ .$

$256/32 = 8$ warps exactly. Clean.
$250/32 = 7.8 \to ⌈ 7.8 ⌉ = 8$ warps. 8th warp mein sirf $250 - 7 \times 32 = 26$ threads hain; baaki $6$ lanes inactive hain lekin warp slot occupy karte hain. Dono cases mein poore warp ki keemat chukani padti hai.

Level 3 — Analysis

Exercise 3.1. Ek SM mein 128 CUDA cores hain. Ek student claim karta hai "32 threads ka ek warp isliye $32/128 = 1/4$ cycle mein khatam ho jaata hai." Error diagnose karo aur ek warp instruction ki sahi minimum latency batao.

Recall Solution 3.1

Sahi minimum 1 cycle hai, quarter cycle nahin.

Student galat kyun hai: ek warp instruction per issue atomic hoti hai — scheduler jo sabse chhota time slice deal mein leta hai woh ek cycle hai. 128 cores hone se woh cycle pieces mein nahin katata.

Figure mein kya observe karo: grid SM ke saare 128 cores ek single tick par dikhata hai. 32 magenta squares woh lanes hain jo is ek warp ko run kar rahi hain — woh us tick par ek saath light up hoti hain aur khatam hoti hain. 96 pale squares is warp ke liye idle hain, lekin orange label yaad dilata hai ki woh kisi doosre warp ko same tick par run karne ke liye free hain. Toh picture point ko visual banati hai: extra cores SM ko sideways wide karte hain (zyada warps ek saath), woh kabhi bhi is ek warp ki single tick ko chhoti nahin karte.

Extra cores ka asli faayda throughput hai, latency kam karna nahin: 128 cores aur multiple schedulers ke saath, SM same cycle mein kai warp-instructions issue kar sakta hai (different warps se), saare 128 cores ko busy rakhta hai — lekin har individual warp ko phir bhi kam se kam ek cycle lagti hai.

Exercise 3.2. Ek warp ke 32 threads mein se har ek ek float (4 bytes) read karta hai. Case A mein woh A[0..31] (consecutive) read karte hain. Case B mein thread $i$ A[i*32] (stride 32) read karta hai. 128-byte transactions use karke har ke liye transaction efficiency $η$ compute karo.

Recall Solution 3.2

Efficiency hai bytes jo tumhe chahiye ÷ bytes hardware ne move kiye: $η = \frac{useful bytes}{transactions \times 128} .$ Dono cases mein useful bytes $= 32 \times 4 = 128$ .

Case A (coalesced): saare 32 consecutive floats exactly $128$ bytes = ek aligned transaction span karte hain. $η_{A} = \frac{128}{1 \times 128} = 1 = 100%.$ Case B (stride 32): consecutive threads $32 \times 4 = 128$ bytes apart land karte hain, isliye har thread ka float apni khud ki 128-byte line mein hota hai → 32 alag transactions. $η_{B} = \frac{128}{32 \times 128} = \frac{1}{32} = 3.125%.$ Case B same useful data ke liye $32 \times$ zyada bytes move karta hai — yahi woh " $32 \times$ slower" hai jiske baare mein tumhe warn kiya gaya tha.

Figure mein kya observe karo: top row (magenta bracket) coalesced case hai — 32 requested floats side by side hain aur ek magenta box dikhata hai ki woh sab ek 128-byte transaction mein fit ho jaate hain, toh koi byte waste nahin. Bottom row (orange) strided case hai — sirf har 8th cell woh float hai jo hume actually chahiye (orange cells), aur har ek apne khud ke transaction box mein wrapped hai, toh hardware same 128 useful bytes deliver karne ke liye 32 full lines kheechta hai. Aankhein instantly dekh leti hain kyun top kuch waste nahin karta aur bottom almost sab waste karta hai.

Exercise 3.3. Registers per SM cap $M = 65, 536$ par hai. Tum $P = 2048$ resident threads chahte ho. Rough per-thread register ceiling kya hai? Agar tumhara kernel actually 64 registers/thread use karta hai toh occupancy ka kya hoga?

Recall Solution 3.3

Ceiling hai total budget evenly split: $\frac{M}{P} = \frac{65 , 536}{2048} = 32 registers/thread .$ Agar kernel ko sach mein 64 registers/thread chahiye, toh tum sirf $P^{'} = \frac{M}{64} = \frac{65 , 536}{64} = 1024 threads$ resident rakh sakte ho — aadhe threads. Kam resident threads matlab latency hide karne ke liye kam warps, toh occupancy drop hoti hai. (Yeh exactly woh hinge hai jo 6.2.04-occupancy-and-performance mein jaata hai.)

Level 4 — Synthesis

Exercise 4.1. Ek SM mein $S = 4$ warp schedulers hain. Har resident warp $C = 10$ cycles useful compute karta hai, phir global-memory read par $L = 200$ cycles stall karta hai. Latency poori tarah hide karne ke liye kitne resident warps $T_{required}$ chahiye? Agar SM 64 resident warps par cap kare, toh kya full hiding possible hai?

Recall Solution 4.1

Bookkeeping banao. Ek warp ke $L = 200$ stall cycles ke dauran, SM ke $S = 4$ schedulers doosre kaam ke $S \cdot L = 4 \times 200 = 800$ warp-instructions issue kar sakte hain. Har warp stall hone se pehle sirf $C = 10$ instructions supply karta hai, isliye pipeline fed rakhne ke liye needed warps ki sankhya hai: $T_{required} = \frac{S \cdot L}{C} = \frac{4 \times 200}{10} = 80 warps .$ SM 64 resident warps par cap karta hai. Kyunki $80 > 64$ , is kernel ke liye full latency hiding impossible hai is SM par. ALUs kuch time idle baithenge — kernel memory-bound hai.

Exercise 4.2. Same SM ke liye ( $S = 4$ , $L = 200$ , cap $= 64$ warps), minimum compute intensity $C$ kya hai jo full latency hiding possible banata hai?

Recall Solution 4.2

$T_{required} \leq 64$ set karo aur $C$ solve karo: $\frac{S \cdot L}{C} \leq 64 ⟹ C \geq \frac{S \cdot L}{64} = \frac{4 \times 200}{64} = \frac{800}{64} = 12.5.$ Kyunki $C$ poore instructions count karta hai, tumhe $C \geq 13$ cycles compute per stall chahiye. Design meaning: compute per memory op badhane se — kaam fuse karna, registers/shared memory se values reuse karna (dekho 6.2.02-GPU-memory-hierarchy), ya 7.3.01-parallel-programming-patterns mein tiling patterns — tum $T_{required}$ ko itna kam kar sakte ho ki 64 warps kaafi ho jaayein.

Exercise 4.3. Tumhe N = 1{,}000{,}000 floats ke do vectors add karne hain. threadsPerBlock aur numBlocks choose karo, batao warps per block kitne hain, aur explain karo kyun 256 ek safer default hai 1000 se.

Recall Solution 4.3

threadsPerBlock = 256 choose karo (32 ka multiple) taaki koi warp partly empty na ho: $256/32 = 8 full warps per block, koi wasted lanes nahin.$ Blocks needed (ceiling, kyunki ek partial block bhi count hoti hai): $numBlocks = ⌈ \frac{1 , 000 , 000}{256} ⌉ = 3907.$ 1000 threads/block kyun nahin? $1000$ 32 ka multiple nahin hai: $1000/32 = 31.25 \to ⌈ 31.25 ⌉ = 32$ warps, aur aakhri warp sirf $1000 - 31 \times 32 = 8$ active lanes run karta hai 32 mein se — har block mein 24 lanes waste. Saath hi 1000, 1024 hard cap ke paas hai, register pressure ke liye koi headroom nahin. 32 ka round multiple (128, 256, 512) warps full rakhta hai aur occupancy tuning flexible.

Level 5 — Mastery

Exercise 5.1. Ek colleague argue karta hai: "Agar main 32 cores wale SM se 128 cores wale SM par switch karun, toh mera 32 threads ka single warp $4 \times$ faster chalega." Unki reasoning ko steel-man karo, phir precisely refute karo, aur batao kya cheez sach mein $4 \times$ better hoti hai.

Recall Solution 5.1

Steel-man (sabse strong honest version): zyada ALUs zyaadatar zyada kaam per unit time karte hain, aur $128/32 = 4$ , toh $4 \times$ speedup table par hona lagta hai.

Refutation: $4 \times$ real hai lekin yeh throughput par apply hota hai, na ki ek single warp ki latency par. Ek warp instruction per issue atomic hoti hai — uska floor ek cycle hai chahe 32 ya 128 cores hon. 32-core SM par, ek warp/cycle cores ko fully occupy karta hai. 128-core SM par, ek warp phir bhi ek cycle mein khatam hota hai, lekin ab (multiple schedulers ke zariye) 4 warps tak ek cycle mein issue ho sakte hain. Toh 128-core SM har cycle mein $4 \times$ zyada warps complete karta hai — lekin koi bhi ek warp exactly utna hi fast hai jitna pehle tha.

Verdict: aggregate warp throughput ~ $4 \times$ scale karta hai; per-warp latency unchanged hai. Yeh fark hai ek race faster khatam karne (latency) aur char races side by side chalane (throughput) mein.

Exercise 5.2. Ek kernel 1000 GB/s peak DRAM bandwidth wale GPU par transaction efficiency $η = 25%$ achieve karta hai. (a) Tumhe actually kitni effective bandwidth milti hai? (b) Ek rewrite accesses ko fully coalesced banata hai ( $η = 100%$ ). Effective bandwidth kitne factor se improve hoti hai, aur kya compute per element change hota hai?

Recall Solution 5.2

(a) Effective bandwidth efficiency times peak hai: $eff BW = η \times peak = 0.25 \times 1000 = 250 GB/s .$ (b) $η = 100%$ par: $eff BW = 1.00 \times 1000 = 1000 GB/s, factor = \frac{1000}{250} = 4 \times .$ Compute per element unchanged hai — coalescing sirf bytes kaise move hote hain fix karta hai, na ki har thread kitne arithmetic ops karta hai. Tumne memory ko $4 \times$ better banaya; ALUs exactly same math karte hain. Isliye memory-bound kernels sirf ek pure access-pattern rewrite se speed mein leap le sakte hain, koi algorithm change nahin. (Seedha relevant hai 9.1.02-neural-network-training-on-GPUs se, jahan huge tensor reads dominate karte hain.)

Exercise 5.3. Prove karo ki "compute intensity $C$ badhana hamesha ek fixed latency hide karne ke liye required warps ghatata hai," aur woh ek degenerate case identify karo jahan formula break karta hai.

Recall Solution 5.3

Shuru karo: $T_{required} (C) = \frac{S \cdot L}{C} .$ Yahaan $S$ (schedulers) aur $L$ (stall cycles) fixed positive numbers hain, aur $C > 0$ . $C$ ki function ke roop mein, yeh $constant / C$ hai, jo $C > 0$ ke liye strictly decreasing hai: $C$ double karne se $T_{required}$ half ho jaata hai. Toh memory op ke beech zyada compute hamesha needed warps ghatata hai. $■$

Degenerate case: $C = 0$ . Agar ek warp stalls ke beech zero useful compute karta hai, toh $T_{required} = S \cdot L /0$ undefined hai (division by zero) — physically, ek warp jo kabhi compute nahin karta woh kuch bhi hide karne mein help nahin kar sakta; tum chahe kitne bhi warps launch karo pure stalls se latency cover nahin ho sakti. Formula assume karta hai $C \geq 1$ (memory ops ke beech kam se kam ek real instruction).