Performance work on ML systems comes down to one repeated question: where does the time go, and what is the hardware waiting for? This chapter gives you the three possible answers (the math units, memory, or overhead), the reason GPUs are built the way they are (throughput instead of latency, which is why they need tens of thousands of threads), and a back-of-the-envelope model of an LLM generating text that already explains why serving systems batch requests and quantize weights. It ends with your toolchain installed and a first measurement from your own card.
An ML model in production is a program like any other: something calls it, it runs, it returns. What makes it unusual is the shape of the work. Almost all the time goes into a few operations (matrix multiplications, attention, a handful of elementwise and normalization steps) applied to billions of numbers. Those operations run on hardware whose peak speed is enormous and whose typical speed, in naive code, is a small fraction of that peak. The job is closing the gap.
The work splits into layers. Each asks the same question about a different part of the stack:
| Layer | The question | Typical artifacts | In this course |
|---|---|---|---|
| Kernels | Is this one GPU function using the hardware well? | cuBLAS, FlashAttention, quantized GEMMs, fused MoE kernels | Parts I–III, V |
| Runtime / serving | Is the GPU busy with useful work, request after request? | vLLM, SGLang, TensorRT-LLM, llama.cpp: batching, KV-cache management, scheduling | Part IV |
| Compilers | Can we generate good kernels instead of writing each by hand? | torch.compile / Inductor, Triton, TVM, XLA | Part V |
| Communication | When the model spans several GPUs, is the data movement hidden behind compute? | NCCL, tensor and expert parallelism, disaggregated prefill/decode | Part VI |
Day to day, the work looks less like “write a clever kernel” and more like detective work. A deployment is too slow or too expensive. You profile it, and the culprit is often not a slow kernel at all. It might be the CPU failing to launch work fast enough, a hidden synchronization between CPU and GPU, requests batched badly, or memory running out and forcing small batches. Knowing how kernels work is what lets you read the evidence, but the fix lives wherever the evidence points.
Never optimize before you know what the hardware is waiting for. Every technique in this course (tiling, fusion, quantization, batching, CUDA graphs, overlap) is an answer to one specific kind of waiting. Applied to the wrong kind, it does nothing.
A chip has a fixed budget of area and power. CPUs and GPUs spend it on opposite goals, and nearly everything strange about GPU programming follows from that choice.
A CPU is built to make one thread finish as soon as possible. Most of its area goes to things that are not arithmetic: out-of-order execution logic that finds independent instructions to run early, branch predictors that guess which way an if goes before it is evaluated, and large multi-level caches that turn slow trips to DRAM into fast hits. All of it serves one goal: reduce latency, the time a single operation takes from start to finish.
A GPU gives up on making any single thread fast. It fills the chip with many simple, in-order arithmetic lanes, a very large register file, and wide memory channels, and it accepts that each individual operation takes a long time. The goal is to maximize throughput, the total work finished per second across all threads. The GPU doesn't shorten latency, it hides it: while some threads wait for memory, others use the arithmetic units.
Because latency is expensive and throughput is cheap. Doubling throughput is roughly linear: build twice as many arithmetic units and pay for twice the area and power. Halving the latency of an operation means higher clocks, higher voltage and more speculative machinery, and power grows much faster than the gain. Once single-thread performance stopped scaling in the mid-2000s (clocks hit a power wall), the cheap way to get more work per second was more, simpler units. GPUs took that path all the way.
For comparison, a current 8-core desktop CPU manages very roughly 1–2 TFLOPS of FP32 with its vector units and 80–100 GB/s from dual-channel DDR5. So your card offers something like 20× the arithmetic and 7× the bandwidth. That gap in ratios is already a hint: the GPU’s arithmetic grew much faster than its memory, so keeping the arithmetic fed is the hard part.
The CPU, by a wide margin, often 10× or more. Both clock at a few GHz, but a GPU lane runs in order with no speculation, its caches are small and shared by thousands of threads, and a trip to memory costs hundreds of nanoseconds that nothing hides when only one thread exists. The GPU is only fast when you give it enough independent work to switch between. One thread on a GPU is the worst program you can write for it. This is also why the sequential parts of a program (the Python that decides what to run next, the scheduler of a serving engine) stay on the CPU: the machine is heterogeneous on purpose.
You already know the key tool here from backend work, possibly without its name. If a database round trip takes 5 ms and you want 2,000 queries per second, you need about 10 queries in flight at all times: a connection pool of 10. That is Little’s law:
work in flight = throughput × latency
Apply it to GPU memory. To stream data at the full 672 GB/s when each request takes, say, 500 ns to come back from DRAM (typical GPU memory latency is several hundred nanoseconds; you will measure your card’s in chapter 3), the chip must keep
672 GB/s × 500 ns ≈ 336 KB in flight
at every moment. Spread over 48 SMs, that is about 7 KB of outstanding loads per SM. If each thread loads one 4-byte float and then waits for it, each SM needs about 1,750 threads waiting simultaneously. An SM on your card can hold at most 1,536 threads. Even a completely full GPU, with one float per thread in flight, cannot quite saturate memory.
That calculation is the whole design of the GPU in miniature:
A GPU is an extremely wide async server with a slow backing store. Threads are requests; a warp stalled on memory is a request awaiting I/O; the scheduler picks any ready one each cycle. Throughput comes from concurrency, never from making one request faster. Your intuition about connection pools, queue depth and tail latency transfers almost directly.
At any moment, a piece of GPU work is waiting on one of three things. Horace He’s essay Making Deep Learning Go Brrrr popularized this split, and it is the most useful diagnostic frame in the field:
A good first model of a kernel’s run time is
time ≈ max( FLOPs / peak, bytes / bandwidth ) + overhead
The ratio of the two inside the max decides which one wins. Divide FLOPs by bytes and you get the arithmetic intensity of the work, in FLOPs per byte. Divide the hardware’s peak FLOP/s by its bandwidth and you get the intensity at which the two limits are equal, the ridge point:
| On your RTX 5070 | Peak | Ridge point (peak ÷ 672 GB/s) |
|---|---|---|
| FP32 on CUDA cores | 30.9 TFLOPS | ≈ 46 FLOPs per byte |
| FP16 tensor cores, FP32 accumulate | 61.7 TFLOPS | ≈ 92 FLOPs per byte |
| FP8 tensor cores, FP32 accumulate | 123.5 TFLOPS | ≈ 184 FLOPs per byte |
| FP4 tensor cores | 493.9 TFLOPS | ≈ 735 FLOPs per byte |
Work below the ridge is memory-bound, and work above it is compute-bound. Two examples you will write yourself:
c[i] = a[i] + b[i]: 1 FLOP for every 12 bytes (two 4-byte reads, one 4-byte write). Intensity ≈ 0.08, more than 500× below the ridge. It is hopelessly memory-bound, and its ceiling is 672 / 12 = 56 billion elements per second no matter how clever the code is.The third regime is the one people forget. Launching a kernel costs a few microseconds of CPU and driver work before the GPU sees it. If the kernel itself only runs for 5 µs, you are overhead-bound, and neither faster math nor faster memory helps. Eager-mode PyTorch running a small model is often in exactly this state.
| Regime | Symptom | What helps | What does not help |
|---|---|---|---|
| Compute | Math-unit utilization near peak | Tensor cores, lower precision, fewer FLOPs (better algorithm) | Faster memory |
| Memory | DRAM throughput near 672 GB/s; math units mostly idle | Moving fewer bytes: fusion, reuse through shared memory, quantization, coalesced access | Faster math, lower-precision math alone |
| Overhead | Gaps between kernels in a timeline; GPU idle while the CPU works | Fewer, bigger kernels (fusion), CUDA graphs, moving CPU work off the critical path | Optimizing any kernel |
Not at all. At 0.08 FLOPs per byte the arithmetic units already sit idle more than 99% of the time. The kernel is waiting on memory, and at 600 of 672 GB/s it is already close to the ceiling. The only ways to make it faster are to move fewer bytes (store in FP16 and you halve the traffic) or to avoid the round trip entirely by fusing the add into whichever kernel produced a or will consume c. Fusion is the single most common optimization in ML systems, and this is why.
If 20% of a request’s time is CPU-side overhead that you cannot remove, then even infinitely fast kernels give at most a 1 / 0.2 = 5× speedup. Measure the whole timeline before optimizing any part of it. That is what the profiler in chapter 5 is for.
You don’t need any machine learning yet for this; chapters 6 and 10 fill it in. Three facts are enough:
Now apply section 4. One token of a 3B FP16 model on your card:
Memory takes 92× longer. Decode for a single user is massively memory-bound: the tensor cores sit idle 99% of the time while the weights stream past. The ceiling is 1 / 8.9 ms ≈ 112 tokens per second, and no kernel, however clever, can beat it without reading fewer bytes.
Serve several users at once. If 64 conversations each need their next token, one pass over the weights can serve all 64: the weights are read once, and each weight is used 64 times. The bytes stay the same, and the FLOPs grow 64×. The time per step barely changes, because memory still dominates, and the total throughput goes up almost 64×. The batch at which compute catches up with memory is the ridge point in disguise: 61.7 TFLOPS × 2 bytes / (2 FLOPs × 672 GB/s) ≈ 92 sequences. This is why every serving engine batches, and why “continuous batching” (Part IV) is one of the biggest wins in LLM serving.
The second lever is to read fewer bytes per weight: store them in 4 bits instead of 16. The same model is now 1.5 GB, and the ceiling rises 4×. That is quantization, also Part IV.
Play with the numbers. Try a 70B model, try INT4, try an H100 at batch 1 and batch 256, and watch which bar dominates.
These ceilings ignore the KV cache, activations and overheads, and real engines typically reach 60–85% of them. The shape of the answer is still right, and it is the kind of estimate an interviewer expects you to do on a whiteboard in two minutes. Chapter 13 makes it exact.
The weights are 124M × 4 B ≈ 0.5 GB, so the memory time is about 0.74 ms per token and the ceiling is around 1,350 tokens/s. But a straightforward implementation launches a couple of hundred small kernels per token (several per layer, 12 layers), and each launch costs a few microseconds of fixed overhead. That adds up to roughly as much time as the useful work. Your first engine will be partly overhead-bound, and you will fix it later with fusion and CUDA graphs. Small models are where overhead hides.
Everything in this course runs on your RTX 5070. Its architecture is consumer Blackwell, compute capability 12.0, which the compiler calls sm_120. Two practical consequences:
sm_120. Any CUDA 13.x is fine too.nvidia-smi prints, top right, the newest CUDA version your driver supports. If it is lower than your nvcc --version, update the driver first.You will work inside WSL2, a real Linux kernel running alongside Windows. It sees your GPU through the Windows driver, and it gives you the Linux command line every lab, codebase and job in this field assumes. One rule matters more than every other step: the only NVIDIA driver is the Windows one. Never install a Linux GPU driver inside WSL, because it would overwrite the bridge to the Windows driver.
wsl --install -d Ubuntu-24.04, then check that wsl -l -v shows VERSION 2.nvidia-smi should list your RTX 5070. That tool comes from the Windows driver; you didn’t install it, and you shouldn’t.cuda-toolkit (never cuda or cuda-drivers, which pull in a driver). Then add /usr/local/cuda/bin to your PATH and check nvcc --version. Also sudo apt install build-essential git.~/), not under /mnt/c/. Files on the Windows drive are reached through a slow bridge, and builds there crawl.nsys and ncu come with the toolkit and run inside WSL. Install the Nsight Systems and Nsight Compute GUIs on Windows and open the report files from WSL with them (Windows sees WSL files at \\wsl$\Ubuntu-24.04\home\…).By default WSL2 may use up to half of your PC’s RAM. The labs need a few GB of host memory at most, so that’s fine. If you ever need more, set memory= in %UserProfile%\.wslconfig.
| Tool | What it answers | First used |
|---|---|---|
nvcc | Compiles .cu files. -Xptxas -v prints registers and shared memory per kernel. | Lab 0 |
nvidia-smi | Driver version, memory in use, clocks, temperature, who is using the GPU. | Lab 0 |
Nsight Systems (nsys) | The timeline: what ran when, on CPU and GPU, and where the gaps are. The overhead detector. | Chapter 5 |
Nsight Compute (ncu) | One kernel under a microscope: achieved bandwidth, FLOP/s, stall reasons, memory transactions, roofline position. | Chapter 5 |
By default, NVIDIA drivers only allow administrators to read the GPU’s performance counters, and ncu fails with ERR_NVGPUCTRPERM. Under WSL2 this is a Windows setting, because the Windows driver owns the counters: NVIDIA Control Panel → enable Developer Settings from the Desktop menu → Developer → Manage GPU Performance Counters → allow access to all users. (On a native Linux machine it is instead the kernel-module option NVreg_RestrictProfilingToAdminUsers=0.)
Your card boosts its clock depending on temperature and power, so the same kernel can vary by several percent between runs. The labs guard against this with warm-up runs and medians. Nsight Compute goes further and locks the clocks while it profiles (to the boost clock by default in current versions; --clock-control base locks them lower), so its timings are stable but can differ from what your benchmarks see. Chapter 5 explains why, and what else the profiler changes.
Two small programs. devquery asks the driver what your card is and derives its peaks. vecadd runs the vector add from section 4 at sizes from a thousand elements to a quarter billion and reports the bandwidth it achieves. You don’t need to understand the code yet. Chapter 2 takes the kernel apart line by line. Today is about the toolchain working and about your first real number.
git pull # the labs live in this repo
cd mlsys/labs/lab00-meet-your-gpu
make test # builds, then checks vecadd against a CPU loop at awkward sizes
make bench # prints your card's properties, then the bandwidth sweepIf your toolkit isn’t on the PATH or you want another architecture: make NVCC=/usr/local/cuda/bin/nvcc ARCH=sm_120. The build prints lines like ptxas info : Used 12 registers for each kernel. Ignore them for now; chapter 2 explains why they matter.
This is the whole kernel you are about to run:
__global__ void vecadd(const float* a, const float* b, float* c, size_t n)
{
size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; // which element is mine?
if (i < n) // the last block may overhang
c[i] = a[i] + b[i];
}One thread per element. Each thread works out which element it owns from its position in the launch and does one addition. There is no loop: the loop has become the grid of threads.
Before running make bench, make three predictions. Write them down, since they are the point of the exercise:
Terrible: a few GB/s, maybe less. The kernel moves 12 KB, which at 672 GB/s would take 18 ns, but a kernel launch and its completion cost several microseconds no matter what it does. The measured time is almost all fixed cost, so “bytes / time” is tiny. This is the overhead-bound regime, measured. The time stays roughly flat from N = 1K up to tens of thousands of elements. That flat floor is your card’s minimum cost of doing anything at all.
Quite possibly, yes. The benchmark runs the same kernel on the same arrays many times. 12 MB fits comfortably in the 48 MB L2 cache, so after the first run the data never leaves the chip, and L2 bandwidth is several times DRAM bandwidth. The rows marked (fits in L2) measure the cache, not the memory. This is the classic benchmarking trap: the working set determines which memory you are actually measuring. Many “impossible” results in blog posts are L2 hits.
Roughly 80–92%. Now every byte truly comes from DRAM, and a streaming kernel with perfect access patterns gets close to, but never reaches, the spec: the DRAM has refresh cycles, switches between reading and writing, and has other fixed costs. That number, not 672, is the practical ceiling for every memory-bound kernel you will ever write on this card. Remember it; you will compare everything against it.
Then run it, and paste the PASTE: lines back: one from make test, two from make bench. Your measured numbers will replace the estimates in this chapter and become the reference line on every roofline in chapter 4. Also note the CPU comparison line near the end: the same loop on one CPU core.
Change float to double in the kernel and the host code. Bytes per element double, FLOPs stay the same. Predict what happens to elements per second and to GB/s before you run it. Consumer cards run FP64 arithmetic at 1/64 of the FP32 rate. Does that matter here? (It shouldn’t, and seeing it not matter is the lesson.)
Performance engineers do arithmetic in their heads constantly. These are the numbers for your card, and the ones worth memorizing:
| Quantity | RTX 5070 | Why you care |
|---|---|---|
| DRAM bandwidth | 672 GB/s spec, ~80–90% achievable | Ceiling for every memory-bound kernel |
| FP32 / FP16 tensor / FP8 / FP4 | 30.9 / 61.7 / 123.5 / 493.9 TFLOPS | Ceilings for compute-bound kernels |
| Ridge point | ≈ 46 FLOP/B (FP32), ≈ 92 (FP16 tensor) | Which side of the roofline you are on |
| SMs · max threads per SM | 48 · 1,536 | How much parallelism the GPU can hold at once (≈ 74K threads) |
| L2 cache | 48 MB | Working sets below this are not measuring DRAM |
| VRAM | 12 GB | Largest model + KV cache you can serve locally |
| Kernel launch overhead | a few µs | Anything shorter than ~10 µs is probably overhead-bound |
A CPU spends it on reducing the latency of one thread: out-of-order logic, branch prediction, big caches. A GPU spends it on throughput: many simple in-order lanes, a huge register file holding thousands of threads’ state, and wide memory interfaces. It hides latency by switching between threads instead of reducing it.
To stream at 672 GB/s with a latency of several hundred nanoseconds you need hundreds of kilobytes in flight. 32 threads loading 4 bytes each put 128 bytes in flight, thousands of times too little. Bandwidth is only reached with enough independent outstanding requests, which means many threads (and often several loads per thread).
Arithmetic intensity is the FLOPs performed per byte moved to or from memory by a piece of work. The ridge point is the hardware’s peak FLOP/s divided by its bandwidth: the intensity above which work becomes compute-bound. For FP32 on the RTX 5070 it is 30.9 TFLOPS / 672 GB/s ≈ 46 FLOPs per byte.
Whether it is overhead-bound. A 4 µs kernel is comparable to launch overhead, so the timeline probably shows the GPU idle between launches. Fusing it with neighbouring kernels or capturing the sequence in a CUDA graph is likely worth more than any change inside it.
Each token requires reading every weight once but only does about 2 FLOPs per weight, an intensity of about 1 FLOP/byte in FP16, far below the ridge. The levers are batching (amortize each weight read over many sequences, raising the intensity) and quantization (fewer bytes per weight).
About 92. Bytes per step are P × 2 and FLOPs are 2 × P × B, so they balance when B = peak × 2 / (2 × bandwidth) = 61.7e12 / 672e9 ≈ 92. The parameter count P cancels. (Attention and the KV cache, ignored here, change this somewhat. See chapter 13.)
No. 12 MB fits in the 48 MB L2, and the benchmark reuses the same arrays, so it is measuring L2 bandwidth. To measure DRAM, the working set must be much larger than L2 (or the cache must be flushed between runs).
sm_120 (compute capability 12.0) and CUDA 12.8 or newer. nvidia-smi shows the highest CUDA version the installed driver supports. It must be at least the toolkit version reported by nvcc --version.