Visual walkthrough — Thread blocks and grids

We are answering ONE question the whole way down:

"I have thousands of workers split into teams. How does each worker figure out which single job in one long to-do list is theirs — with no worker ever guessing the same job as another?"

---

Step 1 — The problem laid flat: one long to-do list

WHAT. Picture our data as one straight row of boxes, numbered $0,1,2,\dots ,N-1$. This is the array C in vector addition (C[i] = A[i] + B[i]). Each box is one job.

WHY start here. Before we talk about teams, we must see the thing being divided. Memory on a GPU is fundamentally a 1D line of addresses — even a 2D image is stored as one long strip. So the target of all our indexing is a single number, an index into this line.

PICTURE. The blue strip below is the to-do list. Each cell is one element.

---

Step 2 — Cutting the list into equal teams (blocks)

WHAT. We chop the long strip into equal chunks. Each chunk is a thread block — a team of workers. Call the team size $B$ (in code: blockDim.x). The teams are numbered $0,1,2,\dots$ and a team's number is $k$ (in code: blockIdx.x).

WHY chop at all? A GPU cannot pour all workers into one pile — hardware forces teams onto separate Streaming Multiprocessors. Chopping the list into equal blocks means every team gets an identical slice of work, and the teams can be handed out to whichever processor is free. This is also why teams must be equal size: equal-size chunks make the arithmetic (coming in Step 4) a clean multiplication.

PICTURE. Same strip, now sliced by yellow dividers into blocks of size $B$. Notice block $k$ starts at a specific place on the strip — hold that thought.

---

Step 3 — Each worker's local number

WHAT. Inside a team, workers count themselves from $0$: worker $0$, worker $1$, up to worker $B-1$. This within-team number is $t$ (in code: threadIdx.x).

WHY a local number and not a global one? Because every team is identical, the hardware gives every worker the same little counter that resets to $0$ in each team. That is cheap and uniform. The cost: worker $t$ in team $0$ and worker $t$ in team $1$ share the same local number but must do different jobs. Steps 4–5 fix exactly this clash.

PICTURE. Zoom into two teams. Both have a worker labelled $t=2$ (red). They point at different boxes on the strip. The local number alone is ambiguous.

---

Step 4 — Counting the jobs that came before you

WHAT. Team $k$ is not the first team. Teams $0,1,\dots ,k-1$ came before it — that is $k$ complete teams, each owning $B$ jobs. So before team $k$ even starts, exactly $k\times B$ jobs are already spoken for. That product is the offset to the front of your team's slice.

WHY multiply? Multiplication is repeated addition of equal groups. We have $k$ equal groups of size $B$. Adding $B+B+\dots$ ($k$ times) is precisely $k\times B$. This clean product is only possible because Step 2 made the blocks equal — it is the payoff for that choice.

$$\underset{\text{how many whole teams are ahead of me}}{\underbrace{k}}\times \underset{\text{jobs each of those teams owns}}{\underbrace{B}}=\underset{\text{jobs claimed before my team begins}}{\underbrace{kB}}$$

PICTURE. The strip with the green bracket spanning boxes $0$ up to $kB-1$ — every job that belongs to an earlier team. Team $k$'s slice begins right after the bracket, at box $kB$.

---

Step 5 — Adding your seat inside the team: the formula is born

WHAT. You are worker $t$ inside team $k$. Your team's slice begins at box $kB$ (Step 4). You sit $t$ seats into that slice. So your job is:

$$i=\underset{\begin{array}{c}\text{start of my}\\ \text{team’s slice}\end{array}}{\underbrace{k\times B}}+\underset{\begin{array}{c}\text{my seat}\\ \text{inside the slice}\end{array}}{\underbrace{t}}$$

Written with the CUDA names:

$$i=\underset{k}{\underbrace{\text{blockIdx.x}}}\times \underset{B}{\underbrace{\text{blockDim.x}}}+\underset{t}{\underbrace{\text{threadIdx.x}}}$$

WHY add $t$? The multiplication got you to the door of your team's slice. The addition walks you $t$ steps in to your own seat. Multiply-then-add is the whole story: jump to your team, then step to your seat.

WHY is this collision-free? Two different $(k,t)$ pairs can never land on the same $i$. Since $0\le t<B$, the value $kB+t$ lands strictly inside $[kB,kB+B-1]$ — that is exactly team $k$'s block of boxes and no other team's. So every worker gets a distinct job, and every job in $[0,GB-1]$ gets exactly one worker. (Here $G$ = gridDim.x is the number of teams.)

PICTURE. Worker $(k=2,t=3)$ with block size $B=4$: the green arrow jumps to box $8=2\times 4$, then the red arrow steps $+3$ to box $11$. That is his job.

---

Step 6 — The leftover-worker edge case (why the if guard exists)

WHAT. Real problems rarely divide evenly. Suppose $N=10$ jobs and $B=4$. How many teams? We cannot use $3$ teams (that covers only $\mathrm{12...}$ wait — $3\times 4=12\ge 10$? Let's use the honest small case): $N=10$, $B=4$ needs $\lceil 10/4\rceil =3$ teams covering $3\times 4=12$ seats. But there are only $10$ jobs. Two workers have no job.

WHY it happens. We launched a whole number of equal teams, and the last team spills past the end of the list. Those extra workers would compute $i=10$ and $i=11$ — boxes that do not exist. Reading/writing there is illegal memory access → crash or garbage.

THE FIX. Every worker checks if (i < N) before touching memory. Workers with $i\ge N$ do nothing. The CUDA launch computes team count with ceiling division:

$$G=\lceil \frac{N}{B}\rceil =\frac{N+B-1}{B}\text{(integer division)}$$

The $+(B-1)$ pushes any nonzero remainder up to the next whole team, so no job is left uncovered — we always round up, never down.

PICTURE. Three teams over $N=10$. Green boxes are real jobs; the two red hatched boxes are phantom seats whose workers are silenced by the guard.

---

Step 7 — Two dimensions: the same idea, twice, then flattened

WHAT. For an image (or a matrix), the natural picture is a 2D field of pixels. We tile it with 2D blocks. Each thread finds its column and row using the exact same jump-then-step rule, once per axis:

$$\text{col}=\underset{k_x}{\underbrace{\text{blockIdx.x}}}\times \underset{B_x}{\underbrace{\text{blockDim.x}}}+\underset{t_x}{\underbrace{\text{threadIdx.x}}}\text{row}=\underset{k_y}{\underbrace{\text{blockIdx.y}}}\times \underset{B_y}{\underbrace{\text{blockDim.y}}}+\underset{t_y}{\underbrace{\text{threadIdx.y}}}$$

WHY two copies of the formula? The 2D grid is just the 1D idea run independently along $x$ (horizontal) and $y$ (vertical). $x$ counts across columns, $y$ counts down rows — the convention matching how we draw matrices (rows go down).

Then flatten. Memory is still a 1D line (Step 1!). A pixel at $(\text{row},\text{col})$ in an image of width $W$ lives at

$$\text{idx}=\underset{\text{skip whole rows above me}}{\underbrace{\text{row}\times W}}+\underset{\text{step across my row}}{\underbrace{\text{col}}}$$

This is row-major order: to reach my box, skip all complete rows above me ($\text{row}$ rows of $W$ boxes each — a multiply), then step $\text{col}$ across my own row (an add). It is Step 4 + Step 5 all over again, one level up.

PICTURE. A 2D pixel field tiled by blocks; one thread's $(\text{row},\text{col})$ found by two arrows, then the row-major flattening to a single index.

Recall Why is flattening also a multiply-then-add?

Because 2D memory is a lie — it is stored as one long line. ::: To reach row $r$ you must skip $r$ whole rows of $W$ elements ($r\times W$, the jump), then move col into your row (the step). Identical shape to the block formula.

---

The one-picture summary

Everything above is one sentence: jump over what came before, then step to your own seat — and do it once per dimension. The figure below stacks all three jump-then-step moves (1D global ID, the leftover guard, and 2D→flat) into a single diagram.

Recall The whole walkthrough in plain words (Feynman retelling)

Imagine a huge to-do list, one job per box, boxes numbered from zero. We can't dump all workers on it, so we cut the list into equal teams and number the teams. Inside each team the workers count themselves from zero — but that local number repeats in every team, so it can't be a job number by itself. Here's the trick: to find your job, first count all the jobs owned by teams ahead of you. Every earlier team owns the same number of jobs, so that count is just (your team number) times (team size) — a multiplication, because it's equal groups added up. That lands you at the door of your team's slice. Now step in by your own local seat number — an addition. Multiply to jump, add to step: that's the global ID, and it never collides because your seat stays inside your team's slice. If the list doesn't split evenly, the last team has empty seats pointing past the end, so each worker checks "is my job real?" before touching memory, and we round the team count up so nothing is missed. For images we just run the same jump-then-step twice — across for columns, down for rows — then flatten to the real 1D memory line by skipping whole rows (multiply) and stepping across our own row (add). Same idea, all the way down.

Related vault topics: Warp Execution · Shared Memory · Thread Synchronization · GPU Occupancy · Memory Coalescing · Parallel Algorithm Design.