Thread blocks and grids: Worked examples — Hinglish deep dive

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

Yeh page parent topic ka drill hall hai. Hum block/grid machinery ke har tarah ke situation ko dhundhenge — clean divisions, ragged remainders, empty inputs, 2D tiling, occupancy walls — aur har ek ko zero se work karenge. Agar neeche koi bhi word naya lagta hai, toh parent note ne usse pehle build kiya hai; yahan hum use use karte hain jab tak yeh muscle memory na ban jaaye.

Kuch bhi karne se pehle, ek reminder us single formula ka jis par sab kuch tika hua hai:

Figure s01 dekho. Har row ek block hai; left side ke row labels ("block 0"…"block 3") unhe naam dete hain. Red seat (labelled globalID = 15) ek specific thread hai. Uska global number nikaalane ke liye hum red vertical arrow follow karte hain — uske upar ke do full rows ko jump karte hue (yaani $blockIdx \times blockDim = 2 \times 6 = 12$ ) — aur phir red horizontal arrow follow karte hain apne column tak step karte hue ( $threadIdx = 3$ ). Yeh jumping-then-stepping ka ek kaam hi poora chapter hai.

Scenario matrix

Jo bhi problem tumhe milegi woh in cells mein se ek mein hogi. Baad ke worked examples mein se har ek mein (Cell A2) jaisi tag hogi taaki tum dekh sako ki poori grid cover ho rahi hai.

#	Case class	Tricky kya hai	Covered by
A1	Exact fit — $N$ block size se divisible	Koi leftover thread nahi; guard kabhi fire nahi karta	Example 1
A2	Ragged fit — $N$ divisible nahi	Last block mein idle threads hain; ceiling + guard chahiye	Example 2
A3	Degenerate: $N = 0$	Empty input	Example 3
A4	Degenerate: $N = 1$	Single element, phir bhi ek poora block	Example 3
A5	$N <$ one block	Grid bilkul 1 block tak round up hoti hai	Example 3
B1	2D tiling — exact	x,y se row/col; row-major flatten karo	Example 4
B2	2D tiling — ragged in ONE dim	Guard sirf ek axis par fire karta hai (exact x, ragged y)	Example 5
B3	2D tiling — ragged in BOTH dims	Guard row aur col dono par corner mein fire karta hai	Example 6
C1	Limiting behaviour — block size $\to$ max (1024)	Kam, mote blocks; occupancy risk	Example 7
C2	Occupancy wall — shared memory blocks/SM ko limit karta hai	Memory, threads nahi, bottleneck hai	Example 8
D1	Word problem — real workload sizing	"2M pixels process karo" ko ek launch mein translate karo	Example 9
D2	Exam twist — reverse the formula	globalID diya hai, block & thread recover karo	Example 10

Worked Examples

Worked example Example 1 — Exact fit

(Cell A1) $N = 1024$ elements ka vector add, block size $B = 256$ . Kitne blocks, aur kya guard if (i < N) kabhi matter karta hai?

Forecast: aage padhne se pehle block count aur kya koi thread waste hogi, yeh guess karo.

Ceiling division se block count compute karo. $blocks = ⌈ \frac{N}{B} ⌉ = \frac{N + B - 1}{B} = \frac{1024 + 255}{256} = \frac{1279}{256} = 4$ Yeh step kyun? Integer division truncate karta hai, isliye hum pehle $B - 1$ add karte hain taaki round-up force ho — warna ek partial block drop ho jaata hai aur kuch data process nahi hota.
Total launched threads count karo. $4 \times 256 = 1024$ . Yeh step kyun? Dekhne ke liye ki humne data se zyada threads launch ki hain ya nahi. Yahan $1024 = N$ exactly hai.
Guard check karo. Last thread ka $globalID = 1023 < 1024$ hai, toh i < N hamesha true hai. Yeh step kyun? Jab perfectly fit hota hai, koi thread idle nahi hoti — guard kabhi fire nahi karta, lekin safety ke liye hum ise rakhte hain.

Verify: launched threads $= N$ , toh waste $= 1024 - 1024 = 0$ . Zero idle threads. ✓

Worked example Example 2 — Ragged fit

(Cell A2) Wahi kernel, ab $N = 1000$ , $B = 256$ . Kitne blocks? Last block mein kitni threads idle hain? Pehli idle thread ka global index kya hai?

Forecast: kya 3 blocks kaafi honge? Wasted threads ki count guess karo.

Ceiling division. $blocks = \frac{1000 + 255}{256} = \frac{1255}{256} = 4$ Yeh step kyun? $3 \times 256 = 768 < 1000$ , toh 3 blocks elements 768–999 miss kar dete hain. Hume 4 tak round up karna hi hai.
Total launched threads. $4 \times 256 = 1024$ . Yeh step kyun? $N = 1000$ se compare karo taaki excess pata chale.
Idle threads. $1024 - 1000 = 24$ threads ke paas koi data nahi hai. Yeh step kyun? Yeh block 3 (last wala) mein rehti hain. Unke global indices $1000..1023$ hain — exactly woh threads jinki guard fire hoti hai.
Pehli idle thread ki identity. globalID $= 1000$ . Formula ko reverse karo: $blockIdx = ⌊ 1000/256 ⌋ = 3, threadIdx = 1000 - 3 \times 256 = 1000 - 768 = 232$ Yeh step kyun? Exactly dikhata hai kahan guard fire karna shuru karta hai: block 3 mein seat 232 aur uske baad sab.

Verify: block 3 globalIDs $768..1023$ cover karta hai; valid wale hain $768..999$ (yeh $232$ threads kaam karti hain) aur $1000..1023$ idle (yeh $24$ hain). $232 + 24 = 256 = B$ . ✓

Worked example Example 3 — Degenerate inputs

$N \in {0, 1, 100}$ (Cells A3, A4, A5) $B = 256$ ke saath, teen tiny/empty cases ke liye launch work out karo.

Forecast: kya $N = 0$ zero blocks launch karta hai ya ek? Step 1 se pehle guess karo.

Case $N = 0$ (empty input). $blocks = \frac{0 + 255}{256} = \frac{255}{256} = 0$ Yeh step kyun? $0/256$ ki ceiling $0$ hai, toh koi block launch nahi hota — kernel kuch nahi karta, jo sahi hai. Koi guard ki zaroorat nahi kyunki koi thread exist hi nahi karti.
Case $N = 1$ (single element). $blocks = \frac{1 + 255}{256} = \frac{256}{256} = 1$ Yeh step kyun? Ek element ke liye bhi 256 threads ka ek poora block chahiye; unमें se 255 idle hain. globalID $0$ waali thread kaam karti hai, threads $1..255$ ki guard fire hoti hai.
Case $N = 100$ (ek block se chhota). $blocks = \frac{100 + 255}{256} = \frac{355}{256} = 1$ Yeh step kyun? $1$ se $256$ elements tak kuch bhi exactly 1 block tak round up hota hai — grid mein kabhi ek block ka fraction nahi hota.

Verify: blocks $(0) = 0$ , blocks $(1) = 1$ , blocks $(100) = 1$ . $N = 1$ ke liye idle threads: $256 - 1 = 255$ ; $N = 100$ ke liye: $256 - 100 = 156$ . ✓

Worked example Example 4 — 2D tiling, exact

(Cell B1) Ek $64 \times 64$ image ( $N = 64$ columns, $M = 64$ rows). Block $16 \times 16$ hai. Grid dimensions nikalo, aur thread ka global (row, col) aur flat memory index nikalo jiska $blockIdx = (x = 1, y = 2)$ aur $threadIdx = (x = 5, y = 3)$ hai. (Yaad raho x = col, y = row.)

Forecast: compute karne se pehle row aur col guess karo.

Har dimension mein grid size. $gridDim.x = \frac{N}{16} = \frac{64}{16} = 4, gridDim.y = \frac{M}{16} = \frac{64}{16} = 4$ Yeh step kyun? Hum 64-wide, 64-tall image ko 16×16 tiles se tile karte hain — 4 tiles across, 4 neeche. Yeh evenly divide hota hai, toh koi ragged edge nahi.
Column (x-direction). Yahan $blockIdx.x = 1$ aur $threadIdx.x = 5$ hai: $col = blockIdx.x \times blockDim.x + threadIdx.x = 1 \times 16 + 5 = 21$ Yeh step kyun? Wahi jump-then-step idea, lekin horizontally. Figure s02 mein red tile dekho: block column 1 matlab 16 pixels skip karo, phir 5 aur step karo. x component column drive karta hai.
Row (y-direction). Yahan $blockIdx.y = 2$ aur $threadIdx.y = 3$ hai: $row = blockIdx.y \times blockDim.y + threadIdx.y = 2 \times 16 + 3 = 35$ Yeh step kyun? Vertical version; block row 2 matlab 32 rows skip karo, phir 3 neeche. y component row drive karta hai — inhe kabhi swap mat karo.
1D memory address par flatten karo (row-major). $idx = row \times width + col = 35 \times 64 + 21 = 2240 + 21 = 2261$ Yeh step kyun? Memory ek seedhi line hai, grid nahi. Row-major matlab "pehle full rows walk karo, phir across." Har full row width $= N = 64$ elements wide hoti hai.

Verify: col $= 21 < 64$ , row $= 35 < 64$ — dono image ke andar hain, guard pass. Flat index $2261 < 64 \times 64 = 4096$ . ✓

Worked example Example 5 — 2D tiling, EK dimension mein ragged

(Cell B2) Image $64 \times 70$ hai ( $N = 64$ columns — 16 ke liye exact, $M = 70$ rows — 16 ka multiple nahi). Block $16 \times 16$ . Grid nikalo, aur blocks ki bottom row mein identify karo ki kaun si threads guard survive karti hain.

Forecast: kaun sa axis ragged hai, x ya y? Kitni rows ki idle threads appear hoti hain?

Ceiling ke saath grid dimensions dono par, chahe sirf ek ragged ho. $gridDim.x = ⌈ \frac{N}{16} ⌉ = \frac{64 + 15}{16} = \frac{79}{16} = 4 (exact), gridDim.y = ⌈ \frac{M}{16} ⌉ = \frac{70 + 15}{16} = \frac{85}{16} = 5$ Yeh step kyun? $64/16 = 4$ cleanly divide hota hai, toh x exact hai; $70/16 = 4.375$ round up hota hai 5 tak, toh y ragged hai. Ceiling dono par apply karna safe hai — exact axis par yeh no-op hai.
Padded region. Grid $64 \times 80$ thread-slots cover karta hai; image $64 \times 70$ hai. Yeh step kyun? Sirf rows 70–79 ki bottom strip padding hai. Guard sirf row condition par fire karta hai yahan — har column valid hai kyunki x exact hai.
Blocks ki bottom row, $blockIdx.y = 4$ . Yeh blocks rows $64..79$ cover karte hain. Valid rows: $64..69$ (yeh 6 hain), invalid rows: $70..79$ (yeh 10 hain). Yeh step kyun? Aisi har block mein, $threadIdx.y \leq 5$ waale threads (rows $64..69$ ) pass karte hain; $threadIdx.y \geq 6$ fail karte hain. col guard kabhi nahi trip karta.
Bottom block per working threads. $16 cols \times 6 rows = 96$ kaam karte hain; $256 - 96 = 160$ ki guard fire hoti hai — sab kuch single ragged (y) axis ki wajah se. Yeh step kyun? Yeh key single-dim lesson hai: guard if (row < M && col < N) ek failing sub-condition se protect kar raha hai, do se nahi.

Verify: gridx $= 4$ , gridy $= 5$ ; bottom block valid rows $= 70 - 64 = 6$ ; working $= 16 \times 6 = 96$ ; idle $= 256 - 96 = 160$ . ✓

Worked example Example 6 — 2D tiling, DONO dimensions mein ragged

(Cell B3) Image $30 \times 20$ hai ( $N = 30$ columns, $M = 20$ rows). Block $16 \times 16$ . Grid nikalo, aur identify karo ki bottom-right corner block mein kaun si threads real kaam karti hain.

Forecast: kya grid $2 \times 2$ blocks ki hogi? Us corner block mein kitni threads waste hogi?

Ceiling ke saath grid dimensions. $gridDim.x = ⌈ \frac{N}{16} ⌉ = \frac{30 + 15}{16} = \frac{45}{16} = 2, gridDim.y = ⌈ \frac{M}{16} ⌉ = \frac{20 + 15}{16} = \frac{35}{16} = 2$ Yeh step kyun? Na $N = 30$ na $M = 20$ 16 se divide hota hai, toh dono directions round up hoti hain — $32 \times 32$ padded region cover karne waale blocks ka $2 \times 2$ grid.
Total covered vs. real. Grid $32 \times 32 = 1024$ thread-slots cover karta hai; image mein $N \times M = 30 \times 20 = 600$ real pixels hain. Yeh step kyun? $1024 - 600 = 424$ baaki thread-slots ko guard if (row < M && col < N) se kill karna hoga.
Bottom-right block $blockIdx = (x = 1, y = 1)$ . Iske columns $col = 16..31$ span karte hain, rows $row = 16..31$ . Yeh step kyun? Hum check karte hain ki in mein se kaun se $30 \times 20$ ke andar hain. Valid cols: $16..29$ (yeh 14 hain), valid rows: $16..19$ (yeh 4 hain).
Us block mein working threads. $14 \times 4 = 56$ kaam karte hain; $256 - 56 = 200$ guard se out ho jaate hain. Yeh step kyun? Yeh corner hai jahan dono guard sub-conditions simultaneously bite karti hain — worst-case ragged cell, Example 5 se contrast jahan sirf ek bit karti thi.

Verify: guard require karta hai col < 30 AND row < 20. Block (x=1, y=1) mein: cols $16..29$ = 14 values, rows $16..19$ = 4 values, product $56$ . Idle $= 256 - 56 = 200$ . ✓

Worked example Example 7 — Limiting behaviour: block size maximum par

(Cell C1) $N = 3000$ elements. $B = 256$ vs. hardware max $B = 1024$ par launches compare karo. Kitne blocks each, aur kitni idle threads each?

Forecast: kya mota 1024 block zyada ya kam threads waste karta hai?

Small blocks $B = 256$ . $blocks = \frac{3000 + 255}{256} = \frac{3255}{256} = 12, launched = 12 \times 256 = 3072, idle = 72$ Yeh step kyun? Compare karne ke liye baseline.
Max blocks $B = 1024$ . $blocks = \frac{3000 + 1023}{1024} = \frac{4023}{1024} = 3, launched = 3 \times 1024 = 3072, idle = 72$ Yeh step kyun? Interesting — yahan same total (3072) aur same waste (72) hai, lekin kahin kam, kahin zyada mote blocks hain.
Limit interpret karo. Sirf 3 giant blocks ke saath, agar 2 memory par stall ho jaayein, latency hide karne ke liye sirf 1 bachta hai — bekar occupancy. 12 small blocks ke saath overlap karne ke liye kaafi zyada hai. Yeh step kyun? Yeh exactly woh "zyada threads per block ≠ tez" galti hai parent note se, jo ab concrete ho gayi.

Verify: blocks(256) $= 12$ , launched $= 3072$ , idle $= 72$ ; blocks(1024) $= 3$ , launched $= 3072$ , idle $= 72$ . ✓

Worked example Example 8 — Occupancy wall: threads nahi, memory

(Cell C2) Ek SM mein 2048 threads aur 96 KB shared memory hai. Tumhara block 256 threads aur 18 KB shared memory use karta hai. SM per kitne blocks fit hote hain, aur occupancy kya hai?

Forecast: kya thread budget bottleneck hai ya memory budget?

Thread budget se allowed blocks. $⌊ \frac{2048}{256} ⌋ = 8 blocks$ Yeh step kyun? Capacity limits floor (round down) use karti hain: tum sirf whole blocks fit kar sakte ho, toh ek fractional block count nahi hoti. Thread budget akele maximum 8 allow karta hai.
Shared-memory budget se allowed blocks. $⌊ \frac{96}{18} ⌋ = 5 blocks$ Yeh step kyun? Shared memory ek alag budget hai, woh bhi floored. $96/18 = 5.33$ , toh sirf 5 whole blocks fit hote hain — fractional 0.33 unusable hai.
Effective blocks = min(8, 5) = 5. Active threads $= 5 \times 256 = 1280$ . Yeh step kyun? Hardware sabhi resource ceilings ka minimum leta hai — memory, threads nahi, hume throttle kar rahi hai.
Occupancy. $occupancy = \frac{1280}{2048} \approx 0.625 = 62.5%$ Yeh step kyun? Occupancy = active threads ÷ max threads. Hum SM ka 37.5% idle chhod rahe hain shared-memory pressure ki wajah se.

Verify: thread limit $= 8$ , memory limit $= 5$ , effective $= 5$ ; active threads $= 1280$ ; occupancy $= 1280/2048 = 0.625$ . ✓

Worked example Example 9 — Word problem: real launch sizing

(Cell D1) Tumhe 1920 × 1080 pixels ( $N = 1920$ columns, $M = 1080$ rows) ki photo par filter apply karna hai. $16 \times 16$ blocks use karke, total kitne blocks launch honge, aur total kitni threads? Idle fraction kya hai?

Forecast: guess karo kya dono dimensions evenly divide hoti hain.

Blocks per dimension (dono ceiling). $gridDim.x = ⌈ \frac{N}{16} ⌉ = 120, gridDim.y = ⌈ \frac{M}{16} ⌉ = \frac{1080 + 15}{16} = \frac{1095}{16} = 68$ Yeh step kyun? $1920/16 = 120$ exactly (clean x), lekin $1080/16 = 67.5$ → $68$ tak rounds hota hai (ragged y).
Total blocks. $120 \times 68 = 8160$ blocks. Yeh step kyun? Grid apne do dimensions ka product hai.
Total launched threads. $8160 \times 256 = 2088960$ . Yeh step kyun? Har block mein $16 \times 16 = 256$ threads hoti hain.
Real pixels & idle threads. Real $= N \times M = 1920 \times 1080 = 2073600$ . Idle $= 2088960 - 2073600 = 15360$ . Yeh step kyun? Idle threads sirf $y$ mein padded strip se aati hain (rows 1080–1087): $8 \times 1920 = 15360$ . Perfectly match karta hai.

Verify: grid $= 120 \times 68$ ; threads $= 2088960$ ; real pixels $= 2073600$ ; idle $= 15360 = 8 \times 1920$ . ✓

Worked example Example 10 — Exam twist: formula ko reverse karo

(Cell D2) Ek thread $globalID = 745$ report karta hai 1D launch mein jahan $blockDim = 128$ hai. Yeh kis block mein hai, aur kaun sa seat (threadIdx)? Phir forward confirm karo.

Forecast: divide karne se pehle block number guess karo.

Integer division se block index recover karo. $blockIdx = ⌊ \frac{745}{128} ⌋ = 5 (kyunki 5 \times 128 = 640 \leq 745 < 768 = 6 \times 128)$ Yeh step kyun? Formula hai $globalID = blockIdx \times B + threadIdx$ . $B$ se divide karke remainder drop karne par threadIdx strip ho jaata hai, sirf block bachta hai.
Remainder (modulo) se thread index recover karo. $threadIdx = 745 mod 128 = 745 - 5 \times 128 = 745 - 640 = 105$ Yeh step kyun? Poore blocks (har $B = 128$ wide) hataane ke baad bacha hua exactly current block ke andar seat hai. Toh yeh thread block 5 ke seat 105 par baitha hai.
Forward-check. Dono recovered values original formula mein plug karo: $blockIdx \times B + threadIdx = 5 \times 128 + 105 = 640 + 105 = 745 ✓$ Yeh step kyun? Agar forward formula original globalID $745$ reproduce karta hai, toh hum ne sahi invert kiya — yeh ek self-consistency proof hai.

Verify: $⌊ 745/128 ⌋ = 5$ , $745 mod 128 = 105$ , aur $5 \times 128 + 105 = 745$ . ✓

Recall

Recall Ceiling division kyun, plain division kyun nahi?

Plain integer division round down karta hai, ek partial last block drop kar deta hai aur tail elements unprocessed chhod deta hai ::: ceiling $⌈ N / B ⌉ = (N + B - 1) / B$ round up karta hai taaki har element cover ho.

Recall "Guard" exactly kise protect karta hai?

if (i < N) / if (row < M && col < N) check last (padded) block ki extra threads ko aisi memory read ya write karne se rokta hai jo exist nahi karti ::: yeh isliye aati hain kyunki grid hamesha whole blocks tak round up hoti hai.

Recall Ragged launch mein idle threads kahan rehti hain?

Hamesha last block mein (1D) ya padded edge/corner blocks mein (2D); unke global indices data size se zyada hote hain, toh guard fire hota hai ::: yeh fully-interior blocks mein nahi aati.

Recall One-dim vs. two-dim ragged: kitni guard sub-conditions fire hoti hain?

Agar sirf ek axis ragged hai, toh sirf us axis ki sub-condition (row < M ya col < N) fire hoti hai ::: agar dono axes ragged hain, toh corner block dono sub-conditions ek saath trip karta hai.

Recall Occupancy threads ki jagah memory se kab limit hoti hai?

Jab $⌊ sharedMem_{S M} / sharedMem_{b l oc k} ⌋$ (ek floor) $⌊ threads_{S M} / threads_{b l oc k} ⌋$ se chhota hota hai ::: hardware sabhi floored resource limits ka minimum leta hai.

Recall globalID ko (block, thread) mein kaise invert karte hain?

$blockIdx = ⌊ globalID / B ⌋$ aur $threadIdx = globalID mod B$ ::: kyunki globalID $= blockIdx \cdot B + threadIdx$ jahan $0 \leq threadIdx < B$ .

Related deep threads: Warp Execution, Shared Memory, Thread Synchronization, Streaming Multiprocessors, Memory Coalescing, Parallel Algorithm Design, CUDA Programming Model.