Chapter 11. Synthesis, after Day 5, 55 minutes

TPU and GPU architecture, side by side

This chapter puts the previous ten side by side. Each section takes one axis — compute, memory, network, software and workload fit — and lines up NVIDIA, Google and AWS on it, with numbers. Earlier chapters' callouts made these comparisons locally; here they form one argument.

Three design bets

GPU: many independent cores, hardware-managed latency hiding. An H100 has 132 SMs, each with 4 Tensor Cores, a large register file and up to 64 resident warps. A hardware L2 cache sits behind them. Independent kernels run on different SMs, and the warp scheduler hides memory latency by switching threads [N-wp] [N-tune, §1.4]. The Scaling Book's summary: GPUs "just work" on new tasks more often and lean far less on the compiler. The price is that their performance is harder to reason about and harder to bring to the roofline [SB12].

TPU: a few large cores, compiler-managed everything. An Ironwood chip has two TensorCores, each a single instruction stream driving two 256×256 systolic arrays. 64 MiB of VMEM per core is filled by DMAs scheduled at compile time, and XLA compiles the whole step. The Scaling Book's summary: TPUs "can often get closer to peak roofline performance with less work", and cost less to build. The price is the compiler: static shapes, whole-program compilation, and poor tolerance for irregular access patterns.

Trainium: in between. A Trainium2 chip has 8 NeuronCores. Each has a 128×128 systolic tensor engine plus separate vector, scalar and general-purpose SIMD engines running concurrently, a 28 MiB SBUF scratchpad in 128 partitions, and a dedicated 2 MiB accumulation buffer, PSUM [A-trn2-nki]. It is systolic and scratchpad-based like a TPU, and multi-core like a GPU, though with 8 cores instead of 132.

Google's own framing, in the PyTorch blog, is the same: "a TPU TensorCore is fundamentally a single-threaded execution processor … On the contrary, a GPU includes many identical compute cores" [T-pt].

Compute and memory, per chip

ChipBF16FP8HBMHBM BWBF16 intensity
TPU7x Ironwood2.31 PFLOP/s4.61 PFLOP/s192 GiB7.38 TB/s313
TPU 8t996 TFLOP/sp5.98 PFLOP/sp216 GiB6.53 TB/s153p
TPU 8i1.1 PFLOP/sp8.81 PFLOP/sp288 GiB8.6 TB/s128p
H100 SXM989 TFLOP/s1.98 PFLOP/s75 GiB3.35 TB/s295
B200 (HGX)2.25 PFLOP/s4.5 PFLOP/s168 GiB8 TB/s281
GB200 (NVL72)2.5 PFLOP/s5 PFLOP/s173 GiB8 TB/s313
Rubin (VR NVL72)4 PFLOP/s17.5 PFLOP/sp268 GiB19.2 TB/s208
Trainium2667 TFLOP/s1.3 PFLOP/s96 GiB2.9 TB/s230
Trainium3671 TFLOP/s2.52 PFLOP/s144 GiB4.9 TB/s137

p provisional: from the JAX Pallas hardware table, not in Google's published spec.

Sources are in scripts/specs.py. NVIDIA per-GPU figures are dense, halved from sparse, and per-GPU values for rack products are derived from rack totals [N-gb200] [N-rubin]. Trainium3's FP8 is MXFP8 [A-trn3]. TPU 8 BF16 and FP8 figures are provisional [J-hw]. Rubin's figures are NVIDIA's product-page claims; CES coverage quoted higher memory bandwidth.

What the table says:

  • Ridge points cluster around 300 for this generation's flagships in BF16: Ironwood 313, GB200 312, H100 295, B200 281. The batch-size intuitions of Chapters 1 and 7 transfer directly between them.
  • Inference-oriented designs push the ridge down. TPU 8i (128, provisional), Trainium3 (137) and Rubin (208) spend more on bandwidth relative to BF16 compute, and put their FLOPs into low precision.
  • HBM capacity per chip: Rubin and 8i at 288 GB, Ironwood at 192 GiB, B200 and GB200 at about 180–186 GB, Trainium3 at 144 GiB.
On-chip structureH100Ironwood (per chip)Trainium2 (per chip)
Independent cores132 SMs2 TensorCores8 NeuronCores
Matrix units528 Tensor Cores4 MXUs, 256×2568 tensor engines, 128×128
Programmer-managed SRAMabout 30 MB shared memory (228 KB per SM)128 MiB VMEM224 MiB SBUF
Hardware cache50 MB L2nonenone
Registersabout 32 MBabout 512 KB of vector registersnot compared

Sources: [N-tune, §1.4], [SB12], [J-hw], [A-trn2].

Network: switch against torus

Scale-up domainTopologyPer-chip bandwidthIn-network reduction
H100 HGX8 GPUsNVSwitch, 1 hop450 GB/s each wayNVLink SHARP
GB200 NVL7272 GPUsNVSwitch, 1 hop900 GB/s each wayNVLink SHARP
Ironwood9,216 chips3D torus of OCS-joined cubes600 GB/s each way over 6 linksnone
TPU 8t9,600 chips3D torus2× Ironwoodnone
TPU 8i1,152 chipsBoardfly, ≤ 7 hops2.4 TB/s totalCollectives Acceleration Engine on chip
Trainium264 chips4×4 torus per instance, rings across 41.28 TB/s totaldedicated collective cores
Trainium3144 chipsNeuronSwitch, all-to-all2.0–2.56 TB/s total—

Sources: [SB12], [N-gb200], [SB2], [G-8dd], [A-trn2-sys], [A-trn3-sys].

Beyond the scale-up domain, GPUs and Trainium use an InfiniBand or Ethernet fat tree. An 8-GPU H100 node has 400 GB/s of egress, 50 GB/s per GPU. A GB200 NVL72 rack has 3.6 TB/s [SB12]. TPUs use DCN between pods: 12.5 GB/s per Ironwood chip, and up to 4× that on 8t.

What each design makes cheap.

  • A switch makes all-to-all traffic cheap, and makes small, latency-bound collectives fast: one hop, plus SHARP for reductions. An AllToAll inside an NVLink node costs about half what it does on a TPU ring [SB12].
  • A torus makes scale cheap. The pod adds links in proportion to chips, so per-chip bandwidth stays constant up to about 9,000 chips with no switch silicon. Ring collectives along an axis get full bandwidth at 10 MB messages.
ChipTimeWhy
TPU7x Ironwood
83.9 µstorus bisection, 3 axes
H100 SXM
336 µscrosses 8 nodes over scale-out
GB200 (NVL72)
18.6 µsone-hop switched domain of 72
Trainium3p
16.8 µsone-hop switched domain of 144

Idealized models from the Scaling Book (TPU ch. 3, GPU ch. 12): no latency floor, no software overhead. The book notes GPU AllReduce reaches about 370 of 450 GB/s only at about 10 GB messages, and multi-node AllToAll may not reach 50%. Matmul: D = 8,192, F = 28,672. Decode uses fp8 weights and KV, fp8 math where the chip has it. Trainium3's per-direction link rate and Rubin's figures are provisional.

Comparison explorer. Choose up to four chips and a workload: a matmul, a decode step, an AllGather or an AllToAll. See each chip's predicted time and what bounds it. Start with an AllToAll of 1 GB, then move N from 8 to 64 to 512: the NVLink domain leads while the job fits inside it, and the TPU's 3D torus catches up once the job spans racks. The models are idealized; see the note below the table.
Worked example: 1 GB collectives, idealized (scripts/ch11_compare.py)
8 chips64 chips512 chips
AllGather, Ironwood5.0 ms (one ring)1.7 ms (3 axes)1.7 ms
AllGather, H1001.9 ms2.5 ms (8 nodes)2.5 ms
AllGather, GB200 NVL721.0 ms1.1 ms1.1 ms (8 racks)
AllToAll, Ironwood1.25 ms0.078 ms0.020 ms
AllToAll, H1000.28 ms0.31 ms0.039 ms
AllToAll, GB200 NVL720.14 ms0.017 ms0.039 ms

At 8 chips, a small TPU slice uses one ring, and the GPU node's one-hop switch wins by 2.5–9×. At 64 chips, a 4×4×4 cube has three axes and a large bisection, but NVL72 still wins the AllToAll by about 4.5×. Past one rack, the GPU falls to its scale-out network, and the torus's bisection, which grows as N2/3N^{2/3}, wins.

Practical reach differs from these formulas. The Scaling Book reports that an 8×H100 AllReduce peaks around 370 GB/s of 450 only near 10 GB messages, and that multi-node AllToAll may not reach 50% of theory [SB12]. A v5e AllGather reaches about 95% of peak at about 10 MB [SB3].

Training rooflines, side by side

The same derivation as Chapter 6, with each platform's bandwidth [SB12]:

α = C / W for collectivesDP/FSDP: tokens per chipTP limit for F = 28,672
v5p, 3 axes2,550850about 11 per axis
Ironwood, 3 axes11,5353,845about 2.5 per axis
H100 in node2,2002,200about 13
H100 across nodes2,4752,475(TP stays in node)
B200 across nodes5,6255,625(TP stays in node)

GPU recipes follow from the 8-GPU boundary: tensor parallelism inside a node, then pipeline, expert and data parallelism across nodes. LLaMA-3 used TP 8 × PP 16 × DP 128; DeepSeek-V3 used EP 64 × PP 16 × DP 2 (Chapter 6). TPU recipes follow from the pod: FSDP and TP inside one ICI domain, data parallelism across pods. The Ironwood row is the outlier, discussed in Chapter 6.

Software and programming model

NVIDIAGoogle TPUAWS Trainium
Kernel languageCUDA, CUTLASS, Triton; threads and blocks; implicit cachingPallas (Mosaic TPU); grid and BlockSpecs; compiler-pipelined DMANKI; explicit DMA to SBUF; programmer picks the engine
Compilationeager kernels; CUDA Graphs replay; torch.compile/Inductor per graphwhole-program XLA; static shapesneuronx-cc to NEFF
PyTorchnativeTorchTPU (PrivateUse1, eager and torch.compile → XLA)TorchNeuron (PrivateUse1, beta)
CollectivesNCCL kernels on SMsHLO ops scheduled by XLA, on DMA enginesdedicated collective cores
Partitioning a chipMIG splits one GPUslices of whole chipsLNC merges NeuronCores

Sources: [N-graphs], [T-pt], [A-torch], [A-nki], Chapters 9–10.

Where each design has the structural advantage

WorkloadAdvantageWhy
Large dense pretrainingTPU (v5p-class α), GPU close behindthousands of chips at full ICI bandwidth; FSDP and TP inside one domain. On Ironwood, α favours larger batches per chip
MoE training with many expertsGPU NVL72 inside 72; TPU torus at multi-rack scaleAllToAll is one hop in a switch; bisection grows with a 3D torus
Low-latency, small-batch decodeGPU NVL72, TPU 8ilatency-bound collectives want few hops: one hop on NVLink, at most 7 on Boardfly
Long-context decode throughputwhichever has the most HBM bandwidth and capacity per dollarKV-read-bound (Chapter 7): Rubin and 8i lead on bandwidth per chip
Embedding-heavy recommendationTPU (SparseCore)dedicated gather/scatter engines, 5–7× on v4
Research code with new ops or dynamic shapesGPUeager execution and a hardware cache tolerate irregularity; XLA does not

Price changes these comparisons, and it is not in any spec sheet. In February 2025, the Scaling Book's Google Cloud prices gave v5e 5.8 × 10¹⁷ FLOPs per dollar, against 3.3 × 10¹⁷ for an H100 [SB8].

Convergence, 2025–26

  • GPUs are adding TPU features.
    • Blackwell added 256 KB of Tensor Memory per SM, because accumulators no longer fit in registers [SB12].
    • Tensor Cores have grown each generation: fed by one warp on Ampere, a whole SM on Hopper and two SMs on Blackwell [SB12].
    • Both move toward larger matrix units with a dedicated accumulator store fed by bulk copies (TMA), closer to an MXU fed from VMEM than to warps in registers.
  • Trainium moved toward the GPU. It replaced Trn2's torus with Trn3's NeuronSwitch [A-trn3-sys], and AWS has previewed NVLink Fusion for Trainium4 [A-trn4].
  • TPU split by workload. 8t keeps the training torus; 8i adopts a low-diameter network and a large SRAM [G-8dd].
  • PyTorch moved off torch-xla on both TPU and Trainium, to native PrivateUse1 backends (Chapter 9).

The evidence against full convergence: GPUs still have over 100 independent SMs, a hardware cache and eager execution. TPUs still have two cores per chip, a compiler-managed scratchpad and whole-program compilation. The design bets of the first section have not changed; the boundaries between them have moved.

Worked problems from the book

Worked example: Counting units and moving bytes

From the Scaling Book's GPU chapter quizzes [SB12].

  • Vector lanes: an H100 has 132⋅4⋅32=16,896132 \cdot 4 \cdot 32 = 16{,}896 fp32 CUDA cores. A TPU v5p has 2⋅4⋅8⋅128=8,1922 \cdot 4 \cdot 8 \cdot 128 = 8{,}192 VPU ALUs, about half, at a similar clock.
  • AllGather inside an 8×H100 node: bf16[4096, 65536] is 537 MB. At 450 GB/s, 537×106⋅7/(8⋅450×109)537 \times 10^6 \cdot 7 / (8 \cdot 450 \times 10^9) = 1.04 ms; in practice about 1.5 ms.
  • AllGather across exactly 2 nodes: the node level is B/514B / 514 GB/s and the cross-node level B/800B / 800 GB/s. The node level is the bottleneck, one reason DeepSeek-V3 used 2-way data parallelism.
  • B200 rooflines: inside a node, TP stays compute-bound up to F/2,500F / 2{,}500 ways. Across nodes, data parallelism needs 5,625 tokens per GPU, because FLOPs grew 2.25× while scale-out bandwidth per node stayed at 400 GB/s. GB200's 3.6 TB/s per rack restores it.

Sources