What is the difference between a CPU and a GPU?

A CPU has a few powerful cores built to finish one stream of instructions as fast as possible, with big caches and branch prediction. A GPU has thousands of simpler lanes built for throughput: slower at any single task, far faster when the same work repeats across millions of data items.

What is a warp in a GPU?

A group of 32 threads that the GPU schedules together and runs in lockstep, one instruction for all of them. If threads in a warp take different sides of an if, the warp runs both sides one after the other, with part of its lanes idle each time.

Why are GPUs used for AI?

Training and running a neural network is mostly matrix multiplication: the same multiply-and-add repeated billions of times on independent numbers. That is exactly the work a GPU is built for, and modern GPUs add tensor cores that multiply small matrices in a single instruction.

notebook № 027 plates · 25 min to read, longer to play

What a GPU is, and how it works

On the patience of crowds

A CPU makes one stream of work fast. A GPU gives up on that and runs thousands of streams at once, slowly, in lockstep, and wins anyway. This notebook takes that bet apart: warps, latency hiding, memory coalescing, tiling, tensor cores and the grid you launch it all with.

plate I · the bet

A thousand slow brushes

The CPU notebook ended with a machine that spends most of its silicon on cleverness: caches, predictors, out-of-order windows, all to make one instruction stream go faster. A GPU spends the same silicon on arithmetic units instead. Lots of them. A current datacentre GPU has more than sixteen thousand 32-bit floating-point lanes, each one simpler and slower than a CPU core.

That is a bet on the shape of the work. Painting a screen is millions of pixels that don’t care about each other, so a thousand slow brushes beat eight fast ones by a mile. Training a neural network is the same story told with matrices.

The bet has a losing side, though. If every step needs the result of the one before, the crowd stands around while one slow lane does all the work. Flip the race to each pixel needs the last and watch the CPU win.

Drop the GPU to 32 lanes. Where does it stop winning? That crossover is the whole question of “should this run on a GPU?”

fig. 2.1~/paint-racerace, then change the work

A toy model: each CPU core paints a pixel per tick, each GPU lane one every four. The clock speeds up as the race goes on; the tick counts are what matter.

plate II · simt

Thirty-two in lockstep

Sixteen thousand lanes can’t each have their own instruction fetcher and decoder; that would eat the very silicon the GPU was trying to save. So lanes are grouped. On NVIDIA hardware, 32 threads form a warp (AMD calls it a wavefront, of 32 or 64), and a warp issues one instruction at a time for all of its threads. NVIDIA calls this SIMT: single instruction, multiple threads.

You write code for one thread; the hardware runs it 32 at a time. That works beautifully until the threads disagree. When an if splits a warp, it can’t go both ways at once. It runs the then side with the other lanes masked off, then the else side with the first group masked off, and only then do they reconverge. The masked lanes are present, powered and doing nothing.

fig. 2.2~/warpchange the condition, or click a lane’s x

Each column is a lane, each row one instruction issued for the whole warp. SIMT efficiency is the share of lane-slots that did useful work.

The condition decides everything. x % 2 == 0 splits every pair of neighbours and runs both paths. x < 100 sends every lane the same way, and the else never runs at all. Real GPU code is full of small rearrangements whose only job is to make neighbouring threads agree: sort the work first, give each warp a uniform chunk, replace a branch with arithmetic.

Pick x == 7: one lane takes the branch and drags the other 31 through it. Then click lanes to make every x odd and see the cost disappear.

III

plate III · latency hiding

Never idle

A GPU has no large caches to hide the trip to memory, and that trip costs hundreds of cycles, just as it does on a CPU. The latency ladder still applies. So how does it avoid sitting idle?

It keeps many warps resident on each streaming multiprocessor (SM), up to 64 on recent datacentre parts, with every warp’s registers living permanently in a huge register file. Switching from one warp to another costs nothing: no state to save, nothing to restore. Each cycle a scheduler picks any warp that is ready and issues its next instruction. When a warp stalls on a load, it simply isn’t picked until the data arrives.

With enough warps, somebody is always ready, and the long wait is hidden behind other warps’ work. The CPU fights latency with caches and prediction; the GPU drowns it in parallelism. That is why GPU programmers obsess over occupancy, the number of warps an SM can keep in flight.

Set the latency to 600 cycles with 20 sums per load. How many warps until the empty (rose) issue slots disappear?

fig. 2.3~/scheduleradd warps until it fills

A simulation of one warp scheduler issuing one instruction per cycle. Each warp does some maths, then waits for a memory load. Time scrolls left.

plate IV · memory

One trip, not thirty-two

When a warp executes a load, all 32 threads ask for an address in the same instant. The memory system doesn’t serve them one by one. It looks at which 32-byte sectors the addresses fall into and fetches those, whole. If thread 0 wants byte 0, thread 1 byte 4 and so on, the 32 floats sit in four neighbouring sectors, one 128-byte line, and a single trip serves the entire warp. That is a coalesced access.

Spread the threads out with a stride and every thread lands in its own sector. The memory system still moves 32 bytes per sector, so most of what it carries is thrown away. With a stride of 8 floats, you pay for eight times the bandwidth you use.

This is why data layout on a GPU matters more than almost anything else. An array of structs, where each thread reads one field, is a strided access in disguise. A struct of arrays turns it back into a coalesced one.

Keep the stride at 1 and slide the offset. Misalignment costs one extra sector, a small tax. Now try a stride of 2.

fig. 2.4~/coalescingchange the stride

The 128 useful bytes, against what the memory system actually moves. Sector and line sizes follow current NVIDIA GPUs; the principle holds everywhere.

plate V · reuse

Bring it in once

Matrix multiplication is the workload GPUs now live for, and it has a lovely property: every value is needed many times. Each element of A gets multiplied by a whole row’s worth of B. The naive kernel, one thread per output element reading straight from global memory, ignores that and fetches the same numbers over and over. It is limited by memory bandwidth long before the arithmetic lanes get warm.

The fix is shared memory: a small scratchpad on each SM, shared by the threads of one block and an order of magnitude quicker to reach than global memory. A block loads a tile of A and a tile of B into it, every thread computes from the fast copy, and then the block slides to the next pair of tiles. Each value fetched from global memory is now used as many times as the tile is wide.

fig. 2.5~/tiled-matmulchange the tile, click C to move the block

A 16 × 16 multiply. Global loads are counted for the whole product: 2 × 16³ ÷ tile width.

Bigger tiles mean more reuse, but shared memory is small (a couple of hundred kilobytes per SM at most) and every block resident on the SM needs its own share. Take too much and fewer blocks fit, occupancy drops, and plate III comes back to bite. Every serious GPU kernel is a negotiation between those two plates. Libraries like cuBLAS and CUTLASS are, at heart, very carefully negotiated tile sizes.

plate VI · specialisation

Tensor cores

Once matrix multiplication became the main event, the obvious next step was hardware that does nothing else. A tensor core takes small matrices as its operands: in its first version (NVIDIA Volta, 2017), it multiplied two 4 × 4 matrices and added a third, 64 multiply-adds, every cycle. An ordinary lane does one.

The race in fig. 2.6 uses Volta’s proportions for one SM: 64 ordinary FP32 lanes against 8 tensor cores. Same matrix, same answer, roughly eight times sooner. Every generation since has made tensor cores wider and added smaller number formats, so the gap has only grown.

The catch is precision and shape. Tensor cores take their inputs in low-precision formats like FP16, BF16 or smaller, usually accumulating in FP32, and they only help if your problem is a matrix multiply. Fortunately for GPU makers, nearly everything in deep learning is.

Make the matrix 64 × 64. The work grows with the cube of the size: 262,144 multiply-adds from a grid that only looks four times bigger than 32 × 32.

fig. 2.6~/tensor-racerace

Ideal throughput, one SM, Volta-era proportions. Tensor cores finish whole 4 × 4 tiles of the output at a time. Real kernels lose some of this to data movement; see plate V.

VII

plate VII · the launch

Grid, block, warp, thread

So how do you actually hand a GPU some work? You write a kernel, the code for one thread, and launch it over a grid of blocks. Each block is a group of threads that run on the same SM, can share its shared memory and wait for each other at barriers. The hardware slices every block into warps of 32. Blocks are independent of each other, so the GPU can spread them across however many SMs it has, in any order.

Every thread finds its piece of the work from its coordinates. The line at the top of nearly every CUDA kernel ever written is i = blockIdx.x * blockDim.x + threadIdx.x. Hover over a thread to see its numbers.

That independence of blocks is the quiet genius of the model. The same program runs on a laptop GPU with a handful of SMs and a datacentre part with more than a hundred; the big one just takes more blocks at a time. It is also why I think the GPU deserves to be a first-class target in a programming language, which is what I’m building Miri to be.

One step further down this road there is a chip that drops the threads altogether and keeps only the matrix multiply. That is the TPU notebook.

fig. 2.7~/launchhover a thread

Blocks coloured by the SM that runs them. Here they are dealt round-robin; the real hardware hands out blocks dynamically as SMs free up.

Questions people ask

What is the difference between a CPU and a GPU?: A CPU has a few powerful cores built to finish one stream of instructions as fast as possible, with big caches and branch prediction. A GPU has thousands of simpler lanes built for throughput: slower at any single task, far faster when the same work repeats across millions of data items.
What is a warp in a GPU?: A group of 32 threads that the GPU schedules together and runs in lockstep, one instruction for all of them. If threads in a warp take different sides of an if, the warp runs both sides one after the other, with part of its lanes idle each time.
Why are GPUs used for AI?: Training and running a neural network is mostly matrix multiplication: the same multiply-and-add repeated billions of times on independent numbers. That is exactly the work a GPU is built for, and modern GPUs add tensor cores that multiply small matrices in a single instruction.

Related material

elsewhere in the notebooks

№ 01 What a CPU is, and how it works On the clock that never stops № 03 What a TPU is, and how it works On a heart that beats in numbers № 06 How a large language model works On guessing the next word