Why GPUs Are Fast — The Execution Model and the Limits of Parallelism

CPUs and GPUs do not mean the same thing by fast. Where the transistor budget goes, how SIMT bundles 32 threads into a warp, why branch divergence costs you, occupancy and register pressure — and finally Amdahl's law as a way to bound the payoff before you start, plus the profiler counters that tell you when the CPU is the bottleneck.

An analogy: one master, or a hundred ordinary cooks

There are two ways to build the same kitchen. In the first, you hire one master who has the entire workflow memorised, anticipates what comes next, and almost never stands idle. In the second, you line up a hundred ordinary cooks. Each is slower than the master, and at any moment several of them are standing in front of an oven doing nothing. But with a hundred of them, somebody is always working while somebody else waits. No waiting has been eliminated, and yet the second kitchen sends out more plates.

That is the difference between a CPU and a GPU. A CPU is built to finish one task in as little time as possible — latency optimisation. A GPU is built to maximise work completed per unit time — throughput optimisation. They do not define "fast" the same way, so measuring one with the other's yardstick guarantees a misreading.

Where the transistor budget goes

A chip has a finite transistor budget, and design is the decision of where to spend it.

A CPU spends most of it keeping a single instruction stream from stalling: branch predictors, speculative execution, reorder buffers, large caches. None of those perform arithmetic. Every one of them is an investment in "do not let this one thread wait".

The GPU took the opposite bet. It barely predicts and barely speculates, and hands that area to arithmetic units and register files. The register file on a GPU compute unit (an SM, streaming multiprocessor, in NVIDIA's terms) runs to hundreds of kilobytes, and it is not unusual for it to be larger than the L1 cache on the same chip. That is not a demand for storage — it is what it costs to keep the state of thousands of threads resident, with nothing saved and nothing restored. Context switching becomes essentially free, which makes "when a thread stalls on memory, run a different one meanwhile" a workable strategy rather than a wish.

Recall Little's law from the memory wall: filling the bandwidth requires $N = \lambda L$ accesses in flight at all times. A GPU's tens of thousands of threads are exactly that $N$. Instead of shortening latency with caches, it buries latency under parallelism. That is the core of the design.

A warning up front: parallelism only buys a constant

Parallelism buys you a constant factor. Using $p$ processors caps the gain at $p$, and it does not lower the complexity class by a single step. $O(n^2)$ stays $O(n^2)$ on a thousand machines; make $n$ ten times bigger and the hundredfold cost comes right back.

So the order never changes: lower the order, then fix how data moves (the cost of matrix multiplication), then parallelise. Amdahl's law, later in this piece, tells you exactly what ceiling that last step will hit before you commit to it.

SIMT — it looks like threads, but it moves in bundles

If those thousands of threads truly ran independently, nothing would be saved on control logic: you would need one instruction fetch and decode unit per thread.

So the GPU compromises. It groups 32 consecutive threads into a bundle and issues one instruction to all of them. NVIDIA calls the bundle a warp, AMD a wavefront (32 or 64 wide depending on generation). Fetch and decode happen once per bundle; 32 arithmetic lanes sit side by side and apply that one instruction to 32 different pieces of data.

This is SIMT — Single Instruction, Multiple Threads. You write ordinary per-thread code, and the hardware runs it as 32-wide SIMD. It is a compromise struck between writability and efficiency, and most GPU programming pitfalls leak out of that seam. The most important one is the moment the premise breaks: when the bundle no longer agrees on which instruction to run.

Branch divergence — when the bundle disagrees

Thirty-two lanes of a warp reach if (x > 0); 17 evaluate true and 15 false. The hardware can only issue one instruction, so the two sides cannot run together. It executes them in sequence: first the true path, with the 15 false lanes masked off and their results discarded, then the false path, throwing away the other 17 lanes' worth of work.

The execution time of a warp is the sum over the paths any of its lanes took. Said in words: it is a bus at a junction where 17 passengers want to go right and 15 want to go left. The bus drives right, comes back, then drives left. $T_p$ is the time to traverse path $p$ once, and it does not matter how few lanes asked for that path — once the bus goes there, all 32 sit through the trip. What would have cost one pass now costs $T_{\text{if}} + T_{\text{else}}$, and in both stretches part of the hardware is computing results that get thrown away. The measure of that waste is warp execution efficiency $E$ — the average number of active lanes per instruction divided by 32, so $E = 1$ with no divergence and $E = 0.5$ on an even split.

The worst case is when lanes scatter into distinct branches. Branch 32 ways on the thread index and the hardware runs 32 paths in sequence with one active lane each: efficiency $1/32$, time 32×. You can lose a factor of 32 without touching the algorithm — purely from how the condition is written.

Loops behave the same way. If the trip count varies per thread, the whole bundle is held until the longest thread finishes. Process variable-length sequences one-per-thread and a single long sequence drags the other 31 along with it. This is where the standard fix — sort by length, then bundle — comes from.

One clarification: since Volta, NVIDIA GPUs give each thread its own program counter, so threads can make progress while diverged. That made locks and fine-grained synchronisation safe to write, but only lanes on the same instruction still execute together. The time cost of divergence did not go away.

Occupancy and register pressure

Back to the latency-hiding inventory. The number of warps that can be resident on a compute unit at once is that inventory, and how full it is is called occupancy. Three resources cap it: registers, shared memory, and fixed hardware limits. From the register side, with $F$ registers in the compute unit's file and $R$ registers used per thread,

where $W$ is the number of resident warps and 32 is the warp width. Read in words: the register file is a shelf of fixed size on each compute unit, every thread walks off with $R$ slots, and one full warp therefore claims 32 threads' worth of them. However many whole warps that shelf can supply is how many are allowed to stay — the division rounds down, so a warp that comes up one slot short does not get in at all. As $R$ grows, $W$ falls in steps. Adding one local variable can halve the resident warps and drop performance off a cliff — that is register pressure. Worse is spilling: when registers run out, the overflow goes to something called "local memory", which despite the name lives in DRAM.

But more occupancy is not automatically better. The goal is hiding latency, not counting warps. If a single thread can issue several independent loads at once (instruction-level parallelism), a modest number of warps already keeps the inventory stocked. Cutting registers to chase occupancy and ending up slower is a common own goal: read occupancy as one symptom, never as the cause.

Amdahl's law — get the ceiling before you start

Suppose a fraction $f$ of a program parallelises perfectly and the remaining $1-f$ is irreducibly serial. The speedup at parallelism $p$ is

$S$ is the factor by which total runtime shrinks, $f$ is the fraction of time spent in the parallelisable part (not the fraction of source lines), and $p$ is the degree of parallelism. Put in words, only the shareable part of the job gets divided by the headcount; the rest still has to be done start to finish by one worker, no matter how many you hire. The $(1-f)$ sitting in the denominator is that unshrinkable floor. Raising $p$ only pushes $f/p$ towards zero, so the ceiling is $1/(1-f)$. Five percent serial means an infinite number of processors still buys you 20×.

Put numbers in: $f = 0.95$ and $p = 100$ gives $S \approx 16.8$. You bought 100 machines and threw away 83 of them. There are two uses for this. One is a go/no-go check before you start — measure what fraction of time the target region actually owns. The other is solving it backwards: measure at two degrees of parallelism, recover $f$, and you know how much headroom is left.

Amdahl assumes the problem size is fixed. If more machines let you grow the problem instead (bigger batches, higher resolution), you have moved to Gustafson's framing and the ceiling is far less brutal — which is usually why training scales cleanly with machine count.

How this shows up on the job

The three counters to read first

1. Why warps are stalled (Warp State Statistics). Nsight Compute (ncu --set full) reports, per instruction, how many cycles were lost and to what. A large Stall Long Scoreboard means waiting on global memory; Stall Barrier means waiting at __syncthreads(), i.e. imbalance inside the block; Stall MIO Throttle means shared memory or the special function units are congested. Do not touch code until you know which one it is.
2. Warp execution efficiency (smsp__thread_inst_executed_per_inst_executed). Near 32 means no divergence; 16 means half your lanes are producing garbage. It is the direct evidence when you suspect divergence.
3. Achieved occupancy (sm__warps_active.avg.pct_of_peak_sustained_active) — measured, not theoretical. A gap between the two points at too few blocks, or at the tail effect.

The compiler talks too. nvcc -Xptxas -v prints per-kernel register counts and spill bytes. If spills are non-zero, start there and ignore everything else.

Choosing a block size

• Use a multiple of the warp width and start at 128 or 256. Anything not a multiple of 32 leaves idle lanes in the last warp; too small and per-block overhead (launch, synchronisation) dominates; too large and one block hogs the resources, cutting how many blocks stay resident.
• When in doubt, let cudaOccupancyMaxPotentialBlockSize compute it from the kernel's actual register and shared-memory usage.
• To raise occupancy, add __launch_bounds__(256). It tells the compiler to restrain register allocation — restrain it too far and you spill, so confirm with -Xptxas -v.
• Keep the grid a whole multiple of the compute-unit count. If you launch 1.1× as many blocks as there are units, most units sit idle through the final ten percent (the tail effect).

Noticing that the CPU is the bottleneck

GPU utilisation in nvidia-smi is not efficiency. It counts whether any kernel was resident at that instant, so a hopelessly inefficient kernel still reads 100%. What you want is a timeline. Lay out the kernel bars in Nsight Systems or torch.profiler and hunt for gaps between kernels. Time in the gaps is time spent on the host.

Three causes cover most of it. (1) The data loader cannot keep up — check num_workers, pin_memory, and moving preprocessing onto the GPU. (2) Implicit synchronisation: .item(), .cpu(), print(tensor) and torch.cuda.synchronize() all drain the pipeline where they appear, and a single logging line inside the loop can erase the benefit of asynchronous execution. (3) Kernels that are too small — each launch carries a few microseconds of fixed cost, so thousands of tens-of-microseconds kernels are dominated by launch overhead. Fuse them (torch.compile) or cut the launch count with CUDA Graphs.

Questions you will get in interviews and design reviews

"Code that is 95% parallel, run on 100 machines — what speedup?" 16.8×, ceiling 20×. "What if a warp branches on the parity of the thread index?" Both paths run in sequence: roughly 2× the time, 50% efficiency. "Is 100% occupancy the fastest?" No — enough to hide the latency is enough. In every case what is being tested is less the arithmetic than whether you can name the resource that is actually limiting you.

Summary

• CPUs optimise latency, GPUs optimise throughput; the difference is whether the transistor budget goes to control or to arithmetic
• GPUs do not shorten latency, they bury it under an inventory of threads — Little's law's $N$, supplied by hand
• SIMT issues one instruction to 32 threads. When the bundle diverges, time becomes the sum over paths — up to 32× in the worst case
• Occupancy is capped by registers, shared memory and hardware limits, but it is a symptom rather than a goal
• Amdahl's $S(p) = 1/((1-f) + f/p)$ is a go/no-go check, with a ceiling of $1/(1-f)$

Next we widen the same parallelism across machines. The new constraint that appears when you add machines is neither arithmetic nor bandwidth — it is communication.