Pidoku

Blackwell: B200, B300 and the Superchip

Module 7 Lesson 7.c 55 min Difficulty 4/5

Prerequisites 02, 2.d, 5.b

The idea in one minute#

Hopper’s die was already as large as a lithography machine can print. To grow, Blackwell put two reticle-sized dies in one package and joined them with a link so fast (10 TB/s) that software sees one GPU. Around them sit eight stacks of HBM3e: 192 GB at about 8 TB/s, more than twice the H100’s bandwidth.

Inside each SM, the fifth-generation tensor cores add 4-bit floating point with block scaling (NVFP4), which doubles math again over FP8 for models that tolerate it. And the GB200 superchip pairs two of these GPUs with a Grace CPU on one board, so a GPU can reach the CPU’s memory at 900 GB/s instead of 64 GB/s over PCIe.

B300 (Blackwell Ultra) is the mid-cycle refresh: 288 GB, more FP4 math, the same bandwidth. Facts checked on 16 October 2026.

A picture#

NVIDIA B200 · TWO DIES, ONE GPUBlackwell · 2024-25 · schematic, not to scale · checked 16 Oct 2026CoWoS-L interposer (silicon bridges in an organic layer)HBM3e24 GB 8-highHBM3e24 GB 8-highHBM3e24 GB 8-highHBM3e24 GB 8-highGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSML2 cacheHBM3e24 GB 8-highHBM3e24 GB 8-highHBM3e24 GB 8-highHBM3e24 GB 8-highGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMGPCSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSMSML2 cache NVHBI123456WHAT YOU ARE LOOKING AT1Two reticle-sized diesEach die is as large as a lithography reticleallows: 104 B transistors on TSMC 4NP. Two of themmake one 208 B-transistor GPU.2NV-HBI, 10 TB/sThe die-to-die link is so fast that software seesone GPU with one memory space. No code change fortwo dies.3HBM3e × 8Four stacks per die. 192 GB physical; 180 GB (DGXB200) or 186 GB (GB200) exposed. ~8 TB/s.4L2 per dieEach die keeps its own L2 near its own HBM; thelink stitches the halves together.5NVLink 518 links, 1.8 TB/s per GPU, twice Hopper. Enough toput 72 GPUs in one NVLink domain.65th-gen tensor coresAdd FP4/FP6 and NVIDIA's NVFP4 scaling format, plusper-SM tensor memory (TMEM) that holdsaccumulators.SIDE VIEW · SECTION THROUGH THE PACKAGE2 logic diesHBM: 8-16 DRAM dies, joined by TSVsHBM base (logic) dieGPU logic die, flip-chip on microbumpsSilicon interposer: short, dense wiresOrganic substrate; bumps to the moduleTRANSISTORS208 B2 dies, TSMC 4NPFP4 DENSE9 PF18 PF with sparsityFP8 DENSE4.5 PFBF16 2.25 PFMEMORY192 GBHBM3e · ~8 TB/sNVLINK 51.8 TB/s18 linksPOWER1,000 WHGX; 1,200 W in GB200Sources: NVIDIA Blackwell architecture technical brief; DGX B200 and GB200 NVL72 datasheets. B300 (Blackwell Ultra): 288 GB HBM3e 12-high, ~15 PF dense FP4, same 8 TB/s.
B200 from above. The orange bar between the dies is the NV-HBI die-to-die link; each die has its own L2 and its own four HBM stacks.
GB200 NVL72 COMPUTE TRAY · TOP VIEW1U, liquid cooled · 2 superchips = 2 Grace + 4 B200 · checked 16 Oct 2026FRONT (cold aisle)REAR → NVLink spineConnectXSuperNIC · 400-800 Gb/sConnectXSuperNIC · 400-800 Gb/sConnectXSuperNIC · 400-800 Gb/sConnectXSuperNIC · 400-800 Gb/sBlueField-3DPU · front-endE1.S NVMe × 8GB200 SUPERCHIP 1Grace CPU72 Arm Neoverse V2coresLPDDR5X · 480 GB · ~512 GB/sB200186GBB200186GBNVLink-C2C900 GB/sGB200 SUPERCHIP 2Grace CPU72 Arm Neoverse V2coresLPDDR5X · 480 GB · ~512 GB/sB200186GBB200186GBNVLink-C2C900 GB/sNVLink 518 links / GPUto spine cartridgeNVLink 518 links / GPUto spine cartridgeNVLink 518 links / GPUto spine cartridgeNVLink 518 links / GPUto spine cartridgeliquid loop: cold plates on every CPU and GPU · supply / return to the rack manifold1234561Network cards at the frontOne NIC per GPU for the scale-out (back-end) network; a DPU for storage andusers.2Grace CPURuns the OS, feeds data and holds spill-over memory. Its LPDDR5X isreachable by both GPUs.3B200 GPUTwo dies, 8 HBM3e stacks; in GB200 each exposes 186 GB.4NVLink-C2CA coherent chip-to-chip link: GPUs read CPU memory at up to 900 GB/s, 7x PCIeGen 5.5NVLink to the spineEvery GPU's 18 links leave the back of the tray through blind-mate connectors tocopper cables.6Liquid coolingAir cannot remove ~5.4 kW from 1U. Cold plates carry heat to the rack's coolantloop.
One 1U compute tray: two superchips, each one Grace CPU and two B200 GPUs. Eighteen of these trays fill an NVL72 rack (lesson 04).

How it really works#

Two dies, one GPU#

Why not simply build a bigger die? Because a scanner exposes at most about 26 × 33 mm per shot, and yield falls off steeply with area. Blackwell’s answer is the same one AMD took with chiplets, but with a much fatter connection:

  • Each die has 104 billion transistors on TSMC 4NP; the package has 208 billion.
  • The dies sit on CoWoS-L: an organic interposer with small silicon bridges embedded only where dense wiring is needed. That is cheaper than one huge silicon interposer and allows a larger package.
  • NV-HBI carries 10 TB/s between the dies. That is more than the 8 TB/s of HBM bandwidth, so a kernel on die 0 reading memory attached to die 1 is not starved. Caches stay coherent and the driver presents one device with one address space. No code changes.
  • Each die keeps its own L2 and four HBM stacks. Data placement still has a small effect: memory on the far die costs a hop.

The fifth-generation tensor core#

Three changes matter to anyone writing or choosing kernels:

ChangeHopperBlackwell
Who issues the MMAa warpgroup (128 threads, wgmma)one thread (tcgen05.mma)
Where accumulators liveregisterstensor memory (TMEM): 256 KB per SM next to the tensor core
Smallest formatFP8FP4 and FP6, with hardware block scaling

TMEM frees the register file for the rest of the kernel, and a pair of SMs can cooperate on one larger tile. Hand-written Hopper kernels therefore need rewriting to reach Blackwell’s peak; the libraries (cuBLAS, CUTLASS, FlashAttention-4 style kernels) already have been.

FP4 and block scaling#

A 4-bit float (E2M1) has only 15 distinct values: ±0, ±0.5, ±1, ±1.5, ±2, ±3, ±4, ±6. To represent real weights, values are grouped into small blocks, and each block stores a scale:

FormatBlock sizeScale formatNotes
MXFP4 (OCP standard)32 values8-bit power of two (E8M0)Portable: AMD, NVIDIA and others support it
NVFP4 (NVIDIA)16 valuesFP8 E4M3, plus one FP32 scale per tensorSmaller blocks and a finer scale: less error

Smaller blocks mean an outlier only spoils 15 neighbours, not 31; a non-power-of-two scale lands closer to each block’s true range. The code below measures the difference.

B300 (Blackwell Ultra)#

  • 288 GB of HBM3e in 12-high stacks (vs 8-high on B200), at the same ~8 TB/s.
  • About 1.5x the dense FP4 math (roughly 14-15 PF, depending on the power limit), and faster exponentials for attention softmax.
  • FP64 and INT8 throughput sharply cut: silicon moved from scientific computing to AI.
  • Up to 1,400 W in GB300.

As 8.a said: a capacity upgrade, not a decode-speed upgrade.

The superchip and the tray#

A GB200 superchip is one Grace CPU and two B200 GPUs on one board:

  • Grace: 72 Arm Neoverse V2 cores and up to 480 GB of LPDDR5X.
  • NVLink-C2C: a coherent chip-to-chip link at 900 GB/s, about 7x PCIe Gen 5 x16 in each direction. The GPUs can read and write Grace’s memory directly, which makes CPU memory a usable overflow tier for KV cache, embeddings or optimizer state (see where AI data lives).
  • In GB200 each GPU exposes 186 GB of its HBM3e.

Two superchips, four NICs (one per GPU, for the scale-out network), a BlueField DPU, local NVMe and cold plates make one 1U, liquid-cooled compute tray. Every GPU’s 18 NVLink links leave the rear of the tray through blind-mate connectors to the rack’s copper spine.

Code#

Quantize a heavy-tailed weight vector to FP4 three ways: one scale for the whole tensor, MXFP4-style blocks of 32 with power-of-two scales, and NVFP4-style blocks of 16 with FP8 scales.

Go
// fp4.go — why block scaling makes 4-bit weights usable.
package main

import (
	"fmt"
	"math"
	"math/rand"
)

// The 8 non-negative values an E2M1 (FP4) number can take.
var fp4 = []float64{0, 0.5, 1, 1.5, 2, 3, 4, 6}

func roundFP4(x float64) float64 {
	best := 0.0
	for _, v := range fp4 {
		if math.Abs(math.Abs(x)-v) < math.Abs(math.Abs(x)-best) {
			best = v
		}
	}
	return math.Copysign(best, x)
}

// roundE4M3 keeps 3 mantissa bits, like an FP8 E4M3 scale.
func roundE4M3(x float64) float64 {
	m, e := math.Frexp(x) // x = m * 2^e, m in [0.5, 1)
	return math.Ldexp(math.Round(m*16)/16, e)
}

// quantize scales each block so its largest value maps to 6, the largest FP4 value.
func quantize(w []float64, block int, scale func(float64) float64) []float64 {
	out := make([]float64, len(w))
	for start := 0; start < len(w); start += block {
		end := min(start+block, len(w))
		amax := 0.0
		for _, v := range w[start:end] {
			amax = math.Max(amax, math.Abs(v))
		}
		s := scale(amax / 6)
		for i := start; i < end; i++ {
			out[i] = roundFP4(w[i]/s) * s
		}
	}
	return out
}

func rmse(a, b []float64) float64 {
	sum := 0.0
	for i := range a {
		sum += (a[i] - b[i]) * (a[i] - b[i])
	}
	return math.Sqrt(sum / float64(len(a)))
}

func main() {
	r := rand.New(rand.NewSource(7))
	w := make([]float64, 1<<16)
	for i := range w {
		w[i] = r.NormFloat64() * 0.02
		if r.Float64() < 0.002 { // rare outliers, as real weight matrices have
			w[i] *= 25
		}
	}
	exact := func(s float64) float64 { return s }
	pow2 := func(s float64) float64 { return math.Pow(2, math.Ceil(math.Log2(s))) }

	perTensor := quantize(w, len(w), exact)
	mx := quantize(w, 32, pow2)
	nv := quantize(w, 16, roundE4M3)

	base := rmse(w, perTensor)
	fmt.Printf("%-28s %10s %10s\n", "scheme", "RMSE", "vs tensor")
	for _, c := range []struct {
		n string
		q []float64
	}{{"one scale per tensor", perTensor}, {"MXFP4: 32, power-of-two", mx}, {"NVFP4: 16, FP8 scale", nv}} {
		e := rmse(w, c.q)
		fmt.Printf("%-28s %10.5f %9.2fx\n", c.n, e, e/base)
	}
	bits := func(block int, scaleBits float64) float64 { return 4 + scaleBits/float64(block) }
	fmt.Printf("\nstorage: MXFP4 %.2f bits/weight, NVFP4 %.2f bits/weight, FP8 8, BF16 16\n",
		bits(32, 8), bits(16, 8))
}

One scale for the whole tensor is useless: the outliers set the scale and most weights round to zero. Block scaling cuts the error several-fold, and NVFP4’s smaller blocks and finer scales cut it further, at a cost of 4.5 bits per weight instead of 4.25. Real FP4 deployment adds calibration and often keeps sensitive layers in FP8, but this is the core of why Blackwell’s FP4 is usable at all.

Remember this#

  • Blackwell = two reticle-sized dies, joined at 10 TB/s, seen by software as one GPU.
  • 192 GB HBM3e at ~8 TB/s on B200; B300 adds capacity (288 GB) and FP4 math, not bandwidth.
  • The fifth-gen tensor core is issued by one thread and accumulates into TMEM; old kernels need rewriting to reach peak.
  • FP4 works because of block scaling; NVFP4 uses 16-value blocks and FP8 scales, MXFP4 uses 32-value blocks and power-of-two scales.
  • GB200 = Grace + 2 B200 joined by NVLink-C2C at 900 GB/s, making CPU memory a fast overflow tier.

Try it#

  1. Run fp4.go. Change the outlier rate from 0.2% to 2%. Which scheme degrades most?
  2. Try block sizes 8 and 64 with the E4M3 scale. Plot error against bits per weight.
  3. A 405B model at NVFP4: how many GB of weights (use 4.5 bits)? Does it fit one GB200 superchip’s two GPUs?

Check yourself#

  1. Why did NVIDIA go to two dies instead of one bigger die?
  2. What does the NV-HBI bandwidth have to exceed for the two dies to behave as one GPU?
  3. What is TMEM, and what does it free up?
  4. Name two differences between MXFP4 and NVFP4.
  5. What can a GB200 GPU do with Grace’s memory that an HGX B200 GPU cannot do as fast?

Sources#

↑↓ navigate↵ openesc close

100% drag to pan · scroll or pinch to zoom · double-click to zoom in