Worked examples — CUDA programming model basics

6.2.13 · D3 · Hardware › GPU Architecture › CUDA programming model basics

Yeh page CUDA programming model basics ka hands-on companion hai. Parent note ne machinery build ki thi: blocks SMs pe map hote hain, threads ek global index carry karte hain, aur memory ke speed tiers hote hain. Yahaan hum us machinery ko use karte hain — un har awkward cases pe jo ek real problem throw kar sakti hai.

Kuch bhi shuru karne se pehle, chalo us ek formula ko dobara anchor karte hain jis par hum poori page rely karenge. Ek lambi line of numbered lockers imagine karo — woh hai tumhari array. Hum us line ko equal blocks (fixed-size chunks) mein kaatте hain. Har block ke andar, threads $0$ se upar number hote hain.

Pehle figure padho. Horizontal strip tumhari array hai; har chhota box ek thread ka element hai, aur har box ke andar ka number us thread ka globalIdx hai. Cyan boxes ordinary threads hain; single amber box (sabse right wala, index $11$) woh thread hai jise hum Example 1 mein track karte hain. Boxes ke upar, white brackets boxes ko blocks mein group karte hain — teen brackets jinka label hai block 0, block 1, block 2. Har box ke neeche chhote white t0 … t3 tags hain jo us block ke andar ki seat number (threadIdx.x) dikhate hain. Pattern notice karo: seats har nayi bracket pe $0$ se restart hoti hain, jabki bada number climb karta rehta hai — woh climb hi global index hai.

Upar har symbol ka ek picture hai. Neeche do aur pieces of notation baar baar aayenge, isliye hum unhe abhi define karte hain, pehli use se pehle:

Ab matrix.

The scenario matrix

Har CUDA-indexing problem jo tumhe milegi woh in cells mein se ek hai. Neeche wale examples us cell ke label ke saath hain jo woh cover karta hai, aur saath mein woh sab mila ke sab cover kar lete hain.

Cell	Kya cheez isse tricky banati hai	Covered by
A. Clean divide	$N$ block size ka exact multiple hai — koi leftovers nahi	Ex 1
B. Ragged tail	$N$ divisible nahi — extra idle threads aate hain	Ex 2
C. Degenerate: tiny $N$	$N$ ek block se chhota, ya $N=0$, ya $N=1$	Ex 3
D. Reverse map	Ek globalIdx diya hai, pata karo kaun sa block & seat tha	Ex 4
E. 2-D grid	Image / matrix — do index axes, row-major flattening	Ex 5
F. Memory-cost word problem	Kya data registers, shared, ya global mein rehna chahiye?	Ex 6
G. Launch-config sizing	Block size choose karo; idle threads aur wasted warp lanes count karo	Ex 7
H. Exam twist: stride loop	Threads se zyada elements — ek thread kaafi handle karta hai	Ex 8
I. 3-D grid	Volumes / batches — teen index axes, depth·height·width flatten	Ex 9

---

Cell A — the clean divide

Forecast: aage padhne se pehle index guess karo. (Hint: last block, last seat — bilkul last element hona chahiye.)

1. Do pieces padho. blockIdx.x = 2, blockDim.x = 4, threadIdx.x = 3. Yeh step kyun? Formula ko exactly yahi teen numbers chahiye; pehle inhe gather karo taaki baad mein kuch guess na karna pade.
2. Mere block se pehle wale threads: $2\times 4=8$. Yeh step kyun? Mere aage $4$ seats ke do poore chunks hain; mujhe saare $8$ skip karne honge. (Figure mein, yeh do full brackets block 0 aur block 1 hain.)
3. Meri seat add karo: $8+3=11$. Yeh step kyun? Mere chunk ke andar seat $3$ mujhe mere chunk ke front door se $3$ aage le jaati hai — figure mein woh amber box hai.
4. Boundary check if (idx < N): kya $11<12$ hai? Haan. Yeh step kyun? Clean divide mein bhi hum test karte hain — acha habit hai, aur baad ke ragged cases ke liye protect karta hai.

Answer: globalIdx $=11$, valid.

Verify: Total threads $=3\times 4=12$, exactly $N$. Jo indices produce hue woh hain $0,1,\dots ,11$ — har element ek baar hit, koi dobara nahi, koi skip nahi. $11$ sabse bada hai, isliye last thread last element hit karta hai. Consistent.

---

Cell B — the ragged tail

Forecast: kya $1000/256$ evenly aayega? Nahi — toh expect karo leftover idle threads. Guess karo kitne.

1. Ceiling division. Hume saare $1000$ elements ke liye kam se kam enough threads chahiye: $\text{numBlocks}=\lceil \frac{1000}{256}\rceil =\frac{1000+256-1}{256}=\frac{1255}{256}=4(\text{integer division}).$ Yeh step kyun? $4$ blocks $1024$ threads dete hain; $3$ blocks sirf $768<1000$ dete hain, jo $232$ elements untouched chhod deta. Hume upar round karna hai — yahi exactly ceiling bracket $\lceil \cdot \rceil$ (upar define kiya) karta hai.
2. $+255$ trick kyun? threadsPerBlock − 1 add karna aur phir integer-divide karna, floating-point math ke bina round-up force karta hai. Yeh step kyun? Integer division truncate karta hai neeche; yeh offset kisi bhi non-zero remainder ko next boundary ke upar push kar deta hai. Yeh sawaal ka jawaab deta hai "kya kuch last full block ke past spill hua?"
3. Total threads launched: $4\times 256=1024$. Yeh step kyun? Blocks fixed sizes mein aate hain — tum partial block launch nahi kar sakte, isliye tum hamesha over-provision karte ho.
4. Idle threads: $1024-1000=24$. Inke idx hain $1000\dots 1023$. Yeh step kyun? Exactly isliye if (idx < N) exist karta hai — iske bina, threads $1000$–$1023$ array ke end se aage read kar lete. Parent note ke mistakes callout mein dekho.

Answer: 4 blocks, 1024 threads launched, 24 idle.

Verify: $4\times 256=1024\ge 1000$ ✓ aur $3\times 256=768<1000$ ✗, isliye $4$ smallest valid count hai. Idle $=24$, aur $1024-24=1000=N$. ✓

---

Cell C — degenerate & edge inputs

Forecast: $N=0$ woh sneaky wala hai — kuch nahi ke liye kitne blocks launch karne chahiye?

1. (a) $N=1$: $\lceil 1/256\rceil =\lceil 0.0039\dots \rceil =1$ block. Kyun? Ek bhi element ke liye ek block chahiye; tum fraction of a block launch nahi kar sakte. Threads launched $=256$, idle $=255$. Sirf thread $0$ idx < 1 pass karta hai.
2. (b) $N=0$: $\lceil 0/256\rceil =0$ blocks. Trick se: $(0+255)/256=255/256=0$ (integer division). Kyun? Kuch karna nahi → ceiling formula $0$ blocks return karta hai.
3. (c) $N=256$: $\lceil 256/256\rceil =1$ block. Trick: $(256+255)/256=511/256=1$. Kyun? Exactly ek full block; $+255$ offset hamein falsely $2$ tak bump nahi karta kyunki $511<512$. Yeh confirm karta hai ki ceiling trick exact multiples pe kabhi over-round nahi karti.

Answer: (a) 1 block, 255 idle; (b) 0 blocks (launch guard karo), 0 idle; (c) 1 block, 0 idle.

Verify: (a) $1\cdot 256-1=255$ ✓. (b) $0\cdot 256-0=0$ ✓. (c) $1\cdot 256-256=0$ ✓, aur trick $1$ deta hai, $2$ nahi, isliye koi phantom extra block nahi.

---

Cell D — the reverse map

Forecast: kaun sa block? Roughly $813/256\approx 3.2$, toh shayad block $3$. Exact ho jaate hain.

1. Block recover karo integer division se (upar define kiya hua floor bracket $\lfloor \cdot \rfloor$ — round down): $\text{blockIdx.x}=\lfloor \frac{813}{256}\rfloor =3.$ Yeh step kyun? Formula tha $\text{globalIdx}=\text{blockIdx}\times 256+\text{seat}$ with $0\le \text{seat}<256$. $256$ se divide karke remainder drop karne se seat strip ho jaati hai aur block bach jaata hai — yeh multiplication ko invert karta hai.
2. Seat recover karo remainder (modulo) se: $\text{threadIdx.x}=813\mathrm{mod}256=813-3\times 256=813-768=45.$ Yeh step kyun? Remainder exactly hai "block ke front door se kitna aage" — woh seat jo step 1 mein skip hui.

Answer: block 3, seat 45.

Verify: Index forward rebuild karo: $3\times 256+45=768+45=813$ ✓. Aur $0\le 45<256$, ek legal seat. Dono directions agree karte hain — division aur modulo index formula ke clean inverse hain.

---

Cell E — the 2-D grid

Images aur matrices ko do axes chahiye. Ek photo pe ek grid imagine karo: threadIdx.x/blockIdx.x columns count karte hain (across), threadIdx.y/blockIdx.y rows count karte hain (down).

Pehle figure padho. Bada square image ka $64\times 64$ corner hai. Faint white lines isse $16\times 16$ blocks mein kaatti hain (chaar across, chaar down). Cyan-shaded square woh ek block hai blockIdx = (2, 3) pe — uske columns $32-47$ hain (yani $2\times 16$ se $2\times 16+15$ tak) aur rows $48-63$ ($3\times 16$ se aage). Us block ke andar, amber square marker hamara target pixel hai: us block ke andar seat (5, 7) column $37$, row $55$ pe land karta hai. Amber label dono chhote index sums spell out karta hai. Notice karo rows downward increase karte hain (image convention) — y-axis bilkul isi wajah se flipped hai.

Forecast: row .y numbers use karta hai, column .x numbers use karta hai. Compute karne se pehle guess karo.

1. Column (x-axis): $\text{col}=\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}=2\times 16+5=37.$ Yeh step kyun? Same 1-D logic jaisi Example 1, sirf horizontal axis pe apply ki.
2. Row (y-axis): $\text{row}=\text{blockIdx.y}\times \text{blockDim.y}+\text{threadIdx.y}=3\times 16+7=55.$ Yeh step kyun? Vertical axis same formula ki ek doosri, independent copy hai.
3. 1-D mein flatten karo (row-major): $\text{globalIdx}=\text{row}\times \text{width}+\text{col}=55\times 512+37=28160+37=28197.$ Yeh step kyun? Memory ek seedhi line hai; $512$ pixels ki ek poori row har nayi row se pehle aati hai. row × width multiply karne se upar ki saari complete rows skip ho jaati hain, phir + col current row mein chalta hai. Yeh exactly wahi 1-D idea hai apne andar nested.

Answer: pixel (row 55, col 37), flat index 28197.

Verify: Bounds: $37<512$ ✓, $55<512$ ✓. Flatten range check: max legal flat index $=512\times 512-1=262143$, aur $28197<262143$ ✓. Flat se col re-derive karo: $28197\mathrm{mod}512=37$ ✓; row $=\lfloor 28197/512\rfloor =55$ ✓.

---

Cell F — the memory-cost word problem

Forecast: 8 baar door wali global memory se re-read karna wasteful lagta hai. Cycle counts guess karo.

1. All-global cost: $8\text{ accesses}\times 300\text{ cycles}=2400$ cycles. Yeh step kyun? Agar kuch bhi cached nahi hai, toh 8 reads mein se har ek full off-chip trip pay karta hai. Memory Hierarchy dekho yeh jaanne ke liye ki global ~$300$ cycles kyun hai (off-chip, ~cm of wire + controller).
2. Shared-memory cost: ek global load $=300$, phir $8$ shared reads $=8\times 5=40$. Total $=300+40=340$ cycles. Yeh step kyun? Data on-chip stage karne ke liye expensive trip ek baar pay karo, phir saste mein re-read karo. Shared memory SM ke andar rehti hai.
3. Speedup: $2400/340\approx 7.06\times$. Yeh step kyun? Yahi ratio woh reason hai jiske liye shared-memory tiling exist karti hai — reuse ek expensive load ko amortise karta hai.

Answer: 2400 vs 340 cycles → shared memory ke saath ≈7.06× faster.

Verify: Break-even reuse count $r$ solve karo: $300r=300+5r\Rightarrow 295r=300\Rightarrow r\approx 1.017$. Toh koi bhi element jo do ya zyada baar read ho, woh already shared memory favour karta hai — aur hum ise $8$ baar read karte hain, comfortably break-even se aage. ✓

---

Cell G — launch-config sizing

Forecast: $1000$ se $1000000$ evenly divide hota hai — tempting! Lekin kya $1000$ allowed hai (max $1024$ hai), aur kya yeh $32$ ka multiple hai? Guess karo hidden waste kahan hai.

1. Option A — 256: $\text{numBlocks}=\lceil 10^6/256\rceil =(10^6+255)/256=3907$. Threads $=3907\times 256=1000192$. Idle (grid-level) $=192$. Yeh step kyun? $256$ threads $256/32=8$ whole warps hain — koi partial warp nahi. Wasted warp lanes per block $=0$, isliye total wasted lanes $=0$.
2. Option B — 1000: $1000\le 1024$, isliye yeh legal hai. $\text{numBlocks}=10^6/1000=1000$ exactly. Threads $=1000000$. Idle (grid-level) $=0$. Yeh step kyun? Exact divide ka matlab hai grid level pe zero leftover threads — koi ragged tail nahi. Ab tak perfect lagta hai.
3. Option B ke liye warp-level waste quantify karo. Warps per block $=\lceil 1000/32\rceil =32$ warps. Woh $32$ warps physically $32\times 32=1024$ lanes occupy karte hain, lekin sirf $1000$ threads real hain. Wasted lanes per block $=1024-1000=24$. Yeh step kyun? Hardware hamesha $32$ lanes ka poora warp schedule karta hai; har block ka last warp $1000-31\times 32=1000-992=8$ live threads ke saath aur $24$ dead lanes ke saath run karta hai. Woh dead lanes phir bhi ek scheduling slot consume karte hain — pure waste.
4. Grid ke across total wasted lanes (Option B): $24\text{ lanes/block}\times 1000\text{ blocks}=24000$ wasted lanes. Yeh step kyun? Poori picture dekhne ke liye per-block waste ko block count se multiply karo. Option A total $192$ idle threads waste karta hai; Option B $24000$ warp lanes waste karta hai — $100\times$ se zyada waste, apne "zero idle threads" headline ke peeche chhupa hua.

Answer: A: 3907 blocks, 1 000 192 threads, 192 idle, 0 wasted lanes. B: 1000 blocks, 1 000 000 threads, 0 grid-idle, lekin 24 wasted lanes/block = 24 000 total — worse choice hai.

Verify: A: $3907\times 256=1000192\ge 10^6$ aur $3906\times 256=999936<10^6$ ✓ (smallest count); $256\mathrm{mod}32=0$ → $0$ waste ✓. B: $1000\times 1000=10^6$ exactly ✓, $1000\le 1024$ ✓ (legal); warps/block $=\lceil 1000/32\rceil =32$, waste/block $=32\times 32-1000=24$ ✓, total $=24\times 1000=24000$ ✓.

---

Cell H — the exam twist (grid-stride loop)

Kabhi kabhi tum purpose se threads se kam threads launch karte ho elements se (jaise ek specific GPU mein fit karne ke liye). Har thread phir loop karta hai, har baar total number of threads ke barabar aage jump karta hua. Woh jump hai stride.

Pehle figure padho. Strip ek $20$-element array hai, boxes $0-19$ number kiye gaye hain. Har box ek chhota white t0…t7 tag carry karta hai jo dikhata hai $8$ threads mein se kaun sa us element ka owner hai — tum dekh sakte ho tags repeat hote hain $0,1,\dots ,7,0,1,\dots$ kyunki har thread $8$ ke stride se aage jump karta hai. Thread 3 ke owned boxes bright amber mein drawn hain; do curved amber arrows uski journey $3\to 11\to 19$ trace karte hain. $19$ ke baad, agla jump $27$ pe land karta, jo strip se bahar hai — isliye thread 3 rok jaata hai. Yahi loop condition i < N apna kaam kar rahi hai.

Forecast: thread 3 element 3 se shuru karta hai aur 8 se jump karta hai. Compute karne se pehle uski list guess karo.

1. Stride compute karo. $\text{stride}=\text{numBlocks}\times \text{blockDim}=2\times 4=8$. Yeh step kyun? Stride total thread count ke barabar hona chahiye taaki 8 threads array ko bina overlap ke tile karen — jaise 8 log "1-8, 9-16, ..." baar baar count off karte hain.
2. Thread 3 march karo. Start $i=3$; phir $3,3+8=11,11+8=19,19+8=27$. Ruk jaao, kyunki $27\nless 20$. Yeh step kyun? Loop condition i < N har thread ko cleanly halt karta hai jab woh end se bahar jaata hai — yeh ek single if (idx < N) check ko ek aise loop se replace karta hai jo self-limit karta hai.
3. Survivors collect karo. Jo values i < 20 pass kar gayi woh hain $3,11,19$. Toh thread 3 $\{3,11,19\}$ handle karta hai — ek thread se teen elements. Yeh step kyun? Grid-stride loop ka poora point yahi hai: $8$ threads $20$ elements cover karte hain har ek thode passes karke.

Answer: thread 3 → elements 3, 11, 19.

Verify: Coverage count: elements $0$–$7$ (first pass, saare 8 threads), $8$–$15$ (second pass), $16$–$19$ (third pass, sirf threads $0$–$3$ pahunche). Total $=8+8+4=20=N$ ✓, har index exactly ek baar hit hua (do threads ek hi start share nahi karte, aur stride thread count ke barabar hai isliye passes kabhi collide nahi karte). ✓

---

Cell I — the 3-D grid

Volumes (medical scans, physics grids) aur batched tensors ko ek teesra axis chahiye: blockIdx.z / threadIdx.z depth count karte hain (kaun sa slice). Yeh same one-axis formula hai, ab teen baar stamp out kiya gaya, phir row-major trick ke nested version se flatten kiya gaya.

Forecast: teen independent sums, phir ek three-term flatten. Pehle $z$ guess karo — woh .z numbers use karta hai.

1. x-coordinate: $x=1\times 8+4=12$. Yeh step kyun? Horizontal axis, base formula ki ek copy.
2. y-coordinate: $y=2\times 8+5=21$. Yeh step kyun? Vertical axis, independent doosri copy.
3. z-coordinate (depth): $z=3\times 8+6=30$. Yeh step kyun? Depth axis, independent teesri copy — yahi naya piece hai jo 3-D add karta hai.
4. Depth-major flatten karo $\text{height}=\text{width}=64$ ke saath: $\text{globalIdx}=30\times (64\times 64)+21\times 64+12=30\times 4096+1344+12=122880+1356=124236.$ Yeh step kyun? $30$ poore $64\times 64$ slices skip karo, phir $21$ poori rows, phir $12$ across.

Answer: $(x,y,z)=(12,21,30)$, flat index 124236.

Verify: Bounds: $12,21,30$ sab $<64$ ✓. Max legal flat index $=64^3-1=262143$, aur $124236<262143$ ✓. Unflatten: $124236\mathrm{mod}64=12=x$ ✓; $(124236//64)\mathrm{mod}64=1941\mathrm{mod}64=21=y$ ✓; $124236//4096=30=z$ ✓.

---

Recall Self-test

Clean vs ragged: +threadsPerBlock-1 trick kya achieve karta hai? ::: Yeh integer division ko upar round (ceiling $\lceil \cdot \rceil$) karne ke liye force karta hai taaki hum kabhi bhi kam blocks launch na karein, floating-point math ke bina. Reverse map: globalIdx aur blockDim diye hue, blockIdx aur threadIdx kaise recover karte ho? ::: blockIdx $=\lfloor$globalIdx / blockDim$\rfloor$ (floor / integer divide), threadIdx $=$ globalIdx mod blockDim. 2-D flatten: row × width + col kyun? ::: Memory 1-D hai; width pixels ki har poori row agli se pehle hoti hai, isliye row × width complete rows skip karta hai aur + col current row mein chalta hai. 3-D flatten: extra term kya hai aur kyun? ::: $z\times (\text{height}\times \text{width})$ — yeh 2-D formula apply karne se pehle $z$ poore 2-D slices skip karta hai. Warp kya hai aur 32 kyun? ::: Warp ek hardware bundle hai 32 threads ka jo lockstep mein run karte hain; 32 ek fixed NVIDIA hardware width hai (ek instruction 32 lanes feed karta hai). Shared vs global break-even: kitne reuses se shared memory jeetni hai? ::: 2 reuses se ($300r=300+5r$ gives $r\approx 1.02$). 256 ko 1000 threads/block pe kyun prefer karte ho, chahe 1000 evenly divide ho? ::: $256$ warp size $32$ ka multiple hai (0 wasted lanes); $1000$ har block mein $24$ dead lanes chodta hai (grid ke across $24000$).

See also: GPU Architecture Overview · Parallel Algorithm Design · OpenCL vs CUDA · Deep Learning Frameworks