Threads, blocks, and the memory hierarchy
The grid/block/thread mapping turns a loop into thousands of lanes. Where a value lives — register, shared, global — then decides how fast those lanes run.
Here is the mental shift CUDA asks for. In a CPU program, a loop over a million elements is one instruction stream that walks the array. In CUDA, you delete the loop. Instead you launch a million threads, give each one an index, and let each compute one element. The kernel body looks like the inside of the loop; the loop itself becomes the launch configuration.
The launch is two nested numbers. You launch a grid of blocks; each block contains some number of threads. Every thread knows its own coordinates and derives a unique global index from them. The hardware schedules whole blocks onto streaming multiprocessors and runs the threads of a block together, which is why block size is a performance decision and not just bookkeeping.
Three coordinates, one index
A thread is addressed by (blockIdx, threadIdx). Flatten them with the
block size and you have the element it owns:
.
Nothing is shared between threads except what they explicitly read or write —
each has its own registers and its own index.
Click threads below to see the flattening, then move the shared-memory slider to feel the latency ladder.
Click any thread — read its global index, then see where its memory lives
grid · 4 blocks × 16 threads = 64 threads
block 0
block 1
block 2
block 3
selected thread
blockIdx = 1, threadIdx = 3
global = blockIdx × 16 + threadIdx = 19
average access latency
171 cycles
= 70% × 30 + 30% × 500 (L2 ignored for simplicity)
Blocks are independent, so their order does not matter — only the per-block work must be correct. That is why the standard move is to stage reused data in fast shared memory before computing: moving a value from global to shared costs one slow trip, but every later read is roughly 16× cheaper. Latencies are illustrative CUDA-era figures and vary by architecture.
The mapping in code
For an array of elements, the canonical launch is:
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
out[i] = a[i] + b[i];
}
Two things matter. First, the guard if (i < N): grids are sized in whole
blocks, so the last block almost always overshoots. Omitting the guard writes
past the end of the array. Second, the index arithmetic is the only difference
between a thread that works and a thread that does not — there is no loop
counter to carry state between iterations.
Block size is a power of two as a rule: 128 or 256 is a common default, capped at 1024 threads per block. Smaller blocks give the scheduler more to interleave (recovering latency); larger blocks amortise more reuse through shared memory. The right point is empirical.
The memory hierarchy
The reason a GPU can hide huge latencies is that it has many warps resident per multiprocessor: while one warp waits on memory, others compute. Latency, in cycles, changes by orders of magnitude depending on where a value sits:
- Registers — private to a thread, effectively free (about one cycle); the compiler allocates them.
- Shared memory / L1 — a small, explicitly managed scratchpad shared by a block (tens of cycles). Fast enough to stage reused tiles.
- L2 — chip-wide cache, shared by all blocks (hundreds of cycles).
- Global memory (DRAM) — the big array (hundreds to over a thousand cycles). Large, but slow.
The standard performance pattern follows directly: load a tile from global to shared once, synchronize, then read it many times from shared. The first read pays the slow trip; every reuse is roughly an order of magnitude cheaper.
Coalescing
Latency is only half of memory cost; the other half is how many bytes move per transaction. When the 32 threads of a warp access consecutive addresses, the hardware serves them in a single wide transaction — coalesced. When their addresses are scattered, the same 32 requests split into many transactions and you waste most of the bandwidth you paid for. This is why data layout (array-of-structs versus struct-of-arrays) is a first-class GPU decision, and why a transposed access pattern can be several times slower for identical arithmetic.
Careful
The classic mistake is assuming threads are cheap to create. Blocks are scheduled, not threads; a block that cannot be resident waits. If you launch far more blocks than fit simultaneously, they queue — correct, but slower. And an out-of-range thread without a guard is not a warning, it is silent memory corruption that may not surface until much later.
Illustrative vs real
The cycle counts (1 / 30 / 200 / 500) are representative CUDA-era figures chosen for the comparison; current architectures differ, and shared memory latency is closer to L1 than the label suggests. The hierarchy is real and worth internalising, but treat the exact numbers as a sketch and read the programming guide for your target GPU.
Check yourself
Eduspheria wiki · Systems for AI, GPU computing
0 / 5 answered
From the exam paper
Modeled on NITJ AI-619, End-Sem June 2025
0 / 5 answered
Where next: synchronization, streams and profiling — keeping blocks correct and finding the bottleneck.