CUDA programming model basics: Visual walkthrough — Hardware deep dive

This page goes deeper than the parent note. If a word here is new, we define it before using it. Prerequisites we lean on live in Thread Warps and SIMT and Streaming Multiprocessors.

Step 1 — One worker per job, laid out in a line

WHAT. Imagine you must add two lists of numbers, position by position: C[i] = A[i] + B[i]. Instead of one person walking down the list, you hire one worker per position. Worker for position 0, worker for position 1, and so on.

WHY. This is the whole point of a GPU: it has thousands of tiny workers. If each worker handles exactly one array slot, the entire addition finishes in one parallel step instead of N sequential steps. But this only works if each worker knows its own slot number. That number is what we must compute.

PICTURE. The bottom row is the array in memory: boxes labelled A[0], A[1], A[2], … Above each box stands one worker. Right now every worker is asking the same question: "which box is mine?"

Step 2 — Why we can't give all workers in one flat line

WHAT. You might think: just number the workers 0, 1, 2, … , N-1 in one long line and be done. The GPU does not let you do that. It forces you to pack workers into fixed-size groups first.

WHY. The hardware (Streaming Multiprocessors) can only manage a limited number of workers at once, and it schedules them in bundles. So CUDA makes you organise workers into equal-sized blocks (a block might hold 256 workers). Blocks are the unit the machine hands out. This is a hardware constraint, not a choice — which is exactly why the naive "one flat line" numbering is unavailable and we need a formula to reconstruct the flat number.

PICTURE. The single long row from Step 1 is now cut into equal chunks. Each chunk is a block. Inside a block the workers are re-numbered starting from 0 again — so worker numbering restarts at every block boundary.

The .x just means "the first dimension" — for a plain list, that's all we need. We meet .y in Step 7.

Step 3 — The two numbers each worker actually knows

WHAT. A CUDA thread does not know its flat global position directly. It is only told two local facts by the hardware:

blockIdx.x — the label on the box its block sits in.
threadIdx.x — how far along it stands inside that block.

WHY. The hardware can cheaply stamp each block with an ID and each worker with a within-block offset. Handing every one of a million threads a pre-computed global number would be wasteful; giving two small local numbers and letting arithmetic do the rest is far cheaper. So the formula's job is: turn (block label, position-in-block) into one flat address.

PICTURE. Zoom on block number 2, where blockDim.x = 4. The four workers wear tags threadIdx.x = 0,1,2,3. Stamped on the block itself: blockIdx.x = 2. The flat address we want is written faintly below each worker — and it is not the same as the tag they wear.

Step 4 — Counting the workers that came before you

WHAT. To find your flat address, first count everyone standing to your left across all earlier blocks. If each block holds blockDim.x workers, and you are in block number blockIdx.x, then the number of workers in all the blocks before you is: $how many whole blocks are before me blockIdx.x \times workers in each block blockDim.x$

WHY. Multiplication is exactly "add the same amount several times." Every earlier block contributed blockDim.x workers. Block 0 contributes one block's worth, block 1 another, … up to block blockIdx.x − 1. That's blockIdx.x blocks, each of size blockDim.x. This product is the starting flat address of your block.

PICTURE. With blockDim.x = 4: block 0 covers flat slots 0–3, block 1 covers 4–7, block 2 starts at $2 \times 4 = 8$ . The pale-yellow bracket spans the 8 workers before block 2; the arrow lands on the first slot of block 2.

Step 5 — Adding your step inside the block

WHAT. Now add how far you personally stand into your own block — that is threadIdx.x. Your flat address is the block's start plus your local offset: $globalIdx = start of my block blockIdx.x \times blockDim.x + my step inside the block threadIdx.x$

WHY. The block's starting slot belongs to the worker with threadIdx.x = 0. The next worker (threadIdx.x = 1) is one slot further right, and so on. So we add the within-block position to the block's start. That single sum is the flat global address — and it is unique, because no two workers share both the same block and the same within-block position.

PICTURE. Same block 2. blockStart = 8 (yellow). The worker with threadIdx.x = 3 gets 8 + 3 = 11. Blue arrow: jump to the block start. Pink arrow: step 3 further in. Landing slot 11 lights up.

Step 6 — The edge case: leftover workers past the end

WHAT. Blocks come in fixed sizes, so the total number of workers is almost never exactly N. If N = 1000 and blockDim.x = 256, we need $numBlocks = ⌈ \frac{1000}{256} ⌉ = 4$ blocks, giving $4 \times 256 = 1024$ workers — but only 1000 array slots. The last 24 workers (global indices 1000–1023) point past the end of the array.

WHY. You cannot launch a fractional block. You must round the block count up (the ceiling ⌈ ⌉), so you always launch at least enough workers — which means sometimes too many. Those extra workers, if they run C[idx] = ..., would read and write memory that isn't yours: a crash or silent garbage. The fix is a single guard line:

if (idx < N) { C[idx] = A[idx] + B[idx]; }

Workers whose globalIdx is 1000 or more simply do nothing.

PICTURE. The line of 1024 workers. The first 1000 stand over real array boxes (blue). The last 24 hang over empty space past the array's edge (pink, hatched "no box here"). A guard rail — the if (idx < N) line — stops them from reaching in.

Step 7 — The 2D case: rows and columns for images

WHAT. An image is a grid of pixels, not a line. Now each worker gets two pairs of numbers — one for the horizontal direction (.x) and one for the vertical (.y) — and we run the same Step-5 formula twice: $col = blockIdx.x \times blockDim.x + threadIdx.x$ $row = blockIdx.y \times blockDim.y + threadIdx.y$

WHY. Memory is still a flat line — pixel (row, col) of a width-wide image lives at flat slot row × width + col. So we compute the column with the exact same reasoning as Step 5 (using the .x numbers), compute the row identically (using the .y numbers), then fold the 2D address into a 1D memory address: $globalIdx = skip whole rows above me row \times width + step along my row col$ Multiplying by width "skips over" every full row above you — identical logic to multiplying by blockDim.x in Step 4, just at the pixel-row scale.

PICTURE. A checkerboard of blocks tiling an image. One worker highlighted: its (col, row) arrows point right and down; the flat address row × width + col is traced by counting whole rows (yellow) then stepping across (pink).

Recall Quick self-checks

For blockDim.x = 8, blockIdx.x = 5, threadIdx.x = 2, what is globalIdx? ::: $5 \times 8 + 2 = 42$ For N = 500, threadsPerBlock = 128, how many blocks? ::: $⌈ 500/128 ⌉ = (500 + 127) /128 = 4$ blocks (512 threads, 12 idle) Which threads must the if (idx < N) guard silence when N=500, launch=512? ::: indices 500–511 (12 extra threads) In a width = 640 image, pixel at row=3, col=10 sits at which flat slot? ::: $3 \times 640 + 10 = 1930$

The one-picture summary

Everything above collapses into a single map: two local numbers in, one unique global address out, with a guard rail for the overshoot.

Recall Feynman retelling — say it like a story

A GPU is a stadium of tiny workers, and they all run the same instruction. To make each one touch a different array slot, we can't just number them one-two-three across the whole stadium, because the hardware forces them into equal-sized blocks and restarts the counting inside every block. So each worker only knows two small facts: which block am I in (blockIdx.x) and how far along am I inside it (threadIdx.x). To rebuild its true seat number we first count everyone in all the blocks before it — that's blockIdx.x × blockDim.x, because every earlier block held exactly blockDim.x workers — and then add its own step inside the block, threadIdx.x. Sum them and you get a unique flat address that never repeats and never skips. Because blocks come in fixed sizes we usually hire a few too many workers, so the last handful point past the end of the array; a single if (idx < N) line tells those extras to sit still. For images we just do the whole trick twice — once sideways for the column, once downward for the row — and then flatten (row, col) into memory with row × width + col, because skipping a whole row of pixels means jumping width slots. That's the entire CUDA addressing scheme: two local numbers, one multiply, one add, one guard.

Related reading: Memory Hierarchy (where these indexed reads land), Parallel Algorithm Design (choosing block sizes), GPU Architecture Overview, and Deep Learning Frameworks which generate exactly these launches for you.