Few, powerful cores; deep out-of-order pipelines; large caches; optimized to run one thread of arbitrary code as fast as possible, including code full of unpredictable branches. Excellent for orchestration, control flow, and anything that isn’t dense regular math.

A GPU packs thousands of simple ALUs organized into Streaming Multiprocessors (SMs). Threads execute in groups of 32 (a "warp") in lockstep — Single Instruction, Multiple Threads (SIMT). Rather than avoiding memory-access latency the way a CPU does (huge caches, branch prediction), a GPU hides it: while one warp waits on a memory read, the scheduler simply runs another of the thousands of resident warps. This only works because GPU workloads are regular and abundantly parallel — exactly what matrix multiplication is.

On top of the general SIMT cores, modern GPUs add Tensor Cores: fixed-function units that perform a small matrix-multiply-accumulate () on a tile of values every cycle, at far higher throughput per watt than doing the same math on general ALUs. NVIDIA’s fifth-generation Tensor Cores (Blackwell) natively support FP4, FP6, FP8, FP16, BF16, TF32 and INT8 — the accelerator literally has different physical circuitry paths for different numeric precisions, which is why the precision choices in Part II are not just a software convenience but a hardware decision.

Google’s TPUs take a different dataflow approach: the Matrix Multiply Unit (MXU) is a systolic array — a grid of processing elements arranged so that data (weights and activations) flows through the array and is reused across many multiply-accumulate steps without being re-read from memory each time. TPU v6e (Trillium) uses a 256×256 MXU (65,536 MACs operating per cycle) delivering roughly 918 peak BF16 TFLOPS per chip with 32 GB of HBM; TPU v7 (Ironwood) keeps the 256×256 MXU but reaches roughly 2,300 peak BF16 TFLOPS and 4.6 PFLOPS at FP8, backed by 192 GiB (about 206 GB) of HBM3e at 7.4 TB/s. Anthropic has committed to deploying over one million Ironwood chips beginning in 2026 ^[1][2]. Try it: Compare TPU and GPU rooflines.

At the small end of the spectrum, Neural Processing Units in phones and laptops (Apple’s Neural Engine, Qualcomm’s Hexagon NPU) are small, extremely power-efficient matrix engines, usually running INT8/INT4-class precision, tightly coupled to a shared pool of on-package memory. They exist to run inference locally within a battery’s power budget, not to train anything. Part VI returns to NPUs, FPGAs, ASICs, and two more exotic architectures — in-memory analog computing and wafer-scale integration — in much greater depth.

It is worth walking that tiled matmul through the hardware once, concretely, because the same execution model underlies essentially everything that runs on a GPU. The work is expressed as a grid of thread blocks; each block is assigned one output tile and runs entirely on one Streaming Multiprocessor (SM) — an H100 has 132 of them. Within a block, threads are grouped into warps of 32 that execute in lockstep (Single Instruction, Multiple Threads — SIMT: all 32 threads run the same instruction on different data). A block’s threads cooperatively copy a tile of A and a tile of B from slow HBM into the SM’s fast shared memory (a software-managed scratchpad, up to 228 KB per SM on Hopper), hit a __syncthreads() barrier so no thread reads the tile before every thread has finished writing it, and then each thread multiply-accumulates its sub-tile with the running sum held in registers (256 KB per SM); on a tensor core, this step is a fused multiply-accumulate over a small matrix fragment (canonically 16×16×16). The block loops over successive K-tiles this way and writes the finished output tile back to HBM ^[12][13].

The reason this is fast is latency hiding. A read from HBM costs hundreds of cycles, and a CPU would fight that latency per-thread with big caches and out-of-order speculation. A GPU does the opposite: it keeps a large number of warps resident on each SM (up to 64 on an H100), and whenever one warp stalls waiting on memory, the scheduler simply issues the next instruction from a different ready warp on the following cycle. With enough resident warps ("occupancy"), the arithmetic units almost never sit idle, and the hundreds-of-cycles memory latency is hidden behind other warps' useful work — which is the whole reason a throughput machine beats a latency-optimized one at this kind of regular, parallel arithmetic (Figure 5) ^[12][13].

FP32 was the historical default for training — full range, full precision, but four bytes per value. TF32 (Ampere-era) is a compute-only compromise: it keeps FP32’s 8 exponent bits (so it never overflows/underflows the way FP16 does) but truncates the mantissa to 10 bits for the matmul itself, transparently, while data still sits in memory as FP32. BF16 keeps FP32’s 8-bit exponent (same dynamic range) but only 7 mantissa bits — it sacrifices precision, not range, which makes it numerically stable for training and is why it became the default 16-bit training format. FP16 instead has 5 exponent bits and 10 mantissa bits — more precise than BF16 but with a much smaller dynamic range, which historically required loss-scaling tricks during training to stop gradients from underflowing to zero.

FP8 comes in two hardware variants because one format can’t serve every tensor well: E4M3 (4 exponent, 3 mantissa bits) has more precision and less range, and is generally used for weights and forward-pass activations; E5M2 (5 exponent, 2 mantissa) trades precision for range and is used where wide dynamic range matters most — chiefly gradients, whose magnitudes vary enormously during training and which are prone to underflow or overflow in a narrower format ^[3][4]. Try it: E4M3 tops out at 448; E5M2 reaches 57,344.

Naively representing a whole tensor at 4 bits per value (plain INT4 or FP4) destroys accuracy, because a single shared scale factor for an entire tensor can’t cope with the outlier activations that real neural networks produce. The fix that made sub-8-bit formats practical is micro-scaling: instead of one scale for a whole tensor, split it into small blocks and give each block its own scale factor, so a block containing an outlier doesn’t force every other block in the tensor to waste dynamic range accommodating it. Try it: One outlier, four 4-bit formats.

This idea has been standardized industry-wide as the Open Compute Project’s Microscaling (MX) format, backed by AMD, Arm, Intel, Meta, Microsoft, NVIDIA, and Qualcomm. MX formats use a block size of 32 elements sharing one 8-bit power-of-two scale (called E8M0). Each stored element is dequantized back to its real value by multiplying it by the block’s shared scale:

where is the low-precision element (an FP4 or FP8 code) and is the block’s shared 8-bit power-of-two exponent — so one scale amortized across 32 values restores the dynamic range each narrow element cannot hold alone.

MXFP8 uses FP8 elements (E4M3 or E5M2) per block; MXFP4 uses FP4 (E2M1: 1 sign, 2 exponent, 1 mantissa bit) elements per block, occupying about 136 bits (17 bytes) for a 32-value block versus 1,024 bits unquantized ^[5][6]. Try it: One 136-bit MXFP4 block.

NVIDIA’s own NVFP4, exclusive to Blackwell-generation GPUs (B200, B300, and the RTX Pro 6000/5000/4000 lines), takes the same E2M1 4-bit element format but refines the scaling scheme with an extra level. Concretely, here is what hardware does to turn one stored 4-bit code back into a usable number:

First, every 4-bit code in a 16-value block is looked up against the E2M1 table — 1 sign bit, 2 exponent bits, 1 mantissa bit — which can represent only eight possible magnitudes: 0, 0.5, 1, 1.5, 2, 3, 4, and 6. A code like 1110 decodes to sign=1, exponent=11, mantissa=0, giving magnitude 4 and signed value −4 ^[57][58]. Second, that raw magnitude is multiplied by a per-block scale: one shared FP8 (E4M3) scale factor covers all 16 values in the block, restoring the dynamic range that a bare 4-bit code cannot hold on its own — this per-block scale is the single feature that separates NVFP4 from plain, much less accurate INT4. Third, that block scale is itself multiplied by one more scale, a single FP32 value shared across the entire tensor, whose job is to keep every block’s FP8 scale well-conditioned no matter how the tensor’s overall magnitude varies. The final dequantized value is therefore a product of three terms at three granularities: Try it: Decode E2M1 code 1110.

This three-level value is what actually enters the Tensor Core’s multiply-accumulate array, with accumulation still happening in higher precision (FP16/FP32) so rounding error doesn’t compound across a GEMM’s thousands of summed terms.

This two-level scaling recovers most of the accuracy lost to plain 4-bit representations — NVIDIA reports accuracy close to FP8 at roughly 4.5 effective bits per value, around 1.8× less memory than FP8 and about 3.5× less than FP16. NVIDIA has shown up to a 10× throughput improvement on some workloads (e.g., the FLUX.2 image model) when moving to NVFP4 inference ^[7][8][9]. Try it: NVFP4 at ≈4.5 bits/value; FP4 weights: 70B in 35 GB.

Why two competing 4-bit formats exist: MXFP4 is the open, cross-vendor standard (AMD, Intel, and others can build hardware for it), using 32-value blocks; NVFP4 is NVIDIA’s tuned variant with a smaller 16-value block and an extra FP32 scaling level, trading a bit of standardization for better accuracy-per-bit on Blackwell hardware specifically. Expect both to matter: MXFP4 for portability, NVFP4 where you’re already committed to NVIDIA silicon.

Fixed-point integer quantization (uniform steps, a zero-point, calibrated ranges) was the original way to shrink edge-inference models and is still common on NPUs. It is gradually losing ground in the datacenter to floating-point block formats (FP8/FP4/MX/NVFP4) because floating-point’s built-in exponent handles the outlier-heavy distributions of real activations far more gracefully than a single fixed integer scale. Try it: INT4 vs FP4 on outliers.

Every transformer layer turns each token’s hidden state into three separate vectors via three learned weight matrices: a Query ("what am I looking for"), a Key ("what do I offer, for others to match against"), and a Value ("what do I actually contribute, if attended to"). Attention scores the current token’s Query against every other token’s Key — a dot product, scaled and passed through softmax — to produce a weight for each past token, then uses those weights to compute a weighted sum of every other token’s Value vector:

Try it: Try the attention equation.

The detail that makes caching possible — and necessary — is this: a given token’s Key and Value vectors depend only on that token’s own hidden state at that layer, computed once via that layer’s fixed, already-trained and weight matrices. They do not depend on which later token is doing the attending. Token 12’s Key vector is exactly the same value whether it is being read by the attention computation for token 13, token 5,000, or token 50,000 — it never needs to be recomputed once produced. Queries are the opposite: a fresh Query is needed for every step, because it represents "what the current token is looking for," which is a new question at every position. This asymmetry — Keys and Values are reusable and static once computed; Queries are single-use and always fresh — is the entire reason a KV cache exists and not, say, a "QKV cache." There is nothing to gain by caching Queries, because each one is used exactly once, in exactly one step, and then discarded.

The clearest way to see the saving is to walk through a single decode step concretely. Suppose five tokens have already been generated ( through ), and the model is about to produce a sixth. At every layer, the KV cache already holds and — five Key vectors and five Value vectors, sitting in HBM exactly as they were computed on earlier steps. Generating requires, at this layer: Try it: One decode step, visually.

Five of the six Key/Value pairs involved in this step’s attention computation were read, not recomputed — and every one of the model’s other layers runs this identical read-five/compute-one pattern in parallel, each against its own layer’s cache.

Try it: Read five, compute one.

Contrast that with what would happen without a cache. Token 7’s step would need and to score against — but rather than reading six already-computed vectors, a cache-less implementation would have to re-run tokens 1 through 6 through this layer’s projections all over again, from their raw hidden states, just to reproduce six vectors that are mathematically identical to what was already sitting in memory a step earlier. Token 8’s step would redo seven projections; token 50’s step would redo forty-nine. The vectors recomputed this way are not approximations or updates — they are bit-for-bit (modulo the non-associative floating-point effects of Part IV) the same values already computed, because a given token’s Key and Value at a given layer never change once that token’s hidden state at that layer is fixed. Recomputing them is pure waste: real GPU cycles and real HBM bandwidth spent reproducing a known answer. Try it: Decoding without a KV cache.

The reason this matters so much is the iterative nature of decoding itself. Generating a response is not one computation — it is one computation per output token, each one waiting on the last, because token N+1 cannot be sampled until token N exists (Part IX, Section 37 covers this data dependency in full). Every one of those steps touches the entire history of the sequence so far: step 50 needs to attend across all 49 prior tokens, and step 5,000 needs to attend across all 4,999 prior tokens. Without a cache, the total Key/Value-projection work summed across an N-token generation grows as , which is — quadratic in sequence length, because every step redoes all the projection work every earlier step already did. With a cache, each step contributes exactly one new projection regardless of how long the sequence already is, so the total work summed across the whole generation grows as — linear. The gap between those two curves widens with every single token generated, which is why the benefit of caching is not a fixed constant-factor speedup but one that grows without bound as conversations, documents, and agentic contexts get longer. Try it: KV cache cumulative work; O(N²) vs O(N) projection work; Sum 1+…+N for 50 tokens.

Try it: Cumulative work, with and without cache.

This is also precisely why the cache is sequence-dependent memory rather than a fixed cost like the model’s weights: the longer the conversation, the more entries accumulate, one Key/Value pair per token per layer per KV-head, forever (until the sequence ends or is evicted). Section 16’s worked example puts real numbers on this — a single 8,000-token conversation on a 70B-parameter model already carries about 2.6GB of KV cache, and a 128,000-token context about 42GB, for one request — and because every concurrent request needs its own cache, this growth is what typically caps how many simultaneous users a server can hold long before raw compute becomes the bottleneck. Try it: 8,000-token cache on 70B; KV cache at 8,000 tokens; KV cache at 128,000 tokens.

Multi-head attention runs several independent copies of this Q/K/V mechanism per layer — each "head" with its own learned projections — letting different heads specialize in different relational patterns, and concatenating their outputs back together. In the original transformer design, every one of those heads gets its own Key and Value projection, so the cache scales with the full head count. Because the cache’s cost is a direct function of how many distinct Key/Value projections exist to store, several architectural variants attack the cache’s size by deliberately reducing that count, at some cost in representational capacity: Try it: MHA cache without sharing.

Try it: MHA, GQA, MQA, MLA compared.

All four variants compute the identical attention equation above — they differ only in how many distinct Key/Value projections exist to be written to, and read back from, the cache. Fewer cached projections directly means fewer bytes the memory system has to move per decode step, which is the exact memory-bandwidth bottleneck this section opened with; the cost is some shared representational capacity across whichever query heads are grouped onto the same K/V head. This is also why the worked example in Section 16 flags GQA specifically: swapping that model’s actual 8-KV-head GQA configuration for full 64-head multi-head attention would make its KV cache eight times larger — the 2.6 GB single-conversation figure above becomes about 21 GB for the exact same conversation, for no gain except representational headroom the model may not even need ^{[192][193][194]}. Try it: Full MHA vs GQA cache; Full MHA: eight times larger.

Everything above concerns the cache’s size — how much data must be stored and moved. PagedAttention, introduced with the vLLM serving engine, attacks a completely different inefficiency: how that memory is allocated, once its size is fixed. A naive implementation reserves one contiguous block of GPU memory per request, sized for the maximum sequence length that request could possibly reach — which wastes enormous amounts of memory twice over: internal fragmentation (a request that only ever generates 60 tokens still reserves space for, say, 4,000) and the impossibility of reclaiming or reusing that reserved-but-unused space for any other request in the meantime. This is precisely the problem operating systems solved decades ago for process memory, and PagedAttention borrows the solution directly: it stores each request’s KV cache not as one contiguous allocation but as a set of fixed-size, non-contiguous "pages" pulled from a shared pool, with a small per-request block table mapping that request’s logical token positions to whichever physical pages actually hold them — the same logical-to-physical indirection a CPU’s page table provides for ordinary process memory. Try it: 60 tokens reserving 4,000.

Try it: PagedAttention allocator simulation.

Two consequences follow directly from that indirection. First, memory is only ever consumed for tokens a request has actually generated — there is no up-front reservation for a worst-case maximum length, which alone recovers most of the wasted memory naive allocation incurs. Second, and just as consequential for serving throughput: because any request’s block table can point at any physical page, two different requests that happen to share a token prefix — a common system prompt, a repeated few-shot preamble, a shared conversation branch — can have their block tables point at the very same physical pages instead of each storing an independent copy. That shared-page mechanism is, in effect, the low-level substrate underneath the higher-level KV-cache-reuse techniques covered in Section 31: vLLM’s automatic prefix caching hashes these fixed-size blocks to detect matching prefixes across otherwise-unrelated requests, and SGLang’s RadixAttention generalizes the same idea into a full prefix tree spanning many concurrent requests. Together, the size-reduction techniques above (GQA/MQA/MLA) and PagedAttention’s allocation strategy are complementary, not competing: one shrinks how much has to be cached per token, the other eliminates the waste in how whatever remains gets stored ^{[16][17][489][490][491]}. Try it: Shared-prefix pages.

FlashAttention addresses a third, related but distinct problem — not the cache’s size or its allocation strategy, but the memory traffic generated by the attention computation itself while it runs. A naive implementation computes the full N×N matrix of attention scores (every token’s Query against every other token’s Key) and materializes that entire matrix in slow HBM before applying softmax and multiplying by Values — a huge, transient blow-up in memory traffic that scales quadratically with sequence length, on top of whatever the KV cache itself already costs. FlashAttention avoids ever forming that full matrix: it tiles the computation into blocks small enough to stay resident in fast on-chip SRAM, and uses an "online softmax" — a running, incrementally-updated version of the softmax normalization — so that each tile’s contribution to the final result can be folded in as it’s computed, without ever needing the full row of scores in memory at once. The N×N score matrix never touches HBM, so the extra memory drops from O(N²) to O(N) and HBM traffic falls several-fold (it is still quadratic in N, but divided by the on-chip block size), pushing attention much closer to compute-bound. The result is mathematically identical, differing from standard attention only by floating-point rounding. Try it: Online softmax, step by step.

Try it: FlashAttention HBM traffic compared.

The three techniques in this section therefore attack three genuinely separate axes of the same underlying problem — attention is expensive to move through memory, at every scale from a single step to an entire multi-tenant server fleet — and production inference stacks run all three simultaneously: FlashAttention (or its variants, e.g. FlashAttention-2/3, FlashDecoding) keeps the per-step attention computation itself from blowing out HBM traffic; GQA/MQA/MLA shrink how much has to be cached and moved per step in the first place; and PagedAttention (plus prefix caching and RadixAttention built on top of it) manages whatever cache remains as efficiently as an operating system manages process memory.

It is worth being concrete about what actually makes 72 GPUs behave like one, because "NVSwitch" names a specific and formidable piece of hardware, not just an abstract fabric. NVSwitch is a dedicated crossbar-switch ASIC — its own ~50-billion-transistor die — that GPUs plug into, rather than being wired point-to-point to one another. The distinction matters: with point-to-point links a GPU can only reach the specific peers it happens to be cabled to, which caps the communication patterns a workload can use; routing every GPU’s NVLink links into a set of NVSwitch crossbar chips instead creates a non-blocking all-to-all fabric, where any GPU can talk to any other at full NVLink bandwidth simultaneously. That is precisely the property that lets a rack of GPUs present itself to software as one enormous GPU with one shared address space ^[392][393].

The Blackwell-generation NVSwitch chip carries 72 NVLink ports and switches roughly 3.6 TB/s of traffic in each direction (7.2 TB/s bidirectional, in NVIDIA’s usual bidirectional convention; 18 such chips give the rack’s ~130 TB/s — the same doubling applies to the 1.8 TB/s-per-GPU figure, which is 900 GB/s each way). Physically, a GB200/GB300 NVL72 rack is built from 18 compute trays (holding the 72 GPUs and their Grace CPUs) and 9 NVLink switch trays, each switch tray carrying 2 NVSwitch chips — 18 switch chips in total. Each GPU’s 18 NVLink links fan out across those switch chips over a passive copper backplane, knitting all 72 GPUs into a single non-blocking domain at the roughly 130 TB/s aggregate quoted above ^{[392][393][394]}. Try it: One NVSwitch chip’s 72 ports; NVL72 trays and links.

Two capabilities extend this beyond raw switching. First, SHARP (Scalable Hierarchical Aggregation and Reduction Protocol) is in-network compute: small arithmetic units inside the NVSwitch itself sum values — gradients, for instance — as they pass through the switch, so a collective all-reduce is computed in the network rather than by shuffling every value between every GPU. Because the switch returns the already-reduced result, this roughly doubles effective all-reduce throughput and offloads that work from the GPUs' own compute units; the same idea exists in NVIDIA’s InfiniBand switches for the scale-out domain ^[395]. Second, the NVLink Switch System tiers these switch chips to extend the coherent NVLink domain beyond one rack: the earlier DGX GH200 reached 256 GPUs this way, the current flagship is NVL72, and NVIDIA has announced an NVL576 configuration (576 Rubin Ultra GPUs across eight racks, joined by a second tier of NVLink switches) for 2027 — a target credible analysts report may slip toward 2028 ^[394][396]. NVIDIA has also begun licensing this fabric: NVLink Fusion (announced 2025) lets third-party CPUs and custom accelerators — with partners including MediaTek, Marvell, Fujitsu, and Qualcomm — connect into the NVLink ecosystem, which is best read as NVIDIA selectively opening its interconnect moat to pull custom silicon toward its fabric rather than toward the open UALink standard (Part XVII) ^[397][398]. Try it: SHARP vs ring all-reduce.

One last piece of the connectivity story points forward. The 130 TB/s NVLink domain rides on a passive copper backplane precisely because copper is cheap, power-free, and reliable at the ~1-meter reach inside a rack — but copper’s usable reach collapses at these signaling rates (Section 39), so the moment a fabric has to cross racks, rows, or halls, it must turn to optics. That is why the scale-out tier (InfiniBand/Ethernet) is where NVIDIA has pushed hardest on co-packaged optics — its Quantum-X and Spectrum-X Photonics switches — and why, as GPU counts climb toward the millions, the interconnect and its power become a first-order constraint. Part XII’s photonics section covers that connectivity shift in depth; the point to carry here is that the switch silicon and the optical link are two halves of the same problem — moving bytes between accelerators fast enough that the math units never wait ^[399]. Try it: Why copper reach collapses.

At cluster scale, the scale-out network itself has a specific, deliberate shape. Rather than a generic flat network, GPU clusters are built as a rail-optimized fat-tree (a Clos/leaf-spine topology): every node’s "GPU 0" NIC connects to "Leaf Switch 0," every node’s "GPU 1" NIC connects to "Leaf Switch 1," and so on, so that each of the eight (or more) GPU rails forms its own dedicated, non-blocking switching plane all the way up to a layer of spine switches that stitch the rails together. A collective operation that needs to talk to "the same GPU index" across many nodes — which is exactly the traffic pattern most all-reduce and all-to-all operations produce — never has to leave its rail, avoiding contention with the other seven rails' traffic entirely. The fat-tree property (spine capacity sized so that any leaf can talk to any other leaf at full line rate) means the network doesn’t degrade as more racks are added, at the cost of needing roughly twice as many switches as a naively oversubscribed design ^[59][62][66]. Try it: Fat-tree switch calculator; Rail-optimised wiring; Non-blocking vs oversubscribed switches.

Try it: Three-tier fat tree, k = 64.

Some large operators use variants such as dragonfly topologies (which trade some worst-case bisection bandwidth for fewer long-haul cables and lower cost at extreme scale) but the rail-optimized fat-tree remains the default reference architecture NVIDIA and most cloud GPU providers ship today for AI training and inference clusters ^[59].

The network topology only matters because of what runs on top of it: collective communication primitives, implemented in libraries like NCCL, that every GPU in a training job calls on every optimizer step. All-reduce — every GPU ends up with the sum (or average) of a tensor that started out different on every GPU — is the single most consequential of these for training throughput, and there are two structurally different ways to implement it.

Try it: Ring vs tree, step by step.

Ring all-reduce arranges every GPU into a logical ring and runs two passes around it: a reduce-scatter pass where each GPU accumulates a different slice of the final result from its neighbor, followed by an all-gather pass where each GPU broadcasts its completed slice around the ring so every GPU ends up with the full result. Every GPU only ever sends to and receives from its two immediate ring neighbors, and — critically — every GPU is sending and receiving simultaneously, so the algorithm keeps every link in the topology saturated at once. This makes ring all-reduce bandwidth-optimal: total time scales with data size and is essentially independent of how many GPUs are in the ring, which is why it’s the default NCCL algorithm for large tensors and large GPU counts.

Tensor parallelism splits an individual layer’s weight matrices across GPUs — e.g., each of 4 GPUs holds one-quarter of a feed-forward layer’s columns — so a single matmul becomes 4 smaller matmuls plus a synchronization step. Because every forward and backward pass through a TP-split layer requires an all-reduce (or all-gather) of activations between the participating GPUs, TP is extremely communication-heavy relative to the compute it does, which is exactly why it is confined to GPUs connected by NVLink within a single node — it should almost never cross the slower scale-out network.

Pipeline parallelism instead assigns different contiguous ranges of layers to different GPUs (or groups of GPUs) — stage 1 holds layers 1–8, stage 2 holds layers 9–16, and so on — and streams micro-batches through the stages like an assembly line. Communication between stages is much lighter (just the activation tensor at the layer boundary, once per micro-batch) than TP’s per-layer synchronization, so pipeline stages can be spread across nodes, tolerating the slower scale-out network far better. The cost is "pipeline bubbles" — idle time while the first and last stages wait for the pipeline to fill and drain — which schedulers try to minimize with many small micro-batches and interleaved stage assignment. Try it: Pipeline bubble vs micro-batches.

Data parallelism is the oldest and conceptually simplest strategy: replicate the entire model on multiple GPUs (or groups of GPUs already running TP+PP internally), split the training batch across the replicas, and all-reduce gradients across replicas after each step. DP scales throughput almost linearly with GPU count as long as the interconnect can keep up with the gradient all-reduce, but it does not reduce per-GPU memory pressure at all — every replica still needs to fit a full copy of the model (or, with ZeRO/FSDP, a carefully sharded copy) — which is why DP is layered on top of TP/PP rather than used as a substitute for them at frontier model scale.

Mixture-of-Experts models add a fourth axis: expert parallelism distributes different experts across different GPUs, and a learned router sends each token to only the small subset of experts (often 2 of many dozens) it needs, activating only a fraction of total parameters per token. This is what lets MoE models like Mixtral or DeepSeek-V3 have enormous total parameter counts while keeping the actual FLOPs-per-token — and therefore inference cost — much closer to that of a far smaller dense model. The tradeoff is an all-to-all communication step (tokens must be routed to whichever GPU holds their assigned expert and results routed back), which is latency-sensitive and, like TP, benefits enormously from staying inside the fast NVLink domain. Try it: Expert parallelism memory and all-to-all; Expert-parallel all-to-all traffic; Total vs active parameters.

Frontier training runs typically combine all four simultaneously — commonly described as "3D" (TP+PP+DP) or "4D" (adding EP) parallelism — with each axis deliberately sized to match a tier of the network hierarchy: tensor-parallel groups sized to fit within one NVLink domain, pipeline stages spanning nodes within a rail-optimized cluster, and data-parallel replication spanning the full scale-out fabric. Getting this mapping right — rather than any single algorithmic trick — is most of what separates a cluster that achieves 40–50% of its theoretical peak FLOPs from one that achieves under 20% ^[64][65]. Try it: Plan Llama 3’s 4D layout.

Putting numbers to all of this makes the memory story concrete. Take a 70-billion-parameter dense model (a Llama-70B-class model) and ask what it takes to serve. The weights alone are parameters × bytes-per-parameter: 140 GB in FP16 (2 bytes), 70 GB in FP8, or about 35 GB in FP4 (0.5 bytes). Since a top-end GPU holds 80–192 GB of HBM, the FP16 version does not fit on one GPU — its weights need at least two, and in practice it is served across a tensor-parallel 8-GPU node — whereas quantized to FP4 (~35 GB) it fits comfortably on a single 80 GB GPU, which is exactly why the compression techniques of Part XVI matter so directly to deployment cost ^[421][427]. Try it: 70B in FP16: 140 GB; 70B in FP8: 70 GB; 70B in FP4: 35 GB; 70B in FP4 fits one GPU.

The KV cache is the second, sequence-dependent memory cost, and it is easy to underestimate. Its size is . For this model (80 layers, and — thanks to grouped-query attention — 8 key/value heads of dimension 128, in FP16) that works out to roughly 0.33 MB per token (327,680 bytes, or 0.31 MiB), so a single 8,000-token conversation carries about 2.6 GB of KV cache, and a 128,000-token context about 42 GB — for one sequence. (Had the model used full multi-head attention with 64 KV heads instead of GQA’s 8, that cache would be eight times larger — which is precisely why the attention-shape tricks of Part IX exist.) Because every concurrent request needs its own KV cache, this is what caps batch size long before compute does ^[428]. Try it: KV cache at 8,000 tokens; 70B KV cache per token; 0.3125 MiB = 0.33 MB; KV cache: 0.33 MB per token; KV cache at 128,000 tokens; 8K tokens: 2.5 GiB; Full MHA vs GQA cache; 128,000-token KV cache; 128K tokens: 40 GiB; MHA instead of GQA: 8×; Batch limit from KV capacity.

Finally, a back-of-envelope on speed: decode is memory-bound, so single-stream tokens per second is roughly HBM bandwidth divided by the bytes read per token (≈ the model size, since every weight is read once per token). For the FP16 model (~140 GB) at one H100’s bandwidth (~3.35 TB/s) — a hypothetical, since 140 GB does not fit in one 80 GB H100 — that ceiling is about tokens/second for a single request — an order-of-magnitude figure that ignores KV-cache reads and overheads, and one that batching does not raise for any single stream even as it multiplies the GPU’s aggregate throughput across many streams ^[428][429]. These three numbers — weight footprint, KV-cache growth, and the bandwidth-bound token rate — are the whole reason memory, not raw FLOPs, dominates the inference story this primer keeps returning to. Try it: Decode ceiling: ≈24 tok/s; Single-stream decode ceiling.

A closing caution on how to read any of the performance figures in this primer. A chip’s headline "TFLOPS" is a peak — the throughput if every arithmetic unit ran flat out with no stalls — and real workloads never reach it. The standard measure of what a training run actually achieves is Model FLOPs Utilization (MFU): the fraction of a chip’s peak FLOPs spent on the useful forward-and-backward math of the model. Even very well-tuned frontier runs land around 30–50% MFU — Google’s PaLM reported 46.2%, and Meta’s Llama 3 405B reported 38–43% on its H100 clusters — with the rest lost to memory stalls, inter-GPU communication, non-matmul operations (normalization, softmax, activations) that don’t hit tensor-core peak, and pipeline bubbles ^[430][431]. Try it: PaLM MFU vs HFU.

NVIDIA’s own superchips make this division explicit rather than incidental. The Grace CPU packs 72 Arm Neoverse V2 cores connected by NVIDIA’s own Scalable Coherency Fabric (3.2 TB/s of internal bisection bandwidth), and uses server-class LPDDR5X memory — unusual for a datacenter CPU — to get up to 500 GB/s of host memory bandwidth at roughly a fifth of the power per bit of conventional DDR memory. In the GB200 superchip, Grace connects to its Blackwell GPU partners over NVLink-C2C, a cache-coherent chip-to-chip interconnect at 900 GB/s — not PCIe. Its job, per NVIDIA’s own technical material, is to handle data loading and preprocessing, kernel launch, orchestration, and — in the successor Vera CPU — increasingly agentic workloads: driving tool calls, sandboxed code execution, and even running the small speculative-decoding "draft" models described in Part IV, so the GPU’s cycles go almost entirely to the matrix math it’s uniquely built for ^[27][28][29]. Try it: NVLink-C2C vs PCIe.

Three concrete reasons explain why an ARM core, not an x86 core, ends up glued to the GPU in these designs. First, power: in a rack with a fixed power and cooling budget, every watt spent on the host CPU is a watt not available to the accelerators actually doing the expensive work, and ARM’s Neoverse cores deliver the "keep the GPU fed" work (data movement, orchestration, light control-flow-heavy code) at meaningfully lower power per unit of that work than a comparable x86 server chip. Second, coherent, tight coupling: NVLink-C2C gives the CPU and GPU a shared, cache-coherent address space at nearly a terabyte per second, letting the two chips exchange data and hand off work with far less latency and software overhead than a PCIe-attached host CPU would ever manage — this only makes sense with a chip NVIDIA can co-package this tightly, which is easier to do with a licensable core architecture like ARM’s than with x86. Third, density: many efficient cores, rather than fewer powerful ones, suit the actual job (moving lots of independent small chunks of data and orchestrating lots of independent small requests) far better than a CPU built to run a single fast thread.

In March 2026, Arm Holdings announced its own self-designed data center processor, developed with Meta as lead partner — its first move from licensing chip designs to shipping its own silicon. The chip packs up to 136 Neoverse V3 cores at 3.5–3.7 GHz, built as a chiplet design on TSMC’s 3nm N3P process, with a 300W TDP and dense per-rack core counts (over 8,000 cores per rack air-cooled, 45,000+ liquid-cooled). Crucially, Arm’s own marketing frames this explicitly as an "agentic AI infrastructure" CPU — meaning it is built to run the orchestration layer of AI systems (managing many parallel agent loops, tool calls, and requests) at scale, not to compete with GPUs or TPUs on the tensor math itself. Even ARM’s own dedicated AI silicon reinforces the front-end/back-end split rather than breaking it ^[30][31][32].

It is easy to read the Grace story and conclude that ARM has taken over the AI datacenter. It has not — not as the host. The large majority of GPU and accelerator servers in the world are still hosted by x86 CPUs, from Intel (Xeon) or AMD (EPYC), and NVIDIA’s own reference platforms make the point: the DGX H100 and H200 are hosted by dual Intel Xeon (Sapphire Rapids) CPUs, and — tellingly — the Blackwell-generation DGX B200 8-GPU box stayed on x86 (dual Emerald Rapids Xeons) rather than moving to ARM, while OEM HGX B200 baseboards ship with a customer’s choice of EPYC or Xeon hosts. ARM (Grace) appears specifically in NVIDIA’s superchip products — GB200 and GB300 — where the CPU and GPU are co-packaged on one coherent module over NVLink-C2C; that is a deliberate, tightly-coupled design point, not the default configuration of the broad install base ^[382][383].

The reasons x86 remains the incumbent host are the mirror image of ARM’s advantages. A big server x86 CPU brings an enormous, mature software ecosystem, a vast installed base, and — critically for a host — a lot of I/O: AMD’s EPYC "Turin" (Zen 5) offers up to 192 cores, 12 channels of DDR5, and up to 128–160 PCIe Gen5 lanes, which is exactly what a GPU host needs to fan out to many accelerators and NICs at once, with the PCIe-attached flexibility to mix and match GPUs from different vendors. AMD explicitly markets certain EPYC parts (e.g. the high-frequency 9575F) as purpose-built GPU-host CPUs, and EPYC is the natural host for AMD’s own Instinct systems (Part XI). x86 CPUs can also do meaningful inference themselves: Intel’s AMX (Advanced Matrix Extensions) is an on-die tile-based matrix unit in recent Xeons (Sapphire Rapids onward, with FP16 added in Xeon 6 "Granite Rapids") that lets the CPU run smaller models without any accelerator at all. The honest summary is that x86 is the default host that feeds most of the world’s GPUs, while ARM is winning the specific, high-value design point where CPU–GPU memory coherence and power efficiency matter most. (Reported market-share figures should be read carefully: ARM is often quoted at a low double-digit percentage of server-CPU units but a much higher share of datacenter revenue, because a Grace superchip costs far more than a commodity host CPU — the two figures measure very different things) ^{[384][385][386]}.

The third camp is RISC-V, an open, royalty-free, extensible ISA governed by RISC-V International — architecturally a peer of ARM and x86, but owned by no one, which is the source of both its appeal and its geopolitics. It is important to be precise about RISC-V’s real role in AI today, because it is easy to overstate: RISC-V is not yet a mainstream datacenter host CPU competing with Xeon, EPYC, or Grace for the GPU-feeding job. Its dominant, shipping role is something quieter and pervasive — control and management cores embedded inside the accelerators themselves. NVIDIA has said it shipped over a billion RISC-V cores in 2024 and runs 30-plus distinct system-management applications on RISC-V cores inside its own GPUs (its on-GPU management processor, the GSP, is RISC-V), having migrated off its proprietary "Falcon" microcontroller starting in 2016; Meta’s MTIA inference accelerators (Part VI) likewise use RISC-V cores as their processing-element and control-core silicon. The chip you think of as an "NVIDIA GPU" or a "Meta ASIC" already contains dozens of RISC-V cores doing the housekeeping ^[387][388].

Beyond embedded control, RISC-V’s larger ambitions are real but still early. Tenstorrent (led by Jim Keller) pairs licensable RISC-V CPU cores (Ascalon) with its AI cores; NVIDIA announced in 2025 that it is bringing CUDA host support to RISC-V (extending the host beyond x86 and ARM), though that is work in progress with no shipping timeline and depends on future datacenter-class RISC-V CPUs that do not yet exist. The most consequential near-term story, though, is strategic: because RISC-V is an open standard no government can revoke a license to, it has become a deliberate hedge against export controls — most prominently in China, where firms such as Alibaba’s T-Head (the XuanTie line) are building increasingly capable server-grade RISC-V chips explicitly to reduce reliance on Western ARM and x86 designs. That has drawn US policy attention: in late 2023 a group of lawmakers urged the Commerce Department to examine export-control options around RISC-V collaboration with China, and in 2024 Commerce confirmed it was assessing the national-security implications — an unresolved tension that connects directly to the export-control dynamics of Part XX. The through-line is that RISC-V matters to AI hardware today less as a CPU you can buy to host a GPU and more as the open, ubiquitous control-plane silicon inside chips and as a geopolitical fault line in who gets to build advanced compute at all ^{[389][390][391]}.

The one place this split doesn’t apply is the edge, where the "host" and the "accelerator" are the same tiny power and thermal budget. Apple’s Neural Engine and Qualcomm’s Hexagon NPU are ARM-ecosystem silicon that genuinely execute matrix multiplications for on-device models — but it is worth being precise about which models run where, because this is a commonly conflated point: fixed-function NPUs like the ANE are built for vision, speech, and other CNN/small-model workloads, while today’s on-device LLMs (Phi-3, Llama 3.2, Gemma 2, Mistral Small, and Apple’s own on-device Foundation Model) actually run on the GPU, not the NPU, on essentially every platform that ships them. Part VI’s NPU and GPU sections go deep on exactly why — the short version is that autoregressive decode is memory-bandwidth-bound in a way that plays to a programmable GPU’s strengths (and a closed, fixed-function NPU’s weaknesses) far more than it plays to raw TOPS ^[33][34].

Apple has shipped a dedicated Neural Engine in every A-series and M-series chip since the A11 Bionic (2017), and for most of that run published an exact peak-throughput figure at each launch:

Two things about this table matter more than the raw numbers. First, Apple published an exact ANE TOPS figure in every phone-chip Newsroom release from the A11 through the A16 — then stopped: A17 Pro, A18, and A18 Pro get only relative claims ("2x faster than A16"), even though M-series chips kept getting exact figures through the M4. The widely-repeated "35 TOPS" figure for A17 Pro/A18 is tech-press extrapolation, not an Apple statement, and shouldn’t be treated as confirmed. Second, the Apple Watch S9 and S10 chips ship a separate, much smaller 4-core Neural Engine (no TOPS figure published for either), used for on-device Siri, the double-tap gesture, dictation, workout/crash detection, and — on S10 — the watchOS Translate app and a voice-isolation network for calls ^{[93][94][102][103]}.

Apple publishes essentially no microarchitecture documentation for the ANE — no public ISA, no programming model outside the black-box Core ML compiler, and no register-level specification. The most rigorous available source is an independent reverse-engineering effort (built from measurement plus static decompilation of Apple’s private compiler, driver, and firmware) covering the A11 through A18 and M1 through M5. Its central, well-evidenced finding is that the ANE is a fixed-function fp16 matrix accelerator with a wide, higher-precision accumulator, fed by streaming DMA — not confirmed to be a systolic array of the TPU/MXU kind covered in Part I, despite that comparison being commonly assumed in secondary coverage. The datapath multiplies in fp16 end to end: it accepts fp32, int32, and bf16 type annotations at the compiler frontend, but the hardware backend only actually implements fp16 arithmetic. INT8/INT4 formats do exist, but only as compressed weight-storage formats that get dequantized to fp16 before the multiply — a materially different role than the native low-precision compute paths described for GPUs and TPUs in Parts I and II, and worth flagging as a common point of confusion when reading ANE "INT8 TOPS" claims ^[104].

The same reverse-engineering work identifies a hard, roughly 2MB on-chip working-set limit: a kernel whose live data tiles fit inside that budget runs at full fixed-function speed, while one that exceeds it stalls waiting on off-chip streaming from the shared unified memory pool. Apple’s own 2022 transformer-deployment paper corroborates this independently, noting that ANE buffers must be contiguous and 64-byte aligned, and that mis-shaped tensors reduce the chance of on-chip residency and increase the chance of falling through to DRAM. Measured on an M1, the ANE delivers roughly 12 fp16 TFLOP/s of compute against the M1’s roughly 68 GB/s of DRAM bandwidth — a very different, far more compute-heavy roofline ridge point than the memory-bound GPU regime that dominates LLM decode ^[104][105]. Try it: The Neural Engine’s roofline.

Across nearly a decade of Apple Newsroom and Apple ML Research material, the confirmed, shipping ANE workload list is consistent: Face ID and TrueDepth authentication, computational photography (Smart HDR, Deep Fusion, Night Mode, portrait segmentation), on-device Core ML models generally (exposed to third-party developers since the A12), whole-Photos-library scene analysis, camera panoptic segmentation via a transformer Apple calls HyperDETR, and — more recently — Live Translation on AirPods. Every one of these is either a convolutional vision model or a small, fixed-shape model whose entire compute graph is known and compiled once ahead of time — exactly the profile that plays to a closed, fixed-function, fp16 accelerator with a small on-chip working set ^{[88][89][92][106][107][108]}.

Why, then, does on-device LLM inference bypass the ANE despite its efficiency on other workloads? The explanation comes primarily from Apple’s own applied-research publications rather than outside speculation. Apple’s 2022 transformer paper states plainly that "many Transformer configurations become bandwidth-bound on the ANE when the sequence length is relatively short… large parameter tensors are being fetched from memory, only to be applied on too few inputs before the next parameter tensor is fetched" — precisely the access pattern of one autoregressive decode step, where sequence length is effectively 1. Apple’s 2024 Core ML LLM deployment paper is even more direct about the consequence: describing an on-device Llama-3.1-8B deployment, Apple writes that they "specifically target the GPU, as the models… are usually constrained by memory bandwidth, and the GPU offers the best combination of compute FLOPS and memory bandwidth on the device." Apple’s own MLX framework — the toolkit Apple’s machine learning research group builds and publishes for exactly this kind of on-device model work — supports only CPU and GPU as compute devices; MLX’s lead maintainer has stated on the project’s public issue tracker that ANE support isn’t planned in part because it’s a closed API with no public interface to target, and community contributors note the ANE has no backpropagation support and computes natively only in 16-bit float ^{[105][109][110]}.

Independent developers hit the same wall from the outside: llama.cpp maintainers describe Apple’s NPU as fast for convolution but without "enough speed for transformers," and note that even Apple’s newest on-device APIs that nominally expose the ANE to developers don’t reliably place transformer workloads there in practice — Core ML’s compiler makes an opaque, take-it-or-leave-it placement decision across CPU/GPU/ANE that developers can’t override. None of this means the ANE is weak in absolute terms: the same reverse-engineering study that documents its limits also measures it at roughly 3.8x the GPU’s speed and 9x its energy efficiency on genuinely compute-bound convolutional work — the ANE is simply built for a different point on the compute-intensity spectrum than autoregressive decode occupies ^[104][111].

Apple’s GPU architecture is what’s called a Tile-Based Deferred Renderer (TBDR): rather than processing an entire frame’s geometry and pixels in one pass the way a traditional "immediate-mode" desktop GPU does, it splits work into two phases — first tiling, where all the geometry for a small screen tile is processed and binned, then rendering, where that tile’s pixels are shaded using a small, dedicated pool of on-chip tile memory, with a fully deferred pixel write-back only once a tile is complete. Apple’s own engineering material describes this directly: "the GPU has a dedicated pool of on-chip memory, which we call Tile Memory… the GPU does not have video memory" in the way a discrete GPU does — everything ultimately lives in the same unified memory pool the CPU and ANE share. A Chips and Cheese teardown of the M2 Pro’s 19-core GPU found a two-level on-chip cache hierarchy (an 8KB L1 and a 3MB L2 delivering over 1 TB/s to the shader cores) sitting in front of a 256-bit LPDDR5 bus ^[113][114].

Starting with the M3 family and A17 Pro (2023), Apple added "Dynamic Caching": rather than statically partitioning on-chip memory ahead of time the way earlier GPU generations (and most competitors) do, the hardware allocates exactly the amount of local memory a given task needs, in real time, which Apple states "dramatically increases the average utilization of the GPU." The same M3 generation added hardware-accelerated ray tracing and mesh shading — the first time either had shipped in Apple Silicon — and M4 roughly doubled ray-tracing throughput again ^[99][101].

Separately from both the ANE and the GPU, Apple has for years shipped an undocumented CPU-adjacent matrix coprocessor the community calls AMX (Apple Matrix Coprocessor), reachable through Apple’s Accelerate framework (BLAS, BNNS, vDSP, LAPACK) rather than any public instruction Apple has ever documented — everything known about it comes from independent reverse-engineering, not Apple disclosure. Community researchers have speculated that AMX is being phased out in favor of ARM’s standardized Scalable Matrix Extension (SME, part of the Armv9-A architecture Apple adopted starting with the M4), since SME accelerates the same class of matrix workloads through an official, public ISA extension — but this is inference from circumstantial evidence (M4’s Armv9/SVE adoption, community SME-exploration projects), not a statement Apple has made, and the primer flags it here as unconfirmed rather than settled fact ^{[115][116][117]}.

Whichever way the CPU-side AMX/SME question resolves, it is architecturally separate from where LLM inference actually runs: the GPU has had its own matrix-multiply instructions in the Metal Shading Language (simdgroup_matrix) since the A14 generation, and the M5 generation (2025) went further, adding dedicated "Neural Accelerator" matmul units directly inside each GPU core, exposed through Metal 4 TensorOps — and it is this GPU-side matrix hardware, not AMX and not SME, that Apple’s own MLX framework and llama.cpp’s Metal backend actually target for LLM inference ^[118][135].

Why this matters comes down to the same memory-bandwidth story that has run through this entire primer, now confirmed in Apple’s own words at consumer-hardware scale. Apple’s Newsroom material gives exact unified-memory bandwidth figures by chip: M2 Pro at 200GB/s, M2 Max at 400GB/s, M2 Ultra at 800GB/s, M4 at 120GB/s, M4 Pro at 273GB/s, and M4 Max at up to 546GB/s — a figure Apple explicitly ties to LLM capability, stating it lets developers "easily interact with large language models that have nearly 200 billion parameters." Apple’s own 2025 MLX research post makes the mechanism explicit and quantitative: "generating subsequent tokens is bounded by memory bandwidth, rather than by compute ability," reporting that the M5’s jump to 153GB/s of bandwidth (from the M4’s 120GB/s) directly produced a 19–27% decode speedup — a first-party confirmation of exactly the prefill-compute-bound-vs-decode-bandwidth-bound framing from Part IV, now measured on a laptop chip rather than a datacenter GPU ^{[97][98][100][101][118]}. Try it: Apple M1–M5 bandwidth table; M5 vs M4 decode speed.

ARM’s flagship GPU tier is branded "Immortalis": the Immortalis-G715 (2022) was ARM’s first mobile GPU with hardware ray-tracing support; the Immortalis-G720 (2023) added a second-generation ray-tracing engine and shipped in MediaTek’s Dimensity 9300 with a 46% GPU performance and ray-tracing uplift over the prior generation; the Immortalis-G925 (2024) scaled up to 24 shader cores (50% more than the previous 16-core maximum) with 37% better performance at 30% less power. ARM’s own developer documentation describes its rendering scheme as "tile-based rendering," explicitly contrasted with "immediate-mode" GPUs — geometry is binned into small tiles and fragment-shaded in a local tile buffer before being written back to memory, a scheme functionally close to Apple’s TBDR, though ARM’s branding notably avoids the word "deferred" that Imagination (PowerVR, the ancestor of Apple’s own GPU lineage) and Apple both use — a real terminology distinction worth noting even though the underlying tile-then-shade mechanism is closely related ^{[119][120][121][122]}.

ARM’s Ethos-U series targets microcontroller-class edge devices paired with Cortex-M cores: the Ethos-U55 (2020) delivers roughly 25.6–512 GOPS depending on configuration and process node; the Ethos-U85 (2024, third generation) scales from 256 GOPS to 4 TOPS and is the first Ethos-U to natively support transformer operators (transpose, gather, matmul) rather than convolution-only workloads. A separate, higher-end Ethos-N78 line targets mobile-class SoCs with a stated scalability range of "1 to 10 TOPS and beyond," though real flagship design wins for Ethos-N78 specifically are genuinely scarce in the public record — widely-repeated claims that chips like MediaTek’s Dimensity 9000 or Samsung’s Exynos 2200 use it directly conflict with those vendors' own confirmed in-house NPU branding, suggesting most major mobile vendors build their own NPU rather than license ARM’s mobile-class design, even when they license ARM’s GPU ^{[123][124][125][126]}.

Real, confirmed licensees give a clearer picture than marketing claims: MediaTek ships Immortalis-G720 in the Dimensity 9300; NXP and Alif Semiconductor both ship Ethos-U55/U85-based microcontrollers for IoT and TinyML use cases, with Alif shipping the first Ethos-U85 silicon. Google’s Tensor chips (covered next) licensed ARM Mali GPUs for four full generations — a useful contrast case, since Google, like Apple, holds only an ARM CPU architecture license, but unlike Apple chose to license ARM’s GPU IP outright rather than design its own ^{[120][127][128][129]}.

Every one of the WSE’s cores (Cerebras calls them PEs, processing elements) is connected to its four cardinal neighbors — north, south, east, west — by a small, dedicated router, forming a strict 2D mesh with no diagonal links and no long-range shortcuts. Each link moves a 32-bit "wavelet" per cycle in each direction (16 bits of data, 16 bits of index, plus routing metadata), and — critically — a wavelet takes exactly one clock cycle to cross from one router to the next, regardless of where on the wafer that hop occurs. Routing is not dynamically arbitrated at runtime the way an NVSwitch or InfiniBand fabric is; instead, each wavelet carries a 5-bit "color" tag, one of 24 statically-configured virtual channels that the Cerebras compiler assigns at compile time, and that color determines both the path the wavelet takes through the mesh and which task on the receiving core consumes it. That static, compiler-scheduled routing is what lets the fabric guarantee single-cycle, non-blocking hops with very little buffering — the same "no runtime surprises" philosophy Part VI’s section on Groq describes, applied here to on-chip communication instead of instruction scheduling ^[72][73][74].

This has a direct architectural consequence worth naming explicitly: execution on the WSE is fundamentally dataflow-triggered rather than instruction-fetch-driven. A core doesn’t poll for work — the arrival of a wavelet on a given color is itself the event that triggers a task to run, with the wavelet’s data and index feeding directly into that task. Cerebras describes the entire fabric as "transporting both the data and the associative control directly in the hardware," which is a meaningfully different execution model from a GPU’s warp scheduler or a CPU’s instruction pipeline ^[78].

Aggregated across the whole wafer, this mesh provides roughly 214–220 petabits per second of core-to-core bandwidth (WSE-2 and WSE-3 report slightly different totals despite WSE-3 having more cores — a real, source-confirmed quirk of the two generations' fabric configuration, not an error). For comparison, a single hop between any two neighboring cores anywhere on the wafer completes in one clock cycle, and Cerebras engineers estimate a full-wafer all-reduce can complete in on the order of a single microsecond — versus multiple hops and microseconds of latency for the scale-up NVLink fabric described in Part III, though the comparison is apples-to-oranges since the WSE’s mesh has no analog to leaving the chip at all ^[73][87].

A standard photolithography reticle can pattern at most about 830–860 mm² of silicon in one exposure — the hard ceiling every conventional chip design lives under. The WSE sidesteps this not by using a bigger reticle (none exists) but by tiling 84 identical reticle-sized dies across the wafer in a 12×7 grid and then, in partnership with TSMC, adding a second, offset lithography exposure whose sole purpose is to lay extra metal wires across the "scribe lines" — the narrow strips between adjacent dies that are normally reserved as a mechanical saw path for cutting wafers into individual chips. Because Cerebras never dices the wafer, those scribe lines instead become another place to route wires, and the extra exposure connects each die’s mesh routers to the matching routers on its neighboring dies, extending the identical 2D mesh across what would otherwise be a hard electrical boundary. Cerebras’s own description of the result is that the "same connectivity between cores and across scribe lines create[s] a homogenous array" — every die is manufactured from the same mask set and lands in the same position relative to its neighbors, so a core’s fabric ports look identical whether its neighbor is on the same die or the adjacent one, and independent academic analysis confirms cores on either side of a scribe line get the same bandwidth as two cores on the same die ^{[75][76][81][84]}.

This is a genuinely hard, previously-failed engineering problem: Trilogy Systems attempted a similar cross-die wafer-scale interconnect in 1980 under Gene Amdahl and could not make it work at production yields. Cerebras’s approach is protected by (and documented in) US Patent 10,777,532, "Apparatus and method for multi-die interconnection," which describes two complementary photoreticle exposures — one for each die, one for the wiring that bridges them — whose patterns overlap at the boundary to form a continuous electrical connection ^[81][84].

Every wafer, however carefully manufactured, contains some density of random defects; a normal chipmaker’s answer is to make dies small enough that a typical wafer yields mostly-good dies and simply bins or discards the rest. That option doesn’t exist when the entire wafer is the product. Cerebras’s fix is to build a modest amount of spare capacity directly into the design: WSE-3 ships with roughly 970,000 physical cores on every wafer but activates only about 900,000 of them — Cerebras states plainly that this means "we enable 93% of our silicon area," with the remaining ~7% held in reserve specifically to route around whatever defects that particular wafer happens to have. Earlier public statements about WSE-1 described a smaller reserved margin, on the order of 1–1.5% extra cores — the two figures reflect different accounting (a conservative sparing budget vs. the realized redundancy actually exercised across a real defect-density distribution), not a contradiction to resolve ^[79][80].

Mechanically, when a defect is identified — during manufacturing test, and in principle during field operation — the compiler’s place-and-route step simply avoids assigning any workload to that specific core, and the mesh’s statically-configured routing paths are laid out around it rather than through it, using nearby spare cores and spare fabric links to preserve full mesh connectivity for everything else. Cerebras describes this as "a sophisticated routing architecture that allows us to dynamically reconfigure connections between cores," backed at the die-to-die link level by "auto-correction state machines that automatically use redundant links when errors are discovered by training and error detection." Public sources do not confirm any specific fuse-based (laser or electrical) disable mechanism, nor a literal dedicated row/column layout for spares — the verified description is that sparing is fine-grained and local to wherever a defect actually falls, not organized into a fixed redundant block ^[79][78].

Each PE carries 48KB of local SRAM (eight 6KB banks, supporting two 64-bit reads and one 64-bit write per cycle, plus a small 256-byte software-managed cache) — a figure that has stayed constant from WSE-2 to WSE-3; the jump from 40GB to 44GB of total on-wafer memory comes entirely from the increase in core count, not from any change to each core’s own memory. Alongside its fabric router, each core has a conventional 6-stage scalar pipeline with 16 general-purpose registers for ordinary control flow, four FP16 fused multiply-accumulate units (8 FLOPs per cycle), and hardware support for 8 simultaneous "microthreads" — independent tensor computation contexts the core can switch between cycle by cycle as wavelets for different tasks arrive ^[74][78].

The fine-grained, unstructured sparsity Cerebras advertises is a direct hardware consequence of the wavelet model rather than a separate feature bolted on top: because a core’s fabric only ever transmits and receives explicit, non-zero wavelets, a zero-valued weight or activation simply never gets sent and never triggers a downstream multiply-accumulate — there is no wasted computation to skip, because the computation for a zero never gets scheduled onto the fabric in the first place. Cerebras describes this as filtering "out all the zero data at the sender, so the receiver doesn’t even see it," and reports roughly a 10x effective utilization gain from this technique versus a dense GPU baseline on models with substantial weight sparsity — the mechanism works at the level of individual scalar values, not fixed 2:4-style structured blocks like the sparsity techniques described in Part II ^[78].

A wafer-scale die creates a power-delivery problem no conventional chip package has to solve: 900,000 active cores drawing current across a roughly 215mm silicon surface add up to more than 20,000 amperes at sub-1-volt core voltages, and delivering that current from the edge of the package — the way a normal chip’s power pins work — would require impractically thick PCB copper layers to avoid ruinous resistive (IR) voltage drop over that distance. Cerebras’s solution, engineered specifically for the WSE and covered by multiple patents, is to deliver power vertically, straight up into the face of the wafer from directly underneath it, rather than laterally from the side.

The physical stack, top to bottom, is: a cold plate carrying a closed water loop pressed against the top of the wafer to remove heat; the wafer itself; a custom elastomeric connector — a silicon-rubber membrane embedded with conductive particles that conducts current only in the vertical direction and, held under several tons of compression, mechanically absorbs the mismatch in thermal expansion between silicon and the PCB beneath it; and a power-delivery PCB carrying over 300 voltage regulator modules distributed across the wafer’s footprint, one cluster per die region, pushing current straight up through the elastomeric connector into the wafer above. Current flows in from the bottom of this "sandwich" exactly as heat flows out the top, and Cerebras’s own description of the design goal is direct: "there is no significant difference in temperature across the surface of the wafer, and all parts of it are effectively and uniformly cooled." The elastomeric connector exists specifically because no rigid PCB material can be manufactured with a thermal expansion coefficient close enough to silicon’s to survive being permanently bonded to a die this large — the compliant connector absorbs that differential expansion instead of fighting it ^[77][82][83].

Cerebras instruments this power and thermal system heavily rather than assuming it behaves uniformly by design alone: WSE-2 alone carries 672 temperature sensors and 10,752 individual voltage sense points spread across its 84 dies specifically to catch localized IR-drop sags or thermal excursions during the bursty, uneven power draw that real AI workloads produce moment to moment — a sharp contrast to a conventional GPU, which needs only a handful of sensors because its die is small enough that power and heat distribute far more evenly on their own ^[86].

At Hot Chips 2026 Cerebras laid out how the company is scaling its wafer-scale platform past the single-system CS-3 into a full rack-scale architecture, and sketched two further steps beyond that. The through-line is that on-wafer SRAM alone has run out of room for frontier-scale models whose weights now measure in the multi-trillion-parameter class, and every new system pushes on a different axis — first bandwidth-per-wafer, then racks-of-wafers with sub-microsecond wafer-to-wafer interconnect, then a third dimension entirely.

CS-4 and the "Nexus" backpack (early access 2026, GA Q3 2026). CS-4 is the first Cerebras system designed as a rack building block rather than a freestanding appliance. Each rack holds three WSE-3 Turbo (WSE-3T) processors — the same 4-trillion-transistor wafer as WSE-3, but with doubled clock speed relative to the WSE-2 baseline, 44 GB of on-chip SRAM per wafer, 750 PFLOPS of AI compute, and 53.5 PB/s of aggregate on-wafer fabric bandwidth — a figure Cerebras notes is roughly 200× the ~260 TB/s of NVLink bandwidth inside an NVIDIA GB200 NVL72 rack, and delivered without the roughly 5,000 internal cables an NVL72 needs to reach that number ^[3]. Across the three wafers a CS-4 rack delivers 129.6 PB/s of aggregate memory bandwidth and 7.2 Tb/s of system I/O; the wafer-to-wafer latency across the internal fabric is ~2 μs, low enough that a three-wafer pipeline behaves much like one larger wafer for the purposes of tensor- and pipeline-parallel decomposition ^[4]. Cerebras publishes CS-4-versus-CS-3 numbers of 2× the per-user token generation speed, 10× the tokens per watt, and — the marketing headline — up to 30× faster than "existing production GPUs" on the same inference workloads ^[5].

The Nexus is the rack-level engineering that makes this practical, and it is where the physical-plant story from Section on the wafer power stack picks up again. Each WSE sits in what Cerebras calls a compute backpack: a modular chassis carrying its own power, cooling, and I/O, built with roughly 50% fewer components than the CS-3 assembly and about 60% automated in manufacture. The DC/DC power converters have moved to within 0.5 mm of the wafer itself — an aggressive shortening of the copper path that Cerebras claims reduces board-level power loss by about 100× versus a conventional GPU carrier board. Each backpack accepts up to 30 hot-swappable PSU modules feeding from 277 V AC, each protected by a 30 A breaker, feeding self-contained liquid loops with automatic safety cutoffs so that a failure of any single backpack — power, cooling, or the wafer itself — is isolated to that one slot rather than taking down the rack ^[86][6].

CS-5 (targeted 2027). CS-5 keeps the WSE-3T-class silicon and Nexus backpack but reorganizes the rack around interactive inference throughput. Cerebras’s public targets are 10,000 tokens per second per user on mid-scale open models (Gemma 4 31B, gpt-oss 120B) and 5,000 tokens per second per user on frontier trillion-parameter models such as DeepSeek, Kimi, and GPT-5.6 Sol, at an efficiency of roughly 3 million tokens per second per megawatt — an order-of-magnitude jump in useful throughput per rack relative to CS-4 that comes almost entirely from software (a decode pipeline redesigned around the wafer’s SRAM-only working set) rather than new silicon ^[7].

CS-6 (R&D, no public GA date). CS-6 is the one architectural jump on the roadmap. The wafer becomes the substrate for a 3D-stacked memory sandwich: on-wafer SRAM and compute on the bottom layer, with wafer-scale SRAM stacked on top through ultra-high-bandwidth vertical connections, and 3D-stacked DRAM integrated into the same package above that. The stated design goal is to keep the sub-nanosecond, uniform-latency access pattern that makes wafer-scale fast today, while lifting the fixed-capacity ceiling that has been the wafer’s biggest weakness — a WSE today can only hold what fits in its 44 GB of SRAM, so frontier models beyond that size have to stream weights in from MemoryX. Cerebras states plainly that CS-6 is aimed at models with more than 50 trillion parameters at interactive speeds, and that the target system footprint is "an order of magnitude smaller" than CS-4 for equivalent performance. Vertical power delivery moves from "0.5 mm away" (CS-4) to integrated directly into the wafer stack, and cooling likewise becomes part of the 3D package rather than a separate backpack layer ^[8][9].

The roadmap read is that Cerebras is betting the wafer-scale approach wins by climbing three orthogonal axes in sequence: WSE-3T pushed the per-wafer compute and bandwidth (CS-3 → CS-4); Nexus turned it into a rack-scale platform with wafer-to-wafer latency competitive with a single GPU’s own NVLink domain (CS-4 → CS-5); and 3D memory integration is the attempt to remove the single hardest structural constraint — fixed SRAM capacity — that has kept wafers a niche for frontier-scale training and inference (CS-5 → CS-6). If any of the three misses — if the 3D stack’s yield is unmanageable, if CUDA-class software ecosystems keep growing on GPU racks while Cerebras’s own stack does not, or if hyperscalers commit to their own captive silicon for the same workloads — the wafer stays a specialized bet. If all three hit, wafer-scale is the first credible alternative to the GPU-rack template for frontier inference specifically.

Groq’s original ISCA 2020 paper describes a "functionally-sliced microarchitecture": instead of tiling many identical general-purpose cores across the die (the GPU/CPU approach — local heterogeneity, global homogeneity), the TSP implements each function once and repeats it as a horizontal slice — MEM (memory), VXM (Vector Execution Module), MXM (Matrix Execution Module), SXM (Switch Execution Module, for data permutation and transpose), and an Instruction Control Unit (ICU) — stacked so that "local functional homogeneity" combines with "chip-wide heterogeneity." The die is organized as two mirrored hemispheres with VXM at the center and MEM/SXM/MXM arranged outward on each side. The real, paper-confirmed counts are substantial: 88 MEM slices (44 per hemisphere) providing 220 MiB (≈230MB decimal) of on-chip SRAM with 176-way bank concurrency; 4 independent 320×320 matrix-multiply-accumulate planes in the MXM, for 409,600 MAC units total, native INT8 with FP16 available via paired byte-planes; 5,120 vector ALUs in the VXM (320 lanes × 16 ALUs/lane); and 144 independent ICUs — one per functional-slice column — occupying under 3% of total die area ^[137][139].

Data moves east-west across the chip in fixed-width "streams" routed through the MEM slices, while instructions flow north-south from each column’s ICU down through that slice’s stack of tiles — a 20-tile pipeline per slice, each tile handling one 16-lane "superlane." Twenty of these superlanes stacked together give the chip’s 320-lane maximum vector width; Hot Chips 34 discloses a 21st, physically-present spare superlane, invisible to the compiler and reserved purely for yield/defect recovery — Groq’s own, much smaller-scale analog to the fail-in-place redundancy Cerebras builds into every wafer. The shipped GroqCard runs this floorplan at 900MHz (the ISCA 2020 paper’s abstract also cites a theoretical 1GHz/820 TOPS design point, but the production spec sheet, the 2022 paper’s own evaluation, and Hot Chips 34 all converge on 900MHz / 750 TOPS INT8 / 188 TFLOPS FP16 as the real, shipped numbers) on a 725mm² die, GlobalFoundries 14nm, roughly 27–29 billion transistors ^{[137][138][139][140]}.

The mechanism that makes this deterministic is what the ISCA 2020 paper formalizes as a timing model with two explicit terms: d_func (the fixed functional delay of a tile) and d_skew (a compiler-computed launch offset for each operand stream). A single 320-wide SIMD instruction issued by one ICU at cycle propagates one tile per cycle down through a slice’s 20-tile pipeline; the compiler pre-computes exactly how many cycles earlier each operand stream must be launched so that data and instruction arrive at the correct tile on the correct cycle, every time, with no exceptions. Hot Chips 34 puts this plainly: the whole chip "can be viewed as fully-pipelined 144-wide VLIW instructions," executing in lockstep with no hardware queue, arbiter, or reorder logic anywhere on the die ^[137][139].

This is also the mechanism behind Groq’s core marketing claim of "deterministic" execution. Groq’s own whitepaper states it directly: "a GroqChip has no control flows, no hardware interlocks, no reactive components like arbiters or replay mechanisms… the compiler schedules all events happening at different times on different functional units… and the hardware simply executes a predetermined script." Because there is no dynamic branch predictor, no cache-miss stall, and no runtime arbiter anywhere on-chip, the same compiled program takes exactly the same number of cycles on every run — which Groq leverages directly to guarantee consistent, reproducible per-token latency for LLM serving, in contrast to GPU-based serving, where dynamic batching, kernel-launch overhead, and cache effects make per-token latency vary from request to request ^{[141][142][143]}.

A single TSP has no HBM and only ~230MB of SRAM, so any model larger than that has to span many chips — and Groq’s answer, "RealScale," is designed to extend the same determinism guarantee across an entire multi-chip system rather than sacrificing it at cluster scale. RealScale explicitly eliminates NICs and any separate switch chip: each TSP has 11 chip-to-chip links — 7 "local" links forming an all-to-all mesh within an 8-chip GroqNode (a dragonfly topology at this level), plus 4 "global" links used for inter-node traffic. Nine GroqNodes combine into a GroqRack: 72 chips total, confirmed by both Groq’s own product briefs and an independent, real-world deployment at Argonne National Laboratory’s AI Testbed, which describes the rack-level network specifically as a "rotational" topology. (Groq’s own documentation is not fully internally consistent on the exact chip-per-rack count — a separate scaling table implies 64, and yet another reference implies 66 — but 72 is the number confirmed by both Groq’s product brief and Argonne’s actually-deployed hardware, so it’s the figure to trust.) The largest published RealScale configuration reaches 145 racks and 10,440 chips, with network diameter held to 5 hops or fewer ^{[144][145][146]}.

Extending determinism across chips whose clocks are never perfectly synchronized requires specific ISA-level machinery, not just a fast network: Hardware and Software Aligned Counters (HAC/SAC) exchange and re-derive a shared notion of time every 256 cycles, and dedicated instructions (SYNC, NOTIFY, DESKEW, RUNTIME_DESKEW) perform initial alignment and ongoing drift correction — Groq’s own term for the result is "plesiosynchronous" (near-synchronous, bounded-drift) alignment, honest about the fact that the underlying clocks are not literally one clock. The SRAM math is the sharpest illustration of why this matters: 72 chips × ~230MB is only about 16.6 GB of total on-chip memory per rack — and a peer-reviewed evaluation running real hardware at Argonne found that a full 72-chip GroqRack is "nearly the largest model that can be fully accommodated" for just a 7-billion-parameter model at FP16 once KV cache is included. Frontier-scale dense models need proportionally more: Llama-2/3-70B is commonly cited (SemiAnalysis, corroborated by independent industry analysts) as needing roughly 576 chips — eight full racks — purely to hold the weights on-chip (at 8-bit weights: 576 × 230 MB ≈ 132 GB, less than the 140 GB the model needs at FP16), since there is no HBM tier to fall back on ^{[138][145][147][148][149]}. Try it: Groq chips to hold 70B.

On GroqCloud, Artificial Analysis (an independent, vendor-neutral benchmark tracker) has repeatedly measured Groq as the fastest hosted inference provider for the models it serves — 276 tokens/sec on Llama 3.3 70B at every tested context length as of December 2024, for example. But the single most rigorous independent architectural comparison available, SemiAnalysis’s "Groq Inference Tokenomics," is careful to state the limit of that advantage directly: "Groq is not competitive architecturally at all for throughput-optimized scenarios" — its real advantage is single-stream, low-batch, low-latency serving, and Nvidia GPU systems retain an order of magnitude better performance-per-dollar for high-batch, throughput-optimized serving. The same analysis finds a real hardware-cost advantage per chip (Groq’s 725mm² GlobalFoundries 14nm die avoids both a leading-edge process and the ~$1,150/chip of HBM a comparable Nvidia part carries) that is partly offset at full-system scale by the sheer amount of host infrastructure a 576-chip deployment requires. A separate, peer-reviewed measurement at Argonne found Groq’s real decode advantage (328–336 tokens/sec at batch size 1, vs. an A100’s 73–84) comes with a genuine, underreported cost: dramatically worse prefill latency than a GPU in the tested configuration, and no support for batching at all once a binary is compiled for one batch size — a real limitation, not a hypothetical one, worth weighing against the headline speed claims ^{[148][149][150]}.

As of this writing, Groq’s trajectory has changed substantially. In December 2025, Nvidia struck a roughly $20 billion deal — structured as a non-exclusive IP license plus the hiring of most of Groq’s leadership team (including CEO Jonathan Ross) rather than a full acquisition — after Groq’s planned second-generation LPU (announced for Samsung’s 4nm process in 2023) was delayed and ultimately folded into the resulting joint effort instead of shipping independently. The chip that emerged, unveiled at GTC in March 2026, is branded "Nvidia Groq 3 LPX" (chip name LP30) and occupies a notably different position than a standalone LPU: Nvidia positions it as a dedicated decode-phase co-processor paired with Vera Rubin GPUs, which handle prefill, while LP30 handles the memory-bandwidth-bound decode step this primer has returned to repeatedly. Nvidia’s own disclosed per-chip figures for LP30 are a substantial jump from the original TSP: roughly 500MB of on-chip SRAM (vs. 230MB), 150 TB/s of SRAM bandwidth, about 1.2 PFLOPS of FP8 per chip (FP8 is now the headline precision, replacing the original chip’s INT8/FP16-only datapath), and 2.5TB/s of bidirectional chip-to-chip bandwidth over 96 links — with a full 256-chip LPX rack rated at 315 PFLOPS aggregate FP8 and 640TB/s of scale-up bandwidth. Nvidia’s own comparative claims ("35x higher inference throughput per megawatt" vs. GB200 NVL72) are the company’s internal modeling, not an independently benchmarked result, and should be read with the same skepticism this primer has applied to other vendors' headline claims throughout ^{[151][152][153][154]}.

CUDA’s core abstraction is a three-level hierarchy: a grid of thread blocks, each block containing up to 1,024 threads, each thread executing the identical kernel function on its own slice of data, identified only by its own coordinates (blockIdx, threadIdx) within that hierarchy. A kernel is launched with a chosen grid/block shape, and the GPU’s scheduler maps blocks onto streaming multiprocessors as hardware becomes available — the same SM/tensor-core hardware Section 5 described. The code below is close to the minimal, canonical CUDA kernel: each of N threads computes exactly one output element, with no loop at all inside the kernel body because the loop has effectively been unrolled across the whole grid.

Real kernels — a tiled matmul, a fused attention kernel — go further and explicitly manage the memory hierarchy from Section 10: a block cooperatively loads a tile of A and B into on-chip shared memory (visible to every thread in that block, roughly two orders of magnitude faster than a round trip to HBM), calls __syncthreads() as a barrier so no thread reads the tile before every thread has finished writing it, and only then does the arithmetic against the fast on-chip copy rather than against global memory directly. Getting that tiling, occupancy, and register allocation right by hand is genuinely difficult, and is exactly the work FlashAttention’s original CUDA implementation (Section 11) had to do to turn attention from memory-bound into compute-bound.

Triton, developed at OpenAI and now the standard way most PyTorch-ecosystem custom kernels get written, raises the unit of programming from a single scalar thread to a block (tile) of data at a time. A Triton kernel still runs across a grid, but inside the kernel the programmer manipulates whole vectors — tl.load and tl.store move a block of contiguous or strided elements at once — while the compiler is responsible for the actual thread mapping, shared-memory allocation, and instruction scheduling that CUDA forces the programmer to write explicitly. This is a large part of why Triton has become the default for writing new fused kernels quickly: PyTorch’s torch.compile emits Triton kernels automatically for fused elementwise and reduction operations it traces, and hand-written Triton kernels (a fused softmax, a custom attention variant) routinely reach 80–95% of a hand-tuned CUDA kernel’s performance at a fraction of the code ^{[176][177][178]}.

Apple’s Metal Shading Language targets the same TBDR GPU architecture described in Section 21, and structurally mirrors CUDA closely: a kernel function, thread positions supplied via attributes such as instead of CUDA’s built-in blockIdx/threadIdx variables, and buffers bound to numbered slots via [[buffer(n)]] rather than passed as ordinary pointer arguments. The most consequential difference is not syntax but memory: because Apple’s GPU and CPU share one physical memory pool (Section 21), there is no separate host-to-device copy step at all — a buffer allocated by the CPU is immediately visible to the GPU kernel, with no PCIe transfer in between. Newer Metal versions add simdgroup_matrix types and a simdgroup_multiply_accumulate intrinsic, giving Metal kernels direct access to Apple’s tile-level matrix-multiply hardware in much the same way CUDA’s WMMA/MMA intrinsics expose NVIDIA’s tensor cores ^[179][180].

CUDA runs only on NVIDIA hardware, so the single largest question for anyone deploying AMD’s Instinct GPUs (Part XI) is what the equivalent software stack looks like — and how much existing CUDA code has to be rewritten to use it. AMD’s answer is ROCm (Radeon Open Compute), an open-source platform spanning drivers, compilers, runtimes, and math libraries that occupies the same layer of the stack CUDA does, and HIP (Heterogeneous-interface for Portability), a C++ runtime and kernel language deliberately built to mirror CUDA’s syntax almost token for token ^[218][219]. The design bet is portability by near-identity: a HIP kernel is written the same way a CUDA kernel is — global functions, the same blockIdx / blockDim / threadIdx built-ins, the same triple-chevron launch — and the HIP compiler can target either AMD GPUs (through ROCm) or NVIDIA GPUs (by lowering straight back to CUDA), so one source tree serves both vendors.

The porting story is where the nuance lives. AMD ships hipify, a tool that mechanically rewrites CUDA source into HIP, and for the language and runtime layer shown above it is largely a find-and-replace (cudaMalloc → hipMalloc, and so on). What hipify does not do automatically is remap calls into NVIDIA’s proprietary libraries — those must be pointed at ROCm’s parallel implementations by hand: cuBLAS maps to rocBLAS (with hipBLAS as the vendor-portable wrapper over both), cuDNN to MIOpen, and TensorRT to MIGraphX, while Composable Kernel (CK) plays the role cuDNN/CUTLASS do for hand-fused operators like a FlashAttention kernel ^[219][220]. Because almost all real model code depends heavily on those libraries, "CUDA-source-compatible" is true at the kernel level but still leaves meaningful integration work at the library level.

Crucially, most practitioners never touch HIP directly, because the framework layer above it now treats ROCm as a first-class backend. PyTorch ships a native ROCm build (a pip-installable wheel as of PyTorch 2.9, bundling a ROCm build of Triton), and torch.compile emits Triton kernels that run on AMD hardware — so the same Triton-and-PyTorch stack described above targets Instinct GPUs largely unchanged, using a hybrid of Composable Kernel for the heaviest monolithic operators and Triton for pointwise fusions ^[220][221]. The remaining gap versus CUDA is one of maturity rather than capability: Triton autotuning and NVIDIA-derived kernel heuristics (tuned for 32-thread warps rather than AMD’s 64-thread wavefronts) can leave AMD performance on the table until they are re-tuned, which is the concrete, present-day form of the much-discussed "CUDA moat" — narrowing steadily, but not yet closed ^[221][222].

Above all three sits the level most model code is actually written at: PyTorch and JAX express a computation as a graph of high-level tensor operations, and a tracing compiler — torch.compile (via TorchInductor, which itself emits Triton) or JAX/XLA — captures that graph and lowers it to fused, hardware-specific kernels automatically, deciding for the programmer which operations to fuse, which to tile, and which existing hand-written kernel (a cuDNN or cuBLAS routine, a FlashAttention kernel) to call directly rather than regenerate. Two portable, vendor-neutral alternatives exist a level below that — OpenCL and its more modern successor SYCL, maintained by the Khronos Group — which express the same thread/block grid model in a single API that can, in principle, target NVIDIA, AMD, and Intel hardware alike, at the cost of usually trailing CUDA and Triton in tooling maturity and achieved performance on NVIDIA’s own hardware, which is why they see far more use in AMD and Intel’s own software stacks than as a genuine cross-vendor default ^{[181][182][183]}.

LLVM, the compiler infrastructure that carried Clang, Swift, and Rust, was built around a single scalar-and-vector IR designed for CPU-style "C-with-vectors" code. That representation loses almost everything a machine learning compiler needs: it doesn’t know what a tensor is, doesn’t preserve loop structure or affine index information, and offers no vocabulary for distributed computation ^[10]. By 2018 Google’s TensorFlow ecosystem had accumulated a stack of graph representations, converters, and codegen paths that shared almost nothing, and every new accelerator vendor was rebuilding the same 80% of compiler infrastructure from scratch — an N-frontends-times-M-targets problem ^[11].

MLIR (Multi-Level Intermediate Representation) is the answer: a single compiler framework that lets many IRs coexist and interoperate. Its central concept is the dialect — a self-contained namespace of operations, types, and attributes that models a particular domain at a particular abstraction level. A modern ML compiler stack might use half a dozen dialects in one program: tosa or stablehlo at the top to represent framework-level tensor ops; linalg and tensor to describe structured loop nests over tensors; affine and scf for loop transformations; vector for SIMD lowering; gpu for kernel and thread abstractions; and llvm at the bottom for scalar code generation. A model is progressively lowered through rewrite passes from the highest dialect toward the lowest, and crucially, information is never recovered — it is preserved from the top so it is still available when a lower pass wants to make a decision (fusion, tiling, vectorization) that depends on it ^[12]. MLIR-based flows routinely reach over 90% of hand-tuned "ninja" GEMM performance on modern accelerators, precisely because the middle-end retains enough structure to autotune without the human writing tile shapes by hand ^[13].

Even with MLIR as the framework, the industry still needed a stable top-level dialect that framework authors could target and accelerator vendors could consume. That’s what StableHLO provides. Extracted from XLA’s internal HLO operation set and hardened for external use, StableHLO defines roughly 100 operations with formal verifiers and type inference, covers dynamism and quantization, and — critically — commits to five years of backward compatibility and two years of forward compatibility, letting you compile a model today and run it on hardware and compilers shipped years apart ^[14].

Under the umbrella of OpenXLA — a multi-vendor project spun out of Google’s XLA in 2023 — StableHLO acts as the narrow waist between many producers (JAX, PyTorch/XLA, TensorFlow, and increasingly others via torch.export) and many consumers: Google TPUs via XLA, AWS Trainium and Inferentia via Neuron, AMD Instinct via a StableHLO-to-ROCm path, Cerebras and Graphcore via their own compilers, and IREE for CPU, mobile, and edge targets ^[15]. This is the concrete engineering answer to the "CUDA moat" question in Part VIII: rather than reimplement CUDA, a challenger implements a StableHLO-consuming backend once and inherits the entire JAX and PyTorch/XLA frontend for free. IREE (Intermediate Representation Execution Environment) is the most visible MLIR-native runtime in that role: it ingests StableHLO, progressively lowers through the dialects above, and emits code for CPU, GPU, and specialized accelerators without relying on hand-written vendor kernels ^[16].

PyTorch’s own compilation story deliberately takes a different route. Its frontend, TorchDynamo, uses CPython frame-evaluation hooks to capture Python code as it runs into an FX graph, gracefully falling back to eager execution on constructs it cannot trace — a design that lets torch.compile work on almost arbitrary user code without a separate export step. The captured graph is handed to TorchInductor, which lowers it through a define-by-run loop-level IR and generates two kinds of code by default: Triton kernels for GPU and C+/OpenMP for CPU footnote:[PyTorch Team, "TorchInductor: A PyTorch-native Compiler with Define-by-Run IR" (2024). https://dev-discuss.pytorch.org/t/torchinductor-a-pytorch-native-compiler-with-define-by-run-ir/747]. Across 180 models from TorchBench, HuggingFace, and TIMM, the Dynamo/Inductor pair delivered a geometric-mean 2.27x inference and 1.41x training speedup in float32 on A100 ^[17]. Because Inductor emits Triton, the same compiled graph runs on NVIDIA and AMD GPUs from one code path — the Triton-on-ROCm story from Part VIII is what makes that work in practice.

Training compilers optimize a static graph once; inference compilers must serve many concurrent requests, of varying lengths, whose KV caches don’t fit any regular tiling pattern. Three stacks dominate that layer in 2026, and they occupy different points on the specialization-versus-flexibility axis.

TensorRT-LLM is NVIDIA’s production compiler for LLM and generative inference. It exposes a high-level Python LLM API but internally builds a compiled engine — fused custom kernels for attention, GEMM, and MoE routing, plus algorithmic optimizations such as prefill/decode disaggregation, wide expert parallelism, and speculative decoding — that takes on the order of tens of minutes to build for a large model but then holds the throughput record on NVIDIA hardware ^[18].

vLLM took the opposite bet: no compile step, keep model onboarding cheap, and win on runtime scheduling. Its landmark contribution is PagedAttention, which stores the KV cache in fixed-size non-contiguous blocks managed like OS virtual memory pages — eliminating fragmentation and enabling flexible sharing across requests ^[19]. That memory model is by now standard (Part XVI covers it in depth), and vLLM is the default first-choice serving engine for most teams because it supports the widest model catalog with no compilation step.

SGLang is the newest of the three and is designed for structured generation — agents, JSON extraction, tree-of-thought, and prefix-heavy workloads. Its core mechanism is RadixAttention, which maintains an LRU cache of KV blocks in a radix tree keyed by token prefix, so any request that shares a system prompt or a few-shot preamble with a prior request reuses that prefix’s cache instead of recomputing it ^[20]. On prefix-heavy Llama 3.1 8B traffic, SGLang has been benchmarked at roughly 16,200 tok/s against vLLM’s ~12,500 tok/s — a ~29% advantage attributable almost entirely to RadixAttention’s prefix reuse ^[21].

Two projects reject the "layer more IR on the existing stack" approach entirely. Modular — founded by Chris Lattner, who also created LLVM, Clang, Swift, and MLIR — builds MAX, a unified inference platform, on top of Mojo, a new systems language that is a superset of Python but compiles through MLIR directly to device code. Mojo hit 1.0 in August 2026 with Apache 2.0 open-source licensing, and Modular was acquired by Qualcomm the month before — a signal that a serious silicon vendor sees the MLIR-based "kill CUDA the language, not just CUDA the runtime" bet as strategic ^[22]. The wager is that a single language, aware of tensors and heterogeneity from the start, will beat the CUDA + Triton + PyTorch + XLA stack on programmer productivity while matching it on generated-code quality across CPUs, GPUs, and custom silicon.

At the other end of the ambition axis, tinygrad (George Hotz) asks how few concepts are needed to run LLaMA and Stable Diffusion, and answers with a codebase in the low tens of thousands of lines. It represents a program as a lazy tensor DAG, fuses aggressively into a small set of ~25 primitive ops (UnaryOps, BinaryOps, ReduceOps, MovementOps), and JITs that lowered form to OpenCL, CUDA, Metal, or WebGPU ^[23]. It won’t beat TensorRT-LLM on throughput, but it is the reference for what a minimal ML compiler has to contain, and it doubles as the software stack for tinygrad’s own AMD-based tinybox hardware.

The distinction matters. Training compilers (XLA, torch.compile, IREE, MAX) see the whole graph once, know the batch and sequence shapes, can autotune fusion and tiling for weeks of runtime, and target a fixed cluster topology. Inference compilers (TensorRT-LLM, vLLM, SGLang) live in a hotter loop: request shapes vary per token, KV caches must be paged, batches form and dissolve continuously, and every millisecond of scheduler latency shows up in tail p99. The IR layers described here — MLIR dialects, StableHLO — are common infrastructure, but the top of the stack forks: training compilers optimize throughput of a known workload; inference compilers optimize throughput and latency of an unknown, adversarial, streaming one. Getting that split right is what separates a research-grade compiler from a production one.

Hopper (H100, 2022) is not just "Ampere with more FLOPs." It introduced four features that fundamentally change what a fast kernel is. The Tensor Memory Accelerator (TMA) is a dedicated DMA engine: a single thread issues a descriptor and the hardware moves a 1D–5D tile between global and shared memory asynchronously, freeing every other thread in the block for arithmetic ^[24]. Warp-group MMA (wgmma.mma_async) issues a matmul across 128 threads (four warps) that runs asynchronously against the tensor cores, so the issuing warp group can continue softmax or epilogue work while the MMA retires ^[25]. Thread-block clusters let up to 16 blocks on H100 run co-scheduled on adjacent SMs within a GPC, connected by a dedicated SM-to-SM network; distributed shared memory makes another SM’s shared memory addressable with load/store/atomics, roughly 7× faster than routing through global memory ^[26]. And asynchronous transaction barriers synchronize on both thread arrivals and byte counts moved, so a producer warp can signal "the tile is here" atomically with the transfer completing.

Triton’s programming model — one implicit warp scheduler, tiles as the programmer’s unit, everything under the tile handled by the compiler — cannot express producer/consumer warp specialization across these asynchronous units without leaking hardware details up into the DSL. In practice, that means the kernels whose speed depends on overlapping TMA loads, wgmma issue, and softmax computation on distinct warp groups — the FlashAttention family, ping-pong GEMMs, MoE gathers with structured expert skew — are still written a layer down. This is the ceiling: Triton is excellent at what it targets (portability across NVIDIA and AMD, fast iteration on new fusions), but its tile-first abstraction erases exactly the scheduling knobs Hopper made programmable ^[27].

CUTLASS is NVIDIA’s C++ template library for GEMM and adjacent operations; CUTLASS 3.0 (2023) rebuilt it around a new core library called CuTe ("Cute Tensor"), which introduces a single Layout vocabulary type and an algebra over layouts that composes tiling, partitioning, and thread-to-data mappings as first-class objects ^[28]. Layouts are checked at compile time — if the code compiles, the thread-to-tensor mapping is correct by construction — which turns the error-prone bookkeeping of manual ldmatrix, swizzling, and tensor-core fragment layout into type-level algebra. The GEMM hierarchy is then factored into an architecture-agnostic collective mainloop (tile-loop structure) and atoms that plug in per-architecture (wgmma on Hopper, tcgen05 on Blackwell), so the same collective can retarget generations.

Two low-level realities are surfaced explicitly. First, shared-memory swizzling: tensor-core loads via ldmatrix do vertical per-lane column reads, and a naïve row-major tile causes multiple lanes to hit the same bank; CuTe encodes the swizzle as part of the layout, and empirically it can drop a Hopper GEMM from ~913k to ~303k cycles for the same math ^[29]. Second, register pressure: a warp-group MMA fills a large accumulator fragment in registers, and if the epilogue and next tile’s operands do not fit, spills to local memory collapse throughput; CUTLASS’s ping-pong scheduler alternates two warp groups so one does the softmax/epilogue while the other issues the next MMA, keeping the tensor cores at ~100% duty and splitting the register budget in half ^[30].

FlashAttention-3 (Shah, Bikshandi, Zhang, Thakkar, Ramani, Dao, 2024) is the load-bearing example. It is implemented on CUTLASS/CuTe and uses three Hopper-specific techniques the Triton-2 version could not: warp specialization with dedicated producer warps issuing TMA and consumer warp groups issuing wgmma; the ping-pong schedule interleaving softmax on one warp group with GEMM on another; and FP8 block quantization with incoherent processing (Hadamard rotation) to control outliers. The result: 1.5–2.0× faster than FlashAttention-2 in FP16, ~75% of H100 peak (up from ~35%), and ~1.2 PFLOPS in FP8 ^[31]. Every one of those wins is a hand-scheduled use of a Hopper primitive Triton does not expose.

CUTLASS is powerful and difficult; a production Hopper GEMM is several thousand lines of template metaprogramming. ThunderKittens (Spector, Arora, Singhal, Fu, Ré; Hazy Research, October 2024; ICLR 2025) is the answer to the question "what if you kept CUTLASS’s access to wgmma/TMA but threw out most of the templates?" ^[32]. It exposes three levels matching the GPU hierarchy: 16×16 register tiles (rt) and shared tiles (st) as the basic data types at warp level, with PyTorch-like tile-wise operations; a thread-block layer with templated producer/consumer overlap primitives that call TMA and wgmma directly; and a grid-level layer that hides launch and teardown. The library targets Hopper (H100, wgmma) and Blackwell (B200, tcgen05) — Ampere is no longer actively supported ^[33].

The performance point is that "close to CUTLASS, in one page of code" is achievable: ThunderKittens matches cuBLAS and FlashAttention-3 on GEMM and attention forward, beats the strongest baselines by 10–40% on attention backward, and delivers 8× on state-space models and 14× on linear attention against prior open kernels — the last two being cases where CUTLASS-quality implementations simply didn’t exist. A ~200-line ThunderKittens H100 GEMM in the README hits ~855 TFLOPS, 86% of theoretical peak. The design philosophy is that a small, opinionated set of tile primitives with correct default swizzling and register layouts is enough to close most of the gap to hand-tuned CUTLASS, for the class of kernels ML researchers actually write (attention variants, custom recurrences, MoE routers).

Blackwell (SM100, 2025) moved the tensor-core interface again. The tcgen05.mma (UMMA) family replaces Hopper’s wgmma with warp-level rather than warp-group-level issue, delivers 2–4× the throughput per instruction, and stores accumulators in a new on-chip Tensor Memory (TMEM) — 256 KB per SM, 512×128 32-bit cells — rather than in registers, dramatically cutting the register pressure that constrained Hopper kernel designs ^[34]. UMMA is single-thread-launch (one thread issues even a 2-CTA cooperative MMA), atoms are larger (128×256×16 versus Hopper’s 64×128×16), and the supported data types now include the microscaling formats introduced in Part II: NVFP4, MXFP4, MXFP6, MXFP8, plus the traditional FP8/BF16/TF32 paths. Block-scaled MXFP4 hits 4× Hopper’s FP8 throughput per SM ^[35].

CUTLASS 3.8 and later ship the Blackwell backend as a set of new atoms plugged into the existing collective mainloop, which is precisely the architectural payoff of CuTe’s layout algebra: the same kernel structure retargets, and only the atoms change. ThunderKittens shipped Blackwell support in the same generation. Triton support for tcgen05 exists but trails the hand-written path — the standard pattern of the last five years, where a new tensor-core generation is first tapped by CUTLASS/ThunderKittens kernels (weeks after silicon), then by Triton (months later), then by torch.compile fusions on top of that.

The practical rule that has emerged is not "one language wins" but a stratified stack.

FlashAttention-3, xformers' Hopper-tuned attention backends, and NVIDIA’s own cuBLAS/cuDNN kernels sit in the bottom row; DeepSpeed and vLLM’s fused paged-attention variants and a growing number of custom SSM and linear- attention kernels sit in the middle row; almost everything else — the long tail of custom fusions that torch.compile emits — is in the top row. The important structural fact for 2026 is that this is not a temporary state. As long as each new NVIDIA generation exposes new asynchronous primitives (TMA on Hopper, TMEM and single-thread UMMA on Blackwell, whatever comes on Vera Rubin), there will be a 6–12 month window where only CUTLASS and ThunderKittens-class libraries can wring the new peak out of the silicon, and Triton catches up after. The kernels that define frontier inference throughput will keep being written below Triton — and read by anyone who wants to understand why one lab’s H100 delivers 75% of peak while another’s delivers 35%.

SYCL is the Khronos Group’s answer to CUDA at the language layer: a royalty-free open standard that expresses heterogeneous compute in single-source modern C, with host and device code compiled from the same translation unit, offloaded to CPUs, GPUs, FPGAs, or DSPs footnote:[Khronos Group, "SYCL 2020 Specification (revision 12)" (2021). https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html]. SYCL 2020, ratified in February 2021 and based on C17, was the first version to decouple SYCL from OpenCL as its only substrate: implementations now target CUDA, HIP, Intel Level Zero, or OpenCL as backends, and SYCL 2020 adds unified shared memory (a single pointer visible to host and device), group algorithms, and sub-groups — features first pioneered inside Intel’s open-source DPC (Data Parallel C) compiler ^[36].

Intel’s oneAPI ships DPC as its production SYCL compiler and bundles math, DNN, and collective libraries (oneMKL, oneDNN, oneCCL) beneath it, positioning SYCL as the Intel-preferred cross-vendor answer to CUDA on Ponte Vecchio, Gaudi, and Arc GPUs footnote:[Khronos Group, "SYCL 2020 Specification" press release (2021). https://www.khronos.org/news/press/khronos-releases-sycl-2020-final-specification]. The independent implementation that most matters, though, is _AdaptiveCpp_ — the community project formerly named hipSYCL and then OpenSYCL — which is the only SYCL toolchain that can target Intel, NVIDIA, _and_ AMD GPUs (plus CPUs) from a single binary and, in the C standard parallelism (std::par) role, has been shown to outperform vendor compilers on several benchmarks ^[37]. llama.cpp uses SYCL as its Intel-GPU path (Arc, Flex, Data Center Max), which is a good illustration of where SYCL has actually stuck: as a first-class backend for vendors that lack an incumbent CUDA-equivalent runtime, rather than as a portable replacement for CUDA on NVIDIA.

OpenCL, SYCL’s own predecessor and once the industry’s flagship portable compute API, is now legacy in AI contexts. Its C-based kernel language, manual buffer management, and lack of a single-source model made it strictly less pleasant than CUDA even before tensor cores existed, and almost every modern portable stack (SYCL, ONNX Runtime’s mobile providers, tinygrad’s fallback path) treats OpenCL as an execution target, not a programming model. It survives because Adreno, Mali, and older AMD/Intel integrated GPUs still expose it, and llama.cpp keeps an OpenCL backend specifically for Adreno and Arm-based devices ^[38].

Two portable runtimes are what practitioners reach for when they want one artifact to run on any hardware. IREE (Intermediate Representation Execution Environment) is Google’s MLIR-native compiler-and-runtime: it ingests StableHLO, ONNX, or torch.export output, progressively lowers through the dialect stack described in the previous section, and emits code for CPU (LLVM), Vulkan, CUDA, ROCm, Metal, and embedded targets from one flow. Binary sizes fall as low as ~30 KB for embedded deployments, and the project is production enough that AMD submitted an IREE-based SDXL implementation to MLPerf in early 2025 ^[39].

ONNX Runtime (ORT) is the more widely deployed of the two — the onnxruntime package sits inside Windows itself, Office, Azure ML, Hugging Face Optimum, and countless mobile apps — and its cross-vendor mechanism is the execution provider (EP) interface. An ONNX model is a graph; an execution provider is a plugin that claims a subset of that graph’s nodes for a particular hardware backend, and the runtime partitions the graph across whichever EPs are installed, with unclaimed nodes falling through to the CPU provider ^[40]. The list of shipping EPs is the map of the industry: CUDA and TensorRT for NVIDIA, ROCm and MIGraphX for AMD, OpenVINO for Intel CPU/GPU/NPU, CoreML for Apple, DirectML for any Windows GPU, QNN for Qualcomm, CANN for Huawei Ascend, WebGPU for browsers, plus XNNPACK, oneDNN, NNAPI, and Arm Compute Library for CPUs and phones. The same C, C++, Python, C#, Java, or JavaScript API drives all of them, and the user chooses a priority order — for example, "try TensorRT, fall back to CUDA, fall back to CPU" — without touching model code.

The two runtimes reflect two philosophies. IREE compiles ahead of time into device-specific binaries and treats MLIR as the source of truth; ONNX Runtime dispatches to whatever vendor library is fastest on the current node and treats the ONNX graph as the source of truth. Both work. ONNX Runtime wins deployment share because its EP list matches the industry one-to-one and because Microsoft has shipped it on every Windows install since Windows 11 as the substrate for DirectML and Windows Copilot Runtime.

Triton, described in the previous section as OpenAI’s tile-level kernel language, has an AMD backend that quietly closed most of the practical CUDA-only gap for custom kernels. Since ROCm 5.7, and much more usably since ROCm 6.2 in 2024, pip install triton on an AMD PyTorch build gives a Triton toolchain that lowers straight to AMDGPU LLVM, and `torch.compile’s Inductor path emits those Triton kernels on MI-series GPUs unchanged ^[41][42]. The remaining performance gap is autotuning: Triton’s default kernel heuristics were tuned for NVIDIA’s 32-thread warps, and AMD’s 64-thread wavefronts benefit from different tile shapes — an issue of tuning, not of the abstraction itself.

Vulkan Compute is the biggest surprise of the last two years. Vulkan was designed as a graphics API, but its compute-shader path is a genuine cross-vendor GPU abstraction — one driver interface, one SPIR-V IR, one binary that runs on any NVIDIA, AMD, Intel, or Apple GPU (via MoltenVK) shipped this decade. llama.cpp’s Vulkan backend has become the default consumer path: it needs no vendor toolkit (no CUDA install, no ROCm install), it runs on Windows (where ROCm for llama.cpp is effectively Linux-only), and its performance on token generation now matches CUDA on NVIDIA and often beats ROCm on AMD on the same hardware ^[43][44]. On Radeon, prompt processing still favors ROCm (~753 vs ~624 pp512 tokens in one benchmark) but token generation on Vulkan reaches ~49.4 vs ROCm’s ~47.0 tg128 — close enough that most single-user deployments prefer the Vulkan path for its portability ^[45]. The lesson is that inference workloads — mostly memory-bandwidth-bound during decode — hide the tensor-core advantage that separates vendor-native from vendor-neutral in training, so a driver-level abstraction can now serve most users without a measurable performance tax.

Apple’s Metal is the counter-example to portable ambitions: it is resolutely vendor-native, exposes Apple’s tile-based deferred rendering GPU and its simdgroup_matrix intrinsics directly, and is the backend that every other portable stack (llama.cpp, ONNX Runtime CoreML EP, tinygrad, IREE) targets when running on Apple silicon. This is the same tension shape as CUDA-vs-portable, just on a smaller ecosystem: Apple publishes MoltenVK to translate Vulkan onto Metal for portability, but serious Apple-silicon performance is always written directly to Metal Shading Language.

WebGPU is the browser’s answer, and by 2026 is the most reach any GPU API has ever had. Chrome, Firefox, Edge, and Safari all ship WebGPU by default, covering roughly 83% of global browser traffic, and WGSL (WebGPU Shading Language) is its portable shader language — a Rust-flavored, statically-validated, explicit-binding language that compiles behind the scenes to Metal Shading Language on macOS/iOS, HLSL on Windows/DirectX, and SPIR-V on Vulkan platforms ^[46]. WebLLM, transformers.js, and ONNX Runtime Web now serve 4-bit-quantized 7B–8B models entirely inside the tab at 25+ tokens/second on consumer laptops, with WGSL kernels implementing PagedAttention and FlashAttention inside the browser’s memory budget ^[47]. The abstraction cost is a 3–10 second one-time WGSL compile on first load, and — as with Vulkan — the near absence of vendor tensor-core intrinsics in the exposed API, which caps peak FLOPs below what the same hardware would reach through Metal or CUDA directly.

The vendor-neutral abstractions (SYCL, ONNX Runtime, Vulkan, WebGPU) and the vendor-native ones (CUDA, ROCm, Metal, TensorRT) don’t compete on the same axis. Native paths win peak-FLOPs performance because they can name tensor-core intrinsics and vendor-specific memory tiers directly; portable paths win every deployment property that isn’t peak performance — one binary across vendors, no toolkit install, one API across programming languages, browser reach, mobile reach, Windows-versus-Linux neutrality. The empirical outcome, visible in every 2026 production inference stack, is a division of labor: training runs on CUDA (with ROCm second) because training is compute-bound and every percent of MFU compounds over weeks; inference runs through ONNX Runtime, IREE, vLLM (often with a Vulkan or ROCm backend), or a WebGPU runtime, because inference is memory-bound and the portable-abstraction tax collapses to single-digit percent while the deployment savings are decisive.

The forcing function that keeps the portable side honest is MLPerf. MLPerf Inference v5.1 (September 2025) drew a record 27 submitting organizations — AMD, Intel, NVIDIA, Google, and a long tail of OEMs, cloud providers, and academic groups — reporting on five newly-available accelerators including AMD MI355X, Intel Arc Pro B60, and NVIDIA GB300, alongside three new workloads (DeepSeek-R1, Llama 3.1 8B, Whisper Large V3) ^[48]. That participation is what forces every vendor’s software stack to cover every current model architecture within a benchmark cycle. Without MLPerf, the "CUDA moat" would compound purely through model authors targeting NVIDIA first; with it, an AMD or Intel submission requires that ROCm or oneAPI have a working, competitive path for DeepSeek-R1 or Llama 3.1 within months of release. The abstractions in this section — ONNX Runtime, IREE, SYCL, Triton-on-AMD — are the plumbing that lets a non-NVIDIA vendor turn up on the leaderboard at all, and the reason they receive continuous investment despite the CUDA-first reality of most training.

Choosing only the top-k experts by itself does not discard information the way, say, cropping an image would — it is closer to the network learning to route each token down whichever path it has learned produces the right answer, the same way a very deep dense network already routes different inputs through effectively different computational paths via its own nonlinearities, just made explicit and sparse. The genuine, measurable information-loss mechanism in MoE is a specific, well-documented systems detail: load-balancing by fixed expert capacity. Because every expert must run on real hardware with a real memory and compute budget, most production MoE implementations (GShard, Switch Transformer, and Expert Choice routing) assign each expert a fixed per-batch token capacity, and if more tokens are routed to a popular expert than its capacity allows, the overflow tokens are dropped from that expert entirely for that layer — typically passed through via a residual/skip connection with no expert computation applied at all rather than being processed by a lower-priority expert. That is a real, non-hypothetical information-loss event, distinct from the routing decision itself, and it is precisely why load-balancing losses exist: Switch Transformer and GShard add an auxiliary loss term during training that penalizes uneven expert utilization, while DeepSeek-V3 replaces the auxiliary loss with a lightweight, dynamically-adjusted per-expert bias added directly to the router’s scores, nudging tokens away from overloaded experts without the auxiliary loss’s tendency to slightly hurt the model’s primary training objective. Expert Choice routing inverts the relationship entirely — instead of each token choosing its top-k experts, each expert chooses its own top-tokens up to a fixed capacity — which guarantees perfect load balance by construction but reintroduces the same drop risk from the opposite direction, for whichever tokens no expert chooses to fill its capacity with ^{[186][187][188][189]}. Try it: Watch tokens overflow capacity; Aux loss vs bias balancing; DeepSeek-V3 sigmoid gating; Expert Choice routing.

Once a token’s chosen experts have each produced their own output vector, reconciliation is simple and lossless by construction: the final output for that token at that layer is a weighted sum of the selected experts' outputs, using the same router probabilities (renormalized over just the selected top-k) as weights:

Try it: Renormalised gates, weighted sum.

where is the token’s hidden state, the -th selected expert, and its (renormalized) router gate weight, summed over just the chosen experts.

Nothing is voted on, truncated, or averaged away; every activated expert’s full output vector contributes to the final result in exact proportion to how strongly the router selected it. In a multi-GPU expert-parallel deployment (Section 14), this combine step is the mirror image of the earlier dispatch step: tokens were scattered to whichever GPU holds their assigned expert via an all-to-all communication primitive, and the combine step gathers each token’s expert outputs back to its original "home" GPU via a second all-to-all, after which the weighted sum is computed locally. Real production architectures vary this same skeleton in specific, named ways: Mixtral routes each token to 2 of 8 experts per layer; DeepSeekMoE and DeepSeek-V3 use many more, much smaller "fine-grained" experts (a finer-grained division of labor than a few large experts allows) alongside one or more always-on "shared experts" that every token passes through regardless of routing, explicitly to capture common knowledge that would otherwise be wastefully duplicated across many specialized experts; and Llama 4 follows the same shared-expert-plus-routed-experts pattern popularized by DeepSeek (gpt-oss, by contrast, routes each token to 4 of 128 experts with no shared expert), reflecting a broad 2025–2026 industry convergence on "a few things every token needs, plus a sparse, learned selection of things only some tokens need" as the working answer to how to split parameters without splitting competence ^{[186][190][191]}. Try it: Dispatch and combine volume; Mixtral: 46.7B total, 12.9B active; DeepSeek-V3: 671B / 37B active; gpt-oss: no shared expert; gpt-oss-120b in 80 GB.

Every transformer layer projects its input hidden states into three separate vectors per token — a Query, a Key, and a Value — via three learned weight matrices; this is exactly the matmul-heavy workload Section 5’s tensor cores exist to accelerate. Attention then scores every token’s Query against every other token’s Key (a dot product, scaled and passed through softmax to produce a probability distribution over "how much should this token attend to every other token"), and uses those probabilities to compute a weighted sum of every other token’s Value vector. Written compactly, that is the single most important equation in the field:

Try it: Try the attention equation.

where , , and are the query, key, and value matrices (one row per token), is the key-vector dimension, and the scaling keeps the dot products from growing too large before the softmax.

Multi-head attention simply runs many independent, smaller copies of this same Q/K/V mechanism in parallel per layer (each "head" with its own, separate learned projections) and concatenates their outputs, letting different heads specialize in different relational patterns; GQA (grouped-query attention) and MQA (multi-query attention) shrink the KV cache specifically by having multiple query heads share a smaller number of key/value heads, and DeepSeek’s Multi-head Latent Attention compresses the cached K/V further into a lower-rank latent representation reconstructed on the fly — both are direct responses to the KV-cache memory-bandwidth pressure Section 11 already described, applied to the attention mechanism’s own shape rather than to the serving system around it ^{[192][193][194]}. Try it: GQA, MQA and MLA visually; GQA: 64 query, 8 KV heads.

When a request first arrives, every token of the prompt is already known in full, so the model computes Query, Key, and Value for every prompt token, and every layer’s attention, in one shot — a large, dense matrix-times-matrix operation across the entire prompt length at once, with essentially no sequential dependency between tokens within this step. This is why prefill is compute-bound and highly parallel: it looks, computationally, like a single big batched matmul, exactly the shape GPUs and TPUs are built to run at close to peak FLOPs, and FlashAttention (Section 11) is what keeps that large attention matrix from ever fully materializing in slow HBM during this step. Try it: FlashAttention tiling, step by step.

Generating the next token is structurally different, and this is the crux of the "surely not left to right" intuition: computing token N+1 requires knowing token N, because token N+1’s own Key and Value vectors don’t exist yet, and — more fundamentally — because autoregressive generation is defined as sampling one token at a time from the model’s own probability distribution over the vocabulary, conditioned on everything generated so far including itself. There is no way to compute token 47 before token 46 exists, because token 46 might change what token 47 should even be. That is the actual dependency forcing sequential, one-token-at-a-time generation — not a left-to-right reading convention, but a genuine data dependency: each step’s output is required as part of the next step’s input. Try it: See the causal mask.

What is not iterative, and is worth separating out explicitly, is the path through the network’s own layers within a single step: a single forward pass for one new token still runs strictly feedforward, once, from layer 1 to the final layer, exactly like prefill’s pass did — the network does not loop back to revisit earlier layers within that one step. The iteration is entirely across time steps (across generated tokens), not across layers within a step. Concretely, the decode loop is: (1) take the newest token, run it through every layer once, where each layer’s attention reads that new token’s own Q against the Key/Value vectors of every previous token pulled straight from the KV cache rather than recomputed; (2) take the final layer’s output, apply a softmax over the vocabulary, and sample the next token; (3) compute that new token’s own Key and Value vectors and append them to the cache; (4) feed the newly sampled token back in as the input to step (1) again, and repeat until an end-of-sequence token or length limit is hit. Because step (1) only has to touch one new token’s worth of fresh compute per layer while still needing to read the entire growing KV cache from HBM, decode is memory-bandwidth-bound rather than compute-bound — the opposite profile from prefill — which is exactly why continuous batching (packing many different requests' individual decode steps together into one batched, GPU-filling operation) and chunked prefill (splitting a large new prompt’s prefill into pieces interleaved with other requests' pending decode steps, so a long prefill doesn’t stall everyone else’s token generation) exist as serving-system techniques, and why Section 16’s claim that inference has "two very different phases" is a hardware-utilization statement as much as an algorithmic one ^{[195][196][197][198]}. Try it: Sampling temperature 0.7; Greedy decoding (T → 0); Hot sampling, T = 1.5.

A GB200 NVL72 compute tray — the basic field-replaceable server unit — pairs 2 Grace CPUs (using power-efficient LPDDR5X rather than conventional DIMMs, consistent with Part V’s framing of the host CPU’s job as getting out of the accelerator’s way as cheaply in power as possible) with 4 Blackwell GPUs, several ConnectX-7/8 NICs (one per one or two GPUs, following the dedicated-rail design from Section 12), one or two BlueField-3 DPUs to offload networking, storage, and security tasks from the host CPU entirely, and local NVMe storage for checkpoints and data staging. Multiple such trays, plus separate NVLink-switch trays, populate one rack.

The GB200 NVL72 rack assembles 18 compute trays and 9 NVLink-switch trays into 72 Blackwell GPUs behaving as one non-blocking NVLink domain (Section 13) over a passive copper backplane, at roughly 120–132 kW per rack — an order of magnitude above a conventional enterprise server rack, delivered via DC bus bars and shared power shelves rather than per-server AC power supplies, and cooled by direct-to-chip liquid cooling manifolds rather than rack-level air alone. Try it: NVL72 rack wiring; One NVL72 rack’s power.

NVIDIA’s reference "Scalable Unit" (SU) groups 8 such racks — 576 GPUs, roughly 1.2 MW including networking (8 × 132 kW ≈ 1.06 MW for the compute racks alone) — behind a shared tier of scale-out networking: per-SU InfiniBand leaf switches (Quantum-2, 64 ports of 400G) for latency-sensitive collective traffic, or Ethernet leaf/spine switches (Spectrum-X, SN5600-class, 64 ports of 800G) for more general multi-tenant traffic, matching the rail-optimized fat-tree topology from Section 13 at physical scale. Cabling technology crosses over at roughly this scale too: passive copper DAC cables stay economical only within about 2–3 meters, which is usually enough to reach within a row, while any run beyond that — between rows, or up to spine switches — needs active optical cables (AOC) instead, since copper’s signal loss at 400G/800G rates becomes unworkable much past that distance. Try it: Size a Scalable Unit.

Multiple Scalable Units combine through a further spine/super-spine switching tier into a full datacenter hall — NVIDIA’s published reference designs describe configurations up to 16 SUs, roughly 128 racks and 9,216 GPUs, as a single fat-tree fabric — with the row-level topology from Section 13 simply repeating one tier higher. Beyond one hall, a campus is mostly a power and cooling problem rather than a networking one: multiple halls sharing substation capacity, chiller plants, and water infrastructure, which is exactly the domain the next section’s compute-per-watt framing and Part VII’s memory-crisis framing both ultimately answer to — at gigawatt campus scale, the available power envelope, not the available silicon, is usually what caps how much compute a given site can actually run ^{[205][206][207][208]}. Try it: Build a 16-SU hall.

Not all "links" between chips are the same kind of thing, and the distinction is fundamental. A packet network — InfiniBand or RoCE over the NICs of Section 12 — moves data as explicit messages between separate memory domains: software posts a send or an RDMA operation, and packets cross the fabric. A coherent load/store fabric is categorically different: it lets a processor issue ordinary memory load / store instructions against memory that physically lives on another chip, with hardware keeping caches coherent so every agent sees one consistent view at cache-line granularity. The first is throughput-oriented messaging across a cluster; the second extends the memory model itself across chips, at far lower latency. Four technologies span this spectrum, and knowing which category each belongs to explains most of the system architecture around a modern accelerator.

PCIe is the baseline every accelerator still speaks. It doubles per-lane rate each generation — PCIe 4.0 at 16 GT/s per lane, 5.0 at 32 GT/s, and PCIe 6.0 at 64 GT/s (up to ~256 GB/s bidirectional at x16), released by PCI-SIG in 2022. PCIe 6.0 hit that rate by switching from simple two-level NRZ signaling to PAM4 (four voltage levels, two bits per symbol) at the same ~16 GHz Nyquist frequency, and — because PAM4 raises the raw error rate — mandating FLIT mode: fixed 256-byte flow-control units carrying lightweight forward-error-correction and CRC inline. The problem is that even a full PCIe 5.0 x16 link (~64 GB/s each way, ~128 GB/s both ways) is roughly an order of magnitude too slow for the GPU-to-GPU all-reduce traffic of Part III, which is the entire reason NVLink exists: to bypass PCIe with a wider, coherent, more energy-efficient path between accelerators ^[444][445]. Try it: PCIe generations compared; PCIe 6.0 PAM4 Nyquist; PCIe 5.0 x16 each way.

NVLink-C2C is NVIDIA’s chip-to-chip coherent link, the one that fuses a Grace CPU and a Hopper GPU into a single Grace Hopper superchip. Its figures make the PCIe gap concrete: up to 900 GB/s total (450 GB/s each way) — quoted as 7× the bandwidth of a PCIe Gen5 x16 link — at 1.3 picojoules per bit, more than 5× more energy-efficient than PCIe Gen5. Critically it is hardware-coherent: CPU and GPU threads address one unified memory space and transparently reach each other’s memory with plain loads and stores, no explicit copy ^[446]. Try it: NVLink-C2C: 7× PCIe 5.0.

CXL (Compute Express Link) is the open coherent standard, and it does something clever: it rides on the PCIe physical layer (CXL 3.0 uses the PCIe 6.0 64 GT/s PHY) but layers three protocols on top — CXL.io (PCIe-style discovery and DMA, the mandatory baseline), CXL.cache (lets a device coherently cache host memory), and CXL.mem (lets the host use device-attached memory as its own cacheable memory). Its killer applications are memory expansion (bolting more capacity and bandwidth onto a host than its DIMM slots allow) and memory pooling (a shared pool of memory dynamically parceled out across many hosts so none has to be provisioned for its worst case); CXL 3.0 added outright memory sharing, a region coherently shared across hosts. This directly targets the memory-capacity wall of Part VII — a way to attach more bytes to an accelerator complex without stacking more HBM ^[447]. Try it: CXL 2.0 vs 3.x bandwidth.

UALink (Ultra Accelerator Link) is the open answer to NVLink’s proprietary lock — the same strategic pressure that produced the AMD and hyperscaler-silicon stories of Part XI, applied to the interconnect. Its 200G 1.0 specification was ratified in April 2025 by a promoter group of AMD, Broadcom, Cisco, Google, HPE, Intel, Meta, and Microsoft; it defines a scale-up accelerator-to-accelerator fabric — memory-semantic read/write/atomic transactions between as many as 1,024 accelerators through UALink switches — explicitly to do inside a rack what NVLink does, without tying buyers to one vendor. It complements Ultra Ethernet (Part XVII), which plays the same open-standard role for the scale-out packet network. First UALink silicon is expected in the 2026–2027 window ^[448]. Try it: UALink port vs NVLink.

Try it: Interconnect bandwidth and pJ/bit.

The compute tray carries a DPU (data processing unit) beside the host CPU; it is worth saying what that part actually does, because it embodies a real architectural principle — keep the host CPU out of the accelerator’s way (Part V) by moving infrastructure work onto dedicated silicon. A DPU such as NVIDIA’s BlueField-3 pairs the NIC with programmable Arm cores and hardware accelerators, and offloads, accelerates, and isolates three domains at once: networking (software-defined networking, overlay tunnels, virtual switching), storage (presenting remote NVMe-over-fabrics drives as if they were local, plus inline compression), and security (line-rate encryption up to 400 Gb/s and a hardware root of trust). AWS built the same idea into its Nitro system years earlier: a family of Nitro cards that offload VPC networking, EBS remote storage, and local NVMe entirely off the host x86 — plus a Nitro security chip and a minimal hypervisor — so instances run at near-bare-metal speed and are hardware-isolated from each other and from AWS itself (the confidential-computing foundation of Part XXII). The through-line is that at datacenter scale the infrastructure tax — virtualization, networking, storage, encryption — is large enough to deserve its own processor, freeing every host CPU cycle to feed the accelerator ^[449][450].

This connects back to the NIC of Section 12 through GPUDirect RDMA, the mechanism that lets the network write directly into GPU memory over PCIe, bypassing a staging copy through host DRAM and the CPU entirely — the same "don’t detour through the host" principle that makes one 400 Gb/s NIC per GPU (as in a DGX H100’s eight ConnectX-7 adapters, 3.2 Tb/s aggregate) actually deliver its bandwidth to the accelerator rather than to the CPU’s memory controller ^[451]. Try it: Eight NICs: 3.2 Tb/s.

Getting power to the die is its own engineering story, and it hinges on one equation. Resistive loss in a conductor is : it scales with the square of the current. For a fixed delivered power , raising the distribution voltage cuts the current proportionally — so moving a rack’s internal distribution from the legacy 12 V to 48 V quarters the current and drops the busbar’s resistive loss by roughly 16× for the same copper. That is why the Open Compute Project’s Open Rack v3 standardizes a 48 V DC busbar (rated on the order of 18 kW), fed by shared power shelves rather than a power supply in every server — and why a GB200 NVL72 rack can distribute ~120 kW down a passive backplane at all. On the board, on-package voltage regulator modules then step that 48 V rail down to the sub-1-volt, hundreds-of-amps supply the GPU die actually runs on, as close to the die as possible to minimize the final high-current path. (For add-in-card GPUs, the analogous story is the 16-pin 12VHPWR connector rated to 600 W — and its well-publicized overheating failures, which forced the revised 12V-2×6 design, a reminder that at these currents the connector is a real engineering limit.) The frontier trend, still emerging as of 2026, is vertical power delivery — feeding current up through the back of the package to shorten that high-current path further ^[452][453]. Try it: Compare 12 V and 48 V loss.

Part XXI covers the cluster-scale data pipeline and parallel filesystems; the piece that belongs here, in the node, is the local NVMe SSD tier. Each training node typically carries local NVMe as a fast scratch layer distinct from the networked filesystem, used for three jobs: checkpoint staging (writing multi-terabyte checkpoints to local flash first, so GPUs stall for seconds rather than minutes while the data drains to remote storage in the background — the asynchronous checkpointing of Part XVII), dataset caching (keeping hot data shards local to avoid re-fetching them over the network every epoch), and spilling (overflow when host DRAM is exhausted). It slots into a clean hierarchy — GPU HBM (multiple TB/s) → host DRAM (hundreds of GB/s) → local NVMe (single-digit GB/s per drive, tens aggregated) → networked/object storage over the NIC. And it has its own direct path: GPUDirect Storage creates a DMA channel straight from NVMe into GPU memory, bypassing the host-DRAM bounce buffer exactly as GPUDirect RDMA does for the network — NVIDIA’s own framing notes this can lift the effective storage-to-GPU ceiling from the ~50 GB/s of a host-mediated path toward ~200 GB/s when many drives and NICs feed the GPU directly ^[454]. Try it: Checkpoint staging to NVMe; HBM to NVMe ladder.

A conventional NIC moves packets between wire and host DRAM. A SmartNIC adds programmable packet processing — matches, actions, rewrites — but still fundamentally runs at the direction of the host. A DPU/IPU crosses one more line: it embeds general-purpose CPU cores of its own, boots its own operating system, and runs infrastructure code (virtual switching, storage protocol termination, security policy, telemetry) that the host cannot see or interfere with, on the tenant-untrusted side of a hardware trust boundary.

NVIDIA’s product line makes the progression concrete. ConnectX-7 and ConnectX-8 are SmartNICs at 400 Gb/s and 800 Gb/s respectively — programmable but stateless-ish. BlueField-3, by contrast, integrates 16 Arm Cortex-A78 cores at 2.0 GHz with 8 MB of shared L3, 32 GB of on-board DDR5, 22 billion transistors, and hardware accelerators for RoCEv2 congestion control, NVMe-over-Fabrics, and line-rate IPsec/TLS/MACsec crypto — all at 400 Gb/s per port, programmable through the DOCA SDK, and equipped with a hardware root of trust and secure boot^[49][50].

Every hyperscaler has now built or bought an equivalent. AWS Nitro, first shipped in the C3 generation in 2013 and completed with C5 in 2017, is the origin of the pattern: a family of Nitro Cards handling VPC networking, EBS remote storage, and local NVMe, plus a Nitro Security Chip acting as an immutable hardware root of trust that validates every other component’s firmware at boot. Werner Vogels notes the Xen-era hypervisor consumed roughly 30 % of instance resources for management overhead; Nitro’s hardware offload recovers essentially all of it, delivering near-bare-metal performance while shrinking the software trust boundary to a minimal purpose-built hypervisor^[51][52].

Google’s C3 machine series runs on the Intel IPU E2000 (codename Mount Evans), an ASIC co-designed by Google Cloud and Intel that pairs up to 16 Arm Neoverse N1 cores at 3 GHz with a P4-programmable packet-processing pipeline and offloads 200 Gb/s of VPC packet processing off the host Xeons^[53]. Google’s broader host-offload stack, Titanium, layers on top of the IPU to accelerate networking, storage, and security for both C-series and AI-optimized instances^[54]. AMD’s Pensando Salina 400, announced October 2024, targets the same envelope: 16 Arm Neoverse N1 cores, 232 P4 MPU engines, up to 128 GB DDR5, and dual 400 GbE ports^[55].

The line between a DPU and a full x86 server is genuinely blurring — a BlueField-3 or Salina 400 has more Arm cores, more DRAM, and more crypto throughput than many production edge servers in 2020.

The offload story has four load-bearing pieces, and they map cleanly onto the four things a host CPU is worst at doing while also feeding a 700–1200 W accelerator:

RDMA and congestion control. RoCEv2 (RDMA over Converged Ethernet v2) and InfiniBand both let a NIC DMA directly into remote memory — including remote GPU memory via GPUDirect RDMA (Part X, Section 40) — bypassing kernel, TCP stack, and CPU. End-to-end latency drops from roughly 50 μs on TCP/IP to about 5 μs on RoCEv2 and 2 μs on InfiniBand^[56]. But RoCEv2 requires careful congestion control (PFC + ECN + DCQCN) to stay lossless, and doing that congestion loop in host software would burn cores at 400 Gb/s. The DPU runs it in hardware.

Storage — NVMe over Fabrics. The DPU presents remote NVMe drives to the host as if they were local PCIe NVMe devices, terminating NVMe-oF (over RoCE or TCP) in silicon. For AI clusters this collapses the checkpoint-and-dataset tier of Part X into a single virtual-drive abstraction, at wire-rate throughput and without CPU involvement.

Tenant isolation and rootless multi-tenancy. Because the DPU sits on the untrusted side of the PCIe boundary and runs its own OS, the cloud operator can enforce virtual-networking policy, encryption, and telemetry entirely off-host — the tenant’s root, and even the tenant’s hypervisor, cannot reach it. This is the architectural foundation of AWS’s confidential-computing posture and Google’s equivalent on Titanium: the infrastructure trust boundary is a physical chip, not a software layer.

Encryption in flight and collective bootstrap. Line-rate IPsec, TLS, and MACsec at 400 Gb/s are handled by dedicated crypto engines on the DPU; a host CPU asked to encrypt 400 Gb/s of AES would need dozens of cores just for that. And for collectives, the DPU handles connection setup, queue-pair state, and multicast-tree bootstrap — the paperwork of an all-reduce, so the GPU can spend its cycles on the arithmetic.

The through-line: GPU cycles are the most expensive silicon in the datacenter by a wide margin, and spending them on network-stack, storage-stack, or crypto-stack code is a categorically bad trade. The DPU exists to make that trade impossible.

For two decades InfiniBand was the default fabric for HPC and, by inheritance, for early AI clusters — sub-500 ns switch latency, credit-based lossless flow control, a centralized subnet manager. NVIDIA’s acquisition of Mellanox in 2019 made that inheritance a single-vendor stack. The industry response, mirroring the UALink story in Part X, was to build an open Ethernet alternative — the Ultra Ethernet Consortium (UEC), hosted under the Linux Foundation’s Joint Development Foundation. UEC released Specification 1.0 on 11 June 2025, a 560-page document spanning transport, congestion control, RDMA semantics, link and PHY layers, and network security, developed by AMD, Broadcom, Cisco, HPE, Intel, Meta, Microsoft, and Oracle among others^[57][58]. Try it: NVLink vs InfiniBand hop.

The most consequential departure from classical Ethernet — and from InfiniBand — is multi-path packet spraying. Traditional ECMP hashes a flow’s 5-tuple to a single path, which for AI workloads is disastrous: an all-reduce collective is a handful of huge, long-lived elephant flows, and hashing them collides them onto shared links while other links sit idle. UEC’s Ultra Ethernet Transport (UET) instead sprays the packets of a single flow across every available path simultaneously, doing per-packet rather than per-flow load balancing; the receiving NIC reassembles the resulting out-of-order stream in hardware before the RDMA verb completes, so applications never see the reordering^[59][60].

NVIDIA’s proprietary answer, Spectrum-X, builds the same idea into its Spectrum-4 Ethernet switches and BlueField-3 SuperNICs: switches select the egress port with the minimum queue occupancy on a per-packet basis, the SuperNIC reorders on ingress, and NVIDIA reports up to 95 % effective bandwidth under load — versus the 60 % or so typical of ECMP-based clusters at scale — and a 1.6× improvement over conventional Ethernet at the Israel-1 supercomputer deployment^[61]. UEC 1.0 references REPS (Recycled Entropy Packet Spraying) as one candidate load-balancing algorithm and pairs spraying with a modern telemetry-driven congestion-control scheme; both stacks agree the endpoint NIC must do the reassembly^[62].

This is where DPUs and the fabric-standards story converge: packet spraying only works because the endpoint has enough silicon to reassemble tens of millions of out-of-order arrivals per second at line rate. Neither a legacy NIC nor a host CPU can do that at 400 Gb/s; a BlueField-3-class SuperNIC or a UEC-conformant Pollara-400-class part can. The fabric standard and the offload silicon co-evolved, and by 2026 they are effectively a single design point: adaptive-routed, sprayed, hardware-reassembled Ethernet or InfiniBand, with the DPU as the load-bearing endpoint.

The strategic implication mirrors Part X and Part XI: as AI clusters grow past tens of thousands of GPUs and the fabric becomes as much of a bottleneck as the accelerator, whoever controls the DPU controls a large fraction of the total system cost, power, and — through the trust boundary the DPU enforces — the multi-tenant security model of the cloud built on top of it. That is why every hyperscaler has now designed one.

CXL is not one protocol but three multiplexed onto the same PCIe PHY. CXL.io is essentially PCIe: discovery, configuration, DMA, interrupts — the mandatory baseline every CXL device speaks. CXL.cache lets a device (typically an accelerator) coherently cache host memory, so a smart NIC or FPGA can hold a copy of a host cache line and participate in the CPU’s coherency protocol. CXL.mem inverts that: it lets the host use device-attached memory as its own cacheable memory, mapped into the CPU’s physical address space and reachable with ordinary loads and stores. Devices declare which subset they implement: Type-1 (accelerator with cache, no local memory — CXL.io + CXL.cache), Type-2 (accelerator with its own memory that the host can also see coherently — all three protocols), and Type-3 (a pure memory expander — CXL.io + CXL.mem). Type-3 has driven every shipping product so far ^[63].

Versions matter because each one changes what topologies are legal. CXL 1.1 (2019) rides PCIe 5.0 and permits only a single direct-attached device per root port — memory expansion, nothing more. CXL 2.0 (2020) adds a single layer of switching, hot-plug, and the crucial notion of memory pooling: one Type-3 device can be partitioned into logical devices and each partition dynamically assigned to a different host, so stranded DIMMs in one server become usable capacity for another. CXL 3.0 (2022) moves to the PCIe 6.0 PHY (64 GT/s, PAM4), allows multi-level switching and full fabric topologies, and — for the first time — permits genuine memory sharing, where multiple hosts coherently access the same region. CXL 3.1 (Nov 2023) adds Port-Based Routing to scale the fabric beyond a single switch domain and introduces the Trusted Execution Environment Security Protocol (TSP) for encrypted CXL memory, plugging into the same confidential-computing story as Part XXII ^[64].

Three distinct problems, one wire. The first is stranded memory: a hyperscale fleet provisions DRAM for each server’s worst-case working set, but the fleet-average utilization sits far below that, so a large fraction of every rack’s DIMMs is idle at any given moment. Meta’s own characterization of its production fleet found that a substantial fraction of allocated DRAM was effectively cold, and that a CXL-attached slow tier plus intelligent page placement recovered nearly all of the performance of an all-local baseline: their TPP paper (Transparent Page Placement, ASPLOS 2023) reports that a tiered system with a CXL far tier came within 1% of an ideal all-local configuration, ran 18% ahead of stock Linux, and beat NUMA Balancing and AutoTiering by 5–17%; most of the patches have shipped in Linux since v5.18 ^[65].

The second is the memory bandwidth vs. capacity tradeoff inside a single server. A modern two-socket host is limited to twelve or sixteen DDR5 channels, and past a certain point adding capacity means adding channels the CPU doesn’t have. CXL.mem lets an accelerator or CPU bolt on additional bandwidth in the form of extra memory controllers behind a Type-3 device, at the cost of one extra hop of latency. Marvell’s Structera line, announced July 2024 and sampling in Q4 2024, splits this into two SKUs: Structera X is a plain memory-expansion controller (up to 12 DDR4 modules per device, targeting DIMM-recycling deployments), while Structera A integrates server-class Arm cores and multiple memory channels as a near-memory accelerator for embedding lookups in recommendation models and similar bandwidth-bound workloads ^[66].

The third — the one changing fastest as of 2026 — is KV-cache offload for LLM inference. Part XVI’s decode-phase analysis showed that a serving node’s throughput is often capped not by its GPU’s FLOPs but by how many concurrent request states (KV caches) fit in HBM. Offloading cold or paged-out KV blocks to CXL memory instead of NVMe changes the arithmetic: at ~200 ns access latency versus ~50 μs for NVMe, a CXL far tier is closer to HBM than to storage. Astera Labs, running their Leo CXL controllers on Azure M-series VMs, reports 67% lower latency, 2× concurrent LLM instances, and 30% larger batch sizes when KV cache is stored on CXL rather than on SSD-backed alternatives ^[67]. A joint XConn/MemVerge demonstration at SC25 measured over 5× improvement versus SSD-based or RDMA-based KV-cache sharing across GPUs ^[68], and a late-2025 arXiv paper on 1M-token inference argues that CXL-backed KV storage expands the effective KV budget by 4–8× versus GPU-only serving ^[69]. Try it: KV cache: CXL vs NVMe.

Because the three big DRAM vendors — Samsung, SK Hynix, Micron — have effectively withdrawn from developing their own CXL controllers and are shipping modules built around third-party silicon, the ASIC market has consolidated onto a handful of fabless designers ^[70]. Astera Labs' Leo is the first-mover for CXL 2.0 pooling silicon, has completed interop with DDR5-5600 RDIMMs from all three memory vendors, and is the controller behind Microsoft’s Azure M-series CXL rollout ^[71]. Marvell’s Structera A/X and MangoBoost’s controllers cover the near-memory-acceleration and expansion niches respectively. Panmnesia is the current front-runner on CXL 3.x switching: it announced sampling of a PCIe 6.0 / CXL 3.x fabric switch, claims sub-100 ns latency for its 3.1 IP, and has publicly demonstrated a CXL-GPU prototype at OCP that attaches CXL memory or CXL-SSDs directly to a GPU’s PCIe root complex, reporting roughly 6× lower inference latency than a storage- or RDMA-based baseline on RAG and DLRM workloads ^[72]. Try it: Figure 68, interactive.

Actual fleet deployment lags the specification. Most shipping hardware in 2025 is CXL 2.0 Type-3 memory expansion: single-host capacity extension, no switching, no true pooling across hosts in production. Meta’s TPP work is deployed in Meta’s own fleet, but hyperscaler-wide CXL 2.0 pooling remains largely at the pilot stage; CXL 3.0 switches are sampling as of late 2025 rather than shipping at volume; and CXL 4.0 (announced Nov 2025, doubling bandwidth to 128 GT/s on the PCIe 7.0 PHY) is still a paper standard ^[73]. The gap between what the specs allow and what a hyperscaler will actually rack up is the usual one: rack-scale pooling only pays off when the switch, controller, memory module, hot-plug software, and page-placement OS all mature together, and the ecosystem is still catching up to the standard.

The reason CXL is a distinct tier — rather than a slower DDR or a faster NVMe — is a specific latency band. Local DDR5 delivers sub-100 ns round-trip latency; a CXL memory controller adds roughly 70–150 ns over local DDR, plus a few tens of nanoseconds per retimer as the device sits further from the CPU, for a typical CXL-attached DIMM access in the 170–250 ns range — comparable to reaching memory on a second CPU one NUMA hop away, and cheaper than the ~350 ns Intel Optane persistent memory once cost ^[74]. HBM3 by contrast is at multi-TB/s bandwidth, though its load latency from an SM is several hundred nanoseconds (roughly 300–500 ns) — GPUs hide it with parallelism rather than avoid it — and NVMe is at tens of microseconds. The right way to think about it is as an extra rung on the ladder: HBM (fast/tiny) → DDR (fast/small) → CXL-attached (near/medium) → CXL-pooled across a switch (medium/large) → NVMe (slow/huge), where each rung is roughly one to two orders of magnitude larger and slower than the one above. Try it: The memory latency ladder; CXL vs local DRAM; HBM latency vs DRAM; CXL latency budget.

The competing approach is proprietary rather than protocol-level: NVIDIA’s NVLink-C2C, introduced earlier in this Part, gives a Grace CPU and Hopper (or Blackwell) GPU a 900 GB/s cache-coherent link over which they share a single address space — 96 GB of HBM3 on the GPU side plus 480 GB of LPDDR5X on the CPU side, all reachable with plain loads and stores ^[75]. That is 7× the bandwidth of PCIe 5.0 x16 at roughly 1.3 pJ/bit, and it delivers some of what a CXL 3.x fabric promises — the "GPU can address CPU memory as if it were its own" property — but only within the specific superchip boundary NVIDIA sells. CXL is the open answer to the same question, one step slower and one order of magnitude more flexible: any host, any device, any vendor, as long as they speak the same three protocols on the same PCIe PHY. Try it: NVLink-C2C vs PCIe 5.0.

The through-line back to Part VII’s memory wall is direct: HBM stacks can only get so tall, DDR channels can only get so wide, and NVMe is a different kind of storage entirely. CXL doesn’t lift the memory wall so much as add a shelf halfway up it — a place to keep KV cache, embedding tables, and cold pages that would otherwise crowd out active weights in HBM or stall waiting on flash. Whether that shelf ends up load-bearing at hyperscale, or whether NVLink-style proprietary fabrics keep swallowing the same problem inside a single vendor’s SKU boundary, is the open question the 2026–2027 hardware cycle is designed to answer.

The problem NoCs solve is the classic all-to-all wiring problem. A crossbar that lets agents talk to others has wire cost and area, which becomes untenable well before reaches the hundred-plus SMs on a modern GPU or the hundreds of thousands of processing elements on a wafer-scale part. A shared bus is cheap in wires but serializes every transaction and does not scale in bandwidth. The NoC replaces both with a switched fabric: each tile has one router with a small, fixed number of ports (four to eight, typically), and packets hop router-to-router until they reach the destination. The cost of the fabric grows linearly with the number of tiles rather than quadratically, while aggregate bandwidth grows with the number of links, not with any single wire’s speed ^[77].

The two figures of merit that dominate NoC design are bisection bandwidth — the aggregate bandwidth across the tightest cut that splits the fabric in half, which upper-bounds any workload that has to shuffle data across the whole chip (an all-reduce, a transposed matmul, a shuffle-heavy attention layer) — and energy per bit moved per millimeter, since on-chip wires above a few hundred microns burn a surprising amount of power relative to a floating-point operation. At 3nm, moving a bit across a millimeter of on-chip wire costs on the order of one picojoule, and a full 32-byte cache line across half a die can cost more energy than the FLOP it enables ^[78]. Every topology choice below is, ultimately, a trade between these two numbers plus area. Try it: Bisection and hop counts.

The classical choices span a spectrum from cheap-and-slow to expensive-and-fast:

Given a topology, a router needs a routing algorithm — a rule for choosing the next hop. The simplest is XY (dimension-order) routing on a mesh: a packet travels along X until it reaches the destination column, then along Y. It is deterministic, cheap, and provably deadlock-free on a mesh because it never introduces a cyclic dependency between channels; it is also oblivious to congestion, so hot spots can starve neighbors. Adaptive routing lets a router choose among several minimal (or non-minimal) paths based on local queue occupancy, improving throughput under skewed traffic at the cost of extra buffering and a genuine risk of deadlock unless the algorithm is carefully constructed (typically by allocating extra virtual channels for "escape" routes) ^[80]. Try it: Topology zoo, interactive.

Flow control is how a router decides whether it can send the next flit (flow-control digit — the smallest unit the network moves in one cycle) into its neighbor. The dominant scheme in modern NoCs is credit-based flow control: the sender maintains a running count of free buffer slots at the receiver, decrements on send, and increments when the receiver returns a credit. It guarantees no packet is ever dropped due to buffer overflow (unlike lossy datacenter Ethernet, which relies on end-to-end retransmit), which is essential for a fabric that carries cache-coherence traffic where retransmit would violate the memory model.

Virtual channels (VCs) are the other load-bearing invention. Each physical link is time-multiplexed among several logical channels, each with its own buffer and credit counter. VCs serve three simultaneous purposes: they break head-of-line blocking (a stalled packet on VC 0 doesn’t hold up VC 1); they provide the extra channel dependencies needed to prove deadlock-freedom on topologies that would otherwise deadlock (rings and tori being the canonical case, where a single VC will always have a cyclic dependency around the ring); and they let coherence protocols run on separate virtual networks — request, response, and snoop each on their own VC — so a full snoop response cannot back-pressure a request queue and cause protocol-level deadlock. The classical result behind almost all of this is Dally’s 1992 paper introducing virtual channels as the general mechanism for deadlock avoidance in k-ary n-cube networks, which is still the direct ancestor of every commercial NoC on the market.

Almost all modern NoCs are packet-switched with wormhole or virtual-cut-through routing: a packet is split into flits, the head flit reserves the route, body flits stream through, and the tail flit releases the resources. Wormhole switching lets a packet occupy multiple routers at once (like a worm), minimizing buffer requirements at the cost of holding channels across hops. Circuit-switched NoCs, which reserve an end-to-end path before sending data, are rare in general-purpose accelerators but reappear in photonic NoCs (below), where setting up an optical path amortizes across a large burst.

Each vendor’s on-chip fabric reflects the workload and physical constraints of its target. NVIDIA’s H100 introduced an SM-to-SM network inside each GPC (Graphics Processing Cluster) that lets a thread block cluster — a group of thread blocks guaranteed to be co-resident — directly access another SM’s shared memory via TMA (Tensor Memory Accelerator) DMA, with hardware asynchronous transaction barriers for synchronization. NVIDIA reports up to a 7× latency improvement for cluster-scale synchronization versus the pre-Hopper global-memory round-trip ^[81]. The topology is a mesh within each GPC and a hierarchical fabric between GPCs and L2 slices.

AMD’s MI300 Infinity Fabric is explicitly a two-level packet-switched network: XCDs and CCDs (compute chiplets) talk to Coherent Master (CM) blocks on the underlying I/O dies, which encapsulate requests as flits and route them via a switch fabric that spans four IODs (with cross-IOD hops over an active interposer layer) to Coherent Slaves (CS) that own the HBM3 memory controllers and probe filters. Each CS also owns a co-located 2 MB Infinity Cache slice, so the NoC is simultaneously the coherence substrate and the LLC access path. Each octa-core Zen 4 CCD attaches with a wide GMI link at roughly 71 GB/s read bandwidth, and cross-socket Infinity Fabric links deliver 128 GB/s bidirectional ^[82][83].

Intel’s Meteor Lake introduces an explicit NoC inside its SoC tile — a low-latency, coherent fabric that connects the tile’s high-bandwidth devices — sitting atop a passive Foveros base die whose Foveros Die Interconnect (FDI) achieves under 0.3 pJ/bit for tile-to-tile hops, roughly an order of magnitude better than off-package PCIe on the same node ^[84]. Apple’s M-series does something similar in spirit but different in mechanism: two M-Max dies join via UltraFusion, an interposer-based die-to-die fabric with more than 10,000 signal lines and, as of the M5 Ultra generation, over 4.4 TB/s of inter-die bandwidth — enough that the OS sees a four-die M5 Ultra as one processor with a unified 512 GB memory pool ^[85].

Tenstorrent Wormhole and Blackhole are pure NoC-first designs: an 80-tile grid of Tensix compute cores on a 10×12 physical layout, connected by two 2D torus NoCs running in opposite directions (NoC 0 clockwise, NoC 1 counter-clockwise). Each Tensix core has two NoC interfaces so it can drive both planes simultaneously, and Ethernet links at the edges extend the same NoC out to neighboring chips, giving the illusion of one arbitrarily large mesh across many devices ^[86]. Cerebras takes the mesh idea to its limit: on WSE-3, ~900,000 processing elements sit on a single wafer in a 2D rectangular mesh, each PE routing 32-bit wavelets to its four immediate neighbors at an aggregate core-to-core bandwidth of 215 PB/s. Every wire is under a millimeter long, which is why the fabric can carry that much bandwidth at a power budget that still fits in a chassis ^[87]. Try it: Wormhole 10×12 torus; Wafer-scale mesh figures.

Once a chip has more than a handful of caches, keeping them consistent becomes an on-chip networking problem, not a memory problem. Snoop-based MESI protocols broadcast on every miss and are viable up to roughly 16 caches; beyond that, a directory-based variant — where a directory (often distributed across LLC slices) tracks which caches hold each line — is the only workable approach, and it is entirely a NoC workload. Each coherence transaction becomes several message classes (request, forward, response, invalidate, acknowledgment) traveling over the NoC on separate virtual networks so that no class can back-pressure another into protocol deadlock. Modern implementations extend MESI to MOESI or MESIF to reduce writeback traffic; the GEMS-style reference implementation of a MESI directory protocol requires on the order of 30 stable and transient states, all of which are ultimately just descriptions of NoC message ordering constraints ^[88]. This is why AMD’s MI300 architecture calls the coherence agents "Coherent Master" and "Coherent Slave" and treats them as first-class endpoints on Infinity Fabric — the coherence protocol is a NoC application.

The scaling law that governs on-chip networks is uncomfortable. Bisection bandwidth on a 2D mesh grows only as while compute grows as , so at wafer scale a plain mesh is bandwidth-starved for any workload that must cross the chip in one direction (a global all-reduce, a transposed matmul). Tori help by a constant factor. Fat trees help asymptotically but at large area cost. And the underlying wires themselves are not getting cheaper: on-chip global wire capacitance scales roughly with wire length, so as die sizes grow toward the reticle limit, energy per bit for a long cross-die hop stays flat or worsens even as transistors shrink. Try it: √N bisection scaling.

The frontier response is to move the long hops off electrons entirely. Photonic NoCs replace long on-chip wires with silicon waveguides carrying light, exploiting wavelength-division multiplexing to run dozens of independent channels over a single waveguide. Lightmatter’s Passage M1000, unveiled at Hot Chips 2025 and shipping in 2025, is a 4,000 mm² photonic interposer that sits under a multi-die compute stack and provides 114 Tbps of total optical bandwidth across 1,024 serial channels at 56 Gbps in each direction, with 256 fibers attached to the package — effectively an all-optical NoC substrate that a conventional electrical die can plug into as its scale-out fabric ^[89]. Whether photonic NoCs displace electrical fabrics inside a single die is a longer bet — the cost of on-die lasers, thermal sensitivity of ring resonators, and reliability of integrated waveguides are all still open — but at the multi-die and rack-adjacent scale, the technology is now shipping rather than hypothetical ^[90]. Either way, the NoC — electrical, hybrid, or photonic — remains what determines whether the compute behind it can be fed, and it is one of the few layers in the AI hardware stack where the classical computer-architecture textbooks of the last two decades still say most of what needs to be said. Try it: Passage M1000 bandwidth arithmetic.

Google’s Tensor Processing Unit is the oldest and by far the most battle-tested custom AI chip, now in its seventh generation. Its core is the systolic-array MXU (Matrix Multiply Unit) described in Section 4.3, but the architecture that makes TPUs formidable at scale is the interconnect: chips are wired into a 3D torus over Google’s Inter-Chip Interconnect (ICI), and pods are stitched together with Optical Circuit Switching (OCS) — a reconfigurable optical fabric that can dynamically re-slice the topology and, crucially, route around a failed chip or tray without taking the pod down, which matters enormously for month-long training runs ^[228][229]. The current generation, TPU v7 (Ironwood), delivers per chip roughly 4,614 TFLOPS of FP8, 192GB of HBM at 7.4 TB/s, and 1.2 TB/s of ICI bandwidth, and scales to a 9,216-chip pod sharing 1.77 PB of HBM ^[230]. The lineage in between — v5e (efficiency-tuned), v5p (performance-tuned), and v6e (Trillium) — is what most of Google’s own model training (Gemini) has run on for years, which is the real proof point: TPUs are not a science project but a production training platform at frontier scale.

The scale of external commitment is now enormous. In October 2025 Anthropic announced it would expand to up to one million TPUs, with well over a gigawatt of capacity coming online in 2026 — a deal reported in the tens of billions of dollars ^[231]. (A separate, later Anthropic–Google–Broadcom arrangement describes multiple additional gigawatts of next-generation TPU capacity starting in 2027; the two should not be conflated ^[232].)

Amazon splits its silicon by task, exactly along the training/inference divide of Part IV: Trainium for training, Inferentia for inference. The building block is the NeuronCore, and the scale-up/scale-out split mirrors NVIDIA’s: NeuronLink is AWS’s chip-to-chip scale-up fabric (analogous to NVLink), while the Elastic Fabric Adapter (EFA) handles scale-out across servers. Trainium2 offers 96GB of HBM per chip, assembled into a 64-chip UltraServer; Trainium3 (announced late 2025, AWS’s first 3nm AI chip) roughly doubles that to 2.52 FP8 PFLOPS and 144GB of HBM3e per chip at 4.9 TB/s, adds MXFP8/MXFP4 and structured sparsity, and scales to 144-chip UltraServers ^[233][234], all programmed through the Neuron SDK. The flagship deployment, Project Rainier, is an AWS-built cluster of more than 500,000 Trainium2 chips brought online in under a year specifically for Anthropic, roughly five times the training capacity Anthropic had previously — with AWS’s UltraCluster design targeting up to a million Trainium chips ^[235].

Microsoft’s Maia 100 (announced 2023) and Meta’s MTIA line are narrower bets, both aimed primarily at inference rather than training. Maia 100 shipped but, by most accounts, has served largely internal and non-production workloads, and its successor (codename Braga / Maia 200) has reportedly slipped its mass-production timeline into 2026 ^[236]. Meta’s MTIA — despite standing for "Meta Training and Inference Accelerator" — has in every shipped generation (v1 on TSMC 7nm, v2 on 5nm) been an inference part, purpose-built for the recommendation and ranking models that dominate Meta’s own compute bill, with a published roadmap of further inference-optimized generations ^[237][238]. Both illustrate the same point from the opposite end of Google’s success: designing a competent inference ASIC is achievable, but building a training-class platform that frontier labs will actually trust with a flagship run — as Google’s TPU and AWS’s Trainium now are — is a much higher bar that most in-house programs have not yet cleared.

The specific pairings map the industry:

Two names recur in that table that most people outside the industry never hear: Broadcom and Marvell. They are the custom-silicon design partners — between them they enable the overwhelming majority of hyperscaler AI ASICs — and what they provide is the hard, non-differentiating engineering a hyperscaler doesn’t want to build in-house: high-speed SerDes (the serializer/deserializer circuits that drive signals between chips), advanced-packaging integration, licensed interface IP, physical design, and the actual tape-out and foundry relationship. When Google or Amazon "designs its own chip," a company like Broadcom or Marvell typically does a large share of the physical implementation and shepherds it through TSMC ^[265][266].

That "arms dealer" role runs deeper than most public discussion captures, because these two companies also build much of the networking silicon that stitches AI clusters together — so they profit from the buildout twice over. Broadcom is the dominant merchant Ethernet-switch vendor: its Tomahawk line of switch ASICs is the workhorse of AI cluster fabrics (Tomahawk 5 at 51.2 Tb/s of switching capacity, and the 2025 Tomahawk 6 at 102.4 Tb/s on a 3 nm process — the first single-chip switch at that bandwidth), while its Jericho line (Jericho3-AI, Jericho4) provides the deep-buffer routing silicon used to connect AI fabrics across longer distances and between datacenters. Broadcom’s co-packaged-optics switch "Bailly" (Section 56) and its SerDes leadership round out a franchise whose AI-related revenue reached roughly $20 billion in 2025, alongside a custom-XPU business publicly described as serving around half a dozen large customers (reported to include Google, Meta, Anthropic, and OpenAI) ^{[265][352][353]}. Marvell plays the same dual game from second position: a custom-compute platform behind accelerators like AWS Trainium and Microsoft Maia, its own Teralynx Ethernet switch silicon (Teralynx 10 at 51.2 Tb/s, T100 at 102.4 Tb/s), a strong 800 Gb/s optical-DSP and SerDes franchise, and — via its late-2025 acquisition of Celestial AI (Section 56) — a photonic-interconnect play aimed at scale-up connectivity between accelerators ^[354][355]. The strategic takeaway is that "custom silicon" has not disrupted the merchant-chip vendors so much as handed them a second business: the hyperscalers design the accelerator’s soul, but Broadcom and Marvell supply the body — the interconnect IP, the switch chips, and the physical implementation — for a large fraction of the non-NVIDIA silicon in the industry.

Two more oligopolies sit behind every chip. EDA (electronic design automation) — the software used to design, simulate, and verify chips — is dominated by just three firms: Synopsys, Cadence, and Siemens EDA, together well over 90% of the market, with individual tools (Synopsys PrimeTime for timing signoff, Siemens Calibre for physical verification) holding near-monopoly shares of their niches. No leading-edge chip is designed or verified without them ^[266]. IP licensing is the adjacent business: rather than design every block from scratch, chip designers license pre-built, pre-verified building blocks — Arm licenses CPU cores (Part V), while Synopsys and Cadence license interface IP (PCIe, DDR/HBM controllers, SerDes, UCIe) that gets integrated into a larger design.

Finally, HBM memory — the stacked DRAM that Part III showed is the lifeblood of every accelerator — comes from just three suppliers: SK Hynix (the clear leader and NVIDIA’s primary HBM4 source), Micron, and Samsung. HBM has been a genuine supply bottleneck, with capacity reported sold out through 2026, which makes the memory makers as strategically important to the AI buildout as the logic foundries — a chip you can’t feed is a chip you can’t ship ^[267][268].

The Nyquist frequency of a 224 Gb/s PAM4 lane sits near 56 GHz. At those speeds copper is dominated by skin-effect loss, dielectric loss, and reflections off every via, bump, and connector on the path; a single misbehaving pad can eat the entire link budget ^[91]. In practice this collapses the useful reach of passive copper Direct Attach Cables (DACs) to roughly one meter at 224 Gb/s per lane — barely enough to stay inside a single rack, and not enough to leave one. Active Electrical Cables (AECs) with a retimer chip in-line stretch that to about 3–5 m at a power cost, but any coherent scale-up domain wider than one rack is now an optical problem by physics, not by preference ^[92]. Try it: 224G PAM4 Nyquist.

The scale-up domain used to sit safely inside copper’s reach: an NVLink-connected 8-GPU HGX baseboard, or Blackwell’s 72-GPU NVL72 rack, is a copper-cabled backplane precisely because everything is within about two meters of everything else. But the next step — a 576-GPU or larger coherent domain, spanning several racks that must share memory as if they were one machine — puts the far GPU at ten or fifteen meters, which is optical territory. This is why every serious "bigger NVLink" proposal in 2026 (NVIDIA’s rumored Rubin-class scale-up, Ayar Labs' UCIe-optical chiplet path, Lightmatter’s Passage, Celestial AI’s Photonic Fabric) is photonic at its core: the electrical version simply does not close timing.

A silicon photonics link is four building blocks welded together on a CMOS-compatible die.

The laser is an external continuous-wave source, usually a distributed-feedback (DFB) diode at ~1310 nm or a comb source producing eight or sixteen wavelengths on a fixed ITU grid. Every CPO product shipping in 2026 — NVIDIA’s, Broadcom’s, Ayar Labs' — keeps the laser off the hot main package, in a pluggable External Laser Source (ELS) module such as Lumentum’s ELSFP, precisely because laser reliability degrades sharply with temperature and lasers age faster than the CMOS around them ^[93].

The modulator imprints data onto that light. Two topologies dominate. Mach–Zehnder modulators split light into two arms, phase-shift one arm electrically, and recombine — robust and thermally tolerant but physically large (millimeters). Ring modulators (used by Ayar Labs, Lightmatter, and TSMC’s COUPE) are 10–20 μm resonators that shift their resonance wavelength under bias, giving ~100× smaller footprint and lower drive voltage at the cost of extreme temperature sensitivity — a ring drifts about 0.1 nm/K, which for a WDM channel spacing of ~0.8 nm means active thermal tuning is not optional ^[94].

The photodetector at the far end is a reverse-biased germanium-on-silicon diode that turns light back into current, followed by a transimpedance amplifier. And wavelength-division multiplexing (WDM, or its dense variant DWDM) packs 8–16 independent lanes onto a single fiber on different colors, so one physical fiber carries 400–800 Gb/s of aggregate traffic. Ayar Labs' SuperNova source and TeraPHY chiplet run 16 wavelengths and deliver 8 Tb/s bidirectional per chiplet ^[95].

The pluggable transceiver in an OSFP or QSFP-DD cage is the incumbent: a fully self-contained module (laser + modulator + DSP + photodetector) that plugs into the switch faceplate. An 800G 2xFR4 pluggable burns roughly 15 W and lands in the 15–20 pJ/bit efficiency band ^[96]. Linear-drive pluggable optics (LPO) rip out the DSP retimer entirely, leaning on the switch ASIC’s SerDes to drive the modulator directly; that cuts per-port power to 8–10 W, about 10–12 pJ/bit at 800G, but sacrifices reach and eye margin. Co-packaged optics take the last step: the optical engine (modulator, driver, TIA, photodetector) is bump-bonded next to the switch die on the same organic substrate or interposer, with only micrometers of electrical trace between them. Broadcom’s Tomahawk 5-Bailly reports 5.4 W per 800G of delivered bandwidth — a 65% reduction versus pluggables — and its 2026 Tomahawk 6 "Davisson" at 102.4 Tbps aims at 3.5 W per 800G port, roughly 70% below pluggable ^[97][98]. Try it: Pluggable 15 W in pJ/bit; Bailly 5.4 W per 800G; Davisson 3.5 W per 800G.

Try it: Link power and pJ/bit.

NVIDIA announced Quantum-X Photonics (InfiniBand) and Spectrum-X Photonics (Ethernet) at GTC in March 2025, both built on TSMC’s SoIC-X 3D hybrid-bonded silicon photonics engine — the first shipping product to stack a 6 nm electronic IC directly on a 65 nm SOI photonic IC via hybrid bonding ^[99]. Quantum-X ships 144 ports of 800 Gb/s on 200 Gb/s SerDes, liquid-cooled, in late 2025; Spectrum-X follows in 2026 at 128 or 512 ports of 800 Gb/s (100 or 400 Tb/s aggregate). NVIDIA claims 3.5× better power efficiency, 4× fewer lasers (because a shared ELS drives multiple engines), and 10× better link resiliency versus a pluggable-based reference design.

Broadcom’s Bailly (Tomahawk 5 + CPO) reached volume production in 2025 at 51.2 Tbps and — per Meta’s ECOC 2025 field data — hit an MTBF of 8.2 million hours in trial deployment, roughly 10× the pluggable-based comparison system, silencing the durability skepticism that dogged early CPO ^[100]. Tomahawk 6 Davisson doubles that to 102.4 Tbps with 200 Gb/s-per-lane CPO for 2026 shipment.

Ayar Labs took the different bet: rather than sell a switch, sell the optical I/O chiplet that any compute die can bolt on over UCIe. Their TeraPHY at 8 Tb/s bidirectional per chiplet, driven by the SuperNova 16-wavelength source, targets scale-up interconnect directly — replacing NVLink-class electrical fabrics with UCIe-optical inside the package ^[101]. Lightmatter’s Passage M1000, unveiled at Hot Chips 2025, goes further still: a 4,000 mm² active photonic interposer that carries 114 Tb/s of optical bandwidth, 256 fiber attachments, and 1.5 kW of power delivery under the compute dies it hosts — effectively a photonic replacement for the metal RDL layers of a conventional interposer ^[102]. Celestial AI’s Photonic Fabric — the technology behind Marvell’s $3.25 B acquisition closed February 2026 — puts 16 Tb/s per chiplet and claims 25× the bandwidth at 10× lower latency of first-generation CPO by disaggregating memory and compute across the fabric ^[103]. Try it: Passage: 57 vs 114 Tb/s.

The manufacturing enabler for all of the above is TSMC’s Compact Universal Photonic Engine (COUPE), which uses SoIC-X hybrid bonding to stack an electronic IC on a photonic IC with sub-micron pitch, and then drops the resulting stack onto a CoWoS interposer next to the compute die and HBM. TSMC has publicly targeted COUPE volume production in 2026 and full CoWoS-L integration on the same 2026 timeline, with claimed 5–10× power-efficiency and 10–20× latency improvements versus the prior generation of edge-of-package optical tiles ^[104][105]. Samsung and Intel Foundry are chasing, targeting 2028–2029 for equivalent turnkey CPO — a two-generation gap that concentrates 2026 CPO capacity at TSMC in the same way HBM concentrates at SK Hynix (Part XII).

A 51.2 Tbps CPO switch fans out through 64 to 128 MPO-class fiber terminations at the package edge. Broadcom’s Bailly uses 128 400 Gbps FR4 interfaces on single-mode fiber (SMF, OS2); NVIDIA’s Quantum-X Photonics runs 144 ports of 800 Gb/s, similarly on SMF. All shipping CPO in 2026 uses single-mode over multi-mode, because ring-modulator WDM channels and cluster-scale reach (500 m to 2 km) both demand SMF’s low dispersion and low loss ^[106]. MPO-16 and dual-MPO-12 connectors dominate the physical panel because they cram 12–16 fibers into a footprint the size of a fingernail; a CPO switch’s faceplate is therefore not the row of pluggable cages it used to be but a dense grid of MPO adapters with an ELSFP laser tray beside them.

The mechanical failure modes shift accordingly. A pluggable transceiver failure is a 30-second field replacement; a CPO engine failure means an in-warranty package swap, since the optical engine is bonded to the switch substrate. This is why every 2026 CPO product externalises the laser (the shortest-lived component) and increasingly moves toward pigtailed fibers with expanded-beam or physical-contact detachable connectors on the outside of the package — trading a small optical-loss budget for the ability to replace a single fiber jumper without touching the CPO engine ^[107].

The strategic question for 2026–2027 is not whether CPO switches replace pluggable switches — that transition is already underway and well-priced — but whether scale-up interconnect (NVLink, NVSwitch, UALink, Infinity Fabric) goes optical inside the package. Copper NVLink at 224 Gb/s per lane cannot leave the rack; a Rubin-class or MI400-class successor that wants a 512-GPU or 1024-GPU coherent domain has no choice but photonic. Ayar Labs, Lightmatter, and Celestial AI (now inside Marvell) are all positioned exactly on that bet, and NVIDIA’s public 2025 statement that Rubin’s scale-up will be "silicon photonics-based" — plus its equity stake in Ayar Labs — reads as internal validation ^[108]. If that bet lands, the 2027 flagship system’s block diagram looks less like "GPU pod plus NIC" and more like the picture in Figure 81: HBM and compute die in the middle, photonic tiles at the edges, and fibers — not copper — carrying every bit that leaves the package. The memory wall and the copper wall are the same wall, viewed from two sides; light punches through both.

The single biggest workload shift of 2025–2026 deserves its own treatment, because it changed what the hardware is being asked to do. For most of the deep-learning era, a model produced its answer in a fixed forward pass: you paid a roughly constant amount of compute per query. Reasoning models — OpenAI’s o1 and o3, DeepSeek-R1, and their successors — break that assumption. They are trained (largely by reinforcement learning) to generate a long internal chain of thought before answering — thinking on paper, checking their own steps, backtracking — and that thinking can run to tens of thousands of internal tokens the user never sees. Crucially, OpenAI disclosed that a model’s accuracy improves smoothly as it is allowed to spend more test-time compute: quality is now something you can buy at inference, not only at training. This is the inference-time (test-time) scaling law, and it sits alongside the training-time scaling laws of Part I as a second axis along which spending more compute reliably buys more capability ^[437][438].

The hardware consequences are direct and large. Reasoning shifts the load overwhelmingly onto decode — the sequential, memory-bandwidth-bound token-at-a-time phase — because the "thinking" is just an enormous number of generated tokens. It inflates per-request latency and compute by roughly one to two orders of magnitude, and it balloons the KV cache (a 30,000-token chain of thought can consume several gigabytes of KV cache for a 70B-class model on its own), which as the batching discussion of Part IV showed is precisely what limits how many requests a GPU can serve at once. The right unit of cost stops being cost per token and becomes cost per answer. And it is the main reason analysts now expect inference, not training, to dominate total AI compute by the end of the decade — reversing the field’s decade-long assumption that training was where the compute went. Every memory-bandwidth, KV-cache, and decode-optimization theme in this primer is amplified by the arrival of models that think before they speak ^[437][439]. Try it: 30,000-token chain-of-thought cache.

Getting power into a rack is only half the problem; getting the heat back out is the other, and it has quietly become one of the hardest constraints on deployment. Per-GPU power climbed from roughly 700 W on an H100 to about 1,000 W on a B200, and a GB200 NVL72 rack draws on the order of 120–132 kW — roughly ten times a conventional 10–15 kW enterprise rack. At the die itself, heat flux has passed the point where moving air can carry it away: an H100 dissipates around 86 W/cm², well above the ~50 W/cm² beyond which forced air stops being adequate ^[433][434]. Try it: H100 heat flux: 86 W/cm²; Compute H100 heat flux.

The response is a ladder of increasingly aggressive cooling, and the industry has been forced up it (Figure 47). Air cooling — fans and heatsinks — tops out around 15 kW per rack conventionally, perhaps 40 kW with aggressive design. Direct-to-chip liquid cooling (DLC) puts cold plates directly on the GPUs, CPUs, and NVSwitches, circulating coolant through them, and extends the ceiling into the ~60–120 kW-plus range; this is the mainstream approach for Blackwell-class systems, and NVIDIA’s GB200 NVL72 ships liquid-cooled by design with no air-cooled variant at all — heat leaves through the cold plates to a Coolant Distribution Unit (CDU) that hands it to the facility water loop. Immersion cooling, where whole servers are submerged in a dielectric fluid, pushes density highest of all (~100–250 kW-plus per rack) but is operationally harder; its two-phase variant, which boils the fluid off the chip, has been set back by 3M’s exit from the fluorochemicals that made those coolants ^{[433][434][435]}. Try it: Place a rack on the ladder; Cold-plate coolant flow.

The through-line is that cooling, not silicon, is now frequently what gates how dense a rack can be — and it feeds straight back into the efficiency story, since better cooling is what allows the tight PUE and high density that make the joules-per-token math work. It also only intensifies from here: NVIDIA’s projected Rubin Ultra "Kyber" rack is expected around 600 kW, roughly five times a current Blackwell rack, which no amount of air could ever hope to cool ^[436]. Try it: A 600 kW rack’s cooling.

Power Usage Effectiveness (PUE, recap from Section 58) captures only facility overhead — total power divided by IT power — and by design ignores water, carbon, and any heat that leaves the building for productive use. The Green Grid published three companion metrics through the early 2010s, all now codified under ISO/IEC 30134, to close those gaps ^{[109] [110]}.

Water Usage Effectiveness (WUE) is annual site water consumption (in litres, for humidification and evaporative cooling) divided by annual IT energy (in kWh). An air-cooled facility can report a site WUE near 0 L/kWh; a facility on evaporative cooling in a warm climate typically sits between 1.4 and 2.5 L/kWh. Equinix reports a fleet-wide 2024 site WUE around 0.91 L/kWh, rising to 1.41 L/kWh at its evaporative sites; the industry average sits closer to 1.8 ^[111]. Source WUE additionally counts the water evaporated by the power plants generating the site’s electricity — often several times larger than site WUE and the number the "Making AI Less Thirsty" paper argues is the honest one to report ^[112]. Try it: Water use from WUE.

Carbon Usage Effectiveness (CUE) multiplies total facility energy by the grid’s carbon intensity (kgCO2e/kWh) and divides by IT energy. A facility on Icelandic geothermal-hydro power can report CUE approaching zero; the same hardware on the Louisiana grid, one of the most fossil-heavy in the United States, reports CUE ten to twenty times higher. Because CUE is essentially PUE × grid carbon intensity, the metric captures the fact that where you build matters more than how well you run it. Try it: CUE from PUE and grid.

Energy Reuse Effectiveness (ERE), defined by ASHRAE and codified in ISO/IEC 30134-6, subtracts reused energy (heat exported to district heating, greenhouses, industrial processes) from total facility energy before dividing by IT energy — so a site can report ERE below 1.0 if it exports more heat than its non-IT overhead consumes. Nordic operators routinely publish ERE numbers in this range: Meta’s Odense (Denmark) campus feeds around 100,000 MWh/year of waste heat into the local district-heating network, and Stockholm Data Parks builds its business model on it ^[113]. Try it: ERE with heat reuse.

The 2025 Uptime Institute Global Data Center Survey reports the sobering industry-wide picture behind these definitions: PUE is tracked by 74% of operators (fleet average 1.54, essentially flat for six years), but water is tracked by only 47% and renewable-energy consumption by only 42% ^[114]. The metrics exist; the reporting has not caught up. Try it: Fleet PUE 1.54 compared.

Every metric in the previous subsection is operational — it counts what the facility does per unit of time. A full Life Cycle Assessment (LCA), by contrast, accounts for the emissions spent once, before the hardware ever runs a token: wafer processing, HBM stacking, substrate and interposer fabrication, PCB manufacture, packaging, and transport to site. These are collectively called embodied (or cradle-to-gate) emissions, and they matter because a GPU’s useful life at hyperscale is short — three to five years, on aggressive refresh cycles — and the embodied share of its total footprint rises steeply as grids decarbonise.

For NVIDIA’s H100, published product-carbon-footprint data and independent LCA studies converge on an embodied footprint on the order of 150–200 kgCO2e per SXM card, or roughly 1.3 tCO2e per eight-GPU HGX H100 baseboard once the substrate, cold plate, and NVLink switches are included ^[115]. Roughly 42% of that impact traces to memory (HBM3 stacks are effectively several dies each, bonded through TSVs, and every die carries its own wafer cost), 25% to the other integrated circuits, and 18% to thermal and mechanical components. NVIDIA’s own comparison of HGX B200 against HGX H100 reports a 24% reduction in embodied carbon per baseboard alongside the well-known headline efficiency gain ^[116]. Try it: H100 embodied vs operational carbon.

The dominant driver of embodied carbon in leading-edge silicon is not the mass of the chip but the energy intensity of the fab. Modern GPUs are printed at 5 nm and 3 nm nodes on EUV lithography, which uses roughly ten times the electricity per wafer of the 28 nm-era processes it displaced; the fabs themselves also consume enormous volumes of ultrapure water and use high-global-warming-potential gases (SF6, C4F8, NF3) for etching and chamber-cleaning. HBM adds a further layer: each stack is a 3D assembly of six to twelve DRAM dies plus a logic base die, and yield losses in the stacking step multiply through the wafer count. Transport (largely air-freight from Taiwan/Korea fabs to US and European hyperscale sites) is a smaller but non-trivial contributor.

For a training run, the operational share still dominates: Patterson et al.'s "Carbon Emissions and Large Neural Network Training" estimated GPT-3 at 552 tCO2e for its single training run and showed that the choice of model, hardware, datacenter, and geographic region can compound to a 100–1000× swing in emissions for the same task ^[117]. For a single H100 amortised over a five-year life at typical utilisation, however, embodied and operational contributions land within a factor of two of each other only on a very clean grid (below roughly 15 gCO₂/kWh at 60% utilisation; at 30 g/kWh operation is already about 4× embodied) — which is why Green Software Foundation’s Software Carbon Intensity (SCI) specification, adopted as ISO/IEC 21031:2024, requires practitioners to report embodied and operational components separately rather than as a single blended number ^[118]. Try it: Embodied vs operational, clean grid.

There is a Jevons-paradox reading of the last decade of AI hardware history that the industry has been reluctant to acknowledge. Each generation of accelerator has delivered a step-change reduction in joules per token — H100 to B200 alone claims a roughly 25× improvement in inference energy per token for frontier models — and yet total datacenter electricity consumption is projected by the IEA to double between 2024 and 2030 (Section 58). Recent empirical work labels this the "AI Jevons Paradox": AI adoption reduces the carbon intensity of a given task while simultaneously producing statistically significant increases in absolute corporate emissions, because the reduced per-unit cost unlocks orders of magnitude more usage ^[119]. The DeepSeek-R1 moment of January 2026 — an order-of-magnitude cheaper reasoning model that expanded the addressable use-cases for reasoning rather than shrinking the compute demand — is the case study most often cited.

Water is the metric AI made unignorable. Li et al. estimated that training GPT-3 in Microsoft’s US datacenters directly evaporated about 700,000 litres of freshwater on-site, and that global AI water withdrawal will reach 4.2–6.6 billion cubic metres by 2027 — comparable to the annual withdrawal of Denmark ^[120]. Meta’s Richland Parish (Louisiana) campus alone is permitted for 23 million gallons per day, with Meta projecting actual use closer to 500–600 million gallons per year once seasonal cooling is netted out ^[121]. The engineering divide is between cooling towers (open evaporative loops — cheap, low PUE, high WUE, water simply gone) and closed-loop chilled-water plants (higher PUE, near-zero site WUE, but more expensive to build and to run in hot climates); the last two years have visibly bent hyperscaler designs — Microsoft and Google both — toward closed-loop and adiabatic alternatives in water-stressed regions.

The reporting story on carbon is thornier, and it centres on the difference between annual matching and hourly matching of clean power. Under the older annual-matching regime, a company can buy enough Renewable Energy Certificates (RECs) or Power Purchase Agreements (PPAs) to match its yearly consumption in aggregate and legitimately claim "100% renewable" — while running at night, in fossil-heavy regions, on whatever mix the grid actually delivers. Google’s 2018 white paper "The Internet is 24x7. Carbon-free energy should be too." argued this was the honesty gap the industry needed to close, and its 2020 24/7 Carbon-Free Energy (CFE) commitment redefined the target as matching each hour of consumption at each site with contemporaneously generated carbon-free electricity — measured through Time-based Energy Attribute Certificates (T-EACs) rather than fungible annual RECs ^[122]. Microsoft’s 2024 disclosure that its Scope 3 emissions had risen ~30% year-on-year was partly a consequence of exactly this shift: after the company stopped purchasing unbundled RECs, its reported footprint jumped without any actual change in operations, simply because the accounting had been quietly inflating the "green" share ^[123].

The geographic dimension follows straight from Patterson’s 4Ms framework: because grid carbon intensity varies 5–10× across regions even within a single country, where a workload runs can dominate every other lever. Nordic sites — Iceland’s geothermal-hydro grid, Sweden and Norway’s hydro-heavy grids — report near-zero CUE and PUE below 1.1 thanks to year-round free cooling, and Verne, atNorth, Nscale, and Google’s Hamina campus have all located AI-capable capacity there specifically to capture that arithmetic ^[124]. US Southern siting — Texas, Louisiana, Georgia — sits at the opposite end: cheap land and fast interconnect, but a grid mix dominated by natural gas and, in Louisiana’s case, considerable coal. Meta’s siting of a $10B AI campus in Richland Parish alongside its own renewable PPAs is the sharp end of that trade-off: the physical electrons feeding the racks come from one of the dirtiest grids in the US, the offset is contractual, and whether that counts as "green" depends entirely on which of the frameworks in this section a reader accepts.

Google, Microsoft, Meta, and AWS all publish annual sustainability reports (Google’s Environmental Report, Microsoft’s Environmental Sustainability Report, Meta’s Sustainability Report, Amazon’s Sustainability Report). The 2024–2025 editions share a common shape: PUE has plateaued near 1.1, absolute emissions are up double-digit percentages driven by AI buildout, water use is disclosed but under conservative accounting, and the 2030 pledges — Google net-zero, Microsoft carbon-negative, Meta and Amazon net-zero across value chain — are described in language that has migrated from "on track" to "will require breakthroughs." The metrics in this section are what will eventually settle whether those pledges were met, missed, or quietly redefined; the current reporting, judged against the definitions rather than the marketing, already tells a clear story.

The macro numbers make the shape of the problem visible. Lawrence Berkeley National Laboratory’s 2024 United States Data Center Energy Usage Report found that US datacenter electricity climbed from 58 TWh in 2014 to 176 TWh in 2023 — 4.4% of national consumption — and projected 325–580 TWh by 2028, a range whose lower bound alone is a doubling in five years ^[125]. The International Energy Agency’s Electricity 2026 update globalises the picture: worldwide datacenter demand from 485 TWh in 2025 to roughly 950 TWh in 2030 in its base case, climbing to about 1,200 TWh by 2035, with accelerated-server (i.e. AI) consumption growing 30% a year against 9% for conventional servers ^{[126] [127]}. Datacenters are now on track to be responsible for roughly half of all electricity-demand growth in the United States through the end of the decade.

Zoom in from those aggregates to specific sites and the shape of the buildout becomes clearer. Meta’s Hyperion campus in Richland Parish, Louisiana is planned at 2 GW of IT load, growing toward 5 GW at build-out; Meta has committed to finance more than 5.2 GW of new natural-gas generation and roughly 240 miles of new 500 kV transmission to support it ^{[128] [129]}. xAI’s Colossus in Memphis went from zero to a 100,000-GPU cluster in 122 days, but its grid connection was capped at 150 MW by the Tennessee Valley Authority in November 2024 and only doubled to 300 MW in February 2026 — so xAI installed roughly 35 on-site Solar SMT-130 gas turbines totalling around 400 MW to run the site while the utility caught up, an approach the local NAACP chapter is now suing over ^{[130] [131]}. SemiAnalysis counts five separate 1 GW-plus campuses coming online in 2026, each operated by a different hyperscaler, and 20 GW of new datacenter capacity energized that year with 30 GW more scheduled in 2027 — and concludes that model capability is still outrunning it ^[132].

Set that pace against how long the physical apparatus takes to build. A GPU cluster of tens of thousands of accelerators can be lit in 4–18 months; xAI’s 122 days is the outlier, but 12 months for a hyperscale fit-out is routine. A leading-edge fab takes 2–3 years from groundbreaking. A combined-cycle gas plant takes 30–48 months. A restart of a shut-down nuclear reactor takes about four years. A greenfield large light-water reactor takes 10–15. And a high-voltage transmission line — the thing that actually moves the electrons to the campus — averages ten years or more from route identification to energisation, and often stretches past 15 ^{[133] [134]}. The gap between how fast AI compute can be built and how fast the wires and generation to run it can be built is now measured in years, and no amount of GPU procurement changes that arithmetic.

Every large load in North America connects to the grid through an Independent System Operator (ISO) or Regional Transmission Organisation (RTO) — PJM in the mid-Atlantic, MISO in the Midwest, ERCOT in Texas, CAISO in California, SPP in the Great Plains, ISO-NE and NYISO in the Northeast. The generator- and large-load-interconnection process is standardised, and its stages are what determine the years. An application enters the queue; the ISO runs a feasibility study (weeks to months) to see whether it can be connected at all; a system impact study (typically 12–18 months) computes what additional transmission, reactive support, and stability work the interconnection triggers; a facilities study (12–24 months) designs those upgrades; then the interconnecting party pays for them and waits for them to be built (2–5 years, sometimes decades if a new high-voltage line is required); finally, the executed Interconnection Service Agreement (ISA) allows energisation ^{[135] [136]}.

The queue itself is a scale of its own. Berkeley Lab’s tracking shows more than 2 TW of proposed generation and storage nationally sitting in interconnection queues at year-end 2025; ERCOT alone has more than 400 GW in queue with an average wait of 4.5 years; PJM has been suspending 46–79% of projects; MISO averages a 20% suspension rate; and Enverus estimates that only about 10% of queued capacity will actually come online in the next three years ^{[137] [138]}. FERC issued Order 2023 in 2023 to force ISOs off first-come-first-served serial studies and onto first-ready cluster-study reforms, and CAISO, ERCOT, SPP, and MISO all signed record capacity of ISAs in 2025 as those reforms took effect — but the backlog is deep enough that even record throughput is not clearing it fast enough to keep pace with datacenter demand growth ^[139].

The most consequential regulatory development of 2025 came at the end of the year, when FERC issued a December 18, 2025 order finding PJM’s tariff "unjust and unreasonable" for co-located large loads and directing PJM to create three new categories of transmission service — Firm Contract Demand, Non-Firm Contract Demand, and standard network service — for datacenters co-locating with generation. The order also revised the rules for Behind-the-Meter Generation (BTMG), clarifying that netting a large co-located load against BTMG is not just and reasonable at scale because it can cost-shift transmission costs onto other ratepayers, and effectively foreclosing that structure above a small threshold ^{[140] [141]}. That ruling — the direct legacy of the November 2024 rejection of the original Amazon-Talen behind-the-meter agreement at Susquehanna — is what pushed the industry off pure behind-the-meter arrangements and toward the grid-connected, front-of-the-meter PPA structures now dominant.

Given those timelines, hyperscalers have effectively split their power strategy in two: a short-term energisation strategy that gets electrons flowing in 2025–2028, dominated by gas and grid, and a long-term decarbonisation and firming strategy that arrives from 2028 onward, dominated by nuclear. The menu itself is illustrated in Figure 47a.

The bindin cooling constraint at datacenter scale is water — for evaporative cooling towers, for adiabatic pre-cooling of air, and for humidification of dry hall air. A hyperscale AI campus at 500 MW IT load with a PUE of 1.15 and a WUE of 1.4 L/kWh (litres per IT kWh) evaporates roughly 6 GL of water per year (~1.6 billion US gallons). At the top end that is the domestic consumption of a small city. Try it: Water for 500 MW.

Meta’s Richland Parish permit illustrates the scale and the discomfort. The site is registered to draw up to 23 million gallons/day (about 8.4 billion gallons/year) from the Mississippi River Alluvial Aquifer under peak conditions, though Meta projects average use of 1.5 million gallons/day (500–600 million gallons/year) once seasonal cooling is netted out ^{[159] [160]}. Google reported over 23 billion litres of freshwater consumed fleet-wide in 2023; Microsoft’s fleet WUE varied from ~0.01 L/kWh in Denmark to ~1.63 L/kWh in Arizona, purely as a function of climate ^[161].

The competitive dynamic in dry regions is with municipal supply, which is what pushed Microsoft in December 2024 to announce zero-water-evaporation designs for all new datacenters (Part XV), applied since August 2024 and saving 125 million litres per site per year — at the cost of a nominal PUE rise. Meta’s Louisiana filings similarly emphasise closed-loop cooling on the highest-density AI halls. The direction of travel is clear: water is becoming the binding cooling constraint in the exact regions where cheap land and fast interconnect made siting economical, and the industry is re-optimising away from evaporation deliberately.

The unique wrinkle AI has added to power engineering is that training clusters do not draw a smooth load. A synchronous training step consists of a compute-heavy phase in which every GPU on the cluster runs at full utilisation, followed by a communication-heavy phase in which they exchange gradients across NVLink and InfiniBand at much lower compute (Part XVII). The result is a square-wave load pattern: for tens of thousands of GPUs running in lockstep, the aggregate power draw can swing from ~30% idle to 100% full and back in tens of milliseconds. At gigawatt scale that is hundreds of megawatts of step change on a sub-second timescale ^{[162] [163] [164]}.

Utilities were not built for this. Conventional industrial loads change on minutes-to-hours timescales; a training cluster changes on the timescale of grid-synchronous machinery’s own mechanical response. If the load’s oscillation frequency happens to line up with a resonant mode of a nearby turbine’s shaft, or with a transmission line’s dominant electromechanical mode, the mechanical coupling is real and damaging — subsynchronous resonance is a known category of turbine failure, and it can be triggered by an oscillating electrical load as easily as by an oscillating source. Voltage and reactive power sag under the compute phase and swing back under the communication phase, producing flicker and harmonic distortion that propagate to other loads on the same feeder.

The engineering response is layered. On the datacenter side, operators install multi-megawatt lithium-ion battery banks — xAI’s Colossus put in 168 Tesla Megapacks specifically as a power-quality device (Part XV) — sized to absorb the compute-phase surge and inject during the communication-phase dip, presenting the utility with a much smoother net draw ^[165]. NVIDIA’s forthcoming 800 VDC datacenter architecture (Rubin generation and beyond) collapses the AC-DC conversion chain and adds high-frequency buffering directly at the rack, further decoupling the utility from the GPU-timescale load. Some sites also deliberately dither the training schedule at the microsecond scale to defeat perfect synchronisation ^{[166] [167]}. On the utility side, the answer is dynamic reactive support — Static Synchronous Compensators (STATCOMs) at the substation, sized in the tens to hundreds of MVAr, absorbing and injecting reactive power on millisecond timescales to hold voltage and suppress harmonics. Southern States and ABB now market STATCOMs specifically as datacenter grid-stability products, and utilities in the PJM footprint have begun requiring them as a condition of interconnection for large training loads ^{[168] [169]}. The economic upshot is that a 500 MW training cluster is not merely a 500 MW load from the utility’s point of view; it is a 500 MW load that demands 50–100 MVAr of fast reactive support and megawatt-scale short-duration storage before the interconnection is granted. Try it: 48 V vs 800 V DC.

The most extreme response to the interconnection queue is not to be on the interconnected grid at all. A true islanded datacenter runs entirely on its own generation with no grid tie: the electrons the fuel makes never touch anyone else’s wire. This is legally simpler than grid-interconnected-with-behind-the-meter — no ISA, no FERC filings, no transmission-cost allocation — because it looks to regulators like a private industrial facility with a private power plant, of which there are tens of thousands in the United States. It is operationally harder, because the site loses the grid as a stabiliser and must self-provide all reactive, frequency, and reserve services.

The middle ground — grid-connected but behind-the-meter, where the datacenter is powered primarily from an on-site plant with a small grid tie for backup and export — was the structure the original Amazon-Talen deal used, and the structure that FERC’s November 2024 rejection and December 2025 order have progressively squeezed. FERC’s core concern is what it calls cost shifting: a large behind-the-meter load pays only for the small quantity of grid power it draws, while relying on the grid’s transmission and reliability services as free options — a service other ratepayers ultimately fund. The December 2025 order institutionalises Firm and Non-Firm Contract Demand transmission services precisely to price those services fairly, and effectively caps pure BTMG netting at scales well below hyperscale training ^[170].

The state layer is separately active. Texas, courting datacenter capital under ERCOT’s less-federalised jurisdiction, has broadly welcomed behind-the-meter arrangements. Louisiana passed 2024 legislation smoothing the Meta/Entergy PPA cost-allocation. Virginia — home to Dominion Energy’s data-center-heavy service territory — has begun requiring large loads to contribute to transmission planning studies. The regulatory picture in 2026 is that pure BTMG at gigawatt scale is dead on the interconnected grid; grid-connected FTM PPAs, ratepayer-cost-neutralised co-location, and true off-grid islanded designs are the three surviving structures.

Epoch AI’s total-cost-of-ownership analysis of a 1 GW AI datacenter is the clearest recent public model. Server hardware dominates: for a facility holding roughly $30 B of accelerators depreciated over five years, compute-side cost runs about $6 B/year, versus roughly $0.6 B/year in electricity at typical grid rates and $0.5 B/year in facility overheads. Under that arithmetic, IT depreciation is roughly 10× energy OPEX — silicon still dominates ^[171]. Try it: Epoch 1 GW datacenter TCO.

But the arithmetic bends as the IT lifespan changes and as power cost shifts by region. Shortening GPU life from 5 to 3 years — plausible given the pace of Blackwell → Rubin → Feynman generational transitions — pushes annual TCO from ~$7 B to ~$11 B and shifts the OPEX share downward. Conversely, extending life to 7 years pulls TCO to ~$5.4 B and pushes energy to about 11% of annual cost. And at siting decision time, industrial power tariffs range from $30/MWh in ERCOT wind-heavy hours to $120/MWh in constrained coastal markets; a 1 GW site drawing 8 TWh/year pays $240 M at the low end and $960 M at the high end. That $720 M/year swing is roughly 2× the amortised cost of the site’s transmission connection over its 20-year facility life. It is why the industry is chasing cheap firm power even when the physical land or fibre routing is otherwise inferior — and why Louisiana, Wyoming, rural Pennsylvania, and Icelandic geothermal have become viable siting choices they were not in 2020. Try it: TCO on a 3-year life; TCO on a 7-year life; 8 TWh energy bill range.

The rule of thumb operators now use is that above roughly $80/MWh delivered firm, power OPEX becomes the primary driver of cluster location choice; below it, latency, fibre, and labour dominate. Every hyperscale AI campus announced since mid-2024 sits at a location where the operator has secured firm power at or below that threshold, usually via a long PPA or by committing to build generation itself. Try it: The $80/MWh siting rule.

The lever hyperscalers have that other large industrial loads do not is that training workloads are latency-insensitive at the hours-to-days scale. A training run does not care whether it finishes at 3 AM or 3 PM; it does care about wall-clock time overall, but it can be paused, migrated, or throttled with no user-visible effect. Google has run this lever in production since 2020 through its carbon-intelligent computing platform, which shifts flexible compute (training, batch analytics, media transcoding) toward the hours and datacenter regions where carbon-free energy is available, and dispatches downward in response to utility demand-response calls ^{[172] [173]}. During Winter Storm Fern in early 2024 Google reduced power at its Pryor, Oklahoma facility at the request of the Grand River Dam Authority — the first documented case of a hyperscale datacenter operating as a dispatchable demand-response asset.

The flexibility asymmetry, though, is stark and worth stating precisely. Training can be time-shifted at the hours-to-days scale, and training plus inference batch can be spatially shifted across regions with modest extra egress cost. Inference for interactive users cannot be shifted at all — the request is here now, the answer must be here in 250 ms. As the compute mix bends decisively toward inference through 2027 (Part XIII), the share of aggregate AI load that is flexibility-monetisable is shrinking, not growing. This is one of the tensions in the "AI as a grid asset" argument: the very workload shift that makes reasoning models economically essential also removes the flexibility that made carbon-intelligent computing possible.

The map of where AI capacity is being built follows the map of where firm power is available cheaply, and the pattern is now clear enough to name.

United States Southeast and Texas. Louisiana, Mississippi, Alabama, and Georgia offer cheap land, cheap gas, favourable state tax treatment, and transmission systems (MISO South, SERC, ERCOT) that — while queued — are faster to build across than the Northeast. Meta Louisiana, xAI Memphis, and Anthropic’s rumoured campus in the region all fall in this pattern. The tradeoff is grid mix (gas- and coal-heavy) and water stress in the drier sub-regions.

US Pacific Northwest and Mountain West. Washington, Idaho, and Wyoming draw on hydro (BPA / Grand Coulee), stranded gas, and — for TerraPower and X-energy — new nuclear siting on retiring coal plants under DOE Advanced Reactor Demonstration incentives. Microsoft Wisconsin fits a similar pattern in MISO North.

Nordics. Sweden, Norway, Iceland, and Finland offer 90%-plus renewable grids (hydro, geothermal, wind), 40–60% lower power cost than continental Europe, and free ambient cooling year-round. Microsoft, CoreWeave, and Brookfield committed more than $15 B to Nordic AI infrastructure in 2025; OpenAI’s Stargate Norway plans 100,000 GPUs entirely on hydroelectric power ^{[174] [175]}.

Middle East. Abu Dhabi (via G42) and Riyadh (via HUMAIN) have moved from oil-export capital to compute-export capital as an explicit sovereign strategy. Stargate UAE is a planned 5 GW campus in Abu Dhabi with a first 200 MW cluster targeted for 2026, powered by a mix of natural gas, solar, and eventual nuclear; HUMAIN has committed roughly $100 B for 11 datacenters totalling 2.2 GW ^{[176] [177]}. The geopolitical bargain is that both governments trade cheap firm power and capital for US-controlled AI infrastructure and export-controlled GPUs — a bargain the US Bureau of Industry and Security explicitly conditioned on end-use monitoring in 2024–2025.

Ireland — the cautionary story. Between 2015 and 2022 Ireland captured a disproportionate share of European hyperscale capacity via Dublin’s cluster of Amazon, Microsoft, Google, and Meta sites. By 2023 those facilities consumed over 20% of national electricity — more than all urban households combined — and EirGrid declared a de facto moratorium on new Dublin-area datacenter connections that lasted through 2025. The Commission for Regulation of Utilities' December 2025 ruling formally ended the freeze but replaced it with strict conditions: any datacenter above 10 MVA now must provide its own on-site dispatchable generation or storage. The moratorium is the clearest political precedent for what happens when a small grid lets AI-scale load run away from it, and every other jurisdiction is watching ^{[178] [179]}.

China. The domestic story is the East Data, West Compute (东数西算) policy, formally approved by the NDRC in 2022, designating eight national computing hubs — with the western clusters (Guizhou, Inner Mongolia, Gansu, Ningxia) explicitly sited to consume otherwise-stranded hydro, wind, and solar generation. Guizhou’s hydro-powered Gui’an New Area anchors the buildout; Huawei, Tencent, and Alibaba have all announced large western-hub AI capacity. Domestic accelerator supply (Huawei Ascend, Cambricon) closes the second half of the sovereign compute loop ^[180].

Pulling the pipelines together produces a fairly stark picture. The credible 2026–2030 US datacenter power pipeline — everything signed, financed, and in construction — adds up to roughly 60–90 GW of new AI-directed capacity across the hyperscalers, with SemiAnalysis’s tracking putting 20 GW energised in 2026 and 30 GW targeted for 2027 ^[181]. The credible generation-addition pipeline against that need — new gas, renewables, restarts, and (past 2030) SMRs — is thinner and more back-loaded. Roughly 15–20 GW of new gas can plausibly enter service in 2025–2028 nationally; the Three Mile Island restart adds 835 MW in 2027–2028; Palisades, if it restarts, adds another 800 MW around the same time; SMR contributions before 2030 are essentially zero.

The mathematics that follows is uncomfortable. In the 2026–2029 window, if model demand grows at the pace analysts are projecting, either (a) hyperscalers accept gas at scale — pushing the sector’s carbon curve further past the pledges Part XIII named — or (b) the buildout slows to the pace at which nuclear restarts and combined-cycle plants can be commissioned. There is no third option that materially changes the generation math on this timescale. From 2030 onward the SMR and expanded nuclear pipeline could plausibly rebalance the mix, but only if the first-of-a-kind builds actually deliver on their announced dates — a proposition the industry’s honest observers rate as 50/50 at best.

The through-line back to the rest of this primer is that every optimisation Parts I–XX described — memory bandwidth, KV-cache paging, speculative decoding, sparsity, quantisation, cooling, PUE — is now competing for value against a constraint that lives entirely outside the datacenter: how many gigawatts of firm power a specific utility, in a specific service territory, can add to a specific substation on a specific timeline. The industry has spent five years learning to optimise for FLOPS/W. It is about to spend the next five learning to optimise for FLOPS per available gigawatt-year, which is a fundamentally different unit — one that integrates land, permitting, transmission, water, and community consent as much as it does silicon.

That is what makes energy the true Layer 1 of the AI stack. Every layer above it — the transformers, the schedulers, the CDUs, the racks — assumes electrons show up on time. In 2026 that assumption is no longer safe, and the industry’s ability to close the gap between announced clusters and flowing electrons will determine what actually gets built.

A D2C loop has three physical stages, and they exist in every mainstream AI rack from a GB200 NVL72 to a Lenovo Neptune to a Supermicro SuperBlade.

The cold plate sits directly on the GPU or NVSwitch die, bolted down over a thin thermal interface material in place of what would have been an air-cooled heatsink. Modern plates for a Blackwell-class GPU are copper (occasionally nickel-plated to arrest galvanic couples — see below) with a thin machined lid; the coolant path inside is either microchannel (an array of parallel fins ~200 µm wide the fluid runs through, the standard approach used by Cooler Master, Delta, Boyd/Aavid, and most tier-one integrators) or microjet impingement (an array of tiny orifices that squirt fluid perpendicularly at the hot silicon, JetCool’s differentiator) ^[183]. Microjet plates report 20–40% lower thermal resistance than microchannel at the same flow, and up to a 3× reduction against older designs, which is the difference that lets a single-die 3,000 W part remain a candidate for D2C at all ^[184]. Boyd, one of the highest-volume vendors, reported shipping its five millionth liquid cold plate to hyperscalers in September 2025 — a useful scale check on how quickly D2C moved from HPC-exotic to commodity ^[185].

The rack manifold is the vertical pair of pipes running the height of the rack — one supply, one return — that fans out coolant to every compute tray. Each tray taps the manifold through a pair of quick disconnects (QDs, or "dripless connectors"): spring-loaded couplings that shut off both halves of the loop the instant the tray is unplugged, so a compute tray can be pulled for service without draining anything or wetting the floor. The Open Compute Project has standardized this interface as the Universal Quick Disconnect (UQD), adopted by every mainstream OEM so that plates, manifolds, and hoses from different vendors interoperate ^[186]. The OCP’s parallel Cold Plate Development and Qualification whitepaper sets thermal-resistance, pressure-drop, and long-term-reliability targets that let a hyperscaler qualify plates from multiple suppliers to the same bar ^[187].

The Coolant Distribution Unit (CDU) closes the loop. It is a liquid-to-liquid heat exchanger, plus pumps, filters, sensors, and control electronics, sitting between the technology (rack-side) loop and the facility (building-side) loop that ultimately dumps heat to a dry cooler or cooling tower on the roof. It performs three jobs: it isolates the clean, precisely-regulated rack coolant from the rough facility water; it pumps and filters the rack loop; and it holds the rack-side supply temperature, flow rate, and pressure at whatever the silicon vendor’s datasheet requires. Two placement patterns are common. An in-row CDU is a rack-sized cabinet placed next to (or between) compute racks and serves a group of them — Vertiv’s Liebert XDU 1350 and CoolChip CDU 1350 handle up to 1,350–1,368 kW, and CoolIT’s CHx1000 handles 1,000 kW at up to 1.5 LPM/kW, enough for on the order of ten GB200 NVL72 racks each ^{[188] [189]}. An in-rack CDU is a 4U unit inside a single rack — Motivair’s In-Rack CDU is one example — used when the facility can’t yet run a shared loop, or for pilot deployments and edge sites ^[190]. The main vendors in this segment are CoolIT Systems, Motivair (recently partnered with Schneider Electric), Vertiv, Nortek, and Boyd/Aavid. Try it: CDU flow per kW.

The mainstream coolant is not pure water. It’s PG25: 75% deionized water mixed with 25% propylene glycol, plus a corrosion-inhibitor and biocide additive package. The 25% PG concentration is the industry sweet spot because it is biostatic (below ~10% PG the algae actually feed on the glycol; at 25% they starve), and because propylene glycol is non-toxic if it leaks near people or drinking water — unlike ethylene glycol, which is acutely poisonous ^{[191] [192]}. Pure DI water actually cools slightly better — it has higher heat capacity — but it demands constant filtration and chemistry monitoring to stay clean, and it freezes on the wrong side of any outdoor pipe.

The plumbing sizing follows straight from the heat load and the coolant’s thermal properties. A Lenovo GB300 NVL72 datasheet spells out what the tradeoff looks like at the rack level: at a 25 °C facility-water supply, the rack needs about 59 L/min (15.6 GPM); at 30 °C it needs 71 L/min; at 35 °C, 89 L/min; and at 45 °C — the top of the ASHRAE W45 warm-water class — 177 L/min (46.8 GPM), roughly triple the cold-supply flow. The pressure drop across the rack rises in the same direction, from about 2.3 psi at 25 °C to 18.4 psi at 45 °C ^[193]. Those flows imply a much larger temperature rise across the rack than the 8–10 °C balance point below: by , about 130 kW at 59 L/min is a rise of roughly 30 K, while a 10 K rise needs around 200 L/min — so read the datasheet figures as minimum flows at each supply temperature, not as the design point. This is the concrete form of the warm-water tradeoff Part XV described: warmer supply water buys more chiller-free hours, but at the cost of larger pumps, thicker pipes, and higher parasitic pump power. Most Blackwell-class deployments land near 30–32 °C supply and 40 °C return — a ΔT across the rack of roughly 8–10 °C — as the balance point ^[194]. Try it: Lenovo GB300 flows and ΔP; 59 L/min: implied rise; 177 L/min: implied rise; Pipe size for 177 L/min; Coolant pressure drop, 59 L/min; Triple flow, ~8× pressure drop; Rack coolant at 10 K rise; Flow at ΔT 8–10 K.

Everything above is single-phase D2C: the coolant stays liquid from cold plate to CDU and back. The alternative is two-phase D2C, where a low-boiling dielectric fluid is allowed to boil on the die and condense in a heat exchanger, exploiting the latent heat of vaporization to move roughly the same heat with about a tenth the mass flow ^[195]. The commercial champion here is ZutaCore’s HyperCool, which uses a pumped, non-conductive refrigerant and has been adopted into Rittal AI racks and, more recently, into single-slot PCIe cold plates for NVIDIA’s RTX PRO 6000 Blackwell server GPUs ^{[196] [197]}.

The tradeoff is well-defined. Two-phase’s advantages are density (higher heat flux tolerated per plate area), a near-isothermal die (the fluid holds temperature at its boiling point as heat is added, so hotspots flatten), and — because it’s dielectric — a fluid that doesn’t short electronics if it leaks. Its disadvantages are cost (dielectric fluids are expensive and, after 3M’s 2022 PFAS phase-out, harder to source cleanly), complexity (managing a boiling front is harder than pushing liquid), and a maturity gap: single-phase D2C is what NVIDIA, AMD, Supermicro, Dell, HPE, and every major OEM ship by default in 2026, while two-phase remains a specialist play used where the density or leak tolerance is worth the surcharge ^[198]. The industry consensus is that single-phase dominates the next 5–10 years, with two-phase as the long-term watch item as die-level heat flux keeps climbing.

Once liquid is in the same rack as live silicon, three failure modes dominate the reliability conversation.

Leak detection and pressure sign. The received-wisdom fear is that a loop of pressurized water above running electronics is a bug waiting to happen. The mainstream mitigation is a layered leak-detection scheme — rope sensors under the manifolds, drip pans, and CDU-integrated pressure and flow monitoring that flags a decaying loop before it drips. Chilldyne’s negative-pressure CDU is the more radical answer: the entire loop runs at sub-atmospheric pressure, so if a fitting cracks, air gets sucked in rather than water leaking out, and the CDU sees the pressure change and can shut the loop down cleanly. The claim is a categorically leak-proof loop at the cost of somewhat more complex pump architecture and lower coolant flow rates per unit of pump power ^{[199] [200]}.

Quick-disconnect reliability. A UQD is rated for on the order of 100,000 connect/disconnect cycles with dry-break sealing, which is well in excess of any realistic service life — but only when clean. Grit or biofilm on the poppet is the field-failure mode, which is why the CDU carries 25 µm filters (50 µm optional) that catch particulate before it reaches the couplings ^[201].

Galvanic corrosion. If a cold plate is copper (or nickel-plated) and somewhere else in the loop an aluminum fitting appears, the two metals sit far apart on the galvanic series and set up an electrochemical cell: the aluminum corrodes preferentially and the loop eventually pits and leaks. The rule of thumb the industry now enforces is: keep the whole wetted path in a single family (typically copper + stainless), and where mixed metals are unavoidable, use dielectric unions to break the electrical circuit and add corrosion inhibitors to the coolant. This is one of the reasons PG25 comes pre-formulated with an inhibitor package rather than as raw glycol ^[202].

The last practical question is whether an existing air-cooled hall can be upgraded to D2C in place. In practice, mostly no. A greenfield AI hall is designed around raised-floor plumbing, structural loads for CDUs and manifolds, a piped facility-water loop from the roof plant down to every row, and a heat-rejection plant (dry coolers or cooling towers) sized to the full rack density from day one. An air-cooled hall built for 10–15 kW/rack has none of those: adding D2C means threading pipe through a slab designed for hot/cold aisles, punching holes for CDU risers, replacing 80% of the in-service servers whose air-cooled heatsinks won’t mate to a cold-plate manifold, and — often the binding constraint — adding a rooftop heat-rejection plant that the building’s roof steel was never designed to carry ^[203]. This is one of the largest hidden reasons for the wave of greenfield AI campus builds after 2023: an old colo shell can be re-tenanted for inference on 40 kW air-assisted racks, but the 120 kW training rack of Part X essentially requires new steel.

The single most important serving technique is continuous batching (also called in-flight or iteration-level batching). Naive static batching gathers a fixed set of requests, runs them together until the last one finishes, and only then admits new work — which means the GPU sits partly idle as short requests complete and wait for long ones. Continuous batching instead schedules at the granularity of a single decode step: every iteration, finished sequences are evicted and freshly-arrived requests are admitted, so the batch is continuously topped up and the accelerator stays saturated through the memory-bound decode phase. This, combined with PagedAttention’s memory efficiency, is what lets a server hold far more concurrent requests than static batching ever could ^[289][291].

Try it: Animate static vs continuous batching.

The description above is easiest to trust by watching a single tick happen. A decode step is the natural unit of scheduling because it is already the unit of work: every request in the batch, regardless of how long it has left to run, contributes exactly one new token’s worth of forward-pass compute per step (Section 37) — there is no notion of a request being "partway through" a step the way it can be partway through its overall generation.

Concretely, at the top of a tick the scheduler walks every occupied slot in the current batch. A request that isn’t finished simply advances: it contributes its newest token to the batched forward pass, and its position in the KV cache grows by one entry per layer as usual (Section 11). A request whose most recently generated token was an end-of-sequence marker (or that has hit its length limit) is evicted: it is removed from the batch entirely, its response is finalized and returned to the caller, and — this is the detail that makes the technique’s memory efficiency work — its KV-cache pages (Section 11’s PagedAttention blocks) are freed back to the shared pool immediately, in the same tick, rather than sitting reserved until some later cleanup pass. Freed slots (and freed KV-cache pages) are then handed to newly-arrived requests whose prefill phase has already completed while they waited in a queue — admission is essentially free from the decode loop’s perspective, because the expensive part (prefill) already happened separately, and slotting a ready request into an empty batch position costs nothing but bookkeeping. Try it: Watch eviction and backfill; Pages freed and reused.

The throughput consequence follows directly from what static batching does instead: a static batch runs every request in lockstep until the last one finishes, so a batch containing one 50-token request and one 2,000-token request holds that first slot idle — occupying a GPU batch position and its KV-cache memory while contributing nothing — for the 1,950 extra steps the long request still needs. Continuous batching’s per-tick eviction-and-backfill means no slot is ever idle for longer than the single tick it takes the scheduler to notice it’s free, which is why real serving engines report throughput gains of several-fold over static batching on realistic, highly variable-length request mixes. Try it: See the 1,950 idle steps.

Decode’s fundamental problem is that it is sequential — one forward pass per token, waiting on the last (Section 37). Speculative decoding breaks that serialization without changing the output distribution. A small, fast draft model (or a set of lightweight prediction heads) proposes several tokens ahead; the large target model then verifies all of them in a single parallel forward pass, accepting the longest correct prefix and discarding the rest. Because verifying tokens costs about the same as generating one, every accepted token is nearly free. This is the same trick that makes prefill fast and decode slow (Section 37) turned into a deliberate technique: batch several tokens' worth of work into one parallel pass instead of paying decode’s one-token-at-a-time tax on all of them.

Part IX’s central observation — that prefill is compute-bound and decode is memory-bandwidth-bound — has a direct architectural consequence at scale: run them on different pools of hardware. Prefill/decode disaggregation sends the prefill of each request to a GPU pool tuned for compute throughput, then hands the resulting KV cache to a separate decode pool tuned for memory bandwidth and latency, each with its own independent parallelism strategy. DeepSeek disclosed doing exactly this in its V3/R1 serving, and NVIDIA’s open-source Dynamo framework (which sits above vLLM/TensorRT-LLM/SGLang and routes work between dedicated prefill and decode workers, moving KV caches between machines over a purpose-built transfer library) has made the pattern close to an industry standard as of mid-2026. Two lighter-weight relatives — chunked prefill (interleaving pieces of a long prompt’s prefill with other requests' ongoing decode steps so one big prompt doesn’t stall everyone) and KV-cache offloading (spilling cache to CPU DRAM or NVMe to extend effective capacity) — attack the same utilization problem without physically splitting the fleet ^{[294][295][296]}.

Disaggregation has a more radical extension that is becoming its own competitive category: splitting prefill and decode not just across separate pools of the same GPU, but across entirely different classes of hardware — each chosen because it is the best possible fit for one phase. The logic follows directly from everything above. Prefill is compute-bound and loves raw FLOPs, so a dense GPU is ideal for it. Decode is memory-bandwidth-bound, and the ceiling on single-stream token rate is bandwidth ÷ model size (the ~24-tokens/second figure from Part IV) — so the way to serve tokens fast is not more FLOPs but more bandwidth to the weights, which is exactly what a huge on-chip SRAM fabric provides. Groq and Cerebras both keep the entire model in on-wafer/on-chip SRAM at tens of petabytes per second — three to four orders of magnitude beyond HBM’s few TB/s — which is why they, not GPUs, hold the records for fastest single-stream generation (Part VI), and why "ultra-low-latency inference" has emerged as a distinct product category rather than a footnote to throughput serving ^[440].

The 2025–2026 pattern is to pair these strengths in one heterogeneous pipeline. NVIDIA’s own move (Part VI) is the clearest sign it is a real architecture and not a niche: its Groq-derived LP30 is positioned as a dedicated decode co-processor beside Vera Rubin GPUs that do prefill. Independently, in July 2026 AMD and Cerebras announced a disaggregated inference solution along the identical seam — AMD Instinct/Helios GPUs handle high-throughput prompt prefill and long context, while a Cerebras Wafer-Scale Engine handles ultra-low-latency decode, operating as "a single disaggregated inference workflow" and modeled (on a 1-trillion-parameter Kimi model) at up to ~5× higher tokens per second per watt versus a Cerebras-only configuration, targeted for Cerebras Cloud in late 2026 ^[441]. Both deals encode the same insight this primer has built up piece by piece: the two phases of inference want opposite hardware, and the fastest, most efficient systems are increasingly the ones willing to use two different accelerators in one request rather than forcing a single chip to be good at both. It also marks Cerebras’s own repositioning — from the wafer-scale training engine of Part VI toward being one of the fastest hosted inference providers, which is where its SRAM-everywhere bet pays off most directly.

An agent harness’s call to an LLM endpoint is not one hop to one GPU. It crosses a frontend tier — stateless, horizontally scaled, CPU-class infrastructure that exists to route, admit, and account for the request — before it ever touches an accelerator, and the accelerator’s own work is itself split across the serving-stack machinery this Part has already built up.

Concretely, in order:

The pattern worth naming is that the frontend tier — steps 1, 2, 3, and 6 — is stateless and cheap to scale: it is ordinary CPU infrastructure, and an operator can add load balancers and gateway instances almost for free relative to the GPUs behind them. Every dollar and every hard constraint lives in steps 4 and 5, on the accelerator. This is the load-bearing fact behind the rest of this section: a "pod" is sized by its GPU count, not its frontend footprint, and scaling the frontend is never the bottleneck — scaling the backend is.

Putting rough numbers to the six stages, for the worked 70B-class model this primer has used throughout (Section 16’s ~140 GB FP16 footprint, H100’s 3.35 TB/s HBM, Section 11’s 0.33 MB/token KV cache): Try it: Bytes per token: MB vs MiB.

Try it: Per-stage latency waterfall.

Two things fall out of this table directly. First, TTFT and TPOT are set by entirely different bottlenecks (Part XIII): TTFT by compute and by whatever the prefix cache didn’t already cover, TPOT by memory bandwidth divided by however much the batch has amortized it. Second, for anything beyond a short answer, decode dominates wall-clock time by roughly an order of magnitude over every frontend and prefill cost combined — which is exactly why every technique in this Part (continuous batching, speculative decoding, disaggregation) targets decode specifically, and why a reasoning agent that emits thousands of hidden tokens before answering (Part XIII) can turn a 5-second turn into a 100-second one without a single frontend millisecond changing. Try it: Add 10,000 reasoning tokens.

"Pod" is used loosely across the industry — anywhere from one 8-GPU node (Section 12) to a full rack (Section 39, a GB200 NVL72’s 72 GPUs) to NVIDIA’s Scalable Unit (Section 39, 8 racks / 576 GPUs). The arithmetic is identical at every scale; only the GPU count changes. Take the rack as the concrete unit, since it is the one this primer has already sized in kilowatts.

Start from the single-GPU numbers this primer has already established. One H100’s worth of bandwidth caps the 70B model at roughly 24 tokens/second (in practice the 140 GB of FP16 weights need two or more 80 GB GPUs) (Section 16’s bandwidth-over-model-size ceiling) — but that figure is for one unbatched stream, and no production system runs that way. Continuous batching (Figure 50) reuses each expensive HBM weight-read across every concurrent request, so aggregate per-GPU throughput at a realistic serving batch (tens to low hundreds of concurrent sequences, depending on sequence length and KV-cache headroom) commonly reaches several thousand tokens/second/GPU on current-generation hardware — the same order of magnitude as the $0.42/million-token illustration in Part XX, which assumed 2,000 output tokens/second per accelerator. Try it: From 24 to thousands tok/s.

Scaling that to a rack:

Try it: Pod requests per second.

A GB200 NVL72 rack (72 Blackwell GPUs, Section 39) sustaining, say, 3,000 output tokens/second/GPU in aggregate gives the rack roughly 216,000 tokens/second in total decode throughput. Against a short chat answer (~300 output tokens) that is on the order of 700 requests/second the rack can sustain in steady-state decode; against a 10,000-token agentic/reasoning turn (Part XIII) the same rack sustains only about 20 requests/second — a 33× drop for the identical hardware, entirely because the tokens-per-answer denominator grew. This is the pod-level restatement of Part XIII’s "cost per answer, not cost per token" point: a pod’s request-per-second capacity is not a fixed hardware spec, it is hardware throughput divided by however much a given workload asks of it, and reasoning-style agentic workloads can cut that capacity by more than an order of magnitude versus short chat. Try it: 720 vs 21.6 req/s.

Two ceilings sit underneath that formula, and either can bind first depending on the workload: Try it: Where compute starts to bind.

Given that model, there are three distinct ways to grow capacity, and they are not interchangeable:

The practical rule these three paths imply: scale out is the default lever, cheap and near-linear, and covers the common case of "more users, same workload mix." Scale up is reached for only when a single request’s own size — model weights or context length — genuinely does not fit in one NVLink domain. Disaggregated specialization is reached for once a workload’s phase mix is skewed enough (heavy agentic tool-use and reasoning traffic, versus short stateless chat) that a uniform pod is provably spending GPU-hours on the wrong bottleneck for a meaningful fraction of its traffic — which, per Part XIII’s own numbers, is an increasing fraction of all traffic as reasoning-style agents become the dominant workload rather than the exception.

Knowledge distillation trains a compact student model to reproduce the behavior of a larger, more capable teacher — and it changes the model’s architecture (the student is a genuinely smaller network), which is exactly what distinguishes it from quantization (same network, fewer bits). The original insight, from Hinton, Vinyals, and Dean, is that a teacher’s full softmax distribution carries far more information than a hard label: the relative probabilities among the wrong answers ("this 7 looks a bit like a 1") — what they called dark knowledge — teach the student the structure the teacher learned. A temperature term softens the distribution to expose those small probabilities, and the student is trained to match the teacher there ^[400].

Distillation comes in several forms. Response-based (or logit) distillation, the original recipe, has the student match the teacher’s output distribution. Feature-based distillation goes deeper, training the student to match the teacher’s intermediate activations as well — the FitNets "hints" idea — which lets a thin, deep student mimic a wide teacher ^[401]. Data (or sequence-level) distillation is the form dominating LLMs today: the teacher simply generates a training corpus, and the student is supervise-fine-tuned on it. DeepSeek-R1 is the canonical example — its full reasoning model produced roughly 800,000 reasoning traces that were used to fine-tune much smaller Qwen and Llama students into capable reasoning models by imitation alone, with no reinforcement-learning stage on the students at all ^[402].

A variant matters specifically for the edge and for autonomous systems, and answers a question that recurs around "lite" on-device models: the privileged teacher. Here the teacher is trained with information available only at training time — extra sensors, future frames, a high-definition map, or ground-truth world state — and the student is distilled to reproduce the teacher’s decisions from only the inputs it will actually have at deployment. The formal basis is "learning using privileged information," unified with distillation by Lopez-Paz and colleagues; the cleanest autonomous-driving instance is Learning by Cheating, where a privileged teacher that sees the true bird’s-eye layout teaches a camera-only student to drive ^[403][404]. This is precisely the shape of a compact on-vehicle driving model: a big, resource-unconstrained teacher trained offboard, distilled into a small network that runs in real time inside the car — and it is distillation, not merely quantization, that produces the smaller network.

The consequences are worth stating plainly. A student is, as a rule of thumb, capped by its teacher — it inherits the teacher’s biases and systematic errors, and tends to generalize worse off the distribution the teacher’s data covered (though narrow weak-to-strong exceptions exist, so this is a strong tendency rather than a law). Distillation reduces inference cost by producing a smaller deployable model, but it adds a training cost: you must have the teacher, and you must train the student on teacher-scale supervision ^[400][402].

Where distillation changes the architecture, quantization keeps the architecture and lowers the bit-width of the weights (and sometimes activations) — the numeric formats of Part II, applied to a trained model. There are two routes. Post-training quantization (PTQ) quantizes an already-trained model with only a little calibration data and no gradient training: fast and cheap, but accuracy falls as bit-width drops. Quantization-aware training (QAT) simulates quantization during a training or fine-tuning run so the weights adapt to the rounding noise, recovering accuracy at the cost of that run — increasingly the practical path to usable FP4-class formats, though rotation-based PTQ methods (below) can now reach 4-bit without full QAT, so "FP4 needs QAT" is a tendency, not an absolute ^[405].

Two design axes shape every scheme. What is quantized: weight-only (e.g. 4-bit weights, 16-bit activations) attacks the memory-bandwidth bottleneck of decode and is the common LLM choice; weight-and-activation (8-bit or 4-bit both) enables integer matmul but is much harder. And granularity: a single scale per tensor is cheapest and least accurate, a scale per channel is better, and a scale per small block or group of weights is the most accurate — the same block-scaling idea that underlies the MX and NVFP4 formats in Part II. The recurring enemy is the outlier activation: a few activation channels have enormous magnitudes that stretch the quantization range and starve everything else of precision. The canonical algorithms each attack the low-bit problem differently — GPTQ uses approximate second-order information to quantize weights layer-by-layer with tiny error, AWQ protects the small fraction of activation-salient weight channels, SmoothQuant migrates the difficulty from activations into weights by an equivalent rescaling to enable 8-bit-everything, and rotation methods (QuaRot, SpinQuant) apply Hadamard/learned rotations that mathematically eliminate the outliers and open up genuine 4-bit weights and activations ^{[406][407][408][409][410]}. A concrete 2025 endpoint: OpenAI’s open-weight gpt-oss models were trained quantization-aware directly in MXFP4, which is what lets the 120B-parameter model run on a single 80 GB GPU ^[297]. Quantization is inference-only — it shrinks footprint, bandwidth, and latency but does nothing for the original training cost (and QAT adds training cost).

Pruning changes the number of weights, deleting the ones that contribute least. Unstructured pruning zeroes individual weights anywhere in a matrix — the highest compression ratio, but the zeros land in arbitrary positions that ordinary dense-matmul hardware cannot skip, so the memory saving is real while the speed-up often is not. Structured pruning removes whole units (channels, attention heads, filters), yielding a genuinely smaller dense model that runs faster on any hardware, at the cost of more accuracy lost per unit removed. The conceptual backdrop is the lottery ticket hypothesis — the finding that a dense network contains a sparse subnetwork which, trained from the right initialization, can match the full model — which is why "most of the weights are removable" is a reasonable prior at all ^[411].

For LLMs specifically, two one-shot methods matter because they prune to roughly 50% sparsity without retraining: SparseGPT frames pruning as a layer-wise sparse-regression problem solved with second-order information, and Wanda prunes each weight by the simple product of its magnitude and the norm of its input activation — no weight update at all ^[412][413]. The hardware-friendly middle ground is the 2:4 semi-structured pattern from Part II (exactly two of every four contiguous weights zero), which NVIDIA’s sparse tensor cores accelerate for a real ~2× on the sparse matmul, capped at 50% sparsity ^[414]. As with quantization, pruning mainly reduces inference cost; structured pruning and low-rank factorization usually want a short fine-tune to recover accuracy, whereas the SparseGPT/Wanda contribution is precisely that they skip it.

Three further techniques round out the family. Low-rank factorization compresses the weights themselves: a large weight matrix is approximated as the product of two much smaller matrices (typically via SVD), cutting both parameters and FLOPs when the matrix is close to low-rank ^[415]. This is easy to confuse with — but is the opposite of — LoRA and the broader parameter-efficient fine-tuning (PEFT) family: LoRA does not compress the base model at all; it freezes the full model and trains only small low-rank adapter matrices, slashing the trainable-parameter and optimizer-state cost of adaptation (reported up to roughly 10,000× fewer trainable parameters than full fine-tuning). QLoRA combines this with a 4-bit frozen base so a 65–70B model fine-tunes on a single 48 GB GPU, and multi-LoRA serving systems (S-LoRA, Punica) swap thousands of adapters over one shared base at inference ^[416][417]. LoRA is thus a training-cost technique, not a model-shrinking one — it is included here because it is constantly conflated with the others.

The accelerator consequence of LoRA is worth drawing out, because it is what makes fine-tuned models economical to serve at scale. A LoRA adapter is only megabytes — the low-rank matrices — against the base model’s tens of gigabytes of weights, so a GPU can hold the (expensive, HBM-resident) base model once and hot-swap thousands of tiny per-customer or per-task adapters in and out around it, rather than dedicating a separate GPU’s worth of memory to each fine-tune. That single fact is why a provider can offer a distinct fine-tuned model to every customer without needing an accelerator per customer: the marginal hardware cost of one more fine-tune collapses from "another copy of the model" to "another few megabytes of adapter," and the base model’s weight reads are amortized across every tenant’s requests exactly the way batching (Part IV) amortizes them across concurrent users.

It is worth being precise about the mechanism, because the naïve version would break batching. There are really two ways to deploy a fine-tune: merge the adapter into the base weights ( computed once), which gives zero inference overhead but produces a dedicated, standalone model that no longer shares anything; or keep the adapter separate and apply it at runtime, which is what multi-tenant serving does. Keeping adapters separate is only efficient because the base model’s weights stay frozen and shared — the adapters are purely additive deltas, not replacement layers — so all tenants read the same base. The hard part is that within a single batched forward pass, different requests use different adapters, which would ordinarily force them into separate, un-batched matmuls; systems like S-LoRA solve this with custom CUDA kernels for heterogeneous batching that apply each row’s own adapter inside one batched operation, plus a "unified paging" memory pool that swaps cold adapters between GPU HBM and host DRAM alongside the KV cache — together letting one GPU serve thousands of adapters at up to ~4× the throughput of naïve approaches ^[442]. Both vLLM and NVIDIA’s TensorRT-LLM now ship this multi-LoRA serving path as a standard feature ^[443]. The result is exactly the architecture the question "can’t many businesses share one base model?" implies: a single foundation model held once in HBM, with a distinct low-rank fine-tune per team, per customer, or per task layered on top per request — the operational backbone of every commercial fine-tuning API and of enterprises running dozens or hundreds of task-specific variants on shared infrastructure. Custom fine-tuning at scale is, in hardware terms, a memory-sharing trick built on the fact that the adapter is negligible next to the weights it adapts.

Finally, the whole family so far assumes you compress a big model; the alternative is to design a small one from the start. Efficient-by-design architectures — MobileNet’s depthwise-separable convolutions, EfficientNet’s compound depth/width/resolution scaling — and neural architecture search (NAS), especially hardware-aware "once-for-all" networks that train one super-network and then extract a specialized sub-network per device without retraining, produce compact models natively rather than by shrinking a large one ^{[418][419][420]}. No giant model needs to exist first.

The empirical case for treating inference compute as a first-class scaling axis was made most cleanly by Snell, Lee, Xu, and Kumar in August 2024. They studied two families of test-time strategies against math-reasoning benchmarks — searching against a dense process-based verifier, and letting the model sequentially revise its own draft — and asked the FLOPs-matched question directly: given a compute budget, are you better off spending it on a larger model or on more inference for a smaller one? Their headline finding is that a compute-optimal allocation of test-time compute (chosen adaptively per prompt) is more than 4× more efficient than naive best-of-N sampling, and that a smaller base model with generous test-time compute can outperform a model 14× its size at equal FLOPs — but only in a regime where the smaller model already has non-trivial success on the problem ^[204]. That regime clause matters: test-time compute is a lever that amplifies an existing signal, not a substitute for the pretraining that produced the signal.

OpenAI’s own o1 announcement in September 2024 made the point visually with two log-linear curves — accuracy on AIME rising with train-time reinforcement learning compute, and, on a separate axis, accuracy rising with test-time compute per problem — and reported that o1 reached ~74% on AIME 2024 with "long thinking" versus roughly 12% for GPT-4o ^[205]. The o1 architecture is not public, but the API surface is the operationally important part: the model emits an internal, hidden stream of reasoning tokens before its user-visible answer, billed at output-token rates and capped by a caller-supplied budget (OpenAI’s guidance is to allow up to 25,000 reasoning tokens for hard prompts) ^[206]. Anthropic exposed a nearly identical knob: Claude’s thinking parameter takes a budget_tokens field with a minimum of 1,024 and no fixed maximum below the request’s max_tokens, and the user is billed for the full thinking-token count regardless of whether the tokens are surfaced in the response ^[207]. DeepSeek-R1, published January 2025, made the picture concrete for the open ecosystem: a 671B-parameter Mixture-of-Experts model trained with a rule-reward RL pipeline (GRPO) that induces chains-of-thought thousands of tokens long, with distilled variants down to Qwen-1.5B and Llama-8B that inherit much of the parent’s reasoning behavior by imitation on ~800,000 generated traces ^[208]. Try it: 671B MoE: 37B active.

The "thinking budget" a caller sees is a single scalar, but underneath it sits a family of test-time strategies the model or its serving system can compose.

Self-consistency (Wang et al., 2022) is the simplest: sample independent chain-of-thought completions at nonzero temperature, then take the majority answer. On GSM8K it lifted PaLM-540B accuracy by roughly 18 percentage points over greedy chain-of-thought — a substantial gain, but one that costs the decode compute per answered question ^[209]. Best-of-N generalizes the pattern by scoring each sample with a separate verifier rather than voting on the final answer.

The verifier itself is the second technique, and there are two flavors that matter operationally. An outcome reward model (ORM) scores a completed solution end-to-end — one score per candidate answer. A process reward model (PRM) scores each intermediate reasoning step, so the search can prune branches early rather than waiting for a wrong argument to reach its wrong conclusion. Lightman et al.'s "Let’s Verify Step by Step" showed that a PRM-selected best-of-N reached 78% on a representative MATH subset, substantially above the ORM-selected baseline — the first strong evidence that process supervision was worth the extra labeling cost it demanded ^[210].

Tree- and graph-search over reasoning is the third layer. Yao et al.'s Tree of Thoughts (2023) framed problem-solving as BFS/DFS over a tree whose nodes are natural-language "thoughts," with the model itself acting as both the successor generator and the value estimator ^[211]. By 2024, several groups had layered Monte-Carlo Tree Search on top — AlphaZero-style rollouts using a PRM as the value network — to collect step-level preference data automatically and to guide decoding at test time ^[212]. The o1/o3/R1 wave folded these ideas together: the model is trained (via RL) to internally emit the search rather than requiring an external harness to run it. The "aha moments" DeepSeek reports — the model spontaneously backtracking, re-checking, trying a different framing — are what a learned in-context tree search looks like when it has been compiled into the weights.

The rest of this Part characterized inference as a two-phase workload: compute-bound prefill sizing time-to-first-token, memory-bandwidth-bound decode sizing sustained throughput, and the ratio between them fairly predictable for a given deployment. Reasoning traffic breaks the predictable part.

The first-order effect is that decode-token counts per request explode. A chat turn against a non-reasoning 70B model emits, typically, a few hundred output tokens; the same model wrapped in an agentic harness with a reasoning frontier model might emit tens of thousands of hidden reasoning tokens before it starts on the visible answer. o3’s ARC-AGI evaluation is the extreme case: at "low-compute" configuration it burned roughly 33.5 million tokens per task at ~$26 per task, and at "high-compute" (~172× more compute per problem) it burned 5.7 billion tokens per task at ~$4,560 per task to reach 87.5% on the semi-private set ^[213]. Whatever the pod-capacity arithmetic said before (per Figure 51d — output tokens/sec per GPU times GPU count, divided by tokens per average answer), reasoning workloads move the denominator by one to two orders of magnitude without touching the numerator. The 700-requests/second rack that served a short-chat mix in Figure 51d is a ~20-requests/second rack against a 10,000-thinking-token workload, and a fraction-of-one-request-per-second rack against a hard o3-style problem. Try it: Rack capacity: chat vs reasoning.

The second effect is that per-request compute now varies wildly within the same tenant, on the same endpoint, using the same model. The budget_tokens parameter is caller-controlled: one request asks for 1,024 thinking tokens, the next asks for 32,000. From the scheduler’s point of view this is worse than long context, because the length is not knowable from the prompt alone — the model decides at generation time how much thinking it needs. Continuous batching (Section 70’s Figure 50) still works, but its steady-state assumption — that the batch turns over on a roughly stationary distribution of completion lengths — no longer holds. Batches now contain both 200-token chat completions and 30,000-token thinking traces, and the tail-heavy KV-cache occupancy of the latter directly compresses the batch size the former can enjoy.

The third effect is a workload-mix problem, and it is the one operators have to actively schedule around. A pod that serves a public API sees reasoning and non-reasoning traffic through the same front door. Reasoning requests are heavy, decode-dominated, and comparatively insensitive to inter-token latency (the user is waiting for the answer, not reading tokens as they stream); chat requests are light, prefill-sensitive, and very sensitive to TTFT. Running them on the same continuous batch means either the reasoning traffic starves the chat traffic of KV-cache pages, or the chat traffic’s TTFT-oriented admission policy leaves reasoning throughput on the floor. The industry response, as of 2026, is to route them to separate pools — the same disaggregation logic that separated prefill from decode (Figure 51), extended one axis further to separate reasoning decode from chat decode.

Part XIII already made the point that the correct denominator for cost is the answer, not the token. Reasoning models make this operational. For a workload where each answered question consumes hidden tokens and visible tokens, at effective decode cost per token, the served cost per answer is

Try it: Cost per reasoning answer.

and — because is itself the quantity the operator is scaling to buy accuracy — the compute-per-correct-answer curve looks nothing like the compute-per-token curve. Doubling roughly doubles per-request cost, but (per Snell et al. and the o3 numbers above) accuracy gains from further thinking flatten well before the cost does. Every operator running a reasoning endpoint at scale ends up building the same control loop: measure accuracy-on-target as a function of thinking budget for this workload, pick the point on the curve where the marginal correct-answer cost matches the alternative (routing to a bigger model, escalating to a human), and cap budget_tokens there. The knob is set per-tenant, per-task, sometimes per-prompt-class, and the calibration is a first-class piece of the production stack rather than an API afterthought.

The final consequence closes the loop back to hardware. Reasoning-heavy workloads are the strongest case yet for the decode-specialized accelerators of Part VI: SRAM-everywhere systems deliver tokens several times faster per stream than an HBM-fed GPU, and when most of a request’s wall-clock is hidden decode, a factor-of-several on that phase is a factor-of-several on end-to-end latency. NVIDIA’s LP30 co-processor and the AMD–Cerebras disaggregation announcement in Section 70 are the productization of that observation. The reason "reasoning-time compute" is now discussed alongside "training-time compute" as a scaling axis is that the dominant workload class of 2026 and onward is not one where the model answers immediately, but one where it thinks first — sometimes for a long time — and the hardware wins go to whoever can serve that thinking cheapest.

A vision-language model is not one model but a pipeline: a modality-specific encoder turns pixels or audio into a sequence of embeddings, a projection layer maps those into the language model’s token space, and the LLM prefill consumes them alongside the text prompt. The encoder is comparatively small — Google’s SigLIP-So400m is a 400M-parameter Vision Transformer used by both Gemma 3 and ColPali ^[214], and OpenAI’s CLIP-L variants sit in the 300–400M range — so encoder compute is dwarfed by the LLM behind it. What is not free is the token bill: a single 448×448 image at a 14×14 patch grid produces ~1,024 patch embeddings, and every one of those tokens lands in the LLM’s prefill just like a text token would. A high-resolution image or a multi-image prompt can easily push a request’s effective prompt length past 10,000 tokens before the user has typed a word, and TTFT (Section 57) grows accordingly. Try it: TTFT for an image-heavy prompt.

Audio follows the same pattern with different numbers. Whisper’s encoder processes fixed 30-second windows into ~1,500 encoder frames each, and its decoder autoregresses transcription tokens out of them; production servers run in-flight batching against the encoder to pack shorter clips together and report roughly 2.8–10× throughput gains over the reference batch-size-1 implementation ^[215]. Video is the case where the stack changes shape entirely. A 30-second clip at 24 fps is 720 frames; decoding it on the SMs would burn a substantial fraction of the GPU’s compute before the model does anything, so production VLM servers route the compressed bitstream through NVDEC — the fixed-function video decode engine that has sat alongside the SMs on every NVIDIA GPU since Kepler — and hand raw YUV frames directly to CUDA for the RGB conversion and patch embedding. NVIDIA’s Dynamo multimodal path and the FlashCodec research prototype both call video decode the dominant contributor to end-to-end multimodal latency when it isn’t offloaded to NVDEC, and schedule NVDEC engines as a first-class resource alongside the SMs ^[216]. NVENC — the encode counterpart — is absent from datacenter SKUs (H100, B200), which matters when the output is text or an image; when the output is a video, encode happens on consumer or workstation parts or through software H.264/H.265, which becomes its own scheduling problem.

Where a language model’s latency is dominated by decode’s per-token loop, a diffusion model’s latency is dominated by its sampler. The denoising-diffusion-probabilistic-model (DDPM) formulation of Ho et al. originally ran 1,000 reverse-diffusion steps per image; DDIM cut that to roughly 50 without retraining by reformulating the reverse process as a deterministic ODE. Everything since has been about cutting further. Latent Consistency Models (Luo et al., 2023) collapse the sampler to 1–4 steps by learning a consistency function that maps any noised latent directly to the trajectory’s origin — a 10–100× wall-clock reduction versus 25–50-step DDIM while matching its FID ^[217]. Consistency Models proper (Song et al., 2023) do the same trick from scratch and reach FID 3.55 on CIFAR-10 in a single step ^[218]. Distribution Matching Distillation (DMD, Yin et al., CVPR 2024; DMD2, NeurIPS 2024) distills a diffusion teacher into a one-step generator by matching the student’s output distribution to the teacher’s via KL divergence, reaching ~30× speedup over Stable Diffusion 1.5 at comparable quality — ~90 ms per image, versus ~2.6 seconds at 50 steps ^[219]. The 2024–2025 production image models embed these ideas directly: Stability’s SD3.5 Large Turbo is an 8B-parameter MMDiT distilled with Adversarial Diffusion Distillation to run at 4 steps with classifier-free guidance disabled, versus 28 steps for the non-turbo variant ^[220]; Black Forest Labs' Flux.1-schnell is a 12B rectified-flow transformer that produces high-resolution output in 1–4 steps.

The serving consequence is that a batched-image inference server looks much more like a bulk-compute throughput system than a latency-oriented text server. Each request is forward passes through a model of tens of billions of parameters — for Flux.1-dev at 12B, that means ~38 GB of weights read times per image — and batching across requests is straightforward because there is no equivalent of KV-cache pressure to constrain it: a diffusion sampler’s state is just the current noised latent (a few hundred kilobytes at 1024×1024), not a sequence-length-dependent cache that grows with every step. The right batch size is set by the model’s HBM footprint plus per-sample activation memory, not by memory-bandwidth-per-token arithmetic (Section 16). The sampler is the whole latency budget: a 4-step Turbo pipeline running on one H100 is ~4× the wall clock of a one-step distilled model at the same resolution, and a full 28-step run is 7× that again. Everything the inference engine does — kernel fusion, fp8 weight formats, flash-attention for the diffusion transformer’s own self-attention (Section 11) — trims per-step time; the number of steps is set at model-training time by whichever distillation recipe produced this checkpoint.

Video generation lives an order of magnitude further out. OpenAI’s Sora (as reported by third-party analyses) takes roughly 12 minutes of H100 time to produce 1 minute of output — about 12× slower than real time, versus the "faster than real-time" ratios diffusion image models routinely hit ^[221]. Tencent’s open HunyuanVideo 1.5 renders a short clip in ~75 seconds on a single RTX 4090 ^[222]. Practitioner rules of thumb put one minute of high-quality 1080p video at 800–1,500 still-image-equivalents of compute, which is why video-generation infrastructure is priced by the second of output rather than the token or the image, and why every large lab’s video-generation traffic is served from dedicated pools rather than squeezed into the same fleet that handles chat completions.

Retrieval-augmented generation adds a step Part XVI did not model at all: before the LLM sees the prompt, a vector database is asked which chunks of a corpus are most relevant, and those chunks are pasted into the context. The latency of that lookup, and the token cost of the chunks it returns, are now first-class components of an end-to-end request budget.

Vector search itself has three canonical index families, each with a different position on the recall-latency curve. HNSW (Hierarchical Navigable Small World graphs) stores neighbors as a multi-layer graph and traverses it greedily; it delivers the lowest single-query latencies — routinely low-single-digit milliseconds at 90%+ recall on tens-of-millions of vectors — at the cost of a memory-heavy graph and slow index builds. IVF-PQ (inverted-file with product quantization) partitions the space into Voronoi cells and stores each vector as a compressed subquantizer codebook lookup; it trades a little accuracy for 4–5× compression, faster builds, and better throughput at very large scale. ScaNN, Google’s anisotropic-quantization variant, sits between them and reports roughly 2× higher recall-at-fixed-latency than either HNSW or IVF-PQ on standard benchmarks. GPU acceleration is the recent shift: NVIDIA’s cuVS (built on RAFT) accelerates all three families and reports index-build speedups of 8–40× over CPU implementations, with CAGRA (its GPU-native graph index) delivering 8× search throughput at low single-query latency; it ships as a drop-in backend for FAISS, Milvus, Weaviate, Elasticsearch, OpenSearch, and Apache Lucene, so a serving-side migration to GPU search is now a configuration change rather than a rewrite ^[223].

Above the index, the vector-DB vendor landscape settles on the same distinction. Weaviate reports sub-3-ms mean single-query latency at 90th percentile on the DBPedia-OpenAI dataset; Pinecone reports 74,000 QPS at 90% recall on the BigANN 2023 filtered-search benchmark, with P50 ~45 ms and P99 ~96 ms sustained at 600 QPS across 135M vectors; Milvus, Qdrant, and Turbopuffer round out the field, with Turbopuffer’s S3-backed serverless design trading raw latency for storage cost by paging index segments on demand rather than keeping the whole index in RAM ^[224].

The end-to-end RAG budget is a stack of components, not a single call. A conservative 2026 production pipeline looks like: (1) embed the query — one forward pass through a small embedding model (BERT-scale or a distilled sentence encoder), single-digit milliseconds on CPU or 1–2 ms on GPU; (2) ANN search — HNSW or CAGRA on the vector index, low single-digit milliseconds for a top-100 or top-200 candidate pull; (3) optional hybrid scoring — a parallel BM25 or SPLADE sparse-index query merged into the dense results via Reciprocal Rank Fusion, adding another few milliseconds and consistently improving recall on identifier-heavy or long-tail queries by covering the cases dense retrieval misses ^[225]; (4) rerank — a cross-encoder (a small BERT-scale model that scores each query-document pair jointly) reduces the top-100 to a top-5 or top-10 by running one forward pass per candidate, which is where retrieval latency actually piles up: reranking commonly accounts for 60–85% of pre-generation latency in production pipelines, precisely because it runs serial inferences where the earlier stages ran one ^[226]; (5) context assembly and prefill — the reranked chunks are concatenated into the prompt and the LLM prefill begins.

Late-interaction rerankers are the interesting middle ground. ColBERT (Khattab & Zaharia, 2020) and ColBERTv2 (Santhanam et al., 2022) encode query and document independently — like a dense retriever — but store a 128-dimensional embedding per token rather than one vector per document, and score with a MaxSim operator that adds up each query token’s best match across the document’s tokens. This gives most of a cross-encoder’s precision at retrieval-time cost, because the expensive per-document encoding runs offline; ColBERTv2’s residual compression brings the per-corpus storage down 6–10× versus v1 (MS MARCO from 154 GB to 16–25 GB). ColPali (Faysse et al., 2024) extends the same trick to document images: a 3B-parameter PaliGemma vision encoder produces 1,024 patch embeddings per PDF page, and MaxSim scores query tokens against patches directly — replacing the OCR-then-chunk pipeline that traditional document RAG requires with a single visual retrieval pass ^[227].

The chunk-size / top-k / rerank triangle is where a RAG operator actually tunes the system. Smaller chunks (say 200 tokens) improve retrieval precision but require larger top-k to cover any given answer, driving both rerank cost and prompt length up. Larger chunks (1,000+ tokens) reduce top-k for the same coverage but bring more irrelevant text into the LLM’s context, extending prefill and diluting attention. Cross-encoder reranking is worth its latency exactly when the top-1 chunk matters — legal, medical, identifier-sensitive retrieval — and worth skipping when the LLM can tolerate a noisier top-10 and the token cost of pasting all of them in. The point Part XVI made about prefill scaling with prompt length lands directly here: every retrieved chunk that ends up in the prompt lengthens the compute-bound prefill in exact proportion to its token count, so retrieval doesn’t just add its own latency — it inflates the LLM’s, too. The best-tuned production RAG systems in 2026 aggressively minimize the number of tokens that reach the LLM, because token count is the one quantity that shows up as latency at every downstream stage.

Recall the classic result from Leviathan et al.'s original speculative decoding paper: for a draft of length with per-token acceptance probability , the expected number of tokens produced per target-model forward pass is ^[228]. At and , that’s about 1.9 tokens per pass; at , about 3.4; at , about 4.1. The curve is steep and the ceiling is , so the whole game is: raise , or raise without collapsing . Everything below is one of those two moves. Try it: Tokens per pass at k = 4.

Medusa (Cai et al., 2024) removes the separate draft model entirely. Instead, it bolts extra feed-forward heads — each a single layer with a residual connection — onto the target model’s own final hidden state ^[229]. Head predicts the token at position directly from the hidden state that produced token . Because every head runs in parallel off the same hidden state, drafting costs one target-model forward pass, not draft-model passes — the draft-side sequential tax that vanilla speculative decoding pays disappears.

Each head produces its top- candidates rather than a single guess, giving a tree of possible continuations. Medusa verifies the whole tree in one target-model pass using tree attention — a masked-attention pattern that lets every candidate branch attend only to its own ancestors, so many candidate paths share the same prefix computation. Reported speedups: 2.2× on a frozen backbone (Medusa-1) and 2.3–3.6× with joint fine-tuning of the backbone (Medusa-2), depending on task. The cost is training: the heads must be trained (or distilled) against the target model’s own outputs. Try it: Medusa’s m^K candidate tree.

EAGLE (Li et al., 2024) attacks from a different direction. Instead of predicting tokens, it predicts the target model’s own second-to-top-layer feature vector — the hidden state that gets fed into the LM head — one step ahead ^[230]. The insight is that features are more structured and lower-entropy than tokens: two different sampled tokens can correspond to nearly identical features, so a feature-space predictor sees a smoother target and its predictions match the real feature more often. A tiny transformer layer (the "EAGLE head") consumes the previous feature plus a one-step-ahead token embedding, emits the next feature, and the target model’s own LM head then decodes that feature to a distribution over tokens. On MT-Bench, EAGLE reported 3× over vanilla decoding, 2× over Lookahead, and 1.6× over Medusa.

EAGLE-2 (2024) adds dynamic draft trees: the shape and depth of the tree adapt per-step based on confidence scores from the EAGLE head, which the authors show correlate well with actual acceptance probability. Contexts where the draft is confident get deeper, wider trees; uncertain contexts get pruned to conserve verification budget. This raises reported speedups to 3.05–4.26×, roughly 20–40% over EAGLE-1 ^[231]. EAGLE-3 (2025) drops the feature-prediction constraint that gave the family its name, going back to direct token prediction but fusing features from multiple layers of the target model via a "training-time test" simulation that exposes the draft head to its own inference-time errors during training. Reported single-task speedups reach 6.47× on HumanEval and average 5.51× across five tasks, about 1.4× over EAGLE-2 ^[232].

Lookahead decoding (Fu et al., 2024) needs no draft model, no extra heads, no training — it runs entirely inside the target model ^[233]. The insight is that autoregressive decoding can be reframed as solving a nonlinear system by Jacobi iteration: given a guessed sequence of future tokens, run one forward pass, and the model’s outputs at each position become an improved guess for the next iteration; converged positions match what serial decoding would have produced. Each lookahead step runs two branches in parallel inside a single target forward pass — a lookahead branch that maintains a 2D window of in-flight Jacobi trajectories to mine promising n-grams, and a verification branch that checks the best n-gram candidates from a pool built up over previous steps. Reported: up to 1.8× on MT-Bench, 4× on code-completion with multi-GPU strong scaling. The appeal is practical: no model changes, no training, no draft-model warmup — drop it into a serving stack and go.

Prompt lookup decoding (PLD, Saxena 2023) is the crudest and, on the right workload, the fastest of all. Skip the learned draft entirely: when generating, search the existing input plus generated output for n-grams that match the current suffix and propose the continuation verbatim ^[234]. On RAG workloads where the model’s answer quotes retrieved passages, summarization where the summary reuses input spans, or code editing where the edit is mostly the original file, PLD routinely delivers 2–4× speedups at zero training cost. It fails silently on generative tasks with no input-output overlap — but the acceptance-rate feedback adjusts automatically, since a miss just means one wasted verification step. REST (He et al., 2023) generalises PLD to retrieve n-grams from an external corpus indexed by suffix array; RASD (2025) plugs retrieval-based drafts into a learned scaffold.

Skeleton-of-Thought (Ning et al., 2023) is not speculative decoding at all — it’s a prompting-level trick that gets a related benefit ^[235]. The model first generates a short skeleton (a bulleted outline of the answer), then each bullet is expanded in an independent, parallel generation call. Reported 2.39× speedup across 11 models. It is exact only in the sense that the final answer is genuinely produced by the target model; there is no rejection-sampling guarantee that the distribution matches a single serial run, and quality on tasks that need cross-bullet coherence can drop.

Every one of these variants shares the same weakness: they trade compute for latency. That trade only wins when the GPU has spare compute — i.e., when decode is memory-bandwidth-bound, one request at a time. As batch size grows, decode moves up the roofline toward the compute-bound regime (Section 37’s argument in reverse), and the "free" verification of tokens stops being free. Measured effect: vLLM’s own benchmarking shows speculative decoding losing to plain decode above a workload-dependent concurrency threshold — the verification pass’s compute overhead grows linearly with while the memory savings vanish ^[236]. The 2025 "Batch Speculative Decoding Done Right" paper documents the same pattern across engines: at high concurrency, speculative decoding underperforms standard decoding, and larger makes it worse ^[237]. Try it: Batch-pressure crossover.

Modern serving stacks handle this by making adaptive — long draft chains for lightly-loaded instances, short chains or no drafting at all when the batch is full. TensorRT-LLM and vLLM both expose knobs for this; SGLang’s scheduler picks per-request. As with everything in Part XVI, the real production question is not "which variant wins on a benchmark" but "which one wins on my traffic mix, at my SLO." A fleet with strict TTFT/TPOT targets and low concurrency wants EAGLE-3; a throughput-optimised MoE fleet at 128-way batches may find any form of speculation costs more than it saves.

Part IX established that the KV cache size grows linearly in sequence length: bytes per request. Plugging in Section 16’s 70B-class illustration (0.33 MB/token) at the context lengths of interest is enough to make the scale visceral: a 128K conversation carries roughly 42 GB of cache, a 1M-token prompt carries roughly 330 GB, and a 10M-token prompt would carry over 3 TB — a sizeable fraction of a whole NVL72 rack’s HBM for a single request. This is why long-context serving is memory-bound in a way the short-context case is not: at these sizes the cache itself no longer fits in one GPU’s HBM, and often not in one NVLink domain’s aggregate HBM either. Try it: KV size against GPU memory; KV cache at 128K tokens; 1M-token KV cache; KV cache at 1M tokens; 10M tokens against a rack.

Prefill compute is the other ceiling, and it stays quadratic even after FlashAttention. FlashAttention (Section 11) tiles attention so the score matrix never materializes in HBM — but the arithmetic itself is unavoidable, because every query still attends to every key. What FlashAttention rescues is memory footprint and memory-bandwidth traffic; it does not shrink the FLOP count. On a 1M-token prompt the attention layers alone need on the order of attention FLOPs per layer per head (, ) — an amount of compute that is absolutely serviceable in aggregate on a modern rack, but that dominates time-to-first-token in a way it never did at 4K. Try it: 4N²d attention FLOPs at 1M.

If the cache for one request doesn’t fit on one GPU, the sequence itself has to be sharded across devices. Ring attention, from Liu, Zaharia, and Abbeel in 2023, does exactly this: it splits the sequence into blocks, pins each block’s queries to its assigned device, and rotates the corresponding key and value blocks around the devices in a ring — every device eventually sees every KV block, computing its local queries' attention against each in turn ^[241]. The trick is that the peer-to-peer KV transfers overlap with the on-device blockwise attention math, so the communication is (approximately) hidden behind the compute. This lets sequence length scale linearly with device count with no approximation — devices support an longer context than a single device — and it is the direct ancestor of the context parallelism implementations now shipping in production training frameworks (Part XVIII). The same primitive serves at inference time for the prefill phase of very long prompts: the prompt is chunked across devices, KV shards circulate, and the resulting cache ends up already sharded across the ring, ready for a sequence-parallel decode. Try it: Cache growth and compression.

Sharding grows aggregate memory linearly; a parallel line of work compresses each device’s share of the cache. The four techniques below attack it differently but agree on the empirical premise: at long context, most cache entries contribute almost nothing to the next token’s output.

StreamingLLM (Xiao et al., 2023) begins from a striking empirical observation: in a trained autoregressive LLM, the first few tokens of the sequence receive an outsized share of attention weight at every subsequent position — regardless of what those tokens semantically are. Xiao and colleagues named these attention sinks and showed that keeping their KV entries pinned in cache, alongside a rolling window of the most recent tokens, is enough to keep a Llama-2, MPT, Falcon, or Pythia model coherent out to 4 million tokens of streaming input — with no fine-tuning, and up to 22.2x speedup over sliding-window-with-recomputation baselines ^[242]. The mechanism explains why naïve sliding-window attention collapses: evicting the sink tokens destroys the softmax’s normalization anchor. Try it: StreamingLLM: constant-size cache.

H2O (Zhang et al., 2023) generalizes the same observation. Attention scores follow a power law — a small fraction of tokens (Zhang’s heavy hitters) capture most of the attention mass at any given step, and this set is stable enough across steps to predict from cumulative attention so far. H2O evicts non-heavy-hitter, non-recent entries and reports up to 29x throughput improvement over DeepSpeed Zero-Inference and Hugging Face Accelerate at 20% cache retention ^[243]. SnapKV (Li et al., 2024) extends the idea to the prefill side: rather than waiting for decode to identify heavy hitters, it uses a small "observation window" at the end of the prompt to vote on which prefix tokens matter, and compresses the prefill cache immediately — often keeping under 50% of the raw cache with minimal accuracy loss ^[244]. CacheGen (Liu et al., SIGCOMM 2024) attacks a different axis: it treats the KV cache as a bitstream to be compressed and streamed over the network — a custom tensor encoder plus adaptive bandwidth-aware compression cuts cache size by 3.5–4.3x and end-to-end context-loading delay by 3.2–3.7x ^[245]. And below all of these sits the crudest lever of all — quantizing the cache itself to INT8 or INT4 (Section 72’s quantization machinery applied to K and V tensors), a 2–4x reduction that composes with every technique above. Try it: H2O at 20% retention; SnapKV at 50%; CacheGen’s 4× encoding; INT8/INT4 KV cache.

Everything above is a serving-time technique layered on top of a stock dense-attention model. The more radical answer is to change the model: give attention a structure that doesn’t cost to compute or to cache in the first place. Longformer (Beltagy, Peters, Cohan) established the sparse-attention template: replace global attention with a fixed-size sliding window (each token attends to neighbors on each side, so compute is rather than ), optionally interleaved with dilated windows that reach further at no extra cost, plus a handful of global-attention tokens for tasks that need long-range dependencies. Mistral’s sliding-window attention layers and the banded/dilated patterns in many 2024–2026 long-context open models are direct descendants.

The more radical departure eliminates the KV cache entirely. Mamba (Part XIX) is the state-space variant this primer has already covered; RWKV (Peng et al., 2023) reformulates the transformer as an RNN with linear-attention math that trains in parallel like a transformer but runs inference with constant-time-per-token, constant-memory state ^[246]; Microsoft’s RetNet (Sun et al., 2023) offers the same trick under a different derivation, exposing three equivalent computation modes — parallel for training, recurrent for inference, chunkwise for long sequences — from a single retention operator ^[247]. Their hardware consequence is the point worth stating: because the per-step "cache" is a fixed-size recurrent state rather than a growing KV cache, decode’s memory-bandwidth-bound token-rate ceiling (Section 16) no longer depends on context length at all — a 1M-token conversation reads the same-sized state per step as a 1K-token one. In practice, 2026 production serving is dominated by hybrid architectures (a few full-attention layers for global routing, most layers linear or state-space) that trade some of transformer quality for this bandwidth invariance.

The single most important long-context optimization in real deployments isn’t any of the above; it is prefix caching, and the reason is that production long-context traffic is overwhelmingly repeated long prefixes, not unique long prompts. An agent harness prepends the same system prompt and tool schemas to every request (Part XVI, Section 71); a RAG pipeline prepends the same retrieved documents to every follow-up question; a coding assistant prepends the same repository summary to every edit. These shared prefixes can be tens or hundreds of thousands of tokens, and their KV caches are identical byte-for-byte across every request that shares them.

SGLang’s RadixAttention organizes the paged KV blocks in a radix tree so any two requests sharing a leading prefix automatically reuse each other’s cached computation (Section 70 already noted this); vLLM ships automatic prefix caching on the same principle. Anthropic’s prompt-caching product exposes the same thing to end users of the Claude API: on a cached read, cost drops by 90% (10% of the base input price) and latency by up to 79% on a 100K-token cached document, at the cost of a 25% write-time premium the first time the prefix is seen ^[248]. Gemini and OpenAI expose comparable features. The hardware picture is exactly the one Part XVI built for prefill/decode disaggregation, generalized: prefill for a shared prefix is done once, its KV cache is written to a shared pool (typically HBM-backed with CPU-DRAM or NVMe overflow, Section 70’s KV-cache offloading pattern), and every subsequent request that shares that prefix skips prefill for those tokens entirely — turning a 30-second TTFT on a cache-miss into a sub-second TTFT on a cache-hit, without changing the model, the batch size, or the number of GPUs at all. Try it: Shared prefix blocks, simulated; Cached reads at 10%, writes 125%.

The deeper lesson connects all five techniques in this section. Every long-context serving win — ring attention’s linear scaling, StreamingLLM’s sinks, H2O’s heavy hitters, RWKV’s constant state, prefix caching’s prefill amortization — is the same shape as every other memory-bandwidth win in this primer: it identifies work that would otherwise be redone or carried unnecessarily and either shares it, compresses it, or eliminates it. At 1M tokens the constants have changed, but the roofline hasn’t.

Part XVII’s failure-rate arithmetic applies unchanged to a serving fleet: one H100 has a mean time between failures on the order of 50,000 hours, so a cluster of GPUs expects a failure roughly every five hours, and a 100,000-GPU serving fleet expects one every 30 minutes ^[249][250]. The failure classes that a serving system must plan for, however, differ from training in an important way: they interact directly with in-flight, streaming state. Try it: 10,000 GPUs: every 5 hours; 100,000 GPUs: every 30 minutes; 10⁴ GPUs: failure every 5 h; 10⁵ GPUs: every 30 minutes.

The cheapest recovery, and the one every production stack ships by default, is request-level retry at the router. The vLLM Router (December 2025) exemplifies the pattern: configurable exponential backoff with jitter, circuit breakers that eject a worker on repeated health-check failure, and re-dispatch of the in-flight requests to a surviving pod — which for a stateless-from-the-router’s-perspective request means simply re-running prefill somewhere else ^[255]. This is fine for short prompts and early failures, and terrible for the case that matters: a request that has already streamed thousands of decoded tokens when its GPU dies. Retry means starting over — the caller sees an eight-second stall and a fresh answer that isn’t the one they were watching arrive.

The 2024–2026 research answer is to treat the KV cache itself as replicated state that survives the loss of any one GPU. DéjàVu (Microsoft/ETH, ICML 2024) streams KV-cache updates from each generation step to a companion replica, so a decode-node crash restarts on the replica with the cache already warm — no re-prefill, and recovery time bounded by the streaming lag rather than the prompt length ^[256]. Splitwise (Microsoft, ISCA 2024) and DistServe (OSDI 2024) — introduced in Part XVI as prefill/decode disaggregation techniques — provide the substrate that makes replicated KV caches practical: because the KV cache is already an explicit, transferable artifact that crosses the wire from prefill to decode pool, it can equally cross the wire to a hot-standby decode replica at negligible incremental cost, riding the same GPUDirect-RDMA / NIXL transfer path ^[257][258]. Mooncake (Moonshot AI’s Kimi production stack, 2024) generalizes this into a KV-cache-centric fabric that pools CPU DRAM, SSD, and RDMA across the cluster as a shared cache tier, from which any surviving decode GPU can pull a lost request’s cache back into HBM — turning what would have been a re-prefill into a memory-hierarchy fault ^[259]. Both vLLM and SGLang now ship Mooncake integration as a standard P/D-disaggregation option ^[260].

These translate to three deployment postures, familiar from any HA system but with numbers set by KV cache size (Section 11) rather than disk:

Which posture is right depends entirely on fault-domain sizing — how many in-flight requests a single failure destroys. A pod (Part XVI, Section 71) that runs one giant tensor-parallel model across all 72 GPUs in a GB200 NVL72 has a blast radius of the entire pod’s continuous batch: hundreds of concurrent requests, all lost together on any hardware fault in the domain. A pod that runs eight independent 8-GPU model replicas (64 of the 72 GPUs; nine would fill the rack) has one-eighth the blast radius per fault, at the cost of not being able to serve models that don’t fit in 8 GPUs. Beyond a certain blast-radius threshold, per-pod retries are simply unacceptable: the retry storm from a single failure saturates the surviving pods and cascades into a fleet- wide outage, which is exactly the failure mode that circuit breakers exist to prevent ^[261]. Try it: Blast radius of one 72-GPU model; NVL72 replica blast radius.

The circuit-breaker pattern in a 2026 inference stack is concrete: the router tracks per-pod error rate and p95 latency, and when either exceeds a threshold it opens the circuit — new requests are refused outright (returning HTTP 429 with a Retry-After header) or routed to a warm secondary, and no new load is added to the failing pod until a cooldown probe succeeds. Retries themselves are constrained by a latency budget, not a fixed count: a request that has already spent 80% of its SLO in queue and prefill cannot afford a five-second exponential-backoff retry, so the router must decide, at the moment of failure, whether the remaining budget is enough for a retry to plausibly succeed or whether the request should fail fast rather than block the socket ^[262]. This is the direct inference-time analog of training’s "lose minutes, not days" checkpoint cadence: the goal is to bound the worst SLO impact per failure, not to eliminate failures.

Multi-region failover extends the same idea one level up. Traffic is normally pinned to the nearest region for network latency and KV-cache locality (a request whose system prompt is in RadixAttention’s prefix tree in us-east-1 is far cheaper served there than in eu-west), and a regional outage triggers a drain — new requests are redirected to another region’s pods while in-flight streams complete or fail in place. Modern locality-aware balancers (SkyLB, and the SkyWalker follow-up) preserve prefix affinity across the failover by replicating a multi-region prefix trie, so drained-away traffic still hits warm caches on the far side instead of paying full prefill cost on every request during the incident ^[263]. In practice, the frequent-but-small failures (single-GPU XIDs, node reboots) are absorbed by per-pod retries and hot standbys; regional failover is reserved for the rare-but-large events (a datacenter power incident, a bad model-serving deploy) whose blast radius no in-pod mechanism can contain — and it is precisely those rare events that dominate the public status-page timelines of Anthropic and OpenAI when incidents do occur ^[264].

Everything above assumes a failure announces itself. SDC does not, and LLM inference is the first mass workload in which stochastic output and corrupted output are literally indistinguishable to the caller. For high-stakes inference — a medical triage answer, a financial trade, a safety-critical robotics decision — the emerging pattern is result validation: run the same request on two independent pools (different racks, ideally different hardware generations) and compare outputs, either exhaustively (dual-run) or on a sampled fraction of traffic (probabilistic screening). Because LLMs are non-deterministic by default, comparison isn’t bit-identical; it uses greedy decoding or a fixed seed, or compares distributions rather than samples. The cost is the obvious one — up to 2× the compute for full dual-run — and this is why the technique is deployed narrowly, at the top of the value stack, rather than fleet-wide. Cheaper approximations exist: activation checksums at layer boundaries, periodic canary-prompt screening that runs a known prompt through every GPU on a schedule and flags any divergence from the reference output, and layer-wise output-hash comparison between hot-standby replicas ^[265][266]. None of these completely close the gap — a plausible-looking wrong token that both replicas agree on because they share the same silicon defect model survives dual-run — which is why SDC in inference is, in 2026, an open problem rather than a solved one, and why the trend is toward out-of-distribution redundancy: pair an NVIDIA decode with a Cerebras or AMD decode on the same request, so that a systematic silicon fault cannot corrupt both sides identically.

The takeaway from all of this is a shift in how the serving stack is budgeted. Every technique in Part XVI was pitched as utilization: keep the accelerator busy, amortize the weight read, drive tokens/sec/GPU up. The techniques in this section are pitched as headroom: keep enough compute, memory, and cache replication in reserve that the mean-time-to- failure of the fleet does not translate into an equal mean-time-to-SLO- breach for the caller. The two are in direct tension — a 2× hot-standby decode pool is a 2× throughput hit — and the 2026 state of the art is mostly about buying the fault-tolerance property that the workload actually needs, at the smallest possible cost in utilization, rather than about eliminating faults themselves.

A complementary lever is activation checkpointing (also called gradient checkpointing or recomputation). The activations saved during the forward pass for use in the backward pass can dominate memory; checkpointing keeps only a subset and recomputes the rest during the backward pass by re-running the forward from the nearest saved point. The classic result (Chen et al., 2016) is that checkpointing every -th layer of an n-layer network cuts activation memory from to at the cost of roughly one extra forward pass (about +33% compute). Try it: √n checkpointing, animated.

The reason this is a real memory problem, and not just a bookkeeping detail, is backpropagation’s own mathematical structure. Computing the gradient of the loss with respect to an early layer’s weights requires the chain rule to pass back through every layer between that weight and the final loss — and each of those intermediate steps needs the activation (the layer’s actual output value on this specific forward pass) that was produced at that point, because the local derivative at a nonlinearity or a matmul depends on the actual input value it saw, not just on the weights. A naive implementation therefore keeps every layer’s activations resident in memory from the moment the forward pass produces them until the backward pass finally reaches that point and consumes them — for an n-layer network this is memory, and for a frontier-scale model with activations dwarfing the parameter count itself at long sequence lengths, this is often the single largest consumer of GPU memory during training, ahead of even the 16-bytes-per-parameter optimizer state quantified above.

Checkpointing breaks that all-or-nothing tradeoff by being deliberately selective about which activations survive the forward pass. Concretely, with a 9-layer network checkpointed roughly every `sqrt(9) = 3`rd layer: during the forward pass, the inputs to layers 1, 4 and 7 and the output of layer 9 are saved, while every other layer’s activation is computed, immediately consumed by the next layer, and then discarded — it does not need to survive until the backward pass starts, because the backward pass will simply reconstruct it on demand. During the backward pass, when gradient computation reaches a layer whose activation was saved, it’s read straight from memory at no extra cost, exactly as in an unmodified backward pass. When it reaches a layer whose activation was not saved — say, layer 5 — it re-runs a short forward pass starting from the nearest earlier checkpoint (layer 4) up through the missing layer, regenerating exactly the activations needed for that segment, uses them immediately for the gradient computation, and discards them again once that segment’s gradients are computed. This is the same "recompute a known value instead of storing it forever" trade this primer’s KV-cache section examined from the opposite direction — there, caching avoided recomputation; here, deliberately not caching (most activations) and recomputing on demand is what saves the memory, because activations, unlike a KV cache’s entries, are only needed briefly during the backward pass rather than for the entire remaining generation. Try it: Step through the 9-layer example.

A minimal illustration of the mechanism (not a production implementation — real frameworks like PyTorch’s torch.utils.checkpoint handle randomness, non-determinism, and autograd-graph integration far more carefully) makes the store/recompute split concrete:

Try it: Run this code step by step.

Selective recomputation refines this further — recomputing only the operations that are cheap in FLOPs but expensive in memory (notably the attention softmax), recovering most of the memory saving for a few percent of overhead ^[460][461]. The choice of which layers to checkpoint is itself a tunable knob rather than a fixed sqrt(n) rule in practice: checkpointing more densely spends more compute on recomputation but shortens each recompute segment, while checkpointing more sparsely saves more memory at the cost of longer recompute segments — the sqrt(n) spacing is simply the point that minimizes total added compute for a given memory budget, not a hard requirement, and production frameworks (DeepSpeed, FSDP) expose this as a configurable policy so it can be tuned against the actual memory ceiling of the hardware in use rather than the asymptotically optimal but workload-blind default.

A pretraining step is one forward, one backward, one optimizer update — the arithmetic is fixed and the whole cluster runs the same kernel in lock-step. An RLHF step is a pipeline. First, the current policy generates a batch of completions from a set of prompts — an autoregressive decode job, hundreds to thousands of tokens per prompt, running on the same weights that are about to be updated. Second, one or more auxiliary networks score the completions: a reward model, a reference model whose log-probabilities feed the KL-penalty that keeps the policy from wandering off the base distribution, and (for PPO) a value/critic head. Third, the policy takes a gradient step against a surrogate objective computed from those scores. The generation step is compute the pretraining loop never had to do at all, and OpenRLHF’s own profiling puts it at roughly 80% of end-to-end RLHF wall time — meaning the training system is now, in effect, an inference system that occasionally does backprop.^[267] Try it: Rollout share of RLHF time.

The algorithm that started the modern era is Proximal Policy Optimization (PPO). Schulman et al. (2017) replaced trust-region policy optimization’s constrained update with a much simpler clipped surrogate objective: the ratio between the new and old policy’s probability of each action is clipped to so that a single update cannot move the policy too far, permitting multiple minibatch epochs per rollout without collapse.^[268] Applied to LLMs by OpenAI’s InstructGPT work, PPO became the RLHF default — but it is heavy: at any moment four LLM-sized networks must be resident (policy, reference, reward, critic), gradients flow through two of them, and the on-policy assumption means each batch of generations is thrown away after one or two update epochs. Try it: PPO’s four resident networks.

Direct Preference Optimization (DPO), from Rafailov et al. (2023), removed the RL loop entirely for the preference-alignment case. Their reparameterization shows that the KL-constrained reward-maximization objective has a closed-form optimal policy in terms of the reward, which can be inverted so that the reward is implicit in the policy itself; the training objective then reduces to a simple pairwise classification loss over preferred/dispreferred completion pairs.^[269] There is no reward model, no critic, no rollout — just two forward passes (policy and frozen reference) per pair and a gradient step, roughly the cost of ordinary SFT. Llama 3 is the industrial datapoint: Meta’s post-training pipeline ran six rounds of reward modeling, rejection-sampled SFT, and DPO — dropping PPO from Llama 2’s recipe because DPO gave comparable quality at a fraction of the compute and engineering surface.^[270] Try it: DPO memory and compute.

Group Relative Policy Optimization (GRPO), introduced in DeepSeekMath (Shao et al., 2024), attacks PPO from the other side: it keeps the online RL loop but discards the critic. For each prompt, the policy samples a group of completions (typically 8–64), each is scored, and the advantage of each completion is computed as its reward normalized against the group’s mean and standard deviation — the group itself serving as the baseline that a learned value function would otherwise provide.^[271] That single change removes one full LLM-sized network (the critic) from the resident set and the entire training-signal pipeline for it, cutting RLHF’s memory footprint sharply — one of four networks, but one of only two that carry optimizer state, so closer to 40–45% of the bytes in a full fine-tune and eliminating the critic’s own instability. GRPO is the algorithm behind DeepSeek-R1’s reasoning training: pure RL on top of a base model, rule-based rewards (accuracy plus format) rather than a learned reward model, and enough compute for chain-of-thought behavior to emerge — AIME 2024 pass@1 rising from 15.6% to 71.0% on the DeepSeek-V3-Base policy.^[272] DeepSeek deliberately kept the reward rule-based rather than neural precisely because scaling neural reward models invites reward hacking — the policy learning to exploit the reward model’s blind spots — and controlled comparisons in the R1 lineage put the hacking rate at 12–16% for RL-from-base runs versus under 1% for SFT-only baselines.^[273] Try it: GRPO memory saving vs PPO; GRPO with rule-based reward.

The infrastructure question for any of these algorithms is where each of the resident networks lives. A naive PPO implementation co-locates all four on the same GPU pool with the same parallelism plan — which wastes memory (each has its own optimizer state, but only the actor needs one) and, more expensively, forces the generation stage to run under a training-shaped weight layout that is poorly suited to autoregressive decode. The two dominant open frameworks solve this in different ways.