Modern SoC Design Series · Presentation 05

AI / ML Accelerator
Architectures

From roofline models and dataflow engines to systolic arrays, sparsity exploitation, and the micro-architecture of today's NPUs — a rigorous tour of how silicon is shaped to maximise AI throughput per watt.
Roofline Model Systolic Arrays Dataflow Engines Sparsity Mixed Precision NPU / DSA Transformer Accelerators Case Studies
02

Why Specialised Hardware?

the end of free scaling

Dennard Scaling — Broken ~2006

  • Voltage no longer scales with feature size — power density rises
  • Clock frequency stalls; multi-core alleviates, not eliminates, the wall
  • Memory bandwidth grows far slower than compute — the "memory wall"
  • ILP extraction in general-purpose CPUs hits diminishing returns

Compute Demand vs Supply Gap

  • GPT-3: ~3.14 × 10²³ FLOP training
  • LLM inference: O(2·N·L) FLOP/token for N-param, L-layer model
  • Video diffusion, RL — orders of magnitude beyond NLP
  • NVIDIA A100 GPU: ~312 TFLOPS (FP16); not enough at scale

ML Workload Characteristics

  • Regularity — repeated MAC operations with fixed structure
  • Data reuse — weights reused across many inputs (batching)
  • Low precision — INT8 or FP8 often sufficient for inference
  • Sparsity — zero activations & pruned weights → skip logic
  • Parallelism — SIMD, tensor, pipeline — all exploitable
Domain-Specific Architectures (DSAs) align micro-architecture with workload structure: fixed datapaths, scratchpad memories, reduced control overhead, and dedicated precision support — yielding 10–1000× efficiency over CPUs.
03

The Roofline Model

arithmetic intensity as the design compass

Attainable FLOPS = min( Peak FLOPS , I × Peak BW )

Arithmetic Intensity I

DefinitionFLOP / byte of DRAM traffic
Dense GEMM (large)~300 FLOP/B
Conv2D (large kernel)~100 FLOP/B
Attention (MHA, seq=2k)~15 FLOP/B
Batch-1 GEMV (inference)~1–4 FLOP/B
Element-wise ops (ReLU)<1 FLOP/B

Ridge Point — NVIDIA H100 SXM

FP16 Tensor TFLOPS (dense)989
HBM3 bandwidth3.35 TB/s
Ridge point (dense FP16)~295 FLOP/B
GEMV inferenceMemory-bound !
GFLOPS/s Arithmetic Intensity (FLOP/B) Ridge GEMV Attn Conv GEMM Memory-bound Compute-bound 2 15 100 ~591
Accelerator design is an exercise in shifting the ridge point: boost compute throughput, reduce DRAM bandwidth demand through tiling & data reuse, or do both — placing target workloads in the compute-bound regime.
04

Compute Primitives: MAC → GEMM → Tensor

the hierarchy of matrix operations

Scalar MAC

c += a × b
  • Single multiply-accumulate per cycle
  • CPU FPU or ALU path
  • 1 FLOP / cycle (scalar)
  • Poor reuse — reload A,B each time

SIMD / Vector

C[i] += A[i] × B[i]
  • AVX-512: 16 FP32 MACs / cycle
  • Arm SVE2: scalable vector length
  • Still 1D — limited tensor reuse
  • Compiler auto-vectorises easily

Tensor / Matrix Unit

C += A × Bᵀ
  • NVIDIA Tensor Core: 16×16×16 MMA
  • Google TPU MXU: 128×128 systolic
  • Apple AMX: 32×32 tiles per cycle
  • Maximises data reuse; bandwidth ↓↓

GEMM Tiling Strategy — Fitting the Memory Hierarchy

  • Tile into L1/register file — keep hot data in the fastest storage; amortise load cost over many MACs
  • K-dimension tiling — partial sums accumulate in accumulators, not DRAM
  • M × N output stationary — output tile stays; stream A and B rows/cols past it
  • NVIDIA cuBLAS/CUTLASS — auto-tunes tile sizes per GPU architecture (warp-level MMA)
  • Shared memory as L1 scratchpad — explicitly managed; programmer loads tiles once
  • Memory-compute overlap — async copy + Tensor Core instructions interleaved
05

Systolic Arrays

data flowing like a heartbeat through a 2D PE mesh

Operating Principle

  • 2D grid of Processing Elements (PEs), each performing MAC
  • Data injected from edges; flows rhythmically (systolically) across the array
  • Each PE accumulates a partial sum before passing data to its neighbour
  • Weight stationary: weights pre-loaded into PE register files; activations stream through
  • Output stationary: partial sums stay; both A and B flow (lower register pressure)
  • Row stationary (Google TPU v1): minimises total data movement energy

Google TPU v1 MXU

Array size256 × 256 PEs
PrecisionINT8 MACs
Peak throughput92 TOPS
Clock700 MHz
Weight FIFO28 MiB on-chip
Key insightNo caches → predictable
MAC MAC MAC MAC Activations → Weights ↓ Partial sums →
Weight-stationary 4×4 systolic array (concept). In practice the TPU v1 MXU is 256×256.
06

Dataflow Architectures

execution driven by data availability, not program counters

Dataflow vs Control-flow

  • Control-flow (von Neumann): sequential PC fetch/decode; memory bottleneck; branch misprediction overhead
  • Dataflow: instructions fire when operands are ready; implicit parallelism; no speculative execution overhead
  • ML graphs are acyclic (mostly) → perfect fit for dataflow scheduling
  • Compiler maps layer graph onto PE array dataflow schedule at compile time — zero runtime overhead

Dataflow Variants

  • Weight-stationary: weights loaded once into PE RFs; minimise weight movement energy
  • Output-stationary: partial sums accumulate locally; maximise accumulator reuse
  • Row-stationary (TPU): each PE computes one convolution row; minimises total data movement
  • No-local-reuse (NLR): all operands from global buffer — less efficient but simpler PE

Spatial Dataflow Engines

  • Cerebras WSE-3: 4 T transistors · 900,000 cores · 44 GB SRAM · 21 PB/s on-wafer BW · 125 PFLOPS peak; CS-3 inference service serves Llama-3.3-70B at >2,000 tok/s
  • Groq LPU (TS-P1): 14 nm deterministic TSP; 188 INT8 TOPS · 230 MB on-chip SRAM · multi-chip serving delivers ~280 tok/s (Llama-3.3 70B) or ~1,665 tok/s with speculative decoding
  • SambaNova SN40L: reconfigurable dataflow units (RDU); spatial compute array with streams; 2-tier DRAM+HBM
  • Tenstorrent Blackhole: mesh of RISC-V tensix cores; compiler-controlled NoC; open-source toolchain
Compiler is the hardware: spatial dataflow chips push scheduling complexity into ahead-of-time compilation. Hardware becomes a static, predictable substrate — runtime control logic shrinks dramatically, freeing area for compute.
07

Numerical Formats & Mixed Precision

matching bit-width to workload sensitivity

FormatBitsSignExpMantDynamic Range
FP646411152~1.8×10³⁰⁸
FP32321823~3.4×10³⁸
TF32191810FP32 range, ↓precision
BF1616187FP32 range (easy cast)
FP16161510±65504 (overflow risk)
FP8 E4M38143Fwd pass preferred
FP8 E5M28152Grad pass preferred
INT881—7Inference standard
INT4 / NF441—3QLoRA, edge inference

NVIDIA Tensor Core Evolution

Volta (V100)FP16 in, FP32 acc
Ampere (A100)TF32, BF16, INT8, INT4, FP64
Hopper (H100)FP8 E4M3/E5M2, FP16, BF16
Blackwell (B200)FP4 NV (2× H100 throughput)
B200 FP4 peak~10 PFLOPS dense / 20 PFLOPS sparse

Mixed-Precision Training (AMP)

  • Master weights: FP32 — guards gradient precision
  • Forward/backward: FP16 or BF16 — feeds Tensor Cores
  • Loss scaling: multiply loss before backward; prevents FP16 underflow
  • Accumulation: FP32 accumulator in Tensor Core MMA instruction
  • Net result: 2–3× throughput vs pure FP32, near-identical accuracy
08

Sparsity Exploitation

skipping zeros — architectural and algorithmic approaches

Sources of Sparsity

  • Activation sparsity: ReLU sets negative activations to zero; 50–80% sparse in practice
  • Weight pruning: iterative magnitude pruning removes small-weight connections
  • Structured pruning: remove entire filters/heads — hardware-friendly but accuracy cost
  • Unstructured pruning: arbitrary zero pattern — needs sparse HW to exploit
  • Dynamic sparsity: varies per input; harder to schedule statically

NVIDIA 2:4 Structured Sparsity

  • Exactly 2 non-zeros in every group of 4 weights
  • Stored in compressed format: 50% weight storage
  • Sparse Tensor Cores decode in hardware — 2× effective throughput
  • Metadata (2-bit indices per group) stored alongside values
  • Ampere+: hardware decompression before MMA; no software overhead
  • Accuracy: <1% loss for most vision/NLP models with retraining

Sparse Format Trade-offs

  • CSR/CSC: row/col pointer + value arrays; irregular memory; good for very sparse (>90%)
  • COO: (row, col, val) tuples; simple but high overhead at <90% sparsity
  • Block-sparse (BSR): 16×16 or 32×32 dense blocks at known positions; GPU-friendly
  • Flash Attention sparse: attends only to local & strided positions; O(N√N) compute
  • MoE routing: dynamic expert selection = structured sparsity in FFN layers

Mixture-of-Experts (MoE) — Sparsity at the Architecture Level

  • Replace dense FFN with N expert sub-networks (e.g. 8, 64, or 256)
  • Router selects top-K experts per token; only those experts execute
  • Parameter count scales with N, compute scales with K — decoupled
  • Mixtral 8×7B: 46.7B params, ~12.9B active; inference cost of 12.9B dense
  • GPT-4: rumoured MoE — explains scale vs serving cost
  • Load balancing aux loss prevents expert collapse
09

Memory Hierarchy in AI Accelerators

keeping compute fed — the bandwidth pyramid

Register File kB · >10 TB/s L1 SRAM / Scratchpad 256–512 kB · 5 TB/s L2 / Unified Buffer SRAM 50–100 MB · 1–2 TB/s HBM3/3E (on-package) 96–192 GB · 3–5 TB/s NVLink / PCIe / CXL pool TB scale · 900 GB/s ↑ Speed / Cost per GB ↓ Capacity

On-Chip SRAM Sizing Considerations

  • Must hold input/output tiles + partial sums simultaneously
  • Larger SRAM → better roofline position; SRAM density limits die area
  • 6T SRAM: ~0.026 µm² / bit in N5; 192 MB SRAM ≈ 40 mm²
  • NVIDIA H100: 50 MB L2; Google TPU v5e: 16 MB unified buffer per chip
  • Cerebras WSE-3: 44 GB on-wafer (no DRAM bottleneck for <700B models)

Flash Attention — Algorithm-Hardware Co-design

  • Standard attention materialises full N×N score matrix — O(N²) HBM reads/writes
  • Flash Attention tiles Q,K,V into SRAM; fuses softmax + matmul kernel
  • HBM reads/writes: O(N) — linear, not quadratic
  • 2–4× wall-clock speedup; enables 100K+ context lengths
  • Flash Attention 3: async WGMMA + softmax pipelining on H100; 75% MFU
10

Transformer Accelerator Micro-Architecture

mapping attention and FFN layers onto silicon

Attention Bottlenecks at Inference

  • Prefill (prompt): large GEMM — compute-bound; exploits Tensor Cores well
  • Decode (autoregressive): GEMV — each token generates one row; memory-bandwidth bound
  • KV Cache: past K,V tensors stored in HBM; grows O(seq_len × layers × heads × d_head)
  • At 70B model, 100K context: KV cache ~200 GB — exceeds single GPU VRAM

KV Cache Optimisations

  • Multi-Query Attention (MQA): share single K,V head across all Q heads — 8× KV cache reduction
  • Grouped Query Attention (GQA): G shared K,V heads (Llama 3 uses G=8); balance between MHA and MQA
  • PagedAttention (vLLM): KV cache in non-contiguous pages; near-zero fragmentation; enables batching
  • Quantised KV: INT8/FP8 KV cache; 2× capacity increase at minor quality cost

Dedicated Transformer Engine (NVIDIA)

  • Hopper H100: Transformer Engine with FP8 training support
  • Per-tensor dynamic scaling maintains numerical stability
  • Async WGMMA instructions: Tensor Core + shared mem overlap
  • TMA (Tensor Memory Accelerator): hardware async data copy engine; overlaps compute with memory
  • 4× H100 Tensor Core throughput vs A100 in FP8 mode

Prefill / Decode Disaggregation

  • Prefill is compute-bound; decode is memory-bound → different optimal hardware
  • Disaggregated serving: separate GPU pools for prefill vs decode (Splitwise, DistServe)
  • Prefill pool: large A/H100; decode pool: cheaper memory-bandwidth GPUs
  • Reduces total cost of serving by 30–40% at scale
11

Quantisation for Inference

reducing bit-width without sacrificing accuracy

Post-Training Quantisation (PTQ)

  • Symmetric: scale = max(|W|) / 127; zero-point = 0; preferred for weights
  • Asymmetric: separate scale + zero-point; better for asymmetric activation distributions
  • Per-tensor: one scale for whole tensor; fastest; least accurate
  • Per-channel: one scale per output channel; standard for INT8 weights
  • Per-group (GPTQ, AWQ): groups of 128 weights share a scale; enables INT4 with <1% PPL loss
  • Calibration dataset: 128–512 samples sufficient

Quantisation-Aware Training (QAT)

  • Simulate quantisation during training — fake quantise nodes in forward pass
  • Gradients flow through fake-quant via straight-through estimator (STE)
  • Model learns to be robust to quantisation error
  • Better accuracy than PTQ, especially at INT4
  • QLoRA: 4-bit base model (NF4) + FP16 LoRA adapters — fine-tune 70B on single A100
  • NF4: normally-distributed 4-bit — optimal for normally-distributed weights

Hardware Quantisation Support

NVIDIA A100INT8, INT4
NVIDIA H100FP8, INT8, INT4
NVIDIA B200FP4, FP6, FP8
Google TPU v5INT8, BF16
Apple ANE (M4)INT8 (CoreML)
Qualcomm HTPINT8, INT4 (QNN)
Samsung Mach-1INT8
  • Dequantisation overhead must be hidden by compute; batch size matters
  • W4A16 (GPTQ/AWQ): INT4 weight, FP16 activation — best decode throughput
12

GPU Architecture: NVIDIA Hopper H100

the current standard bearer for AI compute

H100 SXM5 — Key Numbers

ProcessTSMC N4 (4nm)
Die area814 mm²
Transistors80 billion
SMs132 Streaming Multiprocessors
Tensor Cores / SM4 (4th gen)
FP16 Tensor TFLOPS1,979
FP8 Tensor TFLOPS3,958
HBM3 capacity80 GB
HBM3 bandwidth3.35 TB/s
NVLink 4.0900 GB/s bidirectional
TDP (SXM)700 W

Hopper New Features

  • Transformer Engine: hardware FP8 ↔ FP16/BF16 conversion; dynamic per-tensor scaling
  • WGMMA (Warpgroup MMA): asynchronous Tensor Core with larger 64×16×16 tiles
  • TMA (Tensor Memory Accelerator): independent DMA engine; overlaps data movement with WGMMA
  • DPX instructions: native dynamic programming ops (Smith-Waterman, Floyd-Warshall)
  • NVLink 4 Switch: up to 8× H100 GPUs per NVLink domain (NVL8) at 900 GB/s bidirectional per GPU; GB200 NVL72 extends this to 72 GPUs on NVLink 5.
  • Confidential computing: TEE with hardware attestation for multi-tenant inference

SM Micro-architecture

  • 128 CUDA cores (FP32) per SM + 64 FP64 + 64 INT32
  • 4× 4th-gen Tensor Cores per SM
  • 256 KB register file; 228 KB configurable L1/shared memory
  • 4 warp schedulers; 64 warps (2048 threads) max in flight
13

Google TPU Architecture

purpose-built ASIC for TensorFlow / JAX workloads

TPU v1 (2016)
92 TOPS INT8 · 256×256 systolic MXU · no training support · GDDR3 · 40W · ResNet inference
TPU v2 (2017) — Training
45 TFLOPS BF16 · HBM 16 GB · liquid cooled · 4 cores per chip · TPU Pod via 2D torus ICI
TPU v4 (2021)
275 TFLOPS BF16 · OCS optical reconfigurable interconnect · 4096-chip pods · sparse cores
TPU v5e / v5p (2023)
v5e: 393 TOPS INT8, HBM2 16 GB · v5p: 459 TFLOPS BF16, HBM3 95 GB; Trillium arch
TPU v6 Trillium (2024)
4.7× v5e MXU throughput · SparseCore for embedding · ICI at 1.2 Tb/s per chip · TSMC N3

TPU v4 Inter-Chip Link (ICI)

  • Each chip has 6× bidirectional ICI links at 100 GB/s each = 600 GB/s per chip
  • 4096-chip pod arranged as 3D torus topology
  • Optical Circuit Switch (OCS): software-reconfigurable physical topology — optimise per-workload
  • All-reduce over ICI faster than NVLink at pod scale

TPU SparseCore (v5+)

  • Dedicated accelerator for embedding lookups (recommendation models)
  • Embedding tables: sparse random-access pattern; hostile to systolic array
  • SparseCore: gather/scatter engine + local HBM slice
  • Enables unified chip for both dense (MXU) and sparse (SparseCore) workloads
Google operates the world's largest TPU fleet — over 50 exaFLOPS aggregate (estimated). All Gemini models were trained and served on TPU infrastructure.
14

NPU Integration in SoCs

on-device AI — from smartphone to datacenter SoC

Typical NPU Micro-architecture

DMA Engine Cmd Queue Post-proc MAC Array (systolic / SIMD) Accumulators Unified Buffer SRAM Act/Pool/ Norm Unit SoC Fabric (AMBA AXI / CHI)

Flagship Mobile NPUs (2024–2025)

ChipNPUTOPSProcess
Apple A18 ProANE Gen 538TSMC N3E
Apple M4ANE Gen 538TSMC N3E
Qualcomm SD EliteHexagon HTP75TSMC N4P
Samsung Exynos 2500Mach-134Samsung SF3
MediaTek D9400APU 79050TSMC N3E
Google Tensor G4TPU (Samsung)42Samsung SF4

Datacenter SoC NPU Integration

  • Intel Gaudi 3: 8× matrix multiplication engines (MME) + 24× TPC VLIW cores; 64 GB HBM2e
  • AMD MI300X / MI325X / MI350X: CDNA3/CDNA4; up to 288 GB HBM3e, 8 TB/s (MI350X); native FP4/FP6 on MI350/355X
  • AWS Trainium2 / Trainium3: 8 NeuronCores per chip; Trn2 1.3 PFLOPS FP8 + 96 GB HBM; Trn3 (2025) 2.52 PFLOPS FP8 + 144 GB HBM3e at 3 nm
  • Microsoft Maia 100: N5 SoC; 64 GB HBM2E (1.8 TB/s); rack-scale liquid-cooled; 4.8 Tb/s Ethernet-based accelerator fabric; Azure-native for Copilot workloads
15

Parallelism Strategies for Large Models

distributing compute and memory across accelerators

Data Parallelism (DP)

  • Each device holds a full model copy
  • Different micro-batches processed in parallel
  • Gradient all-reduce after backward pass
  • DDP (PyTorch): overlap all-reduce with backward; efficient at <1B params
  • Scales well; limited by per-device memory (full model must fit)
  • ZeRO (DeepSpeed): shard optimizer states, gradients, params across DP ranks

Tensor Parallelism (TP)

  • Partition weight matrices across devices (column / row sharding)
  • Megatron-LM: column-split Q/K/V projections; row-split output projection
  • Each TP rank computes a partial GEMM; all-reduce/all-gather recombines
  • Requires high-bandwidth interconnect (NVLink) — all-reduce per layer
  • TP=8 typical; beyond 8 communication overhead dominates
  • Sequence parallelism (Ulysses): split along sequence dimension

Pipeline Parallelism (PP)

  • Partition transformer layers across devices: device 0 = layers 0–7, etc.
  • Micro-batching: pipeline bubble reduced by processing B micro-batches simultaneously
  • Interleaved 1F1B schedule: forward-backward alternation; near-zero bubble
  • Communication: P2P activation passing between stages — lower BW requirement than TP
  • GPipe, PipeDream, Interleaved 1F1B — progressive refinements

3D Parallelism — Combining All Three

  • DP × TP × PP: e.g. 512 GPUs = DP=8, TP=8, PP=8 → fits 1T+ parameter models
  • NVLink for TP (high-BW, low-latency); InfiniBand for DP (high-BW); PCIe/NVLink for PP
  • Expert Parallelism (EP): route MoE experts to different devices; all-to-all dispatch/combine
  • GPT-4 / Gemini Ultra estimated: DP+TP+PP+EP combined across >1000 chips
16

CNN Accelerator Design

convolutional workloads — structure, tiling, and reuse

Convolution Loop Nest

for N, C_out, OH, OW, C_in, KH, KW:
  O[n,c_out,oh,ow] +=
    W[c_out,c_in,kh,kw] ×
    I[n, c_in, oh+kh, ow+kw]
  • 7 nested loops → rich tiling space for data reuse
  • im2col transform: maps conv to GEMM (at cost of memory expansion)
  • Winograd: F(2×2, 3×3) reduces MACs by 2.25× for 3×3 kernels
  • FFT-based conv: efficient for large kernels (>7×7)
  • Depthwise separable (MobileNet): factorises 3×3 → depthwise + pointwise; 8–9× fewer MACs

CNN Accelerator Design Choices

  • Line buffer: store H×W rows of input; slide window without re-reading DRAM
  • Double buffering: DMA fills buffer B while PE array consumes buffer A
  • Weight caching: small kernels (3×3×C_in×C_out) fit in on-chip SRAM; reuse across output tiles
  • Activation reuse: output-stationary on OH×OW tile; stream weights past
  • Fusion: Conv + BN + ReLU fused into single kernel — eliminates 2× DRAM round trips

Edge CNN ASIC Examples

Hailo-826 TOPS · 2.5W · dataflow arch
Kneron KL7304 TOPS · INT8 · <3W
Rockchip RK3588 NPU6 TOPS · INT8/INT4
NXP i.MX 9 eIQ4 TOPS · ethos-U65
Arm Ethos-U85up to 4 TOPS · 128–2,048 MACs/cycle (configurable)
STM32N6 NPU600 GOPS · Cortex-M55 + NPU
17

The AI Compiler Stack

bridging high-level ML frameworks to bare silicon

Training Frameworks
PyTorch, JAX, TensorFlow — user-facing Python; define computation graphs eagerly or traced
Graph Capture / IR
torch.compile (TorchDynamo + TorchInductor), XLA, MLIR — capture static computation graph; apply op fusion, constant folding, dead-code elimination
Optimisation Passes
Operator fusion · layout transforms (NCHW↔NHWC) · quantisation insertion · tiling/scheduling (Halide, Ansor) · memory planning
Backend Code Generation
NVCC / PTX for NVIDIA · XLA HLO → LLO for TPU · IREE for mobile · TVM Autoscheduler for edge
Runtime
Kernel launch, memory allocation, stream management, async execution, profiling hooks

MLIR — Multi-Level IR

  • Dialect system: Linalg, Tensor, SCF, Affine, LLVM dialects
  • Progressive lowering: high-level ops → tiled loops → memory-specific ops → ISA
  • Used by: XLA, IREE, Torch-MLIR, OpenAI Triton backend
  • Enables sharing of optimisation passes across hardware backends

NVIDIA CUDA Ecosystem

  • cuDNN: hand-tuned CNN kernels; Winograd, FFT, im2col automatically selected
  • cuBLAS/cuBLASLt: GEMM with Tensor Core support; epilogue fusion
  • CUTLASS: composable CUDA templates; custom GEMM variants
  • Triton: Python DSL for custom GPU kernels; generates near-cuBLAS performance
  • TensorRT: inference optimisation — layer fusion, precision calibration, engine caching

Inference Runtimes

  • vLLM: PagedAttention + continuous batching; de-facto LLM serving standard
  • SGLang: structured generation with radix tree prefix caching
  • ONNX Runtime: cross-platform; EP plugins for TRT, CoreML, QNN
  • llama.cpp: CPU+GPU; GGUF quantised models; ARM NEON + Metal backends
18

Power Efficiency & Performance per Watt

TOPS/W — the metric that matters for edge and datacenter alike

AcceleratorTOPS (INT8)TDP (W)TOPS/WSegment
Apple ANE (M4)38~5~7.6Mobile
Qualcomm HTP (SD Elite)75~8~9.4Mobile
Hailo-8L131.5~8.7Edge ASIC
Google TPU v5e393~170~2.3Cloud
NVIDIA H100 SXM1979700~2.8Cloud
NVIDIA B200 SXM45001000~4.5Cloud
AWS Trainium2~3800~600~6.3Cloud
Arm Ethos-U8540.08~50MCU/IoT

* Estimates from published data-sheets and characterisation reports. TOPS definition varies (dense vs sparse).

Power Breakdown in ML ASICs

  • MAC arrays: 40–60% of total power — scale with utilisation; DVFS critical
  • SRAM access: 25–35% — 6T SRAM reads dominate; scratchpad > cache energy
  • DRAM/HBM PHY: 15–25% — major target for compression and KV cache reduction
  • NoC/interconnect: 5–10% — scales with data movement; locality crucial
  • Control & misc: 3–5% — reduced in DSA vs GPU by removing OoO logic

Energy-Reduction Techniques

  • Clock gating: disable MAC array clocks when inactive; 20–40% dynamic power saved
  • Power gating: shut down entire engine island; longer wake latency
  • DVFS: scale V and f to minimum needed; quadratic V reduction
  • Near-threshold computing: V_DD → V_T; 10× energy reduction; area/speed penalty
  • Data compression: compress zero-heavy activations before DRAM transfer
  • Sparsity gating: detect zero activations early; skip MAC; save dynamic power
19

Case Study: NVIDIA Blackwell B200

the current apex of ML accelerator design (2024)

Architecture Highlights

Die configuration2× reticle dies; NV-HBI bridge
ProcessTSMC N4P (custom)
Total transistors208 billion
Die size2× 462 mm²
SMs192 per GPU (2 dies)
5th-gen Tensor CoresFP4 native; 2× H100 per SM
HBM3e192 GB · 8 TB/s
FP8 Tensor PFLOPS~9 (dense)
FP4 Tensor PFLOPS~10 dense / ~20 sparse
NVLink 51.8 TB/s bidirectional
TDP (SXM)1000 W

New Blackwell Features

  • 5th-gen Tensor Core: native FP4 (MXFP4); doubles throughput vs H100 FP8
  • 2nd-gen Transformer Engine: per-block micro-scaling (MX formats — MXFP8, MXFP4); higher accuracy at lower precision
  • RAS Engine: Reliability/Availability/Serviceability; scrubbing, error injection, live migration support
  • Secure AI: hardware TEE; confidential computing for multi-tenant inference
  • NVLink Switch System (NVL72): 72-GPU all-to-all fabric in a single rack; 1.4 exaFLOPS FP4 aggregate
  • GB200 NVL72: 36× Grace CPUs + 72× B200 GPUs; HBM3e → Grace via NVLink-C2C at 900 GB/s
Blackwell positions NVIDIA to serve frontier model inference: a single NVL72 rack can run a 1.8T-parameter MoE model at interactive latency — a capability that required a full data-centre building just two years prior.
20

Case Study: Apple Neural Engine (ANE)

efficiency-first on-device inference — 38 TOPS at <5W

ANE Architecture (M4 / A18)

  • Proprietary Apple ASIC integrated in SoC alongside CPU clusters and GPU
  • Unified Memory Architecture (UMA): CPU, GPU, ANE all share the same LPDDR5X pool — zero-copy tensor passing
  • 16-bit floating point native (BF16/FP16); INT8 quantisation via CoreML
  • Dedicated convolution and matrix-multiplication datapaths
  • Tightly coupled with CoreML compiler: model → ANEModel → compiled binary at install time
  • M4 ANE: 38 TOPS at <5W total — best TOPS/W in class

Unified Memory — Bandwidth Efficiency

  • M4 LPDDR5X: 120 GB/s shared across CPU, GPU, ANE, Media engines
  • No PCIe hop — GPU/ANE access host tensors at full memory bandwidth
  • Enables on-device Llama 3 8B inference at 30+ tok/s (quantised)
  • Trade-off: memory bandwidth shared across all engines — no dedicated HBM

On-Device LLM Inference Stack

  • Apple Intelligence (iOS 18 / macOS 15): on-device models 3B–7B parameters
  • Private Cloud Compute: larger models on Apple Silicon servers — cryptographic attestation before sending data
  • CoreML: stateless → stateful conversion; KV cache as CoreML state; eliminates re-allocation overhead
  • MLX framework: JAX-like on Apple Silicon; Metal backend; community LLM ecosystem
  • Speculative decoding: draft model on ANE; verify on GPU; 2–3× decode speedup

Evolution of TOPS

A11 (2017) — 1st ANE0.6 TOPS
A14 (2020)11 TOPS
A16 (2022)17 TOPS
A17 Pro (2023)35 TOPS
A18 Pro / M4 (2024)38 TOPS · 63× A11
M5 (Oct 2025)16-core NE, N3P; Apple Intelligence-class
21

Emerging Directions

what comes after the current generation

Processing-in-Memory (PIM)

  • Compute inside or adjacent to DRAM die — eliminates bus bottleneck
  • Samsung HBM-PIM: GEMV engine inside each HBM stack; 2× bandwidth for DNN inference
  • SK Hynix AiM: multiply-accumulate units in DRAM row decoders
  • Key challenge: limited compute density; no caches; programming model
  • Best fit: memory-bound ops — embedding lookup, GEMV decode
  • Prototype systems show 2.4× energy improvement for LLM decode

Analog / Resistive Computing

  • Analog in-memory compute (AIMC): weight stored as conductance; dot product via Kirchhoff's law — O(1) energy per MAC
  • Phase-change memory (PCM), ReRAM, MRAM: non-volatile weight storage — zero boot time
  • IBM PCM chip: 512 analogue columns; INT4 equivalent; 85 TOPS/W demonstrated
  • Challenges: non-idealities (drift, noise, write endurance); ADC/DAC overhead consumes gains
  • Likely path: hybrid digital-analog; analog for weight-stationary inference, digital for accum

Photonic & Neuromorphic

  • Photonic MAC: Mach-Zehnder interferometer arrays; perform GEMM at speed of light, zero dynamic power for multiply; Lightmatter Passage chip
  • Neuromorphic: Intel Loihi 2; event-driven spiking neural network; <1 mW for sparse workloads; not competitive with transformers yet
  • Wafer-scale: Cerebras WSE-3; 44 GB SRAM; eliminates HBM for moderate models
  • 3D die stacking: SRAM-on-logic (TSMC SoIC); put 192 MB SRAM directly atop logic die — bandwidth >100 TB/s
  • Timeline: 3D stacking in production now; photonic & analog: 5–10+ years to volume
22

Summary & Key Takeaways

AI / ML Accelerator Architectures — Presentation 05

Architecture Principles

  • Roofline model guides design: shift ridge point left (more BW) or up (more FLOPS); target workloads in compute-bound regime
  • Systolic arrays maximise data reuse in GEMM; weight/output/row-stationary dataflows trade off different memory resources
  • Dataflow engines (Groq, Cerebras) eliminate runtime scheduling overhead; compiler becomes the hardware
  • Flash Attention: algorithm-hardware co-design moves attention from O(N²) HBM to O(N) — critical for long context
  • Sparsity at multiple levels: activation (ReLU), weight pruning, structured (2:4), MoE

Numerical Formats & Precision

  • Training: BF16/FP16 + FP32 master weights (AMP); FP8 on H100+ with dynamic scaling
  • Inference: INT8 standard; INT4/FP4 for bandwidth-limited decode; FP4 on B200
  • QLoRA/GPTQ/AWQ: PTQ at INT4 with per-group scaling — <1% quality loss

System-Level Themes

  • Memory is the bottleneck — HBM, PagedAttention, KV cache optimisation
  • Parallelism at every level: DP/TP/PP/EP for training; disaggregated serving
  • Compiler stack (MLIR, XLA, TVM, Triton) as competitive differentiator
  • Edge NPUs: 10–50 TOPS/W; UMA avoids bandwidth bottleneck; on-device LLM now real
  • Next frontier: 3D SRAM stacking, PIM, photonic — closing the bandwidth gap

Series Navigation

✅ 01 SoC Packaging ✅ 02 On-Chip Interconnect 📋 03 Memory Hierarchy (next) 04 SerDes & I/O ▶ 05 AI/ML Accelerator Architectures 06 SoCs for ML Inference 07 Security in SoC … 13 Clocks & Resets