Inside Hopper — FP8, TMA, Thread-Block Clusters, NVLink 4

Topics We'll Cover

• Hopper at a Glance
• GH100 Die — Block Diagram
• The Hopper SM — Cluster-Aware
• 4th-Generation Tensor Cores — FP8 Native
• WGMMA — Warp-Group Matrix Multiply
• TMA — The Tensor Memory Accelerator
• Thread-Block Clusters & DSMEM
• DPX — Dynamic Programming Acceleration
• Process, Voltage, Form Factors
• Memory — HBM3, HBM3e, GH200 LPDDR5x
• NVLink 4 & NVSwitch 3 (HGX H100)
• Interactive: Hopper SKU Picker

Hopper at a Glance

Hopper launched in March 2022 (announced) with volume in 2023. Single die family, datacenter only — like Volta, no consumer Hopper. The architecture's job is high-end LLM training and inference. Variants:

GH100 Die — Block Diagram

Density: 80 B / 814 mm² = 98.3 M/mm² — less dense than AD102 (125.5) on the same node because GH100 spends area on FP64, full TC throughput, and HBM3 PHYs.

The Hopper SM — Cluster-Aware

4th-Generation Tensor Cores — FP8 Native

Hopper's tensor cores add FP8 in two formats: E4M3 (1 sign + 4 exp + 3 mantissa, favours mantissa, used for activations and weights after scaling) and E5M2 (1+5+2, favours range, used for gradients). Throughput is 2× BF16.

Software: the Transformer Engine library (open source, PyTorch + JAX bindings) handles per-tensor FP8 scaling automatically. The SM tracks recent activation maxima and applies appropriate scales on the next forward pass — "delayed scaling".

WGMMA — Warp-Group Matrix Multiply

Hopper introduces a new MMA shape: warp-group MMA, operating on 64×N×K tiles where N can be up to 256, K up to 32. Four warps cooperate; the instruction is asynchronous — it returns immediately and signals completion via mbarrier.

wgmma.mma_async.sync.aligned.m64n128k16.f32.bf16.bf16
    {d0, d1, ..., d63},     // 64 FP32 accumulator regs per warp
    desc-a,                          // shared memory descriptor for A tile
    desc-b,                          // shared memory descriptor for B tile
    {scale-d, imm-trans-a, imm-trans-b};
wgmma.commit_group.sync;
wgmma.wait_group.sync 0;

Why this matters: pre-Hopper MMA shapes were small (16×8×16) so each thread had to issue many instructions to fill an output tile. WGMMA does it in one instruction over four warps, freeing the warp schedulers for other work. Used by every tuned cuBLAS / cuDNN / Flash-Attn 3 / CUTLASS 3.x kernel on Hopper.

TMA — The Tensor Memory Accelerator

One TMA unit per SM. The kernel sets up a tensor-map descriptor (base pointer, multi-dim shape, strides, swizzle pattern, OOB handling) and submits a single cp.async.bulk.tensor PTX instruction. The TMA hardware transfers a tile from L2/HBM to shared memory while the SM continues executing other warps.

__shared__ __align__(16) half tile[128][64];
__shared__ __alignas__(8) cuda::barrier<cuda::thread_scope_block> bar;

if (threadIdx.x == 0) cuda::memcpy_async(
    tile, gmem_ptr_with_descriptor,
    cuda::aligned_size_t<16>{128*64*2}, bar);
__syncthreads();
bar.arrive_and_wait();   // blocks here until TMA done

TMA is not just a memcpy — it understands the layout the WGMMA wants. Without TMA, every thread issued global loads, consuming registers as staging buffers. With TMA, kernels become almost-pure compute on shared-memory tiles.

Thread-Block Clusters & DSMEM

New level of the CUDA hierarchy: Grid → Cluster → Block → Warp → Thread. A cluster is up to 16 cooperating thread blocks, all scheduled on the same GPC. They share a cluster-wide barrier and can directly read each other's shared memory: Distributed Shared Memory (DSMEM).

Implementation detail: clusters share a 256 KB L1+shared budget per SM, but the cluster as a whole sees 16×256 KB = up to 4 MB of pooled shared memory at L1 access latency.

DPX — Dynamic Programming Acceleration

Hopper added DPX instructions: hardware-accelerated min/max followed by a fused add, designed for the recurrence patterns in dynamic programming. Smith-Waterman gap-affine alignment in genomics gets a 7× speedup; beam search in NLP and Floyd-Warshall in routing also benefit.

Marginal for LLM training proper, but headline-relevant for biotech / scientific HPC and certain RL tree searches.

Process, Voltage, Form Factors

GH100 is on TSMC 4N — same NVIDIA-customised N4 variant as Ada. 814 mm² is right at the reticle limit; this is a near-impossible-to-yield die. NVIDIA mitigates by binning aggressively: H100 enables 132 of 144 SMs, H800 enables 132 of 144 (export-restricted), H20 enables far fewer at lower clocks for the China market.

Voltage rails: VDD core (~0.85 V at boost), VDDQ-HBM3 (1.1 V), VDDIO-NVLink (~0.75 V). SXM5 carries no aux power connector — the HGX baseboard supplies 700 W via the SXM5 socket itself. PCIe variant uses a CEM 16-pin EPS.

Memory — HBM3, HBM3e, GH200 LPDDR5x

H100 was the first GPU to ship HBM3. H200 then upgraded to HBM3e in mid-2024.

L2 cache: 50 MB on GH100, organised as two 25 MB partitions. ECC SECDED + on-die ECC mandatory.

NVLink 4 & NVSwitch 3 (HGX H100)

NVLink 4 keeps the per-link 25 GB/s/dir of NVLink 3 but doubles links per GPU: 18 links × 50 GB/s/dir = 900 GB/s/dir aggregate. Signalling: 100 Gbps PAM4 per lane, 2 lanes per link.

HGX H100: 8× H100 SXM5 + 4× NVSwitch 3.0 chips (NVSwitch 3 has 64 NVLink 4 ports, ~14 TB/s aggregate per chip). Fully connected at 900 GB/s per pair. NVLink-C2C on GH200 connects Grace CPU to Hopper GPU at 900 GB/s coherent — same protocol family but in-package. Cross-node uses ConnectX-7 NDR 400 Gb/s InfiniBand (8 NICs per HGX H100).

Interactive: Hopper SKU Picker