Hopper — H100, FP8, and the Transformer Engine

Topics We'll Cover

A focused tour of NVIDIA's Hopper generation — the architecture that made FP8 a first-class citizen and turned the SM into something that looks more like a tile-streaming engine than a stack of warp lanes.

• Hopper in One Page
• GH100 — The Die
• The Hopper SM
• 4th-Gen Tensor Cores + FP8
• Transformer Engine — Software Magic for FP8
• TMA — Tensor Memory Accelerator
• Thread-Block Clusters + DSMEM
• DPX — Dynamic Programming Acceleration
• H100 Variants — SXM5, PCIe, NVL, H200, GH200
• Performance — Why H100 Changed Everything
• Interactive: Which Hopper for What?

Hopper in One Page

Hopper, named for Grace Hopper, is the first NVIDIA GPU with native FP8 tensor cores. That single change — together with the Transformer Engine that made FP8 usable without manually rescaling every layer — redefined what "frontier" meant. GPT-4-class models were trained on Hopper.

GH100 — The Die

GH100 is the silicon under every H100 / H200 / GH200 SKU. Two perspectives below: the compute fabric and the memory subsystem. Yields and SKU binning mean none of the shipping products use the full die — the H100 enables 132 of 144 SMs, the H200 keeps the same compute config but swaps to denser memory.

The Hopper SM

The Hopper SM keeps the Ampere-style four-partition layout but adds a pile of new fixed-function units. The headline isn't another increment in CUDA cores — it's that almost everything around the cores has been redesigned for asynchronous, tile-based execution.

What's actually new vs Ampere

• Tensor core: 4th gen, native FP8 + WGMMA (warp-group MMA on 64×N×K tiles, larger than any prior MMA).
• TMA: hardware unit that copies multi-D tiles between L2/HBM and shared memory asynchronously, freeing threads from issuing global loads.
• Thread-block clusters: a new launch-hierarchy level above the block; up to 16 blocks co-scheduled on one GPC.
• DSMEM: blocks within a cluster can access each other's shared memory directly at L1-class latency.
• DPX: hardware min/max+add instructions for dynamic-programming patterns (Smith-Waterman, beam search).
• Async transaction barriers: new mbarrier states that decouple "data arrived" from "all threads ready", enabling deeper overlap.
• 256 KB combined L1+shared, configurable carve-out (vs 192 KB on Ampere). 64 KB register file per partition (unchanged).

4th-Gen Tensor Cores + FP8

The headline numerical change in Hopper is native FP8. Two 8-bit float formats are supported, chosen so that one favours mantissa precision and the other favours dynamic range. They share an exponent bias and convert losslessly to/from FP16/BF16.

Throughput on H100 SXM5

WGMMA — warp-group MMA

The PTX instruction is wgmma.mma_async.sync.aligned.m64nNkK. A warp group (4 warps = 128 threads) collectively issues a single MMA over a 64×N×K tile, with N up to 256. This is far larger than Ampere's m16n8k16 MMA — one instruction now drives a much larger fraction of the matmul. The async suffix is the key: the tensor core proceeds in the background while the warp group does other work (e.g. issues the next TMA or accumulates the previous tile).

// Ampere FP16: one warp issues a 16×8×16 tile, blocking-ish
mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32  d, a, b, c;

// Hopper FP8: a warp-group (4 warps) issues a 64×128×32 tile, async
wgmma.mma_async.sync.aligned.m64n128k32.f32.e4m3.e4m3  d, a-desc, b-desc, scaleD;
wgmma.commit_group.sync;
wgmma.wait_group.sync  0;     // drain when needed

Transformer Engine — Software Magic for FP8

FP8 has only ~256 representable values. Naively casting a transformer's activations to FP8 destroys accuracy unless every tensor is scaled into FP8's narrow dynamic range, then unscaled back. The Transformer Engine (TE) is the open-source library that does this for you.

import transformer_engine.pytorch as te
from transformer_engine.common import recipe

fp8_recipe = recipe.DelayedScaling(
    margin=0, interval=1,
    fp8_format=recipe.Format.HYBRID,    # E4M3 fwd, E5M2 bwd
    amax_history_len=16, amax_compute_algo="max",
)

# Replace nn.Linear with te.Linear — same API, FP8 internals
class Block(nn.Module):
    def __init__(self, d):
        super().__init__()
        self.fc1 = te.Linear(d, 4*d, bias=True)
        self.fc2 = te.Linear(4*d, d, bias=True)

# Wrap the forward; everything between is FP8-accelerated
with te.fp8_autocast(enabled=True, fp8_recipe=fp8_recipe):
    out = block(x)
    loss = out.sum()
loss.backward()                       # grads computed in E5M2 path

Delayed scaling — the trick that hides the rescale stall

"Online" scaling would have to read the activation, compute its max, broadcast that max, then scale — a multi-pass dependency on every tensor. Delayed scaling instead remembers a short rolling history of recent amax values per tensor (default 16 steps), takes the running max, and uses that to scale the next forward pass. The cost is one stale step at start-up; the win is no synchronisation barrier in the hot path.

TMA — Tensor Memory Accelerator

Before Hopper, every thread that participated in a tile-based matmul also had to participate in loading that tile from global memory: address arithmetic, predication, vectorised LDG, register pressure. That overhead competed with the tensor cores for issue slots and consumed registers and ALU bandwidth. TMA removes it.

What it is

• A dedicated copy engine inside each SM that moves multi-dimensional tiles between L2/HBM and shared memory based on a tensor map descriptor.
• The descriptor is a small structure encoding base pointer, tile shape (up to 5D), strides, swizzle pattern, OOB handling, and L2 hint.
• An SM submits one instruction (cp.async.bulk.tensor) and the TMA does the rest asynchronously.
• Completion is signalled via an mbarrier with transaction-count semantics — the consuming threads wait until all the tile bytes have arrived.

// Build a tensor map once on the host
CUtensorMap tmap_a;
cuTensorMapEncodeTiled(&tmap_a, CU_TENSOR_MAP_DATA_TYPE_FLOAT8_E4M3,
    2, dev_a, dim, stride, box, elem_stride,
    CU_TENSOR_MAP_INTERLEAVE_NONE, CU_TENSOR_MAP_SWIZZLE_128B,
    CU_TENSOR_MAP_L2_PROMOTION_L2_128B, CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);

// In the kernel: one warp issues the TMA load
__shared__ alignas(128) __nv_fp8_e4m3 sA[128][64];
__shared__ uint64_t bar;
if (threadIdx.x == 0) {
    cde::cp_async_bulk_tensor_2d_global_to_shared(&sA, &tmap_a, ki*64, mi*128, &bar);
    cde::cp_async_bulk_commit_group();
}
cde::cp_async_bulk_wait_group<0>();   // or wait on bar with mbarrier
// sA is now populated; warp group can WGMMA on it.

Thread-Block Clusters + DSMEM

Hopper adds a new level to the launch hierarchy. Before Hopper, the levels were Grid → Block → Warp → Thread, with the SM as the natural hardware boundary: blocks ran on one SM, no SM-to-SM cooperation existed except via global memory. Hopper inserts a Cluster level above the block.

Co-scheduling & DSMEM

• A cluster contains up to 16 blocks (8 portable, 16 with opt-in).
• All blocks of a cluster are guaranteed to be co-resident on a single GPC — the GPC's local crossbar is the substrate for cross-block traffic.
• Cluster blocks can read and write each other's shared memory directly using DSMEM (Distributed Shared Memory).
• DSMEM access goes through the GPC's switch fabric — latency comparable to L1, much lower than going to L2 or HBM.
• Cluster-wide barriers (cluster.sync()) and atomic ops are provided.

// Compile-time cluster shape, opt-in launch attribute
__cluster_dims__(2, 2, 1)        // 4 blocks per cluster
__global__ void matmul_cluster(const half* A, const half* B, float* C) {
    namespace cg = cooperative_groups;
    auto cluster = cg::this_cluster();
    __shared__ half tile[128][64];

    // Each block loads its own slice via TMA into its own SMEM
    tma_load_my_slice(tile);

    cluster.sync();                        // all blocks have their slice

    // Read peer block's tile via DSMEM — one pointer cast away
    half* peer_tile = cluster.map_shared_rank(tile, 1);
    // peer_tile lives on another SM; load latency is L1-ish, no HBM hop.
}

Practical use — collaborative matmul

The CUTLASS 3.x and cuDNN 9 GEMM kernels split a large output tile across cluster blocks. Each block holds a slice of the working tile and exchanges tile edges with its cluster neighbours via DSMEM as the K-loop progresses. Because peer SMEM reads stay inside the GPC, this avoids the L2 round-trip an Ampere kernel would need for the same data. Net effect: more on-chip bandwidth available to the tensor cores, larger effective tile sizes, better arithmetic intensity.

DPX — Dynamic Programming Acceleration

The smallest of Hopper's headline features but the most domain-specific. DPX adds a family of fused min/max + add instructions that cover the inner loop of dynamic-programming algorithms. The canonical examples are not LLM workloads at all — they are bioinformatics and graph algorithms.

What it actually wins you

• NVIDIA quotes up to 7× speedup for Smith-Waterman gap-affine alignment vs an A100 implementation; ~5× for Floyd-Warshall on dense graphs.
• The exposed API lives in <cuda/std/__bit/dpx> (CUDA 12+) and in cuOpt / NVIDIA Parabricks / NCCL collective gather paths.
• For LLM training and inference, DPX is essentially unused — it's a marketing-relevant feature for the biotech and HPC verticals NVIDIA also wanted to capture.
• Beam search in classical decoding does benefit; for transformer decoding with KV cache it's irrelevant.

// Gap-affine recurrence, one diagonal cell, traditional form:
//   H[i,j] = max( H[i-1,j-1] + s(a,b),  E[i,j],  F[i,j],  0 )
// E[i,j] = max( H[i,j-1] - g_open,  E[i,j-1] - g_extend )
// F[i,j] = max( H[i-1,j] - g_open,  F[i-1,j] - g_extend )
int h = __viaddmax_s32(H_diag, score_match, 0);    // (H_diag + s) ↑ 0
h     = __vimax3_s32(h, E_left, F_up);                       // 3-way max in 1 inst
int e = __viaddmax_s32(H_left, -g_open, E_left - g_extend);
int f = __viaddmax_s32(H_up,   -g_open, F_up   - g_extend);
H[i][j] = h; E[i][j] = e; F[i][j] = f;

H100 Variants — SXM5, PCIe, NVL, H200, GH200

"Hopper" on its own doesn't tell you what's actually in the box. The same GH100 silicon ships in five very different products with widely different memory, power, and link configurations. Pick wrong and you discover at deploy time that NVLink isn't there, or that bandwidth is half what you thought.

DGX H100 — the reference node

Picking the right Hopper

• Training a frontier model: SXM5 in HGX/DGX, period. PCIe variants don't have the NVLink fabric for tensor-parallel scaling.
• Inference of a 70B-class model: H100 NVL or H200; the extra memory matters more than peak compute.
• Mixed enterprise server, occasional ML: H100 PCIe; don't pretend you'll do TP.
• Massive context, KV-cache-bound serving: H200, then GH200 if context truly exceeds HBM.
• Sparse MoE or recommender models with huge embeddings: GH200 — the 480 GB of LPDDR5x as EGM is the right tool.

Performance — Why H100 Changed Everything

The headline NVIDIA pushed at GTC 2022 was 9× training throughput vs A100 on large transformer workloads. The number is real but it's a stack of independent wins; understanding which contributes how much tells you what you actually buy when you upgrade.

Where the 9× comes from

The four wins compound because they are independent: doubling FP-throughput while halving memory reads while reducing kernel-launch overhead while increasing tile size all multiply rather than add. None of them on its own buys the headline; together they do.

For LLM serving

• FP8 weights halve the per-token bytes read from HBM → decode tok/s roughly doubles on bandwidth-bound shapes.
• FP8 KV cache halves the per-token KV bytes too → longer context fits, batch sizes grow.
• WGMMA + TMA make prefill compute-bound rather than overhead-bound → prefill throughput jumps 3–4× over A100 on the same shapes.
• Decode is bandwidth-bound either way — the H100 vs A100 gap there is closer to bandwidth ratio (3.35 / 2.0 = 1.7×) than to FLOPS ratio.

Worked example — Llama-3-70B in vLLM

Interactive: Which Hopper for What?

Pick a Hopper SKU. Get the headline numbers and a one-line use-case recommendation.