Thread blocks and grids: Question bank — Hinglish deep dive

6.2.6 · D5 · Hardware › GPU Architecture › Thread blocks and grids

True ya false — justify karo

Ek block hamesha exactly ek SM par run karta hai

True. Block ke har thread us SM ki on-chip shared memory aur registers share karte hain, isliye block do SMs mein split nahi ho sakta — warna shared memory ka promise toot jaata.

Ek SM ek time mein exactly ek block run karta hai

False. Ek SM utne blocks hold karta hai jitna uska budget allow kare (threads, registers, shared memory), aksar 4–8 ek saath. Unhe interleave karna hi woh tarika hai jisse SM memory latency hide karta hai.

Grid ke andar blocks numerical order mein execute hote hain (block 0, phir 1, phir 2…)

False. Runtime blocks ko kisi bhi order mein schedule karta hai aur concurrently bhi run kar sakta hai; block 5 khatam ho sakta hai block 1 ke shuru hone se pehle bhi. Numbering sirf identity hai, timeline nahi.

__syncthreads() poore grid ko synchronize karta hai

False. Yeh sirf ek block ke andar threads ke liye barrier hai. Ek single ordinary kernel launch ke andar koi built-in whole-grid barrier nahi hota (dekho Thread Synchronization).

Threads-per-block double karne se throughput hamesha double ho jaata hai

False. Bade blocks SM ke fixed budget ka zyada hissa consume karte hain, isliye unke kam blocks fit hote hain; occupancy gir sakti hai aur kam latency hide hoti hai. Throughput aksar 128–256 threads ke aas-paas peak karta hai.

Same warp mein do threads ek if ke alag branches le sakte hain aur dono full speed par run karte hain

False. Ek warp sabhi 32 lanes ke liye ek instruction issue karta hai; divergence par hardware lanes ko mask off karta hai aur dono paths serially chalata hai, isliye divergence se cycles ki cost lagti hai.

Alag-alag blocks ke threads __syncthreads() call karke ek doosre ka wait kar sakte hain

False. Yeh barrier block-scoped hai. Ek kernel ke andar cross-block waiting se deadlock ka risk hai kyunki jinhe tum wait kar rahe ho woh abhi resident bhi nahi ho sakte.

Ek grid sirf programmer ki convenience ke liye 1D, 2D, ya 3D ho sakta hai

True. Dimensionality sirf index arithmetic ko shape karta hai; hardware ultimately sab kuch flatten kar deta hai. 2D grid sirf image/matrix mapping ko naturally padhne laayak banata hai.

Global thread ID threadIdx + blockIdx hai

False. Yeh hai blockIdx * blockDim + threadIdx. Tumhe apne local offset add karne se pehle saare complete blocks (har ek blockDim size ke) ko skip karna padega.

Agar block size 32 ka multiple nahi hai, toh kuch warp lanes waste ho jaate hain

True. Warps hamesha 32 lanes ke hote hain; ek 40-thread block 2 warps (64 lanes) ban jaata hai jismein 24 idle lanes hoti hain jo phir bhi scheduling slots occupy karti hain.

Error dhundo

int i = threadIdx.x; ko multi-block launch ke global index ke roop mein use karna

Error: yeh sirf block ke andar local position deta hai, isliye har block elements 0…blockDim-1 ko re-index karta aur ek doosre ko overwrite karta. Sahi: blockIdx.x * blockDim.x + threadIdx.x.

N = 1000, B = 256 ke liye blocks = N / threadsPerBlock launch karna

Error: integer division 3 blocks (768 threads) deta hai aur elements 768–999 drop kar deta hai. Ceiling division use karo (N + B - 1) / B → 4 blocks.

if (i < N) guard hatana "kyunki math theek hai"

Error: last block mein aksar extra threads hote hain jinki global ID N se zyada hoti hai; guard ke bina woh out-of-bounds memory read/write karte hain. Guard mandatory hai jab bhi N block size ka exact multiple na ho.

Block 0 flag[0] = 1 set karta hai aur baaki blocks ek kernel ke andar while(flag[0]==0) par spin karte hain

Error: blocks ke co-reside hone ki guarantee nahi hai, isliye spinning blocks har SM slot occupy kar sakte hain jabki block 0 kabhi schedule hi na ho → deadlock. Alag kernel launches ya cooperative groups use karo.

2D image ke liye 16×16 blocks se tile karte waqt dim3 blocks((N)/16, (M)/16)

Error: plain division ragged edge tiles truncate kar deta hai. (N+15)/16 aur (M+15)/16 use karo taaki partial border tiles ko bhi ek block mile (in-kernel bounds checks ke saath).

Matrix row ke liye blockIdx.x aur column ke liye blockIdx.y use karna

Error (convention ke hisaab se): x horizontal axis hai → column, y vertical hai → row. Unhe swap karne se coalesced access toot jaati hai kyunki adjacent threadIdx.x ko adjacent columns se map hona chahiye.

Block size 1000 threads set karna ek aise GPU par jiska max 1024 hai "max parallelism ke liye"

Judgment mein error: 1000, 32 ka multiple nahi hai (lanes waste hote hain) aur itna bada block doosron ko crowd out karta hai, occupancy hurt karta hai. 256 jaisa koi 32-multiple choose karo.

Why questions

Blocks independent kyun hone chahiye?

Taaki runtime unhe kisi bhi order mein kisi bhi number of SMs par distribute kar sake — yahi cheez identical code ko 4-SM laptop chip se 100-SM datacenter GPU tak scale karaati hai (yeh CUDA Programming Model ka core promise hai).

Fast shared memory per block kyun hoti hai, per grid kyun nahi?

Sirf wahi threads jo physically ek SM par co-located hain ek on-chip scratchpad cheaply share kar sakte hain; ek grid kai SMs par span karta hai, isliye grid-wide shared memory ke liye slow global-memory traffic ki zaroorat padti.

Kai chhote blocks few giant ones se behtar kyun hote hain?

Zyada independent blocks SM ko zyada kaam dete hain taaki jab bhi ek block memory par stall ho, woh swap in kar sake, isliye latency hidden rehti hai aur compute units busy rehte hain.

Warp size block size ko influence kyun karna chahiye?

Hardware 32-lane warp per instructions issue karta hai, isliye ek block jo 32 ka multiple nahi hai woh apne last warp mein dead lanes chhodta hai, scheduling capacity waste karta hai.

2D indexing row * width + col mein flatten kyun hoti hai?

Memory physically 1D aur row-major stored hoti hai, isliye flattened index linear address reconstruct karta hai; formula mein width (na ki height) daalna hi ek row ko contiguous rakhta hai.

Hum sirf ek bada block launch karke poora problem kyun nahi cover kar sakte?

Ek block hardware-capped hota hai (jaise 1024 threads) aur ek single SM par fit hona chahiye, isliye ek block kabhi million-element array span nahi kar sakta — grid isliye exist karta hai taaki ek SM ki reach se pare bhi tile ho sake.

Do kernel launches mein split karna global barrier ki tarah kyun kaam karta hai?

Ek kernel launch tab complete nahi hota jab tak uske saare blocks khatam na ho jaayein, isliye agla launch guarantee ke saath saare results dekh sakta hai — koi in-kernel spinning ke bina implicit whole-grid synchronization.

Occupancy sirf thread count se nahi, shared memory se bhi limited kyun hai?

Ek SM ko har resource simultaneously satisfy karna hota hai; agar har block shared memory hog kare, toh SM ka woh budget thread slots khatam hone se pehle hi khatam ho jaata hai, isliye tightest constraint jeet jaati hai (dekho GPU Occupancy).

Edge cases

Kya hota hai jab N ek block ke thread count se chhota ho (N = 50, B = 256)?

Tum phir bhi 1 block launch karte ho; if (i < N) guard baaki 206 threads ko idle kar deta hai. Wasteful hai lekin correct hai — guard hi ise safe rakhta hai.

Agar N block size ka exact multiple ho — kya guard phir bhi zaroori hai?

Us specific launch ke liye strictly correctness ke liye nahi, lekin phir bhi rakho: iska cost almost kuch nahi aur yeh tumhe us pal protect karta hai jab N non-multiple mein change ho.

G blocks × B threads ki ek launch mein largest global ID kya ho sakta hai?

(G-1)*B + (B-1) = G*B - 1, yani total threads minus one. Agar yeh N-1 se zyada ho, toh extra threads exactly wohi hain jinhe guard reject kare.

32 threads ka ek block jahan 31 if hit karein aur 1 else hit kare — cost kya hai?

Single warp ke liye dono branches serially execute hote hain: taken lanes run karte hain jabki baaki mask hote hain, phir swap hote hain. Tum dono paths ki cost pay karte ho chahe sirf ek lane ko else ki zaroorat thi.

Agar ek grid un blocks launch kare jo ek SM ek saath hold kar sake se zyada hain, toh launch fail ho jaata hai?

Nahi — extra blocks simply queue ho jaate hain aur schedule hote hain jab earlier blocks retire ho jaayein aur resources free karein. Resource limits concurrent residency cap karte hain, launch hone wale total blocks ki count nahi.

Zero-size launch: kya hoga agar tum blocks = 0 compute karo?

Kernel kuch nahi karta (koi thread run nahi karta), jo usually N = 0 se bug hota hai jo slip through ho gaya — launch se pehle hamesha ceiling division aur nonzero problem size ensure karo.

Recall Quick self-test

Ek sabse gehri idea jo saath le jaao ::: globalID = blockIdx*blockDim + threadIdx, blocks independent aur unordered hote hain, aur __syncthreads() kabhi bhi sirf ek block ke andar threads ko bind karta hai.