Visual walkthrough — ROCm - OpenCL alternatives

This page derives one central result of the parent note in pictures: a formula that gives every single worker on a GPU its own unique home address. We will not write that formula down yet — every symbol in it has to be earned first. By the end you will see why the whole thing is just "counting boxes inside boxes", and you will know what happens at every edge — the last box, the leftover box, the empty case.

Read the parent first if you haven't: ROCm / OpenCL alternatives.

Step 1 — The problem: a huge pile of identical jobs

WHAT. Imagine we have a giant list of numbers and we want to do the same tiny operation to each one — say, add two arrays element by element: C[i] = A[i] + B[i]. Suppose the list has $16$ elements (we use a small number so it fits in a picture; the real note uses over a million).

WHY. A GPU is not one fast worker. It is thousands of slow-ish workers. To use it, we must hand each worker exactly one job and make sure no two workers grab the same job and no job is skipped. So the very first question is: how do we hand out job numbers $0,1,2,\dots ,15$ so every worker knows which one is theirs?

PICTURE. Here is the raw pile of jobs — a single row of numbered boxes.

Step 2 — The first symbol: the global ID

WHAT. You might say: "Just number them $0$ to $15$ in one long line and we're done." That flat position — the worker's address in the whole problem — is important enough to earn a name and a symbol right now.

WHY. We give it a symbol now because it is the quantity we are ultimately hunting for. Everything else on this page exists only to reconstruct $\text{global\_id}$ from the smaller pieces the hardware actually gives us. Hold onto this box-number picture; every later symbol will be measured against it.

PICTURE. The same row again, now with the name $\text{global\_id}$ written under it.

Step 3 — Why we can't just number them all at once

WHAT. So why not stop here and use $\text{global\_id}$ directly? Because the hardware won't let us hand out those flat numbers. It insists on chopping the line into equal chunks first. Each chunk is called a work-group.

WHY. Real silicon executes workers in fixed-size batches — NVIDIA runs them $32$ at a time (a warp), AMD runs them $64$ at a time (a wavefront). So we are forced to slice the flat line into equal groups that match the hardware's batch size.

PICTURE. Same $16$ boxes, but now sliced into groups of $4$. (We pick $4$ instead of $64$ only so it fits on screen — the logic is identical.)

$\underset{\text{total jobs}=16}{\underbrace{G}}÷\underset{\text{jobs per group}=4}{\underbrace{L}}=\underset{\text{how many groups}}{\underbrace{4}}$

Step 4 — Two rulers instead of one

WHAT. The hardware never tells a worker its $\text{global\_id}$ directly. Instead it gives each worker two smaller addresses:

• Its group ID — which box it lives in ($0,1,2,3$).
• Its local ID — its seat inside that box ($0,1,2,3$).

WHY. The hardware only ever knows "you are worker #2 inside group #1". That is a natural consequence of the batching from Step 3. So we are handed the two small rulers and must reconstruct the one big ruler ($\text{global\_id}$) ourselves.

PICTURE. The same strip, labelled twice: the bottom ruler counts groups, the top ruler restarts at $0$ inside every group.

Notice the top ruler resets to 0 at the start of every group. That reset is the whole reason we need a formula: local ID $2$ appears four different times, so it cannot be an address by itself.

Step 5 — Building the formula: skip past full boxes, then step in

WHAT. Now we finally assemble $\text{global\_id}$ from the two rulers. Take worker local ID 2 in group 1. To find its global number we do two things:

1. Skip over every box before ours. There is $1$ full box in front of us, each holding $L=4$ workers, so we skip $1\times 4=4$ addresses.
2. Walk in to our seat. Add our local seat number, $2$.

That gives $4+2=6$. Check against Step 4: yes, the box labelled global $6$ is exactly local-2-of-group-1.

WHY. This is nothing more than reading a number written in a mixed base. "Group $1$, seat $2$, with $4$ seats per group" is the same idea as "$1$ ten and $2$ ones $=12$", except each place is worth $L$ instead of $10$.

PICTURE. The red arrow first jumps over the full groups (that's the multiply), then steps to the seat (that's the add).

$\text{global\_id}=1\times 4+2=6✓$

Step 6 — Going backwards: recovering box and seat from a global number

WHAT. Sometimes we have the global number and want the other direction: which box, which seat? That is just the formula solved backwards using integer division and remainder:

$\text{group\_id}=\lfloor \frac{\text{global\_id}}{L}\rfloor ,\text{local\_id}=\text{global\_id}\mathrm{mod}L$

WHY. Integer division $\lfloor \cdot \rfloor$ answers "how many whole boxes fit below me?" and the remainder $\mathrm{mod}$ answers "how far past the last full box am I?" — the exact inverse of skip-then-step.

PICTURE. Worker global $6$: dividing by $4$ gives quotient $1$ (the box) and remainder $2$ (the seat). The two operations point back to the two parts of Step 5's arrow.

$\lfloor \frac{6}{4}\rfloor =1(\text{group}),6\mathrm{mod}4=2(\text{local})$

Step 7 — Edge case A: the very first and very last worker

WHAT. Check the extremes so nothing is ever undefined.

• First worker: group $0$, local $0$ → $0\times 4+0=0$. Correct: the front of the line.
• Last worker: group $3$, local $3$ → $3\times 4+3=15$. Correct: the final box, final seat, which is $G-1=16-1=15$.

WHY. We must prove the formula covers the boundaries, because off-by-one errors here are the classic GPU bug: read one box too far and you touch memory you don't own.

PICTURE. Both ends lit up — the global range runs exactly $0$ through $15$, no gap, no overshoot.

Step 8 — Edge case B: when the total doesn't divide evenly (the "guard" clause)

WHAT. Suppose we have $G=14$ real jobs but the hardware forces groups of $L=4$. Then $14/4=3.5$, which is not a whole number. The runtime rounds up the launch to $4$ groups $=16$ workers — but only $14$ jobs exist. Workers $14$ and $15$ have no real job.

WHY. The hardware can only launch whole groups. So we always spawn a few extra workers and must switch them off before they read past the end of the array. That is exactly the if (i < n) guard in every kernel:

int i = blockIdx.x * blockDim.x + threadIdx.x; // = global_id
if (i < n) y[i] = a * x[i] + y[i];             // <-- the guard

PICTURE. The two extra workers (global $14,15$) are drawn as ghosts; the red guard line drops them so they do nothing.

Step 9 — Edge case C: degenerate launches, including the empty one

WHAT. Three limiting shapes to check.

• One giant group ($L=G=16$): then group_id is always $0$, and $\text{global\_id}=0\times 16+\text{local\_id}=\text{local\_id}$. The two rulers collapse into one — the formula still holds.
• Groups of one ($L=1$): then local_id is always $0$, and $\text{global\_id}=\text{group\_id}\times 1+0=\text{group\_id}$. Again one ruler; formula still holds.
• No jobs at all ($G=0$): the line is empty, $\text{Num Groups}=\lceil 0/L\rceil =0$, so zero workers are launched. The formula is never evaluated — there is no worker to hold a $\text{global\_id}$ — and nothing goes wrong. An empty workload is a valid, quietly-does-nothing launch.

WHY. A formula you trust must survive its extremes. At $L=G$ and $L=1$ the "boxes inside boxes" picture flattens to a single ruler; at $G=0$ the picture is blank and the guard from Step 8 (if (i < n) with $n=0$) means no worker ever runs. In every case the arithmetic agrees.

PICTURE. Top: one box holding all $16$. Middle: $16$ boxes holding one each. Bottom: the empty line — no boxes, no workers.

The one-picture summary

Everything above is one idea: a flat line of jobs is secretly a grid of boxes, and the address formula is how you fold between the two views. Multiply to skip full boxes, add to step into your seat; divide and remainder to go back.