The Hardware Matrix Multiplier Unit

What this deck covers

Twenty-three slides, broadly grouped:

1. Why a dedicated MMUL exists — energy and the roofline
2. A short history: from MAC to TPU
3. The systolic principle
4. Output-stationary array — TPU style
5. Weight-stationary array — NVDLA / edge style
6. Input-stationary & row-stationary (Eyeriss)
7. Dot-product trees — NVIDIA Tensor Cores
8. What runs on an MMUL? Mapping to the Transformer
9. Number formats — FP32 to FP4
10. Block-scaled microformats — MXFP, NVFP4
11. Stochastic rounding & mixed accumulation
12. Real systems — TPU lineage
13. Real systems — NVIDIA Tensor Cores
14. Real systems — Cerebras / Groq / Dojo / Trainium / ANE / SME / Sohu
15. Memory hierarchy — what feeds the MMUL
16. Are caches useful?
17. Performance: peak vs sustained
18. Power: every joule comes from somewhere
19. Thermals & power delivery
20. Tradeoffs: array size, format, fill bubbles
21. Interactive: dataflow comparator
22. Interactive: number-format explorer
23. Where this is going (2025 → 2027)

Why a dedicated MMUL exists

A modern LLM forward pass is more than 99% multiply-accumulate (MAC). On a general-purpose CPU, almost all the energy spent on that workload is not the math — it's fetching operands from DRAM and shuttling them through register files.

Horowitz, ISSCC 2014 — the numbers that built an industry

An MMUL only wins if every operand fetched from DRAM is reused O(N) times across many MACs. That reuse — temporal in registers and scratchpads, spatial across thousands of PEs — is the entire reason the architecture exists.

The roofline argument

Williams et al. (2009): performance is bounded by min(peak FLOP/s, AI · peak BW), where arithmetic intensity AI = FLOPs per byte of DRAM traffic.

• Vector add: AI ≈ 1/12. Firmly memory-bound.
• 1024×1024 matmul: AI ≈ 2N/3 ≈ 683 FLOPs/byte. Compute-bound on every modern device.
• That's the only common kernel with this property — and it's the one the MMUL is built to exploit.

A short history

The throughline: spatial reuse via locally-connected PEs. Forty-eight years separate Kung's 1978 paper from the Blackwell datasheet — the slide deck has changed, the principle hasn't.

The systolic principle

Kung's metaphor: blood through a heart. Data pumps through a regular array of locally-connected processing elements. Operands enter at the edges, results emerge after a fixed number of cycles, no element ever stalls waiting on global memory.

Three properties that fall out for free

• Regular wiring. Every PE talks only to nearest neighbours — short wires, easy floorplan, no crossbars. Critical at 5 nm where wires dominate energy.
• Pipelining. One MAC per PE per cycle once the array is full — peak throughput = ROWS · COLS MACs per cycle.
• Bandwidth amortisation. Each operand entering the edge is reused O(N) times before leaving. DRAM bandwidth requirements scale as N, not N² or N³.

A 4×4 output-stationary mesh. A streams left→right (blue), B streams top→bottom (purple), accumulators stay put (green).

Output-stationary array — the TPU pattern

Used by every Google TPU (v1 onwards), AWS Trainium / Inferentia, and most "MXU-class" datacentre accelerators.

Cycle-by-cycle

• A flows left → right. Each PE forwards a_in to its eastern neighbour on every cycle.
• B flows top → bottom. Each PE forwards w_in to its southern neighbour.
• The accumulator stays. PE(r, c) computes C[r,c] += A[r,k] * B[k,c] over K cycles — hence "output stationary".

This repo's mmul_pe.sv is the canonical PE: one signed multiply, one signed add, three flops (a-skew, w-skew, accumulator). mmul_systolic_os.sv instances ROWS·COLS of them with edge wiring.

always_ff @(posedge clk or negedge rst_n) begin
  if (!rst_n)         { a_out, w_out, acc_out } <= '0;
  else if (clear)     { a_out, w_out, acc_out } <= '0;
  else if (enable) begin
    a_out   <= a_in;                     // skew right
    w_out   <= w_in;                     // skew down
    acc_out <= acc_out + acc_t'(a_in * w_in);
  end
end

Latency

• Fill: ROWS + COLS − 1 cycles before every PE is computing.
• Steady: K cycles producing R·C MACs each cycle.
• Drain: same again. Result: total = K + R + C − 2.

The "fill bubble" is the cost. For 256×256 with K=64 (typical attention head), 511 cycles of fill+drain wrap 64 productive cycles — 87% bubble. Real TPU code amortises this by chaining many K-tiles back-to-back without draining.

Weight-stationary array — the edge / inference pattern

Used by NVDLA, Apple Neural Engine, ARM Ethos-N, Hailo-8, Qualcomm Hexagon NPU, and most mobile / edge accelerators.

Idea

For inference, weights are fixed and reused across many activations (especially across a batch). So load weights once, hold them, stream activations. Partial sums spill to a separate buffer.

The cost ledger

• Pro: Weights never re-fetched — energy directly on the budget. Lower latency (no fill-skew).
• Con: Higher operand-bandwidth per cycle: M·N reads of activations and acc-spills per step.
• Con: Weight reload is paid every fresh tile — bad for very small networks but invisible for layer-by-layer streaming.

This repo's mmul_weight_stationary.sv implements an M × N register-file held over K compute steps; a separate load_w handshake fills it from the off-chip weight stream.

Input-stationary & row-stationary (Eyeriss)

Two more dataflows for completeness — both more specialised, both interesting.

Input-stationary

Activations pinned, weights stream, partial sums flow. Less common in dense matmul, but appears in some training accelerators where the activation reuse pattern of ∂L/∂W = A^T · ∂L/∂Y dominates the gradient pass.

Row-stationary — Eyeriss (MIT, ISCA 2016)

The most thoughtful of the four. Maps a 1-D row of CNN filter weights onto a single PE; ifmap rows stream horizontally; partial sums flow vertically. Sze and Chen showed this maximises all three reuse patterns simultaneously — weights, ifmaps, and psums.

• 168 PEs, 12×14 array, 108 KB GLB
• 1.4–2.5× more energy-efficient than weight- or output-stationary on AlexNet
• Idea generalises to any K-loop where the inner operand is a small spatial filter

Dot-product trees — NVIDIA Tensor Cores

NVIDIA's Tensor Core is conceptually not a Kung-style systolic. It's a parallel reduction tree: K multipliers feeding an adder tree, with operands fetched per warp from the register file.

1st-gen V100 Tensor Core (2017)

• Computes D = A·B + C on 4×4 matrices.
• A, B FP16; C, D FP16 or FP32.
• 64 FMAs per clock per TC, 8 TCs per SM, 84 SMs ⇒ 125 TFLOPS FP16 / V100.

Why a tree, not an array?

• Lower latency. O(log K) adder-tree depth vs O(K + R + C) in a systolic. Latency matters for warp-level scheduling.
• Flexible shapes. NVIDIA exposes mma.sync shapes like 16×8×16 — easy to compose at the warp level.
• Higher SRAM bandwidth required. The tree consumes K operand pairs per cycle from the register file, with no skew amortisation.

This repo's mmul_dotproduct.sv is a single-cycle K-wide combinational dot product — the simplest model of one Tensor-Core lane. With K = 8 it computes a Q8.8 inner product into a 32-bit accumulator combinationally.

Generation by generation

What does an MMUL run, in a Transformer?

Almost all of it. A typical decoder-only LLM forward pass spends ≥99% of its FLOPs in dense matrix multiplications — so the MMUL is the part of the chip that decides how fast a Llama or GPT runs.

Per-block FLOP map (decoder layer, sequence length S, model dim D, head dim d, FFN dim 4D)

What this looks like on H100 / B200 silicon

• Tensor Cores run all the GEMMs above (highlighted "Yes"). Tensor-Core lanes are idle during softmax, RMSNorm, RoPE, and SiLU — those run on the SM's CUDA cores or Special Function Units.
• FlashAttention 2/3 fuses Q·K^T, softmax, and scores·V into a single kernel that streams partial maxes — keeping data in registers and SRAM and avoiding the materialised S×S score matrix. The MMUL is still doing the GEMMs, but they're now tiled and chained instead of serialised through HBM.
• KV-cache: the K and V tensors of every previous token are held in HBM (often quantised to FP8 or INT4). At decode time, generating one new token reads MS·D bytes per layer — for a 70B model with 100k context, that's tens of GB per token. Decode is memory-bound; the MMUL runs at 1–3% utilisation.

Two regimes that matter for hardware

Number formats: from FP32 to FP4

Halving the bit-width of the multiplier is approximately 4× fewer transistors (multiplier area is quadratic in mantissa width) and ~2× lower energy per MAC. Every generation has chased this curve.

The two design choices that matter

• Range vs precision. BF16 keeps FP32's 8-bit exponent (handles gradient outliers) at the cost of a 7-bit mantissa. FP16 trades the other way — better mantissa, narrow range, requires loss-scaling for training.
• What to do at FP8 and below. 8-bit mantissa is too small for whole-tensor scaling. The fix: per-block scales — see next slide.

Block-scaled microformats — MXFP & NVFP4

Below FP8, the dynamic range of any single element is too small to cover a full tensor. The industry's answer: group elements into a small block, give each block a shared exponent.

OCP MX (Microscaling) — September 2023

• Block size: k = 32 elements.
• Scale: E8M0 — 8-bit unsigned biased exponent only, no sign or mantissa. Value = 2^{scale − 127}.
• Element types: MXFP8 (E4M3 or E5M2), MXFP6 (E3M2 or E2M3), MXFP4 (E2M1), MXINT8.
• Storage: MXFP8 block = 32×8 + 8 = 264 bits (33 B). MXFP4 block = 136 bits (17 B).
• Backed by AMD, ARM, Intel, Meta, Microsoft, NVIDIA, Qualcomm.

NVFP4 — Blackwell-specific

• Block size: k = 16 (vs OCP's 32) — finer-grained scaling.
• Scale type: FP8 E4M3 — mantissa bits in the scale itself, much better SNR near the boundary.
• Plus a per-tensor FP32 global scale.
• NVIDIA's data shows NVFP4 within ~0.5% of FP8 perplexity on Llama-class models.

Stochastic rounding & mixed accumulation

The systematic-bias problem

Round-to-nearest applied many times is biased. Adding 1.0 + 0.49 a million times in FP4 gives 1.0, not 491,490. In low-precision training, this kills convergence.

Stochastic rounding

• Round up with probability proportional to the fractional part.
• Statistically unbiased — the systematic error goes away even at FP4.
• Native on Graphcore IPU (since Colossus); supported in NVIDIA Hopper/Blackwell on the FP8 / FP4 paths.

Two-stage accumulation

Keep the multiplier narrow, keep the accumulator wide:

• TPU MXU: BF16 multiply, FP32 accumulate.
• H100 FP8 path: FP8 mul, FP32 accum, with a periodic FP32 → BF16 spill.
• NVFP4: FP4 mul, FP32 accum, with the per-block E4M3 scale applied once per K=16 step.

Without this, every halving of bit-width would silently halve model quality. With it, the chip behaves numerically like a much wider design at a fraction of the area.

Real systems — Google TPU lineage

Things to notice

• TPU v1 was INT8-only. Designed strictly for inference under a 99th-percentile latency SLA.
• v2 → v6e all stayed at 128×128. Then Trillium doubled the MXU side to 256×256 — first resize since 2017.
• The MXU is one part of the chip. v4 has 32 MiB CMEM, 32 GiB HBM, and the optical-circuit-switched (OCS) datacentre interconnect that lets you reshape pods at runtime.
• Software is half the architecture. XLA has been doing tile-shape optimisation for nine years — without it Trillium's bigger MXU would routinely run below 50% utilisation.

Real systems — NVIDIA Tensor Cores

What changed across generations

• Volta → Ampere: format expansion (FP16 → BF16 / TF32 / INT8 / FP64) and 2:4 sparsity.
• Ampere → Hopper: FP8 + Transformer Engine + TMA — async data movement decoupled from warps.
• Hopper → Blackwell: FP4/FP6/MXFP + TMEM (decoupled MMA inputs from shared memory).
• The pattern: each generation adds a narrower datatype, plus one new memory primitive.

Real systems — the rest of the field

Three philosophies on display

• Datacentre generality (NVIDIA, TPU, AMD, Trainium): big array + hierarchical memory, run anything.
• Determinism & in-memory (Groq, Cerebras): banish DRAM, schedule everything ahead of time. Spectacular for inference latency on supported workloads, brittle elsewhere.
• Workload specialisation (Etched Sohu, Apple ANE, MatX): hardwire one model class, claw back area-efficiency. The asymptote of "do less, faster".

Memory hierarchy — what feeds the MMUL

An MMUL is a giant operand-consumer. A 256×256 BF16 array at 1 GHz needs 256 KB/cycle of A operands and 256 KB/cycle of B. That's ~256 TB/s of pure operand bandwidth. No HBM stack delivers that — it has to come from on-die SRAM.

The hierarchy you actually find on a modern AI chip

The structures you see on every modern AI chip

• Activation buffer / unified buffer: feeds A operands. TPU v1's "UB" was 24 MiB; H100 SMEM is 228 KB/SM × 132 SMs ≈ 30 MB.
• Weight buffer / weight queue: feeds B operands. Often double-buffered so the next tile's weights load while the current tile computes.
• Accumulator buffer: holds partial sums, often 32-bit even when operands are FP4/INT8. TPU v1 had 4 MiB of accumulators.
• Async DMA engines: TMA on Hopper, the TPU v4 "VMEM async" — move tiles between HBM and SRAM without consuming compute lanes.
• Tensor-private memory: TMEM on Blackwell — 256 KB/SM dedicated to tensor-core operands so SMEM pressure doesn't strand the array.

Are caches useful?

Short answer: not in the CPU sense — and modern accelerators have largely abandoned them in favour of explicitly-managed scratchpads.

Why CPU-style caches don't fit an MMUL

• Access patterns are statically known. The compiler knows exactly which weight tile and which activation tile feed the next K-step. Tag arrays, replacement policies, and miss-handling logic are dead silicon — pay the area for the data, not the bookkeeping.
• Tag energy is comparable to data energy at small block sizes. A traditional cache spends ~30% of access energy on tag lookup. A scratchpad is just data.
• Coherence is unwanted. Weights are read-only during a forward pass. There's no producer-consumer race to track.
• Capacity efficiency. A 50 MB L2 with 25% tag/coherence overhead becomes 50 MB of usable scratchpad in an accelerator design.

Where caches DO appear on modern AI chips

Why GPUs keep an L2

The H100's 50 MB L2 isn't really there to serve the Tensor Core MMUL — it's there to serve the rest of the SM (CUDA cores, texture units, atomics, indirect memory accesses). On dedicated AI chips that don't run general code, the L2 budget gets reallocated as larger user-visible scratchpad.

Performance: peak vs sustained

The peak number on the datasheet is the multiplier count × clock frequency × ops per multiplier × 2 (for FMA). Real workloads achieve 30–70% of that.

What eats the gap

• Tile-shape mismatch. A 256×256 MXU running a 192-dim attention head wastes 64 columns. Trillium's bigger MXU buys peak FLOPS at the cost of small-tile utilisation.
• Fill / drain bubbles. A 64-step K loop on a 256×256 array spends 511 of its 575 cycles filling or draining unless multiple tiles are chained.
• Memory-bound phases. Attention's K/V projection is GEMM-heavy; the softmax is bandwidth-bound; the output projection is GEMM-heavy again. Peak FLOPS is invisible to the bandwidth-bound stretches.
• Quantization overhead. NVFP4 dequantise + per-block scale-apply is a non-trivial fraction of the FP4 path's pipeline.
• Synchronisation. Cross-array reductions and all-reduces between chips set the floor.

Indicative numbers from public benchmarks

The 5× gap between "MoEFLOPS marketing slide" and "actual throughput on your model" is not bad design — it's the cost of running real workloads on real memory hierarchies. Compiler engineering closes most of it.

Power: every joule comes from somewhere

Where the watts go on an H100

• Tensor Cores: ~30–40% of die power at full FP8 dense GEMM.
• Memory hierarchy: ~25–35% — HBM PHYs, L2, SMEM banks, DMA engines.
• NVLink / interconnect: ~10–15% on a SXM5 (~700 W package).
• CUDA cores, schedulers, fixed-function: ~10–15%.
• Static / leakage: ~5–10% at modern process nodes.

Energy per useful operation — order of magnitude

Two things to notice: (1) HBM access does not scale with process node — it's PHY/wire-dominated. So as compute energy per MAC drops, the relative cost of operand movement becomes higher. (2) The whole point of MX/FP4 is to get the multiplier energy below the operand-movement energy that flanks it.

Package power is the hard limit

• H100 SXM5: 700 W TDP
• B200 SXM6: 1000 W TDP
• B200 NVL72 rack: 120 kW per rack
• TPU v5p tray: ~15 kW liquid-cooled
• Tesla Dojo training tile (25 chips): 15 kW liquid-cooled

Above ~700 W per package, air cooling gives up. Hence the industry-wide pivot to liquid-cooled cabinets in 2024–2026.

Thermals & power delivery

A modern MMUL package consumes more current than an early-2000s consumer CPU's whole motherboard. The interesting engineering isn't in the array — it's in getting the joules in and the heat out.

Power delivery network (PDN)

• Voltage: typical core rail 0.75–0.85 V.
• Current: at 1000 W and 0.8 V → 1250 A through the package socket. NVIDIA B200 spec'd at >2000 A peak.
• I²R loss: with 1 mΩ socket resistance, 1250 A dissipates 1.5 kW just in the socket. That's why VRMs are placed millimetres from the package and why integrated voltage regulators (IVR) are standard on B200 / TPU v5p.
• di/dt: turning a 256×256 array on/off in one cycle is ~250 A/ns. Decoupling caps span four orders of magnitude (bulk → on-die MIM caps → finFET MOS caps).
• Wafer-scale (Cerebras): power delivered through a vertical stack — the WSE-3 backplane uses thousands of vertical pins per dm² to bypass copper-resistance limits at wafer scale.
• Bow IPU (Graphcore 2022) takes a different route: wafer-on-wafer bonded a power-delivery die under the compute die — first commercial WoW stacked PDN.

Thermal envelope

Hotspot density

A 256×256 array at FP8 dissipates roughly 100 W in ~50 mm² of silicon — that's a flux of ~200 W/cm², several times higher than a kitchen hotplate. Most of the silicon area in a Tensor-Core SM is not multipliers; it's the operand-fan-in network, the result-spill paths, and the local SRAM banks that surround them — partly because this spreads the heat out so hotspot temperatures stay below the 105 °C electromigration limit.

Tradeoffs & design space

Array size

• Small (4×4×4 V100 TC, 8×8 ARM SME ZA): high utilisation on small/odd shapes; underexploits operand reuse.
• Medium (128×128 TPU v2–v5p): sweet spot for transformer-class workloads.
• Large (256×256 Trillium, 256×256 TPU v1): best peak FLOPS, worst utilisation on small batches; needs careful tile fusion.
• Wafer (Cerebras 900K cores): minimises operand movement across an entire model; needs a dedicated compiler.

Format vs accuracy

• FP32 → BF16 in 2018 was effectively free for inference and training (Google's BF16 paper).
• BF16 → FP8 in 2022 needed Transformer Engine and per-tile scaling — not free, but practically lossless on most LLMs.
• FP8 → FP4 in 2024 needs MXFP/NVFP4 + stochastic rounding + careful loss-scaling. Ongoing.
• Below FP4: very active research, no shipping silicon.

Specialisation vs generality

• NVIDIA Tensor Cores: any reasonable shape, any reasonable format. Pays in silicon area.
• TPU MXU: fixed shape, BF16 / INT8 / FP8 only. Half the area for the same throughput.
• Etched Sohu: hardwired transformer dataflow. ~10× the throughput per area, runs only one model class.

Reconfiguration cost

• Switching numeric mode (e.g. FP16 ↔ FP8) on Hopper: free per-kernel.
• Switching weights between layers: real cost — hence weight-stationary on edge, output-stationary for training.
• Switching workloads (inference vs training): the whole compiler stack changes.

Interactive — dataflow comparator

Choose an array geometry and see the resulting cycle count and operand bandwidth for each dataflow.

Notice how the OS array's utilisation collapses for small K (fill bubble dominates) and the dot-product tree wins on latency but loses on per-cycle bandwidth. Real chips combine both — TPUs use OS-systolic for the MXU, NVIDIA TCs use trees, and both schedule many K-tiles back-to-back to hide the bubble.

Interactive — number-format explorer

Pick a format and a real number; see what gets stored.

Try 1e-4 in FP4 (you'll round to zero) versus BF16 (full precision). Try 50000 in FP8 E4M3 (saturates to ±448) versus FP8 E5M2 (representable). This is exactly why H100 pairs E4M3 (small range, more precision — for forward) with E5M2 (large range — for gradients).

Where this is going (2025 → 2027)

• Block-scaled microformats become default. NVFP4 (Blackwell), MXFP4 (AMD MI355X, Microsoft Maia 100), MXFP8/6 across the rest of the field. By 2026, "FP8" will mean MXFP8, not raw OFP8.
• Training in FP4. NVIDIA's late-2024 papers showed convergence parity at FP4 with stochastic rounding + per-tensor scales. Expect this to be production by 2026.
• 3D-stacked SRAM. AMD MI300X's 3D V-Cache lineage points the way; rumoured Blackwell-NeXT brings SRAM-on-logic to NVIDIA. Closing the gap to wafer-scale's per-mm² SRAM.
• Optical I/O. TPU v4's Optical Circuit Switch and Lightmatter Passage demonstrate what TB/s of cabinet bandwidth at single-pJ/bit looks like.
• Transformer-only silicon. Etched, MatX, Cerebras CS-3 dense-KV-cache mode — all bet that the cost of generality is too high relative to hardwiring one model class.
• ISA-level matrix instructions. ARM SME, Intel AMX, RISC-V's nascent matrix extension — every CPU vendor is now adding a small MMUL into the core. The line between "CPU" and "AI chip" continues to blur.
• Power delivery, not transistor density, is the limit. Backside power delivery (Intel PowerVia, TSMC SPR) is the 2026–27 process feature that buys ~5% performance at iso-power — and matters precisely because the array is current-limited, not transistor-limited.