Exercises — Thread blocks and grids

WHAT: Ek formula mein plug in karo. WHY: block index hamare se pehle complete blocks count karta hai (har ek mein 128 seats hain), phir hum apna seat add karte hain. $\text{globalID}=3\times 128+10=384+10=394$ Answer: 394. Picture: hum blocks 0, 1, 2 ke upar se kood gaye (wo hai $3\times 128=384$ elements) aur block 3 mein 10 seats andar land kiye.

• threadIdx.x ::: thread ki apne block ke andar position (0 se blockDim.x - 1 tak).
• blockDim.x ::: har block mein threads ki sankhya.
• blockIdx.x ::: grid ke andar block ki position (0 se gridDim.x - 1 tak).
• gridDim.x ::: grid mein blocks ki sankhya.

Memory hook: thread words "thread/block position" pe khatam hote hain; dim words sizes hain.

Jhooth. __syncthreads() sirf ek block ke andar ek barrier hai — dekho Thread Synchronization. Blocks autonomous hain; ek single kernel launch ke andar koi cross-block barrier nahi hota. Alag blocks ke threads is tarah kabhi ek doosre ka intezaar nahi karte.

WHAT: ceiling division, kyunki hum har element ko cover karna chahte hain chahe aakhri block partly empty ho. $\text{blocks}=\lceil \frac{N}{B}\rceil =\frac{N+B-1}{B}=\frac{1000+255}{256}=\frac{1255}{256}=4\text{ (integer)}$ WHY ceiling: $1000/256=3.9$, aur 3 blocks sirf $3\times 256=768$ threads dete hain — 232 elements abhi bhi uncovered. Toh upar 4 tak round karo.

• Total threads launched: $4\times 256=1024$.
• Idle (guarded-out) threads: $1024-1000=24$.

Un 24 threads ka globalID $[1000,1023]$ mein hai aur if (i < N) guard unhe kuch nahi karne deta.

$\text{globalID}=2\times 256+200=512+200=712$ Kyunki $712<1000$, guard if (i < N) pass ho jaata hai — ye zaroor element 712 process karta hai.

WHAT: har axis pe wahi formula apply karo, phir flatten karo. WHY flatten: memory ek 1D line hai; row-major matlab "ek full row ke saath chalo, phir next pe drop karo" — dekho Memory Coalescing. $\text{col}=5\times 16+3=83$ $\text{row}=7\times 16+9=121$ $\text{idx}=\text{row}\times \text{width}+\text{col}=121\times 512+83=61952+83=62035$ Answer: col 83, row 121, memory index 62035.

WHY occupancy ek minimum hai: ek block tabhi resident ban sakta hai jab har resource jo use chahiye ek saath available ho — threads, registers, aur shared memory. Agar thread budget 8 blocks allow karta hai lekin memory budget sirf 5 allow karta hai, toh sirf 5 SM pe actually baith sakte hain: sabse tight resource real ceiling hai. Toh hum har limit alag se compute karte hain aur sabse chhota lete hain — dekho GPU Occupancy.

• Thread limit: $\lfloor 2048/256\rfloor =8$ blocks.
• Shared-memory limit: $\lfloor 96/12\rfloor =8$ blocks.

WHY yahan floor: ek block jo apna poora memory footprint fit nahi kar sakta simply admit nahi hota — "0.7 of a block" resident jaisi koi cheez nahi hoti — toh koi bhi fractional part neeche discard ho jaata hai.

Dono limits 8 kehte hain, toh 8 blocks fit hote hain $\Rightarrow 8\times 256=2048$ resident threads $=2048/2048=$ 100% occupancy.

WHY phir min: occupancy us resource se cap hoti hai jo pehle khatam ho (dekho 3.1). Dono compute karo, chhota rakho.

• Thread limit: $\lfloor 2048/256\rfloor =8$ blocks.
• Shared-memory limit: $\lfloor 96/20\rfloor =4$ blocks. WHY floor: $96/20=4.8$, lekin 5th block ko $5\times 20=100$ KB $>96$ KB chahiye — wo fit nahi hota, toh hum fractional 0.8 drop karte hain aur 4 rakhte hain.

Binding constraint chhota wala hai: $\min (8,4)=4$ blocks. $\text{threads}=4\times 256=1024,\text{occupancy}=\frac{1024}{2048}=50\%$ Lesson: memory hunger — thread count nahi — ne occupancy aadhi kar di.

WHAT: ek warp 32 threads ka fixed bundle hai jo ek instruction saath run karta hai (SIMT) — dekho Warp Execution. Ek lane ek thread ki fixed slot hai us bundle ke andar, 0–31 numbered; jab warp ek instruction execute karta hai, saare 32 lanes ise saath fire karte hain (jab tak kuch mask off na hon).

• Warps per block: $256/32=8$ warps.
• Jab ek warp do paths mein split hota hai, hardware if path lanes 16–31 masked off ke saath run karta hai, phir else path lanes 0–15 masked off ke saath. Har pass masked half waste karta hai — wo lanes abhi bhi kuch na karte hue ek cycle consume karte hain.
• Toh har divergent warp roughly 2× cycles spend karta hai; divergent region mein roughly 50% slots idle mask-outs hain.

Practice mein fix: branches is tarah arrange karo ki ek poora warp ek path le (branch on globalID / 32, lane number pe nahi), taki ek warp ke andar koi lane kabhi doosre ke against mask na ho.

Choice: B = 256 — 32 ka multiple (whole warps, no wasted lanes), aur 128–512 sweet spot mein jo enough blocks per SM rakhta hai memory latency hide karne ke liye occupancy starve kiye bina. $\text{blocks}=\lceil \frac{1000000}{256}\rceil =\frac{1000000+255}{256}=\frac{1000255}{256}=3907\text{ (integer)}$ Check: $3906\times 256=999936<1000000$, toh 3906 kam hai; 3907 sahi hai.

• Total threads: $3907\times 256=1000192$.
• Idle (guarded) threads: $1000192-1000000=192$.

Launch: vecAdd<<<3907, 256>>>(...) with if (i < N) guard.

WHAT: x ke along columns cover karo, y ke along rows (convention: x = horizontal = columns) — dekho Parallel Algorithm Design. $\text{gridDim.x}=\lceil \frac{N}{16}\rceil =\lceil \frac{500}{16}\rceil =32(500/16=31.25\to 32)$ $\text{gridDim.y}=\lceil \frac{M}{16}\rceil =\lceil \frac{1000}{16}\rceil =63(1000/16=62.5\to 63)$

• Threads launched: $(32\times 16)\times (63\times 16)=512\times 1008=516096$.
• Useful outputs: $M\times N=1000\times 500=500000$.
• Guarded-out threads: $516096-500000=16096$.
• Fraction wasted: $16096/516096\approx 0.0312=3.12\%$.

Thodi waste normal hai — ye non-multiple grid ko tile karne ki keemat hai.

WHAT: index rule teen axes mein se har ek pe independently apply hota hai — kuch naya nahi, bas x, y, z side by side. $\text{gridDim.x}=\lceil 64/8\rceil =8,\text{gridDim.y}=8,\text{gridDim.z}=8$ $x=2\times 8+1=17$ $y=3\times 8+5=29$ $z=4\times 8+7=39$ Answer: grid $8\times 8\times 8$ blocks hai; thread voxel $(x,y,z)=(17,29,39)$ pe map karta hai.

1D memory index mein flatten karo — aur WHY exactly yeh nesting: memory ek lambi line hai, toh hum decide karna chahte hain kaun sa coordinate sabse tez chalta hai. Convention x ko fastest-changing banati hai (neighbouring x memory mein adjacent hain, jo Memory Coalescing help karta hai), phir y, phir z slowest. Toh hum address layers mein banate hain: z-slice chuno (har slice $Y\times X$ voxels hold karta hai), phir us mein y-row (har row $X$ voxels hold karta hai), phir us row ke along x step karo: $\text{idx}=(z\times Y+y)\times X+x=(39\times 64+29)\times 64+17=2525\times 64+17=161617$ WHY kaam karta hai: $(z\times Y+y)$ row number hai volume ke bilkul shuru se count kiya hua, aur $X$ se multiply karna hamare se pehle ke saare full rows skip karta hai — phir $+x$ current row mein humari seat pe land karta hai. Ye wahi "apne se pehle ke complete groups skip karo, phir apna offset add karo" logic hai jaise 1D global-index formula, teen deep nested.

WHAT: resident threads $=(\text{block size})\times \min (\text{thread-limit blocks},\text{resource-limit blocks})$, jahan thread-limit blocks $=\lfloor 2048/\text{block size}\rfloor$.

128: thread limit $=2048/128=16$; resources 12 allow karte hain; $\min (16,12)=12$. $\text{threads}=128\times 12=1536,\text{occ}=\frac{1536}{2048}=75\%$

256: thread limit $=2048/256=8$; resources 8 allow karte hain; $\min (8,8)=8$. $\text{threads}=256\times 8=2048,\text{occ}=\frac{2048}{2048}=100\%$

1024: thread limit $=2048/1024=2$; resources 2 allow karte hain; $\min (2,2)=2$. $\text{threads}=1024\times 2=2048,\text{occ}=\frac{2048}{2048}=100\%$

Winner: 256-thread block. Dono 256 aur 1024 100% occupancy reach karte hain, lekin 256 safer choice hai — 8 blocks resident hone ke saath, ek stalled block abhi bhi 7 aur scheduler ke liye rehne deta hai latency run aur hide karne ke liye; 1024 ke saath tumhare paas sirf 2 blocks hain, toh ek stall tumhari latency-hiding capacity half kar deta hai. 128-thread block eliminate ho jaata hai kyunki resources (threads nahi) ise 12 blocks pe cap karte hain $\Rightarrow$ sirf 75% occupancy. 256 chuno.

Residency limit jo ise fail karti hai: har SM at most 8 blocks at a time hold kar sakta hai, aur sirf 4 SMs hain, toh at most $4\times 8=32$ of the 40 blocks simultaneously resident ho sakte hain. Baaki 8 blocks ko ek resident block ke finish hone aur apna slot free karne ka intezaar karna padega pehle shuru hone se.

Kyun deadlock hota hai: blocks SMs pe arbitrary order mein schedule hote hain. Maan lo 32 resident slots saare blocks 1–32 se bhar gaye hain, har ek flag pe spin kar raha hai. Spinning blocks kabhi finish nahi hote, toh unke slots kabhi free nahi hote, toh not-yet-resident blocks — jinme block 0 bhi ho sakta hai — kabhi schedule nahi hote. Agar block 0 waiting 8 mein se ek hai, flag kabhi set nahi hota, toh resident blocks forever spin karte hain: ek circular wait. Koi guarantee nahi ki block 0 pehle run kare; block indices koi execution order imply nahi karte.

Sahi alternatives (Thread Synchronization se):

1. Do kernel launches mein split karo. Kernel boundary ek implicit global barrier hai — launch 1 ki sab cheez launch 2 shuru hone se pehle complete ho jaati hai.
2. Cooperative groups (CUDA 9+) grid-level sync, cudaLaunchCooperativeKernel ke saath launch kiya, jo guarantee karta hai ki saare blocks co-resident hain toh grid barrier legal hai.
3. (Last resort) coordination ke liye global memory pe atomic operations — sahi lekin slow, aur phir bhi yeh assume karne se bachna hoga ki koi block pehle run kare.

"No waste" matlab N block size ka exact multiple ho (remainder 0). Har ek check karo:

• $4096/64=64$ exactly → 64 blocks, no waste. ✅
• $4096/128=32$ exactly → 32 blocks, no waste. ✅
• $4096/256=16$ exactly → 16 blocks, no waste. ✅
• $4096/512=8$ exactly → 8 blocks, no waste. ✅

Charon $4096=2^{12}$ ko evenly divide karte hain (ye saare powers of two hain jo ise divide karte hain), aur saare warp multiples hain (har ek 32 ka multiple hai). Toh har listed size zero guarded-out threads deti hai. Unme se, 256 Exercise 5.1 mein argue kiye gaye occupancy/latency-hiding balance ke liye usual pick hai.

• Universal 1D index formula ::: $\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}$
• Teen axes mein se har ek pe wahi rule ::: x, y, aur z pe independently apply karo (jaise $\text{blockIdx.z}\times \text{blockDim.z}+\text{threadIdx.z}$)
• N elements cover karne ke liye block count (round up) ::: $\lceil N/B\rceil =(N+B-1)/B$
• Fixed resource budget fit karne wale blocks (round down) ::: $\lfloor \text{budget}/\text{per-block cost}\rfloor$
• Occupancy ::: resident threads ÷ SM max threads, saare resource limits ke minimum se capped
• Kyun blocks ek kernel mein ek doosre ko message nahi kar sakte ::: arbitrary scheduling order + limited residency → deadlock risk