gpu

How it is built

Architecture overview

A data-center GPU is a throughput machine: instead of making one thread fast, it keeps tens of thousands of threads in flight and switches between them every cycle, so the arithmetic units rarely wait for memory. The chip is built from many identical streaming multiprocessors (SMs; AMD calls them compute units, CUs), each with its own register file, a small fast scratchpad and arithmetic pipelines, plus dedicated matrix units that do most of the floating-point work of machine learning. The SMs share a large L2 cache and a stack of high-bandwidth memory (HBM) in the same package, and GPUs talk to each other over a dedicated link fabric instead of through the host. The one idea that makes it work is latency hiding by oversubscription: a thread block stays resident on its SM until it finishes, so switching between warps is free, and the programmer's job is to supply enough parallelism and to move data in patterns the memory system can serve. NVIDIA H100 (Hopper) and B200/B300 (Blackwell) and AMD Instinct MI300X (CDNA 3) are the reference designs; they differ in packaging and details, not in this shape.

Parts:

• Execution cores — SMs/CUs running warps (32 threads on NVIDIA, 64 on AMD CDNA) in lockstep with per-thread masks — simt-warp-scheduling (no page yet), CUDA and GPU languages, SIMD programming
• Matrix units — tensor cores / matrix cores that multiply small tiles in one instruction — systolic-matrix-units (no page yet), low-precision-number-formats (no page yet), Tiled matrix multiplication
• On-chip memory — register file, shared memory/LDS, L1, shared L2 — gpu-memory-coalescing-banking (no page yet), CPU caches and the memory hierarchy, Memory ordering
• Chip fabric and coherence — crossbar between SMs, L2 slices and memory controllers — network-on-chip (no page yet), cache-coherence-protocols (no page yet)
• Device memory — HBM stacks on an interposer beside the die — dram-organization (no page yet)
• Scale-up links — NVLink/NVSwitch or Infinity Fabric between GPUs — Collective communication, Ring all-reduce
• Graphics path — raster, texture and ray-tracing units, absent on pure compute parts — triangle-rasterization (no page yet), mipmapping-texture-filtering (no page yet), Bounding volume hierarchy
• Power and reliability — clock/voltage control, memory error correction — power-management-dvfs (no page yet), hamming-codes (no page yet), Reed-Solomon codes
• Performance model — which resource bounds a kernel — Roofline model

 host CPU ── PCIe Gen5 / NVLink-C2C ──┐
                                      ▼
 ┌─────────────────────────── GPU package ──────────────────────────────┐
 │  grid scheduler ──► GPC 0 … GPC n ──► SM 0 … SM m                    │
 │   SM:  warp schedulers │ FP32/INT32 pipes │ tensor cores │ registers │
 │        shared memory + L1 (one SRAM pool) ◄── async copy engine (TMA)│
 │                 │                                                    │
 │      crossbar / on-chip network                                      │
 │                 ▼                                                    │
 │            L2 cache (partitioned, shared by all SMs)                 │
 │                 ▼                                                    │
 │      memory controllers ──► HBM stacks (same package, interposer)    │
 └───────────────┬──────────────────────────────────────────────────────┘
                 │ NVLink ports / Infinity Fabric links
                 ▼
         NVSwitch or direct links ──► other GPUs in the node and rack

H100 is one large die plus HBM stacks; Blackwell joins two reticle-sized dies into one logical GPU; MI300X stacks eight accelerator dies on four I/O dies. Software sees one device in each case, so the programming model (grid, block, warp, thread) stays stable while the packaging changes underneath.

How the parts fit together

Execution cores: warps, blocks and occupancy

A kernel launches a grid of thread blocks; each block is placed on one SM and stays there. Inside a block, threads are grouped into warps of 32 that execute one instruction together; if threads branch differently, the inactive lanes are masked off while the others run, so divergence costs time but not correctness (CUDA Programming Guide, "Warps and SIMT"). Unlike SIMD programming on a CPU, the vector width is not part of the program. The mechanism is on simt-warp-scheduling (no page yet); two consequences follow.

• Resident state is the budget. An H100 SM holds up to 64 warps (2,048 threads) and 32 blocks, has 64K 32-bit registers, and lets one thread use at most 255 (Hopper Tuning Guide). Full thread occupancy leaves 32 registers per thread; a kernel that needs 128 registers per thread runs at a quarter of it. That is the register-versus- parallelism trade every kernel makes.
• Latency is hidden by other warps, not by caches or speculation. The Best Practices Guide states that executing other warps while one is stalled is the only way to hide latency. A CPU core spends silicon on prediction and reordering; a GPU spends it on arithmetic, which AMD's HIP documentation names as the source of its throughput density.

AMD CDNA differs in detail: wavefronts have 64 threads, a CU has four SIMD units, and workgroups are mapped to CUs non-deterministically, so code must not assume a placement (HIP hardware-implementation documentation). The language layer is CUDA and GPU languages.

Matrix units

The tensor core (NVIDIA) or matrix core (AMD) multiplies and accumulates a small tile per instruction, so registers and issue slots are spent on tiles instead of scalars (systolic-matrix-units (no page yet)). The H100 SM has four fourth-generation tensor cores, with twice the MMA rate of an A100 SM on the same types and four times with FP8 (Hopper in-depth post). The format list is the other half of the design: FP8, FP16, BF16, TF32 and INT8 on Hopper; FP6 and block-scaled NVFP4 on Blackwell, where 16-value blocks carry an FP8 scale and the tensor an FP32 scale (Blackwell Ultra post); FP8, TF32, FP16, BF16, INT8 and sparsity on CDNA 3 (white paper). Formats are introduced in low-precision-number-formats (no page yet); their use in training and inference is in Mixed-precision training and Weight quantization.

The contract between instruction and kernel is a tile shape. Hopper's wgmma is asynchronous and issued by a warpgroup of 128 threads; Blackwell's tcgen05.mma is issued by a single thread, can span a CTA pair of cooperating blocks, and accumulates into tensor memory, a per-SM scratchpad of 128 lanes by 512 columns of 32-bit cells (256 KB) outside the register file (PTX ISA). The software that cuts a large product into these tiles is Tiled matrix multiplication; CUTLASS is the open library that does it for Hopper and Blackwell.

On-chip memory

Memory is a ladder whose rungs are explicit in the programming model.

• Registers are private to a thread. Running out spills to local memory, which is off-chip and costs as much as global memory (Best Practices Guide).
• Shared memory is a programmer-managed scratchpad per block in the same SRAM as L1: 256 KB per SM on H100, of which up to 228 KB can be shared memory (Hopper Tuning Guide). It has 32 banks of 4 bytes; lanes that hit different banks are served together, lanes that hit one bank are serialised (Best Practices Guide). AMD's equivalent is the 64 KB local data share.
• L2 is shared by all SMs: 50 MB on H100, 126 MB on GB200 (tuning guides). On MI300X each accelerator die has a 4 MB L2, and a 256 MB Infinity Cache on the I/O dies sits between L2 and HBM as a memory-side cache (CDNA 3 white paper).
• Global memory is HBM behind L2. A warp's contiguous accesses are merged into as few 32-byte transactions as possible, which is why lane k should read word k (Best Practices Guide; gpu-memory-coalescing-banking (no page yet)). The CPU analogue, with caches instead of threads hiding latency, is CPU caches and the memory hierarchy.

Hopper moves data between levels with the Tensor Memory Accelerator (TMA): one thread posts a descriptor with the tensor shape and block coordinates, and hardware generates the addresses and copies a block between global and shared memory asynchronously, tracked by a barrier that counts bytes as well as arrivals (Hopper in-depth post). Kernels become pipelines in which some warps only move data and others only compute; FlashAttention-3 uses this on H100 and reports a 1.5-2.0x speedup with FP16 attention (Shah et al. 2024). The idea it builds on is Kernel fusion, and the algorithm that avoids writing the full score matrix to HBM is FlashAttention.

Thread block clusters add a level above the block: up to 8 blocks (16 with an opt-in on H100 and B200) are scheduled together on one GPC and can address each other's shared memory, with about 7x faster block-to-block exchange than via global memory (Hopper post). Coherence is deliberately weak: CDNA 3's vector L1 has a "very relaxed coherency model", and L2 is where hardware coherence starts (white paper). See Memory ordering and cache-coherence-protocols (no page yet).

Chip fabric and device memory

Between the SMs and memory sit a crossbar (network-on-chip (no page yet)) and a partitioned L2. HBM is a stack of DRAM dies on an interposer, and bandwidth is its headline: about 3 TB/s from five HBM3 stacks on H100 SXM5, 1.55 TB/s on A100-40GB (Hopper Tuning Guide); 5.3 TB/s and 192 GB on MI300X, 6.0 TB/s and 256 GB on MI325X (white paper). Rows, banks and refresh, and why wide-and-slow beats narrow-and-fast, are dram-organization (no page yet).

Packaging is where vendors diverge. MI300X splits the GPU into eight accelerator dies and four I/O dies and presents one device or, in partition modes, up to eight independent GPUs. Blackwell puts two reticle-sized dies behind a die-to-die link rated at 10 TB/s with one coherent L2 (Blackwell Ultra post). The reason is the reticle limit: a bigger GPU needs several dies, joined fast enough that software need not care which die holds the data.

Scale-up links

Hopper's fourth-generation NVLink has 18 links of 50 GB/s bidirectional, 900 GB/s total, against 12 links and 600 GB/s on A100 (Hopper Tuning Guide); Blackwell Ultra quotes 1.8 TB/s and domains up to 576 GPUs (Blackwell Ultra post). NVSwitch adds in-network reduction (Hopper post). AMD connects an eight-GPU node all-to-all with Infinity Fabric and attaches each GPU to the host by x16 PCIe Gen 5 (white paper). What runs over these links is Collective communication, mostly Ring all-reduce or tree variants.

Graphics path and power

H100 has no RT cores or video encoder, and only two TPCs can run vertex, geometry and pixel shaders (Hopper post). Graphics parts keep the fixed-function units that compute parts drop: triangle-rasterization (no page yet), mipmapping-texture-filtering (no page yet), and RT cores that accelerate BVH traversal and ray-triangle tests (Turing post; structure in Bounding volume hierarchy).

The H100 SXM5 module has a 700 W power budget and the PCIe card 350 W (Hopper post, preliminary), so clocks follow power and temperature (power-management-dvfs (no page yet)). Device memory is error-corrected, and the Best Practices Guide notes the cost of ECC on scattered accesses. SECDED codes of the Hamming family are the textbook way to correct a flipped bit (hamming-codes (no page yet)); the vendor documents I opened do not say which code HBM uses, so treat that as an assumption. Reed-Solomon codes are the alternative for burst errors.

Data and control flow

One matrix-multiply kernel (Hopper style):

1. The grid scheduler splits the launch into blocks or clusters and places them on SMs with enough free registers and shared memory.
2. A producer warp issues TMA loads of the next A and B tiles from HBM through L2 into shared memory, bypassing registers.
3. A transaction barrier flips when the bytes have arrived; a consumer warpgroup issues asynchronous wgmma instructions that read the tiles from shared memory and accumulate in registers.
4. While tensor cores work, the producer fetches the next tile (double buffering): the overlap the hardware is designed for.
5. After the last K step the warpgroup converts accumulators (for example to BF16 or scaled FP8), applies the epilogue and writes the tile back through shared memory and TMA. Time goes to step 3 when tiles are large enough to be compute-bound, otherwise to steps 2 and 5. The Roofline model names the line: roughly 1 PFLOP/s of dense FP16 against 3 TB/s gives a ridge point order of 300 FLOP per byte (derived from the Hopper post's preliminary figures). A batch-1 decoding step reads each FP16 weight once for one multiply-add, about 1 FLOP/byte, and is memory-bound by more than two orders of magnitude.

One training step on eight GPUs:

1. Each GPU runs forward and backward on its own batch with fused kernels.
2. Gradients are all-reduced over NVLink or Infinity Fabric, overlapped with the remaining backward kernels.
3. Each GPU applies the optimizer update to its weights (or shard); the next step starts. The step is fabric-bound when the all-reduce does not hide behind the backward pass.

Scaling

Limits in the order they usually bite, with the lever for each:

1. Memory bandwidth. Most kernels outside large matrix products are HBM-bound. Levers: reuse in shared memory and registers (Tiled matrix multiplication), fuse operators (Kernel fusion), shrink the data (Weight quantization).
2. Matrix-unit utilisation. Peak needs large tiles, asynchronous copies and the right format. Levers: better kernels, then a smaller format.
3. Occupancy and instruction-level parallelism. Levers: block size, register and shared-memory budgets.
4. Memory capacity. Weights, activations and caches must fit in HBM. Levers: more GPUs, Gradient checkpointing, a larger-HBM part.
5. Interconnect. Beyond one GPU, link bandwidth decides efficiency. Levers: larger NVLink domains, overlapping communication, a different parallelism split.
6. Power. Rack power caps full-clock GPUs; performance per watt comes from lower precision and reuse. Compute grows faster than bandwidth, so kernels drift toward memory-bound; vendors answer with bigger caches, more HBM stacks and smaller formats.

Failure modes

Real implementations

The alternative: a CPU core wins on branchy, latency-bound work; a custom accelerator (ai-accelerator (no page yet), the closest system in the graph) fixes the dataflow in hardware and trades flexibility for efficiency. The GPU is the programmable middle.

Sources

• NVIDIA, NVIDIA Hopper Architecture In-Depth, technical blog, 2022: SM count, tensor cores, TMA, clusters, HBM3, NVLink, power. post
• NVIDIA, Hopper Tuning Guide and Blackwell Tuning Guide, CUDA Toolkit docs: occupancy limits, shared memory, L2, NVLink. Hopper, Blackwell
• NVIDIA, Inside NVIDIA Blackwell Ultra, technical blog, 2025: two-die design, tensor memory, NVFP4, NVLink 5. post
• NVIDIA, CUDA C++ Best Practices Guide: coalescing, banks, occupancy, local memory. guide
• NVIDIA, CUDA Programming Guide, programming model: warps, SIMT, clusters. guide
• NVIDIA, PTX ISA: warpgroup MMA, fifth-generation TensorCore, tensor memory. PTX ISA
• NVIDIA, NVIDIA Turing Architecture In-Depth, technical blog: RT cores and BVH traversal. post
• AMD, Introducing AMD CDNA 3 Architecture, white paper: chiplets, caches, HBM, partitioning, formats. white paper
• AMD, HIP documentation: Hardware implementation: wavefronts, SIMD units, scheduling. docs
• J. Shah et al., FlashAttention-3, 2024. paper
• A. Bakhoda et al., Analyzing CUDA workloads using a detailed GPU simulator, ISPASS 2009. DOI
• M. Khairy et al., Accel-Sim, ISCA 2020. DOI
• B. Tine et al., Vortex: Extending the RISC-V ISA for GPGPU and 3D-Graphics, MICRO 2021. DOI