← Chapter 17: A kernel is not a function Chapter 19: Why threads group into blocks →
date 06 / jul / 26
mtwtfss
THE TRANSLATION BARRIER — Chapter 18 index home

Why thread indexing exists

*the gps formula*

Mass production is useless without organization. Hire ten thousand workers, hand every one of them the same blueprint and the same tool, and they all build the exact same front door. A GPU has the same problem. It runs one piece of code across millions of threads, so unless each thread knows which piece of the data is its own, they all pile onto the same spot and do the same redundant work.

Thread indexing is the small formula that fixes this, the GPS that gives every thread a unique address and turns a chaotic mob into an orderly assembly line.

10,000 workerssame blueprint + tooleveryone builds the same doorGPU PROBLEMMillions of threads run thesame code. With no indexing,they all hit the same data.→ redundant work / data racesTHE SOLUTIONgive every thread a uniqueaddress with a small formula.chaotic mob → assembly line

The formula itself is short. A thread's index is its position inside its block, plus its block number times the size of a block. The hardware hands every thread three values for free, baked into special registers it can read instantly: which seat it holds inside its block, which block it belongs to, and how many threads a block contains. From those three numbers, each thread computes its own slot with pure arithmetic, no searching, no sorting, no operating system involved.

↳ Now the two questions that actually make this click, because the formula looks arbitrary until you try to break it.

THE GPS FORMULA (1D)global_id = blockIdx.x × blockDim.x + threadIdx.xthe hardware hands every thread these three — for free:threadIdx.xseat number inside the blockblockIdx.xblock number in the gridblockDim.xhow many threads per blockBLOCK 00123BLOCK 10123BLOCK 2012301234567891011GLOBAL MEMORY — one flat 1D tape, every slot uniqueblock × size = the big jump to your gate · seat = the small step to your chair

First Question: Why not just use the seat number?

The hardware already gives each thread a seat number inside its block. Why not just use that seat number as the memory address? Why all the extra math?

Because seat numbers repeat across blocks. Inside block 0 the threads are numbered 0, 1, 2, 3. But inside block 1 they are also numbered 0, 1, 2, 3. And block 2, again 0, 1, 2, 3. So if you used the seat number alone, the first thread of every block would all target slot 0 at the same time. Block 0's first thread, block 1's first thread, block 2's first thread, all crashing onto the same element.

↳ That is a data race, pure chaos. The seat number is only unique inside a block, never across the whole grid.

Use the seat number alone, and every block's first thread targets slot 0.BLOCK 00123BLOCK 10123BLOCK 201230ELEMENT 0 — all three crash hereDATA RACEThe seat number is unique only INSIDE a block — never across the whole grid.

Second Question: Why multiply by block size?

Fine, so mix in the block number. But why multiply the block number by the block size? Why not just add the raw block number?

Walk it through and the bug appears instantly. Say each block holds 4 threads. If the formula were just seat plus block number, then block 1's first thread computes 0 plus 1, which is 1. But slot 1 already belongs to block 0's second thread. Block 1's second thread computes 1 plus 1, which is 2, also already taken. The blocks overlap and stomp on each other's data. The raw block number only nudges you forward by one, when block 1 actually needs to start a full 4 elements later, past everything block 0 owns.

↳ That multiplication is the fix, and it's worth seeing what it physically means. Think of the memory as one long tape divided into zones, one zone per block.

Multiplying the block number by the block size is the stride, the big jump that lands you exactly at your block's starting gate. Block 0 starts at 0 × 4, which is 0. Block 1 starts at 1 × 4, which is 4. Block 2 at 2 × 4, which is 8. Each block lands cleanly at the front of its own zone with no overlap.

And then the seat number is the small local step from that gate to your exact chair. Block 1, seat 2, lands at gate 4 plus 2, which is element 6. The multiplication gets you to the right neighborhood, the addition gets you to the right house.

Add the raw block number and zones overlap. Multiply by block size and they don't.WRONGglobal = threadIdx + blockIdxBlock 00123Block 11234slots 1·2·3 collideOVERLAP — threadsstomp on each otherCORRECTglobal = blockIdx × blockDim + threadIdxBlock 00123Block 14567no overlapeach block getsits own clean zoneMEMORY = ONE TAPE, ONE ZONE PER BLOCK01234567891011Block 0 zoneBlock 1 zoneBlock 2 zoneBlock 1, seat 2 → gate 1×4 = 4 + seat 2 = element 6Multiply = jump to your zone's gate. Add = step to your seat.

Higher Dimensions: Collapsing onto a flat tape

So far this is all one-dimensional, a single line of data. But real data is rarely a line. An image is a 2D grid of pixels. An AI tensor can be 3D or higher. And yet physical memory is always a flat 1D tape, no matter how many dimensions your data pretends to have. So the real job of indexing is to collapse those dimensions down onto the flat tape.

For a 2D image, each thread now computes two coordinates instead of one, a column and a row, each with the same seat-plus-block-times-size pattern but along its own axis. Then to find the single flat address, you multiply the row by the full width of the grid and add the column. The row number times the width skips over all the complete rows above you in one jump, and the column slides you across the current row.

↳ A 3D volume just adds one more layer to this, where the depth coordinate multiplies by the area of a full 2D slice, skipping entire planes of data at a time before the row and column place you within the slice. Each higher dimension is just a bigger stride that leaps over a bigger chunk of the tape.

Physical memory is always a flat 1D tape — so indexing collapses every dimension onto it.2D IMAGE (width W = 8)★x (col)y (row) ↓flat index = y × W + x(3, 2) → 2×8 + 3 = 191617181920212223skip 2 full rows of 8, then step 3 acrossEACH DIMENSION = ONE STRIDEdimensionstridex (column)1y (row)Wz (slice)W × Hw (volume)W × H × Da 3D volume adds a depth stridethat leaps over whole 2D slicesHigher dimensions are just bigger strides leaping over bigger chunks of the tape.

The Performance Trap

And here is the trap that turns this from a math exercise into a performance issue, the part most people miss. The order in which you wire these coordinates into the formula decides your memory speed, because of how warps read memory.

Remember a warp is 32 threads moving in lockstep, and the hardware loves it when those 32 threads read 32 neighboring slots, because it can grab them all in a single memory transaction. But if you build your formula so that consecutive threads in a warp land on different rows instead of neighboring columns, those 32 threads now reach into 32 scattered locations, and the hardware has to do 32 separate transactions instead of one. Same data, but memory throughput collapses by an order of magnitude.

A warp of 32 reads fastest when its 32 threads hit 32 neighbouring slots.GOODthreadIdx.x → columns (contiguous)0123⋯3031warp0123⋯303132 neighbours→ 1 transactionFAST ✓BADthreadIdx.x → rows (strided)0123⋯3031warp32 scattered slots→ 32 transactionsSLOW ✗Map the fastest-changing index (threadIdx.x) to the most contiguous axis — the column.

↳ The rule that prevents this is simple:

The fastest-changing index must map to contiguous memory.

The thread's seat number, the thing that varies fastest across a warp, must be attached to the most contiguous part of the address, the column. Keep the fastest-changing thread index mapped to neighboring memory, and the 32 threads of a warp always read one clean contiguous block. Flip it, and you slaughter your bandwidth.

So the indexing formula is not just bookkeeping to keep threads from colliding. The exact order you arrange it in is the difference between feeding the warp in one gulp and starving it with 32 scattered fetches.

↳ The General Takeaway

Indexing is what converts a mob of identical threads into an assembly line where each one owns a unique slot. The block number times the block size jumps you to your zone, the seat number steps you to your spot, and for multi-dimensional data each extra dimension is just a larger stride leaping over rows or whole slices of the flat tape.

Get the formula right and you avoid collisions. Get the order of the formula right, with the fastest thread index on the most contiguous axis, and you also keep the memory fast.

ONE FORMULA, TWO PAYOFFSRIGHT FORMULAblock × size + seat→ no collisionsCORRECTRIGHT ORDERfast index on thecontiguous axis → coalescedFASTThe same little equation decides whether your code is correct AND whether it's quick.