Visual walkthrough — Thread blocks and grids

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

Hum poore raaste sirf EK sawaal ka jawaab de rahe hain:

"Mere paas hazaaron workers hain jo teams mein bante hain. Har worker kaise pata lagaata hai ki ek lambi to-do list mein se kaun sa ek job sirf uska hai — bina kisi doosre worker ke wahi job choose kiye?"

---

Step 1 — Problem ko seedha bichchhaao: ek lambi to-do list

KYA. Apne data ko ek seedhi row of boxes ki tarah socho, numbered $0,1,2,\dots ,N-1$. Yahi array C hai vector addition mein (C[i] = A[i] + B[i]). Har box ek job hai.

YEH YAHIN SE KYUN SHURU KAREIN. Teams ki baat karne se pehle, hume woh cheez dekhni hai jo divide ho rahi hai. GPU ki memory fundamentally ek 1D line of addresses hai — ek 2D image bhi ek lambi strip ki tarah store hoti hai. Isliye haari saari indexing ka target ek single number hota hai, is line mein ek index.

PICTURE. Neeche ki blue strip to-do list hai. Har cell ek element hai.

---

Step 2 — List ko equal teams (blocks) mein kaatna

KYA. Hum lambi strip ko equal chunks mein kaatenge. Har chunk ek thread block hai — workers ki ek team. Team size ko $B$ bulao (code mein: blockDim.x). Teams ko $0,1,2,\dots$ se number kiya jaata hai aur team ka number $k$ hai (code mein: blockIdx.x).

KAATEIN KYUN? GPU saare workers ko ek hi pile mein nahi daal sakta — hardware teams ko alag-alag Streaming Multiprocessors par force karta hai. List ko equal blocks mein kaatne ka matlab hai ki har team ko kaam ka ek identical slice milta hai, aur teams ko jo bhi processor free ho usse de sakte hain. Yahi wajah hai ki teams ka equal size hona zaroori hai: equal-size chunks agle Step 4 ki arithmetic ko ek clean multiplication banate hain.

PICTURE. Wahi strip, ab yellow dividers se size $B$ ke blocks mein kaati gayi. Notice karo ki block $k$ strip par ek specific jagah se start hota hai — yeh baat yaad rakhna.

---

Step 3 — Har worker ka local number

KYA. Team ke andar, workers khud ko $0$ se count karte hain: worker $0$, worker $1$, worker $B-1$ tak. Yeh within-team number $t$ hai (code mein: threadIdx.x).

EK LOCAL NUMBER KYUN, GLOBAL KYUN NAHI? Kyunki har team identical hoti hai, hardware har worker ko wahi chhota counter deta hai jo har team mein $0$ par reset ho jaata hai. Yeh sasta aur uniform hai. Cost yeh hai: team $0$ ka worker $t$ aur team $1$ ka worker $t$ ka same local number hai lekin unhe alag jobs karni hain. Steps 4–5 bilkul isi clash ko fix karte hain.

PICTURE. Do teams mein zoom karo. Dono mein $t=2$ label waala worker hai (red). Woh strip par alag-alag boxes point karte hain. Akela local number ambiguous hai.

---

Step 4 — Apne se pehle ke jobs count karo

KYA. Team $k$ pehli team nahi hai. Teams $0,1,\dots ,k-1$ usse pehle aayi hain — woh $k$ complete teams hain, har ek $B$ jobs ki malik. Toh team $k$ ke shuru hone se pehle, exactly $k\times B$ jobs pehle hi le li gayi hain. Woh product aapki team ke slice ke front ka offset hai.

MULTIPLY KYUN? Multiplication hai hi equal groups ki repeated addition. Hamare paas $k$ equal groups of size $B$ hain. $B+B+\dots$ ($k$ baar) add karna exactly $k\times B$ hai. Yeh clean product tabhi possible hai jab Step 2 ne blocks ko equal banaya — woh choice ka yahi payoff hai.

$$\underset{\text{mere aage kitni puri teams hain}}{\underbrace{k}}\times \underset{\text{un teams mein se har ek ke jobs}}{\underbrace{B}}=\underset{\text{meri team ke shuru se pehle claim kiye gaye jobs}}{\underbrace{kB}}$$

PICTURE. Strip jisme green bracket boxes $0$ se $kB-1$ tak span karta hai — woh har job jo pehli teams ka hai. Team $k$ ka slice bracket ke bilkul baad, box $kB$ par shuru hota hai.

---

Step 5 — Team ke andar apni seat add karo: formula ka janam

KYA. Aap team $k$ ke andar worker $t$ ho. Aapki team ka slice box $kB$ par shuru hota hai (Step 4). Aap us slice mein $t$ seats andar baithe ho. Toh aapka job hai:

$$i=\underset{\begin{array}{c}\text{meri}\\ \text{team ke slice ki shuruaat}\end{array}}{\underbrace{k\times B}}+\underset{\begin{array}{c}\text{slice ke andar}\\ \text{meri seat}\end{array}}{\underbrace{t}}$$

CUDA ke naamon se likha:

$$i=\underset{k}{\underbrace{\text{blockIdx.x}}}\times \underset{B}{\underbrace{\text{blockDim.x}}}+\underset{t}{\underbrace{\text{threadIdx.x}}}$$

$t$ ADD KYUN KAREIN? Multiplication aapko aapki team ke slice ke darwaze tak pahunchaa deta hai. Addition aapko $t$ steps andar le jaata hai aapki apni seat tak. Multiply-phir-add poori kahani hai: apni team par jump karo, phir apni seat tak step karo.

YEH COLLISION-FREE KYUN HAI? Do alag $(k,t)$ pairs kabhi same $i$ par nahi land kar sakte. Kyunki $0\le t<B$ hai, isliye $kB+t$ ki value strictly $[kB,kB+B-1]$ ke andar aati hai — yeh exactly team $k$ ke boxes ka block hai aur kisi doosri team ka nahi. Toh har worker ko ek distinct job milta hai, aur $[0,GB-1]$ mein har job ko exactly ek worker milta hai. (Yahan $G$ = gridDim.x teams ki sankhya hai.)

PICTURE. Worker $(k=2,t=3)$ block size $B=4$ ke saath: green arrow box $8=2\times 4$ par jump karta hai, phir red arrow $+3$ step karke box $11$ par pahunchta hai. Yahi uska job hai.

---

Step 6 — Leftover-worker edge case (kyun if guard exist karta hai)

KYA. Real problems rarely evenly divide hoti hain. Maano $N=10$ jobs hain aur $B=4$. Kitni teams? Hum $3$ teams use nahi kar sakte (woh sirf $\mathrm{12...}$ wait — $3\times 4=12\ge 10$? Chalo honest small case use karein): $N=10$, $B=4$ ko $\lceil 10/4\rceil =3$ teams chahiye $3\times 4=12$ seats cover karte hue. Lekin sirf $10$ jobs hain. Do workers ka koi job nahi hai.

YEH KYUN HOTA HAI. Humne equal teams ki poori sankhya launch ki, aur aakhri team list ke end ke baad chali gayi. Woh extra workers $i=10$ aur $i=11$ calculate kar lenge — woh boxes exist hi nahi karte. Wahan read/write karna illegal memory access hai → crash ya garbage.

FIX. Har worker memory touch karne se pehle if (i < N) check karta hai. $i\ge N$ waale workers kuch nahi karte. CUDA launch team count ceiling division se calculate karta hai:

$$G=\lceil \frac{N}{B}\rceil =\frac{N+B-1}{B}\text{(integer division)}$$

$+(B-1)$ kisi bhi nonzero remainder ko agli poori team tak push kar deta hai, toh koi bhi job uncovered nahi rehta — hum hamesha upar round karte hain, neeche nahi.

PICTURE. $N=10$ ke upar teen teams. Green boxes real jobs hain; do red hatched boxes phantom seats hain jinke workers guard ke zariye chup kara diye jaate hain.

---

Step 7 — Do dimensions: wahi idea, do baar, phir flatten

KYA. Ek image (ya matrix) ke liye, natural picture pixels ka ek 2D field hai. Hum use 2D blocks se tile karte hain. Har thread apna column aur row dhundta hai bilkul usi jump-then-step rule se, har axis ke liye ek baar:

$$\text{col}=\underset{k_x}{\underbrace{\text{blockIdx.x}}}\times \underset{B_x}{\underbrace{\text{blockDim.x}}}+\underset{t_x}{\underbrace{\text{threadIdx.x}}}\text{row}=\underset{k_y}{\underbrace{\text{blockIdx.y}}}\times \underset{B_y}{\underbrace{\text{blockDim.y}}}+\underset{t_y}{\underbrace{\text{threadIdx.y}}}$$

DO COPIES KYUN? 2D grid sirf 1D idea hai jo independently $x$ (horizontal) aur $y$ (vertical) dono mein chalta hai. $x$ columns ke across count karta hai, $y$ rows ke neeche count karta hai — yeh convention match karta hai ki hum matrices kaise draw karte hain (rows neeche jaati hain).

Phir flatten karo. Memory abhi bhi ek 1D line hai (Step 1!). Width $W$ ki image mein $(\text{row},\text{col})$ par ek pixel is jagah par rehta hai:

$$\text{idx}=\underset{\text{mere upar ki puri rows skip karo}}{\underbrace{\text{row}\times W}}+\underset{\text{apni row ke across step karo}}{\underbrace{\text{col}}}$$

Yeh row-major order hai: apne box tak pahunchne ke liye, mere upar ki saari complete rows skip karo ($\text{row}$ rows of $W$ boxes har ek — ek multiply), phir apni khud ki row mein $\text{col}$ step karo (ek add). Yeh Step 4 + Step 5 hi hai, ek level upar.

PICTURE. Ek 2D pixel field blocks se tiled; ek thread ka $(\text{row},\text{col})$ do arrows se milta hai, phir ek single index mein row-major flattening.

Recall Flattening bhi ek multiply-then-add kyun hai?

Kyunki 2D memory ek jhooth hai — yeh ek lambi line ki tarah store hoti hai. ::: Row $r$ tak pahunchne ke liye aapko $r$ poori rows of $W$ elements skip karni hongi ($r\times W$, the jump), phir apni row mein col move karo (the step). Block formula jaisi bilkul wahi shape.

---

Ek-picture summary

Upar sab kuch ek sentence hai: jo pehle aaya usse jump karo, phir apni seat tak step karo — aur har dimension mein ek baar karo. Neeche ka figure teeno jump-then-step moves ko (1D global ID, leftover guard, aur 2D→flat) ek single diagram mein stack karta hai.

Recall Saari walkthrough simple words mein (Feynman retelling)

Socho ek badi to-do list hai, ek job per box, boxes zero se numbered hain. Hum saare workers ko ek saath nahi daal sakte, isliye hum list ko equal teams mein kaatate hain aur teams ko number karte hain. Har team ke andar workers khud ko zero se count karte hain — lekin woh local number har team mein repeat hota hai, isliye woh akele job number nahi ban sakta. Yahi trick hai: apna job dhundhne ke liye, pehle woh saare jobs count karo jo aapke aage waali teams ke hain. Har pehli team ke jobs ki sankhya same hai, isliye woh count sirf (aapki team number) times (team size) hai — ek multiplication, kyunki woh equal groups add ho rahe hain. Yeh aapko aapki team ke slice ke darwaze par pahunchaata hai. Ab apne local seat number se step andar karo — ek addition. Multiply to jump, add to step: yahi global ID hai, aur yeh kabhi collide nahi karta kyunki aapki seat aapki team ke slice ke andar rehti hai. Agar list evenly split nahi hoti, toh aakhri team mein khaali seats hain jo end ke baad point karti hain, isliye har worker check karta hai "kya mera job real hai?" memory touch karne se pehle, aur hum team count ko upar round karte hain taaki kuch miss na ho. Images ke liye hum wohi jump-then-step do baar chalate hain — columns ke liye across, rows ke liye down — phir real 1D memory line par flatten karte hain poori rows skip karke (multiply) aur apni khud ki row mein step karke (add). Wahi idea, poore raaste.

Related vault topics: Warp Execution · Shared Memory · Thread Synchronization · GPU Occupancy · Memory Coalescing · Parallel Algorithm Design.