CUDA cores and execution model: Worked examples — Hinglish deep dive

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

Yeh parent topic ka "haath se kaam karo" wala child note hai. Parent ne tumhe machinery di; yahan hum har tarah ka input us machinery se guzaarte hain — best case, worst case, zero, degenerate, ek real word problem, aur ek exam trap. Kuch bhi naya assume nahi kiya gaya: neeche use kiya gaya har symbol wahan pe re-anchor hota hai jahan woh pehli baar aata hai.

Definition Woh symbols jo hum har jagah reuse karte hain (pehle yeh padho)

Kisi bhi arithmetic se pehle — scenario table se bhi pehle — inhe simple words mein pakad lo:

$N$ = data elements ki sankhya jo tum process karna chahte ho (jaise array ki length). Ek plain count, koi units nahi.
$B$ = block size = threads per block (tum choose karte ho; max 1024).
Ek warp = 32 threads ka ek fixed group jise hardware lockstep mein run karta hai. Hamesha 32.
SM = Streaming Multiprocessor = GPU chip par ek self-contained "engine" jo CUDA cores ka ek batch, unke warp schedulers, aur ek shared register file ka maalik hota hai. Ek GPU mein bahut saare SMs hote hain; har ek kai warps run karta hai. Jab bhi tum padhte ho "SM caps at 64 warps", iska matlab hai ek aisa engine ek waqt mein sirf 64 warps resident rakh sakta hai.
$η$ (Greek letter "eta") = transaction efficiency = (bytes jo tumhe actually chahiye the) ÷ (bytes jo hardware ne physically move kiye). 0 aur 1 ke beech ek pure fraction. Yahi symbol neeche matrix ke Cases D aur E mein use hota hai.

Scenario matrix

Neeche har cell ek alag case class hai jo yeh topic tumpe throw kar sakta hai. Worked examples us cell ke saath labeled hain jise woh hit karte hain. (Is table mein har symbol — $N$ , $η$ , SM — upar wale block mein define kiya gaya tha.)

#	Case class	Woh extreme / corner jo test hota hai	Example
A	Perfect launch	threads evenly warps mein divide ho jaate hain, sab cores ko kaam milta hai	Ex 1
B	Ragged tail	`N` block size ka multiple nahi → ek partial final warp	Ex 2
C	Degenerate: N ≤ 0 ya N bahut chhota	ek warp se bhi kam threads; idle cores; negative/zero	Ex 3
D	Memory best case	coalesced access, efficiency $η = 1$	Ex 4
E	Memory worst case	strided access, efficiency $η$ minimum	Ex 4
F	Branch divergence	warp do paths mein split ho jaata hai, serialized	Ex 5
G	Latency hiding — succeeds	stall cover karne ke liye kaafi warps hain	Ex 6
H	Latency hiding — fails (limit)	required warps SM cap se zyada ho jaate hain	Ex 6
I	Real-world word problem	image processing, launch config choose karo	Ex 7
J	Exam twist	"128 cores ÷ 32 = ¼ cycle" trap + register ceiling	Ex 8

Agar koi bhi row shaky lage to prerequisites: 6.1.03-SIMD-vs-SIMT, 6.2.02-GPU-memory-hierarchy, aur baad mein 6.2.04-occupancy-and-performance.

Ex 1 — Cell A: perfect launch

Worked example Perfect launch

Tum $N = 2048$ floats process karte ho block size $B = 256$ ke saath. Kitne blocks, total kitne warps, aur kya koi threads waste hue?

Forecast: aage padhne se pehle warps ki sankhya guess karo. (Hint: 2048 ÷ 32.)

Step 1 — Blocks ki zaroorat. $numBlocks = ⌈ \frac{N}{B} ⌉ = ⌈ \frac{2048}{256} ⌉ = 8.$ Yeh step kyun? Tum sirf whole blocks launch karte ho; ceiling $⌈ \cdot ⌉$ ("round up") guarantee karta hai ki har data element ke liye kam se kam ek thread ho. Yahan yeh evenly divide ho jaata hai isliye koi rounding nahi hoti.

Step 2 — Warps per block. $warps/block = \frac{B}{32} = \frac{256}{32} = 8.$ Yeh step kyun? Ek block 32 ke warps mein kata jaata hai. $256/32$ ek whole number hai, isliye koi partial warp nahi.

Step 3 — Total warps. $8 blocks \times 8 warps/block = 64 warps .$ Yeh step kyun? Total warps = poori grid kitne lockstep groups produce karti hai.

Verify: $64 \times 32 = 2048$ threads = exactly $N$ . Zero wasted threads. Har core jo run karta hai use real kaam milta hai. ✅

Yeh "textbook happy path" hai — sab kuch divide ho jaata hai. Har baad wala example is se ek deliberate break hai.

Ex 2 — Cell B: ragged tail

Worked example Ragged tail (N block size ka multiple nahi)

Ab $N = 1000$ same $B = 256$ ke saath. Kitne blocks? Kitne idle threads exist karte hain, aur kahan?

Forecast: 4 se zyada ya kam blocks? Aur kya last warp full hai?

Step 1 — Blocks (ceiling ab real kaam kar raha hai). $numBlocks = ⌈ \frac{1000}{256} ⌉ = ⌈ 3.906 ⌉ = 4.$ Yeh step kyun? $3$ blocks sirf $768$ threads cover karte hain $< 1000$ . Baaki $232$ elements ke liye tumhe 4th block chahiye. Ceiling exactly woh "hamesha round up karo taaki koi chhoot na jaaye" rule hai.

Step 2 — Total threads launched vs needed. $4 \times 256 = 1024 threads launched, 1024 - 1000 = 24 idle threads .$ Yeh step kyun? Un 24 extra threads ka global index $i \geq N$ hai; kernel mein if (i < N) guard unhe kuch bhi nahi karne deta.

Step 3 — Kaun sa warp ragged hai? (index by index build kiya gaya). Pehle last block locate karo. Block indices $0, 1, 2, 3$ run karte hain; 4th block (blockIdx $= 3$ ) global thread indices $3 \times 256 = 768$ se lekar $768 + 256 - 1 = 1023$ tak ka maalik hai. Ab us block ko local index se 32 ke warps mein split karo. Local index $0$ global $768$ hai; local index $31$ global $799$ hai; aur aise hi aage. Block ke andar warp number $w$ local indices $32 w$ se $32 w + 31$ tak cover karta hai. Last active global index $N - 1 = 999$ hai. Block 3 ke andar uska local index $999 - 768 = 231$ hai. Local index $231$ kis warp mein hai? Divide karo: $⌊ 231/32 ⌋ = 7$ , isliye block ka warp 7 ragged wala hai. Warp 7 local indices $32 \times 7 = 224$ se $255$ tak, yaani global indices $768 + 224 = 992$ se $1023$ tak cover karta hai. Warp 7 mein active threads count karo: woh global indices $992$ se last active $999$ tak hain. Woh count hai (last active − first) + 1 $= 999 - 992 + 1 = 8$ active threads. Baaki $32 - 8 = 24$ lanes idle wale hain (global $1000..1023$ , sab $\geq N$ ). Yeh step kyun? Humne global → local → warp explicitly walk kiya taaki " $999 - 992 + 1 = 8$ " magic na lage: $992$ woh jagah hai jahan warp 7 global space mein shuru hota hai, $999$ last real element hai, aur "+1" isliye kyunki dono endpoints inclusive hain.

Verify: active threads $= 4 \times 256 - 24 = 1000 = N$ . ✅ 24 idle lanes sab last block ke warp 7 mein hain; baaki sab warps full hain.

Ex 3 — Cell C: degenerate & tiny inputs (N ≤ 0 aur N tiny)

Worked example Ek warp se bhi kam threads — aur illegal case

$B = 256$ ke saath teen sub-cases: (a) $N = 5$ ; (b) $N = 0$ ; (c) $N = - 3$ .

Forecast: $N = 5$ ke liye, warp 0 mein 32 cores mein se actually kitne useful kaam karte hain? Aur negative $N$ ke liye kya hona chahiye?

Step 1 — $N = 5$ : blocks. $⌈ \frac{5}{256} ⌉ = 1 block .$ Yeh step kyun? Ek single element ke liye bhi ek poora block launch karna padta hai.

Step 2 — $N = 5$ ke liye warp utilisation. Ek warp (32 threads) issue hota hai; sirf threads $0..4$ i < 5 pass karte hain. Toh $5/32 \approx 15.6%$ warp useful hai. Yeh step kyun? Yeh SIMT ka tax hai: warp atomic hai — tum "sirf 5 lanes" run nahi kar sakte; tum 32 run karte ho aur 27 ko mask off kar dete ho. Dekhte hain 6.1.03-SIMD-vs-SIMT kyun masking, skipping nahi, mechanism hai.

Step 3 — $N = 0$ : empty launch. $⌈ \frac{0}{256} ⌉ = 0 blocks .$ Yeh step kyun? Zero blocks ki grid kuch nahi launch karti — kernel turant return kar deta hai. Koi core run nahi karta. Yeh true degenerate floor hai.

Step 4 — $N = - 3$ : illegal input. Elements ki count negative nahi ho sakti, isliye $N = - 3$ meaningless data hai, valid scenario nahi. Do cheezein tumhe protect karti hain: (i) launch-config formula $⌈ N / B ⌉$ with $N = - 3, B = 256$ gives $⌈ - 0.0117 ⌉ = 0$ blocks — toh agar yeh slip bhi ho jaata, kuch launch nahi hota; (ii) zyada zaroori, kernel guard if (i < N) har thread ke index $i \geq 0$ ko negative $N$ se compare karta hai, jo kabhi bhi true nahi hota, isliye har thread fall through kar deta hai aur kuch nahi karta. Yeh step kyun? Matrix ne " $N \leq 0$ " ka vaada kiya tha, isliye hume sign case dikhana hi hoga, sirf zero nahi. Takeaway: negative $N$ "kuch mat karo" mein collapse ho jaata hai, exactly $N = 0$ ki tarah — guard ka < comparison hi dono ko safe banata hai.

Verify: $N = 5$ → active/total $= 5/32 = 0.15625$ . $N = 0$ → 0 blocks, 0 warps. $N = - 3$ → $⌈ - 3/256 ⌉ = 0$ blocks, guard kabhi true nahi → 0 active threads, koi crash nahi. ✅

Ex 4 — Cells D & E: memory best case vs worst case

Worked example Coalesced vs strided access

32 threads ka ek warp har ek ek float (4 bytes) read karta hai. (D) thread $i$ A[i] read karta hai, aur (E) thread $i$ A[i*32] read karta hai — dono ke liye transaction efficiency $η$ calculate karo.

Forecast: case E — kya $η$ 50% ke kareeb hoga ya 3% ke kareeb?

Recall the definition (parent se, aur hamare symbol block se): $η = \frac{useful bytes requested}{transactions \times 128 bytes/transaction} .$ Useful bytes dono cases mein same hain: $32 \times 4 = 128$ bytes. Sirf transaction count badlata hai.

Step 1 — Case D, useful bytes. $32 threads \times 4 bytes = 128 useful bytes .$ Yeh step kyun? Consecutive addresses A[0..31] exactly ek 128-byte aligned line span karte hain (figure mein amber band dekho — saare 32 arrows us mein land karte hain).

Step 2 — Case D, transactions & $η$ . $transactions = 1, η = \frac{128}{1 \times 128} = 1 = 100%.$ Yeh step kyun? Ek 128-byte transaction har zaroorat ka byte carry karta hai; kuch bhi waste nahi.

Step 3 — Case E, transactions. Addresses hain A[0], A[32], A[64], … — har ek 128 bytes apart, isliye har ek ek alag 128-byte line mein land karta hai (figure mein scattered cyan arrows). Yeh $32$ alag transactions hain. Yeh step kyun? Hardware sirf un threads ko merge kar sakta hai jo ek 128-byte line share karte hain; stride-32 merging ko completely defeat kar deta hai.

Step 4 — Case E, $η$ . $η = \frac{128}{32 \times 128} = \frac{1}{32} = 0.03125 = 3.125%.$ Yeh step kyun? Har transaction 128 bytes haul karta hai lekin tum sirf 4 rakhte ho → $4/128 = 1/32$ .

Verify (units): $η$ bytes ÷ bytes = dimensionless hai, aur dono mein $0 < η \leq 1$ . ✅ Slowdown factor = $η_{D} / η_{E} = 1/ (1/32) = 32 \times$ . Yeh parent ke "up to 32× slower" se match karta hai. ✅

Ex 5 — Cell F: branch divergence

Worked example Ek warp jo beech mein split ho jaata hai

32 threads ke ek warp ke andar, aadhe if lete hain aur aadhe else:

if (threadIdx.x % 2 == 0) { heavy_path(); }  // 20 cycles
else                       { light_path(); }  // 8 cycles

Warp kitne cycles spend karta hai, aur utilisation kya hai?

Forecast: kya yeh $20$ , $8$ , $20 + 8$ , ya $max (20, 8)$ cycles leta hai?

Step 1 — Ek warp ek saath dono kaam kyun nahi kar sakta. Saare 32 lanes ek instruction pointer share karte hain (yahi SIMT mein "SI" hai). Jab lanes agree nahi karte, hardware ek path run karta hai doosre lanes ko masked off (idle) karke, phir doosra path — serialized. Figure mein do stacked timelines dekho.

Step 2 — Total cycles. $cycles = 20 (even lanes) + 8 (odd lanes) = 28.$ Yeh step kyun? Paths ek ke baad ek run karte hain, isliye costs add hote hain, max nahi.

Step 3 — Utilisation. 20-cycle heavy path ke dauran, sirf 16 lanes active hain; 8-cycle light path ke dauran, doosre 16. Useful lane-cycles $= 16 \times 20 + 16 \times 8 = 448$ . Total lane-cycles jo warp ne occupy kiye $= 32 \times 28 = 896$ . $utilisation = \frac{448}{896} = 0.5 = 50%.$ Yeh step kyun? Two-way split mein aadhe lanes har waqt idle baithte hain — classic divergence penalty.

Verify: agar instead sab 32 lanes heavy path lete (no divergence), cost $= 20$ cycles, utilisation $= (32 \times 20) / (32 \times 20) = 100%$ . Divergence ne 20 cycles ko 28 kar diya aur 100% ko 50% kar diya. ✅

Ex 6 — Cells G & H: latency hiding, success aur failure

Worked example Kya memory latency hide karne ke liye kaafi warps hain?

Parent ka formula use karte hue $T_{required} = \frac{S \cdot L}{C}$ jahan $S$ =schedulers, $L$ =stall cycles, $C$ =useful compute cycles per warp between stalls. SM caps at $64$ resident warps. (G) $S = 4, L = 200, C = 20$ . (H) $S = 4, L = 200, C = 10$ .

Forecast: kaun sa case latency fully hide kar sakta hai?

Step 1 — Yeh formula kyun. Ek warp ke $L$ stall cycles ke dauran, SM $S \cdot L$ warp-instructions issue kar sakta hai; har doosra warp $\approx C$ supply karta hai phir woh bhi stall ho jaata hai, isliye pipe full rakhne ke liye $S \cdot L / C$ warps chahiye. (Full reasoning parent mein hai.)

Step 2 — Case G. $T_{required} = \frac{4 \times 200}{20} = 40 warps .$ Yeh step kyun? Zyada compute intensity $C = 20$ matlab har warp har stall mein zyada coverage deta hai.

Step 3 — Cap se compare karo. $40 \leq 64$ ⇒ latency fully hidden. Scheduler hamesha ek ready warp dhundh leta hai. Yeh step kyun? Required warps actually SM pe fit bhi hone chahiye; $T_{required}$ ko 64-warp cap se compare karna hi akela cheez hai jo batata hai ki hiding physically possible hai ya nahi.

Step 4 — Case H. $T_{required} = \frac{4 \times 200}{10} = 80 warps > 64 \Rightarrow cannot hide fully.$ Yeh step kyun? $C$ ko aadha karne se requirement hardware ceiling se zyada ho gayi — SM stall ho jaata hai, memory latency expose ho jaati hai. Isliye compute intensity badhana (dekho 6.2.04-occupancy-and-performance) important hai.

Verify: $40$ needed vs $64$ available → $24$ warps ka margin (G succeeds). $80 > 64$ → $16$ warps ka deficit (H fails). ✅ $C$ ko 10→20 double karne se $T_{required}$ 80→40 aadha ho gaya, inverse relationship confirm karta hai.

Ex 7 — Cell I: real-world word problem

Worked example Image ki grayscale conversion

Tum ek $1920 \times 1080$ RGB image ko grayscale mein convert karte ho, ek thread per pixel, $16 \times 16$ threads ke 2D blocks use karke. Har dimension mein kitne blocks, total threads, aur kitne waste hue?

Forecast: kya $1080$ evenly $16$ se divide hota hai?

Step 1 — Total pixels (real $N$ ). $N = 1920 \times 1080 = 2, 073, 600 pixels .$ Yeh step kyun? Ek thread per pixel matlab $N$ = threads ki sankhya jo tumhe chahiye.

Step 2 — Blocks per dimension (2D mein ceiling). $blocks_{x} = ⌈ \frac{1920}{16} ⌉ = 120, blocks_{y} = ⌈ \frac{1080}{16} ⌉ = ⌈ 67.5 ⌉ = 68.$ Yeh step kyun? $1920/16 = 120$ exactly, lekin $1080/16 = 67.5$ → 68 tak round up, ek ragged bottom edge create karta hai (cell B 2D mein wapas aata hai!).

Step 3 — Threads launched vs wasted. $threads = (120 \times 16) \times (68 \times 16) = 1920 \times 1088 = 2, 088, 960.$ $wasted = 2, 088, 960 - 2, 073, 600 = 15, 360.$ Yeh step kyun? Extra $1088 - 1080 = 8$ pixel rows × $1920$ columns $= 15, 360$ threads image ke bahar fall karte hain aur if (x<W && y<H) guard se mask ho jaate hain.

Verify: wasted $= 8 \times 1920 = 15, 360$ . ✅ Aur $2, 088, 960 - 15, 360 = 2, 073, 600 = N$ . ✅ Waste fraction $= 15360/2088960 \approx 0.74%$ — bahut chhoti, kyunki sirf ek edge ragged hai.

Ex 8 — Cell J: exam twist (ek mein do traps)

Worked example "Quarter cycle" aur register ceiling

Part 1: Ek SM mein 128 CUDA cores hain. Ek student claim karta hai ki ek warp instruction $32/128 = 1/4$ cycle mein khatam ho jaati hai. Sach ya jhooth? Part 2: SM mein $M = 65, 536$ registers hain aur tum $P = 2048$ resident threads chahte ho. Per-thread register ceiling kya hai, aur kya hoga agar tumhare kernel ko 40 registers/thread chahiye?

Forecast: guess karo Part 1 sach hai ya nahi, aur compute karne se pehle register ceiling guess karo.

Step 1 — Part 1, kyun yeh jhooth hai. Ek warp instruction per issue atomic hai: uska minimum 1 cycle hai, kabhi kam nahi. Extra cores throughput dete hain (zyada warps parallel mein issue), ek warp ke liye kam latency nahi. 128 cores aur 32-thread warps ke saath, SM ek saath $128/32 = 4$ warp-instructions issue kar sakta hai — yeh 4 warps per cycle hai, not ek warp ek quarter cycle mein. Yeh step kyun? Latency (ek warp ke liye time) aur throughput (warps per unit time) alag axes hain. Trap dono ko swap karta hai.

Step 2 — Part 2, register ceiling. Yahan $M = 65, 536$ SM ki register file mein 32-bit registers ki total sankhya hai, aur $P = 2048$ woh threads hain jo tum resident (alive) ek saath rakhna chahte ho. Woh saare live threads ek register file share karte hain, isliye har thread zyada se zyada use kar sakta hai: $\frac{M}{P} = \frac{65 , 536}{2048} = 32 registers/thread .$ Yeh step kyun? Agar $P$ live threads $M$ registers ko bina kuch bacha ke baantein, to har thread ka fair share exactly $M / P$ hai. Isse zyada use karo aur tum saare $P$ threads resident nahi rakh sakte.

Step 3 — Kernel ko 40 registers/thread chahiye (> 32). Tumhara kernel 40 registers per thread maangta hai, lekin upar ceiling sirf 32 hai full occupancy par — isliye full occupancy impossible hai. Iski jagah, resident threads drop karne chahiye jab tak $40 \times P \leq M$ : $P_{m a x} = ⌊ \frac{M}{40} ⌋ = ⌊ \frac{65 , 536}{40} ⌋ = ⌊ 1638.4 ⌋ = 1638 threads .$ Yeh step kyun? Tum kabhi $M$ total registers exceed nahi kar sakte, isliye zyada registers per thread matlab SM par kam threads. Kam resident threads → kam resident warps → kam latency hiding (yeh seedha Ex 6 mein wapas jaata hai, jahan required warp count se neeche jaana memory stalls expose karta hai).

Step 4 — Warps mein kya cost hai. $warps lost = \frac{2048 - 1638}{32} = \frac{410}{32} \approx 12.8 warps fewer resident.$ Yeh step kyun? Occupancy actually warps mein measure hoti hai; thread drop ko warps mein convert karna ( $\div 32$ ) scheduler ke pool par real hit dikhata hai. ~13 warps kho dena tumhe "hides latency" (Ex 6 Case G, needs 40) se "cannot" mein push kar sakta hai agar margins tight the.

Verify: $32 \times 2048 = 65, 536 = M$ exactly. ✅ Aur $40 \times 1638 = 65, 520 \leq 65, 536$ ✅ jabki $40 \times 1639 = 65, 560 > 65, 536$ ✗ — confirm karta hai 1638 true cap hai. Part 1: $128/32 = 4$ warps/cycle throughput, latency abhi bhi 1 cycle. ✅

Recall — har cell ke liye ek line

Recall Ragged tail abhi bhi ek poore warp ki cost kyun leta hai?

Kyunki warp atomic hai — 32 lanes ek saath issue hote hain; idle lanes mask hote hain, skip nahi; woh warp slot occupy karte hain. ::: Warp hamesha 32 lanes ke roop mein issue hota hai; guard if (i<N) extra walo ko mask karta hai lekin woh warp slot phir bhi lete hain.

Recall N = 0 ya negative N ke liye kya hota hai?

Dono "kuch mat karo" mein collapse ho jaate hain: $⌈ N / B ⌉ = 0$ blocks aur i < N guard kabhi true nahi hota. ::: Zero ya negative N 0 blocks launch karta hai aur har guard comparison fail ho jaata hai, isliye koi thread kaam nahi karta — koi crash nahi.

Recall Case E: stride-32 access ki efficiency?

$η = 1/32 = 3.125%$ — sirf 128 useful bytes ke liye 128 bytes ke 32 transactions. ::: $η = 128/ (32 \times 128) = 3.125%$ .

Recall Two-way branch split: cost aur utilisation?

Costs add hote hain ( $20 + 8 = 28$ cycles), utilisation 50% tak gir jaati hai kyunki aadhe lanes har waqt idle hain. ::: Paths serialize hote hain, isliye cycles add hote hain aur even split mein utilisation aadhi ho jaati hai.

Recall Latency hiding formula aur kab fail hoti hai?

$T_{required} = S L / C$ ; yeh fail hoti hai jab yeh ~64-warp SM cap se zyada ho jaaye. ::: Jab $S L / C$ > resident-warp cap, latency expose ho jaati hai; requirement kam karne ke liye $C$ badhao.

Recall Register ceiling: 65,536-register SM par 40 registers/thread?

Full occupancy sirf 32/thread allow karta hai, isliye tum $⌊ 65536/40 ⌋ = 1638$ resident threads tak gir jaate ho. ::: Zyada registers per thread matlab kam resident threads ( $P_{m a x} = ⌊ M /40 ⌋ = 1638$ ), occupancy kam ho jaati hai.

Recall "Quarter cycle" trap?

Jhooth — ek warp instruction $\geq 1$ cycle hai; 128 cores 4 warps/cycle throughput dete hain, sub-cycle latency nahi. ::: Ek warp instruction minimum 1 cycle par atomic hai; extra cores throughput add karte hain, latency kam nahi karte.

Dekho bhi 7.3.01-parallel-programming-patterns aur 9.1.02-neural-network-training-on-GPUs jahan yahi exact corners real performance decide karte hain.