Two answers to one question
Suppose you have to compute a billion independent fused multiply–adds, and a fixed budget to spend on silicon: some area on a wafer, some watts, some fabrication process. How should you spend it?
There are two coherent answers, and they pull in opposite directions.
The first answer is make each operation finish as soon as possible. The quantity to minimise is latency: the interval between issuing an instruction and having its result available for the next instruction. This is the right objective when the work arrives as a single chain in which each step needs the previous step's answer — parsing a file, walking a linked list, running an interpreter loop, deciding what to do next. There is only ever one thing to do, so the only way to go faster is to do that one thing faster.
The second answer is make as many operations finish per second as possible, and accept that any individual one may take a long time. The quantity to maximise is throughput. This is the right objective when a million independent operations are sitting in a queue and nobody is waiting on any particular one of them.
A CPU is the first answer built in silicon. A GPU is the second. Almost every architectural difference between them — cache sizes, core counts, clock speeds, the existence of branch predictors, the number of registers, the width of the memory bus — falls out of that single divergence of objective.
A note on the numbers before we start
This article is full of quantities, and they are not all equally trustworthy. Some are printed on a datasheet. Some I did not measure at all and have taken from someone else’s published microbenchmarks. Some I derived here from two other numbers. And some are the output of a deliberately simplified model. Conflating those would be the easiest way to mislead you, so every figure and table below is tagged:
- spec
- Published by the vendor. Bandwidths, peak FLOP rates, capacities, architectural limits. Trust these to the precision given, while remembering that "peak" is a ceiling nobody reaches.
- measured
- From published microbenchmarks, not specifications. Every cache and memory latency in this article is of this kind — vendors do not publish them. The shape of the ladder is robust across studies; individual cycle counts vary with clock, access pattern and measurement method. Treat them as within a factor of about 1.5.
- derived
- Computed here from spec figures. Ridge points, arithmetic intensities, traffic totals, collective timings. The arithmetic is shown wherever it matters so you can check it; if an input is wrong, the output moves with it.
- model
- The output of a simplified model, not a measurement. Every interactive lab is one of these. They reproduce the shape of real behaviour — where the cliffs are, which direction things move — while ignoring prefetchers, caches, overlap and instruction-level parallelism. Each one lists what it ignores.
- schematic
- The drawing is not to scale. Used where a true-to-scale rendering would be unreadable, such as a die floorplan or a six-order-of-magnitude capacity range.
Little's Law, the law that governs both
The two objectives are not independent. They are tied together by a relationship so general that it holds for queues at supermarkets and for load instructions in flight:
Read it as a constraint on the machine designer. Solve for throughput and you get T = C / L. There are exactly two ways to make a machine complete more work per second: shrink the latency of each operation, or increase the number of operations you keep in flight at once. That is the whole design space.
A CPU attacks L. It spends enormous resources making a single instruction stream advance quickly: multi-level caches so that operands are close by, branch predictors so that control decisions do not stall the pipeline, out-of-order execution to squeeze whatever concurrency exists out of a program that was never written to be parallel. That last mechanism is the tell — the CPU has to discover concurrency, at runtime, in hardware, because the programmer did not supply any.
A GPU attacks C. It gives up on making any individual operation fast — a global memory load can take five hundred cycles and nobody apologises — and instead demands that you hand it tens of thousands of independent operations up front. Then it keeps them all in flight, so that the long latencies overlap and the arithmetic units never go hungry.
This is why "is a GPU faster than a CPU?" is not a well-formed question. A GPU running a strictly sequential dependent chain is dramatically slower than a CPU — lower clock, longer latencies, no speculation. The GPU only wins when you can supply the concurrency it needs, and the amount it needs is set by Little's Law, not by preference.
modelSimplification: a real CPU core is superscalar and deeply pipelined, so it also keeps dozens of operations in flight. That machinery is the subject of §4 — here one core is drawn as a single slot to isolate the latency-versus-width argument.
Carry this forward
Throughput is bought with concurrency, and concurrency is something you must supply. Every GPU performance problem in the rest of this article is, at bottom, a failure to keep enough independent work in flight — or a failure to feed that work fast enough from memory.
The transistor budget
Silicon is a zero-sum allocation. Every square millimetre spent on a branch predictor is a square millimetre not spent on a multiplier. So the clearest way to understand what a processor is for is to open it up and ask what fraction of it does arithmetic.
Open a modern CPU core and the answer is uncomfortable: very little of it. The bulk is machinery whose only job is to keep a small number of arithmetic units busy despite a program that keeps branching, keeps missing in cache, and keeps expressing dependencies. Instruction fetch and decode. A branch predictor with tens of kilobytes of history tables. Register renaming. A reorder buffer holding hundreds of in-flight instructions. Load and store queues that track memory ordering. Hardware prefetchers guessing the next cache line. Two or three levels of cache. All of this exists to reduce latency — the CPU's stated objective — and none of it computes anything.
Open a GPU streaming multiprocessor and the proportions invert. There is no branch predictor, no register renaming, no reorder buffer, no speculation. In their place: a very large register file, a wide bank of arithmetic lanes, and a scheduler whose entire job is to pick which of the resident warps has an instruction ready this cycle. The GPU does not try to hide latency by predicting the future. It hides latency by having something else to do.
schematicBlock sizes are schematic and chosen for legibility; they are not measured die areas. The quantities in the final stage are published figures for an NVIDIA H100 SXM and a 64-core server CPU with twelve DDR5-4800 channels and two 256-bit FMA pipes per core.
The numbers at the end of that figure are worth sitting with. The GPU has roughly sixteen times as many fp32 arithmetic lanes, but only about eleven times the fp32 throughput, because it clocks lower. And it has only about seven times the memory bandwidth. Width scales more easily than bandwidth — which is the entire subject of the next section.
Note also what the comparison omits. The H100's tensor cores deliver on the order of 989 TFLOP/s of dense BF16, roughly 150× the CPU's general fp32 rate. Specialised matrix hardware is where the gap becomes enormous, and it is also where the bandwidth problem becomes acute: units that fast will starve on anything but the most carefully staged data.
What "parallelism" actually means
The word is used as if it named one thing. It names at least five, they live at different levels of the stack, and they are bought with completely different hardware. Most confusion about why GPUs are good at deep learning comes from collapsing them together.
So before going any further into either machine, it is worth separating them properly. The question each kind answers is the same — what else can happen at the same time? — but the answers differ in where the independent work comes from.
schematicCell widths are uniform so the shapes are comparable; real instruction latencies vary enormously. The point of the figure is the source of the concurrency in each case, not the timings.
Two of these deserve extra attention now, because they are the ones people conflate.
ILP is found; DLP and SIMT are supplied
Instruction-level parallelism is extracted by the hardware from code you wrote sequentially. You do not ask for it. A CPU core spends an extraordinary amount of silicon looking for it — reordering hundreds of instructions, predicting branches so it can keep looking past them, renaming registers so false dependencies do not block the search. That machinery is the reason a CPU core is large, and §4 is about what it buys.
Data-level and SIMT parallelism are the opposite. The hardware supplies width, and the burden is on you to hand it work that fits: contiguous data, uniform control flow, no dependency between iterations. Hand it something that does not fit and the width sits idle. Nothing in the machine will rescue you, because there is nothing in the machine whose job that is.
The sentence to carry through the rest of the article
Parallelism exists at several levels at once, and the engineering question is never "is this parallel?" but "at which level is the independence, and does the hardware I have exploit that level?" A workload can be gloriously parallel at one level and hopeless at another — a task that is embarrassingly thread-parallel may still be unvectorisable, and code that vectorises beautifully may diverge catastrophically on a GPU.
Inside a CPU core: the pipeline
Everything a CPU does to fight latency starts from one observation: executing an instruction is not a single act. It is a sequence — fetch the bytes, decode them, read the operands, do the arithmetic, touch memory, write the result back. If you build one lump of logic that does all of it, the clock period must be long enough for the slowest path through the whole lump, and only one instruction can be in the machine at a time.
Split it into stages separated by latches and two things happen at once. The clock can run faster, because each stage is shorter. And each stage can hold a different instruction, so five instructions are in flight simultaneously. The latency of any one instruction did not improve — it may have got slightly worse — but throughput went up fivefold. This is the same latency/throughput trade from §1, applied inside a single core, and it is the reason a modern core has fourteen to twenty pipeline stages rather than five.
Pipelining only pays if the pipe stays full, and three things empty it. A structural hazard is two instructions wanting the same unit. A data hazard is an instruction needing a result that is not ready. A control hazard is not knowing which instruction comes next, because the previous one was a branch. Most of the CPU core you saw in Figure 02 is dedicated to making these three problems disappear.
modelInstruction latencies here are illustrative round numbers chosen to make the hazard visible, not measurements of any specific core. Real x86 cores issue up to four to six instructions per cycle from a reorder window of several hundred entries, so the effects shown sequentially here overlap heavily in practice. The ordering of the effects is the claim; the cycle counts are not.
What all of that machinery is really doing
Step back from the mechanisms and they share a single purpose: finding concurrency in a program that does not declare any. The reorder buffer is a hardware search for independent instructions. The branch predictor is a way of extending that search across control-flow boundaries, by guessing which instructions will be needed and starting them early. Register renaming exists because the search keeps hitting false dependencies created by the compiler's reuse of sixteen architectural register names.
This search has hard limits. Empirically, general-purpose code contains only a few instructions' worth of extractable parallelism per cycle — sustained IPC of one to three is typical, against a peak issue width of four to eight. Doubling the reorder buffer buys very little; the window is already larger than the available independence. That diminishing return is precisely why the industry stopped making cores wider and started making more of them.
SMT is a two-way GPU
Simultaneous multithreading — Hyper-Threading, in Intel's branding — makes one physical core hold the architectural state of two threads and interleave them. When thread A stalls on a cache miss, thread B's instructions fill the empty issue slots. That is exactly the GPU's latency-hiding strategy: when this thread cannot proceed, run a different one. The difference is scale and honesty. A CPU keeps two threads resident and pretends it has twice the cores. An SM keeps up to sixty-four warps resident and builds its entire performance model around the fact.
SIMD: the CPU's answer to width
Out-of-order execution extracts concurrency the programmer did not declare. SIMD takes the opposite approach: it lets the programmer declare it, and then executes it almost for free.
The idea is to widen the registers and make one instruction operate on all of it. An x86 ymm register is 256 bits — eight fp32 values — and a single vfmadd multiplies and adds all eight in one go. The instruction fetch, decode, rename, schedule and retire costs are paid once for eight results instead of eight times for one. All that expensive front-end machinery from Figure 02 is amortised across the vector, which is why SIMD is so much more energy-efficient per FLOP than scalar code.
The width has grown in steps: 128 bits with SSE (4 fp32 lanes), 256 with AVX/AVX2 (8 lanes), 512 with AVX-512 (16 lanes). Arm's NEON is 128-bit; SVE is width-agnostic by design. Every step doubles peak throughput without touching the clock.
The catch is that a vector instruction imposes one control flow on all its lanes. If the loop body has an if, you cannot have half the lanes take the branch. You compute both sides and select with a bitmask — so the work is done whether or not it is needed. And if the array length is not a multiple of the vector width, the last iteration is partly wasted, either masked off or handled by a scalar clean-up loop. Both of those problems reappear in exactly the same shape on a GPU, where they are called divergence and tail effects.
- Vector instructions
- —
- Lanes issued
- —
- Wasted lanes
- —
- Lane utilisation
- —
- Speed-up vs scalar
- —
—
modelCounts vector instructions and masked lanes for a single loop. It assumes every element is contiguous and already in L1, so it isolates the width effect alone — a real gather, an unaligned base, or a cache miss would each dominate the numbers shown here.
One more property of SIMD matters for the comparison ahead. The eight lanes of an AVX register are positions inside one register belonging to one thread. They are not independent threads. They share a program counter, they share a stack, and — critically — a vector load expects them to come from consecutive addresses. Loading eight values from eight unrelated addresses requires a gather instruction, which is substantially slower than a contiguous load, because it is really eight separate accesses wearing a trench coat.
A GPU's thirty-two lanes are different animals: each is a thread with its own registers, its own address, and its own view of the program. The hardware then notices when their addresses happen to be contiguous and fuses them into one transaction. Same physical goal, opposite default. That difference is the subject of §10.
The GPU execution model
Suppose you have a million independent elementwise operations. How would you describe them to a machine, and how would that machine schedule them?
You could write a million threads and let a scheduler sort it out, but §3 already told us why that fails: a million independent instruction streams means a million program counters, a million sets of registers, and coordination costs that dwarf the work. You could write one vector instruction a million lanes wide, but no such instruction exists and no register file could feed it.
The answer both hardware and software converged on is a hierarchy of groups, where each level is the largest unit that can afford to share something. Get that idea and the vocabulary — grid, block, warp, thread, SM — stops being terminology to memorise and becomes the only sensible answer to the question.
Start with the hardware, because it is the thing that actually exists. A GPU is organised as a hierarchy of replication, and almost every rule you will meet later is a statement about which level something is shared at.
At the top is the whole device: a few graphics processing clusters, a large shared L2, a stack of HBM around the edges, and the links that connect it to other GPUs and to the host. Inside a cluster are texture processing clusters, and inside those are the streaming multiprocessors — the unit that actually runs your code. An SM is itself divided into four processing blocks, each with its own warp scheduler, its own slice of the register file, and its own arithmetic lanes.
The names are NVIDIA's, and AMD and Intel use different ones — compute unit, Xe-core — but the shape is the same everywhere, because it falls out of the same constraints.
specDrawn for an NVIDIA H100 SXM5: 8 GPCs, up to 9 TPCs each, 2 SMs per TPC — 144 SMs on the full die, 132 enabled on the SXM part.
Two numbers from the last stage deserve emphasis. 270,336 threads resident at once is the concurrency the device is built to hold — Little's Law's C, in hardware. And 33 MB of register file is larger than the last-level cache of many CPUs, which tells you where the GPU decided to spend its memory budget: not on caching things that might be reused, but on holding the state of an enormous number of threads so that none of them ever has to be swapped out.
That last point is the deepest architectural difference between the two machines. On a CPU, a context switch means saving registers to memory and restoring another thread's — hundreds of cycles, so the OS does it every few milliseconds. On a GPU, every resident thread's registers are already in the register file, permanently, for the lifetime of the thread. Switching warps costs nothing. It happens every cycle.
And the software hierarchy that maps onto it
CUDA gives you three levels of software abstraction — thread, block, grid — plus a fourth on Hopper, the cluster. Each maps onto a hardware level from Figure 06, and each level's mapping determines exactly one thing: what can be shared, and what can be synchronised.
The pivotal one is the block. A block is guaranteed to be resident on a single SM for its entire lifetime. That guarantee is what makes shared memory and __syncthreads() possible: threads can cooperate because the hardware promises they are neighbours. The flip side is that the guarantee stops at the block boundary. Two blocks may run on different SMs, at different times, in either order, or one may not start until the other has finished — so there is no way to synchronise them mid-kernel without deadlocking.
schematicEight SMs and 24 blocks, shrunk from the real 132 for legibility. The dispatch order shown is one arbitrary order among many, which is exactly the point — the hardware makes no ordering guarantee, so no particular order drawn here is more correct than another. Cluster sizes up to 8 blocks are portable; 16 is H100-specific.
Written out, the correspondence is the whole mental model on one page:
| Software | Hardware | Shares | Synchronises with |
|---|---|---|---|
| thread | lane | its own registers | — |
| warp (32 threads) | processing block | a program counter, in practice | __syncwarp(), shuffles, votes |
| block | one SM | shared memory, L1 | __syncthreads() |
| cluster Hopper+ | one GPC | distributed shared memory | cluster.sync() |
| grid | the whole GPU | L2, global memory | end of kernel |
| multiple grids | multiple GPUs | nothing, without an explicit copy | NCCL collectives, stream events |
Read the "shares" column downward and you are reading a memory hierarchy. That is not a coincidence — it is the point. The execution hierarchy exists to name the levels of the memory hierarchy at which cooperation is affordable. Which is the subject of §9.
SIMT: one instruction, thirty-two threads
You write CUDA as if every thread were an independent program. The hardware does not run it that way. Threads are bundled into groups of 32 called warps, and a warp is the smallest thing a scheduler can schedule. One instruction is fetched, decoded once, and issued to 32 lanes at once.
This is where the amortisation from §2 actually happens. The fetch, the decode, the dependency check, the scheduling decision — all of that control machinery runs once and pays for 32 results. That is the whole trick, and it is why a GPU can afford to have 16,896 arithmetic lanes but only 528 schedulers.
It is close to SIMD, but not identical, and the difference matters. In SIMD you write one instruction operating on a vector, and the lane index is implicit in the data layout. In SIMT you write scalar code, and each lane has its own registers, its own addresses, and — since Volta — its own program counter. Lanes can do genuinely different things. They just cannot do them at the same time.
if are executed whenever any lane needs either side, and the cost of the branch is the sum of the two bodies, not the maximum.modelThe instruction timeline is schematic: bodies are drawn as two slots each so the serialisation is visible, not because a branch body takes two cycles. Since Volta each thread has its own program counter, so the hardware may interleave the two sides rather than running one strictly after the other — the total work, and therefore the cost, is the same either way.
The practical consequence is a rule you can apply while writing code, not just while profiling it: divergence costs are determined by disagreement inside a warp, and warps are aligned to consecutive thread IDs. A branch on tid / 32 is free. A branch on tid % 2 halves your throughput. A loop whose trip count varies per thread costs the maximum trip count in the warp, for every warp.
The lab below makes that concrete. Pick how many distinct paths the code has, and how those paths are distributed across threads.
- Cycles
- —
- If nothing diverged
- —
- Slowdown
- —
- Lane utilisation
- —
—
modelCounts serialised passes only. It ignores instruction-level parallelism, memory stalls, and the fact that other warps can hide some of this — all of which soften the effect on real hardware without changing its direction.
Why warps are 32 and not 8 or 128
Wider warps amortise control logic over more lanes, which is the whole point. But they also make disagreement more likely — a 128-wide warp diverges whenever any of 128 threads disagrees — and they coarsen the granularity of everything: memory requests, scheduling, occupancy. Thirty-two is a compromise that has survived fifteen years and is unlikely to be an accident. AMD, notably, uses 64 (a "wavefront") on CDNA and offers both 32 and 64 on RDNA, which tells you the optimum is shallow.
Latency hiding, and the occupancy trap
A CPU core hides latency by finding other work inside one instruction stream: reorder buffers, speculation, register renaming. All of that machinery exists to extract a few independent instructions from code that was written sequentially.
The GPU does not bother. It has thousands of instruction streams already, so when a warp stalls, the scheduler simply picks a different warp. There is no prediction, no rollback, no reordering — just replacement. This is the payoff for that 33 MB register file: because every resident warp's state is already in registers, switching between them costs nothing at all.
Which turns performance into an inventory question. How many warps do you need resident to keep the pipes full? Little's Law answers it, exactly as it did in §1.
modelA single-scheduler model with a fixed 4-cycle issue burst and a 28-cycle memory wait, chosen so that eight warps exactly cover the gap. Real latencies are far longer and far more variable, several schedulers share an SM, and a warp can be blocked for reasons other than memory. The quantity being illustrated is the ratio latency ÷ issue interval, which is what sets the warp requirement; the absolute numbers are chosen for legibility.
That last stage is the important one, and it is the reason for the word trap in this section's title.
Occupancy — resident warps divided by the 64 slots an SM has — is the most quoted GPU metric and the most misread. It is not a performance target. It is one of two ways to buy the concurrency that Little's Law demands, and the other way is to give each thread more independent work. Volkov's demonstration that a well-unrolled kernel can reach peak bandwidth at 25% occupancy, while a naive one stalls at 100%, is fifteen years old and still surprises people every week.
What occupancy is good for is diagnosis. If a kernel is latency-bound and has neither warps nor instruction-level parallelism, occupancy tells you which resource is standing in the way. Usually it is one of three: registers, shared memory, or a badly chosen block size.
- Occupancy
- —
- Blocks per SM
- —
- Resident warps
- —
- Binding limit
- —
—
__launch_bounds__ exists to prevent.modelModels compute capability 9.0: 64 warp slots, 32 blocks and 65,536 registers per SM, 228 KB of shared memory per SM with a 227 KB per-block cap, registers allocated at 256-per-warp granularity and shared memory at 128-byte granularity. It does not model the L1/shared carveout, which can lower the shared memory budget further.
The rule that actually holds
Do not tune for occupancy. Tune for enough concurrency to cover the latency you have, and measure. Occupancy and instruction-level parallelism are substitutes, and the register pressure that lowers the first is very often what bought you the second.
The GPU memory hierarchy
The GPU's ladder has the same shape as the CPU's — small and fast at the top, large and slow at the bottom — but the rungs are in strange places, and one of them is not a cache at all.
Start at the top and the first surprise is the size. The register file is 33 MB, larger than the L3 of many server CPUs, and it is at the fastest level. That inversion is deliberate: registers are where the state of 270,336 threads lives, and the whole latency-hiding scheme of §8 depends on none of it ever being spilled.
The second surprise is the rung labelled shared memory. It is not a cache. It is a scratchpad — the same SRAM as L1, but with software-controlled placement and no tags, no eviction policy, and no automatic anything. You decide what goes in it and when. The GPU took the transistors a CPU spends on guessing what you will need, and handed you the keys instead.
measuredCapacities and bandwidths are H100 SXM5 published figures. Latencies are approximate, taken from the shape reported by published microbenchmark studies rather than from any vendor specification, and vary with access pattern and clock.
Two rows of that last comparison deserve to be stared at.
GPU memory latency is worse than a CPU's, not better. Roughly 350 nanoseconds against roughly 80. HBM is optimised for bandwidth, and bandwidth is bought partly with latency. The GPU is not fast at getting one thing; it is fast at getting a hundred thousand things at once.
Placement is your job. On a CPU, moving data between L3 and L1 is something you influence — through access order, prefetch hints, blocking — but never command. On a GPU, moving data from HBM into shared memory is an instruction you write. That is more work and vastly more control, and it is the reason a hand-tuned GPU kernel can beat a compiler by an order of magnitude while a hand-tuned CPU kernel usually cannot.
That is as far as we need to go here. This section exists so that the next three — coalescing, banks and tiling — have somewhere to stand: each is a statement about how to use one rung of this ladder well. What the ladder costs you, and how to predict which rung is limiting a given kernel, is a whole model of its own, and it is what the next article in this series is built around.
Coalescing: the single most important access pattern
When a warp executes a load, it does not issue one request. It issues 32 addresses at once, and the memory system's job is to satisfy all of them. It does this by working out which sectors — aligned 32-byte chunks — those addresses fall into, and fetching each distinct sector once.
That is the entire mechanism, and everything follows from it. If the 32 addresses are consecutive 4-byte words, they span 128 bytes: four sectors, four transactions, every byte fetched is a byte used. If the same 32 threads read with a stride of 32 words, they touch 32 different sectors: 1,024 bytes moved to deliver 128 bytes of useful data. Same instruction, same number of threads, same total useful work — eight times the memory traffic.
Notice what this is not. It is not a cache effect, and it is not about reuse over time. It happens within a single instruction, before any cache has an opinion. A perfectly coalesced kernel and a perfectly scattered one differ in the number of transactions the memory controller has to perform, and on a machine whose bottleneck is bandwidth, transactions are the currency.
- Sectors fetched
- —
- Bytes moved
- —
- Bytes used
- —
- Bandwidth efficiency
- —
—
modelCounts distinct 32-byte sectors per warp-level request, which is what the L1 and L2 actually transact in. It assumes a cold cache; in practice L2 may already hold a sector another warp fetched, which is precisely why access patterns that share sectors between warps win twice.
The rule, and why it is not the same as SIMD
The rule to carry is short: consecutive threads should touch consecutive addresses. Not consecutive per thread — consecutive across the warp. A loop where thread t processes elements t*N through t*N+N-1 looks tidy and is catastrophic. The version where thread t processes elements t, t+32, t+64 looks worse and runs many times faster.
This is where the SIMT/SIMD distinction from §5 becomes practical rather than pedantic. A CPU vector unit loading 16 floats issues one address and gets 64 contiguous bytes; a scattered access needs an explicit gather instruction, and gathers are slow and obviously so. A GPU warp always issues 32 independent addresses. Coalescing is not a special path — it is what happens when those 32 addresses happen to be adjacent. The hardware gives you gather semantics for free and charges you for the bandwidth afterwards.
Where this bites in practice
Almost every real GPU performance bug is one of three things: an array-of-structures layout that spreads each thread's fields across sectors; a transpose or reduction whose natural indexing walks columns; or an index vector that turns a contiguous read into a gather. The first two are fixed by changing layout or staging through shared memory. The third is usually fixed by sorting — which is why high-performance sparse and mixture-of-experts kernels spend real time permuting data before they touch it.
Shared memory has thirty-two doors
Shared memory is fast, but it is not magic, and it has one failure mode you will meet the first time you write a transpose.
The SRAM is divided into 32 banks, each four bytes wide, with addresses interleaved across them: word 0 in bank 0, word 1 in bank 1, word 32 back in bank 0. Each bank can serve one address per cycle. So a warp's 32 accesses complete in a single cycle if — and only if — they land in 32 different banks.
If several threads want different addresses in the same bank, the hardware serialises them: an n-way conflict takes n cycles. There is one important exception. If several threads want the same address, the value is broadcast to all of them for free. Conflicts come from different words in the same bank, never from the same word.
Which sets up the classic disaster. Store a 32×32 tile of floats in shared memory and have each thread read one element of a column. Thread t reads element t*32 + c, whose bank is (t*32 + c) mod 32 = c — the same bank, for every thread, with 32 different addresses. One instruction, 32 cycles.
- Conflict degree
- —
- Cycles for this access
- —
- Banks used
- —
- Shared memory cost
- —
—
modelModels 32 four-byte banks with same-address broadcast, which is the behaviour from Kepler onward. Eight-byte and sixteen-byte accesses are handled in multiple phases and are not modelled here.
The padding trick is worth understanding rather than memorising, because the reasoning generalises. Bank index is address mod 32. Any stride sharing a factor of two with 32 folds threads onto a subset of banks; the more factors it shares, the worse the fold. A row width of 32 shares all five, so every thread lands in one bank. A width of 33 is coprime with 32, so consecutive threads walk through all 32 banks in turn.
The same reasoning explains the strided cases in the lab: stride 2 gives a 2-way conflict, stride 4 a 4-way, stride 16 a 16-way, and any odd stride none at all. It is the same modular arithmetic that produces conflict misses in a CPU cache, whose sets are chosen by address bits in exactly the same way — a hardware structure indexed by a slice of the address, meeting an access pattern aligned to a power of two.
Two power-of-two traps, one cause
Cache sets are indexed by middle address bits; shared memory banks by low ones. In both cases a power-of-two stride collapses a wide structure onto a narrow one. In both cases the fix is to break the alignment — pad the row, offset the array, choose a leading dimension that is not a round number. Round numbers are convenient for humans and pathological for hardware.
Tiling: manufacturing reuse
Everything so far has been about doing arithmetic faster. This section is about doing less fetching — and it is worth noticing that the FLOPs are fixed by the mathematics and the data is as large as it is. So where can an improvement possibly come from?
From reuse. How much traffic a computation generates is not a property of the algorithm — it is a property of the implementation, and specifically of how many times each element fetched from slow memory is used before it is thrown away. Matrix multiplication is the canonical demonstration, because its naive form re-reads everything constantly and its tiled form barely re-reads at all, and the difference is nothing but bookkeeping.
derivedn = 16 and T = 4 are chosen so the grids are readable; the reuse argument is independent of both. The reuse factor in the later stages is simply the tile dimension T. The three-level reuse factors are representative of production GEMM kernels rather than measurements of a specific one.
The quantity that matters is the reuse factor: how many times each element fetched from slow memory is used before it is discarded. In the naive loop it is one. With a T × T tile it is T, because every element loaded participates in T of the products the tile supports.
That is a striking result on its own. Reuse grows with the tile dimension and does not depend on the size of the problem, which means the right tile size is a property of the machine rather than of the matrix. It is why tile sizes in fast GEMM kernels cluster around a hundred-odd elements rather than scaling with the input.
Bandwidth is not the only thing choosing that number, though, and it would be tidy but wrong to suggest otherwise. Tensor core instructions accept operands only in fixed tile shapes, so the tile must be a multiple of those. Shared memory capacity caps it from above, and occupancy — §8 — pulls it down further, since a bigger tile means fewer resident blocks. The tile must also divide the problem, or the ragged edge costs a separate code path.
How large a tile you actually need is a question we cannot answer yet, because answering it requires putting a number on the machine's appetite for reuse. That number arrives in §14, and the next article in this series is largely about its consequences.
The general principle, stated once
Tiling is not a matrix multiplication trick. It is the general answer to a bandwidth-bound problem: restructure the computation so that data brought into a fast level of the hierarchy is fully exploited before it leaves. Convolution tiles over spatial windows, FFT over butterflies, attention over key–value blocks, graph kernels over vertex neighbourhoods. Same idea, different index arithmetic — and the next section is the case where it changed a whole field.
Tensor cores, and the pipeline that feeds them
By now the amortisation argument should feel familiar. A CPU amortises instruction fetch over one operation. SIMD amortises it over a vector. SIMT amortises it over 32 lanes. Each step spends less control per result.
Tensor cores apply the identical argument one level further down, but to a different scarce resource: register file bandwidth. An ordinary fp32 FMA reads two operands and writes one result — three register accesses for two FLOPs. Once an SM has 128 such lanes, feeding them becomes as hard as building them. A matrix instruction breaks that ratio, because a small matrix multiply reuses each operand many times inside the unit, where no register file is involved at all.
specPeak rates are H100 SXM5 dense published figures; vendor marketing often quotes the structured-sparsity rate, which is double and requires a 2:4 sparsity pattern. The operand counts for the matrix instruction are derived from a 16×16×16 tile. The pipeline timelines are schematic: real load and math phases do not divide into equal cycle blocks, and the achievable overlap depends on shared memory capacity and instruction scheduling.
Two consequences of tensor cores are worth stating plainly, because they explain a great deal of what modern kernels look like.
They widened a gap that was already there. Peak fp16 arithmetic went up roughly 15× relative to the fp32 lanes on the same die, and memory bandwidth did not go up 15× to match. Every generation of matrix hardware has made the arithmetic cheaper without making the data cheaper to fetch, and §14 is about what that does to the operations you actually run.
They made data movement the hard part of the kernel. When the arithmetic instruction consumes a 16×16 tile, delivering that tile in the right layout at the right cycle becomes most of the engineering. This is why cp.async, the TMA, and Hopper's warp specialisation exist, and why writing a competitive GEMM is now mostly a scheduling exercise rather than an arithmetic one.
Two details that decide whether you can use them
The inputs are narrow; the accumulator is not. A tensor core multiplying fp16 operands does not accumulate in fp16 — it accumulates in fp32. This is not a detail, it is the reason mixed precision works at all. A dot product of length 4,096 in pure fp16 would lose catastrophically to rounding, because each partial sum is rounded to eleven bits of mantissa. Keeping the running sum in fp32 while the multiplies happen in fp16 gives you the bandwidth and throughput of the narrow format with error behaviour close to the wide one. When you read that a model "trains in fp16", the accumulation is almost always fp32, and bf16 exists precisely to make the input side of that trade safer by keeping fp32's exponent range at the cost of mantissa bits.
The shapes are not negotiable. A matrix instruction consumes a fixed tile — 16×8×16 for the Ampere-generation mma, larger warp-group shapes on Hopper — and the operands must be distributed across the warp's registers in exactly the layout the hardware expects. You do not get to choose. Two consequences follow: a tile dimension that is not a multiple of the instruction shape wastes part of every instruction, and getting data from shared memory into the right registers needs specific swizzled layouts and dedicated instructions such as ldmatrix. This is why a dimension of 4,096 runs beautifully and 4,090 does not, and why libraries pad aggressively rather than handle the ragged edge.
A note on "peak"
Published tensor-core figures often quote the structured-sparsity rate, which is double the dense rate and requires a 2:4 sparsity pattern the model must actually have. The dense numbers in Figure 16 are the ones to reason with. And peak assumes every cycle issues a matrix instruction with operands already in place — a well-tuned GEMM reaching 70–80% of dense peak is doing very well.
When does more compute stop helping?
We can now describe, from first principles, how a GPU exposes an enormous amount of arithmetic parallelism. Sixteen thousand lanes, grouped into warps, grouped into blocks, resident by the hundreds of thousands, switched between at zero cost, fed by tensor cores that consume a whole tile per instruction.
And yet. A machine capable of nearly a thousand teraFLOP/s spends most of its life waiting.
Here is the observation that should be nagging by now. In §9 the memory ladder ended at HBM, and HBM had a number attached to it: 3.35 TB/s. In §13 the tensor cores had a number too: 989 TFLOP/s in fp16. Put them beside each other and divide.
To keep the arithmetic units busy, every single byte fetched from memory must feed roughly 295 floating-point operations. Not 3. Not 30. Two hundred and ninety-five.
Now think about what a neural network actually does most of the time. An elementwise addition reads two numbers and writes one, performing a single operation: well under one FLOP per byte. A GELU, a LayerNorm, a softmax, a dropout mask — all of them are in the same range. Even the attention score computation, before any of the tricks, is nowhere near 295.
Every one of those operations runs at a percent or two of what the machine can do, and no amount of the parallelism we spent this article building will change it. The lanes are not the constraint. They were never the constraint.
derivedIntensities are computed from each operation’s minimum traffic in fp16 — the bytes it must read and write if nothing is cached or fused — against its FLOP count. Real implementations can do worse. The 295 line is 989 TFLOP/s ÷ 3.35 TB/s, both H100 SXM5 published figures. This is a first look at a model developed properly in the next article.
This is the wall the rest of the series is about, and it reframes almost everything we just built.
Coalescing in §10 was not a micro-optimisation; it was a way of not wasting the scarce resource. Bank conflicts in §11 were the same thing one level up. Tiling in §12 looked like a way to organise a loop, and it is really a way of manufacturing reuse so that a byte fetched once does more work before it is discarded. Even the tensor cores of §13, read carefully, are an amortisation of operand traffic rather than a faster multiplier.
Every one of those was a data-movement technique wearing the costume of a parallelism technique. Once you have the right model, they stop looking like a list of tricks and start looking like one idea applied at different levels.
Next in this series
Blog 02 — Why Fast GPUs Still Wait for Memory. Arithmetic intensity defined properly; the roofline as a way to predict which of the two ceilings you are under; the CPU cache hierarchy, set conflicts and the TLB; why the compute-to-bandwidth gap has widened with every hardware generation; and FlashAttention as the case study where taking all of it seriously changed what models could be trained. The question it answers: when is your program limited by arithmetic, and when by the distance to its data?
What to carry with you
If you remember nothing else from this article, remember that every question about performance on either machine reduces to the same two: how far does the data have to travel, and is there enough other work to do while it travels? Everything below is a special case.
The five ideas, in the order they were introduced
| Idea | On a CPU | On a GPU | Where it goes next |
|---|---|---|---|
| Amortise control spend less overhead per result |
pipelining, superscalar issue, SIMD | one instruction per 32 lanes; tensor cores per tile | workload-level parallelism, across devices |
| Cover latency you cannot remove it, so hide it |
out-of-order execution, speculation, prefetch | switch warps; more independent work per thread | overlapping communication with the backward pass |
| Keep data close reuse before you release |
cache blocking, loop interchange | shared memory tiling, register tiling, fusion | the memory model, and then the network |
| Move whole lines the granularity is not one byte |
64-byte cache lines; stride destroys them | 32-byte sectors; coalesce or waste bandwidth | the same story at every level of the ladder |
| Beware powers of two round numbers alias onto hardware |
cache set conflicts at large strides | shared memory bank conflicts | shapes that do not divide the hardware |
Numbers worth knowing by heart
| Quantity | Value | Why it matters |
|---|---|---|
| fp16 tensor peak | 989 TFLOP/s spec dense; the sparsity figure is double | the number that makes everything else look slow |
| fp32 lane peak | 67 TFLOP/s spec 16,896 lanes, no tensor cores | what the same silicon does without matrix instructions |
| HBM3 bandwidth | 3.35 TB/s spec | divide the line above by this one, and read §14 |
| HBM latency | ~350 ns · ~550 cycles measured microbenchmarked, ±50% | worse than a CPU's; hidden, never removed |
| Warp size | 32 threads spec | the granularity of divergence, coalescing and scheduling |
| Memory sector | 32 bytes spec | the unit the memory system actually transacts in |
| Shared memory banks | 32 × 4 bytes spec | conflict-free only when addresses spread across all of them |
| Threads resident per SM | 2,048 · 64 warps spec | your entire latency-hiding budget |
| Registers per SM | 65,536 · 256 KB spec | why context switching is free, and why it can bind occupancy |
| Shared memory per SM | up to 228 KB spec | the size of the tile you are allowed to build |
| SM shared memory | up to 228 KB spec | the scratchpad you place data in yourself |
A diagnostic order of operations
When a GPU kernel is slower than it should be, this sequence resolves it faster than guessing. It stops one step short of where it should, and the next article supplies the missing first step.
- Check that the work is actually parallel at the level the hardware exploits (§3). A loop with a carried dependency is not made parallel by launching a million threads at it.
- Look for divergence (§7). Ask what fraction of a warp is active in the hot branch. A branch on
tid % 2halves you and looks completely innocent in the source. - Check the access pattern before anything else in memory (§10). A strided or gathered load can cost 8× the traffic of the coalesced version while looking identical in the source.
- Then shared memory conflicts (§11). A column-major access into a 32-wide tile is a 32-way conflict and a one-word padding fix.
- Then ask whether data is being re-read that could have stayed close (§12). Reuse is manufactured by restructuring loops, not by hoping the cache helps.
- Check whether the tensor cores are being used at all (§13), then whether they are being fed — an un-pipelined loop leaves them idle half the time.
- Only then look at occupancy (§8), and treat it as a diagnosis rather than a target. Ask whether there is enough concurrency, from warps or from independent work per thread, to cover the latency you have.
- And before all of it, in the next article: find out which ceiling you are under. Most of the list above is wasted effort if you have not established whether the kernel is limited by arithmetic or by data movement.
The one-sentence version
A CPU is a machine built to make one thing happen as soon as possible, and it spends most of its silicon finding parallelism you did not write down. A GPU is a machine built to make an enormous number of things happen per unit time, and it spends almost none — which is why it is faster, and why it makes you do the work.
References
Grouped by the section they support. Where a claim in this article is a microbenchmark result rather than a published specification, the source is noted as such.
Foundations: latency, throughput, and the memory wall
- Roofline: an insightful visual performance model for multicore architectures — Williams, Waterman & Patterson, 2009. The model developed in the next article of this series, and still the clearest single page on why arithmetic intensity is the right axis.
- Computing's energy problem (and what we can do about it) — Horowitz, ISSCC 2014. Why moving a byte costs far more energy than computing on it — the physical root of the constraint §14 points at.
- What every programmer should know about memory — Drepper, 2007. Long, and worth it. Background for the CPU side of §§4–5, and the reference for the memory material in the next article.
- Software optimization resources — Agner Fog. Measured instruction latencies and throughputs for x86 microarchitectures; the empirical backing for §4.
- Hitting the memory wall: implications of the obvious — Wulf & McKee, 1995. Named the problem thirty years ago and got the trajectory right.
GPU architecture and execution
- GPU Glossary — Modal. The best short reference for the vocabulary in §§6–13, with the hardware and software hierarchies cross-linked.
- CUDA C++ Programming Guide — NVIDIA. Authoritative for the occupancy limits, allocation granularities and bank behaviour modelled in Figures 15, 17 and 18. Appendix on compute capabilities is the one to bookmark.
- Dissecting the NVIDIA Volta GPU architecture via microbenchmarking — Jia et al., 2018. Where the memory latencies quoted in §9 come from. Vendor specifications do not publish these; this is measurement.
- Understanding latency hiding on GPUs — Volkov, 2016. The definitive treatment of §8, including the demonstration that high occupancy is neither necessary nor sufficient for peak throughput.
- How to optimize a CUDA matmul kernel for cuBLAS-like performance — Boehm, 2022. A worked, measured walk through every tiling level in §12, with the speedup attributable to each.
- CUTLASS — NVIDIA. The production answer to §§12–13: templated tiling hierarchies, async pipelines and warp specialisation, readable as source.
Attention, fusion and long context
- FlashAttention: fast and memory-efficient exact attention with IO-awareness — Dao, Fu, Ermon, Rudra & Ré, 2022. The case study the next article builds to, and the paper that made "IO-aware" a design principle rather than an implementation detail.
- FlashAttention-2: faster attention with better parallelism and work partitioning — Dao, 2023. Reworks the loop order and work partitioning; the reason the traffic model in Figure 23 puts the query block in the outer loop.
- Online normalizer calculation for softmax — Milakov & Gimelshein, 2018. The recurrence in Figure 23, published four years before it became load-bearing for the whole field.
- Attention is all you need — Vaswani et al., 2017. For the operation whose quadratic intermediate makes the next article’s case study necessary.
- Triton — OpenAI. Makes the tiling of §12 expressible without writing the index arithmetic by hand.
Multi-GPU and distributed training
- ZeRO: memory optimizations toward training trillion parameter models — Rajbhandari, Rasley, Ruwase & He, 2019. The workload-level parallelism foreshadowed in §3, with the communication volume derived per stage.
- Megatron-LM: training multi-billion parameter language models using model parallelism — Shoeybi et al., 2019. Why tensor parallelism stays inside a node: it communicates activations twice per layer.
- NCCL — NVIDIA. The implementation of Figure 24, including the hierarchical algorithms that split a collective by link speed.
- Ring Attention with blockwise transformers for near-infinite context — Liu, Zaharia & Abbeel, 2023. Extends blockwise attention across devices, making the sequence dimension itself a level of the hierarchy.
- Anatomy of high-performance matrix multiplication — Goto & van de Geijn, 2008. The multi-level blocking argument of §12, worked out for CPUs a decade before it was reinvented for GPUs.
Two further microbenchmark studies of Hopper were consulted for the latency figures in §9 but are not linked here, because their identifiers could not be verified while this article was written offline. Treat all cache and memory latencies in this article as approximate, measured quantities rather than specifications — the shape of the ladder is robust, the individual numbers are not.