A GPU is an enormous calculator fed by a narrow pipe. This chapter shows how the chip works, how to predict a kernel's speed before writing it, and how to write the kernels (in CUDA, Triton, and NumPy simulations) that keep the calculator fed.

Here is a puzzle to start the chapter. An NVIDIA H100 can perform about $10^{15}$ bfloat16 floating-point operations per second. Ask it to add two vectors of a billion float32 numbers each, a billion additions in total. How long does that take?

If you divide a billion additions by the GPU's 67 trillion float32 operations per second, you get 15 microseconds. The real answer is about 3.6 milliseconds, 240 times longer. The additions are not the work. The work is moving the data. The two inputs are 4 GB each and the output is another 4 GB, and the GPU's memory delivers about 3.35 TB per second, so just moving 12 GB takes $12 / 3350 \approx 3.6$ ms. During that time the arithmetic units are idle more than 99 percent of the time.

That one calculation contains most of this chapter. A GPU is two machines bolted together, an enormous calculator and a comparatively narrow pipe that feeds it. Whether your kernel runs at 1 percent or 80 percent of the chip's peak depends almost entirely on how you use the pipe. FlashAttention (§12.3), fused kernels, the batching that makes inference affordable, and the model FLOPs utilization that every pretraining team tracks (§6.7) are all consequences of it.

This section builds the mental model. We look at how the chip is organized (streaming multiprocessors, warps, and the SIMT execution model), walk the memory hierarchy with real numbers for an H100, see how a GPU hides the long wait for memory by juggling thousands of threads, and meet the tensor cores that supply almost all of its arithmetic. The next section turns this into a quantitative tool, the roofline model.

Two ways to spend a billion transistors

A CPU core is built to make one stream of instructions finish as soon as possible. Most of its silicon goes to things that reduce the latency of that one stream. Large caches keep data close, branch predictors guess which way an if will go, and out-of-order execution finds independent instructions to run while an earlier one waits. A modern server CPU has a few dozen such cores.

A GPU makes the opposite bet. It assumes you have millions of independent pieces of work, such as the elements of a tensor, and it only cares how many of them finish per second. So it spends its silicon on arithmetic units and on the storage needed to keep many threads alive at once, and it spends very little on making any single thread fast. When a thread has to wait for memory, the GPU does not try to predict or reorder anything. It switches to another thread that is ready.

In plain words, a CPU is a latency machine and a GPU is a throughput machine. The same program can be 50 times faster on one or the other depending on whether it has enough independent work to fill the GPU.

The chip, from the outside in

An H100 in its SXM form (the version on server boards) is organized as a set of 132 identical streaming multiprocessors, or SMs. The SM is the unit that actually runs your code. Everything you launch is split into chunks of work called thread blocks, and each block runs entirely on one SM. The full GH100 die has 144 SMs, and the shipping product enables 132 of them, which improves manufacturing yield.

Each SM is divided into four partitions (NVIDIA calls them processing blocks). A partition has its own warp scheduler, which can issue one instruction per clock cycle, 32 lanes of float32 arithmetic, one tensor core, a slice of the register file, and units for loads, stores, and special functions such as exp. The four partitions share 256 KB of fast on-chip memory that serves as both the L1 cache and the programmer-managed shared memory. Outside the SMs sit a 50 MB L2 cache shared by the whole chip and 80 GB of HBM3, the stacked DRAM that NVIDIA calls GPU memory.

Look at the last row. One H100 can hold over a quarter of a million threads resident at the same time, each with its own registers. This is the most important number in the table, and the rest of the section explains why a GPU needs it.

Warps and the SIMT model

When you write a CUDA kernel, you write the program for one thread (§11.3 shows the syntax). The hardware never schedules threads one at a time, though. It groups 32 consecutive threads of a block into a warp, and the warp is the unit of execution. Each cycle, a partition's scheduler picks one warp that is ready and issues one instruction for all 32 of its threads, which each apply it to their own registers. NVIDIA calls this single instruction, multiple threads (SIMT).

SIMT looks like SIMD (the vector instructions of a CPU), but the programming model is different. In SIMD you write vector code explicitly. In SIMT you write scalar code, and the hardware discovers the vector parallelism by running 32 copies in lockstep. The catch is what happens when those 32 copies disagree.

Suppose a kernel contains if (x > 0) { A } else { B } and, within one warp, some threads have positive x and others do not. The warp cannot issue A and B at the same time, because it issues one instruction at a time. So it runs A with the "else" threads masked off, then B with the "if" threads masked off. This is warp divergence, and its cost is simple to state. A warp pays for every path that at least one of its threads takes.

In plain words, if any lane in a warp needs a branch, the whole warp waits while that branch runs.

Since the Volta generation, threads in a warp can be scheduled more independently, so diverged threads can interleave and even synchronize with each other. That fixes some deadlocks, but it does not make the two paths run simultaneously, so the cost model above still holds. The cell simulates it. The same data either arrives in random order, so that nearly every warp sees both signs, or sorted, so that each warp sees only one.

With random data almost every warp diverges and pays 30 cycles, twice the 15 it would pay with no divergence. Sorting the same data makes the cost almost ideal, with only the one warp that straddles zero paying for both paths. Real kernels avoid divergence by making branch conditions depend on the warp index rather than the thread index where possible, and by using arithmetic (max, where-style selects) instead of branches for short paths.

import numpy as np

def divergent_cycles(path, costs):
    path, costs = np.asarray(path), np.asarray(costs)
    return int(costs[path].max())          # placeholder: pretends there is only one warp

The memory hierarchy, with real numbers

Every byte a kernel uses lives somewhere in a hierarchy of memories, and each level trades size for speed. The numbers below are for an H100 SXM. Capacities and HBM bandwidth come from NVIDIA's datasheet and tuning guide. The latencies are empirical measurements, taken with pointer-chasing microbenchmarks on an H800 (a Hopper GPU with the same SM design) by Luo et al. (2024), and they are best-case figures for an idle chip.

The last two rows leave the chip. NVLink joins the GPUs of one server, and its bandwidth is quoted as a total over both directions, 600 GB/s for an A100, 900 GB/s for an H100, and 1.8 TB/s for a B200, so half of each in either direction. Even the fastest is several times slower than HBM, which is why distributed training works so hard to overlap communication with compute.

Three things in this table are worth committing to memory.

The bandwidth gap. Shared memory delivers about nine times the bandwidth of HBM. The total comes from 32 banks, each 4 bytes wide, on every SM, so $32 \times 4 = 128$ bytes per cycle per SM, times 132 SMs, times a clock of about 1.83 GHz, gives $128 \times 132 \times 1.83\times10^{9} \approx 3.1\times 10^{13}$ bytes per second. (Banks matter for performance too, and §11.3 counts bank conflicts.) Registers are faster still.

The pyramid is upside down near the top. On a CPU, the register file is tiny and the caches are big. On a GPU, each SM's register file (256 KB) is as large as its L1 and shared memory combined, because it has to hold the private state of up to 2,048 resident threads. Kernels that compute a lot per byte, like matrix multiplication, keep their working set in registers.

Latency is long, and it is measured in cycles. A load from HBM takes about 479 cycles before the data arrives even on an idle chip, and more under load. At about 1.83 GHz that is about 260 nanoseconds, enough time for one SM's tensor cores to perform about two million floating-point operations.

The cell below turns the table into the numbers that drive kernel design.

The first line is the whole problem in two numbers. Each cycle, an SM can perform 4,096 bfloat16 FLOPs on its tensor cores, but it receives only about 14 bytes, seven bfloat16 numbers, from HBM. A kernel that does not reuse each number it loads hundreds of times cannot possibly keep the tensor cores busy. That ratio, 295 FLOPs per byte, is the ridge point of the H100, and §11.2 is built around it.

Hiding latency with many warps

A warp that issues a load from HBM cannot use the result for about 500 cycles. If that warp were the only one on its partition, the partition would sit idle for those cycles. The GPU's answer is to keep many warps resident and to switch between them for free. Each warp's registers stay in the register file the whole time, so a "context switch" is just the scheduler choosing a different warp in the next cycle. While some warps wait for memory, others compute.

How many warps are enough? Suppose each warp alternates between issuing one load, waiting $\ell$ cycles for it, and then issuing $c$ instructions that use the data. The scheduler can issue one instruction per cycle. With $W$ warps, the partition has $W(c + 1)$ instructions to issue in every window of $c + 1 + \ell$ cycles, so the fraction of cycles in which it issues anything is about

In plain words, you need enough warps that their combined work covers one memory latency. The longer the latency and the less work each warp does per load, the more warps you need.

Try it in the widget. With a latency of 480 cycles and only 8 instructions per load, a partition needs about 55 warps to stay busy, far more than the 16 it can hold (64 per SM, 4 partitions). That is why memory-bound kernels never reach the issue rate, and why their speed is set by bandwidth instead.

The same reasoning, applied to bytes instead of instructions, is Little's law from queueing theory. The amount of data in flight must equal the bandwidth times the latency. To sustain 3.35 TB/s with a latency of about 260 ns, the chip must have about $3.35\times10^{12} \times 2.6\times10^{-7} \approx 0.9$ MB of loads outstanding at every moment, about 6.6 KB per SM, or about 52 warp-wide loads of 128 bytes each. Under load the latency grows, so in practice it is more. There are three ways to get there.

• More warps. Raise occupancy, the number of resident warps per SM divided by the maximum (64 on an H100).
• More independent loads per warp. Each thread issues several loads before using any of them (instruction-level parallelism), or uses 16-byte vector loads, which move 512 bytes per warp instruction.
• Asynchronous copies. Since Ampere, cp.async copies global memory into shared memory without tying up registers, and Hopper's Tensor Memory Accelerator (TMA) moves whole tiles with one instruction. A kernel can then load the next tile while computing on the current one.

Occupancy and what limits it

How many warps fit on an SM depends on what each block asks for. An SM has a fixed budget of resident threads (2,048), registers (65,536), shared memory (228 KB on an H100), and blocks (32), and a block is only scheduled when all of its resources fit. Whichever budget runs out first limits occupancy. Registers are allocated per warp in chunks of 256, and on Ampere and Hopper each block also reserves 1 KB of shared memory for the system.

A worked example. A kernel with 256 threads per block (8 warps) that uses 64 registers per thread needs $64 \times 32 = 2{,}048$ registers per warp, so 16,384 per block, and $65{,}536/16{,}384 = 4$ blocks fit, 32 warps, 50 percent occupancy. Reduce register use to 32 per thread and 8 blocks fit, 64 warps, 100 percent. The widget lets you try other kernels.

Registers have a hard ceiling too. A thread can address at most 255 of them, and if a kernel needs more, or if you cap the count to raise occupancy (with __launch_bounds__ or the -maxrregcount compiler flag), the compiler spills the excess to local memory. Despite the name, local memory lives in HBM, cached in L1 and L2, so every spilled value turns a free register access into a memory access. The compiler's -Xptxas -v report and Nsight Compute both show spills, and a kernel that suddenly slows down after a small change often has one.

def occupancy(threads_per_block, regs_per_thread, smem_per_block):
    warps = -(-threads_per_block // 32)
    blocks = 64 // warps                  # placeholder: only the warp limit
    return {"blocks": blocks, "warps": blocks * warps, "occupancy": blocks * warps / 64, "limiter": "warps"}

Tensor cores, where the FLOPs are

The float32 lanes are not where an H100's headline number comes from. Each SM's four tensor cores each compute a small matrix multiply-accumulate, $\mathbf{D} = \mathbf{A}\mathbf{B} + \mathbf{C}$, where the four matrices are tiles a few elements on a side, as a single operation. On an H100 each tensor core completes 512 bfloat16 multiply-adds, 1,024 FLOPs, per cycle. Across the chip that is

against 67 TFLOP/s from all 16,896 float32 lanes. The tensor cores supply about 15 times as much arithmetic as everything else combined, and they only do one thing, multiply matrices. (NVIDIA does not publish the clock behind its tensor-core peak, so we inferred 1.83 GHz from it. The float32 and FP64 peaks correspond instead to the maximum boost clock of 1.98 GHz, $16{,}896 \times 2 \times 1.98\times10^{9} \approx 67$ TFLOP/s, so treat every per-cycle figure in this section as approximate to within about 8 percent. The A100's 312 TFLOP/s checks the same way, $108 \times 4 \times 512 \times 1.41\times10^{9}$, with half the per-SM rate and its 1.41 GHz boost clock.)

The rate depends on the number format, roughly doubling each time the operands halve in size. These are dense peaks from NVIDIA's datasheets. The datasheets print numbers twice as large for most rows, which assume 2:4 structured sparsity (two zeros in every four weights), and you should halve those for ordinary dense work. Number formats themselves are the subject of §8.5. The B200 column is one GPU of an HGX B200 server (eight GPUs at about 1,000 W each). Blackwell parts run at other power limits in other systems, with somewhat different peaks.

Read the last row against the others. From A100 to B200 the 16-bit tensor-core rate grew about sevenfold while bandwidth grew about fourfold, so each generation needs more FLOPs per byte to stay busy, a trend §11.2 makes precise.

Two consequences matter for machine learning. First, anything that is not a matrix multiplication runs on hardware that is 15 times slower, so the non-matmul parts of a model can take a surprising share of the time unless they are cheap in FLOPs and fused (§11.7). Second, special functions are slower still. The H100 has 16 special-function units per SM, which evaluate exp and friends at about 3.9 trillion per second in total, about 250 times less than the matmul rate. Shah et al. (2024) point out that this is why the exponentials in attention's softmax become a real cost in FlashAttention-3 (§12.4).

Why sizes should be multiples of 64

A tensor-core kernel does not compute $\mathbf{C} = \mathbf{A}\mathbf{B}$ element by element. It divides $\mathbf{C}$ into tiles, for example $128\times128$, gives each tile to one thread block, and loads operands in 16-byte vectors. Matrix shapes that do not fit this structure waste work in three ways.

Alignment. If a bfloat16 matrix has a row length that is not a multiple of 8 elements (16 bytes), rows do not start on 16-byte boundaries, and the fastest kernels, which assume aligned vector loads, cannot be used. NVIDIA's performance guide recommends dimensions that are multiples of 8 for 16-bit formats, and larger multiples help further.

Tile quantization. A dimension of 129 needs two tiles of 128, and the second does almost no useful work but costs as much as the first.

Wave quantization. Tiles are spread over the 132 SMs in waves. A matrix product with 132 tiles takes one wave, and one with 133 tiles takes two, the second wave using one SM while 131 sit idle.

This is why Karpathy padded GPT-2's vocabulary from 50,257 to 50,304 tokens (a multiple of 128) when reproducing it, and measured a training step falling from about 96 to 93 milliseconds, a few percent for free. The 47 extra rows of the embedding table are never produced by the tokenizer (§2.5).

The $1536\times1408$ product has exactly 132 tiles and runs in one perfect wave, while $1536\times1536$, which is only 9 percent more work, needs 144 tiles and therefore two waves, wasting almost half the GPU. The 129-row product shows tile and wave quantization at once, with half its tiles nearly empty and half the SMs idle. Libraries soften this with smaller tiles for awkward shapes and with schemes such as split-K and stream-K that divide the long inner dimension, but the effect is real and shows up as a sawtooth when you plot matmul speed against size.

def gemm_waves(M, N, tile_m, tile_n, n_sm):
    tiles = (M // tile_m) * (N // tile_n)          # placeholder: forgets the partial tiles
    waves = tiles // n_sm
    return tiles, waves, 1.0

Host, device, and the cost of launching work

The CPU (the host) and the GPU (the device) have separate memories joined by PCIe, which on an H100 moves about 64 GB/s in each direction, 50 times less than HBM. So tensors should be moved to the GPU once and stay there, and data loading should overlap with compute (§0.9).

The host also drives everything. Each operation in PyTorch's eager mode is a kernel launch, a request the CPU places in a queue (a CUDA stream) for the GPU to run. Launches are asynchronous. The CPU queues work and moves on, and it waits only at a synchronization point such as .item(), .cpu(), or printing a GPU tensor. Each launch costs the CPU a few microseconds, and PyTorch's own dispatch adds more. A kernel that runs for less time than that leaves the GPU idle between kernels, a third regime besides compute-bound and memory-bound that Horace He calls overhead-bound. Fusing kernels and capturing CUDA graphs are the cures (§11.7).

In PyTorch

You can see the two machines on any computer. The block below queries the GPU if one is present, then measures the same contrast on the CPU, a bandwidth-bound elementwise add against a compute-bound matrix product. The CPU's numbers are much smaller than an H100's, but their ratio tells the same story.

Read more

• Kirk, Hwu & El Hajj (2022). Programming Massively Parallel Processors: A Hands-on Approach, 4th edition. The standard textbook. Chapters 4 and 5 cover the execution model, warps, divergence, occupancy, and the memory hierarchy.
• NVIDIA (2022). NVIDIA Hopper Architecture In-Depth. The H100's SM, tensor cores, TMA, thread block clusters, and the A100 comparison table.
• NVIDIA. NVIDIA Hopper Tuning Guide. The authoritative limits per SM (registers, warps, blocks, 228 KB of shared memory, the 50 MB L2).
• NVIDIA. GPU Performance Background User's Guide. A short official explanation of SMs, tensor cores, and why shapes and batch sizes matter.
• Luo et al. (2024). Benchmarking and Dissecting the Nvidia Hopper GPU Architecture. Microbenchmarked latencies and bandwidths for every level of the hierarchy, the source of our latency column.
• Volkov (2010). Better Performance at Lower Occupancy. The classic talk showing that instruction-level parallelism can replace occupancy.
• Lindholm et al. (2008). NVIDIA Tesla: A Unified Graphics and Computing Architecture. The paper that introduced the SIMT design.

Watch next

▶

▶

▶