AI Accelerators — Presentation 01

The Design Space and the Memory Wall

Every AI chip does the same sums. What separates the families is where the weights live, who decides when data moves, and how big the engines are. GPUs and systolic designs as the baseline, the decode memory wall, and every family run on the same decoder model.

Families Weights in HBM vs SRAM Static vs dynamic schedule Roofline Decode Capacity Bounds vs published
Where weights live → What moves → Who decides → Same model, every family
00

Topics We’ll Cover

Live companion and sources

This deck follows chapter 1, the design space, chapter 2, chapter 3 and chapter 10, choosing of AI Accelerators Explained, where every number here is animated. Specifications, sources and access dates: the specifications page.

01

One Question for Every Chip

Every AI chip does the same sums: multiply activations by matrices of weights, layer after layer. What separates the families is what happens around the arithmetic. Ask the same three questions of each:

Where do the weights live?

Off-chip DRAM (HBM, GDDR, LPDDR), a tier of SRAM beside DRAM, on-chip SRAM only, or inside the memory array itself.

What moves on each step?

Every weight once per decode step, plus the KV cache. At batch 1 that is about two flops per weight read.

Who decides when it moves?

The hardware at run time (warps and caches), the compiler ahead of time (a static schedule), or a mix.

A fourth axis, how big the engines are, runs from a few large matrix units (a TPU) through a hundred or so cores (a GPU’s SMs) to 900,000 tiny cores on one wafer.

02

The Map: Eight Families on Three Axes

Each family’s position is a category, with a sourced basis on the site’s specifications page; a real chip can sit between two.

FamilyWeights live inWho schedulesGrainSite
GPUs
NVIDIA H100, B200
off-chip DRAMhardware, at run timetens to hundreds of coresch 2
Systolic designs
Google TPU, AWS Trainium
off-chip DRAMcompiler, ahead of timea few large enginesch 3
Wafer-scale
Cerebras WSE-3
on-chip SRAMmixedhundreds of thousands of coresch 4
Compiler-scheduled
Groq LPU
on-chip SRAMcompiler, ahead of timea few large enginesch 5
Reconfigurable dataflow
SambaNova RDU, Graphcore IPU
tiered SRAM + DRAMcompiler, ahead of timetens to hundreds of coresch 6
Meshes of small cores
Tenstorrent Blackhole
tiered SRAM + DRAMmixedtens to hundreds of coresch 7
FPGAs and edge NPUs
AMD Versal AI Engines, Apple Neural Engine
tiered SRAM + DRAMcompiler, ahead of timetens to hundreds of coresch 8
In-memory and analogue
phase-change crossbars (research)
in the memory arraycompiler, ahead of timetens to hundreds of coresch 9
03

The Running Workload: a Decoder Layer

Every comparison runs the same models: Llama 3.1 8B and 70B, with the layer shapes of the Llama 3 paper. One layer of the 70B model holds 855,638,016 weights in four matrices (QKV, attention out, the gated MLP’s up-and-gate and down projections); the model has 70,548,193,280.

$$\frac{\text{flops}}{\text{byte}} \approx \frac{2\,N_w\,B}{N_w\,b/8} = \frac{16\,B}{b}$$

for a batch of \(B\) sequences and \(b\)-bit weights: about two flops per byte for one user at 8 bits. A 2,048-token prefill reuses each weight 2,048 times.

2.03
flop per byte, 70B FP8 decode, batch 1
591
H100 ridge point, FP8 (flop per byte)
70.7 GB
70B in FP8 plus a 1,024-token KV cache

A decode step at batch 1 sits hundreds of times below the ridge: it is a memory read with a little arithmetic attached. That single fact decides most of what follows.

04

GPUs: the Baseline

NVIDIA H100 SXMNVIDIA B200 (HGX B200)
FP8 dense1.98 PFLOP/s h100-page4.5 PFLOP/s hgx-page
HBM80 GB h100-page180 GB dgx-b200-page
HBM bandwidth3.35 TB/s h100-page8 TB/s dgx-b200-page
NVLink (both ways)900 GB/s h100-page1.8 TB/s hgx-page
Power700 W h100-page1.79 kW dgx-b200-page

One H100 holds Llama 3.1 70B in FP8. A decode step for one user is bounded by reading it from HBM: 20.8 ms, 48.1 tokens per second, using 0.34% of the arithmetic. Batching fixes that: on eight H100s the same step becomes compute-bound at batch 1,024 (chapter 2 animates the climb).

Notes

B200 figures are per GPU from the HGX B200 totals (dense is half the sparse figure stated); its power is a share of the DGX B200 system’s maximum, not a chip TDP.

05

Systolic Designs

Google’s TPUs and AWS’s Trainium build their matrix engines as weight-stationary systolic arrays (256 × 256 MXUs on TPU v6e and TPU7x, a 128 × 128 grid on Trainium2), fed from HBM by a compiler.

Google TPU7x (Ironwood)AWS Trainium2
FP8 dense4.61 PFLOP/s tpu-7x1.3 PFLOP/s trn2-arch
HBM206 GB tpu-7x103 GB trn2-arch
HBM bandwidth7.38 TB/s tpu-7x2.9 TB/s trn2-arch
Array rows256 tpu-arch128 nki-trn2
TensorCores / NeuronCores2 tpu-7x8 trn2-arch
0.39%
one 70B layer on a 256² array, batch 1
99.9%
the same layer at batch 256

A large array needs a batch as wide as itself to fill it. More on the arrays themselves: the Google TPU series and Systolic Arrays Explained.

06

Capacity Is the First Trade

Where the weights live sets how many chips it takes to hold a model. The fewest chips of each family’s representative part that hold Llama 3.1 70B (weights plus one KV cache, no room for anything else):

1101001,00010,000100,000GPU1Systolic1Wafer-scale4LPU chain308Dataflow3Mesh3Edge NPU3In-memory16,570
Chips of each family’s representative part needed to hold Llama 3.1 70B and one 1,024-token KV cache (log scale). Colour: where the weights live.

One H100 or TPU7x; 4 Cerebras wafers in 16-bit (the number Cerebras publishes); 308 first-generation GroqChips; 16,570 research crossbar chips.

07

Every Family on the Same Model

Llama 3.1 70B, one user, a 1,024-token context and a 2,048-token prompt, on the fewest chips of each family’s representative part (chapter 10 animates it):

FamilyPart, formatChipsPer tokenTokens/s, one userPrefill (2,048)J/tokenDecode utilisation
GPUsNVIDIA H100 SXM, FP8120.8 ms48.1147 ms14.60.34%
Systolic designsGoogle TPU7x, FP819.44 ms10662.9 msnot disclosed0.33%
Wafer-scaleCerebras WSE-3, BF1646.63 µs151,0002.32 ms0.1794.27%
Compiler-scheduledGroqChip, INT83081.09 ms915387 ms0.2350.06%
Reconfigurable dataflowSambaNova SN40L RDU, BF16325.8 ms38.8152 msnot disclosed0.29%
Meshes of small coresTenstorrent Blackhole, FP8345.4 ms22146 ms40.80.16%
FPGAs and edge NPUsApple M4, INT83194 ms5.172.55 snot disclosed0.64%
In-memory and analogueIBM HERMES project chip, INT816,57042.7 µs*23,400315 µs0.01430.31%

* crossbar time only (attention, digital partial sums and chip-to-chip transfers are not counted). Energy is TDP × the bound’s time, only where a power figure is given.

1/s100/s10,000/s1,000,000/sGPU48.1/sSystolic106/sWafer-scale151,000/sLPU chain915/sDataflow38.8/sMesh22/sEdge NPU5.17/sIn-memory23,400/s
The model’s bound on one user’s tokens per second, Llama 3.1 70B at batch 1 (log scale). Bounds, not measured speeds: the SRAM designs reach a small share of theirs.
08

Bounds Against Published Figures

The model takes no published speed as an input; it is checked against them. Ratio = published / model:

CheckKindPublishedModelRatioSource
Llama 3.1 70B, 16-bit: CS-3 systems neededcapacity441cerebras-inference
a 20B model, 16-bit: CS-3 systems needed (weights only)capacity111cerebras-inference
Llama 2 70B, INT8: GroqChips used vs the fewest that hold itcapacity5763011.91nextplatform-groq
8x H100, Llama 2 70B, FP8, MLPerf Offline tokens/s vs the compute ceilingthroughput21,806115,2000.189mlperf-h100
576 GroqChips, Llama 2 70B: tokens/s per user vs the boundlatency3003,6440.0823nextplatform-groq
4x CS-3, Llama 3.1 70B: tokens/s per user vs the boundlatency450150,7000.00299cerebras-inference
1x CS-3, Llama 3.1 8B: tokens/s per user vs the boundlatency1,8001,387,0000.0013cerebras-inference
16 SN40L sockets, Llama 3.1 8B, BF16: tokens/s per user vs the boundlatency1,0421,7910.582sn40l
16 SN40L sockets, Llama 3.1 70B, BF16, speculative decoding: tokens/s per user vs the one-token-per-step boundspeculative457203.32.25sn40l
How to read it

Chip counts match. For designs that read their weights from DRAM the bound is close (MLPerf’s H100s reach 19% of the compute ceiling; the SN40L paper’s 8B figure is 58% of its bound). For SRAM designs it is not: once weights are in SRAM, per-layer latency the vendors do not break down limits one user. The “speculative” row beats a one-token-per-step bound because speculative decoding emits several tokens per weight read.

09

What the Families Are Converging On

Fast memory beside capacious memory

SRAM, HBM and DDR tiers on SambaNova’s RDUs; NVIDIA’s announced Groq 3 LPX racks pair SRAM-only LPUs with GPUs; Cerebras’s announced CS-4 raises SRAM bandwidth to 43.2 PB/s per wafer.

Narrower numbers

FP8 on H100, B200, TPU7x and Trainium2; block FP8 on Tenstorrent; INT4 on edge NPUs. Halving the weight width halves the bytes of a decode step.

The compiler owns the data movement

XLA on TPUs, whole-layer fusion on RDUs, a static schedule on LPUs, per-core data-movement programs on meshes.

The interconnect moves onto the chip

NVLink, TPU inter-chip links, NeuronLink, and Ethernet ports on Tenstorrent’s cards.

Next: deck 02, the SRAM machines, and deck 03, dataflow, meshes, edge devices and in-memory compute.