Advanced architecture: memory systems, accelerators, and the roofline in practice

Why this subject matters now

Processor microarchitecture has been the glamorous half of computer architecture for decades, but the binding constraints of the last ten years live elsewhere. Compute throughput per chip grew roughly 1000x between a 2010 CPU socket and a 2022 H100 doing bf16 matrix math, while DRAM latency barely moved and per-pin bandwidth improved by small integer factors. Everything interesting in modern systems design is a response to that widening gap, whether HBM stacks bonded onto silicon interposers, 50 MB last-level caches on GPUs, prefetchers of increasing ambition, tensor units that extract quadratic compute from linear data movement, or precision reductions from fp32 to fp8 that exist mostly to move fewer bytes. A systems practitioner today is expected to hold the whole memory-side story quantitatively, to know why a random pointer chase on this machine costs 134.6 ns while sequential streaming costs 4.5 ns per line, why one thread of that chase can never exceed half a gigabyte per second no matter how fast the core is, why an H100 needs on the order of a megabyte of loads in flight to keep its memory system busy, and why a matmul at 25 percent occupancy can outrun one at 100 percent.

The subject also has a second audience now. Training and serving large models is a memory-system and interconnect problem before it is a compute problem, and the people who make clusters fast spend their days with exactly the tools this page works through, rooflines calibrated with measured bandwidth, occupancy and register-pressure arithmetic, tensor-core utilization counters, NVLink bus-bandwidth accounting, and the precision ladder. The theory here is anchored on Hennessy and Patterson's quantitative treatment, on Jacob, Ng, and Wang's memory-systems text, on Williams, Waterman, and Patterson's roofline paper, and on the accelerator dataflow literature that runs from Chen, Emer, and Sze's Eyeriss to Jouppi's TPU. The numbers are anchored on this machine, an Intel Xeon Platinum 8480+ guest with 52 vCPUs and 442 GiB of RAM, and two NVIDIA H100 80GB HBM3 GPUs joined by 18 NVLink links. Where the machine hides something behind virtualization, the page says so and measures what it can.

Inside the DRAM chip

Banks, rows, and columns

A DRAM die is a two-dimensional grid of capacitors, each storing one bit as charge, each read destructively. The grid is organized as banks, and each bank as rows and columns. A modern DDR5 x8 device of 16 Gb has 32 banks arranged as 8 bank groups of 4 banks, and each bank's row is 1 KiB wide on an x8 part. A 32-bit DDR5 subchannel built from four x8 devices therefore exposes an effective 4 KiB row. The bank is the unit of parallelism inside the chip. Different banks can be working on different rows at once, which is the property every memory-controller scheduler exploits.

Access is a three-act protocol, and the acts have names that appear on every datasheet. ACTIVATE (ACT) reads an entire row out of the array into the bank's row buffer, a row of sense amplifiers. Because reading the capacitors drains them, the sense amplifiers also restore the charge, which is why a row must stay open a minimum time (\(t_{RAS}\)) before it may be closed. READ or WRITE (the CAS command) then moves a column, one burst of data, between the row buffer and the pins. PRECHARGE (PRE) writes the row back and returns the bitlines to their reference voltage so a different row can be activated. Three timing parameters govern the acts. \(t_{RCD}\) is the delay from ACT to the first CAS, \(t_{CL}\) (also written CAS latency) is the delay from READ to the first data beat, and \(t_{RP}\) is the delay from PRE until the next ACT may issue. They are quoted in controller clock cycles on module labels, which is why a label like 40-39-39 means nothing until multiplied by the clock period.

one DDR5 bank (of 32 per device)

            columns ─────────────────────►
   rows │  ┌──────────────────────────────┐
        │  │        capacitor array       │
        ▼  │  (one row = 1 KiB on an x8)   │
           │  ..............................│
           │  ███ activated row ███████████│ ─┐ ACT: row → sense amps
           └──────────────────────────────┘  │ (destructive, so also restore)
           ┌──────────────────────────────┐ ◄┘
           │  row buffer (sense amps)      │ ── CAS: one column burst ─► pins
           └──────────────────────────────┘
                     │
                     └─ PRE: restore + equalize bitlines, row now closed

timing:  ACT ──t_RCD──► READ ──t_CL──► data (burst)
         PRE ──t_RP───► next ACT       t_RAS: min ACT→PRE

The row buffer is a cache of exactly one row per bank, and it creates the three-case latency structure that dominates DRAM behavior. A row-buffer hit (the wanted row is already open) pays only \(t_{CL}\) plus the burst. A closed bank (no row open) pays \(t_{RCD} + t_{CL}\). A row-buffer conflict (a different row is open) pays the full \(t_{RP} + t_{RCD} + t_{CL}\). Memory controllers reorder requests aggressively to convert conflicts into hits, batching requests to the same row before closing it. The scheduling problem was formalized by Rixner, Dally, and colleagues, whose FR-FCFS (first-ready, first-come-first-served) policy, prioritize row hits then oldest, is still the baseline every DRAM scheduling paper compares against. Row-buffer locality is thus a fourth kind of locality, invisible to the cache hierarchy. Two misses that land in the same DRAM row are far cheaper than two that alternate between rows of one bank.

Refresh, the tax on every capacitor

The charge on a DRAM cell leaks, so every row must be rewritten periodically, 32 or 64 ms retention at normal temperature depending on generation and density. The device handles this through REFRESH commands issued by the controller every \(t_{REFI}\) (refresh interval). Each command occupies the device for \(t_{RFC}\) (refresh cycle time), during which the refreshing bank set can serve nothing. For a 16 Gb DDR5 device the standard values are \(t_{REFI} = 3.9\ \mu\text{s}\) and \(t_{RFC} = 295\) ns for all-bank refresh.

DDR5 versus HBM3: two answers to the pin problem

Bus width, stacking, and what each interface is for

A DRAM interface is pins times per-pin rate. DDR5 keeps the pin count small and civilized because its signals must survive centimeters of motherboard trace and a socketed connector. A DIMM presents two independent 32-bit subchannels, 64 data bits total, at 4800 to 8800 MT/s. One DDR5-4800 DIMM therefore moves \(64 \times 4.8\,\text{Gb/s} / 8 = 38.4\) GB/s, and a server socket gets to hundreds of GB/s only by replicating channels. The Xeon 8480+ under this page has 8 memory channels for a theoretical \(8 \times 38.4 = 307\) GB/s per socket. The virtue of this design is capacity and serviceability, terabytes per socket of replaceable, individually cheap modules.

HBM inverts every one of those choices. Instead of few pins driven fast over long traces, it uses an enormous number of pins driven moderately over sub-millimeter wires. DRAM dies are thinned, stacked 8 or 12 high, connected vertically by through-silicon vias (TSVs), and the stack sits on a silicon interposer next to the GPU die, close enough that a 1024-bit bus per stack is routable. An HBM3 stack exposes 16 independent 64-bit channels (each split into two 32-bit pseudo-channels), and at HBM3 rates of up to 6.4 Gb/s per pin one stack delivers up to 819 GB/s. The price is that the memory is soldered forever to the package, capacity per stack is modest, and the interposer packaging (TSMC's CoWoS and its competitors) is expensive and supply-constrained. The reward is an order of magnitude more bandwidth per package than any DIMM system. In short, DDR gives capacity and replaceability at hundreds of GB/s, and HBM gives thousands of GB/s at fixed, modest capacity.

DDR5 (few fast pins, far away)        HBM3 (many moderate pins, adjacent)

 CPU ──── cm of PCB trace ──── DIMM     ┌────────┐  ┌─ DRAM die ─┐
  8 channels × 64 bits                  │  GPU   │  ├────────────┤ ×8-12 dies
  4.8-8.8 Gb/s per pin                  │  die   │  ├────────────┤ stacked on
  38.4 GB/s per DIMM                    └───┬────┘  ├────────────┤ TSVs
  replaceable, TB-scale capacity            │ 1024 bits └──┬─────┘
                                        ════╪══ silicon ═══╪════ interposer
 H100: 5 stacks × 1024 = 5120 pins          └── < 1 mm ────┘
        at 5.24 Gb/s per pin

Why HBM trades latency for bandwidth

A common misreading is that HBM is "faster memory." At the device level it is not. The banks inside an HBM stack are ordinary DRAM with \(t_{RCD}\), \(t_{CL}\), and \(t_{RP}\) in the same 14-18 ns range as a DIMM, because the limiting physics, sense-amplifier settling on long bitlines, is the same. End to end, a GPU's memory latency is actually several times worse than a CPU's, at hundreds of nanoseconds, as microbenchmark studies of NVIDIA GPUs (Jia et al.'s dissection series is the standard reference) consistently measure, versus the 134.6 ns this host's CPU achieves. The latency is spent in the GPU's deep queueing. Requests from 132 SMs funnel through a crossbar into dozens of memory-controller queues sized to keep 5120 pins busy, and a deep queue is precisely a latency amplifier. The design is coherent, not contradictory. A GPU is built to hide latency with parallelism (the next section quantifies how much), so it happily accepts worse latency in exchange for an interface that moves ten times the bytes. A CPU makes the opposite bet for the opposite reason. The two memory systems are not on a better-and-worse axis. They are the two endpoints of Little's law.

Memory-level parallelism and Little's law

The law, derived, and what it says about memory

Little's law is queueing theory's one free lunch. In any system in steady state, the average number of items inside equals the arrival rate times the average time each item spends inside, \(N = \lambda \, W\), with no assumptions about distributions, ordering, or independence (Little, 1961). Apply it to a memory system. If a machine is to sustain bandwidth \(B\) bytes per second against a memory latency of \(W\) seconds, and each in-flight request carries \(g\) bytes (one cache line or sector), then the number of requests simultaneously outstanding must average

The product \(B \times W\), the bytes that must be in flight to hide the latency, is the memory system's bandwidth-delay product, the exact analogue of the network quantity of the same name. This single equation organizes everything on this page. Prefetchers, out-of-order windows, miss-status-holding registers, GPU occupancy, and unrolled loads are all mechanisms for raising \(N\). The derivation is worth doing once because it is a counting argument, not a stochastic one.

Measuring the concurrency limit directly on this machine

The theory predicts that bandwidth from one thread should scale linearly with the number of independent pointer chases it interleaves, until some hardware resource that tracks outstanding misses saturates. A microbenchmark that runs \(k\) independent Sattolo-cycle chases over a 2 GiB region from a single thread, each chase contributing exactly one line in flight, shows the plateau directly. Scaling is close to linear for the first several chains, then collapses as requests queue behind a full fill-buffer structure. The plateau bandwidth divided into \(g/W\) recovers the number of miss-tracking entries the core exposes, on the order of 10-16 across recent Intel cores. This is why single-thread random-access workloads, pointer-heavy graph traversals, hash joins, garbage collectors, cannot come close to memory bandwidth, and why the same 4 KiB-page chase costs 255.0 ns against 134.6 ns with huge pages (os.json). With a 32 GiB working set, the 4 KiB TLB misses add a page walk to nearly every hop, and the walk itself is more dependent memory traffic.

The memory hierarchy as a bandwidth cascade

The single most useful mental model for a modern accelerator is not the classic latency pyramid but a bandwidth cascade. Each level down the hierarchy delivers roughly an order of magnitude fewer bytes per second than the one above, and the entire craft of a fast kernel is keeping the working set as high in the cascade as it will fit. Assembling one table for this machine, with each row labeled measured, derived, or spec, makes the drop concrete. The register and shared-memory rows are derived from the SM's issue arithmetic (shared memory serves 32 banks of 4 bytes per clock per SM, so \(32 \times 4 \times 1.98\,\text{GHz} \times 132 \approx 33\) TB/s aggregate). The L2 row is the whitepaper's modeled figure, and HBM, NVLink, and PCIe are measured in this repository's JSON.

The cascade spans nearly four orders of magnitude, from >100 TB/s in the register file to ~50 GB/s at the NIC. Three cliffs matter most. The first is SMEM to HBM, roughly 10x, which is the entire reason for tiling. A matmul or attention kernel that streams operands from HBM is bandwidth-bound, whereas one that loads a tile into shared memory and reuses it is not (the parallel-computing page works the tiled-matmul and flash-attention kernels that do this). The second is HBM to NVLink, about 9x, which is why data-parallel training overlaps gradient all-reduce with the backward pass rather than serializing it. The third is NVLink to PCIe, about 6x, which is why an offloading scheme that spills optimizer state to host memory is throttled at PCIe rates no matter how fast the GPU is. Every serious performance decision in this stack is a choice about which cliff a given byte has to cross and how many times.

Prefetching: manufacturing concurrency from patterns

Stream, stride, and spatial prefetchers

A prefetcher is hardware that converts predictability into memory-level parallelism. The taxonomy, surveyed thoroughly by Mittal (2016), runs from simple to speculative. Next-line and stream prefetchers detect monotone sequences of miss addresses and run ahead by a configurable depth. They are why the sequential chase achieves 4.5 ns per line. Stride prefetchers track, per load instruction (indexed by PC), the difference between successive addresses, and issue ahead when the delta is stable, covering column-major walks and strided array code. Spatial-pattern prefetchers learn which offsets within a region tend to be touched together and replay the bit pattern on first touch of a new region. Beyond these sit correlation and temporal-stream prefetchers that record miss-address sequences and replay them, effectively memoizing pointer chases, at substantial metadata cost. Every prefetcher trades accuracy against coverage and timeliness. A wrong prefetch wastes bandwidth and evicts useful lines, so real designs throttle themselves when accuracy drops, and all of them are defeated by a Sattolo cycle, which is constructed to have no pattern at any granularity.

Why strided access decays: the measured table

The GPU makes the cost of non-unit stride unusually clean to see, because its memory system is explicitly organized around 128-byte cache lines fetched as four 32-byte sectors. The measured strided-read bandwidths on this repository's H100 (parallel.json, idle GPU, useful bytes per second) appear in the table below.

The limiting arithmetic is simple. A stride of \(s\) 4-byte floats means each 32-byte sector fetched from HBM contains at most \(\max(1, 8/s)\) useful elements, so once \(s \geq 8\) each fetched sector carries a single useful float and the useful-bandwidth waste factor saturates at 8. The measured asymptote of \(2380.5 / 325.2 = 7.3\) agrees, within noise of the predicted 8. At small strides the decay is gentler than the naive \(1/s\) because the L2 cache recaptures sectors that a neighboring warp will use shortly, and at stride 16 versus 8 the further drop comes from touching twice as many 128-byte lines (and DRAM rows) per useful byte even though sector waste has already saturated. The CUDA tab in the implementation section reproduces the shape of this table. The general lesson transfers directly to CPUs. Hardware moves lines and sectors, so the effective bandwidth of any access pattern is pin bandwidth times the fraction of each fetched line the program actually uses. Array-of-structures layouts where a loop touches one field are stride-\(s\) reads in disguise, which is the entire case for structure-of-arrays layouts in performance code.

NUMA on this machine, honestly

NUMA, non-uniform memory access, is what happens when "memory" stops being one place. Each socket (or, with sub-NUMA clustering, each quadrant of a socket) has local DRAM channels, and touching another node's memory crosses a socket-to-socket link (UPI on Intel) at higher latency and lower bandwidth. This host is the right place to demonstrate epistemic honesty rather than the effect itself, because it is a KVM guest, and the hypervisor flattens whatever physical topology exists. The guest exposes a single fictitious socket with a single NUMA node covering all 52 vCPUs and a synthetic 16 MiB L3, whereas a physical Xeon Platinum 8480+ has 56 cores and about 105 MB of L3, so this guest is a carve-out of something larger presented as one node. No cross-node measurement is possible from in here, which is itself the most practical NUMA lesson this machine can teach. Measured latency here already includes whatever remote-socket penalty the hypervisor's placement imposes, invisibly and variably.

On bare metal the effects are large and well documented (Hennessy and Patterson treat the architecture). Local versus remote DRAM latency on two-socket servers typically differs by 1.5-2x (order of 100 ns local, 150-200 ns remote), remote bandwidth is capped by the UPI links rather than the remote DIMMs, and Sapphire Rapids adds sub-NUMA clustering (SNC), which splits even one socket into four NUMA domains so that each quadrant's cores favor their nearest memory controllers. The operating system's default policy, first touch, allocates a page on the node of the thread that first writes it, which is why parallel programs that initialize arrays from thread 0 and then process them from all threads quietly serialize every access through one node's controllers. The fixes are mechanical. Initialize in parallel with the same partitioning as the compute loop, pin threads, and bind memory where the topology is real.

Interconnects: PCIe gen5, NVLink, CXL

PCIe lane arithmetic against a measured copy

PCIe gen5 signals at 32 GT/s per lane per direction with 128b/130b encoding, so an x16 link's raw data rate is

before packet overheads (TLP headers, flow control) shave several more percent. Both H100s here report gen5 x16 links. The measured host-to-device copy on this machine (os.json, 256 MiB buffers, torch) is 54.84 GB/s from pinned memory, 87 percent of the lane arithmetic and about as good as PCIe copies get, but only 21.01 GB/s from ordinary pageable memory, a measured 2.61x penalty, because a pageable copy is really host-side staging into a pinned bounce buffer, and the staging memcpy steals the bandwidth. Pinning itself cost a measured 95.7 ms per 256 MiB, roughly 25 copies' worth, which is the arithmetic behind the standard advice. Pin once, reuse the buffer, overlap copies with compute, and never pin inside a loop.

NVLink, measured with NCCL, and the coherence question

The two H100s are joined by NVLink, and nvidia-smi reports 18 links (the NV18 topology), for a raw \(18 \times 26.562 = 478.1\) GB/s per direction, marketed as 450 GB/s effective after protocol framing. What a training job actually experiences is the collective bandwidth, and this repository measured it (networks.json, torch.distributed all_reduce over NCCL 2.26.2). A raw peer-to-peer device copy sustains 388.28 GB/s, and all-reduce bus bandwidth climbs from 41.68 GB/s at 1 MiB messages to 300.14 GB/s at 256 MiB and 330.15 GB/s at 1 GiB. Small messages are latency-dominated. At 1 MiB the transfer spends its time in kernel launches and synchronization, achieving 12 percent of what 1 GiB achieves. For two ranks bus bandwidth equals algorithmic bandwidth (the \(2(N-1)/N\) ring factor is 1 at \(N = 2\)), so 330.15 GB/s is the honest number, 73 percent of the 450 GB/s link rating, with the gap going to protocol, the reduction arithmetic itself, and the fact that a ring all-reduce must both read and write HBM at each step. All of it is an order of magnitude above the 54.84 GB/s PCIe ceiling that the same data would face without NVLink, which is the entire reason the links exist.

The coherence distinction is worth stating precisely, because it separates NVLink from PCIe more than raw bandwidth does. Plain PCIe is non-coherent. A device DMA into host memory does not participate in the CPU's cache-coherence protocol, so software must flush and fence explicitly, and there is no notion of the GPU and CPU sharing a cache line. NVLink between GPUs, and the NVLink-C2C link that binds a Grace CPU to a Hopper GPU, carry coherence traffic. A load on one side can observe a store on the other with hardware-maintained ordering, which is what makes unified memory and fine-grained peer access practical rather than a correctness hazard. CXL brings the same idea to commodity hardware. CXL.cache lets a device coherently cache host memory and CXL.mem lets the host issue plain loads and stores to device-attached DRAM at a latency comparable to one extra NUMA hop. CXL.mem is the one reshaping memory systems, because it turns DRAM into a pooled, fabric-attached resource. A CXL expander is, in the language of the NUMA section, a memory node with no CPUs on it, addressable with ordinary load/store semantics.

Tensor cores and systolic arrays

Everything so far has been about moving bytes. The reason the byte-moving matters so much is that a modern accelerator's compute unit is so voracious that only a carefully amortized memory system can keep it busy. The compute unit in question is a two-dimensional array of multiply-accumulate (MAC) cells. NVIDIA calls its version a tensor core, Google's TPU is one large version of it, and the design principle is the systolic array that H. T. Kung and Charles Leiserson described in 1978. This section derives why such an array amortizes bandwidth, works the reuse arithmetic of the three canonical dataflows, and then derives the arithmetic intensity a tensor-core array demands, connecting it to the roofline.

The systolic dataflow, derived

Consider computing \(C = AB\) with \(A\) of shape \(M \times K\) and \(B\) of shape \(K \times N\). A systolic array is a \(P \times P\) grid of MAC cells, each holding one accumulator. In the weight-stationary formulation, one \(P \times P\) tile of \(B\) is preloaded so that cell \((i,j)\) holds \(B_{ij}\) and never moves. Rows of \(A\) then march in from the left edge, one column of the array per clock, while partial sums march down from the top. At each clock, cell \((i,j)\) receives an input element from its left neighbor and a partial sum from above, computes \(\text{psum} + a \cdot B_{ij}\), passes the input right and the new partial sum down. After the operands have rippled through, the bottom edge emits one row of \(C\) per clock. The name systolic is Kung's. Data pulses through the array like blood through a heart, each cell touching each datum once.

weight-stationary P×P systolic array (P = 4 shown)

  inputs (rows of A) stream in ►      B tile preloaded, stationary
        a30 a20 a10 a00 ─►┌────┬────┬────┬────┐
        a31 a21 a11 a01 ─►│B00 │B01 │B02 │B03 │  each cell: psum += a·B
        a32 a22 a12 a02 ─►│B10 │B11 │B12 │B13 │  pass a right, psum down
        a33 a23 a13 a03 ─►│B20 │B21 │B22 │B23 │
                          │B30 │B31 │B32 │B33 │
                          └─┬──┴─┬──┴─┬──┴─┬──┘
                            ▼    ▼    ▼    ▼
                          columns of C emerge (partial sums)

  P² MACs per clock, fed by ~P new inputs per clock  ⇒ operand reuse ≈ P/2

The bandwidth argument is now immediate and is the whole point of the design. The array performs \(P^2\) MACs every clock, one per cell. It is fed by roughly \(P\) new input elements streaming in the left edge each clock (the weights were loaded once and amortize over the whole tile pass, and the partial sums are internal). The ratio of compute to fresh operands is therefore

counting both the input stream and the eventual output drain as operand traffic. Each operand brought to the array edge participates in \(\Theta(P)\) MACs before it leaves. This is the amortization. A naive matmul that read both operands of every MAC from memory would need two operand fetches per MAC, an operand-reuse of \(\tfrac12\). The array raises that by a factor of \(P\). For the TPUv1's \(256 \times 256\) array that is a reuse of 128, and for a Hopper tensor core doing a warp-level \(16 \times 16 \times 16\) MMA the reuse per issued instruction is on the same order. The MAC array does not make arithmetic cheaper. It makes each byte that crosses the memory interface do \(\Theta(P)\) times as much arithmetic, which is exactly what a bandwidth-starved system needs.

The dataflow taxonomy: weight-, output-, and row-stationary

Which datum is held stationary in the array, and which stream, is the defining design choice of a spatial accelerator, and Chen, Emer, and Sze's Eyeriss work (ISCA 2016) gave the taxonomy its standard names by classifying designs according to which data type is kept local to maximize its reuse. The three canonical choices are stated below for a convolution or matmul with weights \(W\), inputs (activations) \(I\), and partial sums \(P\!s\).

• Weight-stationary keeps weights pinned in each PE's register and streams activations past them, maximizing weight reuse. It is what the TPU and most inference accelerators use, because in inference the weights are reused across every element of a large batch or feature map while activations are seen once.
• Output-stationary keeps each partial sum pinned in its PE and streams weights and activations in, maximizing partial-sum reuse by never spilling an accumulator to memory until the reduction over \(K\) is complete. It minimizes the most expensive traffic, the read-modify-write of partial sums, which is why deep reductions favor it.
• Row-stationary, the dataflow Eyeriss introduced, keeps a row of the convolution's computation, a one-dimensional weight row and the corresponding input row, resident in each PE, and reuses all three data types at once by mapping the 2D convolution's reuse structure onto the 2D PE array. Chen et al. showed it minimizes total data movement energy across the whole array rather than optimizing any single data type, which under a realistic energy model is what actually matters.

The reason the taxonomy is worth memorizing is that it maps directly onto an energy model. Chen and Sze report, and it is the number every architect carries, that in a 65 nm process a single 32-bit DRAM access costs roughly 200 times the energy of one MAC, an on-chip global-buffer (SRAM) access about 6 times, a neighbor or array access about 2 times, and a local register-file access about 1 time. Data movement, not arithmetic, dominates the energy budget, so the winning dataflow is the one that keeps each datum in the cheapest level for the longest, and that is a per-workload question. The next problem quantifies exactly how much a reuse dataflow saves.

How much arithmetic intensity keeps the array fed

The array amortizes bandwidth, but only if the kernel supplies enough arithmetic per byte. The precise requirement is the roofline ridge point. Define arithmetic intensity \(I = \text{FLOPs} / \text{bytes moved from HBM}\) (units FLOP/byte). The roofline model (Williams, Waterman, Patterson, 2009) says attainable performance is

where \(\pi\) is peak compute and \(\beta\) is peak memory bandwidth. The two regimes meet at the ridge point \(I^\star = \pi / \beta\). Below it a kernel is memory-bound and above it compute-bound. For this H100, using the measured bf16 tensor-core peak of 744.6 TFLOP/s (h100.json, \(n = 4096\)) and the pin bandwidth 3.35 TB/s,

and using the whitepaper's dense bf16 peak of 989.5 TFLOP/s it is 295 FLOP/byte. Either way the number is large. To keep the tensor cores saturated, a kernel must do on the order of 220-300 floating-point operations for every byte it reads from HBM. In bf16, two bytes per element, that is 440-590 FLOPs per element loaded. No elementwise or reduction kernel comes close. Only matmul-shaped work, whose intensity grows with tile size, can. This is the arithmetic that forces the whole tensor-core design. The array is worth building only because matmul's intensity can exceed the ridge, and it is worth feeding from a bandwidth cascade only because the ridge is so high.

The precision ladder: fp32, tf32, bf16, fp8

The tensor core is not one datapath but a ladder of them, each trading numerical range or precision for throughput and, just as importantly, for bytes moved. Every rung halves or doubles something, and the measured throughputs on this machine (h100.json, square matmul at \(n = 4096\)) make the ladder concrete.

The design of each rung follows a single idea. For neural-network math the exponent range matters more than the mantissa. bf16 keeps fp32's full 8-bit exponent and throws away 16 mantissa bits, so it never overflows or underflows where fp32 would not, which is why it trains large models stably where fp16, with its 5-bit exponent, needs loss scaling to keep gradients in range (Micikevicius et al., 2018). tf32 is the subtle one. It keeps the 8-bit exponent and 10 mantissa bits (19 bits total), is stored in a 32-bit container so code sees fp32 tensors, but the tensor core internally rounds the inputs to tf32 before multiplying, which is how a matmul that looks like fp32 runs 7.5x faster on this card. fp8, standardized by Micikevicius et al. (2022) in two variants, E4M3 for weights and activations and E5M2 for gradients, halves bytes again and doubles the spec throughput, and it is now the default for large-model inference and increasingly for training with per-tensor or per-block scaling to place the tiny dynamic range where the values live. The through-line for this page is that every rung down the ladder is a memory-system decision as much as a compute one. fp8 doubles tensor-core throughput, but it also halves the bytes read from HBM against bf16, which for a bandwidth-bound stage such as attention's value load or an embedding gather is the larger win.

The accuracy cost is real but bounded, and the standard reference on why it is bounded is the observation that matmul accumulation should happen in higher precision than the inputs. Markidis et al. (2018) measured that NVIDIA tensor cores multiply fp16 inputs but accumulate in fp32, so the error of a length-\(K\) dot product grows like the fp16 rounding of the inputs plus fp32 accumulation error, far better than pure fp16 would give. The same split-accumulate principle carries to bf16 and fp8. The practical rule that falls out is to store and multiply in the lowest precision the range tolerates, accumulate in fp32, and keep a master copy of weights in fp32 or bf16 for the optimizer.

Structured sparsity and quantization hardware

Two further hardware features extend the ladder past dense low precision. The first is structured 2:4 sparsity, introduced with the A100's Sparse Tensor Cores and carried into Hopper. The constraint is deliberately rigid. In every contiguous group of four weights, at most two may be nonzero. A network is trained dense, pruned to satisfy the 2:4 pattern, and fine-tuned. Mishra et al. (2021) showed this recovers full accuracy across a wide range of models. The hardware payoff is that the sparse tensor core stores only the two nonzeros per group plus a 2-bit index selecting their positions, so it skips half the MACs and reads half the weight bytes, delivering up to a nominal 2x over the dense rung. The rigidity is the point. Unstructured sparsity gives no speedup on such a MAC array because the array cannot skip arbitrary zeros without gather logic that would cost more than it saves, whereas the fixed 2:4 pattern maps to a small multiplexer in front of each PE row. Sparsity here is a memory-bandwidth technique wearing a compute costume. Half the weight bytes cross the cascade, which for the weight-bound decode phase of LLM inference is exactly the bytes that matter.

The second is native quantization support, integer tensor cores (int8, int4) for inference, and the fp8 scaling machinery just described. The hardware provides fast per-tensor and, increasingly, per-block scale application so that a low-precision matmul can be dequantized in the epilogue at negligible cost. The systems consequence, again, is bytes. An int8 or fp8 weight is one byte, a quarter of an fp32 weight, so a 70B-parameter model that needs 280 GB in fp32 fits in 70 GB in int8 and streams from HBM four times faster, which is why quantized inference is bandwidth-bound-friendly and why the accelerator landscape below competes so hard on low-precision throughput and on-package memory.

The roofline, applied end to end

The roofline is only useful when a real kernel is placed on it with measured coordinates. Do that for three kernels on this H100, using the measured effective HBM bandwidth of 3063.5 GB/s as \(\beta\) and the measured bf16 tensor-core peak of 744.6 TFLOP/s as \(\pi\), so the ridge sits at \(I^\star = 744.6\times10^{12} / 3.0635\times10^{12} = 243\) FLOP/byte.

Each row tells the same story from a different corner of the plane. Vector add has intensity \(1/12\), one add per three fp32 accesses of 4 bytes. Its roofline cap is \(\beta I = 3063.5 \times 10^{9} \times 0.083 = 255\) GFLOP/s, a factor of 2900 below the tensor-core peak, so the only question about a vector add is whether it reaches memory bandwidth, and at 2988 GB/s it does, 97 percent of copy peak. The bf16 \(n = 2048\) matmul sits far to the right of the ridge, so its ceiling is the compute roof, and its measured 453.4 TFLOP/s is 61 percent of that ceiling, with the gap being tile efficiency at that middling size rather than any bandwidth problem. The two attention rows are the most instructive, because they are the same math at two intensities. Naive attention writes the entire \(L \times L\) score matrix to HBM and reads it back, which for \(L = 2048\) is 1.07 GB of traffic (h100.json) that pins it at memory-bound 25.9 TFLOP/s, whereas flash attention keeps the scores in SRAM and never spills them, raising the intensity enough to reach 460.0 TFLOP/s, a measured 17.8x. Flash attention is a roofline move. It does not change the FLOPs, it changes the bytes, and by changing the bytes it slides the kernel from the memory roof to the compute roof. The parallel-computing page derives and benchmarks the kernel itself. Here the point is that the roofline predicted the win before a line of it was written.

Reading a profile

A profiler turns the roofline into three numbers a practitioner reads in order. The first is achieved memory bandwidth, reported by Nsight Compute as a percentage of peak DRAM throughput (dram__throughput.avg.pct_of_peak_sustained_elapsed) and by the vendors' bandwidth counters. For a memory-bound kernel this is the whole story. The vector add above at 97 percent of copy peak is done, and no restructuring will help. The second is occupancy, the ratio of resident warps to the hardware maximum, which sets \(N\) in Little's law and therefore whether the kernel supplies enough outstanding requests to hide latency. Occupancy is bounded by whichever resource runs out first, registers per thread, shared memory per block, or block slots. The register-pressure arithmetic is exact and worth doing, so it is the next problem. The third is tensor-core utilization (sm__pipe_tensor_op_hmma.avg.pct_of_peak_sustained_active on recent tools), which for a matmul-shaped kernel says what fraction of the MMA pipeline's cycles issued a tensor instruction. A compute-bound matmul stuck at 40 percent tensor utilization is starved somewhere upstream, in the shared-memory load pipeline or the operand-collector, not at the array. The diagnostic discipline is to read them in that order. If bandwidth is saturated, stop. Otherwise, if occupancy is low, find the limiting resource, and if tensor utilization is low, look at the feeding pipeline.

The accelerator landscape

The systolic array and the memory cascade are the two axes along which every AI accelerator is positioned, and reading the landscape as points in that plane makes the field legible. The attribution matters, so each is named with the group that built it.

• Google TPU (Jouppi et al., ISCA 2017) is the reference systolic machine, a single large weight-stationary \(256 \times 256\) MAC array in the v1, fed by a large software-managed on-chip buffer, with the explicit design thesis that a big deterministic array beats a cache-heavy CPU or GPU on inference energy per operation. Later TPU generations added bf16 training, HBM, and the inter-chip optical interconnect that makes TPU pods a single-fabric alternative to NVLink-plus-InfiniBand clusters. The TPU is the datapoint that proved the systolic bet at datacenter scale.
• Cerebras (Lauterbach and colleagues, IEEE Micro 2021) takes the opposite view of the memory cascade, building the entire accelerator on one wafer so the SRAM and the compute never cross a package boundary. The Wafer-Scale Engine puts hundreds of thousands of cores and tens of gigabytes of on-wafer SRAM on a single die-sized chip, eliminating the HBM rung entirely for models that fit and collapsing the cascade's most expensive cliff into on-chip bandwidth. The tradeoff is capacity. What does not fit in wafer SRAM must stream from an external memory service.
• Groq (Abts et al., ISCA 2020) attacks latency and determinism with a tensor-streaming processor that has no caches and no dynamic scheduling. The compiler places every operation on a statically scheduled datapath so that latency is a compile-time constant. It is the accelerator built for the low-latency inference corner, trading the GPU's flexibility for a memory system whose timing is known exactly.
• AWS Trainium (Annapurna Labs, presented at Hot Chips) is a systolic training accelerator paired with high-bandwidth memory and a custom collective interconnect, positioned as a cost-per-training-token alternative to GPUs within one cloud. It is the datapoint that the weight-stationary array plus HBM plus a good collective fabric is now a reproducible recipe rather than a research artifact.
• Tenstorrent (Vasiljevic et al., IEEE Micro 2021) builds a grid of small general-purpose cores each with a matrix engine, connected by an on-chip network, and leans on conditional execution and a software 2.0 compiler rather than one monolithic array, betting that a programmable mesh generalizes better than a fixed systolic array as model architectures churn.

The common thread is that none of these competes primarily on peak MAC count. They compete on where in the bandwidth cascade the working set lives and on how deterministically the array can be fed. Cerebras removes the HBM rung, Groq removes the cache nondeterminism, the TPU and Trainium optimize the array-plus-HBM-plus- collective recipe, and Tenstorrent trades a bit of array efficiency for programmability. Every one of them is a different answer to the same question this page has been asking, how to keep a voracious MAC array fed across an increasingly steep memory cascade.

Implementation

Three self-contained microbenchmarks make the page's claims runnable. The first is a CUDA kernel that reproduces the strided-read bandwidth collapse of the prefetching section, reading one float every stride elements and reporting useful bandwidth. Because a strided load wastes the unused bytes of each 32-byte sector, useful bandwidth falls toward one eighth of the stride-1 rate as the stride passes 8, the shape measured in parallel.json.

// strided_read.cu  --  measure useful bandwidth vs stride on an H100
// nvcc -arch=sm_90 -O3 strided_read.cu -o strided_read
#include <cstdio>
#include <cstdint>
#include <cuda_runtime.h>

// each thread reads N/stride elements, stride floats apart, and accumulates.
// the accumulate keeps the compiler from deleting the loads.
__global__ void strided_sum(const float* __restrict__ x, float* __restrict__ out,
                            int64_t n, int stride) {
  int64_t tid = blockIdx.x * (int64_t)blockDim.x + threadIdx.x;
  int64_t nthreads = gridDim.x * (int64_t)blockDim.x;
  float acc = 0.0f;
  for (int64_t i = tid * stride; i < n; i += nthreads * (int64_t)stride) {
    acc += x[i];
  }
  // one write per thread; negligible next to the reads
  out[tid % 1024] = acc;
}

int main() {
  const int64_t n = 1LL << 28;           // 256 M floats = 1 GiB
  float *x, *out;
  cudaMalloc(&x, n * sizeof(float));
  cudaMalloc(&out, 1024 * sizeof(float));
  cudaMemset(x, 1, n * sizeof(float));

  const int strides[] = {1, 2, 4, 8, 16, 32};
  cudaEvent_t a, b; cudaEventCreate(&a); cudaEventCreate(&b);

  for (int s : strides) {
    int64_t touched = n / s;               // elements actually read
    dim3 block(256), grid(4096);
    strided_sum<<<grid, block>>>(x, out, n, s);   // warmup
    cudaDeviceSynchronize();
    cudaEventRecord(a);
    for (int r = 0; r < 30; ++r) strided_sum<<<grid, block>>>(x, out, n, s);
    cudaEventRecord(b); cudaEventSynchronize(b);
    float ms = 0; cudaEventElapsedTime(&ms, a, b); ms /= 30.0f;
    double gbs = (touched * (double)sizeof(float)) / (ms * 1e-3) / 1e9;
    printf("stride=%2d  useful GB/s = %8.1f\n", s, gbs);
  }
  cudaFree(x); cudaFree(out);
  return 0;
}

The second is a Triton bf16 matmul, the kernel that actually exercises the tensor cores. It tiles the output into BLOCK_M × BLOCK_N blocks, streams BLOCK_K-deep slabs of the operands through shared memory, and accumulates in fp32, exactly the reuse scheme of Problem 6. On this H100 a Triton fp16 matmul at \(n = 4096\) reaches 511.6 TFLOP/s against cuBLAS's 702.4 (parallel.json), which is the tile-efficiency gap of Problem 7 made concrete.

# triton bf16 matmul: C = A @ B, tensor-core MMA with fp32 accumulate
# A: (M, K) bf16   B: (K, N) bf16   C: (M, N) fp32
import torch, triton, triton.language as tl

@triton.jit
def matmul_kernel(A, B, C, M, N, K,
                  sam, sak, sbk, sbn, scm, scn,
                  BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):
    pid_m = tl.program_id(0)
    pid_n = tl.program_id(1)
    offs_m = pid_m * BM + tl.arange(0, BM)          # rows of this C tile
    offs_n = pid_n * BN + tl.arange(0, BN)          # cols of this C tile
    offs_k = tl.arange(0, BK)
    a_ptr = A + offs_m[:, None] * sam + offs_k[None, :] * sak
    b_ptr = B + offs_k[:, None] * sbk + offs_n[None, :] * sbn
    acc = tl.zeros((BM, BN), dtype=tl.float32)       # fp32 accumulator
    for k in range(0, K, BK):
        a = tl.load(a_ptr, mask=offs_m[:, None] < M, other=0.0)
        b = tl.load(b_ptr, mask=offs_n[None, :] < N, other=0.0)
        acc += tl.dot(a, b)                          # issues tensor-core MMA
        a_ptr += BK * sak
        b_ptr += BK * sbk
    c_ptr = C + offs_m[:, None] * scm + offs_n[None, :] * scn
    tl.store(c_ptr, acc, mask=(offs_m[:, None] < M) & (offs_n[None, :] < N))

def matmul(a, b):
    M, K = a.shape; K2, N = b.shape; assert K == K2
    c = torch.empty((M, N), device=a.device, dtype=torch.float32)
    grid = (triton.cdiv(M, 128), triton.cdiv(N, 128))
    matmul_kernel[grid](a, b, c, M, N, K,
                        a.stride(0), a.stride(1), b.stride(0), b.stride(1),
                        c.stride(0), c.stride(1),
                        BM=128, BN=128, BK=64, num_warps=8)
    return c

a = torch.randn(4096, 4096, device="cuda", dtype=torch.bfloat16)
b = torch.randn(4096, 4096, device="cuda", dtype=torch.bfloat16)
c = matmul(a, b)
print("max abs err:", (c - (a.float() @ b.float())).abs().max().item())

The third is the roofline arithmetic itself in plain Python, reading the measured constants and reporting, for a kernel's FLOPs and HBM bytes, which roof binds and how far the measurement is from it. This is the calculation behind every table on the page.

# roofline.py -- place a kernel on this H100's roofline
# measured constants (h100.json / parallel.json), this repository
PEAK_BF16 = 744.6e12     # FLOP/s, measured n=4096 bf16 matmul
BETA_HBM  = 3063.5e9     # B/s, measured effective HBM bandwidth (fp32 add)
RIDGE = PEAK_BF16 / BETA_HBM
print(f"ridge point I* = {RIDGE:.0f} FLOP/byte")

def roofline(name, flops, bytes_moved, measured_flops):
    I = flops / bytes_moved
    cap = min(PEAK_BF16, BETA_HBM * I)
    regime = "compute-bound" if I >= RIDGE else "memory-bound"
    print(f"{name:22s} I={I:8.2f}  {regime:14s} "
          f"cap={cap/1e12:7.1f} TFLOP/s  meas={measured_flops/1e12:7.1f} "
          f"({100*measured_flops/cap:4.0f}% of cap)")

# vector add: n adds, 3n fp32 accesses of 4 bytes
n = 268_435_456
roofline("vadd fp32",        n,        3*n*4,           255e9)
# bf16 matmul n=2048: 2 n^3 flops, 3 n^2 * 2 bytes
m = 2048
roofline("matmul bf16 2048", 2*m**3,   3*m*m*2,         453.4e12)
# naive vs flash attention already have measured tflops; show intensity via bytes

How it is done in practice

The gap between these derivations and a deployed system is mostly bookkeeping about which byte crosses which cliff. A production training run for a 70B-parameter model, sharded data-parallel across hundreds of GPUs, spends its time in exactly the quantities this page derived. The forward and backward passes are matmuls the roofline says are compute-bound at large tile sizes, the gradient all-reduce moves the model's worth of bytes across NVLink and the network every step, and the whole thing is fast only if the all-reduce overlaps the backward pass so its bandwidth cost hides under compute. This repository's networks.json models exactly that overlap. For a 70B model at data-parallel width 256 with a 4M-token global batch, the compute is 16.4 s and the NVLink-class all-reduce 5.6 s per step, a 0.34 communication-to-compute ratio that overlaps cleanly, whereas at a 500k-token batch the compute drops to 2.05 s and the same 5.6 s of communication now dominates at a 2.72 ratio, the batch-size cliff that governs how small a step can be before the interconnect becomes the bottleneck.

Inference is the mirror image. The prefill phase is compute-bound matmul, well above the roofline ridge, but the decode phase generates one token at a time and is weight-bound. Each token reads the entire model's weights from HBM to do a single small matmul, so decode throughput is set by HBM bandwidth divided by model bytes, not by tensor-core peak. This is why every serving stack reaches for the memory-side levers this page derived, fp8 and int8 quantization to quarter the weight bytes, 2:4 sparsity to halve them again, KV-cache paging to keep the attention state in HBM rather than recomputing it, and speculative decoding to convert the memory-bound single-token step into a compute-bound multi-token verification. Each is a move on the roofline, sliding decode from the memory roof toward the compute roof, and each is chosen by the same arithmetic that placed the kernels on the roofline above.

The current research frontier

The last few years have pushed on every rung of this page at once. On precision, the frontier is sub-8-bit training and inference. Microscaling (MX) formats, standardized across NVIDIA, AMD, Intel, ARM, and Microsoft in 2023, attach a shared block exponent to short vectors so that fp4 and fp6 become numerically usable, and NVIDIA's Blackwell generation ships fp4 tensor cores that double the fp8 rung again. On sparsity, the debate is whether the rigid 2:4 pattern is the right structure or whether coarser block sparsity and mixture-of-experts routing, which is sparsity at the layer granularity rather than the weight, is the more durable lever. The MoE line, pursued at scale by Google, Mistral, DeepSeek, and others, is sparsity that the memory cascade rewards because only the routed experts' weights cross the HBM cliff per token. On the memory system itself, CXL memory pooling and disaggregation are moving from specification to deployment, and processing-in-memory, computing inside the DRAM or on the HBM base die rather than moving data to a separate compute chip, is an active line at the SAFARI group, Samsung, and SK Hynix, aimed squarely at the bandwidth cliff that this page argues is the binding constraint. On interconnect, the competition is between NVIDIA's NVLink plus InfiniBand fabric, Google's optically-switched TPU pods, and Ethernet-based alternatives standardized through the Ultra Ethernet Consortium, all trying to keep the collective bandwidth of the bandwidth cascade's lowest rungs from capping the largest training runs. The unifying thesis of the frontier is the thesis of this page. Compute has run ahead of memory, and the interesting work is now in the cascade.

Open source to read

• NVIDIA/cutlass is the reference for how a real tensor-core matmul is tiled across the bandwidth cascade. Open media/docs/efficient_gemm.md for the pipelining and shared-memory staging that turn the roofline ceiling into an achieved number.
• triton-lang/triton is the compiler behind the matmul tab. Read python/tutorials/03-matrix-multiplication.py, which is the autotuned version of the kernel above, then 06-fused-attention.py for the flash-attention roofline move.
• NVIDIA/nccl-tests is how the NVLink and collective bandwidth numbers in networks.json are measured. Start with src/all_reduce.cu and the bus-bandwidth definition in the README.
• stas00/ml-engineering is the practitioner's field guide to the whole cascade at cluster scale. The network/ and compute/ chapters connect this page's arithmetic to production training.
• cyanguwa/nersc-roofline is the NERSC roofline toolkit. It shows how to collect the FLOP and byte counters and plot a real kernel on the roofline, the empirical version of the Python tab.
• CMU-SAFARI/ramulator2 is the standard DRAM simulator for the timing and refresh model of the first sections. Its configs expose exactly the \(t_{RCD}\), \(t_{RP}\), and \(t_{RFC}\) parameters worked in Problems 1 and 2.
• NVIDIA/nvbench is a small, honest microbenchmark harness for the kind of bandwidth and occupancy sweeps this page relies on, with correct warmup and timing built in.

Common misconceptions

"HBM is faster memory than DDR." Only in bandwidth. The DRAM banks inside an HBM stack have the same 14-18 ns row and column latencies as a DIMM, and a GPU's end-to-end memory latency is several times worse than a CPU's because of its deep request queues. HBM buys bandwidth and pays with latency, and it is worth it only because a GPU hides latency with the parallelism Little's law quantifies.

"Higher occupancy is always faster." Occupancy sets \(N\) in Little's law, which matters only when the kernel is latency-bound. A compute-bound matmul already issuing a tensor instruction every cycle gains nothing from more warps, and the fastest matmul kernels often run at modest occupancy with high register counts. Read bandwidth first, occupancy second, tensor utilization third.

"A faster clock or a better compiler can speed up a pointer chase." A single dependent chase has one request in flight by construction, so its bandwidth is fixed at one cache line per memory latency, 0.475 GB/s on this host, regardless of clock or flags. The only cure is concurrency, more independent chases, prefetching, or a different data structure.

"Sparsity gives a speedup on any hardware." Only structured sparsity does, and only where the hardware supports it. Unstructured zeros give no speedup on a MAC array because skipping arbitrary positions costs more than the skipped work. The rigid 2:4 pattern exists precisely so the skip maps to a small multiplexer.

"Arithmetic intensity above the ridge means the kernel hits peak." It is a necessary condition, not a sufficient one. An \(n = 1024\) bf16 matmul has intensity 341, well above the 222 ridge, yet reaches only 15 percent of peak because the tiles are too small to fill the machine. The roofline sets the ceiling. Occupancy and tile size decide whether you reach it.

"Lower precision is purely a compute optimization." It is at least as much a memory optimization. fp8 halves the bytes an operand costs against bf16, and for a bandwidth-bound stage such as LLM decode, where each token reads the whole model from HBM, halving the bytes is the entire win, independent of any change in tensor-core throughput.

"PCIe and NVLink differ only in bandwidth." They also differ in coherence. Plain PCIe is non-coherent, requiring explicit flushes and fences, whereas NVLink and NVLink-C2C carry hardware coherence traffic, which is what makes unified memory and fine-grained peer access correct rather than a hazard.

Self-check

References

1. Hennessy, J. and Patterson, D. Computer Architecture: A Quantitative Approach, 6th edition. Morgan Kaufmann, 2019. The standard treatment of memory hierarchy, DRAM, NUMA, and the roofline.
2. Jacob, B., Ng, S., and Wang, D. Memory Systems: Cache, DRAM, Disk. Morgan Kaufmann, 2007. The reference on DRAM device organization, timing parameters, and controller scheduling.
3. Williams, S., Waterman, A., and Patterson, D. "Roofline: an insightful visual performance model for multicore architectures." Communications of the ACM, 52(4), 2009. doi:10.1145/1498765.1498785.
4. Little, J. D. C. "A Proof for the Queuing Formula L = λW." Operations Research, 9(3), 1961. doi:10.1287/opre.9.3.383.
5. Kung, H. T. and Leiserson, C. E. "Systolic Arrays (for VLSI)." Sparse Matrix Proceedings, 1978. The origin of the systolic dataflow.
6. Chen, Y.-H., Emer, J., and Sze, V. "Eyeriss: A Spatial Architecture for Energy-Efficient Dataflow for Convolutional Neural Networks." ISCA, 2016. The dataflow taxonomy and the row-stationary design.
7. Sze, V., Chen, Y.-H., Yang, T.-J., and Emer, J. Efficient Processing of Deep Neural Networks. Morgan & Claypool, 2020. The survey with the normalized data-movement energy model.
8. Jouppi, N. P. et al. "In-Datacenter Performance Analysis of a Tensor Processing Unit." ISCA, 2017. arXiv:1704.04760.
9. NVIDIA. "NVIDIA H100 Tensor Core GPU Architecture." Whitepaper, 2022. Source for the SM, tensor-core, and HBM3 figures compared against measurement.
10. Markidis, S., Der Chien, S. W., Laure, E., Peng, I. B., and Vetter, J. S. "NVIDIA Tensor Core Programmability, Performance & Precision." IPDPSW, 2018. arXiv:1803.04014.
11. Micikevicius, P. et al. "Mixed Precision Training." ICLR, 2018. arXiv:1710.03740.
12. Micikevicius, P. et al. "FP8 Formats for Deep Learning." 2022. arXiv:2209.05433.
13. Mishra, A. et al. "Accelerating Sparse Deep Neural Networks." NVIDIA, 2021. arXiv:2104.08378. The 2:4 structured-sparsity recipe.
14. Rixner, S., Dally, W. J., Kapasi, U. J., Mattson, P., and Owens, J. D. "Memory Access Scheduling." ISCA, 2000. FR-FCFS and the DRAM scheduling baseline.
15. Mittal, S. "A Survey of Recent Prefetching Techniques for Processor Caches." ACM Computing Surveys, 49(2), 2016. doi:10.1145/2907071.
16. Kim, Y., Yang, W., and Mutlu, O. "Ramulator: A Fast and Extensible DRAM Simulator." IEEE Computer Architecture Letters, 2015. The standard DRAM timing/refresh simulator (SAFARI group).
17. Kim, Y. et al. "Flipping Bits in Memory Without Accessing Them: An Experimental Study of DRAM Disturbance Errors (RowHammer)." ISCA, 2014.
18. Jia, Z., Maggioni, M., Staiger, B., and Scarpazza, D. P. "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking." 2018. arXiv:1804.06826. GPU memory-latency methodology.
19. Volkov, V. "Understanding Latency Hiding on GPUs." PhD thesis, University of California, Berkeley, 2016. The GPU statement of Little's law for occupancy.
20. Abts, D. et al. "Think Fast: A Tensor Streaming Processor (TSP) for Accelerating Deep Learning Workloads (Groq)." ISCA, 2020.
21. Lauterbach, G. "The Path to Successful Wafer-Scale Integration: The Cerebras Story." IEEE Micro, 41(6), 2021.
22. Vasiljevic, J. et al. "Compute Substrate for Software 2.0 (Tenstorrent)." IEEE Micro, 41(2), 2021.
23. JEDEC. JESD238 High Bandwidth Memory (HBM3) DRAM standard, 2022, and the earlier JESD235 HBM standard. Source for channel, pseudo-channel, and per-pin-rate definitions.