Worked examples — CUDA programming model basics

This page is the hands-on companion to CUDA programming model basics. The parent note built the machinery: blocks map to SMs, threads carry a global index, and memory has speed tiers. Here we use that machinery on every awkward case a real problem can throw at you.

Before anything, let us re-anchor the one formula we will lean on the whole page. Picture a long line of numbered lockers — that is your array. We chop the line into equal blocks (fixed-size chunks). Inside each block, threads are numbered from $0$ upward.

Read the figure first. The horizontal strip is your array; each little box is one thread's element, and the number inside each box is that thread's globalIdx. The cyan boxes are ordinary threads; the single amber box (rightmost, index $11$) is the thread we track in Example 1. Above the boxes, the white brackets group boxes into blocks — three brackets labelled block 0, block 1, block 2. Below each box the small white t0 … t3 tags are the seat number inside that block (threadIdx.x). Notice the pattern: seats restart at $0$ at every new bracket, while the big number keeps climbing — that climb is the global index.

Every symbol above now has a picture. Two more pieces of notation appear repeatedly below, so we define them now, before first use:

Now the matrix.

The scenario matrix

Every CUDA-indexing problem you will meet is one of these cells. The examples below are labelled with the cell they cover, and together they cover all of them.

Cell	What makes it tricky	Covered by
A. Clean divide	$N$ is an exact multiple of block size — no leftovers	Ex 1
B. Ragged tail	$N$ not divisible — extra idle threads appear	Ex 2
C. Degenerate: tiny $N$	$N$ smaller than one block, or $N=0$, or $N=1$	Ex 3
D. Reverse map	Given a globalIdx, find which block & seat produced it	Ex 4
E. 2-D grid	Image / matrix — two index axes, row-major flattening	Ex 5
F. Memory-cost word problem	Should data live in registers, shared, or global?	Ex 6
G. Launch-config sizing	Choose block size; count idle threads and wasted warp lanes	Ex 7
H. Exam twist: stride loop	More elements than threads — one thread handles many	Ex 8
I. 3-D grid	Volumes / batches — three index axes, depth·height·width flatten	Ex 9

---

Cell A — the clean divide

Forecast: guess the index before reading on. (Hint: last block, last seat — should be the very last element.)

1. Read the two pieces. blockIdx.x = 2, blockDim.x = 4, threadIdx.x = 3. Why this step? The formula needs exactly these three numbers; gather them first so nothing is guessed later.
2. Threads before my block: $2\times 4=8$. Why this step? Two whole chunks of $4$ seats sit ahead of me; I must skip all $8$. (In the figure, that is the two full brackets block 0 and block 1.)
3. Add my seat: $8+3=11$. Why this step? Seat $3$ inside my chunk lands me $3$ past my chunk's front door — the amber box in the figure.
4. Boundary check if (idx < N): is $11<12$? Yes. Why this step? Even in a clean divide we still test — good habit, and it protects the ragged cases later.

Answer: globalIdx $=11$, valid.

Verify: Total threads $=3\times 4=12$, exactly $N$. The indices produced are $0,1,\dots ,11$ — every element hit once, none twice, none skipped. $11$ is the largest, so the last thread hits the last element. Consistent.

---

Cell B — the ragged tail

Forecast: will $1000/256$ come out even? No — so expect leftover idle threads. Guess how many.

1. Ceiling division. We need at least enough threads for all $1000$ elements: $\text{numBlocks}=\lceil \frac{1000}{256}\rceil =\frac{1000+256-1}{256}=\frac{1255}{256}=4(\text{integer division}).$ Why this step? $4$ blocks give $1024$ threads; $3$ blocks give only $768<1000$, which would leave $232$ elements untouched. We must round up — that is exactly what the ceiling bracket $\lceil \cdot \rceil$ (defined above) does.
2. Why the $+255$ trick? Adding threadsPerBlock − 1 before integer-dividing forces a round-up without floating-point math. Why this step? Integer division truncates down; the offset pushes any non-zero remainder over the next boundary. It answers the question "did anything spill past the last full block?"
3. Total threads launched: $4\times 256=1024$. Why this step? Blocks come in fixed sizes — you cannot launch a partial block, so you always over-provision.
4. Idle threads: $1024-1000=24$. These have idx $=1000\dots 1023$. Why this step? This is exactly why if (idx < N) exists — without it, threads $1000$–$1023$ read past the array end. See the mistakes callout in the parent note.

Answer: 4 blocks, 1024 threads launched, 24 idle.

Verify: $4\times 256=1024\ge 1000$ ✓ and $3\times 256=768<1000$ ✗, so $4$ is the smallest valid count. Idle $=24$, and $1024-24=1000=N$. ✓

---

Cell C — degenerate & edge inputs

Forecast: $N=0$ is the sneaky one — how many blocks should you launch for nothing?

1. (a) $N=1$: $\lceil 1/256\rceil =\lceil 0.0039\dots \rceil =1$ block. Why? Even one element needs one block to hold it; you cannot launch a fraction of a block. Threads launched $=256$, idle $=255$. Only thread $0$ passes idx < 1.
2. (b) $N=0$: $\lceil 0/256\rceil =0$ blocks. Via the trick: $(0+255)/256=255/256=0$ (integer division). Why? Nothing to do → the ceiling formula returns $0$ blocks.
3. (c) $N=256$: $\lceil 256/256\rceil =1$ block. Trick: $(256+255)/256=511/256=1$. Why? Exactly one full block; the $+255$ offset does not falsely bump us to $2$ because $511<512$. This confirms the ceiling trick never over-rounds on exact multiples.

Answer: (a) 1 block, 255 idle; (b) 0 blocks (guard the launch), 0 idle; (c) 1 block, 0 idle.

Verify: (a) $1\cdot 256-1=255$ ✓. (b) $0\cdot 256-0=0$ ✓. (c) $1\cdot 256-256=0$ ✓, and the trick gives $1$, not $2$, so no phantom extra block.

---

Cell D — the reverse map

Forecast: which block? Roughly $813/256\approx 3.2$, so block $3$ maybe. Let's be exact.

1. Recover the block by integer division (the floor bracket $\lfloor \cdot \rfloor$ defined above — round down): $\text{blockIdx.x}=\lfloor \frac{813}{256}\rfloor =3.$ Why this step? The formula was $\text{globalIdx}=\text{blockIdx}\times 256+\text{seat}$ with $0\le \text{seat}<256$. Dividing by $256$ and dropping the remainder strips off the seat and leaves the block — it inverts the multiplication.
2. Recover the seat with the remainder (modulo): $\text{threadIdx.x}=813\mathrm{mod}256=813-3\times 256=813-768=45.$ Why this step? The remainder is exactly "how far past the block's front door" — the seat we skipped over in step 1.

Answer: block 3, seat 45.

Verify: Rebuild the index forward: $3\times 256+45=768+45=813$ ✓. And $0\le 45<256$, a legal seat. Both directions agree — division and modulo are the clean inverse of the index formula.

---

Cell E — the 2-D grid

Images and matrices need two axes. Picture a grid laid over a photo: threadIdx.x/blockIdx.x count columns (across), threadIdx.y/blockIdx.y count rows (down).

Read the figure first. The big square is a $64\times 64$ corner of the image. The faint white lines chop it into $16\times 16$ blocks (four across, four down). The cyan-shaded square is the one block at blockIdx = (2, 3) — its columns run $32-47$ (that is $2\times 16$ to $2\times 16+15$) and its rows run $48-63$ ($3\times 16$ onward). Inside it, the amber square marker is our target pixel: seat (5, 7) inside that block lands at column $37$, row $55$. The amber label spells out both little index sums. Notice rows increase downward (image convention) — the y-axis is flipped for exactly that reason.

Forecast: row uses the .y numbers, column uses the .x numbers. Guess before computing.

1. Column (x-axis): $\text{col}=\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}=2\times 16+5=37.$ Why this step? Same 1-D logic as Example 1, applied along the horizontal axis only.
2. Row (y-axis): $\text{row}=\text{blockIdx.y}\times \text{blockDim.y}+\text{threadIdx.y}=3\times 16+7=55.$ Why this step? The vertical axis is a second, independent copy of the same formula.
3. Flatten to 1-D (row-major): $\text{globalIdx}=\text{row}\times \text{width}+\text{col}=55\times 512+37=28160+37=28197.$ Why this step? Memory is a straight line; a full row of $512$ pixels sits before each new row. Multiplying row × width skips all complete rows above, then + col walks across the current one. This is exactly the 1-D idea nested inside itself.

Answer: pixel (row 55, col 37), flat index 28197.

Verify: Bounds: $37<512$ ✓, $55<512$ ✓. Flatten range check: max legal flat index is $512\times 512-1=262143$, and $28197<262143$ ✓. Re-derive col from flat: $28197\mathrm{mod}512=37$ ✓; row $=\lfloor 28197/512\rfloor =55$ ✓.

---

Cell F — the memory-cost word problem

Forecast: re-reading from far-away global memory 8 times sounds wasteful. Guess the cycle counts.

1. All-global cost: $8\text{ accesses}\times 300\text{ cycles}=2400$ cycles. Why this step? If nothing is cached, every one of the 8 reads pays the full off-chip trip. See Memory Hierarchy for why global is ~$300$ cycles (off-chip, ~cm of wire + controller).
2. Shared-memory cost: one global load $=300$, then $8$ shared reads $=8\times 5=40$. Total $=300+40=340$ cycles. Why this step? Pay the expensive trip once to stage the data on-chip, then re-read cheaply. Shared memory lives inside the SM.
3. Speedup: $2400/340\approx 7.06\times$. Why this step? This ratio is the reason shared-memory tiling exists — reuse amortises the one expensive load.

Answer: 2400 vs 340 cycles → ≈7.06× faster with shared memory.

Verify: Break-even reuse count $r$ solves $300r=300+5r\Rightarrow 295r=300\Rightarrow r\approx 1.017$. So any element read twice or more already favours shared memory — and we read it $8$ times, comfortably past break-even. ✓

---

Cell G — launch-config sizing

Forecast: $1000$ divides $1000000$ evenly — tempting! But is $1000$ even allowed (max is $1024$), and is it a multiple of $32$? Guess where the hidden waste is.

1. Option A — 256: $\text{numBlocks}=\lceil 10^6/256\rceil =(10^6+255)/256=3907$. Threads $=3907\times 256=1000192$. Idle (grid-level) $=192$. Why this step? $256$ threads is $256/32=8$ whole warps — no partial warp. Wasted warp lanes per block $=0$, so total wasted lanes $=0$.
2. Option B — 1000: $1000\le 1024$, so it is legal. $\text{numBlocks}=10^6/1000=1000$ exactly. Threads $=1000000$. Idle (grid-level) $=0$. Why this step? An exact divide means zero leftover threads at the grid level — no ragged tail. So far it looks perfect.
3. Quantify the warp-level waste for Option B. Warps per block $=\lceil 1000/32\rceil =32$ warps. Those $32$ warps physically occupy $32\times 32=1024$ lanes, but only $1000$ threads are real. Wasted lanes per block $=1024-1000=24$. Why this step? The hardware always schedules a full warp of $32$ lanes; the last warp of each block runs with $1000-31\times 32=1000-992=8$ live threads and $24$ dead lanes. Those dead lanes still consume a scheduling slot — pure waste.
4. Total wasted lanes across the grid (Option B): $24\text{ lanes/block}\times 1000\text{ blocks}=24000$ wasted lanes. Why this step? Multiply per-block waste by the block count to see the full picture. Option A wastes $192$ idle threads total; Option B wastes $24000$ warp lanes total — over $100\times$ more waste, hidden behind its "zero idle threads" headline.

Answer: A: 3907 blocks, 1 000 192 threads, 192 idle, 0 wasted lanes. B: 1000 blocks, 1 000 000 threads, 0 grid-idle, but 24 wasted lanes/block = 24 000 total — the worse choice.

Verify: A: $3907\times 256=1000192\ge 10^6$ and $3906\times 256=999936<10^6$ ✓ (smallest count); $256\mathrm{mod}32=0$ → $0$ waste ✓. B: $1000\times 1000=10^6$ exactly ✓, $1000\le 1024$ ✓ (legal); warps/block $=\lceil 1000/32\rceil =32$, waste/block $=32\times 32-1000=24$ ✓, total $=24\times 1000=24000$ ✓.

---

Cell H — the exam twist (grid-stride loop)

Sometimes you launch fewer threads than elements on purpose (e.g. to fit a specific GPU). Each thread then loops, jumping forward by the total number of threads each time. That jump is the stride.

Read the figure first. The strip is a $20$-element array, boxes numbered $0-19$. Each box carries a small white t0…t7 tag showing which of the $8$ threads owns it — you can see the tags repeat $0,1,\dots ,7,0,1,\dots$ because every thread jumps ahead by the stride of $8$. The boxes owned by thread 3 are drawn in bright amber; the two curved amber arrows trace its journey $3\to 11\to 19$. After $19$, the next jump would land on $27$, which is off the strip — so thread 3 stops. That is the loop condition i < N doing its job.

Forecast: thread 3 starts at element 3 and jumps by 8. Guess its list before computing.

1. Compute the stride. $\text{stride}=\text{numBlocks}\times \text{blockDim}=2\times 4=8$. Why this step? The stride must equal the total thread count so the 8 threads tile the array without overlap — like 8 people counting off "1-8, 9-16, ..." repeatedly.
2. March thread 3. Start $i=3$; then $3,3+8=11,11+8=19,19+8=27$. Stop, because $27\nless 20$. Why this step? The loop condition i < N cleanly halts each thread when it runs off the end — this replaces the single if (idx < N) check with a loop that self-limits.
3. Collect the survivors. The values that passed i < 20 are $3,11,19$. So thread 3 handles $\{3,11,19\}$ — three elements from one thread. Why this step? This is the whole point of the grid-stride loop: $8$ threads cover $20$ elements by each doing a few passes.

Answer: thread 3 → elements 3, 11, 19.

Verify: Coverage count: elements $0$–$7$ (first pass, all 8 threads), $8$–$15$ (second pass), $16$–$19$ (third pass, only threads $0$–$3$ reach). Total $=8+8+4=20=N$ ✓, every index hit exactly once (no two threads share a start, and the stride equals the thread count so passes never collide). ✓

---

Cell I — the 3-D grid

Volumes (medical scans, physics grids) and batched tensors need a third axis: blockIdx.z / threadIdx.z count depth (which slice). It is the same one-axis formula, now stamped out three times, then flattened with a nested version of the row-major trick.

Forecast: three independent sums, then one three-term flatten. Guess $z$ first — it uses the .z numbers.

1. x-coordinate: $x=1\times 8+4=12$. Why this step? Horizontal axis, one copy of the base formula.
2. y-coordinate: $y=2\times 8+5=21$. Why this step? Vertical axis, independent second copy.
3. z-coordinate (depth): $z=3\times 8+6=30$. Why this step? Depth axis, independent third copy — this is the new piece 3-D adds.
4. Flatten depth-major with $\text{height}=\text{width}=64$: $\text{globalIdx}=30\times (64\times 64)+21\times 64+12=30\times 4096+1344+12=122880+1356=124236.$ Why this step? Skip $30$ whole $64\times 64$ slices, then $21$ whole rows, then $12$ across.

Answer: $(x,y,z)=(12,21,30)$, flat index 124236.

Verify: Bounds: $12,21,30$ all $<64$ ✓. Max legal flat index $=64^3-1=262143$, and $124236<262143$ ✓. Unflatten: $124236\mathrm{mod}64=12=x$ ✓; $(124236//64)\mathrm{mod}64=1941\mathrm{mod}64=21=y$ ✓; $124236//4096=30=z$ ✓.

---

Recall Self-test

Clean vs ragged: what does the +threadsPerBlock-1 trick achieve? ::: It forces integer division to round up (ceiling $\lceil \cdot \rceil$) so we never launch too few blocks, without floating-point math. Reverse map: given globalIdx and blockDim, how do you recover blockIdx and threadIdx? ::: blockIdx $=\lfloor$globalIdx / blockDim$\rfloor$ (floor / integer divide), threadIdx $=$ globalIdx mod blockDim. 2-D flatten: why row × width + col? ::: Memory is 1-D; each full row of width pixels sits before the next, so row × width skips complete rows and + col walks the current one. 3-D flatten: what is the extra term and why? ::: $z\times (\text{height}\times \text{width})$ — it skips $z$ whole 2-D slices before applying the 2-D formula. What is a warp and why 32? ::: A warp is a hardware bundle of 32 threads run in lockstep; 32 is a fixed NVIDIA hardware width (one instruction feeds 32 lanes). Shared vs global break-even: from how many reuses does shared memory win? ::: From 2 reuses ($300r=300+5r$ gives $r\approx 1.02$). Why prefer 256 over 1000 threads/block even though 1000 divides evenly? ::: $256$ is a multiple of the warp size $32$ (0 wasted lanes); $1000$ leaves $24$ dead lanes per block ($24000$ across the grid).

See also: GPU Architecture Overview · Parallel Algorithm Design · OpenCL vs CUDA · Deep Learning Frameworks