Foundations — Thread blocks and grids

Before you can read the parent note, you need to own every word and symbol it throws at you. This page builds them one at a time, from nothing, each one earning the next.

1. Thread — the single worker

The picture: imagine one person at a desk with a tiny task card that says "add A[?] + B[?]". The ? is not filled in yet — that's the whole point of the indexing we build later. Every thread runs the same card; only the number it plugs into ? differs. Figure s01 shows exactly this: one desk, one card, and a pink arrow pointing at the empty ? we still have to fill.

Why the topic needs it: the GPU's power is that it has thousands of these desks working at the same instant. But thousands of identical workers all running the same card would all do the same number — useless. So we need a way to give each worker a different index. That need drives everything below.

2. Block — a team of threads

The picture: take the desks from before and group them into a room. Everyone in one room:

• sits close together,
• shares one whiteboard on the wall (shared memory),
• can shout "everybody wait here!" and pause until all roommates catch up (`__syncthreads()`).

Why the topic needs it: hardware cannot physically give a million workers one giant shared whiteboard — wires would be too long and too slow. Grouping into small rooms keeps the shared whiteboard tiny and fast. The block is the boundary of "who can cheaply cooperate."

3. Grid — the map of all blocks

The picture: now zoom out from one room to the whole building, a floor plan tiled with identical rooms. The building covers the entire warehouse floor (the whole array, the whole image). Rooms do not talk to each other during the job — each is autonomous. Figure s02 stacks the two levels: blue dots are threads, each yellow box is a block, and the whole picture is the grid.

Why the topic needs it: one block can only hold ~1024 threads, but your array might have a million elements. You need many blocks. The grid is simply "all the blocks, arranged to blanket the whole data set."

4. The four built-in coordinates (and their .x .y .z fields)

Every thread, the moment it starts, is handed four little rulers. These are the raw symbols the parent note uses.

The picture (a hotel): blockIdx = which floor you're on, blockDim = rooms-per-floor (a constant), threadIdx = your room number on that floor, gridDim = how many floors the hotel has. Figure s03 draws this hotel and circles the "you" room so each of the four rulers has a physical meaning.

Why the topic needs them: these four are the only facts a thread knows about itself. Everything the thread does — which array element it touches — must be computed from these four numbers.

5. The multiply-and-add: globalID

Now every symbol is defined, so we can read the parent note's central formula honestly. We start with the simplest case, then extend to 2D and 3D using the .x .y .z fields.

What it does, in plain words: "count all the workers in the rooms before me, then add my seat number in this room."

Why multiplication? Each earlier room contributed exactly blockDim.x workers. If I'm in room number blockIdx.x, then blockIdx.x complete rooms came before me. To count them all: (number of rooms before me) × (workers per room) = blockIdx.x × blockDim.x. Multiplication is repeated addition — it's the shortcut for "add blockDim.x, blockIdx.x times."

Why then add threadIdx.x? That count only reaches the front door of my room. To find my exact seat, I step in threadIdx.x more places.

What it looks like: a number line where the rooms are equal-length fences and globalID is your absolute position on the whole fence. Figure s04 draws that fence and marks one worker's absolute seat.

The same idea in 2D and 3D. In higher dimensions you compute one coordinate per axis using exactly the same multiply-and-add, once for .x, once for .y, (once for .z):

6. The bracket symbols ⌈ ⌉ — ceiling (and what N and B mean)

The parent note writes $\lceil N/B\rceil =(N+B-1)/B$. Before touching the brackets, pin down the two letters.

Why we round up, not down: if you have N = 1000 elements and blocks of B = 256 threads (i.e. blockDim.x = 256), then 1000/256 = 3.9. Three blocks hold only 768 elements — 232 are left with no worker! You must add a fourth (partly-empty) block. You can never round down and leave data untouched, so ceiling is forced by the problem.

That extra empty block is exactly why the parent's kernel needs an if-guard if (i < N): the leftover workers in block 4 have no valid element and must sit still instead of touching memory that isn't theirs.

7. Warp — the marching squad of 32

One more word the parent leans on.

The picture: inside each room the workers don't move freely — they march in rows of 32, and the whole row must take the same step together. If half a row is told "turn left" and half "turn right," they can't do both at once, so the row does left while the right half stands frozen, then vice versa — wasted steps (divergence).

Why the topic needs it: this is why the parent says "block size should be a multiple of 32." A block of 40 threads still marches as two rows of 32 (64 slots), leaving 24 slots idle every step. Rounding block sizes to multiples of 32 wastes nothing.

Prerequisite map

Everything upstream flows into the parent topic Thread blocks and grids. Once these are yours, the parent's CUDA examples read as plain English.

Equipment checklist

Cover the right side and answer aloud. Every one must be a reflex before you go on.