Part I · Chapter 3

GPU Architecture: From Silicon to FLOPS

How modern accelerators actually work — and how architecture shapes every design choice

Book: AI/ML Infrastructure from Silicon to Scale | Author: Venkat Vinjam | Cross-ref: Ch 3A: Accelerator landscape Ch 3B: GPU roadmap Ch 4: HBM deep dive Ch 5: Thermal/TCO

"The GPU is not a magic compute box. It is a collection of physical tradeoffs — transistor density, wire length, heat dissipation, memory bandwidth — that determine, with mathematical precision, exactly how fast your model can go and why it cannot go faster." — Principal Performance Architect's GPU Principle

This chapter builds the mental model of the GPU that every subsequent chapter depends on. Chapters 3A and 3B go deeper into the accelerator landscape and specific GPU generations respectively. This chapter covers the architectural concepts that are stable across generations — the compute hierarchy, tensor cores, the memory subsystem, and the transformer engine — so you have the vocabulary and mental model to reason about any GPU, current or future.

3.1 The GPU Compute Hierarchy

H100 SXM5 COMPUTE HIERARCHY [SHIPPED]: GPU (full chip): 132 Streaming Multiprocessors (SMs) Peak BF16: 989 TFLOPS | Peak FP8: 1,979 TFLOPS Peak FP32: 67 TFLOPS (non-Tensor Core) TDP: 700W Streaming Multiprocessor (SM): 4 warp schedulers (can issue 4 independent warps/cycle) 4 × 16 = 64 CUDA Cores (FP32 units) 4 × 1 Tensor Core (4th gen) — handles 16×16×16 BF16 GEMM per cycle 228 KB L1 Cache / Shared Memory (SMEM) — programmer-configurable 65,536 × 32-bit Registers Warp (fundamental execution unit): 32 threads executing in lockstep (SIMT — Single Instruction, Multiple Thread) Thread divergence causes serialization (branches in GPU kernels are expensive) Memory access: 32 threads issue coalesced 128-byte HBM transactions Key occupancy formula: Occupancy = min(active_warps, max_warps) / max_warps At 64 regs/thread: 65536/64 = 1024 threads = 32 warps → 50% occupancy At 32 regs/thread: 65536/32 = 2048 threads = 64 warps → 100% occupancy Register spilling: overflow → L2 → 200-800× latency penalty

Figure 3.1 — H100 SXM5 compute hierarchy. 132 SMs each contain 64 CUDA Cores, one 4th-gen Tensor Core, 228 KB SMEM, and 65K registers. All SMs share a 50 MB L2 cache and 80 GB HBM3 at 3.35 TB/s. NVLink 4.0 provides 900 GB/s for intra-node tensor parallelism.

3.2 Tensor Cores and the GEMM Engine

TENSOR CORE GENERATIONS: V100 (1st) A100 (3rd) H100 (4th) B200 (5th) Peak BF16 (TC): 112 TFLOPS 312 TFLOPS 989 TFLOPS ~4,500 TFLOPS Min FP format: FP16 BF16 FP8 NVFP4 Tile size: 16×16×16 16×16×16 16×16×16 WGMMA (larger) Sparse support: No 2:4 2:4 Structured TENSOR CORE OPERATION (4th gen, H100): Input: A [16×16 BF16] × B [16×16 BF16] Output: C [16×16 FP32] accumulated Throughput: one 16×16×16 GEMM per warp per cycle CRITICAL ALIGNMENT REQUIREMENTS (the silent 30% tax): BF16 GEMM: matrix dimensions must be multiples of 16 FP8 GEMM: matrix dimensions must be multiples of 16 If M=4097 instead of 4096: padded to 4112 → 0.4% waste (negligible) If M=13 instead of 16: padded to 16 → 18.75% compute waste (significant!) Production consequence: attention head count × d_head must be ≡ 0 mod 16. LLaMA-3 70B: H_q=64, d_head=128 → 64×128=8192 ✓ (multiple of 16)

3.3 The Memory Subsystem

The H100 has a six-tier memory hierarchy. See Chapter 4 (§4.2) for the complete five-tier breakdown from registers through HBM. The sixth tier in multi-GPU systems is peer GPU HBM accessed via NVLink.

Tier	H100 Capacity	Bandwidth	Latency	Programming Model
Registers	256 KB/SM (~33 MB total)	~unlimited (0 cycles)	1 cycle	Compiler-managed; register pressure → spill
L1 / SMEM	228 KB/SM (configurable)	~80–100 TB/s aggregate	20–30 cycles	Explicit: __shared__ memory
L2 Cache	50 MB (unified)	~12 TB/s	200–400 cycles	Hardware-managed; cannot be bypassed
HBM3	80 GB	3.35 TB/s	400–800 cycles	Global memory; primary bottleneck
NVLink peer	Neighbor GPU HBM	~900 GB/s (NVSwitch fabric)	~1–5 µs	NCCL collective; transparent to framework
Host DRAM	1–2 TB via PCIe	128 GB/s (PCIe Gen5)	~10–50 µs	cudaMemcpy; avoid during inference

3.4 The Transformer Engine (H100 Innovation)

THE TRANSFORMER ENGINE — the key H100 architectural differentiator: Problem: FP16/BF16 is accurate but has 2× the bytes of FP8. FP8 is fast but requires per-tensor scaling to prevent overflow. Solution: Hardware unit that: 1. Monitors activation magnitudes per tensor (automatic, no user code) 2. Computes per-tensor scale factors dynamically 3. Casts to FP8 E4M3 (weights) or FP8 E5M2 (gradients) for GEMM 4. Accumulates in FP32 (prevents numerical degradation) 5. Outputs in BF16 for next layer FP8 E4M3: 1 sign + 4 exponent + 3 mantissa bits → range ±448, more precision FP8 E5M2: 1 sign + 5 exponent + 2 mantissa bits → range ±57344, more range RESULT: ~1.9× GEMM throughput over BF16 with automatic quality preservation. This is the primary reason H100 delivers ~3× training speedup over A100 (FP8 TC throughput × NVLink BW improvement × architectural improvements). FP8 on H100 [SHIPPED]: 1,979 TFLOPS (2× BF16 989 TFLOPS) NVFP4 on B200 [SHIPPED]: 9,000 TFLOPS (FP4 doubles over FP8)

3.5 GPU Generations at a Glance

GPU	Year	Peak BF16	HBM	HBM BW	NVLink BW	TDP	Key Innovation
V100 SXM2	2018	112 TFLOPS	32 GB HBM2	900 GB/s	300 GB/s	300W	First Tensor Cores (FP16)
A100 SXM4	2020	312 TFLOPS	80 GB HBM2e	2.0 TB/s	600 GB/s	400W	3rd gen TC, BF16 native, MIG
H100 SXM5	2022	989 TFLOPS	80 GB HBM3	3.35 TB/s	900 GB/s	700W	Transformer Engine, FP8, NVLink 4.0
H200 SXM5	2024	989 TFLOPS	141 GB HBM3e	4.8 TB/s	900 GB/s	700W	HBM3e (memory upgrade only)
B200 SXM	2024	~2,250 TFLOPS (FP8: 4,500)	192 GB HBM3e	8.0 TB/s	1,800 GB/s	1,000W	Dual-die MCM, NVFP4, NVLink 5.0
GB300	2025	~2,800 TFLOPS (FP4: 9,000)	288 GB HBM4	12.0 TB/s	1,800 GB/s	1,400W	HBM4, FP4 training (MLPerf v5.1)

[SHIPPED] for V100–B200. [ROADMAP] for GB300 specs except MLPerf-verified performance.

3.6 SXM vs PCIe: The NVLink Question

SXM vs PCIe FORM FACTOR: SXM (Server eXchange Module): Direct die-to-die connection to NVSwitch via copper traces (no PCIe) NVLink BW: 900 GB/s (H100 SXM) vs 400 GB/s (H100 PCIe) TDP: 700W (SXM) vs 350W (PCIe) — SXM requires liquid or active air Deployment: In HGX/DGX/OVX servers — not in standard PCIe servers PCIe: Standard GPU card with PCIe x16 host interface Half the NVLink bandwidth → TP communication is 2× slower Lower power envelope → fits in standard datacenter racks WHEN SXM MATTERS: TP=2 or TP=4 within a node: 900 vs 400 GB/s (SXM 2.25× faster) A 70B model at TP=2 with one AllReduce per layer (80 layers, ~0.5 GB each): SXM: 0.5 GB / 900 GB/s = 0.56 ms overhead per layer PCIe: 0.5 GB / 400 GB/s = 1.25 ms overhead per layer Difference: 0.69 ms × 80 layers = 55 ms per sequence → significant at short output WHEN PCIe IS ACCEPTABLE: DP-only (no TP/PP): no NVLink communication during inference Small models fitting on single GPU: no inter-GPU communication needed

3.7 AMD MI300X — The Memory Moat

Architectural approach: MI300X uses a chiplet design — 6 compute XCDs (CDNA 3) + 4 memory controller MCDs on a passive silicon interposer — to break the reticle-limit constraint that forces NVIDIA to use dual-die MCM on Blackwell. The payoff is 192 GB HBM3 in a single package.

Property	H100 SXM5	MI300X	H200 SXM5
Peak BF16	989 TFLOPS	1,307 TFLOPS	989 TFLOPS
HBM Capacity	80 GB HBM3	192 GB HBM3	141 GB HBM3e
HBM Bandwidth	3.35 TB/s	5.3 TB/s	4.8 TB/s
TDP	700W	750W	700W
LLaMA-3 70B TPOT (batch=1)	42 ms	27 ms	29 ms
70B single-GPU fit (BF16)?	No (140 GB)	Yes	No (140 GB)
Ridge point	295 FLOP/byte	247 FLOP/byte	206 FLOP/byte

The Memory Moat — cost consequence: MI300X can serve LLaMA-3 70B on a single GPU (192 GB HBM3). H100 requires TP=2 (two GPUs, one AllReduce per layer). At $25K/GPU and $3/GPU-hr: one fewer GPU per serving instance saves $25K capex and $13K/year in OpEx on a 5,000-hour annual utilization. At 100 serving instances: $2.5M capex and $1.3M/year OpEx savings.

3.8 Custom Accelerators — When to Look Beyond GPUs

Accelerator	Architecture	Key Strength	When to Use	When NOT to Use
Google TPU v5p	Systolic array (256×256)	Gemini training at trillion-token scale	Large-scale training on GCP, JAX ecosystem	Irregular kernels, non-GCP deployment
AWS Trainium2	Matrix engine + NeuronLink	60% cost reduction vs H100 for LLM training on AWS	AWS-native training at scale	Non-AWS infrastructure; non-Neuron models
Intel Gaudi 3	MME + HBM3e, IBFabric	Open software stack, OAM form factor	Open ecosystem, heterogeneous clusters	CUDA kernel dependency, NVIDIA-only stack
Groq LPU	SRAM-only, streaming dataflow	Sub-1ms TTFT for small models (<13B)	Ultra-low latency, real-time inference	Large models, high-throughput batching

Key Takeaways

Core Concepts

The GPU is organized as: GPU → SMs → Warps → Threads. Tensor Cores operate on 16×16×16 tiles; SM occupancy is limited by register pressure (65K regs per SM). At 128 regs/thread, max 512 threads → 25% occupancy. Register spilling adds 200-800× latency per spilled value.
The Transformer Engine (H100) automatically manages FP8 ↔ BF16 casting per tensor — no user code required. This delivers 1.9× GEMM throughput over BF16 and is the primary driver of H100's 3× improvement over A100 for transformer workloads.
SXM and PCIe are the same compute die with different host interfaces. SXM provides full NVLink bandwidth (900 GB/s H100). PCIe provides half (400 GB/s). For TP=2 inference on 70B, SXM saves ~55ms per sequence vs PCIe — significant for latency-sensitive applications.

Cost & Productivity Impact

MI300X's 192 GB HBM3 eliminates the need for TP=2 to serve LLaMA-3 70B (140 GB BF16), saving one GPU per serving instance. At 100 concurrent serving deployments: $2.5M capex savings and $1.3M/year OpEx reduction — purely from fitting the model on fewer GPUs.
H100 → H200 upgrade provides zero compute improvement (same TFLOPS) but 43% bandwidth improvement (3.35 → 4.8 TB/s). ROI is positive only for bandwidth-bound workloads: long-context inference, decode at small batches, models that barely fit in memory. Zero ROI for compute-bound prefill workloads.
Tensor Core alignment requirements (dimensions must be multiples of 16) are a 'silent tax'. A model with d_head=130 instead of 128 wastes 1.5% of Tensor Core throughput. At 1,000 GPUs running 24/7 at $3/GPU-hr: $39K/year wasted on misalignment alone.

Principal-Level Insight

The H100 ridge point (295 FLOP/byte) represents the regime boundary — not a target. Principal architects design workloads to land on the right side of the ridge for their optimization goal: compute-side for throughput (prefill), memory-side for latency (decode). The same model crosses sides at different batch sizes, requiring different optimization strategies.
Custom accelerator evaluation requires answering three questions before building a business case: (1) What fraction of your operations are compatible with the accelerator's ISA? (2) What is the team migration cost in engineer-months? (3) At what GPU-spend breakeven does the custom accelerator's per-TFLOP cost advantage amortize the engineering investment? The answer to (3) is typically >$10M/year GPU spend.
GPU generation transitions change the regime classification of workloads. An operation at AI=200 FLOP/byte is compute-bound on V100 (ridge 124), near-ridge on H200 (ridge 206), and memory-bound on H100 (ridge 295). When your fleet upgrades GPUs, re-evaluate every kernel's regime classification — optimizations valid for V100 may be wrong on H100.

Review Questions

FoundationalBuild the mental model — no math required

Explain what an SM is and how the relationship between SM count, warp count, and register pressure affects GPU occupancy. Why does low occupancy not always mean low performance?
What is the Transformer Engine in H100, and why does it enable FP8 training without requiring manual scaling from the user? What does it actually do at the hardware level?
Why does the MI300X's 192 GB HBM matter more for inference than for training? What specific serving scenario does it enable that H100 cannot serve without multi-GPU setup?

IntermediateApply and calculate — production context

An H100 cluster is running LLaMA-3 70B training with TP=8. Each AllReduce communicates a tensor of 256 MB (one weight matrix gradient). (a) Calculate the time for this AllReduce over NVLink (900 GB/s, ring algorithm). (b) If the GEMM for that layer takes 8 ms, what is the communication overhead as a percentage? (c) At what tensor size does the AllReduce become the dominant cost?
A kernel uses 128 registers per thread on H100. (a) What is the maximum SM occupancy? (b) If you reduce to 64 registers by restructuring the kernel (without changing FLOP count), what happens to occupancy and why might throughput increase even if the kernel does more L2 accesses? (c) What Nsight Compute metric would you use to diagnose this?
Compare TPOT for LLaMA-3 70B BF16 at batch=1 on H100 (80 GB, 3.35 TB/s), H200 (141 GB, 4.8 TB/s), and MI300X (192 GB, 5.3 TB/s). Which GPU is fastest, and by how much? Does the ridge point change matter for batch=1?

Principal LevelDesign, decide, and defend — architecture reviews

Your company is buying 1,000 new GPUs for a 70B inference serving fleet (BF16, P99 TPOT target 20ms, 50K concurrent requests, average 2048-token context). Create a detailed comparison of: H100 (TP=2, $25K/GPU), H200 (TP=1, $35K/GPU), and MI300X (TP=1, $30K/GPU). Include: total GPUs needed, total capex, annual OpEx at $3/GPU-hr, and the breakeven analysis. Show all calculations.
An engineering team proposes migrating from NVIDIA H100 to Google Cloud TPU v5p for GPT-4 class model training. Design the evaluation framework: what metrics must you measure, what is the minimum trial scale, what quality signals are non-negotiable, and what is the decision timeline? What hidden costs is the team likely underestimating?