GPU Architecture.

A CPU is latency-optimized: a few fat cores, deep pipelines, big out-of-order windows and huge caches, all to make one thread finish as fast as possible. A GPU is the opposite trade — throughput-optimized: thousands of small, in-order lanes that hide latency by keeping enormous numbers of threads in flight and switching between them for free. Everything your instinct learned tuning a JVM inverts here.

CUDACUDANVIDIA's parallel programming model and runtime for writing GPU kernels. exposes a nested hierarchy. You write a kernel — a function run by many threads in parallel — and launch it over a grid of thread blocks. A thread has private registers and an index; a warp is exactly 32 threads that execute in lock-step (the true hardware scheduling unit — treat 32 as a hard constant); a block is up to 1024 threads guaranteed to run on a single SMStreaming Multiprocessor (SM)The GPU's core compute unit — holds CUDA cores, Tensor Cores, registers and shared memory., able to cooperate via shared memory and barriers; a grid is all the blocks of one launch, independent and free to run in any order on any SM.

SIMT (Single Instruction, Multiple Threads) is the model: the 32 threads of a warp share instruction issue and each runs the same instruction on different data. If threads in a warp take different branches — if (threadIdx.x % 2) — the warp diverges: the hardware runs each path serially with the other lanes masked off, throttling you to a fraction of peak. It is the GPU analog of a mispredicted branch, except the cost is structural, not statistical.

One systems fact to internalize: kernel launches are asynchronous (like submitting to an executor and getting a future) and cost ~5–20 µs each — and a real transformer forward pass is thousands of launches. That overhead is why CUDA Graphs (capture the whole sequence once, replay it) matter for DecodeDecodeThe memory-bound second phase: generate one token at a time, autoregressively., where each step is tiny.

The SMStreaming Multiprocessor (SM)The GPU's core compute unit — holds CUDA cores, Tensor Cores, registers and shared memory. is the GPU's core — the rough analog of a CPU core, except one SM runs dozens of warps concurrently. An H100 (GH100) has 132 active SMs; an RTX 4090 (AD102) has 128. Every block of your grid is distributed across them by a hardware block scheduler (the 'GigaThread engine').

A modern SM is partitioned into 4 sub-partitions, each with its own warp scheduler, a slice of the register file, CUDA cores (scalar ALUs for the elementwise work — activations, normalization, residual adds, softmax) and one Tensor Core. Each cycle a scheduler picks a warp whose operands are ready and issues an instruction; when one warp stalls on a memory load it simply issues from another. That warp-swapping is latency hiding — the entire reason a GPU tolerates ~500 ns HBM latency without dying.

Two on-chip resources shape performance. The register file (~256 KB / SM) is per-thread private state, partitioned across all resident threads — more registers per thread means fewer resident warps. Shared memory / L1 (up to ~228 KB / SM on Hopper) is a software-managed scratchpad that makes tiled matmul and FlashAttention fast. Hopper adds the TMA (Tensor Memory Accelerator) to asynchronously copy large tiles into shared memory, keeping the Tensor Cores fed.

Occupancy = resident warps ÷ the SM's maximum — a proxy for how much latency-hiding parallelism you gave the scheduler. It is limited by whichever resource runs out first (registers/thread, shared mem/block, threads/block). The nuance every optimizer learns: higher occupancy is not always faster. A tiled GEMM or FlashAttention kernel deliberately burns registers and shared memory for data reuse, running at modest occupancy but near-peak throughput. You want enough warps to hide latency, not the maximum.

If you remember one thing about why LLMs run on GPUs specifically, it is Tensor CoresTensor CoreSpecialized unit doing small matrix multiply-accumulates — the workhorse for LLMs.. A Tensor Core performs a matrix-multiply-accumulate (MMA) on small tiles in a single hardware operation — $D=A\times B+C$ — versus a CUDA core doing one scalar fused multiply-add at a time.

A transformer is, at the FLOP level, a stack of large matrix multiplications: the QKV projections, the attention-score matmul, the output projection and the big FFN/MLP GEMMs. Tensor Cores execute these an order of magnitude faster per watt because MMA amortizes instruction overhead and data movement across a whole tile — on a modern datacenter GPU, ~10–30× the FP32 CUDA-core throughput. This is the reason LLM training and PrefillPrefillThe compute-bound first phase: process the whole prompt in parallel, fill the KV cache. are feasible.

Tensor Cores accept progressively narrower inputs, and narrower = more throughput and less memory traffic: TF32 → FP16/BF16 → FP8 → FP4, roughly doubling throughput at each step. BF16 keeps FP32's 8-bit exponent (wide dynamic range, fewer overflow headaches) at the cost of mantissa bits, which is why it became the default LLM dtype. The accumulation is done in higher precision (typically FP32) internally, so you get narrow-input speed with acceptable error — the core trick, and the hardware side of QuantizationQuantizationStoring weights (and sometimes activations/KV) in fewer bits to cut memory and bandwidth..

Each generation added a format: Volta (V100) introduced Tensor Cores (FP16); Turing added INT8/INT4; Ampere (A100) added TF32, BF16 and sparsity; Hopper (H100) added FP8 and the Transformer Engine; Blackwell (B200/GB200) added FP4. The Transformer Engine dynamically picks per-layer precision and manages FP8 scaling factors automatically — a big part of why H100 crushed A100 on transformer workloads beyond raw spec deltas.

This is the section that matters most for inference. DecodeDecodeThe memory-bound second phase: generate one token at a time, autoregressively. is memory-boundMemory-boundLimited by memory bandwidth, not compute — the GPU waits on HBM, not the math units.: each generated token requires reading the entire model's weights (and the growing KV cacheKV cacheStored key/value vectors for all past tokens, so each new token needs only one forward pass.) from HBMHBMHigh-Bandwidth Memory — the stacked DRAM on a datacenter GPU (the 'VRAM'). while doing comparatively little math. Your token/sec ceiling is set by HBM bandwidth, not FLOPs — the arithmetic-intensityArithmetic intensityFLOPs performed per byte moved from memory — decides compute- vs memory-bound. argument from Module 1.

Read the bandwidth column of the hierarchy top to bottom and you fall off a cliff at every level: on-chip SRAM (registers/shared) is ~10× the bandwidth of HBM, and HBM is ~50× the bandwidth of PCIe to host. The whole game of a fast kernel is keep data as high up this pyramid as possible and touch HBM as few times as you can — which is literally what FlashAttention does, tiling attention so the score matrix never round-trips to HBM.

HBMHBMHigh-Bandwidth Memory — the stacked DRAM on a datacenter GPU (the 'VRAM'). (High Bandwidth Memory) is DRAM dies stacked vertically and wired to the GPU through a silicon interposer with an extremely wide bus (thousands of bits) rather than the narrow-fast bus of GDDR. Wide-and-moderate-clock beats narrow-and-fast for aggregate bandwidth and energy per bit. Generations climb steeply: HBM2e (~2.0 TB/s, A100) → HBM3 (~3.35 TB/s, H100 SXM) → HBM3e (~4.8 TB/s H200, ~8 TB/s B200).

Why this dominates decode: emitting one token from a 70B FP16 model streams ~140 GB of weights through the math units, so $\text{tokens/sec}\approx \text{HBM bandwidth}/\text{bytes read per token}$. Double the bandwidth (H100 → B200) and single-stream decode roughly doubles — without any FLOP increase mattering. The RTX 4090's GDDR6X (24 GB at ~1.0 TB/s) is excellent value for a 7B–13B model but ~3× less bandwidth than H100. The roofline below makes the memory-bound reality visual.

A single GPU has finite HBM (80 GB on H100). Big models don't fit, and even when they do you want multiple GPUs for Throughput vs latencyThroughput vs latencySystem-wide tokens/sec across all requests vs the speed felt by one request.. How fast GPUs talk to each other and to the host sets the ceiling on the parallelism strategies of Module 5.

PCIePCIeThe general-purpose bus linking GPU, CPU and NIC; the slow path for GPU-GPU traffic. is the baseline bus every GPU sits on: PCIe Gen4 x16 gives ~32 GB/s per direction, Gen5 x16 ~64 GB/s. That is fine for loading weights from host once, but far too slow for the per-layer all-reduce traffic of tensor parallelism — which is exactly why NVLink exists.

NVLinkNVLink / NVSwitchHigh-bandwidth direct GPU-to-GPU interconnect within a server (far faster than PCIe). is NVIDIA's proprietary point-to-point GPU interconnect, roughly an order of magnitude faster: ~600 GB/s (A100, 3rd gen), ~900 GB/s (H100, 4th gen), ~1.8 TB/s (B200, 5th gen). So H100↔H100 over NVLink is ~900 GB/s vs ~128 GB/s bidirectional over PCIe 5 — the difference between tensor parallelism being viable or being throttled by communication. NVSwitch then adds a switch chip connecting all GPUs in a node (typically 8) in a non-blocking all-to-all fabric, so every GPU talks to every other at full NVLink bandwidth at once.

That is what makes an 8-GPU HGX/DGX box behave like one big accelerator for collective ops. Cross-node you drop to InfiniBand/Ethernet (~400 Gb/s-class per NIC), a large step down — hence the rule 'keep a model within one NVLink domain if you can.' The parallelism strategies that ride on this are Module 5.

MIGMIG (Multi-Instance GPU)Hardware partitioning of one GPU into isolated smaller GPU slices. (Multi-Instance GPU, on A100/H100 and newer) hardware-partitions one GPU into up to 7 isolated instances, each with a dedicated slice of SMs, L2 and HBM — with real fault and performance isolation, not just time-slicing. Each instance looks like a smaller standalone GPU to CUDA.

Reach for it for small models and multi-tenancy — serving several 1–7B models or many low-QPS endpoints where a whole H100 per model wastes silicon — and for predictable latency SLAs, since there is no noisy-neighbor interference between instances. It is the GPU analog of K8s resource requests/limits with real isolation. Skip it for any single model that needs more memory or compute than one slice provides; partitions are static-ish (reconfiguring requires draining), so MIG suits steady multi-tenant serving, not spiky single-tenant bursts. Contrast plain time-slicing (software context-switch, no isolation) and MPS (concurrent processes sharing SMs, no hard isolation) — MIG is the only one with hardware-enforced boundaries.

On the concurrency side: by default kernels on the default stream run serially in issue order. A CUDA stream is an ordered queue — operations within a stream run in order, operations in different streams may overlap. The classic win is overlapping a host→device copy on one stream with a compute kernel on another, so PCIe transfer hides behind computation; small DecodeDecodeThe memory-bound second phase: generate one token at a time, autoregressively. kernels can also be co-scheduled on different SMs when neither fills the GPU.

Scheduling happens at three layers: the driver places stream operations onto hardware queues; the block scheduler distributes blocks to SMs; each SM's warp schedulers pick ready warps cycle-by-cycle. As an engineer you mostly control layer 1 (streams, graphs, batching) and let the hardware handle the rest — and CUDA Graphs collapse a whole stream of launches into one replayable unit, killing per-kernel overhead in autoregressive decode.

It is not one thing; it is a stack of reinforcing decisions, all pointed at feeding matmuls and moving bytes. (1) HBM3 at ~3.35 TB/s directly raises the decode token/sec ceiling because decode is bandwidth-bound — the biggest lever. (2) 4th-gen Tensor Cores + FP8 roughly double usable math throughput vs BF16 for prefill and training, and FP8 also halves the bytes moved per weight, easing the bandwidth wall from the other side.

(3) The Transformer Engine automates FP8 scaling per layer so you actually realize that speedup without numerical blowups. (4) NVLink 4 (~900 GB/s) + NVSwitch make 8-GPU tensor parallelism efficient, so large models run without communication becoming the bottleneck. (5) A large L2 (~50 MB) + TMA + big shared memory keep more data on-chip, meaning fewer HBM round-trips and more useful work per byte of bandwidth.

Every item either adds bandwidth or reduces the bytes that must cross it. That is the throughline of the whole module: since inference decode is memory-bandwidth-boundMemory-boundLimited by memory bandwidth, not compute — the GPU waits on HBM, not the math units., NVIDIA's fastest inference gains come from HBM-generation jumps and precision reduction (FP16 → FP8 → FP4), not from piling on FP32 FLOPs.

The figures below are widely documented, approximate and dtype-dependent (dense FP16 Tensor Core throughput is quoted; vendor 'PFLOPS' headlines often add sparsity and lower precision, so confirm the footnote before capacity planning). Read the table as an inference engineer and a few rules fall out.

Lead with bandwidth, then capacity. Decode throughput tracks memory bandwidth, so the 4090's ~1.0 TB/s vs H100's ~3.35 TB/s is a bigger practical gap for single-stream generation than any FLOP comparison, and B200's ~8 TB/s HBM3e is the headline reason it roughly doubles inference over H100. Capacity decides whether the model fits at all — a 70B FP16 model (~140 GB) doesn't fit on one 4090 or even one H100, but fits comfortably on one B200 (~192 GB), collapsing an 8-GPU-with-NVSwitch problem into a single card. That is often a larger cost win than raw speed.

The 4090 has no NVLinkNVLink / NVSwitchHigh-bandwidth direct GPU-to-GPU interconnect within a server (far faster than PCIe)., so multi-4090 setups are capped at PCIe (~32–64 GB/s) — great for one model on one card, poor for scaling a big model across cards. FP8/FP4 support (H100 onward; FP4 on Blackwell) means newer parts read fewer bytes per weight, hitting the bandwidth wall from both sides. And TDP is your real TCO input: GB200's ~2.7 kW is a whole module (2 GPUs + Grace CPU), so perf-per-watt — not peak PFLOPS — is what shows up on the bill.

GB200 is not a single GPU. It is a superchip: 2× B200 + 1 Grace CPU joined by NVLink-C2C, a coherent ~900 GB/s CPU↔GPU link that lets the GPUs treat Grace's large LPDDR5X pool as fast, coherent extended memory. NVIDIA packages these into GB200 NVL72 racks (72 GPUs in one NVLink domain), so 'GB200' can mean the 2-GPU superchip or the 72-GPU rack — numbers differ by ~36×. Net: for a home lab or one small model, the 4090 is unbeatable value; for production serving, choose on memory bandwidth and capacity within an NVLink domain — FLOPs are rarely your binding constraint at decode time.