Thread blocks and grids: Worked examples — Hardware deep dive

This page is the drill hall for the parent topic. We will hunt down every kind of situation the block/grid machinery can throw at you — clean divisions, ragged remainders, empty inputs, 2D tiling, occupancy walls — and work each one from zero. If any word below feels new, the parent note built it first; here we use it until it is muscle memory.

Before anything else, one reminder of the single formula everything rests on:

Look at figure s01. Each row is one block; the row labels on the left ("block 0"…"block 3") name them. The red seat (labelled globalID = 15) is one specific thread. To find its global number we follow the red vertical arrow — jumping over the two full rows above it (that is $blockIdx \times blockDim = 2 \times 6 = 12$ ) — and then follow the red horizontal arrow stepping across to its column ( $threadIdx = 3$ ). That single act of jumping-then-stepping is the entire chapter.

The scenario matrix

Every problem you will ever meet lives in one of these cells. The worked examples afterwards each carry a tag like (Cell A2) so you can see the whole grid gets covered.

#	Case class	What makes it tricky	Covered by
A1	Exact fit — $N$ divisible by block size	No leftover threads; guard never fires	Example 1
A2	Ragged fit — $N$ not divisible	Last block has idle threads; need ceiling + guard	Example 2
A3	Degenerate: $N = 0$	Empty input	Example 3
A4	Degenerate: $N = 1$	Single element, still one whole block	Example 3
A5	$N <$ one block	Grid rounds up to exactly 1 block	Example 3
B1	2D tiling — exact	Row/col from x,y; flatten row-major	Example 4
B2	2D tiling — ragged in ONE dim	Guard fires on only one axis (exact x, ragged y)	Example 5
B3	2D tiling — ragged in BOTH dims	Guard fires on row and col at the corner	Example 6
C1	Limiting behaviour — block size $\to$ max (1024)	Fewer, fatter blocks; occupancy risk	Example 7
C2	Occupancy wall — shared memory limits blocks/SM	Memory, not threads, is the bottleneck	Example 8
D1	Word problem — real workload sizing	Translate "process 2M pixels" into a launch	Example 9
D2	Exam twist — reverse the formula	Given globalID, recover block & thread	Example 10

Worked Examples

Worked example Example 1 — Exact fit

(Cell A1) Vector add of $N = 1024$ elements, block size $B = 256$ . How many blocks, and does the guard if (i < N) ever matter?

Forecast: guess the block count and whether any thread is wasted before reading on.

Compute the block count with ceiling division. $blocks = ⌈ \frac{N}{B} ⌉ = \frac{N + B - 1}{B} = \frac{1024 + 255}{256} = \frac{1279}{256} = 4$ Why this step? Integer division truncates, so we add $B - 1$ first to force a round-up — otherwise a partial block gets dropped and some data goes unprocessed.
Count total launched threads. $4 \times 256 = 1024$ . Why this step? To see if we launched more threads than we have data. Here $1024 = N$ exactly.
Check the guard. The last thread has $globalID = 1023 < 1024$ , so i < N is always true. Why this step? When it fits perfectly, no thread is idle — the guard never fires, but we keep it anyway for safety.

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

Worked example Example 2 — Ragged fit

(Cell A2) Same kernel, now $N = 1000$ , $B = 256$ . How many blocks? How many threads in the last block are idle? What is the global index of the first idle thread?

Forecast: will 3 blocks be enough? Guess the number of wasted threads.

Ceiling division. $blocks = \frac{1000 + 255}{256} = \frac{1255}{256} = 4$ Why this step? $3 \times 256 = 768 < 1000$ , so 3 blocks miss elements 768–999. We must round up to 4.
Total launched threads. $4 \times 256 = 1024$ . Why this step? Compare against $N = 1000$ to find the excess.
Idle threads. $1024 - 1000 = 24$ threads have no data. Why this step? These live in block 3 (the last one). Their global indices are $1000..1023$ — exactly the threads whose guard fires.
First idle thread's identity. globalID $= 1000$ . Reverse the formula: $blockIdx = ⌊ 1000/256 ⌋ = 3, threadIdx = 1000 - 3 \times 256 = 1000 - 768 = 232$ Why this step? Shows exactly where the guard starts firing: seat 232 in block 3 and everyone after.

Verify: block 3 covers globalIDs $768..1023$ ; the valid ones are $768..999$ (that's $232$ threads doing work) and $1000..1023$ idle (that's $24$ ). $232 + 24 = 256 = B$ . ✓

Worked example Example 3 — Degenerate inputs

$N \in {0, 1, 100}$ (Cells A3, A4, A5) With $B = 256$ , work out the launch for three tiny/empty cases.

Forecast: does $N = 0$ launch zero blocks or one? Guess before step 1.

Case $N = 0$ (empty input). $blocks = \frac{0 + 255}{256} = \frac{255}{256} = 0$ Why this step? Ceiling of $0/256$ is $0$ , so no blocks launch — the kernel does nothing, which is correct. No guard needed because no thread exists.
Case $N = 1$ (single element). $blocks = \frac{1 + 255}{256} = \frac{256}{256} = 1$ Why this step? One element still needs one whole block of 256 threads; 255 of them are idle. Thread with globalID $0$ works, threads $1..255$ have their guard fire.
Case $N = 100$ (smaller than one block). $blocks = \frac{100 + 255}{256} = \frac{355}{256} = 1$ Why this step? Anything from $1$ to $256$ elements rounds up to exactly 1 block — the grid never has a fraction of a block.

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

Worked example Example 4 — 2D tiling, exact

(Cell B1) A $64 \times 64$ image ( $N = 64$ columns, $M = 64$ rows). Block is $16 \times 16$ . Find the grid dimensions, and the global (row, col) plus flat memory index of the thread with $blockIdx = (x = 1, y = 2)$ and $threadIdx = (x = 5, y = 3)$ . (Recall x = col, y = row.)

Forecast: guess row and col before computing.

Grid size in each dimension. $gridDim.x = \frac{N}{16} = \frac{64}{16} = 4, gridDim.y = \frac{M}{16} = \frac{64}{16} = 4$ Why this step? We tile the 64-wide, 64-tall image with 16×16 tiles — 4 tiles across, 4 down. It divides evenly, so no ragged edge.
Column (x-direction). Here $blockIdx.x = 1$ and $threadIdx.x = 5$ : $col = blockIdx.x \times blockDim.x + threadIdx.x = 1 \times 16 + 5 = 21$ Why this step? Same jump-then-step idea, but horizontally. Look at the red tile in figure s02: block column 1 means skip 16 pixels, then step 5 more. The x component drives the column.
Row (y-direction). Here $blockIdx.y = 2$ and $threadIdx.y = 3$ : $row = blockIdx.y \times blockDim.y + threadIdx.y = 2 \times 16 + 3 = 35$ Why this step? Vertical version; block row 2 means skip 32 rows, then 3 down. The y component drives the row — never swap them.
Flatten to a 1D memory address (row-major). $idx = row \times width + col = 35 \times 64 + 21 = 2240 + 21 = 2261$ Why this step? Memory is a straight line, not a grid. Row-major means "walk full rows first, then across." Each full row is width $= N = 64$ elements wide.

Verify: col $= 21 < 64$ , row $= 35 < 64$ — both inside the image, guard passes. Flat index $2261 < 64 \times 64 = 4096$ . ✓

Worked example Example 5 — 2D tiling, ragged in ONE dimension

(Cell B2) Image is $64 \times 70$ ( $N = 64$ columns — exact for 16, $M = 70$ rows — not a multiple of 16). Block $16 \times 16$ . Find the grid, and identify in the bottom row of blocks which threads survive the guard.

Forecast: which axis is ragged, x or y? How many rows of idle threads appear?

Grid dimensions with ceiling on both, even though only one is ragged. $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$ Why this step? $64/16 = 4$ divides cleanly, so x is exact; $70/16 = 4.375$ rounds up to $5$ , so y is ragged. Ceiling is safe to apply to both — it is a no-op on the exact axis.
The padded region. The grid covers $64 \times 80$ thread-slots; the image is $64 \times 70$ . Why this step? Only the bottom strip of rows 70–79 is padding. The guard fires only on the row condition here — every column is valid because x is exact.
Bottom row of blocks, $blockIdx.y = 4$ . These blocks cover rows $64..79$ . Valid rows: $64..69$ (that's 6), invalid rows: $70..79$ (that's 10). Why this step? In each such block, threads with $threadIdx.y \leq 5$ (rows $64..69$ ) pass; $threadIdx.y \geq 6$ fail. The col guard never trips.
Working threads per bottom block. $16 cols \times 6 rows = 96$ work; $256 - 96 = 160$ have their guard fire — all because of the single ragged (y) axis. Why this step? This is the key single-dim lesson: the guard if (row < M && col < N) is protecting against one failing sub-condition, not two.

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, ragged in BOTH dimensions

(Cell B3) Image is $30 \times 20$ ( $N = 30$ columns, $M = 20$ rows). Block $16 \times 16$ . Find the grid, and identify which threads in the bottom-right corner block do real work.

Forecast: will the grid be $2 \times 2$ blocks? How many threads in that corner block are wasted?

Grid dimensions with ceiling. $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$ Why this step? Neither $N = 30$ nor $M = 20$ divides by 16, so both directions round up — a $2 \times 2$ grid of blocks covering a $32 \times 32$ padded region.
Total covered vs. real. Grid covers $32 \times 32 = 1024$ thread-slots; the image has $N \times M = 30 \times 20 = 600$ real pixels. Why this step? The $1024 - 600 = 424$ leftover threads must all be killed by the guard if (row < M && col < N).
Bottom-right block $blockIdx = (x = 1, y = 1)$ . Its columns span $col = 16..31$ , rows $row = 16..31$ . Why this step? We check which of these are still inside $30 \times 20$ . Valid cols: $16..29$ (that's 14), valid rows: $16..19$ (that's 4).
Working threads in that block. $14 \times 4 = 56$ do work; $256 - 56 = 200$ are guarded out. Why this step? This is the corner where both guard sub-conditions bite simultaneously — the worst-case ragged cell, contrasting Example 5 where only one bit.

Verify: the guard requires col < 30 AND row < 20. In block (x=1, y=1): cols $16..29$ = 14 values, rows $16..19$ = 4 values, product $56$ . Idle $= 256 - 56 = 200$ . ✓

Worked example Example 7 — Limiting behaviour: block size at the maximum

(Cell C1) $N = 3000$ elements. Compare launches at $B = 256$ vs. the hardware max $B = 1024$ . How many blocks each, and how many idle threads each?

Forecast: does the fat 1024 block waste more or fewer threads?

Small blocks $B = 256$ . $blocks = \frac{3000 + 255}{256} = \frac{3255}{256} = 12, launched = 12 \times 256 = 3072, idle = 72$ Why this step? Baseline to compare against.
Max blocks $B = 1024$ . $blocks = \frac{3000 + 1023}{1024} = \frac{4023}{1024} = 3, launched = 3 \times 1024 = 3072, idle = 72$ Why this step? Interesting — same total (3072) and same waste (72) here, but far fewer, fatter blocks.
Interpret the limit. With only 3 giant blocks, if 2 stall on memory, only 1 remains to hide latency — poor occupancy. With 12 small blocks there is much more to overlap. Why this step? This is exactly the "more threads per block ≠ faster" mistake from the parent note, now made concrete.

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

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

(Cell C2) An SM holds 2048 threads and 96 KB shared memory. Your block uses 256 threads and 18 KB shared memory. How many blocks fit per SM, and what is the occupancy?

Forecast: is the thread budget or the memory budget the bottleneck?

Blocks allowed by the thread budget. $⌊ \frac{2048}{256} ⌋ = 8 blocks$ Why this step? Capacity limits use the floor (round down): you can only fit whole blocks, so a fractional block does not count. The thread budget alone permits at most 8.
Blocks allowed by the shared-memory budget. $⌊ \frac{96}{18} ⌋ = 5 blocks$ Why this step? Shared memory is a separate budget, also floored. $96/18 = 5.33$ , so only 5 whole blocks fit — the fractional 0.33 is unusable.
Effective blocks = min(8, 5) = 5. Active threads $= 5 \times 256 = 1280$ . Why this step? Hardware takes the minimum of all resource ceilings — memory, not threads, throttles us here.
Occupancy. $occupancy = \frac{1280}{2048} \approx 0.625 = 62.5%$ Why this step? Occupancy = active threads ÷ max threads. We're leaving 37.5% of the SM idle because of shared-memory pressure.

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

Worked example Example 9 — Word problem: sizing a real launch

(Cell D1) You must apply a filter to a photo of 1920 × 1080 pixels ( $N = 1920$ columns, $M = 1080$ rows). Using $16 \times 16$ blocks, how many blocks total launch, and how many total threads? What fraction are idle?

Forecast: guess whether both dimensions divide evenly.

Blocks per dimension (ceiling both). $gridDim.x = ⌈ \frac{N}{16} ⌉ = 120, gridDim.y = ⌈ \frac{M}{16} ⌉ = \frac{1080 + 15}{16} = \frac{1095}{16} = 68$ Why this step? $1920/16 = 120$ exactly (clean x), but $1080/16 = 67.5$ → rounds to $68$ (ragged y).
Total blocks. $120 \times 68 = 8160$ blocks. Why this step? The grid is the product of its two dimensions.
Total threads launched. $8160 \times 256 = 2088960$ . Why this step? Every block has $16 \times 16 = 256$ threads.
Real pixels & idle threads. Real $= N \times M = 1920 \times 1080 = 2073600$ . Idle $= 2088960 - 2073600 = 15360$ . Why this step? The idle come entirely from the padded strip in $y$ (rows 1080–1087): $8 \times 1920 = 15360$ . Perfectly matches.

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

Worked example Example 10 — Exam twist: reverse the formula

(Cell D2) A thread reports $globalID = 745$ in a 1D launch with $blockDim = 128$ . Which block is it in, and which seat (threadIdx)? Then confirm forward.

Forecast: guess the block number before dividing.

Recover the block index by integer division. $blockIdx = ⌊ \frac{745}{128} ⌋ = 5 (since 5 \times 128 = 640 \leq 745 < 768 = 6 \times 128)$ Why this step? The formula is $globalID = blockIdx \times B + threadIdx$ . Dividing by $B$ and dropping the remainder strips off threadIdx, leaving the block.
Recover the thread index by remainder (modulo). $threadIdx = 745 mod 128 = 745 - 5 \times 128 = 745 - 640 = 105$ Why this step? The leftover after removing whole blocks (each $B = 128$ wide) is exactly the seat inside the current block. So this thread sits at seat 105 of block 5.
Forward-check. Plug both recovered values back into the original formula: $blockIdx \times B + threadIdx = 5 \times 128 + 105 = 640 + 105 = 745 ✓$ Why this step? If the forward formula reproduces the original globalID $745$ , we inverted correctly — a self-consistency proof.

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

Recall

Recall Why ceiling division and never plain division?

Plain integer division rounds down, dropping a partial last block and leaving tail elements unprocessed ::: ceiling $⌈ N / B ⌉ = (N + B - 1) / B$ rounds up so every element is covered.

Recall What exactly does the "guard" protect against?

The if (i < N) / if (row < M && col < N) check stops the extra threads in the last (padded) block from reading or writing memory that does not exist ::: they arise because the grid always rounds up to whole blocks.

Recall In a ragged launch, where do the idle threads live?

Always in the last block (1D) or the padded edge/corner blocks (2D); their global indices exceed the data size, so the guard fires ::: none appear in fully-interior blocks.

Recall One-dim vs. two-dim ragged: how many guard sub-conditions fire?

If only one axis is ragged, only that axis's sub-condition (row < M or col < N) fires ::: if both axes are ragged, the corner block trips both sub-conditions at once.

Recall When is occupancy limited by memory instead of threads?

When $⌊ sharedMem_{S M} / sharedMem_{b l oc k} ⌋$ (a floor) is smaller than $⌊ threads_{S M} / threads_{b l oc k} ⌋$ ::: the hardware takes the minimum of all floored resource limits.

Recall How do you invert globalID back to (block, thread)?

$blockIdx = ⌊ globalID / B ⌋$ and $threadIdx = globalID mod B$ ::: because globalID $= blockIdx \cdot B + threadIdx$ with $0 \leq threadIdx < B$ .

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