Memory: capacity, bandwidth, and who decides placement

The decision turns on exactly one quantity: how much uncertainty remains at runtime.

A cache is insurance. You pay a continuous premium — tag arrays, comparators, replacement metadata, coherence traffic, and typically 5–10% of the array’s area plus energy on every access — in exchange for a policy that adapts to an access pattern nobody predicted. The payoff is that unanticipated behaviour degrades rather than fails. That is worth a great deal when the access stream depends on data (pointer chasing, hash lookups, branchy control flow), and it is worth almost nothing when the stream is known at compile time.

A scratchpad is the opposite bet. No tags, so no tag energy and no tag area; every byte is usable rather than statistically usable, which matters more than it sounds — a 128 KB cache holds substantially less than 128 KB of your working set once conflict misses and replacement policy are accounted for, while a 128 KB scratchpad holds 128 KB. Timing is deterministic, which is a prerequisite for static scheduling and for any hard-deadline system. And the compiler can implement policies a cache cannot express at all: keep this tile resident across the whole loop nest, evict that one immediately because it is streaming and will never be reused, double-buffer this one against the next iteration.

ML is unusually favourable to the scratchpad because its three preconditions all hold at once. The loop bounds are known before execution; the index expressions are affine, so the exact access sequence is computable; and the reuse is both large and statically derivable from the loop structure. Very little else in computing has all three. That is why the TPU has no data caches — VMEM is programmer-controlled scratchpad, and XLA schedules every transfer — and why the GPU keeps adding scratchpad-shaped machinery: shared memory has always been one, cp.async made explicit staging cheap on Ampere, and Hopper’s TMA is a dedicated engine for compiler-issued bulk tile copies with address generation in hardware.

Two honest qualifications:

• Real designs are hybrids, deliberately. A GPU has scratchpad (shared memory) and cache (L1, L2) at nearly the same level, so a kernel can use explicit staging for the tensors it understands and let the cache absorb everything else. This is the right engineering answer, because a real model has some tensors with statically-known reuse and some without.
• The scratchpad’s failure mode is not slowness. When the compiler cannot prove a shape or an index, it does not run slowly — it cannot generate the schedule at all, and must fall back. That cliff is exactly Part Q, and it is the price of the whole strategy.

Ask what the SRAM would be for, because the two purchases substitute for each other only through the mapping, and the exchange rate is bad.

The relevant closed form, derived properly in Module 7: for a GEMM tiled at a level of capacity $C$, the traffic through that level scales as $O(N^3/\sqrt{C})$. Traffic falls as the square root of capacity. So doubling SRAM cuts traffic by a factor of $\sqrt 2 \approx 1.41$, whereas doubling HBM bandwidth cuts memory time by 2× directly. On that comparison alone, bandwidth wins for any workload already tiled sensibly, and it is the reason vendors spend so aggressively on HBM despite its cost.

But the square-root law describes the continuous regime, and the decision is usually made on thresholds, where SRAM wins enormously:

• A working set that now fits entirely. If doubling SRAM lets the whole weight tile, or the whole KV block, or an entire small model live on-die, traffic does not fall by 1.41× — it falls to roughly zero, because the data is fetched once instead of once per tile. Groq’s architecture is this threshold argument taken to its conclusion: enough SRAM to hold the model, and no HBM at all.
• You cannot buy HBM bandwidth freely. SRAM bandwidth scales with banks and area, which is a design decision. HBM bandwidth scales with stacks and beachfront, which is a packaging constraint with a hard ceiling and a large cost per increment. Frequently “double the HBM bandwidth” is simply not on the menu at any price, which settles the question.
• Energy. ~1 pJ/byte against ~50 pJ/byte. In a power-limited design — which every large accelerator is — moving traffic on-die is not merely faster, it is what lets the clocks stay up.

So the decision rule: if a threshold is within reach, take the SRAM; otherwise take the bandwidth. And the codesign-specific version, which is the point of this module — you cannot answer the question without knowing the mapping, because the mapping decides both what the working set is and whether a threshold exists. That is the circular dependency from Module 1, made concrete enough to cost.

Because latency is hidable and wasted granularity is not.

Latency has a standard defence available at every level: get more work in flight. Warps, out-of-order windows, prefetchers, double buffering, asynchronous copies — every one of these overlaps a long latency with useful work, and if you have enough independent work, latency effectively disappears. Little’s law makes this quantitative: to sustain bandwidth $B$ at latency $L$ you need $B \times L$ bytes in flight, and on an H100 that is roughly 3.35 TB/s × 400 ns ≈ 1.3 MB of memory requests outstanding at all times. Large, but achievable, and it is what the whole SIMT machine is for.

Wasted granularity has no such defence. If you request 4 useful bytes and the memory system delivers a 32-byte sector, you have consumed 32 bytes of the bandwidth you paid for and used 4. No amount of concurrency recovers the other 28 — the bandwidth is gone, spent on bytes that were fetched and discarded. More parallelism makes you waste it faster. The only fix is changing the access pattern itself, which means changing the layout or the mapping.

This has consequences that show up all over the stack, and it is worth seeing them as one phenomenon:

• Layout is a first-class mapping decision, not an afterthought. NCHW versus NHWC, row versus column major, blocked and tiled layouts, and the padding of leading dimensions to avoid bank conflicts and channel aliasing are all granularity management.
• Gathers are expensive for a structural reason. An embedding lookup or an unstructured sparse access requests a handful of bytes per index from unrelated addresses. Each one costs a full sector and possibly a DRAM row activation. This is why unstructured sparsity underdelivers so badly (Module 9) — the FLOPs vanish but the traffic barely moves.
• Granularity compounds up the ladder. 32-byte sector, then a DRAM row of a few kilobytes, then a PCIe TLP, then a network MTU. A pattern that is merely inefficient at one level can be catastrophic once it crosses a level with much coarser granularity, which is why the same code can be 20% slow on-node and 10× slow across a network.

The framing to keep: latency is a concurrency problem, granularity is a layout problem, and only one of them has a general solution.

Because they scale by different physical mechanisms, and only one of them tracks Moore’s law.

On-die SRAM capacity and bandwidth scale with area: more macros, more banks, more ports, all of which come from transistors that keep getting smaller and cheaper. HBM bandwidth scales with package beachfront — the physical perimeter available for stacks, the interposer routing, the number of TSVs and the achievable signalling rate. That does not shrink with the process node. The result is the numbers from Module 2: H100 at 989 BF16 TFLOP/s against 3.35 TB/s (~295 FLOP/byte), B200 at ~2.25 PFLOP/s against 8 TB/s (~280), TPU v6e at 920 TFLOP/s against 1.6 TB/s (~575). Compute grows faster; the ridge point climbs; and the required arithmetic intensity climbs with it.

What software has to do in response is a coherent programme, not a list of tricks:

• Tile larger, and tile more levels. A higher ridge point means the tile that saturates the machine is bigger, which means it needs the bigger on-die capacity that the same trend is providing. These two facts are the same fact.
• Fuse aggressively. Every intermediate that reaches HBM and comes back is traffic charged against a shrinking budget. Fusion converts it to SRAM traffic at ~1/50 the energy. This is why the compiler modules in the ML Software series spend so much effort on fusion — it is a direct response to this trend.
• Recompute rather than store. When arithmetic is nearly free and bandwidth is scarce, recomputing a value is cheaper than saving and reloading it. This is why AOTAutograd’s min-cut partition is a genuine optimisation rather than a memory-saving hack, and why FlashAttention recomputes attention weights in the backward pass instead of materialising them.
• Send fewer bytes. Module 10.

The unifying statement: the trend converts problems from “not enough FLOPs” into “not enough locality”, and locality is a software property. That is precisely why the mapping deserves a whole Part of this series, and why hardware teams increasingly measure their compilers rather than their peak numbers.