The GPU in One Picture

The idea in one sentence

A GPU runs one function as thousands of threads nested four levels deep, and the rule that decides how fast they run is that 32 threads sitting side by side should ask for 32 addresses sitting side by side.

The other bargain

Lecture 24 took one matrix multiply on one CPU and drove it from seconds to milliseconds. Every rung of that ladder spent its effort the same way: keeping a few fast cores fed out of a little fast memory.

A graphics processor makes the other bargain. It gives up on making any single core clever and buys thousands of plain ones instead. The chip this chapter measures against is NVIDIA’s RTX A6000. Its silicon is divided into 84 streaming multiprocessors, each carrying 128 arithmetic units, so 10752 multiplications can start in the same clock tick. NVIDIA’s Ampere GA102 whitepaper is where every hardware number in this chapter comes from.

Ten thousand arithmetic units are worth nothing on their own. Memory feeds them, and memory is where the design shows. So the picture you need is not a picture of the arithmetic. It is a picture of how the threads are grouped, because the grouping is what the memory system sees.

One function, four levels

You write one function. It is called a kernel, and you write it as though it will run once, for one piece of the work. Then you launch it, and the hardware runs that one function in many thousands of copies at the same time.

The launch says how many copies and how they are grouped. You give two numbers: how many blocks, and how many threads in each block. That pair is the grid. Every thread runs the same code and sees a different pair of indices, and the indices are the only thing that tells one copy from another.

Underneath your two numbers the hardware adds two levels of its own. A block is assigned to one streaming multiprocessor and stays there until it finishes, so the threads of a block can use that unit’s private fast memory and can be told to wait for each other. Inside the block the threads are cut into groups of 32, in order, and each group is called a warp. The warp is what actually issues instructions.

Each level owns one thing, and the rest of this chapter is those four sentences unpacked. The grid says how much work there is. The block says what shares a desk and can be made to wait. The warp says what runs in lockstep. The thread says what owns registers.

The index is an address again

A thread knows three things about itself. Its position inside its block, the position of its block inside the grid, and the size of a block. Write $t$ for the first, $b$ for the second, and $n_b$ for the third. Then the piece of work this thread owns is number

That is Lecture 23’s address formula wearing different clothes. There the entry in row $i$ and column $j$ sat at $p + i\,s_0 + j\,s_1$, an index times a stride plus an index times a stride. Here the block index carries the stride $n_b$ and the thread index carries a stride of one. A GPU launch is a strided walk over the work, and you are the one choosing the strides.

In CUDA the three names are threadIdx.x, blockIdx.x and blockDim.x, and the smallest useful kernel is the one that adds two vectors.

__global__ void add(const float *a, const float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

// one thread per element, 256 threads to a block
int threads = 256;
int blocks  = (n + threads - 1) / threads;
add<<<blocks, threads>>>(a, b, c, n);

Two details in those lines carry the model. The triple angle brackets are the launch, and they are the only syntax CUDA adds to C++ that you cannot look past. The guard if (i < n) is there because the block count was rounded up, so the last block runs off the end of the array and its spare threads have to do nothing. Round up and mask off is how nearly every kernel handles a size that is not a multiple of the block.

One instruction, thirty-two lanes

The warp is the level people skip, and skipping it is why their kernels are slow.

NVIDIA’s CUDA C++ Best Practices Guide puts it plainly: the smallest executable unit of parallelism on a CUDA device is 32 threads. A warp issues one instruction and all 32 lanes execute it. No arrangement of your code lets lane 5 be doing a multiplication while lane 6 does an addition.

So what happens at an if? The hardware runs both sides. It executes the true side with the false lanes switched off, then the false side with the true lanes switched off, and afterwards the warp is whole again. The two branches cost the sum of their times instead of the larger of them.

A branch that all 32 lanes take together is free. A branch that depends on the thread index, or on the data, is what quietly halves a kernel. Divergence is legal, nothing warns you about it, and that is exactly why it earns a drawing.

Four places a number can live

Now the memory. A GPU has four stores and they differ by more than speed. They differ by who can see them.

Registers are private to one thread. Shared memory belongs to one block and is invisible to every other block. The L2 cache is one pool for the whole chip. Global memory is the card’s own DRAM, which is where your data arrives and where your answer has to end up.

Read that top bar again, because it is the fact that shapes the next chapter. Each multiprocessor has a 256 KB register file and 128 KB of memory that it splits between L1 and shared memory. Across 84 multiprocessors that is 21 MB of registers against 10.5 MB of L1 and shared memory, and 6 MB of L2. On a CPU the register file is a rounding error next to the caches. On a GPU it is the largest fast store on the chip.

The speeds are as lopsided as the capacities. The whitepaper gives shared memory a bandwidth of 128 bytes per clock per multiprocessor, which is exactly the 32 banks of 4 bytes the programming guide describes. At the A6000’s boost clock of 1.80 GHz, across 84 multiprocessors, that peak is

in bytes per second, or 19.4 TB/s, against 768 GB/s from global memory. A factor of 25 for data you moved onto the chip once and then used many times.

Shared memory has thirty-two doors

That 19.4 TB/s is a peak, and the peak has a shape.

Shared memory is 32 banks, each four bytes wide, and each bank hands out one word per clock. The words are dealt to the banks in turn, the way cards go round a table: word 0 to bank 0, word 1 to bank 1, and after word 31 the deal starts again, so word 32 is back at bank 0. The four-byte word at index $w$ lives in bank $w \bmod 32$, and no arrangement of your code changes that map.

A warp issues one shared-memory instruction for all 32 of its lanes, and the banks decide what that instruction costs. If the 32 addresses land on 32 different banks, all 32 doors open together and the warp is served in one cycle. That is the case the 128 bytes per clock describes.

If two lanes want two different words from one bank, the bank can hand out only one of them this cycle. The best practices guide says what happens then: the hardware splits a request that has bank conflicts into as many separate conflict-free requests as it needs, and the effective bandwidth falls by that factor. Two lanes on one bank cost two cycles. Thirty-two lanes wanting 32 different words from one bank cost 32 cycles, and that is not a rare accident. A warp reading down a column of a tile that is 32 floats wide asks for words 0, 32, 64 and so on, every one of them in bank 0.

There is one exception, and it works in your favour. When several lanes ask for the same address, the bank reads that word once and broadcasts it to all of them, so a warp whose 32 lanes all want the same number costs one cycle, exactly like the perfectly spread case. An operand that every thread needs is free to share.

Nothing in the instruction says which of the three you got, and no compiler warning appears. The fix is to move the numbers. Padding a 32-wide tile to 33 floats a row shifts every row along by one bank, so a column walks the doors instead of hammering one, at the price of one wasted float per row. Lecture 26 pads. Lecture 26.5 does it without the waste, with a permutation of the addresses called swizzling that clears the rows and the columns at once.

Thirty-two threads, one transaction

One more rule, and it is the one the chapter’s single sentence is about.

Global memory is not delivered a float at a time. When a warp issues a load, the hardware works out which 32-byte pieces of memory its 32 lanes touched, and fetches those pieces. The best practices guide states it exactly: the accesses of a warp coalesce into as many transactions as there are 32-byte transactions needed to service all of its threads.

Do the arithmetic both ways. Thirty-two lanes reading 32 consecutive floats touch $32 \times 4 = 128$ bytes, which is four transactions, and every byte fetched is a byte somebody wanted. Thirty-two lanes reading 32 floats that are far apart touch 32 separate pieces, so the machine moves $32 \times 32 = 1024$ bytes to deliver the 128 that were asked for. Same instruction, same count of useful numbers, eight times the traffic.

This is Lecture 23’s cache line arriving whether you wanted all of it or not, moved up one level. There the wasteful walk was one core stepping down a column. Here it is 32 lanes of one warp stepping down 32 different rows, and the waste is the same waste.

The price has been measured. Simon Boehm’s CUDA matmul worklog profiles a naive matrix multiply on an RTX A6000 at 15 GB/s of global memory throughput, then changes which thread coordinate is attached to which matrix index, so that threads side by side read addresses side by side. Nothing else about the kernel changes. Throughput goes to 110 GB/s and the kernel from 309 to 1986.5 GFLOP/s, a factor of 6.4. That one swap is the second rung of the ladder in the next chapter.

Enough warps to hide the wait

A load from global memory takes a long time to come back. Coalescing reduces how much traffic there is and does nothing about the wait for any one load. The wait is handled another way, and that way is the reason a GPU wants thousands of threads rather than dozens.

Each multiprocessor is built as four partitions, and each partition has its own warp scheduler. Every cycle a scheduler looks at the warps living on its partition and issues an instruction from one whose operands are ready. A warp stalled on memory is not chosen that cycle. If some other warp is ready, the arithmetic units never notice that anybody was waiting.

So the question becomes how many warps a multiprocessor can hold at once. On this generation the ceiling is 48 resident warps, a number NVIDIA’s Ampere tuning guide lists along with the 65536 registers and the 128 KB of L1 and shared memory. The fraction of that ceiling a kernel reaches is called its occupancy. A block asks the multiprocessor for warp slots, for registers and for shared memory, and the multiprocessor keeps accepting blocks until one of those three runs out.

Boehm’s third kernel is the worked example, and the arithmetic is worth doing yourself. Its block is $32 \times 32$, so 1024 threads, which is 32 warps. The compiler gives each thread 37 registers, so the block needs $1024 \times 37 = 37888$ of the multiprocessor’s 65536 registers. Two blocks would need 75776 and do not fit. One block it is: 32 warps out of 48, an occupancy of 67%, which is what his profiler reports. Its shared memory, two $32 \times 32$ tiles of floats at 8192 bytes, is nowhere near the 99 KB a block is allowed, so shared memory is not what binds here.

Where this is going

You now have the model: a grid of blocks, a block of warps, a warp of 32 lanes, and four memory stores whose ratios are the whole performance story. Lecture 5’s batched forward pass, a batch of 32 images by 784 pixels meeting a $784 \times 128$ weight matrix, is one launch on this machine.

The next chapter takes that multiply seriously. It walks the ladder Lecture 24 walked on a CPU, rebuilt for thousands of threads, from a naive kernel at 1.3% of NVIDIA’s own library to a warp-tiled one at 93.7%. Every rung is one of the four levels above being used properly for the first time, and the largest single rung is the coalescing swap you just read about.

When you want to write kernels rather than read about them, the Library has the practice grounds: LeetGPU for a browser with a real GPU behind it, Triton Puzzles for the language most new kernels are written in, and the GPU MODE reference kernels for problems people compete on.