Worked examples — GPU architecture — SIMT, warps, CUDA model

This page drills the whole GPU architecture — SIMT, warps, CUDA model toolbox through worked cases. Before any code: we treat the GPU as a machine that runs threads in fixed bundles of 32 (a warp), organized into blocks, organized into a grid. Every formula we use here was derived in the parent note — we now put it under stress.

---

Two tools we use constantly — define them first

---

The scenario matrix

Every case class this topic can throw at you, and which example covers it:

#	Case class	Concrete trigger	Example
A	Clean divisible launch	$N$ is a multiple of block size $B$	Ex 1
B	Partial last block (rounding up)	$N$ not divisible by $B$; needs guard	Ex 2
C	Degenerate size	$N=0$ or $N<B$ (one tiny block)	Ex 3
D	Non-multiple-of-32 block	wasted lanes in the tail warp	Ex 4
E	Warp divergence (worst case)	intra-warp if/else, both paths run	Ex 5
F	Zero divergence (warp-uniform)	branch on threadIdx.x / 32	Ex 5
G	Coalesced memory (best case)	stride-1 access, efficiency $=1$	Ex 6
H	Scattered memory (worst case)	stride-$s$ access, efficiency $=1/s$	Ex 6
I	Latency hiding limit	how many warps to hide 400 cycles	Ex 7
J	Word problem (real world)	shade $10000$ pixels, choose launch	Ex 8
K	Exam twist (2D index)	blockDim vs gridDim trap in 2D	Ex 9

Prerequisite threads if any cell feels shaky: SIMD and Vector Instructions (AVX, SSE), Memory Coalescing and Bandwidth, CUDA Shared Memory and Tiling, Amdahl's and Gustafson's Laws.

---

Example 1 — Cell A: the clean divisible launch

Step 1 — Count the blocks. $\text{numBlocks}=\lceil \frac{N}{B}\rceil =\lceil \frac{1024}{256}\rceil =4$ Why this step? $N$ divides evenly, so the ceiling (round-up) does nothing here — this is the baseline against which the messy cases (B, C) look different.

Step 2 — Total threads launched. $4\times 256=1024=N$ Why this step? When $N$ is a multiple of $B$, launched threads equal data elements exactly — no wasted threads, so the bounds guard if (i >= N) return; never fires.

Step 3 — The specific global index. $i=\text{blockIdx.x}\cdot \text{blockDim.x}+\text{threadIdx.x}=3\cdot 256+5=773$ Why this step? "Skip 3 whole blocks ($3\times 256=768$), then step 5 in." This is the derived index formula, not a memorized one.

Verify: $773<1024$, so it is a valid element. Blocks $\times$ threads $=1024$ matches $N$, so every element is covered exactly once. Units: (blocks)(threads/block) + (threads) = threads. ✓

---

Example 2 — Cell B: the partial last block

Step 1 — Round up with the integer trick. $\text{numBlocks}=\lceil \frac{N}{B}\rceil =\frac{N+B-1}{B}=\frac{1000+255}{256}=\frac{1255}{256}=4(\text{integer division})$ Why this step? Plain integer division truncates $1000/256=3.9\dots$ down to $3$, which would drop elements $768\dots 999$. Adding $B-1$ pushes the remainder past the next integer so truncation rounds up to 4 — this is precisely the $\lceil N/B\rceil$ identity defined above.

Step 2 — Threads launched. $4\times 256=1024\text{ threads for }1000\text{ elements}$ Why this step? We deliberately launch more threads than data. The surplus is $1024-1000=24$.

Step 3 — The guard fires for exactly the surplus. The guard is if (i >= N) return;. Threads with $i\in \{1000,\dots ,1023\}$ have $i\ge 1000$, so 24 threads exit early; the other $1000$ do real work. Why this step? Guard on the bad condition ($i\ge N$), because the statement after is the early exit — guarding on i < N would kill the good threads (see the parent's steel-manned mistake).

Verify: working threads $=1024-24=1000=N$. ✓ Every valid element handled exactly once; no valid index dropped, no invalid index writing out of bounds.

---

Example 3 — Cell C: degenerate sizes ($N=0$ and $N<B$)

Step 1 — $N=0$. $\frac{0+255}{256}=\frac{255}{256}=0(\text{integer division})$ Why this step? $255/256=0.996\dots$ truncates to $0$, so the round-up trick correctly gives zero blocks for empty input — the kernel launches nothing, no thread runs, nothing crashes. This is the safe degenerate floor.

Step 2 — $N=10$, count blocks. $\frac{10+255}{256}=\frac{265}{256}=1(\text{integer division})$ Why this step? $265/256=1.035\dots$ truncates to $1$: one block still holds all 256 lanes; you cannot launch "10 threads" as a fractional block — the block size is fixed at 256.

Step 3 — How many exit early for $N=10$. Launched $=256$, valid $=10$, so $256-10=246$ threads hit i >= N and return. Why this step? This shows most of the block is wasted when $N\ll B$ — a hint that tiny problems belong on the CPU (relates to Amdahl's and Gustafson's Laws: not enough parallel work to amortize launch overhead).

Verify: (a) $0$ blocks $\times$ anything $=0$ threads for $0$ elements ✓. (b) working threads $=256-246=10=N$ ✓.

---

Example 4 — Cell D: a block that isn't a multiple of 32

Step 1 — Warps needed (round up to 32). $\text{warps}=\lceil \frac{B}{32}\rceil =\lceil \frac{100}{32}\rceil =\lceil 3.125\rceil =4$ Why this step? A warp is an indivisible bundle of 32 lanes, so a fractional warp count must round up — that is why we use ceiling. 100 threads need 4 warps because 3 warps hold only 96.

Step 2 — Lanes provisioned vs used. $4\times 32=128\text{ lanes provisioned},100\text{ used}$ Why this step? The 4th warp holds threads $96\dots 99$ (4 active) and $100\dots 127$ (28 idle, masked lanes) — look at the amber ghosted lanes in the figure.

Step 3 — Waste fraction. $\text{wasted}=\frac{128-100}{128}=\frac{28}{128}=0.21875=21.875\%$ Why this step? Roughly a fifth of the tail warp's issue slots are burned. Choosing $B=128$ (a multiple of 32) would waste zero lanes — the fix from the parent's "Why 32?" callout.

Verify: $4\times 32=128\ge 100$ and $3\times 32=96<100$, so 4 is the minimal warp count ✓. Waste $=28/128=0.21875$ ✓.

---

Example 5 — Cells E & F: divergence worst case vs zero divergence

Step 1 — Fix the baseline metric: total cycles to finish all 64 lanes. Both layouts do the same amount of work (all 64 lanes run their path). The fair comparison is wall-clock cycles until the last lane is done, assuming the two warps run on the same scheduler back-to-back (worst case for both — no second scheduler). This makes E and F directly comparable. Why this step? Comparing "one warp's cost" is ambiguous; comparing time to finish the whole job is not. We now measure both layouts against this single clock.

Step 2 — Case E (divergent): both paths serialize inside each warp. There is only one instruction fetch unit per warp, so a warp holding both A and B lanes runs A (B-lanes masked) then B (A-lanes masked): $T_{\text{warp}}^{(E)}=T_A+T_B=30+50=80\text{ cycles}$ Both warps here are identical (16A+16B each), so run back-to-back on one scheduler: $T_{\text{job}}^{(E)}=2\times 80=160\text{ cycles}$ Why this step? Divergence forces every warp to pay $T_A+T_B$; there are 2 warps, so $2(T_A+T_B)$.

Step 3 — Case F (uniform): each warp takes exactly one path, no masking waste. Warp 0 (all A) costs $T_A=30$; warp 1 (all B) costs $T_B=50$. Run back-to-back on the same scheduler: $T_{\text{job}}^{(F)}=T_A+T_B=30+50=80\text{ cycles}$ Why this step? When all 32 lanes of a warp agree, that warp pays only its own path — no idle-masked path. The job is just the sum of the two single-path warps.

Step 4 — The speedup, on one honest metric. $\frac{T_{\text{job}}^{(E)}}{T_{\text{job}}^{(F)}}=\frac{160}{80}=2\times$ Why this step? Same total useful work, same hardware, same clock — the only difference is divergence. Reordering lanes so warps are uniform halves the time here. Divergence is per-warp, not per-thread: correctness never changes, only speed.

Verify: $T_A+T_B=80$ (divergent per warp) ✓; $T_{job}^{(E)}=2\cdot 80=160$ ✓; $T_{job}^{(F)}=30+50=80$ ✓; ratio $160/80=2$ ✓.

---

Example 6 — Cells G & H: coalesced vs scattered memory

Step 1 — Bytes actually requested (same for both). $32\text{ threads}\times 4\text{ B}=128\text{ B of useful data}$ Why this step? The numerator of efficiency is fixed — the useful bytes. Only the denominator (bytes transferred) changes with the access pattern.

Step 2 — Case G, stride 1. 32 consecutive floats $=128$ B fall inside one 128-byte transaction. $\text{efficiency}=\frac{128}{128}=1$ Why this step? Consecutive addresses coalesce into a single line — the best case (top row of the figure, one solid cyan line).

Step 3 — Case H, stride 32. Addresses are $0,128,256,\dots$ bytes apart — each of the 32 reads lands in a different 128-byte line: $\text{bytes transferred}=32\times 128=4096\text{ B},\text{efficiency}=\frac{128}{4096}=\frac{1}{32}$ Why this step? Only 4 useful bytes are taken from each 128-byte line — the other 124 are hauled across the bus and discarded (scattered amber dots in the figure). This is the ~$32\times$ slowdown.

Verify: $128/128=1$ ✓; $128/(32\cdot 128)=1/32\approx 0.03125$ ✓. See also Memory Coalescing and Bandwidth.

---

Example 7 — Cell I: how many warps hide the latency?

Step 1 — The hiding condition. While one warp waits $L$ cycles, the scheduler must fill those $L$ cycles with other warps' work. Each other warp supplies $W$ useful cycles before it too stalls. To keep every cycle busy: $(n-1)\cdot W\ge L\Rightarrow n\ge \frac{L}{W}+1$ Why this step? This is why occupancy matters (parent's latency-hiding example): warps are "spare tires." We need enough of them that the ALU never idles.

Step 2 — Plug in. $n\ge \frac{400}{8}+1=50+1=51\text{ warps}$ Why this step? Fewer than 51 leaves gap cycles where the SM stalls with no ready warp; more than 51 gives diminishing returns (parent's "more threads isn't always faster" mistake).

Step 3 — Sanity against SM capacity. Many SMs cap at 64 resident warps; $51\le 64$, so this workload is hideable. If $L$ were $800$, we'd need $101>64$ — latency only partly hidden. Why this step? Occupancy is bounded by hardware limits (registers, shared memory, max warps).

Verify: $400/8+1=51$ ✓, and $51\le 64$ ✓ so hideable; with $L=800$: $800/8+1=101>64$ ✓ not fully hideable.

---

Example 8 — Cell J: real-world word problem

Step 1 — Blocks (round up). $\text{numBlocks}=\lceil \frac{N}{B}\rceil =\frac{10000+127}{128}=\frac{10127}{128}=79(\text{integer division})$ Why this step? $10000/128=78.1\dots$; ceiling rounds up to $79$. $78\times 128=9984<10000$, so 78 blocks would drop 16 pixels. 79 covers them.

Step 2 — Threads launched and surplus. $79\times 128=10112\text{ threads},\text{surplus}=10112-10000=112$ Why this step? 112 threads hit if (i >= 10000) return;. Every pixel gets exactly one working thread.

Step 3 — Warps. $\frac{128}{32}=4\text{ warps/block},79\times 4=316\text{ warps total}$ Why this step? Because $B=128$ is a multiple of 32, no lane is wasted inside any block (contrast Ex 4). Only the tail block has idle threads, not tail warps full of idle lanes.

Verify: $79\times 128=10112\ge 10000$ and $78\times 128=9984<10000$ ✓; surplus $=112$ ✓; working $=10112-112=10000$ ✓; total warps $=316$ ✓.

---

Example 9 — Cell K: the 2D exam twist (blockDim vs gridDim)

Step 1 — Nail the definitions (the trap). blockDim.x = 16 = threads per block in x; gridDim.x = number of blocks in x $=\lceil 256/16\rceil =16$. (Parent mistake: blockDim is NOT the block count.) Why this step? Every 2D index bug traces to swapping these. Units check: threads/block vs blocks — different units, so they can never be the same quantity.

Step 2 — Locate the block index and the in-block thread index, per axis. The block that owns coordinate $x$ is "how many whole blocks of width 16 fit below $x$", i.e. integer division; the in-block offset is the remainder: $\text{blockIdx.x}=\lfloor \frac{x}{16}\rfloor =\lfloor \frac{200}{16}\rfloor =12,\text{threadIdx.x}=200-12\cdot 16=8$ $\text{blockIdx.y}=\lfloor \frac{y}{16}\rfloor =\lfloor \frac{100}{16}\rfloor =6,\text{threadIdx.y}=100-6\cdot 16=4$ Here $\lfloor \cdot \rfloor$ (the floor, round-down partner of ceiling) is exactly the truncating integer division we defined. Why this step? This is the 1D index formula $x=\text{blockIdx.x}\cdot \text{blockDim.x}+\text{threadIdx.x}$ run backwards, once per axis — block index is the quotient, thread index is the remainder.

Step 3 — Rebuild the coordinate as a self-check. $x=\text{blockIdx.x}\cdot \text{blockDim.x}+\text{threadIdx.x}=12\cdot 16+8=200✓$ $y=\text{blockIdx.y}\cdot \text{blockDim.y}+\text{threadIdx.y}=6\cdot 16+4=100✓$ Why this step? Confirms our quotient/remainder split reconstructs the original pixel — no off-by-one.

Step 4 — Flatten $(x,y)$ to one linear memory index (row-major). Row-major means "skip whole rows above me, then step across my row": $i=y\cdot \text{width}+x=100\cdot 256+200=25600+200=25800$ Why this step? This is the address a coalescing check (Ex 6) would use. Because consecutive x gives consecutive i, a warp scanning across a row ($x,x+1,\dots$) is coalesced — a direct link back to Cell G.

Verify: $\lfloor 200/16\rfloor =12$, $200-12\cdot 16=8$; $\lfloor 100/16\rfloor =6$, $100-6\cdot 16=4$; rebuilt $x=200$, $y=100$; flattened $i=100\cdot 256+200=25800$ ✓.

---

Recall Quick self-test — cover the answers

What does $\lceil 3.1\rceil$ equal, and why can it never be $3$? ::: $4$ — ceiling always rounds up on any fractional part Guard for $N=1000$, $B=256$: how many threads exit early? ::: $1024-1000=24$ Warps for a block of 100 threads, and wasted lanes? ::: $4$ warps, $28$ idle lanes ($21.875\%$) End-to-end cost of divergent vs uniform layout ($T_A=30$, $T_B=50$, 64 lanes)? ::: $160$ vs $80$ cycles → $2\times$ Coalescing efficiency for stride-32 4-byte reads? ::: $1/32$ Warps needed to hide $L=400$ with $W=8$ cycles/warp? ::: $51$ Blocks to shade $10000$ pixels with $B=128$? ::: $79$ 2D pixel $(x=200,y=100)$, $16\times 16$ blocks: block index x? flat index? ::: $\text{blockIdx.x}=12$; flat $=25800$