GPU architecture — SIMT, warps, CUDA model: Visual walkthrough

This is the visual companion to the parent GPU note. Read it slowly; each picture is a step of the argument.

Step 0 — The problem we are solving

WHY start here: notice that computing $c [5]$ needs nothing from computing $c [6]$ . Every entry is independent. That independence is the entire reason a GPU can help — we can do all $N$ additions at the same time if we had $N$ workers.

Figure — GPU architecture — SIMT, warps, CUDA model

Look at the picture: three shelves ( $a$ , $b$ , $c$ ), each with $N$ boxes. The arrow shows one worker grabbing box $i$ from $a$ , box $i$ from $b$ , adding them, and dropping the result in box $i$ of $c$ . There are no arrows between the workers — that is what "independent" looks like.

Step 1 — Give each worker a name (the thread index)

WHAT: we hire one thread (a tiny worker) per array slot, and we must tell each worker which slot is mine.

WHY: the GPU runs the same line of code in every thread. If every thread ran c[0]=a[0]+b[0], they'd all fight over slot 0 and ignore the rest. So each thread needs a different number $i$ . That number is computed from the thread's position.

The workers are handed out in blocks — think of a block as one classroom of workers. The GPU gives each thread two coordinates:

threadIdx.x = my seat number inside my classroom (starts at 0).
blockIdx.x = which classroom I am in (starts at 0).
blockDim.x = how many seats are in one classroom (the same for all).

PICTURE: two classrooms of 4 seats. Worker in classroom 1, seat 2 computes $i = 1 \cdot 4 + 2 = 6$ — and the arrow lands on global box 6. Every worker lands on a different box. That is the whole trick.

Step 2 — We usually hire too many workers (the guard)

WHAT: classrooms come in fixed sizes (say $B$ seats). If $N$ isn't a perfect multiple of $B$ , the last classroom has some empty-handed workers whose index $i$ points past the end of the array.

WHY: we must round the number of classrooms up so no data box is left uncovered, then tell the extra workers to sit down and do nothing.

The extra workers are switched off with:

if (i >= N) return;   // exit ONLY the out-of-range workers
c[i] = a[i] + b[i];   // everyone valid does the work

PICTURE: $N = 10$ boxes, block size $B = 4$ . We need $⌈ 10/4 ⌉ = 3$ classrooms = 12 workers. The last two workers (indices 10, 11) have no box — the red X shows them returning early. Boxes 0–9 all get covered.

Step 3 — The hardware bundles workers into warps of 32

WHAT: you wrote code for one scalar worker. The hardware does not schedule workers one at a time — it glues 32 consecutive threads into one rigid bundle called a warp and runs them together.

WHY: fetching and decoding an instruction is expensive. Instead of decoding "add" 32 times, the GPU decodes it once and broadcasts it to 32 arithmetic lanes. This is the "Single Instruction" in SIMT — Single Instruction, Multiple Threads.

PICTURE: one instruction ("ADD") sits at the top; a single fetch/decode box splits into 32 parallel lanes below it, each lane holding its own $a [i]$ , $b [i]$ and producing its own $c [i]$ . One instruction, 32 private data sets. Each lane keeps its own registers — that's the "Multiple Threads" part.

Step 4 — Memory is slow, so warps take turns (latency hiding)

WHAT: a warp asks for $a [i]$ and $b [i]$ from global memory. That load takes roughly 400 cycles to arrive. An addition takes 1 cycle. If the warp just waited, the 32 lanes would sit idle for 400 cycles.

WHY: the GPU refuses to wait. The moment a warp stalls on a load, the scheduler on the Streaming Multiprocessor (SM) instantly switches to another warp that already has its data. With enough warps queued, there is always one ready — so the arithmetic units never idle. This is latency hiding, and it is the reason GPUs keep so many warps alive.

PICTURE: a timeline. Warp 0 issues a load and grays out for 400 cycles (waiting). While it waits, Warps 1, 2, 3, … each issue their own work, filling the gap. By the time Warp 0's data returns, its slot in the timeline was never wasted. The 400-cycle latency is hidden underneath other warps' work.

Step 5 — Edge case: when the warp disagrees (divergence)

WHAT: our add is uniform — every lane does the same "+". But suppose the code had an if/else where some lanes go one way and some the other. There is only one fetch/decode unit per warp, so the two paths cannot run at the same time.

WHY it matters: the hardware runs the then path with the else lanes masked off (switched dark), then runs the else path with the then lanes masked off. The two paths are serialized.

PICTURE: a warp of 8 lanes (shrunk for clarity). In phase 1, the even lanes are lit doing heavy_A while the odd lanes are dark (masked). In phase 2, the odd lanes light up for heavy_B while the evens go dark. Total time = both bars end-to-end.

Step 6 — Edge case: how the 32 addresses are fetched (coalescing)

WHAT: when a warp loads $a [i]$ for its 32 lanes, the hardware collects those 32 addresses and packs them into as few 128-byte memory transactions as possible.

WHY it matters: if lane $i$ reads $a [i]$ (consecutive), the 32 reads of 4 bytes each = 128 bytes fall in one 128-byte line → one transaction. If lane $i$ reads $a [i \cdot 32]$ (scattered), each read lands in a different line → up to 32 transactions for the same 128 useful bytes.

PICTURE: top row — 32 arrows land neatly inside a single 128-byte line (one green transaction). Bottom row — 32 arrows scatter across 32 separate lines (32 red transactions) for the same 128 useful bytes. Same math, 32× the memory traffic.

The one-picture summary

This final board compresses the whole walkthrough: one scalar line $\to$ many threads indexed by $i = blockIdx \cdot blockDim + threadIdx$ $\to$ rounded-up blocks + guard so nothing is dropped $\to$ warps of 32 running one instruction across 32 lanes $\to$ the scheduler swaps stalled warps to hide 400-cycle loads $\to$ with two performance traps drawn on the side: divergence (warp disagrees $\Rightarrow$ paths add) and coalescing (addresses scatter $\Rightarrow$ transactions multiply).

Recall Feynman: tell the whole story to a 12-year-old

You have two shelves of numbers and want to add them box-by-box into a third shelf. Instead of one fast helper, you hire thousands of tiny helpers — one per box. Each helper figures out "which box is mine?" from classroom × seats-per-classroom + my-seat. Because classrooms come in fixed sizes, you round up the number of classrooms so no box is missed, then tell the leftover empty-handed helpers to sit down (if i >= N return). The GPU never runs helpers one at a time — it staples them into bundles of 32 called warps and shouts ONE instruction to all 32 at once. Grabbing numbers from memory is slow (400 heartbeats), so while one bundle waits, the boss instantly puts another ready bundle to work — nobody twiddles thumbs. Two things can wreck the speed: (1) if helpers in the same bundle disagree at an if/else, the bundle must do both sides one after the other — so keep whole bundles agreeing; (2) if the 32 helpers grab numbers that are scattered instead of next to each other, the machine makes 32 trips instead of 1 — so keep the reads neighborly. Same tiny "+" line, but the hardware turned it into a screaming-fast parallel machine.

Related vault topics: SIMD and Vector Instructions (AVX, SSE) · CPU Cache Hierarchy and Latency · Memory Coalescing and Bandwidth · CUDA Shared Memory and Tiling · Amdahl's and Gustafson's Laws · Matrix Multiplication on GPUs

Recall Quick self-test

Global index formula ::: $i = blockIdx.x \cdot blockDim.x + threadIdx.x$ Blocks needed for N=10, B=4 ::: $⌈ 10/4 ⌉ = 3$ Why + B - 1 ::: integer division rounds down; adding $B - 1$ rounds up so the last partial block survives Correct bounds guard ::: if (i >= N) return; Worst-case divergence cost of a 2-way branch ::: $T_{then} + T_{else}$ Coalescing efficiency for stride-32, 4-byte reads ::: $1/32$