Chapter 1.0 — The GPU Execution Model¶

In one line: A GPU is 108 small independent processors that go fast only when thousands of threads are in flight, and the most expensive thing in a decode step is often not arithmetic — it is the CPU asking the GPU to do something.


Part	I — The Machine
Chapter	1.0
Time	~90 min
Prereqs	None — this is the entry point
Notation	$P_{peak}$, $B_{mem}$, $L$, $d_h$, $H_{kv}$
Status	draft

Where we are

Every later chapter reasons about bytes, FLOPs, and time, and none of that arithmetic means anything until you know what the hardware is made of and how work reaches it. This chapter defines the six words the rest of the book uses without ceremony — SM, warp, occupancy, coalescing, tensor core, kernel launch. Chapter 1.1 then asks where the bytes live; Chapter 1.2 asks whether you have enough work per byte.

Why this matters¶

You will read a spec sheet claiming an A100 80 GB delivers 312 TFLOP/s and 2.039 TB/s. You will compute a bandwidth floor for a Llama-3.1-8B decode step — 16.06 GB of weights, ≈11 ms at an achievable 1.5 TB/s — then measure 18 ms and have no idea where the other 7 ms went.

It went into the CPU. A decode step for a 32-layer model issues several hundred kernel launches, each costing single-digit microseconds of host-side work, and at batch 1 there is not enough GPU work per kernel to hide them. Add tensor parallelism across four GPUs and it gets worse: you cut per-rank device work 4×, the launch count stays identical, and the extra hardware buys nothing. That failure is invisible from the roofline and from nvidia-smi, which reports 90-something percent "utilisation" while the SMs idle between kernels.

The sentence to remember

A GPU has two throughputs: how fast it computes, and how fast you can ask it to compute. Optimising the first while the second binds is the most common wasted quarter in inference engineering.

The mental model¶

Stop picturing a processor. Picture a factory floor of 108 identical independent workshops. That is what an A100 is: 108 streaming multiprocessors (SMs), each with its own registers, its own 192 KB of combined L1/shared memory, its own arithmetic units, and its own scheduler. They share exactly two things — the L2 cache and the HBM behind it — with no shared control flow and no coordination except through memory. A single large monolithic processing block on the left contrasted with a dense grid of many small identical units on the right, each holding its own small queue of work.

Figure 1.0.1 — There is no "the processor" to make faster. There are 108 of them, and a kernel that fills three leaves 97% of the machine dark.

Each workshop is bad at waiting. An HBM access costs on the order of 400–600 ns (published order-of-magnitude figure for Ampere, not measured here) while an arithmetic instruction retires in a few nanoseconds. A CPU hides that with caches and out-of-order execution; a GPU hides it with brute-force concurrency. An A100 SM keeps 2,048 threads resident at once — 221,184 across the device, not queued but resident, registers allocated and state on chip — so whenever one stalls, another issues, at zero switching cost. Everything below follows from that one decision.

The mechanism¶

Warps: the thread is not the scheduling unit¶

Threads are a programming abstraction. The hardware schedules warps: groups of exactly 32 threads sharing one instruction pointer, executing the same instruction on different data in the same cycle. Ask for 100 threads and you get 4 warps — 128 slots — with 28 masked off.

When threads inside one warp take different branches, the hardware serialises them: it runs the if path with the else-threads masked off, then the else path with the if-threads masked off. Both cost full time. This is warp divergence. Divergence between warps is free.

Where it bites LLM inference	The divergent branch	Consequence
MoE routing	Tokens in one warp route to different experts	Serialised expert paths — why MoE kernels group tokens by expert first
Sampling	Top-$k$ / top-$p$ / penalties per sequence	Cheap absolutely, but it runs every step, per sequence
Ragged batches	Mixed sequence lengths in one batch	Why engines pad to block boundaries

The rule: make the branch condition uniform across a warp, or make both paths cheap. Branching on a per-batch flag is free; branching on token_id % 2 is a 2× tax.

Occupancy, and the myth that you want 100% of it¶

Occupancy is resident warps per SM over the hardware maximum:

\[ \text{occupancy} = \frac{W_{\text{active}}}{W_{\max}}, \qquad W_{\max} = 64 \text{ on A100 (SM 8.0)} \tag{1.1} \]

Three resources cap $W_{\text{active}}$ and the tightest wins: the register file (65,536 32-bit registers per SM), shared memory (164 KB usable per SM), and the block limit (32 per SM). Use 128 registers per thread and you cap at $65{,}536/(128 \times 32) = 16$ warps — 25% — whatever else you do. Occupancy exists for one purpose: enough resident work to cover memory latency. Once covered, more buys nothing. It is a threshold, not a score.

Worked example — how much occupancy is enough

Little's Law gives the bytes that must be in flight to sustain bandwidth $B_{mem}$ at latency $t_{lat}$:

\[ \text{bytes in flight} = B_{mem} \cdot t_{lat} \tag{1.2} \]

With $B_{mem} = 2.039 \times 10^{12}$ byte/s and $t_{lat} \approx 500$ ns that is 1.02 MB device-wide, or ≈ 9.4 KB per SM. A warp issuing one coalesced 32-bit load per thread puts $32 \times 4 = 128$ bytes in flight, so each SM needs $9{,}440 / 128 \approx 74$ outstanding loads. At 100% occupancy with one load per warp you get 64 — short. At 25% occupancy (16 warps) with 5 independent loads per warp, which any unrolled loop produces, you get 80, and you are done. Occupancy and instruction-level parallelism are substitutes.

The classic occupancy mistake

You cut register usage to lift occupancy from 50% to 75%, the compiler spills the evicted values into local memory — which is physically HBM — and the kernel gets slower while the metric gets better. Occupancy is an input to latency hiding, never an output. Measure achieved bandwidth.

Coalescing: layout is a performance decision¶

The memory system services transactions against 32-byte sectors grouped into 128-byte cache lines, not threads. When a warp's 32 threads read 32 consecutive 4-byte words, that is one cache line, four sectors, one trip — coalesced access. When those same threads read words more than 128 bytes apart, each lands in its own sector: 32 transactions fetching 1,024 bytes to deliver the 128 you asked for.

Left: a warp of aligned arrows landing on one contiguous band of memory. Right: the same arrows fanning out to widely separated bands, each pulling a full oversized block.

Figure 1.0.2 — Same 32 loads, same useful bytes. The right-hand pattern moves 8× the traffic and the memory system charges you for all of it.

Worked example — what a strided KV read costs

A warp reads 32 bf16 elements: 64 bytes of useful data. Coalesced, those 64 contiguous bytes touch 2 sectors → 64 bytes fetched, efficiency 100%.

Strided by one head: in a [B, T, H_kv, d_h] layout, consecutive tokens for a fixed head sit $H_{kv} \cdot d_h = 8 \times 128 = 1{,}024$ elements $= 2{,}048$ bytes apart. Every thread lands in a distinct sector → $32 \times 32 = 1{,}024$ bytes fetched for 64 useful bytes. Efficiency 6.25%, a 16× traffic multiplier; counting full 128-byte lines the worst case reaches 32×.

So attention kernels use [B, H_kv, T, d_h], where one head's whole key history is contiguous — and PagedAttention blocks hold 16 tokens rather than 1, because a 16-token block is 4,096 contiguous bytes per head. The paging scheme in Chapter 2.2 is a coalescing decision wearing an allocator's costume.

Tensor cores: the 312 TFLOP/s has conditions attached¶

An A100 SM contains two unrelated arithmetic resources. CUDA cores are scalar lanes — 64 FP32 per SM, 19.5 TFLOP/s device-wide — executing everything: activations, normalisation, elementwise adds, sampling. Tensor cores consume a small matrix tile and emit a matrix product in one instruction: 312 TFLOP/s dense bf16, 16× the FP32 vector rate.

Every headline GPU number is a tensor-core number. It applies to matrix multiplication only, and only when three conditions hold together. Dtype must be supported — on A100 that means bf16, fp16, tf32, int8, int4, not fp32 (tf32 substitutes) and not fp8, which Ampere lacks entirely. Shapes must align: the bf16 MMA primitive is 16×8×16, so cuBLAS wants contracted and output dimensions in multiples of 8 and prefers 64 or 128, and a dimension of 4,095 instead of 4,096 forces a padded copy or a slower path. Pointers must align to 16 bytes; a tensor sliced at an odd offset loses the fast path silently.

Llama-3.1-8B satisfies all of it by construction: $d = 4096$, $d_{ff} = 14336$, $d_h = 128$, and $V = 128{,}256 = 128 \times 1002$ — padded to stay both tensor-core-aligned and evenly shardable across tensor-parallel widths of 2, 4, and 8.

Where you will actually break alignment

Not in the model — in your code. LoRA ranks of 6 or 12, a classifier head with 7 classes, a vocabulary trimmed to exactly 50,001 entries. Each drops a GEMM off the tensor-core path, and the penalty is bounded by the spec ratio: up to 16×. Pad to a multiple of 8.

On newer silicon

Hopper (SM 9.0) adds wgmma — warpgroup-level asynchronous matrix multiply across 128 threads — plus thread block clusters and distributed shared memory, letting blocks on different SMs read each other's shared memory. Blackwell adds fifth-generation tensor cores with native fp4. FlashAttention-3 is built on wgmma and the Hopper Tensor Memory Accelerator and will not run on an A100. None of this exists on SM 8.0 and no success criterion here depends on it.

Kernel launch overhead: the part that decides your decode loop¶

Here is the payoff.

A kernel is one GPU program launch. Your Python calls torch.matmul; the dispatcher resolves the operator, selects a cuBLAS routine, and the driver writes a launch packet into a stream — an ordered queue the GPU consumes asynchronously. The CPU does not wait; it runs ahead.

sequenceDiagram
    autonumber
    participant Py as Python / PyTorch
    participant Drv as CUDA driver (CPU)
    participant Q as Stream queue
    participant SM as SMs (GPU)
    Py->>Drv: dispatch op, resolve kernel
    Drv->>Q: enqueue launch packet
    Note over Py,Drv: CPU cost per launch,<br/>paid whether or not the GPU is busy
    Py->>Drv: dispatch next op
    Drv->>Q: enqueue launch packet
    Q->>SM: launch grid of blocks
    SM-->>Q: retire
    Q->>SM: launch grid of blocks
    Note over Q,SM: GPU idles whenever the queue<br/>drains faster than the CPU refills it
    Py->>SM: synchronize
    SM-->>Py: all prior work complete

Figure 1.0.3 — The CPU is a producer, the GPU a consumer. Nothing goes wrong until the producer is slower, and then nothing you do to the kernels helps.

One launch through the PyTorch eager path — Python frame, dispatcher, driver call — costs 5–10 µs (an order-of-magnitude figure from common practice, not a measurement from this repo; you will produce your own in Do it). Because CPU and GPU run concurrently, step time is a maximum, not a sum:

\[ t_{\text{step}} = \max\!\big(\underbrace{N_k \cdot t_{\text{launch}}}_{\text{CPU floor}},\ \underbrace{t_{\text{GPU}}}_{\text{device floor}}\big) \tag{1.3} \]

Everything now turns on $N_k$, the kernel count per step. One decoder layer needs RMSNorm, QKV projection, RoPE, attention, output projection, residual add, second RMSNorm, gate and up projections, SiLU-and-multiply, down projection, and a second residual add — a dozen fused, 20–35 unfused. Multiply by $L = 32$.

Worked example — when a Llama-3.1-8B decode step becomes launch-bound

Device floor. Batch-1 decode reads every weight once: 14.96 GiB = 16.06 GB. At an achievable 1.5 TB/s, $t_{\text{GPU}} \approx 10.7$ ms.

CPU floor, fused path. $N_k = 500$ (≈15 kernels per layer), $t_{\text{launch}} = 7\ \mu$s: $500 \times 7\ \mu\text{s} = 3.5$ ms. So $\max(3.5,\ 10.7) = 10.7$ ms — launch overhead completely hidden. This is why people believe it does not matter.

CPU floor, eager path. Unfused execution reaches $N_k \approx 1{,}200$ at $9\ \mu$s, so $1{,}200 \times 9\ \mu\text{s} = 10.8$ ms. The floors tie, the queue runs dry constantly, and jitter in either one lands directly in TPOT.

Tensor parallelism across 4 A100s. Each rank streams a quarter of the weights — 4.0 GB, so $t_{\text{GPU}} \approx 2.7$ ms — but the kernel count is unchanged, so the CPU floor stays at 3.5 ms:

\[ \max(3.5,\ 2.7) = 3.5\ \text{ms} \quad \text{— launch-bound} \]

You bought four GPUs, cut the device floor 4×, and got 3.1×. At TP=8 the device floor drops to 1.35 ms while step time stays pinned at 3.5 ms: the extra hardware does nothing.

A CUDA graph is the fix. Execute the step once in capture mode and the driver records the whole DAG of launches into one replayable object. Replay costs one launch instead of 500, collapsing $N_k \cdot t_{\text{launch}}$ from 3.5 ms to roughly 10 µs, after which Equation 1.3 always resolves to the device floor — the only floor worth attacking. The price is rigidity: a graph freezes shapes and memory addresses, which is why vLLM captures a discrete set of batch sizes and pads up to the nearest, and why graphs cover decode but not prefill.

A horizontal execution timeline with wide idle gaps between short work blocks above, and the same blocks packed edge to edge without gaps below.

Figure 1.0.4 — The gaps are not slow kernels. They are the GPU waiting for the CPU to ask.

Streams and the synchronisation you keep forgetting¶

The asynchrony that makes launch overhead hideable also makes naive benchmarking meaningless.

# [1] WRONG — measures how fast Python enqueues work, not how fast the GPU does it.
t0 = time.perf_counter()
for _ in range(100):
    y = model(x)
print(time.perf_counter() - t0)   # microseconds. Impossible. You timed the queue.

# [2] RIGHT — block until the GPU has actually finished.
torch.cuda.synchronize()
t0 = time.perf_counter()
for _ in range(100):
    y = model(x)
torch.cuda.synchronize()
print(time.perf_counter() - t0)

The first version is not merely optimistic, it is absurd — 50 µs for work that must move gigabytes. Any result implying you beat the HBM bandwidth floor is a missing synchronize() until proven otherwise. Two rules follow: warm up before timing, discarding 10–20 iterations, and never call .item(), .cpu(), or print() on a GPU tensor inside a timed loop, because each forces an implicit synchronisation that destroys the overlap.

In practice¶

flowchart TD
    A[Measure step time<br/>after synchronize] --> B[Count kernel launches<br/>with torch.profiler]
    B --> C[Estimate CPU floor<br/>launches x per-launch cost]
    A --> D[Estimate GPU floor<br/>bytes moved / bandwidth]
    C --> E{{CPU floor > GPU floor?}}
    D --> E
    E -- Yes --> F[Launch-bound]
    E -- No --> G[Bandwidth-bound]
    F --> H[Capture CUDA graphs<br/>fuse ops, cut launch count]
    G --> I[Attack bytes<br/>quantise, batch, reuse]
    H --> J[Re-measure]
    I --> J

Figure 1.0.5 — Two floors, two completely different fixes. Compute both before optimising anything.

Flag / tool	What it controls	How to reason about it
`--enforce-eager`	Disables CUDA graph capture	A debugging switch, not a setting. Leaving it on costs the entire launch-overhead story above
`--cuda-graph-sizes`	Which batch sizes get captured	Capture the sizes your traffic produces. Each costs memory for its private pool; uncaptured sizes fall back to eager
`CUDA_LAUNCH_BLOCKING=1`	Forces every launch synchronous	Makes error messages point at the right kernel. Never benchmark with it set
`torch.profiler` (CUDA activity)	Ground truth on kernel count and gaps	Wide gaps on the GPU track with a busy CPU track means launch-bound
`nsys profile`	CPU and GPU tracks on one timeline	The only tool that shows the producer/consumer relationship directly
`torch.compile(mode="reduce-overhead")`	Fusion plus automatic CUDA graphs	The one-line version of this chapter's fix for research code

Two habits. nvidia-smi utilisation is not utilisation — it reports the fraction of sample intervals containing any resident kernel, which a gappy launch-bound loop trivially satisfies. And report kernel count alongside step time, which turns "the step got slower" into "the step issues 300 more kernels."

Failure modes¶

Symptom	Cause	Fix
Benchmark reports a step time below the bytes-over-bandwidth floor	Missing `torch.cuda.synchronize()` — you timed the enqueue, not the execution	Synchronize before starting and before stopping the timer; discard 10–20 warm-up iterations
Adding GPUs via tensor parallelism gives sublinear or zero speedup at low batch	Launch count is constant per rank while device work shrinks, so the CPU floor in Equation 1.3 binds	Capture CUDA graphs; confirm with a timeline that GPU gaps closed
`nvidia-smi` shows 95%+ utilisation but throughput is far below the roofline	Utilisation counts intervals with any resident kernel, not SM busy-ness	Measure achieved occupancy and bandwidth; read gap structure in `nsys`
Raising occupancy makes the kernel slower	Register pressure relief caused spills to local memory, which is physically HBM	Revert; optimise memory-level parallelism (unrolled independent loads) instead of the occupancy number
A GEMM runs an order of magnitude below its dtype peak	A dimension is not a multiple of 8, or a slice broke 16-byte alignment, so cuBLAS took the non-tensor-core path	Pad dimensions to a multiple of 8 (64 or 128 is better); check the selected kernel name in the profiler
Attention throughput collapses after a KV layout change	Head-major became token-major, so a warp's reads stride by $H_{kv} d_h$ and each thread pulls its own sector	Restore head-contiguous layout; confirm with sector-efficiency counters
An MoE layer is far slower than its FLOP count predicts	Tokens in one warp route to different experts, serialising the expert paths	Group tokens by expert before the expert GEMMs so each warp is routing-uniform

Do it¶

There is no committed code artifact for this chapter yet. Build the measurement yourself — forty lines, and it makes Equation 1.3 real.

The experiment. Allocate two small CUDA tensors, small enough that GPU work per operation is negligible (a 1,024-element add_ is ideal). Time three things over 10,000 iterations each, discarding 20 warm-up iterations, with torch.cuda.synchronize() before and after each timed region: (1) the loop without a trailing synchronize — the wrong measurement, produced deliberately; (2) the same loop correctly bracketed; (3) 100 of those operations captured with torch.cuda.graph(g) and replayed 100 times.

Success criteria, all objective: (1) is at least 5× smaller than (2), reproducing the fake-speedup failure mode on your own hardware; (2) lands in the 2–15 µs range — that is your $t_{\text{launch}}$; (3) is at least 5× faster than (2).

Then close the loop: count the kernels in one decode step of any 7–8B model with torch.profiler and evaluate Equation 1.3 against the bandwidth floor from Chapter 1.1. You will then be able to state, before running anything else, at what tensor-parallel width your decode loop stops being a GPU problem.

Summary¶

A GPU is 108 independent SMs, not one processor. It hides ~500 ns memory latency with 221,184 resident threads.
The scheduling unit is the warp — 32 threads in lockstep. Divergence inside a warp serialises both paths; divergence between warps is free.
Occupancy is a threshold, not a score. Memory-level parallelism substitutes for it one-for-one.
Coalescing decides effective bandwidth. Contiguous reads cost one transaction; striding by a head dimension costs 8–32×, which is why KV cache is head-contiguous and paged in 16-token blocks.
Step time is $\max(\text{CPU floor}, \text{GPU floor})$. Several hundred launches at ~7 µs is a 3.5 ms CPU floor — invisible at batch 1 on one GPU, dominant at TP=4. CUDA graphs collapse it to a single launch, and that is the whole reason they exist.

Key terms¶

SM · warp · occupancy · coalesced access · tensor core · CUDA graph · bandwidth-bound

Exercises¶

Recall

How many threads execute in lockstep in a warp, and what does divergence inside one cost?
State Equation 1.3 in words, and name one action that reduces each of its two terms.

Derive

A kernel uses 96 registers per thread on an A100 SM with 65,536 registers. What is the maximum occupancy from register pressure alone? Is that necessarily a problem?
Your decode step issues 640 kernels at $t_{\text{launch}} = 8\ \mu$s, with a device floor of 4.2 ms. Compute the step time at TP=1, 2, and 4. At which width does the loop become launch-bound?

Design

You serve Llama-3.1-8B on 4× A100 with TP=4. TPOT is 3.6 ms against a computed device floor of 2.7 ms, nvidia-smi reads 97%, and batch sizes range from 1 to 48. Diagnose it, propose an ordered set of fixes, and state what each costs.

Worked solutions

**1.** 32 threads. On divergence the hardware runs both paths sequentially with non-participating threads masked off, so the warp pays the *sum* of both durations, not the maximum. Divergence between different warps costs nothing. **2.** $t_{\text{step}} = \max(N_k \cdot t_{\text{launch}},\ t_{\text{GPU}})$: a step takes as long as the slower of "the CPU asking for the work" and "the GPU doing it", because the two run concurrently. Cut the first term by cutting $N_k$ — fusion or CUDA graph capture. Cut the second by cutting bytes or FLOPs — quantisation, batching, better reuse. **3.** $65{,}536 / (96 \times 32) = 21.3$, so **21 warps**, and occupancy is $21/64 \approx 33\%$. Not necessarily a problem: an SM needs ~74 outstanding loads to saturate HBM, and 21 warps supply that if each sustains about 4 independent loads in flight — and the 96 registers are probably *why* it can. Cutting registers to raise occupancy would reduce memory-level parallelism and risk spills to local memory: a better number, a worse kernel. **4.** CPU floor $= 640 \times 8\ \mu\text{s} = 5.12$ ms. Device floor 4.2 ms at TP=1, 2.1 ms at TP=2, 1.05 ms at TP=4 — and the step stays at **5.12 ms** in all three. The loop is launch-bound at every width including TP=1, so scaling out cannot help until $N_k$ falls. Capture a CUDA graph and the CPU floor drops to ~10 µs, after which TP=1 gives 4.2 ms and TP=2 gives 2.1 ms. A flat tensor-parallel scaling curve signals a CPU-side bottleneck, not a communication one. **5. Diagnosis.** The 0.9 ms gap above the 2.7 ms floor is, at ~7 µs per launch, about 130 kernels' worth of enqueue that is not being hidden. The `nvidia-smi` reading is worthless — it reports intervals containing any resident kernel. Confirm with `nsys` first: wide regular gaps on the GPU track with a saturated CPU track means launch-bound, while gaps clustered around NCCL all-reduces mean communication and this plan is wrong. **Ordered fixes.** 1. **Verify CUDA graphs are on.** One unsupported op silently forces the eager path for the whole step. Free. 2. **Widen `--cuda-graph-sizes`.** Batches from 1 to 48 mean uncaptured sizes fall back to eager, producing bimodal TPOT. Capture a ladder (1, 2, 4, 8, 16, 24, 32, 48). Price: memory per pool, plus padding waste when 17 is padded to 24. 3. **Reconsider TP=4**, which also pays four all-reduces per layer. If the SLO is met at TP=2, two TP=2 replicas double throughput on the same hardware. Price: 15 GiB of duplicated weights per replica, straight out of the KV cache budget in [Chapter 2.2](M2T2-attention-kv-cache.md). 4. **Raise batch size through admission control.** Launch cost is per step, not per sequence. Price: queueing delay, which lands in TTFT. The ordering is the lesson: the top two are configuration rather than engineering, and they recover most of the gap. The kernels were never the problem.

Going deeper¶

CUDA C++ Programming Guide — authoritative on warps, occupancy limits, and coalescing rules
NVIDIA A100 architecture whitepaper — source of the 108 SMs, 192 KB L1/shared, and 312 TFLOP/s figures
Getting Started with CUDA Graphs — the launch-overhead argument and the capture/replay API
Volkov, Better Performance at Lower Occupancy — the talk that killed the 100%-occupancy myth
PyTorch — CUDA semantics — streams, asynchrony, and why timing needs synchronize()

Next: Chapter 1.1 — The GPU Memory Hierarchy, which takes the SMs you now understand and asks where their operands actually live.