CUDA cores and execution model: Worked examples — Hardware deep dive

This is the "get your hands dirty" child of the parent topic. The parent gave you the machinery; here we run every kind of input through it — best case, worst case, zero, degenerate, a real word problem, and an exam trap. Nothing new is assumed: every symbol used below is re-anchored the moment it appears.

Definition The symbols we reuse everywhere (read this first)

Before any arithmetic — even before the scenario table — pin these down in plain words:

$N$ = the number of data elements you want processed (e.g. array length). A plain count, no units.
$B$ = block size = threads per block (you choose it; max 1024).
A warp = a fixed group of 32 threads the hardware runs in lockstep. Always 32.
SM = Streaming Multiprocessor = one self-contained "engine" on the GPU chip that owns a batch of CUDA cores, their warp schedulers, and a shared register file. A GPU has many SMs; each runs many warps. Whenever you read "the SM caps at 64 warps", it means one such engine can only keep 64 warps resident at once.
$η$ (Greek letter "eta") = transaction efficiency = (bytes you actually needed) ÷ (bytes the hardware physically moved). A pure fraction between 0 and 1. This is the symbol used in the matrix's Cases D and E below.

The scenario matrix

Every cell below is a distinct case class this topic can throw at you. The worked examples are labelled with the cell they hit. (Every symbol in this table — $N$ , $η$ , SM — was defined in the block just above.)

#	Case class	The extreme / corner it tests	Example
A	Perfect launch	threads divide evenly into warps, all cores fed	Ex 1
B	Ragged tail	`N` not a multiple of block size → a partial final warp	Ex 2
C	Degenerate: N ≤ 0 or N tiny	fewer threads than one warp; idle cores; negative/zero	Ex 3
D	Memory best case	coalesced access, efficiency $η = 1$	Ex 4
E	Memory worst case	strided access, efficiency $η$ minimal	Ex 4
F	Branch divergence	warp splits into two paths, serialized	Ex 5
G	Latency hiding — succeeds	enough warps to cover the stall	Ex 6
H	Latency hiding — fails (limit)	required warps exceed the SM cap	Ex 6
I	Real-world word problem	image processing, choose a launch config	Ex 7
J	Exam twist	"128 cores ÷ 32 = ¼ cycle" trap + register ceiling	Ex 8

Prerequisites if any row feels shaky: 6.1.03-SIMD-vs-SIMT, 6.2.02-GPU-memory-hierarchy, and later 6.2.04-occupancy-and-performance.

Ex 1 — Cell A: the perfect launch

Worked example Perfect launch

You process $N = 2048$ floats with block size $B = 256$ . How many blocks, how many warps total, and are any threads wasted?

Forecast: guess the number of warps before reading on. (Hint: 2048 ÷ 32.)

Step 1 — Blocks needed. $numBlocks = ⌈ \frac{N}{B} ⌉ = ⌈ \frac{2048}{256} ⌉ = 8.$ Why this step? You launch whole blocks only; the ceiling $⌈ \cdot ⌉$ ("round up") guarantees at least one thread per data element. Here it divides evenly so no rounding happens.

Step 2 — Warps per block. $warps/block = \frac{B}{32} = \frac{256}{32} = 8.$ Why this step? A block is chopped into warps of 32. $256/32$ is a whole number, so no partial warp.

Step 3 — Total warps. $8 blocks \times 8 warps/block = 64 warps .$ Why this step? Total warps = how many lockstep groups the whole grid produces.

Verify: $64 \times 32 = 2048$ threads = exactly $N$ . Zero wasted threads. Every core that runs gets real work. ✅

This is the "textbook happy path" — everything divides. Every later example is a deliberate break from this.

Ex 2 — Cell B: the ragged tail

Worked example Ragged tail (N not a multiple of B)

Now $N = 1000$ with the same $B = 256$ . How many blocks? How many idle threads exist, and where?

Forecast: more or fewer than 4 blocks? And is the last warp full?

Step 1 — Blocks (with the ceiling doing real work now). $numBlocks = ⌈ \frac{1000}{256} ⌉ = ⌈ 3.906 ⌉ = 4.$ Why this step? $3$ blocks cover only $768$ threads $< 1000$ . You need a 4th block for the remaining $232$ elements. The ceiling is exactly the "always round up so nobody is left out" rule.

Step 2 — Total threads launched vs needed. $4 \times 256 = 1024 threads launched, 1024 - 1000 = 24 idle threads .$ Why this step? Those 24 extra threads have global index $i \geq N$ ; the if (i < N) guard in the kernel makes them do nothing.

Step 3 — Which warp is ragged? (built up index by index). First, locate the last block. Block indices run $0, 1, 2, 3$ ; the 4th block (blockIdx $= 3$ ) owns global thread indices from $3 \times 256 = 768$ up to $768 + 256 - 1 = 1023$ . Now split that block into warps of 32 by local index. Local index $0$ is global $768$ ; local index $31$ is global $799$ ; and so on. Warp number $w$ inside the block covers local indices $32 w$ to $32 w + 31$ . The last active global index is $N - 1 = 999$ . Its local index inside block 3 is $999 - 768 = 231$ . Which warp holds local index $231$ ? Divide: $⌊ 231/32 ⌋ = 7$ , so warp 7 of the block is the ragged one. Warp 7 covers local indices $32 \times 7 = 224$ up to $255$ , i.e. global indices $768 + 224 = 992$ up to $1023$ . Count the active threads in warp 7: they are the global indices from $992$ up to the last active one $999$ . That count is (last active − first) + 1 $= 999 - 992 + 1 = 8$ active threads. The remaining $32 - 8 = 24$ lanes are the idle ones (global $1000..1023$ , all $\geq N$ ). Why this step? We walked global → local → warp explicitly so the " $999 - 992 + 1 = 8$ " is not magic: $992$ is where warp 7 begins in global space, $999$ is the last real element, and the "+1" is because both endpoints are inclusive.

Verify: active threads $= 4 \times 256 - 24 = 1000 = N$ . ✅ The 24 idle lanes all sit in warp 7 of the last block; all other warps are full.

Ex 3 — Cell C: degenerate & tiny inputs (N ≤ 0 and N tiny)

Worked example Fewer threads than one warp — and the illegal case

Three sub-cases with $B = 256$ : (a) $N = 5$ ; (b) $N = 0$ ; (c) $N = - 3$ .

Forecast: for $N = 5$ , how many of the 32 cores in warp 0 actually do useful work? And what should happen for a negative $N$ ?

Step 1 — $N = 5$ : blocks. $⌈ \frac{5}{256} ⌉ = 1 block .$ Why this step? Even a single element needs one whole block launched.

Step 2 — Warp utilisation for $N = 5$ . One warp (32 threads) is issued; only threads $0..4$ pass i < 5. So $5/32 \approx 15.6%$ of the warp is useful. Why this step? This is the SIMT tax: the warp is atomic — you cannot run "just 5 lanes"; you run 32 and mask 27 off. See 6.1.03-SIMD-vs-SIMT for why masking, not skipping, is the mechanism.

Step 3 — $N = 0$ : the empty launch. $⌈ \frac{0}{256} ⌉ = 0 blocks .$ Why this step? A grid of zero blocks launches nothing — the kernel returns immediately. No cores run. This is the true degenerate floor.

Step 4 — $N = - 3$ : the illegal input. A count of elements cannot be negative, so $N = - 3$ is meaningless data, not a valid scenario. Two things protect you: (i) the launch-config formula $⌈ N / B ⌉$ with $N = - 3, B = 256$ gives $⌈ - 0.0117 ⌉ = 0$ blocks — so even if it slipped through, nothing launches; (ii) more importantly, the kernel guard if (i < N) compares each thread's index $i \geq 0$ against a negative $N$ , which is never true, so every thread falls through and does nothing. Why this step? The matrix promised " $N \leq 0$ ", so we must show the sign case, not just zero. The takeaway: negative $N$ collapses to "do nothing", exactly like $N = 0$ — the guard's < comparison is what makes both safe.

Verify: $N = 5$ → active/total $= 5/32 = 0.15625$ . $N = 0$ → 0 blocks, 0 warps. $N = - 3$ → $⌈ - 3/256 ⌉ = 0$ blocks, guard never true → 0 active threads, no crash. ✅

Ex 4 — Cells D & E: memory best case vs worst case

Worked example Coalesced vs strided access

A warp of 32 threads each reads one float (4 bytes). Compute the transaction efficiency $η$ for (D) thread $i$ reads A[i], and (E) thread $i$ reads A[i*32].

Forecast: case E — will $η$ be closer to 50% or to 3%?

Recall the definition (from the parent, and from our symbol block): $η = \frac{useful bytes requested}{transactions \times 128 bytes/transaction} .$ Useful bytes are the same in both cases: $32 \times 4 = 128$ bytes. Only the transaction count changes.

Step 1 — Case D, useful bytes. $32 threads \times 4 bytes = 128 useful bytes .$ Why this step? Consecutive addresses A[0..31] span exactly one 128-byte aligned line (look at the amber band in the figure — all 32 arrows land inside it).

Step 2 — Case D, transactions & $η$ . $transactions = 1, η = \frac{128}{1 \times 128} = 1 = 100%.$ Why this step? One 128-byte transaction carries every needed byte; nothing wasted.

Step 3 — Case E, transactions. Addresses are A[0], A[32], A[64], … — each 128 bytes apart, so each lands in a different 128-byte line (the scattered cyan arrows in the figure). That is $32$ separate transactions. Why this step? The hardware can only merge threads that share a 128-byte line; stride-32 defeats merging completely.

Step 4 — Case E, $η$ . $η = \frac{128}{32 \times 128} = \frac{1}{32} = 0.03125 = 3.125%.$ Why this step? Every transaction hauls 128 bytes but you keep only 4 of them → $4/128 = 1/32$ .

Verify (units): $η$ is bytes ÷ bytes = dimensionless, and $0 < η \leq 1$ in both. ✅ Slowdown factor = $η_{D} / η_{E} = 1/ (1/32) = 32 \times$ . That matches the parent's "up to 32× slower". ✅

Ex 5 — Cell F: branch divergence

Worked example A warp that splits down the middle

Inside a warp of 32 threads, half take the if and half take the else:

if (threadIdx.x % 2 == 0) { heavy_path(); }  // 20 cycles
else                       { light_path(); }  // 8 cycles

How many cycles does the warp spend, and what is the utilisation?

Forecast: does it take $20$ , $8$ , $20 + 8$ , or $max (20, 8)$ cycles?

Step 1 — Why a warp cannot do both at once. All 32 lanes share one instruction pointer (that is the "SI" in SIMT). When lanes disagree, the hardware runs one path with the other lanes masked off (idle), then the other path — serialized. See the two stacked timelines in the figure.

Step 2 — Total cycles. $cycles = 20 (even lanes) + 8 (odd lanes) = 28.$ Why this step? Paths run one after another, so costs add, not max.

Step 3 — Utilisation. During the 20-cycle heavy path, only 16 lanes are active; during the 8-cycle light path, the other 16. Useful lane-cycles $= 16 \times 20 + 16 \times 8 = 448$ . Total lane-cycles the warp occupied $= 32 \times 28 = 896$ . $utilisation = \frac{448}{896} = 0.5 = 50%.$ Why this step? Half the lanes sit idle at every moment of a two-way split — the classic divergence penalty.

Verify: if instead all 32 lanes took the heavy path (no divergence), cost $= 20$ cycles, utilisation $= (32 \times 20) / (32 \times 20) = 100%$ . Divergence turned 20 cycles into 28 and 100% into 50%. ✅

Ex 6 — Cells G & H: latency hiding, success and failure

Worked example Do we have enough warps to hide memory latency?

Using the parent's formula $T_{required} = \frac{S \cdot L}{C}$ where $S$ =schedulers, $L$ =stall cycles, $C$ =useful compute cycles per warp between stalls. The SM caps at $64$ resident warps. (G) $S = 4, L = 200, C = 20$ . (H) $S = 4, L = 200, C = 10$ .

Forecast: which case can fully hide the latency?

Step 1 — Why this formula. Over the $L$ stall cycles of one warp, the SM can issue $S \cdot L$ warp-instructions; each other warp supplies $\approx C$ before it too stalls, so you need $S \cdot L / C$ warps to keep the pipe full. (Full reasoning is in the parent.)

Step 2 — Case G. $T_{required} = \frac{4 \times 200}{20} = 40 warps .$ Why this step? Higher compute intensity $C = 20$ means each warp buys more coverage per stall.

Step 3 — Compare to the cap. $40 \leq 64$ ⇒ latency fully hidden. The scheduler always finds a ready warp. Why this step? The required warps must actually fit on the SM; comparing $T_{required}$ against the 64-warp cap is the only thing that tells you whether hiding is physically possible.

Step 4 — Case H. $T_{required} = \frac{4 \times 200}{10} = 80 warps > 64 \Rightarrow cannot hide fully.$ Why this step? Halving $C$ doubled the requirement past the hardware ceiling — the SM stalls, exposing memory latency. This is exactly why raising compute intensity (see 6.2.04-occupancy-and-performance) matters.

Verify: $40$ needed vs $64$ available → margin of $24$ warps (G succeeds). $80 > 64$ → deficit of $16$ warps (H fails). ✅ Doubling $C$ from 10→20 halved $T_{required}$ from 80→40, confirming the inverse relationship.

Ex 7 — Cell I: real-world word problem

Worked example Grayscale conversion of an image

You convert a $1920 \times 1080$ RGB image to grayscale, one thread per pixel, using 2D blocks of $16 \times 16$ threads. How many blocks in each dimension, total threads, and how many are wasted?

Forecast: does $1080$ divide evenly by $16$ ?

Step 1 — Total pixels (the real $N$ ). $N = 1920 \times 1080 = 2, 073, 600 pixels .$ Why this step? One thread per pixel means $N$ = number of threads you need.

Step 2 — Blocks per dimension (ceiling in 2D). $blocks_{x} = ⌈ \frac{1920}{16} ⌉ = 120, blocks_{y} = ⌈ \frac{1080}{16} ⌉ = ⌈ 67.5 ⌉ = 68.$ Why this step? $1920/16 = 120$ exactly, but $1080/16 = 67.5$ → round up to 68, creating a ragged bottom edge (cell B reappears, in 2D!).

Step 3 — Threads launched vs wasted. $threads = (120 \times 16) \times (68 \times 16) = 1920 \times 1088 = 2, 088, 960.$ $wasted = 2, 088, 960 - 2, 073, 600 = 15, 360.$ Why this step? The extra $1088 - 1080 = 8$ pixel rows × $1920$ columns $= 15, 360$ threads fall outside the image and are masked by the if (x<W && y<H) guard.

Verify: wasted $= 8 \times 1920 = 15, 360$ . ✅ And $2, 088, 960 - 15, 360 = 2, 073, 600 = N$ . ✅ Waste fraction $= 15360/2088960 \approx 0.74%$ — tiny, because only one edge is ragged.

Ex 8 — Cell J: the exam twist (two traps in one)

Worked example "A quarter cycle" and the register ceiling

Part 1: An SM has 128 CUDA cores. A student claims one warp instruction finishes in $32/128 = 1/4$ cycle. True or false? Part 2: The SM has $M = 65, 536$ registers and you want $P = 2048$ resident threads. What is the per-thread register ceiling, and what happens if your kernel needs 40 registers/thread?

Forecast: guess whether Part 1 is true, and guess the register ceiling before computing.

Step 1 — Part 1, why it's false. A warp instruction is atomic per issue: its minimum is 1 cycle, never less. Extra cores give throughput (more warps issued in parallel), not lower latency for one warp. With 128 cores and 32-thread warps, the SM can issue up to $128/32 = 4$ warp-instructions simultaneously — that's 4 warps per cycle, not one warp in a quarter cycle. Why this step? Latency (time for one warp) and throughput (warps per unit time) are different axes. The trap swaps them.

Step 2 — Part 2, the register ceiling. Here $M = 65, 536$ is the total number of 32-bit registers the SM's register file holds, and $P = 2048$ is how many threads you want resident (alive) at once. All those live threads share the one register file, so each thread may use at most: $\frac{M}{P} = \frac{65 , 536}{2048} = 32 registers/thread .$ Why this step? If $P$ live threads must split $M$ registers with none left over, the fair share per thread is exactly $M / P$ . Exceed it and you cannot keep all $P$ threads resident.

Step 3 — Kernel needs 40 registers/thread (> 32). Your kernel demands 40 registers per thread, but the ceiling above is only 32 at full occupancy — so full occupancy is impossible. Instead, resident threads must drop until $40 \times P \leq M$ : $P_{m a x} = ⌊ \frac{M}{40} ⌋ = ⌊ \frac{65 , 536}{40} ⌋ = ⌊ 1638.4 ⌋ = 1638 threads .$ Why this step? You can never exceed $M$ total registers, so more registers per thread forces fewer threads on the SM. Fewer resident threads → fewer resident warps → less latency hiding (this feeds straight back into Ex 6, where dropping below the required warp count exposes memory stalls).

Step 4 — What it costs in warps. $warps lost = \frac{2048 - 1638}{32} = \frac{410}{32} \approx 12.8 warps fewer resident.$ Why this step? Occupancy is really measured in warps; converting the thread drop to warps ( $\div 32$ ) shows the real hit to the scheduler's pool. Losing ~13 warps can push you from "hides latency" (Ex 6 Case G, needs 40) to "cannot" if margins were tight.

Verify: $32 \times 2048 = 65, 536 = M$ exactly. ✅ And $40 \times 1638 = 65, 520 \leq 65, 536$ ✅ while $40 \times 1639 = 65, 560 > 65, 536$ ✗ — confirming 1638 is the true cap. Part 1: $128/32 = 4$ warps/cycle throughput, latency still 1 cycle. ✅

Recall — one line per cell

Recall Why does the ragged tail still cost a full warp?

Because a warp is atomic — 32 lanes are issued together; idle lanes are masked, not skipped. ::: A warp is always issued as 32 lanes; the guard if (i<N) masks the extra ones but they still occupy the warp slot.

Recall What happens for N = 0 or negative N?

Both collapse to "do nothing": $⌈ N / B ⌉ = 0$ blocks and the i < N guard is never true. ::: Zero or negative N launches 0 blocks and every guard comparison fails, so no thread does work — no crash.

Recall Case E: efficiency of stride-32 access?

$η = 1/32 = 3.125%$ — 32 transactions of 128 bytes each for only 128 useful bytes. ::: $η = 128/ (32 \times 128) = 3.125%$ .

Recall A two-way branch split: cost and utilisation?

Costs add ( $20 + 8 = 28$ cycles), utilisation drops to 50% because half the lanes idle at all times. ::: Paths serialize, so cycles add and utilisation halves for an even split.

Recall Latency hiding formula and when it fails?

$T_{required} = S L / C$ ; it fails when this exceeds the ~64-warp SM cap. ::: When $S L / C$ > resident-warp cap, latency is exposed; raise $C$ to lower the requirement.

Recall Register ceiling: 40 registers/thread on a 65,536-register SM?

Full occupancy allows only 32/thread, so you drop to $⌊ 65536/40 ⌋ = 1638$ resident threads. ::: More registers per thread forces fewer resident threads ( $P_{m a x} = ⌊ M /40 ⌋ = 1638$ ), lowering occupancy.

Recall The "quarter cycle" trap?

False — a warp instruction is $\geq 1$ cycle; 128 cores give 4 warps/cycle of throughput, not sub-cycle latency. ::: One warp instruction is atomic at 1 cycle minimum; extra cores add throughput, not lower latency.

See also 7.3.01-parallel-programming-patterns and 9.1.02-neural-network-training-on-GPUs where these exact corners decide real performance.