Exercises — Thread blocks and grids

WHAT: Plug into the one formula. WHY: the block index counts complete blocks before us (each holds 128 seats), then we add our seat. $\text{globalID}=3\times 128+10=384+10=394$ Answer: 394. Picture: we jumped over blocks 0, 1, 2 (that's $3\times 128=384$ elements) and landed 10 seats into block 3.

• threadIdx.x ::: position of the thread within its block (0 up to blockDim.x - 1).
• blockDim.x ::: number of threads per block.
• blockIdx.x ::: position of the block within the grid (0 up to gridDim.x - 1).
• gridDim.x ::: number of blocks in the grid.

Memory hook: thread words end in "thread/block position"; dim words are sizes.

False. __syncthreads() is a barrier inside one block only — see Thread Synchronization. Blocks are autonomous; there is no cross-block barrier inside a single kernel launch. Threads across blocks never wait for one another this way.

WHAT: ceiling division, because we must cover every element even if the last block is partly empty. $\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$, and 3 blocks give only $3\times 256=768$ threads — 232 elements left uncovered. So round up to 4.

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

Those 24 threads have globalID in $[1000,1023]$ and the if (i < N) guard makes them do nothing.

$\text{globalID}=2\times 256+200=512+200=712$ Since $712<1000$, the guard if (i < N) passes — it does process element 712.

WHAT: apply the same formula per axis, then flatten. WHY flatten: memory is a 1D line; row-major means "walk along a full row, then drop to the next" — see 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 is a minimum: a block can only become resident if every resource it needs is available at once — threads, registers, and shared memory. If the thread budget allows 8 blocks but the memory budget allows only 5, then only 5 can actually sit on the SM: the tightest resource is the real ceiling. So we compute each limit separately and take the smallest — see GPU Occupancy.

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

WHY floor here: a block that cannot fit its whole memory footprint is simply not admitted — there is no such thing as a "0.7 of a block" resident — so any fractional part is discarded downward.

Both limits say 8, so 8 blocks fit $\Rightarrow 8\times 256=2048$ resident threads $=2048/2048=$ 100% occupancy.

WHY the min again: occupancy is capped by whichever resource runs out first (see 3.1). Compute both, keep the smaller.

• Thread limit: $\lfloor 2048/256\rfloor =8$ blocks.
• Shared-memory limit: $\lfloor 96/20\rfloor =4$ blocks. WHY floor: $96/20=4.8$, but a 5th block would need $5\times 20=100$ KB $>96$ KB — it does not fit, so we drop the fractional 0.8 and keep 4.

The binding constraint is the smaller one: $\min (8,4)=4$ blocks. $\text{threads}=4\times 256=1024,\text{occupancy}=\frac{1024}{2048}=50\%$ The lesson: memory hunger — not thread count — cut occupancy in half.

WHAT: a warp is a fixed bundle of 32 threads running one instruction together (SIMT) — see Warp Execution. A lane is a thread's fixed slot inside that bundle, numbered 0–31; when the warp executes one instruction, all 32 lanes fire it together (unless some are masked off).

• Warps per block: $256/32=8$ warps.
• When a warp splits into two paths, the hardware runs the if path with lanes 16–31 masked off, then the else path with lanes 0–15 masked off. Each pass wastes the masked half — those lanes still consume a cycle doing nothing.
• So each divergent warp spends roughly 2× the cycles; about 50% of the slots in the divergent region are idle mask-outs.

Fix in practice: arrange branches so a whole warp takes one path (branch on globalID / 32, not on the lane number), so no lane inside a warp is ever masked against another.

Choice: B = 256 — a multiple of 32 (whole warps, no wasted lanes), and in the 128–512 sweet spot that keeps enough blocks per SM to hide memory latency without starving occupancy. $\text{blocks}=\lceil \frac{1000000}{256}\rceil =\frac{1000000+255}{256}=\frac{1000255}{256}=3907\text{ (integer)}$ Check: $3906\times 256=999936<1000000$, so 3906 is too few; 3907 is correct.

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

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

WHAT: cover columns along x, rows along y (convention: x = horizontal = columns) — see 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\%$.

Small waste is normal — it is the price of tiling a non-multiple grid.

WHAT: the index rule applies independently on each of the three axes — nothing new, just 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 is $8\times 8\times 8$ blocks; the thread maps to voxel $(x,y,z)=(17,29,39)$.

Flatten to a 1D memory index — and WHY this exact nesting: memory is one long line, so we must decide which coordinate walks fastest. Convention makes x the fastest-changing (neighbouring x are adjacent in memory, which helps Memory Coalescing), then y, then z slowest. So we build the address in layers: pick the z-slice (each slice holds $Y\times X$ voxels), then the y-row inside it (each row holds $X$ voxels), then step x along that row: $\text{idx}=(z\times Y+y)\times X+x=(39\times 64+29)\times 64+17=2525\times 64+17=161617$ WHY it works: $(z\times Y+y)$ is the row number counted from the very start of the volume, and multiplying by $X$ skips over all the full rows before us — then $+x$ lands on our seat in the current row. It is the same "skip the complete groups before you, then add your offset" logic as the 1D global-index formula, nested three deep.

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

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

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

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

Winner: the 256-thread block. Both 256 and 1024 reach 100% occupancy, but 256 is the safer choice — with 8 blocks resident, one stalled block still leaves 7 others for the scheduler to run and hide the latency; with 1024 you have only 2 blocks, so a single stall halves your latency-hiding capacity. The 128-thread block is eliminated because resources (not threads) cap it at 12 blocks $\Rightarrow$ only 75% occupancy. Pick 256.

The residency limit that makes this fail: each SM can hold at most 8 blocks at a time, and there are only 4 SMs, so at most $4\times 8=32$ of the 40 blocks can be resident simultaneously. The other 8 blocks must wait for a resident block to finish and free its slot before they can even start.

Why it deadlocks: blocks are scheduled onto SMs in arbitrary order. Suppose the 32 resident slots are all filled with blocks 1–32, every one of them spinning on flag. Spinning blocks never finish, so their slots never free up, so the not-yet-resident blocks — which might include block 0 — never get scheduled. If block 0 is among the waiting 8, flag is never set, so the resident blocks spin forever: a circular wait. There is no guarantee block 0 runs first; block indices imply no execution order.

Correct alternatives (from Thread Synchronization):

1. Split into two kernel launches. The kernel boundary is an implicit global barrier — everything from launch 1 completes before launch 2 starts.
2. Cooperative groups (CUDA 9+) grid-level sync, launched with cudaLaunchCooperativeKernel, which guarantees all blocks are co-resident so a grid barrier is legal.
3. (Last resort) atomic operations on global memory for coordination — correct but slow, and still must avoid assuming any block runs first.

"No waste" means N is an exact multiple of the block size (remainder 0). Check each:

• $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. ✅

All four divide $4096=2^{12}$ evenly (they are all powers of two dividing it), and all are warp multiples (each is a multiple of 32). So every listed size gives zero guarded-out threads. Among them, 256 is the usual pick for the occupancy/latency-hiding balance argued in Exercise 5.1.

• The universal 1D index formula ::: $\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}$
• The same rule on each of three axes ::: apply independently to x, y, and z (e.g. $\text{blockIdx.z}\times \text{blockDim.z}+\text{threadIdx.z}$)
• Block count to cover N elements (round up) ::: $\lceil N/B\rceil =(N+B-1)/B$
• Blocks that fit a fixed resource budget (round down) ::: $\lfloor \text{budget}/\text{per-block cost}\rfloor$
• Occupancy ::: resident threads ÷ SM max threads, capped by the minimum over all resource limits
• Why blocks must not message each other in one kernel ::: arbitrary scheduling order + limited residency → deadlock risk