Hopper: H100 and H200

The idea in one minute#

The H100 is the best-documented AI chip ever made and still the largest installed base, so it is the one to know by heart. It is one large die (GH100) in the middle of a package, surrounded by stacks of HBM, sitting on a silicon interposer. The die is a grid of 8 GPCs, each holding 18 SMs; 132 of the 144 SMs are switched on. Each SM is split into 4 partitions, each with its own scheduler, registers, 32 FP32 lanes and one tensor core.

Learn this one floorplan and every later NVIDIA GPU reads as a variation of it: more SMs, a second die, a different memory, a faster link. Facts checked on 16 October 2026.

A picture#

How it really works#

From package to lane#

Read the first drawing from the outside in.

1. The module. An SXM5 board carries the package, power delivery for 700 W and the NVLink connectors to the HGX baseboard.
2. The package. A CoWoS-S silicon interposer holds the die and the HBM stacks side by side. The interposer exists because a 5,120-bit memory bus needs thousands of wires a few millimetres long; no circuit board can route that.
3. The die. 814 mm² and 80 billion transistors on TSMC 4N, close to the largest die a lithography machine can expose in one shot (the reticle limit, about 830 mm²).
4. GPCs. Eight graphics processing clusters. A GPC groups 9 TPCs, each of two SMs. The GigaThread engine hands thread blocks to GPCs.
5. SMs. 144 on the full die; 132 on an H100 SXM5 (114 on the PCIe card). A die that large almost always has defects; designing in spare SMs and fusing off the bad ones turns a defective die into a sellable one. This is why every H100 has a slightly different set of working SMs.
6. Partitions and lanes. Inside each SM, four schedulers each issue one warp instruction per clock to their own 32 FP32 lanes, 16 INT32, 16 FP64, 8 load/store units, 4 special-function units and one tensor core (second drawing).

What Hopper added#

Hopper is the generation where the GPU was redesigned around transformers:

FlashAttention-3 and the fast FP8 GEMMs in cuBLAS and CUTLASS are written around wgmma and TMA. Code written for A100-style mma.sync runs on Hopper, but at a fraction of peak.

H100 variants and H200#

The H200 is the same GH100 die. The only change is the memory: all six HBM sites filled, with faster HBM3e. That alone makes it about 1.4x faster at memory-bound decoding and lets a 70B model at FP8 fit on one GPU with room for a KV cache. It is the cleanest real-world proof that, for LLM inference, memory is the product.

The numbers, derived#

NVIDIA’s 989 TFLOP/s of dense BF16 is not magic; it is a product of four counts:

132 SMs × 4 tensor cores × 512 FMA per clock per tensor core × 2 FLOP per FMA × 1.83 GHz
= 989 TFLOP/s

FP8 doubles the FMAs per clock: 1,979 TFLOP/s. The “with sparsity” figures double again, but only for weights pruned to a 2-of-4 pattern, which almost no production LLM uses. Quote dense numbers.

Divide peak math by bandwidth and you get the machine balance: the arithmetic intensity a kernel needs to be compute-bound (4.a).

Code#

// hopper.go — derive H100's peak numbers and the arithmetic intensity it needs.
package main

import "fmt"

type GPU struct {
	Name                 string
	SMs                  int
	TensorFMAPerClkPerSM float64 // dense BF16 FMAs per clock per SM
	ClockGHz             float64 // clock NVIDIA's tensor figures assume
	BandwidthTBs         float64
}

func main() {
	gpus := []GPU{
		{"H100 PCIe", 114, 2048, 1.62, 2.0},
		{"H100 SXM5", 132, 2048, 1.83, 3.35},
		{"H200", 132, 2048, 1.83, 4.8},
	}
	fmt.Println("GPU          BF16 TF   FP8 TF   balance BF16   balance FP8  (FLOP per byte)")
	for _, g := range gpus {
		bf16 := float64(g.SMs) * g.TensorFMAPerClkPerSM * 2 * g.ClockGHz / 1000 // TFLOP/s
		fp8 := 2 * bf16
		fmt.Printf("%-11s %8.0f %8.0f %14.0f %13.0f\n",
			g.Name, bf16, fp8, bf16/g.BandwidthTBs, fp8/g.BandwidthTBs)
	}

	// Decode at batch 1 multiplies a weight matrix by one vector: 2 FLOPs per weight,
	// and each FP8 weight is 1 byte. Intensity 2 FLOP/byte, far below the balance.
	fmt.Println("\nbatch  intensity (FLOP/byte, FP8 weights)  regime on H100 SXM5")
	for _, b := range []int{1, 8, 64, 256, 512, 1024} {
		ai := 2.0 * float64(b)
		regime := "memory-bound"
		if ai >= 1979/3.35 {
			regime = "compute-bound"
		}
		fmt.Printf("%5d  %35.0f  %s\n", b, ai, regime)
	}
}

The first table reproduces NVIDIA’s 989 and 1,979 TFLOP/s for the SXM5 (and the PCIe card’s 756, which runs at a lower clock). The second shows why decoding needs batching: one user’s token uses the H100 at well under 1% of its FP8 peak, and only around 300 tokens decoded together does the work become compute-bound.

Remember this#

• H100 = one 814 mm² die, 8 GPCs × 18 SMs, 132 enabled, 4 partitions per SM, one tensor core per partition, 5 HBM3 stacks on an interposer.
• Spare SMs are fused off for yield; that is normal, not a defect.
• Hopper’s real innovations were FP8, wgmma, TMA and clusters: asynchronous, tile-level hardware for transformers.
• H200 is H100 with all six HBM sites filled with HBM3e: same compute, 1.4x the bandwidth, 1.8x the memory.

Try it#

1. Run hopper.go. Change the H100 SXM5 bandwidth to the H200’s 4.8 TB/s and find the new batch at which decode turns compute-bound.
2. Using the drawing, estimate how many SMs a 7-way MIG split gives each instance (6.c).
3. On a machine with an H100, run nvidia-smi -q | grep -i -A2 clocks under load and compare the clock to the 1.83 GHz used above.

Check yourself#

1. Why does the H100 package need a silicon interposer?
2. What changed between H100 and H200, and what did not?
3. Derive 989 TFLOP/s from the SM count and the clock.
4. Why is a single user’s decode memory-bound on any H100?

Sources#

• NVIDIA H100 Tensor Core GPU Architecture whitepaper
• NVIDIA Hopper architecture in depth
• NVIDIA H200 product page
• FlashAttention-3 (Shah et al., 2024) — how wgmma and TMA are used