Computer architecture, from logic gates to out-of-order superscalars and GPUs

Why this subject matters now

For roughly forty years, software got faster by waiting. Dennard scaling let each process generation shrink transistors, drop their voltage, and raise the clock while keeping power density constant, so a binary compiled in 1995 ran several times faster in 2000 on the same source code. That regime ended around 2005 when threshold voltages stopped scaling and leakage current made further voltage reduction impossible. Clocks froze near 3 to 5 GHz and the industry pivoted to multicore, then to wide SIMD, then to GPUs, and now to domain-specific accelerators. Each pivot moved work that hardware used to do implicitly onto the programmer. Cores require parallel algorithms, SIMD requires data layout discipline, GPUs require explicit management of a memory hierarchy, and accelerators require casting the computation as the one dataflow the silicon implements. A practitioner today is expected to reason quantitatively about all of it, to read a profile and know whether a kernel is bound by compute or by memory, to predict when a branch-heavy loop will fall off a cliff, to explain why a transformer trains at 74 percent of a GPU's peak while a memory-bound elementwise pass reaches under one percent, and to know which of those two numbers is actually good.

The subject also matters in the opposite direction. Machine learning has become the dominant customer of new silicon, and the architectures being built for it (tensor cores, systolic arrays, HBM stacks, low-precision datapaths, structured sparsity) are design points chosen from exactly the theory on this page. The "hardware lottery" observation, that research directions win partly because they map onto existing hardware, makes architectural literacy a research skill, not only an engineering one. The good news is that the field rests on a small number of load-bearing ideas, namely locality, pipelining, prediction, renaming, and replication. Everything below is one of those five wearing different clothes.

Digital logic, where the clock period comes from

Combinational versus sequential circuits

A combinational circuit is a pure function of its current inputs. Gates, muxes, adders, and ALUs are all of this kind. Change the inputs and, after the gates' propagation delays settle, the outputs are the new function value, with no memory involved. A sequential circuit adds state, almost always as edge-triggered D flip-flops. On the rising edge of the clock, each flip-flop samples its input D and holds it on its output Q until the next edge. A synchronous machine is nothing more than these two elements alternating. Registers hold the state, a cloud of combinational logic computes the next state and the outputs from the current state, and the clock edge commits the result.

        ┌───────────────┐
 clk ──► │  registers    │──── current state ────┐
         │  (D flip-flops)│                      ▼
         └───────▲───────┘        ┌────────────────────────┐
                 │                │  combinational logic   │
                 └── next state ──│  (gates, no memory)    │
                                  └────────────────────────┘

Three timing parameters of the flip-flop govern everything. The clock-to-Q delay \(t_{cq}\) is how long after the edge the stored value appears at Q. The setup time \(t_{setup}\) is how long before the next edge the input D must be stable to be captured reliably. The hold time \(t_{hold}\) is how long after the edge D must remain stable so the edge does not capture a value that is already changing. From these, two constraints fall out. The setup constraint bounds the clock period from below,

where \(t_{pd}\) is the propagation delay of the longest combinational path between any two registers, the critical path. The signal must leave the source register (\(t_{cq}\)), traverse the slowest logic (\(t_{pd}\)), and arrive at the destination early enough to satisfy setup. The hold constraint is independent of the clock period entirely,

where \(t_{cd}\) is the contamination delay, the shortest path through the logic. It says the fastest new value racing out of a register after an edge must not reach the next register before that register's hold window closes. A hold violation cannot be fixed by slowing the clock. It is fixed by inserting delay buffers, which is why it is the more feared of the two in physical design. Clock skew, the difference in edge arrival time between two registers, tightens both constraints. Skew toward the receiver eats into the setup margin's benefit and directly worsens the hold requirement.

Why pipelining raises throughput and not latency

The setup constraint says frequency is the reciprocal of the worst register-to-register path. Pipelining is the act of cutting a long combinational path into k shorter segments by inserting registers, so the clock can run roughly k times faster and a new input can be accepted every (shorter) cycle. Two facts follow immediately from the timing algebra and are worth internalizing as algebra, not slogans. First, throughput improves by at most k, and in practice by less, because the cut points are never perfectly balanced (the clock is set by the slowest stage, not the average) and because every added register level charges the fixed toll \(t_{cq} + t_{setup}\) per stage. Second, the latency of any single input through the machine gets worse, not better. The work is the same, the register toll is now paid k times, and imbalance rounds every stage up to the slowest one. Pipelining is purely a throughput transformation. If one result is needed as fast as possible and there is no stream of independent work behind it, pipelining is a loss.

ISA design, the contract between software and silicon

RISC versus CISC, stated honestly

An instruction set architecture is the durable interface, the registers, instructions, memory model, and exception behavior that software may rely on across implementations. The classic RISC argument (fixed-length instructions, load/store architecture, many general-purpose registers, simple addressing modes) was that a simple, regular ISA lets the implementation be pipelined aggressively and lets the compiler do the scheduling. The classic CISC counterargument was code density and fewer instructions per task. The honest modern accounting is more nuanced than either camp's slogans. Every high-performance x86 since the mid-1990s has been a RISC-like out-of-order core behind a decoder that cracks x86 instructions into micro-ops, so the execution engines of a modern Intel, AMD, and Arm core look strikingly alike. The ISA war was settled by translation, not victory. What variable-length encoding still costs is the front end. Finding instruction boundaries in an x86 byte stream requires either predecode bits, parallel speculative decoders, or a micro-op cache that bypasses decode entirely, and this is real area, power, and design effort. Fixed 4-byte instructions are one reason Apple could build a very wide in-house decoder for its M-series cores with comparative ease. The cost is bounded, though. On large out-of-order cores the decode tax is a few percent of power, which is why x86 remains competitive. On small cores the tax is proportionally larger, which is why embedded and mobile went RISC decades ago and never returned.

RISC-V is the current center of gravity for both teaching and new silicon because it is modular and unencumbered. The base integer ISA (RV32I or RV64I) is small enough to learn in an afternoon, with 32 registers (x0 hardwired to zero), load/store only, and one memory addressing mode, base register plus 12-bit signed immediate. Everything else is an extension, M (multiply/divide), A (atomics), F/D (floating point), C (16-bit compressed encodings, which recover most of CISC's density advantage), V (vectors), and the privileged architecture for operating systems. The compressed extension is the quiet workhorse. RV64GC code density is comparable to x86-64, removing the last practical CISC advantage while keeping decode trivial, since a 16-bit or 32-bit boundary is determined by two bits of the first parcel.

Addressing modes and calling conventions

Addressing modes are where ISAs reveal their philosophy. x86-64 computes addresses as base + index × scale + displacement in a single instruction, and instructions can take one memory operand directly, so a compiler can fold an array access into an add. Classic Arm adds pre- and post-increment modes that update the base register as a side effect, which shrinks loops over arrays. RISC-V offers only base + immediate, on the argument that the fancier modes cost decode complexity and a wider adder in a critical path while the compiler can synthesize them in one extra instruction that is usually hoisted or strength-reduced anyway. Measurement mostly vindicated this. The fused compare-and-branch and the compressed extension bought back far more than rich addressing modes would have.

The calling convention is the ISA's social contract, and reading disassembly is impossible without it. In the standard RISC-V convention, arguments and return values travel in a0 through a7, ra holds the return address written by jal, sp is the stack pointer, t-registers are caller-saved scratch (the callee may destroy them), and s-registers are callee-saved (a function that uses s1 must save and restore it). A minimal function that needs one saved register compiles to a prologue that pushes ra and s0, a body, and an epilogue that pops and returns.

sum_sq:                        # int sum_sq(int *p, int n)
    addi  sp, sp, -16          # prologue: make frame
    sd    ra, 8(sp)            # save return address (callee will call)
    sd    s0, 0(sp)            # save callee-saved s0
    li    s0, 0                # acc = 0
loop:
    beqz  a1, done             # while (n != 0)
    lw    t0, 0(a0)            # t0 = *p        (base + imm addressing)
    mulw  t0, t0, t0           # t0 = t0 * t0
    addw  s0, s0, t0           # acc += t0
    addi  a0, a0, 4            # p++
    addiw a1, a1, -1           # n--
    j     loop
done:
    mv    a0, s0               # return value in a0
    ld    s0, 0(sp)            # epilogue: restore
    ld    ra, 8(sp)
    addi  sp, sp, 16
    ret                        # jalr x0, 0(ra)

The landscape today looks like this. x86-64 (Intel and AMD) still owns desktops and the majority of server sockets. Arm owns mobile entirely, owns Apple's product line, and has serious server share through AWS Graviton, Ampere, and NVIDIA Grace. RISC-V dominates new embedded designs and management cores (billions of units, including controllers inside NVIDIA GPUs and Western Digital drives) and is climbing toward application processors. The interesting structural fact is that all three now ship the same microarchitectural playbook, so ISA choice is increasingly about licensing, ecosystems, and software, not about achievable performance.

The pipelined datapath and its hazards

From single-cycle to five stages

The single-cycle datapath executes one instruction per clock, with the period stretched to cover the slowest instruction (a load, which needs fetch, decode, address arithmetic, memory, and writeback in one period). Problem 1 already showed why this wastes time, since every instruction pays for the slowest one's path. The classic fix is the five-stage pipeline. IF fetches from instruction memory, ID decodes and reads the register file, EX runs the ALU, MEM accesses data memory, WB writes the register file. Pipeline registers between stages carry both data and the control signals each instruction will need downstream, so five instructions are in flight at once, each owning one stage.

cycle:      1     2     3     4     5     6     7     8
lw   x2,0(x1)   IF    ID    EX    MEM   WB
add  x3,x4,x5         IF    ID    EX    MEM   WB
sub  x6,x7,x8               IF    ID    EX    MEM   WB
and  x9,x2,x3                     IF    ID    EX    MEM   WB

Hazards are the situations where the overlap lies about program order. A structural hazard is two instructions needing the same hardware in the same cycle. The canonical one, instruction fetch and data access colliding on a single memory, is designed away by split instruction and data caches, and the register file is written in the first half of a cycle and read in the second so WB and ID can share it. A data hazard is an instruction needing a value an older in-flight instruction has not yet written. A control hazard is not knowing what to fetch after a branch until the branch resolves.

Forwarding, derived rather than asserted

Consider add x3, x1, x2 followed immediately by sub x5, x3, x4. The add computes x3 at the end of its EX cycle but does not write the register file until WB, two cycles later. The sub reads registers in ID, one cycle after the add's EX. Without help the sub would read a stale x3 and the pipeline would need two stall cycles. But the needed value already exists inside the machine, sitting in the EX/MEM pipeline register. Forwarding (bypassing) adds muxes in front of each ALU input that can select, instead of the register-file read, either the EX/MEM register (result of the instruction one ahead) or the MEM/WB register (two ahead). The control equation is exactly the hazard condition read off the pipeline diagram. Forward from EX/MEM to an ALU source if that older instruction writes a register (RegWrite is set), its destination is not x0, and its destination equals the source register number. Forward likewise from MEM/WB, but only if the EX/MEM case did not already match, so the most recent value wins when both are candidates.

            ┌──────────── forwarding paths ────────────┐
            │                                          │
   ID/EX ──►│  ┌─────┐    EX/MEM         MEM/WB        │
  regfile ─►├─►│ ALU │──►[R]────────────►[R]──► WB     │
   values ─►│  └─────┘     │               │           │
            │      ▲       │               │           │
            │      └───────┴◄──────────────┘           │
            │      (mux selects newest matching value) │
            └──────────────────────────────────────────┘

One data hazard survives forwarding. A load produces its value at the end of MEM, but an immediately following dependent instruction needs it at the start of its own EX, which is the same cycle as the load's MEM. The value does not exist yet anywhere in the machine at the moment it is needed, and no wire can forward backward in time. The hardware must stall one cycle (a hazard detection unit in ID freezes IF and ID and injects a bubble into EX), after which the MEM/WB forwarding path covers it. This is the load-use hazard, and it is why compilers and out-of-order schedulers try to put an independent instruction between a load and its consumer.

Control hazards and CPI arithmetic

A conditional branch resolved in EX means two younger instructions are already in IF and ID when the outcome is known. If the prediction was wrong they must be squashed, costing two cycles. Moving the comparison to ID (feasible for simple equality tests) cuts the penalty to one at the price of a tighter critical path and harder forwarding into ID. Real machines instead predict, which is the next section. Here it is enough to do the accounting. The pipelined CPI is 1 plus the average stalls per instruction,

Branch prediction

From static hints to saturating counters

The cheapest predictor is static. Predict backward branches taken (they are usually loop bottoms) and forward branches not taken. This reaches roughly 60 to 70 percent on integer code and costs nothing, but a pipeline that fetches many instructions per cycle from a deep pipe needs far better. The first dynamic step is a table of 2-bit saturating counters indexed by low PC bits, with states strongly-not-taken (00), weakly-not-taken (01), weakly-taken (10), and strongly-taken (11). Predict from the high bit, increment on taken, decrement on not-taken, saturating at the ends. The second bit exists for hysteresis. A loop branch that is taken nine times then falls through once would, with a 1-bit predictor, mispredict twice per loop visit, once at the exit, and again at the next entry because the exit flipped the bit. The 2-bit counter mispredicts only the exit. The single not-taken decrements strongly-taken to weakly-taken, which still predicts taken at the next entry. For an N-iteration loop the accuracy is \((N-1)/N\) instead of \((N-2)/N\), and for nested short loops the difference compounds.

Correlation, gshare, and TAGE

Per-branch counters cannot see patterns that span branches. The branch guarding if (x == 0) is perfectly predictable from the branch that just set x, and an alternating pattern T,N,T,N defeats a saturating counter entirely (it oscillates between the weak states, mispredicting half the time). Two-level predictors, introduced by Yeh and Patt, add a global history register (GHR), a shift register of the outcomes of the last k branches, and use it to index the counter table. Different recent histories get different counters, so the alternating branch trains two counters, one for "last was T" and one for "last was N", each of which saturates correctly. gshare, McFarling's refinement, XORs the history with the PC bits rather than concatenating them, so the table is shared across branches with less systematic aliasing,

The state of the art in this line is TAGE (tagged geometric history length), due to Seznec. TAGE keeps several tables, each indexed by a hash of the PC with a different history length, and the lengths form a geometric series (for example 4, 8, 16, 32, 64, 128 bits), so the predictor can capture both very short and very long correlations without paying long-history storage for every branch. Entries are tagged, so a lookup knows whether the entry actually belongs to this (PC, history) pair rather than being an alias. Prediction comes from the matching table with the longest history (the provider). Allocation on a misprediction trains a longer-history table, so branches migrate to exactly the history length they need. TAGE variants with a statistical corrector and a loop predictor (TAGE-SC-L) have won every championship branch prediction contest since 2006 and are understood to approximate what ships in current high-end cores, which reach roughly 3 to 6 mispredictions per thousand instructions on general code.

Targets, the BTB and the return address stack

Direction is only half the problem. Fetch needs the target address in the same cycle. The branch target buffer (BTB) is a cache indexed by fetch PC that remembers, for recently taken branches, where they went, so the front end can redirect fetch with zero bubbles on a BTB hit. Returns defeat a BTB because one return site has many callers. The return address stack (RAS) is a small hardware stack that pushes on call and pops on return, predicting returns nearly perfectly until the stack overflows or speculation corrupts it (real designs checkpoint the RAS top pointer per predicted branch for repair). Indirect branches (virtual calls, switch tables, interpreter dispatch) get their own history-indexed target predictor (the ITTAGE variant of TAGE). Interpreters are famous branch-predictor stress tests precisely because one indirect jump's target encodes the entire bytecode stream.

Caches and the memory hierarchy

The mapping spectrum and the address breakdown

DRAM latency has improved by small constant factors while processor cycle times improved by orders of magnitude, so a load that misses everything costs on the order of 200 to 400 cycles on a current server part. The entire memory hierarchy exists to hide this behind two empirical regularities, temporal locality (recently used data is likely to be used again) and spatial locality (data near recently used data is likely to be used soon). A cache stores fixed-size blocks (lines), essentially always 64 bytes on CPUs, and the design space is where a given line is allowed to live. In a direct-mapped cache each address maps to exactly one location, found by indexing with middle address bits. Lookup is a single tag compare, fast and cheap, but two hot addresses that share an index evict each other forever. In a fully associative cache a line can live anywhere, so there are no conflicts, but every lookup compares every tag, which only works for small structures like TLBs. Set-associative is the compromise. The index selects a set, and the line may occupy any of the W ways in it, with W tag compares in parallel. Hill and Smith's classic measurement, still roughly true, is that doubling associativity cuts the miss rate about as much as doubling capacity, up to around 8 ways, after which conflict misses are mostly gone.

The address arithmetic is mechanical and worth being fast at. With a capacity of C bytes, line size B, and associativity W, the offset bits \(= \log_2 B\), the number of sets \(S = C/(B \times W)\), the index bits \(= \log_2 S\), and the tag is every remaining high bit. A concrete example that recurs below is a 32 KiB, 8-way L1 with 64-byte lines, which has \(S = 32768/(64 \times 8) = 64\) sets, so 6 offset bits, 6 index bits, and for 48-bit physical addresses a \(48 - 12 = 36\)-bit tag.

48-bit address, 32 KiB / 8-way / 64 B lines:

  47                          12 11        6 5         0
  ┌──────────────────────────────┬───────────┬───────────┐
  │            tag (36)          │ index (6) │ offset (6)│
  └──────────────────────────────┴───────────┴───────────┘
                                  64 sets      64 B line
  index+offset = 12 bits = 4 KiB = one way = one page (this
  coincidence is engineered; see VIPT below)

The three Cs, write policies, and replacement

Misses are classified by what would have prevented them. Compulsory misses are first-ever touches, and only larger lines or prefetching help. Capacity misses would occur even in a fully associative cache of the same size, because the working set is simply too big. Conflict misses are the remainder, artifacts of limited associativity, and more ways or better index hashing help. On multiprocessors a fourth C, coherence misses, appears, lines invalidated by another core's writes, the subject of the multicore section. Write policy is a second axis. Write-through sends every store to the next level (simple, always-clean lines, but bandwidth-hungry, so it survives mainly in L1s paired with coalescing write buffers). Write-back marks the line dirty and writes it out only on eviction, which is what every modern data cache does. Orthogonally, write-allocate fetches a line on a store miss (betting on locality of subsequent accesses) while no-write-allocate forwards the store onward without filling. Write-back caches essentially always allocate. Replacement within a set is LRU or an approximation (tree-pseudo-LRU) at low associativity. Last-level caches increasingly use scan-resistant policies like RRIP, because a single streaming pass over a large array should not be allowed to flush the entire cache, which is exactly what true LRU permits.

AMAT, the hierarchy's figure of merit

Average memory access time composes recursively. The time at each level is its hit time plus its local miss rate times the time of the level below.

The local miss rate (misses at this level divided by accesses that reached this level) is the right quantity for this formula. The global miss rate (misses divided by all CPU accesses) is the product of the local rates above and including the level, and confusing the two is the most common error in hierarchy arithmetic.

Cache blocking, making arithmetic intensity rather than finding it

The naive triple loop for \(C = AB\) with \(n \times n\) matrices performs \(2n^3\) floating point operations. If the matrices do not fit in cache, the inner product formulation streams a row of A against a column of B for every output element, and each element of B is fetched from memory \(n\) times over the whole computation. Blocking (tiling) partitions the matrices into \(b \times b\) tiles chosen so that three tiles fit in cache, and computes tile-by-tile. Each tile of A and B is then loaded once per tile-row/column of the result, i.e. \(n/b\) times, so total traffic falls from \(O(n^3)\) words to roughly \(2n^3/b + n^2\) words. The flops are unchanged. The arithmetic intensity, flops per byte of memory traffic, for single precision (4-byte words) is therefore

growing linearly in the tile size. The constraint is \(3b^2 \times 4 \le C\). For the 32 KiB L1, \(b \le \sqrt{32768/12} \approx 52\), so \(b = 48\) gives AI \(\approx 12\) flops/byte against roughly 0.17 flops/byte for the unblocked streaming case (2 flops per 12 bytes moved). This is the single most important transformation in high-performance computing, and it is worth seeing that it changes no arithmetic at all. It only reorders the same operations so the reuse that was always present algebraically becomes reuse the cache can see within its capacity. Hong and Kung proved the matching lower bound, that any schedule of the classical algorithm with a fast memory of size M must move \(\Omega(n^3/\sqrt{M})\) words, so blocking with \(b \sim \sqrt{M/3}\) is asymptotically optimal, not just a good trick. The C code in the implementation section measures this effect directly, and the same idea reappears twice more on this page. FlashAttention is blocking applied to the attention score matrix, and the CUDA tiled matmul is blocking applied to GPU shared memory.

Virtual memory

Pages, page tables, and a concrete multi-level walk

Virtual memory gives every process a private address space, enforces protection, and lets physical memory be allocated lazily and non-contiguously. Both are divided into 4 KiB pages, and a per-process page table maps virtual page numbers (VPN) to physical page numbers (PPN). A flat table is impossibly large. A 48-bit space of 4 KiB pages has \(2^{36}\) entries, 512 GiB of table at 8 bytes each, almost all of it describing unmapped holes. Multi-level (radix-tree) tables fix this by making the table sparse. The VPN is split into fields, each indexing one level, and subtrees for unmapped regions are simply absent. RISC-V Sv39 is the cleanest concrete example. A 39-bit virtual address is a 12-bit page offset and three 9-bit VPN fields, and each level is one 4 KiB page holding 512 8-byte page table entries (PTEs). Nine bits indexing 512 entries of 8 bytes is exactly 4 KiB, since the radix is chosen so each table node is itself one page.

Walking a concrete address makes the mechanism stick. Take virtual address 0x16582A7C4 with the root page table at physical page 0x80123 (held in the satp register). The offset is the low 12 bits, 0x7C4. The VPN is the address shifted right by 12, 0x16582A = 1,464,362. Splitting into 9-bit fields from the top gives VPN[2] \(= 1464362 \gg 18 = 5\), VPN[1] \(= (1464362 \gg 9) \,\&\, 511 = 300\), VPN[0] \(= 1464362 \,\&\, 511 = 42\).

VA 0x1_6582_A7C4  (Sv39: 9 | 9 | 9 | 12)

   000000101 | 100101100 | 000101010 | 0111 1100 0100
   VPN[2]=5    VPN[1]=300  VPN[0]=42    offset=0x7C4

walk (each step is one physical memory read):
 1. root:  PTE at 0x80123000 + 5*8   = 0x80123028 → PPN of level-1 table
 2. mid:   PTE at (that PPN «12) + 300*8         → PPN of level-0 table
 3. leaf:  PTE at (that PPN «12) + 42*8          → PPN of the data page,
           plus V/R/W/X/U/A/D permission bits
 PA = (leaf PPN « 12) | 0x7C4

Three loads to translate one load, a 4x cost before the data access itself, and x86-64 with four levels (or five with 57-bit addressing) is worse. If a non-leaf PTE has its valid bit clear the walk stops and the hardware raises a page fault, which is how demand paging, copy-on-write, and mmap all work. The OS handles the fault, fixes the tables, and restarts the instruction, which requires the precise exceptions the out-of-order section builds. A leaf at level 1 instead of level 0 maps a 2 MiB megapage (512 x 4 KiB), and at level 2 a 1 GiB page. Large pages exist mostly to fix the TLB economics computed next.

TLBs and reach

No machine walks tables on every access. The translation lookaside buffer caches recent VPN-to-PPN mappings. A typical current design has a small fully associative L1 TLB (tens of entries, checked in parallel with the L1 cache) and an L2 TLB of one to two thousand entries. The figure of merit is reach. A 1536-entry L2 TLB of 4 KiB pages covers \(1536 \times 4\,\text{KiB} = 6\) MiB. Any working set beyond 6 MiB, which describes essentially every database, JVM heap, and ML workload, misses the TLB even though the data may be sitting in the last-level cache, and each miss costs a multi-level walk (mitigated by dedicated page-walk caches for the upper levels). The same TLB with 2 MiB pages reaches 3 GiB, which is why transparent huge pages and explicit hugetlbfs are worth 10 to 40 percent on TLB-bound workloads and why every serious database and allocator has a huge-page story.

VIPT, why L1 size, ways, and page size are entangled

The L1 lookup wants to start before translation finishes. The virtually-indexed, physically-tagged (VIPT) trick is to index the cache with address bits that are identical in the virtual and physical address, namely the page-offset bits, while the walk or TLB produces the physical tag in parallel. The tag compare then uses physical bits, so no aliasing or homonym problems arise. The constraint is that index + offset bits must fit within the 12-bit page offset, i.e. one way of the cache can be at most one page. The 32 KiB 8-way L1 from the cache section is exactly at the boundary, \(32768 / 8 = 4096\) bytes per way, 6 index + 6 offset = 12 bits, all untranslated. This is not a coincidence but a constraint that has pinned x86 L1 data caches at 32 to 48 KiB for two decades. Growing the L1 requires either more ways (Apple's M-series uses large, higher-associativity L1s enabled partly by the 16 KiB base page of its OS, which frees two more untranslated bits) or giving up pure VIPT.

Out-of-order execution

Why in-order stalls are unacceptable

An in-order pipeline stalls every instruction behind a long-latency one. A single L2 miss freezes hundreds of cycles of issue slots even when abundant independent work sits just behind it in program order. The dependences that actually constrain correctness are read-after-write (RAW), true dataflow. The other two hazard classes, write-after-read (WAR) and write-after-write (WAW), are artifacts of reusing a finite set of register names, not of the computation, and they can be removed by giving each new write a fresh physical location. Out-of-order execution is exactly this. Rename away the false dependences, then let instructions execute in dataflow order, constrained only by RAW edges and structural resources, while a separate in-order mechanism preserves the illusion of sequential execution for exceptions and the outside world.

Tomasulo's algorithm, worked step by step

Tomasulo's 1967 design for the IBM 360/91 floating-point unit contains the whole idea. Each functional unit has reservation stations (RS), buffers holding an operation plus, for each source operand, either the value itself or the tag of the RS that will produce it. A register alias table (RAT, called the register status table in the original) records, for each architectural register, whether its latest value is in the register file or is being produced by some RS. Results broadcast on a common data bus (CDB) carrying (tag, value), and every waiting RS and the register file snoop it and capture matching values. Renaming happens at issue. An instruction reads its sources as values if ready, otherwise as tags, and then overwrites the RAT entry of its destination with its own tag. Because a consumer captures the value or the tag of the specific producer it read at issue, later writes to the same architectural register cannot disturb it, so WAR is gone. Because the RAT always points to the youngest writer, only the youngest write updates the register, so WAW is gone.

The standard six-instruction sequence exercises every case. Assume two load buffers with pipelined 2-cycle loads, latencies of 2 for add/sub, 10 for multiply, 40 for divide, one instruction issued per cycle, one CDB, and execution may start the cycle after the last operand arrives.

I1: fld  f6, 0(x2)      # load
I2: fld  f2, 8(x3)      # load
I3: fmul f0, f2, f4     # RAW on f2 (I2)
I4: fsub f8, f2, f6     # RAW on f2 (I2), f6 (I1)
I5: fdiv f10, f0, f6    # RAW on f0 (I3); reads f6
I6: fadd f6, f8, f2     # RAW on f8 (I4); WRITES f6: WAR vs I5,
                        #                            WAW vs I1

Reading the table teaches four things. First, I4 finishes at cycle 8 while I3, older in program order, finishes at 16. Completion is out of order, driven purely by dataflow. Second, the WAR hazard between I6 (writes f6 at cycle 11) and I5 (reads f6, does not execute until 17) is a non-event, because I5 captured the f6 value into its reservation station at issue in cycle 5. The architectural register f6 is irrelevant to it afterward. Third, the WAW between I1 and I6 on f6 resolves because after cycle 6 the RAT maps f6 to I6's station, so I1's broadcast at cycle 4 updated f6 only up to the moment it was re-renamed. The final architectural value is I6's, as program order demands. Fourth, the critical path of the whole block is I2 to I3 to I5, where \(5 + 11 + 41 = 57\) cycles are dictated by true dependences and latencies. Renaming removed everything else.

The reorder buffer and precise exceptions

Pure Tomasulo updates registers the moment the CDB fires, and the table above shows I6 writing architectural f6 at cycle 11, 46 cycles before the older I5 completes. If I5 raises a divide-by-zero exception at cycle 57, the register file already holds the future, because f6 has a value from an instruction that, architecturally, should never have executed. This is an imprecise exception, and it makes demand paging (which must restart a faulting load exactly) and debuggability impossible. The fix, standard since the late 1980s, is the reorder buffer (ROB). Instructions allocate a ROB entry in program order at issue, execute out of order exactly as above, but write only speculative state (the ROB entry or a physical register). A separate commit stage retires instructions from the ROB head strictly in order, making each one architectural only when all older ones have committed. An exception or branch misprediction is handled by draining or flushing. Everything younger than the faulting instruction is discarded before it can commit, so the architectural state seen by the handler is exactly the sequential-semantics state. Modern designs merge the value storage into a unified physical register file with a checkpointed rename map, but the in-order-allocate, out-of-order-execute, in-order-commit skeleton is universal, and ROB capacity (roughly 300 to 600 entries in current high-end cores, with Apple's big cores measured at the top of that range) is the window within which the machine can find independent work to hide a cache miss.

Load-store queues and memory disambiguation

Registers rename cleanly because names are static. Memory does not, because addresses are computed at runtime. A load may not bypass an older store to the same address, but whether the addresses collide is unknown until both are computed. The load-store queue (LSQ) holds all in-flight memory operations in program order. A store writes the data cache only at commit (stores must never be undone). Until then its address and data wait in the store queue. An executing load searches the store queue for older stores. On an address match with data ready, the value is forwarded directly (store-to-load forwarding, and this path is why a store followed by a load of the same location is fast). On a match without data, the load waits. With no match, it goes to the cache. The remaining question is whether a load may issue while an older store's address is still unknown. Conservative designs wait. Aggressive designs predict independence and speculate, checking later when the store's address arrives and squashing the load and its dependents on a violation. Memory dependence predictors (store sets and descendants of Moshovos and Sohi's work) make this speculation right often enough that all high-end cores do it, and this speculation is also the mechanism behind several of the transient-execution security findings, in which the microarchitecture briefly computed with values the architecture would forbid.

Superscalar, SMT, and the end of the free lunch

Wide issue and where ILP runs out

A superscalar machine fetches, decodes, renames, issues, executes, and commits several instructions per cycle. Current high-end cores are 6- to 10-wide at decode with ROBs in the hundreds of entries. The costs grow faster than linearly. Dependence checks among W co-issued instructions are \(O(W^2)\) comparators, the register file needs roughly \(3W\) ports (port count grows area and delay superlinearly), and wakeup-select scheduling logic over a large window is a wire- and power-dense structure whose delay resists pipelining. Meanwhile the returns diminish. Limit studies going back to Wall's 1991 measurements show that real integer programs, under realistic prediction and finite windows, sustain modest usable instruction-level parallelism, with branches and memory latency, not raw width, as the binding constraints. Doubling width without also improving prediction accuracy and the memory hierarchy buys little, which is why the industry's width growth has been slow and always accompanied by predictor and cache growth.

Simultaneous multithreading

If one thread cannot fill an 8-wide machine, several can. Simultaneous multithreading (SMT, Tullsen, Eggers, and Levy's 1995 design, marketed by Intel as Hyper-Threading) duplicates only the architectural state (PCs, rename maps, architectural registers) and lets instructions from multiple threads coexist in the same out-of-order window, sharing the caches, execution ports, and most queues. When thread A stalls on a miss, thread B's instructions issue into the slots A would have wasted. The gains are workload-dependent, with 10 to 30 percent aggregate throughput on mixed server workloads typical for 2-way SMT, while cache-thrashing pairs can produce losses, and shared-state side channels have caused cloud providers to disable SMT for some tenancy models. GPUs, in the section below, take this exact idea to its logical extreme. Instead of 2 threads hiding occasional misses, tens of warps hide latency that is assumed to be everywhere.

Dennard scaling, the power wall, and dark silicon

The reason all of this reorganization happened is thermal. Dynamic power is \(P = \alpha C V^2 f\), activity factor times switched capacitance times voltage squared times frequency. Dennard's 1974 observation was that scaling feature size by \(1/\kappa\) scales C by \(1/\kappa\) and permits V to scale by \(1/\kappa\) while f rises by \(\kappa\). Power per transistor falls as \(1/\kappa^2\) while transistor density rises as \(\kappa^2\), so power density is constant. Chips could get faster and denser at fixed watts, indefinitely. The regime broke around the 90 nm node because voltage stopped scaling. V cannot drop below a few multiples of the threshold voltage without transistors failing to switch, and lowering the threshold instead raises subthreshold leakage exponentially. With V pinned near 1 volt, every further density increase now increases power density, and frequency became the sacrificial variable. Clocks have sat between 3 and 6 GHz since roughly 2005 while transistor budgets kept doubling. The consequences cascade, first multicore (spend transistors on cores, not frequency), then the dark-silicon analysis of Esmaeilzadeh, Blem, St. Amant, Sankaralingam, and Burger (ISCA 2011) showing that at fixed power budgets a growing fraction of a chip must be idle or clocked down at any moment, and finally the domain-specific turn. If transistors are cheap but watts are scarce, the winning move is specialized circuits that do more useful work per joule, lit up only when their workload runs. Horowitz's ISSCC 2014 energy ledger quantifies the pressure at 45 nm. A 32-bit integer add costs about 0.1 pJ, a 32-bit floating multiply 3.7 pJ, a 32 KiB SRAM read 5 pJ, and a DRAM access roughly 640 pJ. Fetching and decoding an instruction to do a 0.1 pJ add costs orders of magnitude more than the add, and moving the operands from DRAM costs four orders of magnitude more. Every architecture on the rest of this page, SIMD, GPUs, and systolic arrays alike, is a strategy for amortizing instruction and data-movement energy over more arithmetic.

Coherence, false sharing, and consistency on multicores

Cache coherence and MESI, with a worked trace

Private caches replicate data, and replication plus writes requires a protocol. Coherence is the single-location contract. For each memory location, there is a total order of writes, every read returns the latest write in that order, and at any moment a line has either one writer or many readers, never both. The invalidation-based MESI protocol enforces this with four stable states per cache line. Modified means this cache holds the only copy, it is dirty, and memory is stale. Exclusive means the only copy, clean, and the crucial optimization is that E can be silently upgraded to M on a write with no bus traffic, so private data is written for free. Shared means possibly one of several clean read copies. Invalid means not present. Controllers snoop the interconnect (or, at core counts beyond a handful, consult a directory that tracks sharers explicitly) and react to each other's requests. A two-core trace on one line X covers every interesting transition.

Step 1 is why E exists. A subsequent private write would cost nothing, whereas MSI would charge an invalidation broadcast for every first write to private data. Step 4 shows the expensive case, a dirty intervention, where the owning cache must supply the line. Extensions address its costs, with MOESI (AMD) adding an Owned state so the dirty line can be shared without an immediate memory writeback, and MESIF (Intel) adding Forward to pick which of several clean sharers responds. Steps 3 and 5 alternating are the pathology called ping-ponging, a line that two cores take turns writing migrates back and forth on every access, each transfer costing a cross-core round trip on the order of 40 to 100 ns, hundreds of times slower than an L1 hit.

False sharing

Coherence operates on whole 64-byte lines, but programs think in variables. If two threads on different cores each update their own counter, and the two counters happen to occupy the same line, the protocol cannot tell. Every increment by one thread invalidates the line in the other's cache, and the trace above runs in a loop at memory-system speed even though there is no logical sharing at all. This is false sharing, and it turns a perfectly parallel loop into one serialized on the coherence fabric. The fix is layout. Pad or align per-thread hot data to line boundaries (alignas(64) in C++, and the standard exposes std::hardware_destructive_interference_size for exactly this), or restructure to thread-local accumulation with a final reduction. The microbenchmark in the implementation section makes the effect measurable in a few lines. The padded and unpadded versions differ only in layout, execute the identical instruction stream per thread, and typically differ by several-fold in wall time on any current multicore part.

Memory consistency, from SC and TSO to weak models

Coherence orders accesses to one location. Consistency defines what orders hold across locations, and it is the contract that lock-free code lives or dies by. Sequential consistency (SC), Lamport's definition, requires that all memory operations appear to interleave in some single total order consistent with each thread's program order. It is what programmers implicitly assume, and essentially no hardware provides it by default, because the easiest performance optimization in a core breaks it, the store buffer. Letting a core continue past a store while the store waits (for the M-state acquisition of its line) means the core's own later loads can complete before its earlier store is globally visible. This yields total store order (TSO), the x86 model, formalized by Sewell, Owens, and colleagues as x86-TSO. Stores are seen by all cores in one order, loads may forward from the local store buffer, and the one reordering visible to software is store-then-load. The litmus test is this. With x and y initially 0, thread A does x = 1 then reads y, and thread B does y = 1 then reads x. Under SC at least one thread must read 1. Under TSO both can read 0, both stores parked in their local buffers, and this outcome is routinely observable on any x86 machine. Arm and RISC-V (RVWMO) go further. They permit load-load, load-store, and store-store reordering too, constrained only by dependences and explicit barriers, which buys hardware simplicity and power at the cost of requiring software to say what it means.

The practical discipline is to program to the language model, not the hardware. C++ and Rust atomics with acquire/release semantics express exactly the orderings needed. A release store makes all prior writes visible to any thread whose acquire load observes it, which is precisely the producer-consumer handoff, while seq_cst additionally enforces a global order among the marked operations themselves (needed for patterns like the litmus test above, and compiled on x86 to an mfence or locked instruction after the store, the one expensive case on TSO). On Arm, acquire/release map to the dedicated ldar/stlr instructions. The classic failure mode is publishing a pointer with a plain store. On a weak model the consumer can observe the pointer before the pointed-to initialization, a bug that will pass every test on an x86 workstation and fire on an Arm server. Data-race- free programs, in which all conflicting accesses are ordered by synchronization, get SC semantics on every model. Everything outside that discipline requires reading Nagarajan, Sorin, Hill, and Wood's primer, which is the reference this section compresses.

SIMD, data parallelism inside the core

The energy ledger said instruction overhead dwarfs arithmetic, and SIMD is the in-core answer. One fetched, decoded, scheduled instruction operates on a whole vector register of lanes. The x86 line widened from 128-bit SSE to 256-bit AVX/AVX2 to 512-bit AVX-512 (16 fp32 or 32 bf16 lanes), the latter adding per-lane predication through mask registers, which lets conditional loops vectorize without branches. AVX-512's early reputation for frequency throttling on 14 nm parts has largely faded on current cores, and its successor AVX10 unifies the feature sets. Arm's NEON is fixed at 128 bits, while SVE and SVE2 are vector-length agnostic. The ISA exposes a width the implementation chooses (128 to 2048 bits), and code written with predicates and increment-by-vector-length runs correctly on any of them, an idea RISC-V's V extension pushes further with a dynamically set vector length register. Fujitsu's A64FX (512-bit SVE) and AWS Graviton and NVIDIA Grace (SVE2 at narrower widths) are the production proof points.

Getting SIMD performance out of a compiler is mostly a matter of not preventing it. Auto-vectorization fails, silently, on possible pointer aliasing (two arrays that might overlap force scalar order, and restrict or __restrict__ removes the doubt), on loop-carried dependences, on early exits, and on floating-point reductions, where reassociating the sum changes rounding and so requires an explicit flag (-ffast-math or OpenMP simd reduction) to permit. The flags -O3 -march=native -fopt-info-vec-missed on GCC, or -Rpass-missed= loop-vectorize on Clang, make the compiler explain itself, and reading that output is the highest-leverage low-effort optimization step that exists. When the compiler cannot be persuaded, intrinsics (the _mm512_* and vld1q_* families) give manual control at the cost of per-ISA code, which is why production kernels in XNNPACK or oneDNN are generated or hand-written per microarchitecture, and libraries like Google's Highway abstract over the ISAs. The correct mental model is that SIMD is cheap ILP with a layout tax. Structure-of-arrays data, unit-stride access, and branch-free inner loops are the price of admission, and paying it is also exactly what the GPU will demand next, at 32 lanes instead of 16.

The GPU as a throughput architecture

Latency machines and throughput machines

A CPU core is a latency machine. It spends its transistor and power budget making one instruction stream finish as soon as possible, via out-of-order windows, branch predictors, and a deep cache hierarchy, all of which are mechanisms for tolerating or avoiding stalls on behalf of a single thread. A GPU inverts every one of those decisions. It assumes tens of thousands of threads are available, strips the per-thread machinery to almost nothing (in-order issue, no speculation, small per-thread cache share), and hides latency not by avoiding it but by switching. When a warp stalls on memory, the scheduler issues another warp, and with enough warps resident the arithmetic units never notice DRAM latency at all. SMT taken to its limit, as promised above. The measured machine for everything in this section is the NVIDIA H100 80GB HBM3 in this repository, with 132 streaming multiprocessors (SMs), 2048 resident threads per SM, 65536 32-bit registers per SM, compute capability 9.0, measured with PyTorch 2.7 and CUDA 12.8.

SMs, warps, and SIMT

The CUDA software hierarchy (threads, blocks of up to 1024 threads, a grid of blocks) maps onto hardware as follows. Blocks are assigned to SMs, and each block's threads are partitioned into warps of 32 that issue in lockstep, one instruction for 32 lanes. This is SIMT, SIMD with a thread-programming face. Each lane has its own registers and may branch independently, but divergent branches within a warp are executed by running both paths with lanes masked off, so a fully divergent warp runs at 1/32 efficiency. Each H100 SM has four scheduler partitions, each choosing one ready warp per cycle among its resident warps. With up to \(2048/32 = 64\) warps resident per SM, the scheduler nearly always finds issuable work. Chip-wide that is \(132 \times 2048 = 270{,}336\) resident threads, against the dozens a CPU socket holds, which is the clearest single number for the design-point difference.

grid ──► blocks ──► warps (32 threads, lockstep issue)

H100 (measured device properties):
  132 SMs
   └─ per SM: ≤ 2048 threads = 64 warps, 4 warp schedulers
              65,536 × 32-bit registers (256 KiB register file)
              L1/shared memory (48 KiB usable per block by default,
              opt-in to ~227 KiB), 4 tensor cores
  chip: 50 MB L2, HBM3 measured at ~2,992 GB/s (copy, fp32)

register file per SM (256 KiB) is LARGER than L1: the GPU's
"architectural state" for 2048 threads is the biggest memory
closest to the ALUs, the exact inverse of a CPU.

Occupancy and latency hiding by arithmetic

Occupancy is resident warps as a fraction of the 64-warp maximum, and it is limited by whichever resource a block exhausts first, registers (65536 per SM, allocated per thread at compile time), shared memory, or the thread and block caps. Its purpose is Little's law. To keep a pipe of bandwidth B busy at latency L, the outstanding traffic must be at least \(B \times L\). At the measured 2992.4 GB/s of HBM bandwidth and roughly 500 ns of DRAM latency, the chip needs \(2992.4\ \text{GB/s} \times 500\ \text{ns} \approx 1.50\) MB in flight at all times, about 11.3 KiB per SM, or 91 outstanding 128-byte lines per SM. A single warp issuing one load at a time contributes 128 bytes. The machine therefore needs on the order of many tens of warps per SM with loads in flight (fewer if each warp has several independent loads unrolled, which is why unrolling and vectorized 128-bit loads matter as much as occupancy). This is the quantitative sense in which oversubscription replaces caches and speculation. The CPU hides one thread's 500 ns behind a predictor and an out-of-order window, while the GPU hides it behind 63 other warps.

The memory hierarchy on an H100, with measured numbers

The hierarchy is explicit. Per-thread registers (65536 32-bit registers per SM) come first, then shared memory/L1 within an SM, programmer- managed, 48 KiB per block by default and up to about 227 KiB by opt-in, at roughly 100x lower latency than DRAM, then a 50 MB L2 shared by all SMs, and finally 80 GB of HBM3. The measured stream-style bandwidths from this machine (h100.json in this repository) are fp32 copy 2992.4 GB/s, fp32 add 3063.5 GB/s, fp32 reduction 2995.2 GB/s, and bf16 copy 2930.3 GB/s, all within a few percent of the nominal 3.35 TB/s HBM spec after refresh and ECC overheads. Measured matmul throughput at n = 8192 is 51.4 TFLOPS in fp32 (no tensor cores), 409.7 in tf32, 728.7 in bf16, and 700.8 in fp16. At n = 4096 bf16 peaks at 744.6 TFLOPS, about 75 percent of the 989 TFLOPS dense bf16 figure in NVIDIA's Hopper whitepaper. Two structural lessons sit in those numbers. First, bf16 tensor cores deliver \(728.7/51.4 = 14.2\times\) the fp32 SIMT datapath, so on this machine "use the tensor cores" is not an optimization but a change of machine. Second, the bf16 advantage over fp32 input at equal multiply count (728.7 vs 409.7 for tf32) is bandwidth and register economics. Half-width operands double the data per register and per byte moved.

The measured attention numbers show the hierarchy acting on a real kernel. Naive attention at sequence length 4096 materializes a 4.29 GB score matrix in HBM and achieves 23.8 TFLOPS. FlashAttention at the same size keeps tiles in shared memory, moves 0.30 GB peak, and achieves 563.0 TFLOPS, a measured 23.6x. At L = 16384 the naive kernel cannot run at all (the 68.7 GB score matrix approaches the 80 GB card) while the tiled kernel proceeds at 640.1 TFLOPS. Same arithmetic, same silicon. The entire difference is which level of the hierarchy the intermediate lives in. That is cache blocking from the matmul section, transplanted to a programmer-managed hierarchy.

Domain-specific accelerators and systolic arrays

Deriving the systolic array

Even a GPU pays instruction overhead. Every tensor-core MMA is still fetched, scheduled, and fed through a register file. If the workload is known to be matrix multiplication, the logical end point is to delete the instructions entirely and lay the dataflow down in wires. Kung's 1982 formulation is that a systolic array is a grid of processing elements (PEs) through which data flows rhythmically, each PE consuming operands from its neighbors, performing one multiply-accumulate, and passing operands along. Memory is touched only at the array edges, and every value fetched once is reused across an entire row or column of PEs, so the data-movement energy per MAC approaches the wire energy between neighbors, and the 640 pJ DRAM number from Horowitz's ledger is amortized over hundreds of operations.

The output-stationary version for \(C = AB\) with 2x2 matrices makes the choreography concrete. PE(i,j) accumulates \(C_{ij}\) in place. Row i of A streams in from the left, delayed by i cycles. Column j of B streams from the top, delayed by j cycles. A values step right and B values step down each cycle. The skew is what aligns \(A_{ik}\) and \(B_{kj}\) at PE(i,j) at the same instant. With \(A = \begin{bmatrix}1&2\\3&4\end{bmatrix}\), \(B = \begin{bmatrix}5&6\\7&8\end{bmatrix}\), the trace runs as follows.

feed (skewed):        B00=5  B01=6            A moves →, B moves ↓
                      B10=7  B11=8 (delayed)

cycle 1: PE(0,0): 1·5 = 5
cycle 2: PE(0,0): 5 + 2·7 = 19 ✓ C00 done
         PE(0,1): 1·6 = 6            (a=1 arrived from PE(0,0))
         PE(1,0): 3·5 = 15           (b=5 arrived from PE(0,0))
cycle 3: PE(0,1): 6 + 2·8 = 22 ✓ C01
         PE(1,0): 15 + 4·7 = 43 ✓ C10
         PE(1,1): 3·6 = 18
cycle 4: PE(1,1): 18 + 4·8 = 50 ✓ C11

C = [[19, 22], [43, 50]], in 3N−2 = 4 cycles for N = 2,
with every input read from memory exactly once.

A single \(N \times N\) product takes \(3N - 2\) cycles, but the fill and drain skew is pipelined away when products stream back to back, so steady-state throughput is \(N^2\) MACs per cycle from \(2N\) input words per cycle, an arithmetic intensity proportional to N by construction, which is the blocking analysis again, now frozen into silicon geometry.

The TPU and the dataflow taxonomy

Google's first Tensor Processing Unit (Jouppi et al., ISCA 2017) is the canonical production instance, a 256x256 weight-stationary systolic array of 8-bit MACs (65,536 PEs) at 700 MHz, giving \(65536 \times 2 \times 700\ \text{MHz} = 92\) TOPS in about 75 W, driven by a small in-order CISC front end over PCIe. Weight-stationary means the roles differ from the trace above. Weights are preloaded into the PEs and sit still, activations stream in from one edge, and partial sums flow through the array into accumulators at the far edge. The paper's honest accounting is as instructive as the design. On real production inference the MXU often sat far below peak because memory bandwidth (34 GB/s DDR3 in TPUv1) starved it, a roofline lesson its successors fixed with HBM.

Which operand sits still is the taxonomy of accelerator design, named by Chen, Emer, and Sze in the Eyeriss work. Weight-stationary (TPU) amortizes weight fetches, ideal when weights are reused across many inputs. Output-stationary (the worked example above) keeps partial sums local, minimizing accumulator traffic, which grows with deep reduction dimensions. Row-stationary, Eyeriss's contribution (MIT, ISCA 2016), maps rows of filter and input so that all three reuse patterns (weights, inputs, partial sums) are exploited jointly within a PE's register file and a NoC, and was derived by explicitly optimizing energy per MAC over the loop-nest mapping space rather than picking a dataflow by intuition. The general lesson, elaborated by the Eyeriss group's later Timeloop-style mapping tools and by accelerator generators like Berkeley's Gemmini and ETH Zurich's Snitch/Occamy clusters, is that an accelerator is a loop nest with a floorplan. Choosing tiling, ordering, and spatial unrolling of the same six-deep convolution loops determines every property of the machine.

Sparsity and low precision in hardware

Two further levers multiply accelerator efficiency, both by shrinking the work rather than speeding it. The first is low precision. Energy and area of a multiplier scale roughly quadratically with mantissa width, so int8 or fp8 arithmetic is an order of magnitude cheaper than fp32, and the measured 14.2x tensor-core advantage above already includes this effect. Hopper's fp8 formats (E4M3 for forward values, E5M2 for gradients) double bf16 throughput again, and Blackwell extends to block-scaled fp4 with per-32-element scale factors, pushing precision management into the hardware. The second is sparsity. NVIDIA's 2:4 structured sparsity (at most 2 nonzeros in every 4 weights, enforced at pruning time) lets the tensor core skip half the multiplies with a tiny metadata index, doubling peak, and is the pragmatic middle ground between dense hardware and the fully sparse accelerators (MIT's Eyeriss v2 and the SCNN line from NVIDIA research) that gate or skip zero operands dynamically at finer granularity but pay for it in control complexity and load imbalance. Both levers work exactly because the customer workload, deep networks, tolerates aggressive quantization and pruning, a co-design freedom that general-purpose architecture never had.

The roofline model

Derivation and the machine balance of a measured H100

The roofline model (Williams, Waterman, and Patterson, CACM 2009) is the two-parameter summary of every machine on this page. A kernel that performs F flops and moves Q bytes to and from memory has arithmetic intensity \(I = F/Q\) flops per byte. The machine has peak compute \(P_{peak}\) (flops/s) and peak memory bandwidth \(B\) (bytes/s). Execution time is at least the larger of compute time \(F/P_{peak}\) and transfer time \(Q/B\) (assuming perfect overlap), so attainable performance is

Plotted on log-log axes against I, this is a rising line of slope 1 (the memory roof) meeting a horizontal line (the compute roof) at the ridge point \(I^* = P_{peak}/B\), the machine balance. Kernels with \(I < I^*\) are memory-bound and kernels with \(I > I^*\) are compute-bound, and no kernel may sit above either roof. For the measured H100 numbers, using bf16 tensor-core peak as measured at n = 8192 (728.7 TFLOPS) and measured copy bandwidth (2992.4 GB/s),

and for the non-tensor fp32 datapath, \(51.4 \times 10^{12} / 2992.4 \times 10^9 \approx 17.2\) flops/byte. These two numbers explain most H100 performance mysteries. Any elementwise operation has \(I\) below 1 and is doomed to under one percent of tensor-core peak regardless of engineering effort. The only fix is to stop going to memory, i.e. fusion. Big matmuls have \(I\) in the thousands and are limited purely by the compute roof. Attention sits in between and is movable. The algorithm's flops are fixed but its Q depends on whether the score matrix spills to HBM, so FlashAttention is precisely an intensity-raising transformation, sliding the kernel rightward along the roofline from the memory-bound region toward the ridge. The measured numbers used throughout the worked problem below confirm each regime on this machine.

TFLOPS (log)                      H100 80GB, measured
 728.7 ┤· · · · · · · · · ·┌────────────────────────
       │        compute roof│   ▲ matmul bf16 n=8192 (728.7)
       │                    │  ▲ flash attn L=16384 (640.1)
       │      memory roof   │
       │      slope =       │▲ flash attn L=512 (170.2)
       │      2992.4 GB/s  ╱│
  51.4 ┤· · · · · · · ·╱· ·│· fp32 matmul roof
       │             ╱      │
       │           ╱ ▲ naive attn L=4096 (23.8: score matrix
       │         ╱     round-trips through HBM)
  0.25 ┤· · ·╱▲ fp32 add (I = 1/12: bound at 0.25 TFLOPS)
       │   ╱
       └──┬───────┬─────────┬──────────┬────────── I (flops/byte)
        1/12      17.2     243.5      2731
                 (fp32     (bf16      (bf16 matmul
                  ridge)    ridge)     n=8192)

The model's limits are worth stating as plainly as its power. It assumes perfect overlap of compute and transfer, counts only one bandwidth (extensions add rooflines for L2, shared memory, and NVLink), and says nothing about latency-bound kernels with too little parallelism to reach either roof, which is where Little's law from the GPU section takes over. The measured n = 1024 bf16 matmul makes the point. Its algorithmic intensity (about 341 flops/byte) predicts compute-bound operation near peak, yet the measurement is 108.9 TFLOPS, 15 percent of peak, because a 1024-cube matmul simply does not contain enough parallel tile-work to fill 132 SMs and hide its own dependencies. Roofline bounds what is attainable. It does not promise attaining it.

Worked problems

Implementation

Cache blocking, measured in C

The first benchmark realizes the blocking analysis, identical arithmetic, reordered. The naive version walks B down a column with stride N, touching a new cache line per element. The blocked version keeps 48x48 tiles resident and makes the inner loop unit-stride, which also hands the compiler a vectorizable loop. Compile with gcc -O2 -o gemm gemm.c (deliberately -O2 without -march=native, to measure the memory effect rather than the vectorizer). At N = 1024 the blocked version is typically 3 to 8 times faster on current x86 parts, and the gap widens with N as the naive version falls out of successive cache levels.

// gemm.c: naive vs cache-blocked sgemm, single-threaded.
// build: gcc -O2 -o gemm gemm.c   run: ./gemm
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>

#define N  1024
#define BS 48              // 3 * 48*48 * 4 B = 27 KiB: fits a 32 KiB L1

static float A[N*N], B[N*N], C[N*N];   // row-major, C = A * B

static double now_s(void) {
    struct timespec ts;
    clock_gettime(CLOCK_MONOTONIC, &ts);
    return ts.tv_sec + 1e-9 * ts.tv_nsec;
}

static void matmul_naive(void) {
    memset(C, 0, sizeof C);
    for (int i = 0; i < N; i++)
        for (int j = 0; j < N; j++) {
            float acc = 0.0f;
            for (int k = 0; k < N; k++)          // B[k*N+j]: stride-N walk,
                acc += A[i*N + k] * B[k*N + j];  // one cache line per element
            C[i*N + j] = acc;
        }
}

static void matmul_blocked(void) {
    memset(C, 0, sizeof C);
    for (int ii = 0; ii < N; ii += BS)
      for (int kk = 0; kk < N; kk += BS)         // tiles of A, B, C stay hot
        for (int jj = 0; jj < N; jj += BS)
          for (int i = ii; i < ii + BS; i++)
            for (int k = kk; k < kk + BS; k++) {
                float a = A[i*N + k];            // scalar hoisted to register
                for (int j = jj; j < jj + BS; j++)
                    C[i*N + j] += a * B[k*N + j];  // unit stride: line fully used
            }
}

int main(void) {
    for (int i = 0; i < N*N; i++) {
        A[i] = (float)rand() / RAND_MAX;
        B[i] = (float)rand() / RAND_MAX;
    }
    double flops = 2.0 * N * N * N;              // 2n^3 = 2.147e9 flops

    double t0 = now_s(); matmul_naive();   double tn = now_s() - t0;
    t0 = now_s();        matmul_blocked(); double tb = now_s() - t0;

    printf("naive:   %.3f s  %.2f GFLOP/s\n", tn, flops / tn / 1e9);
    printf("blocked: %.3f s  %.2f GFLOP/s  (%.2fx)\n",
           tb, flops / tb / 1e9, tn / tb);
    return 0;
}

A runnable false-sharing microbenchmark

The second benchmark makes coherence traffic visible from user space. Two threads increment disjoint counters. The only variable is whether the counters share a cache line. The counters are volatile so the compiler must issue a store per iteration rather than accumulating in a register (which would measure nothing). Build with gcc -O2 -pthread -o fs false_sharing.c. Expect the shared-line case to run several times slower, per Problem 8's arithmetic, and try perf stat -e cache-misses on Linux to watch the machinery directly.

// false_sharing.c: measure cache-line ping-pong between two cores.
// build: gcc -O2 -pthread -o fs false_sharing.c   run: ./fs
#include <pthread.h>
#include <stdio.h>
#include <stdint.h>
#include <time.h>

#define ITERS 100000000UL

struct shared_line {                 // both counters in ONE 64 B line
    volatile uint64_t a;
    volatile uint64_t b;
};
struct padded {                      // one counter per line
    _Alignas(64) volatile uint64_t a;
    _Alignas(64) volatile uint64_t b;
};

static void *bump(void *p) {
    volatile uint64_t *ctr = (volatile uint64_t *)p;
    for (uint64_t i = 0; i < ITERS; i++)
        (*ctr)++;                    // load + add + store, every iteration
    return NULL;
}

static double run_pair(volatile uint64_t *x, volatile uint64_t *y) {
    struct timespec t0, t1;
    pthread_t ta, tb;
    clock_gettime(CLOCK_MONOTONIC, &t0);
    pthread_create(&ta, NULL, bump, (void *)x);
    pthread_create(&tb, NULL, bump, (void *)y);
    pthread_join(ta, NULL);
    pthread_join(tb, NULL);
    clock_gettime(CLOCK_MONOTONIC, &t1);
    return (t1.tv_sec - t0.tv_sec) + 1e-9 * (t1.tv_nsec - t0.tv_nsec);
}

int main(void) {
    static struct shared_line s;     // s.a and s.b: same line
    static struct padded p;          // p.a and p.b: different lines

    double t_shared = run_pair(&s.a, &s.b);
    double t_padded = run_pair(&p.a, &p.b);

    printf("same line:      %.3f s\n", t_shared);
    printf("padded (64 B):  %.3f s\n", t_padded);
    printf("slowdown from false sharing: %.2fx\n", t_shared / t_padded);
    return 0;
}

The GPU hierarchy in CUDA, naive versus tiled matmul

The CUDA version of blocking targets the programmer-managed level of the hierarchy. Each block stages 32x32 tiles of A and B into shared memory, synchronizes, and computes 32 multiply-adds per element loaded, cutting global-memory traffic by the tile width. The naive kernel's global loads of B are coalesced (adjacent threadIdx.x reads adjacent columns) but every element of A and B is re-fetched from L2 or HBM once per output element it touches. Neither kernel approaches cuBLAS, which adds register blocking, double-buffered asynchronous copies, and tensor cores on top of exactly this structure. The point is the traffic ratio, not peak.

// tiled.cu: shared-memory tiling on the GPU. build: nvcc -O3 -o tiled tiled.cu
#include <cstdio>
#include <cuda_runtime.h>

#define TILE 32                       // 32x32 tile = one warp-wide square

__global__ void sgemm_naive(const float *A, const float *B, float *C, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;   // adjacent x -> adjacent col:
    if (row >= n || col >= n) return;                  // loads of B, C coalesce
    float acc = 0.0f;
    for (int k = 0; k < n; ++k)
        acc += A[row * n + k] * B[k * n + col];        // 2 global reads per FMA
    C[row * n + col] = acc;
}

__global__ void sgemm_tiled(const float *A, const float *B, float *C, int n) {
    __shared__ float As[TILE][TILE];                   // staged in shared memory:
    __shared__ float Bs[TILE][TILE];                   // ~20-30 cycle latency vs
    int row = blockIdx.y * TILE + threadIdx.y;         // ~500 ns for HBM
    int col = blockIdx.x * TILE + threadIdx.x;
    float acc = 0.0f;
    for (int t = 0; t < n / TILE; ++t) {
        As[threadIdx.y][threadIdx.x] = A[row * n + t * TILE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE + threadIdx.y) * n + col];
        __syncthreads();                               // tile fully staged
        #pragma unroll
        for (int k = 0; k < TILE; ++k)                 // 32 FMAs per element
            acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];  // loaded: AI x32
        __syncthreads();                               // done before overwrite
    }
    C[row * n + col] = acc;
}

int main() {
    const int n = 4096;
    size_t bytes = (size_t)n * n * sizeof(float);
    float *A, *B, *C;                                  // (n, n) each
    cudaMalloc(&A, bytes); cudaMalloc(&B, bytes); cudaMalloc(&C, bytes);

    dim3 block(TILE, TILE), grid(n / TILE, n / TILE);
    cudaEvent_t t0, t1; cudaEventCreate(&t0); cudaEventCreate(&t1);
    float ms; double flops = 2.0 * n * n * (double)n;

    sgemm_naive<<<grid, block>>>(A, B, C, n);          // warm up
    cudaEventRecord(t0);
    sgemm_naive<<<grid, block>>>(A, B, C, n);
    cudaEventRecord(t1); cudaEventSynchronize(t1);
    cudaEventElapsedTime(&ms, t0, t1);
    printf("naive: %7.2f ms  %6.1f GFLOP/s\n", ms, flops / ms / 1e6);

    sgemm_tiled<<<grid, block>>>(A, B, C, n);
    cudaEventRecord(t0);
    sgemm_tiled<<<grid, block>>>(A, B, C, n);
    cudaEventRecord(t1); cudaEventSynchronize(t1);
    cudaEventElapsedTime(&ms, t0, t1);
    printf("tiled: %7.2f ms  %6.1f GFLOP/s\n", ms, flops / ms / 1e6);
    return 0;
}

A roofline measurement harness

The harness below reproduces the two coordinates of the roofline on whatever GPU it runs on, streaming bandwidth from a large fp32 add and compute peak from a large bf16 matmul, then derives the machine balance and places both kernels. Run on the H100 in this repository, it produces the numbers quoted throughout this page (3063.5 GB/s for the add, 728.7 TFLOPS at n = 8192, balance 243.5 flops/byte). The same script on a different part maps that machine instead. Note the two essentials of honest GPU timing, warmup iterations to exclude compilation and clock ramp, and a device synchronize before reading the host clock, since kernel launches are asynchronous.

import time
import torch

assert torch.cuda.is_available()
dev = torch.device("cuda")

def time_fn(f, iters=50, warmup=10):
    for _ in range(warmup):
        f()
    torch.cuda.synchronize()               # launches are async: fence first
    t0 = time.perf_counter()
    for _ in range(iters):
        f()
    torch.cuda.synchronize()
    return (time.perf_counter() - t0) / iters

# 1) memory roof: fp32 streaming add, 1 flop per 12 bytes
n = 1 << 28                            # 2^28 elems = 1 GiB per tensor
x = torch.randn(n, device=dev)         # (n,)
y = torch.randn(n, device=dev)         # (n,)
z = torch.empty_like(x)                # (n,)
t = time_fn(lambda: torch.add(x, y, out=z))
gbs = 3 * 4 * n / t / 1e9              # read x, read y, write z
print(f"add : {gbs:7.1f} GB/s   AI = 1/12 flop/B")

# 2) compute roof: bf16 matmul on tensor cores
m = 8192
a = torch.randn(m, m, device=dev, dtype=torch.bfloat16)  # (m, m)
b = torch.randn(m, m, device=dev, dtype=torch.bfloat16)  # (m, m)
t = time_fn(lambda: a @ b)
tflops = 2 * m**3 / t / 1e12
ai_mm = 2 * m**3 / (3 * m * m * 2)     # flops / min bytes (bf16 = 2 B)
print(f"gemm: {tflops:7.1f} TFLOPS  AI = {ai_mm:.0f} flop/B")

# 3) machine balance and kernel placement
balance = tflops * 1e12 / (gbs * 1e9)
print(f"machine balance ~ {balance:.1f} flop/B")
print(f"add is memory-bound:  bound = {gbs/12:.1f} GFLOP/s")
print(f"gemm is compute-bound ({ai_mm/balance:.1f}x past the ridge)")

import time
import jax
import jax.numpy as jnp

key = jax.random.PRNGKey(0)

def time_fn(f, iters=50, warmup=10):
    for _ in range(warmup):
        out = f()
    jax.block_until_ready(out)             # dispatch is async: fence first
    t0 = time.perf_counter()
    for _ in range(iters):
        out = f()
    jax.block_until_ready(out)
    return (time.perf_counter() - t0) / iters

# 1) memory roof: fp32 streaming add, 1 flop per 12 bytes
n = 1 << 28                            # 2^28 elems = 1 GiB per tensor
k1, k2, k3, k4 = jax.random.split(key, 4)
x = jax.random.normal(k1, (n,))        # (n,)
y = jax.random.normal(k2, (n,))        # (n,)
add = jax.jit(lambda x, y: x + y)
t = time_fn(lambda: add(x, y))
gbs = 3 * 4 * n / t / 1e9              # read x, read y, write z
print(f"add : {gbs:7.1f} GB/s   AI = 1/12 flop/B")

# 2) compute roof: bf16 matmul on tensor cores
m = 8192
a = jax.random.normal(k3, (m, m), dtype=jnp.bfloat16)  # (m, m)
b = jax.random.normal(k4, (m, m), dtype=jnp.bfloat16)  # (m, m)
mm = jax.jit(lambda a, b: a @ b)
t = time_fn(lambda: mm(a, b))
tflops = 2 * m**3 / t / 1e12
ai_mm = 2 * m**3 / (3 * m * m * 2)     # flops / min bytes (bf16 = 2 B)
print(f"gemm: {tflops:7.1f} TFLOPS  AI = {ai_mm:.0f} flop/B")

# 3) machine balance and kernel placement
balance = tflops * 1e12 / (gbs * 1e9)
print(f"machine balance ~ {balance:.1f} flop/B")
print(f"add is memory-bound:  bound = {gbs/12:.1f} GFLOP/s")
print(f"gemm is compute-bound ({ai_mm/balance:.1f}x past the ridge)")

How it is done in practice

What a shipping CPU core actually contains

The distance between the five-stage teaching pipeline and a shipped high-end core is roughly a factor of a thousand in engineering effort, but the shipped core is recognizably the same drawing with every box grown and replicated. A current big core (Intel's performance cores, AMD Zen 5, Apple's M-series big cores, Arm's X-series) decodes 6 to 10 instructions per cycle, holds an out-of-order window of roughly 300 to 600 plus reorder-buffer entries backed by physical register files with several hundred registers of each class, sustains 6 to 8 micro-ops issued per cycle across a dozen or so execution ports, and predicts with a TAGE-class direction predictor plus BTB hierarchies and a return stack, at a few mispredictions per kilo-instruction. The cache hierarchy is 32 to 128 KiB of L1 data, roughly 1 to 3 MiB of private L2, and a shared L3 of a few MiB per core, with several concurrent hardware prefetchers (next-line, stride, region-based) that in practice hide a large fraction of what would otherwise be DRAM stalls. Two facts about the process rarely make it into textbooks. First, verification, not design, is the dominant cost. Confirming that renaming, the LSQ, coherence, and the consistency model compose correctly across corner cases consumes more engineer-years than the RTL itself, which is a large part of why formal memory-model specifications (the x86-TSO and Arm axiomatic models, RVWMO for RISC-V) migrated from academia into vendor documentation. Second, the microarchitecture is increasingly a security surface. Spectre class attacks weaponized branch prediction and speculative loads, and every core since has shipped with predictor isolation modes and speculation barriers that architects must now budget for like any other structure.

The GPU software stack and measured production kernels

On the GPU side, the practice gap is bridged by kernel libraries and compilers rather than by hardware alone. cuBLAS and CUTLASS implement the tiled matmul of the implementation section with two more levels of blocking (registers below shared memory, and warp-level tensor-core tiles below that), double-buffered asynchronous copies so tiles stream while arithmetic proceeds, and autotuned tile shapes per problem size. That stack is what turns the naive kernel's few percent of peak into the measured 728.7 bf16 TFLOPS. The measured transformer block numbers from this repository's H100 show what production inference arithmetic looks like end to end. A full multi-head attention block at batch 8, sequence 2048, model width 1024, 16 heads runs in 0.503 ms at 546.6 TFLOPS, with the four projection matmuls contributing half the flops, and the fused-versus-unfused elementwise measurement (0.911 ms to 0.214 ms, 4.26x) is precisely the win that torch.compile, XLA, and Triton chase automatically across every non-matmul op in a model. The division of labor has stabilized. Matmuls go to vendor libraries or CUTLASS-generated kernels, attention goes to purpose-built fused kernels, and the long tail of normalization, activation, and reshaping is fused by compilers, with the roofline as the shared accounting system deciding which bucket each op belongs in.

Fleet economics

At datacenter scale the objective function is performance per watt per dollar over a fleet, not single-kernel latency, and that objective explains the current hardware landscape better than any benchmark chart. Hyperscalers design in-house silicon (Google's TPU line, Amazon's Graviton CPUs and Trainium accelerators, Microsoft's Maia, Meta's MTIA) not to beat merchant parts at peak flops but to delete the margins and tune the memory hierarchy for their actual traffic. The TPU paper's cost-performance framing, rather than its architecture, is its most imitated feature. Power delivery and cooling now bound accelerator deployment (a single H100-class part dissipates 700 W, a training pod megawatts), so low-precision arithmetic is as much a facilities decision as a numerical one. The consistent through-line from the Horowitz ledger to a training cluster invoice is that architecture at every scale is the art of not moving data.

The current research frontier

Packaging as the new scaling axis

With transistor cost improvements slowing, the last four years of headroom have come mostly from packaging. Chiplet designs (AMD's CPU line since Zen 2, Intel's tiled Xeons, NVIDIA's two-die Blackwell) split a system across smaller dies with better yield, joined by dense interconnect bridges. AMD's 3D V-Cache stacks an extra 64 MiB of SRAM directly on top of the compute die, a pure exploitation of the AMAT arithmetic above. And 2.5D CoWoS-style integration of HBM stacks beside the GPU die is the enabling technology behind the measured 3 TB/s on this page's H100. The UCIe consortium is standardizing die-to-die links so chiplets from different vendors can compose, and CXL extends coherence over PCIe physical layers to enable memory pooling and tiering across a rack, which drags the consistency-model questions of the multicore section out of the socket and into the datacenter.

Accelerator diversity and near-memory computing

The accelerator space is unusually pluralistic right now. Cerebras builds wafer-scale engines (the WSE-3 spans an entire 300 mm wafer, with on-wafer SRAM standing in for HBM). Groq's LPU removes caches, branch prediction, and dynamic scheduling entirely in favor of compiler-scheduled, deterministic dataflow, a bet that ML inference is regular enough to software-pipeline the whole chip. Tenstorrent combines RISC-V cores with matrix units and bets on open ISAs. Academic generators keep the design space honest. Berkeley's Gemmini produces systolic accelerators parameterized over the dataflow taxonomy above. ETH Zurich's Snitch and Occamy clusters explore tiny cores marshalling wide FPUs. MIT's Eyeriss v2 and the Timeloop/Accelergy tooling from the same group turned dataflow choice into a searchable optimization problem. Cornell and Tsinghua groups, among others, push high-level synthesis and processing-in-memory prototypes respectively. Near-memory computing has finally shipped in small volumes (UPMEM's DRAM-integrated processors, Samsung's HBM-PIM prototypes), attacking the 640 pJ DRAM-access line of the energy ledger directly, though programming models remain the obstacle.

Precision, sparsity, and learned microarchitecture

Numerics keep descending. Training in fp8 is production practice on Hopper-class hardware, the OCP microscaling (MX) formats standardize block-scaled 4- and 6-bit types across vendors, and Blackwell implements them in silicon. The open question, actively contested across industrial and academic labs, is where training stops tolerating quantization noise, with per-block scaling and stochastic rounding as the current frontier tools. Structured sparsity beyond 2:4 remains an arms race between pruning algorithms and hardware generality. On the CPU side, research keeps probing learned components. Perceptron and TAGE hybrids already ship, ML-guided prefetching and cache replacement appear in top venues yearly, and the practical question is whether learned structures can meet timing and verification budgets rather than whether they can predict. And the RISC-V wave is reaching servers, with vector extension (RVV 1.0) silicon, profiles standardizing feature sets, and multiple vendors sampling out-of-order application cores, which will test whether an open ISA can sustain the software-ecosystem network effects that x86 and Arm built over decades.

Open source to read

Architecture is unusually well served by readable production code. Each entry gives the file to open first.

• llvm-project, where ISA design meets the compiler. Open llvm/lib/Target/RISCV/RISCVInstrInfo.td to see an entire ISA expressed as structured data (encodings, scheduling info, compression patterns), then llvm/lib/CodeGen/MachineScheduler.cpp to see software instruction scheduling against a machine model, the compiler's answer to the hazards on this page.
• riscv/riscv-isa-manual, the ISA specification itself, maintained as a living document. Open the RV32I base chapter (src/rv32.adoc) and read the design commentary boxes. They are a running seminar on why each encoding decision was made.
• gem5, the standard cycle-level research simulator. Open src/cpu/o3/rename.cc to see register renaming implemented exactly as the Tomasulo section describes, then wander into the LSQ and commit stages in the same directory. This is the closest thing to an executable version of this page.
• chipsalliance/rocket-chip, a real, taped-out RISC-V core generator in Chisel. Open src/main/scala/rocket/RocketCore.scala and find the five-stage pipeline, the hazard logic, and the bypass muxes as a few hundred lines of parameterized hardware.
• verilator, the open-source Verilog simulator that compiles RTL to C++. Open examples/make_tracing_c/sim_main.cpp to see the simulation driver pattern, then simulate a small core (rocket or a picorv32) and watch pipeline registers change cycle by cycle.
• NVIDIA/cutlass, the open implementation of GPU matmul at full performance. Open media/docs/efficient_gemm.md first. It is the best short document on hierarchical tiling available, and maps one-to-one onto the blocking and roofline sections above.
• triton-lang/triton, the tile-level GPU kernel language behind much of PyTorch 2. Open python/tutorials/03-matrix-multiplication.py and compare it line by line with the CUDA tiled kernel above. The tile abstraction is the same, with the compiler owning shared memory and scheduling.
• google/XNNPACK, production CPU SIMD microkernels for neural inference. Open any AVX-512 GEMM microkernel under src/f32-gemm/gen/ to see register blocking and per-ISA specialization. The directory structure itself is a map of the SIMD landscape.
• ARM-software/ComputeLibrary, Arm's counterpart, with NEON and SVE kernels for the same workloads. Open examples/neon_sgemm.cpp and trace the call down into the kernel layer to see how a vendor library organizes microkernel dispatch across microarchitectures.

Common misconceptions

"Clock frequency measures performance." The iron law is time = instructions x CPI x cycle time, and all three factors trade against each other. Deep pipelines raise frequency while raising branch penalties and hence CPI, and ISAs differ in instruction counts for the same work. A 3.5 GHz core with IPC 4 beats a 5 GHz core with IPC 2 by 40 percent. Frequency comparisons are meaningful only within a single microarchitecture.

"x86 is doomed by CISC decode, and Arm wins because RISC is faster." The execution engines of high-end x86 and Arm cores are nearly identical out-of-order machines, and x86 pays a real but bounded front-end tax (predecode, micro-op caches), a few percent of core power at the high end. Apple's efficiency advantage came from process leadership, very large caches and windows, and tight system integration far more than from instruction encoding. At the low end the decode tax matters proportionally more, which is where RISC's win is real.

"Pipelining makes instructions execute faster." Each instruction gets slower, the same logic plus per-stage register overhead, rounded up to the slowest stage. Only throughput improves, as Problem 1's arithmetic showed (latency 1050 to 1750 ps, throughput 3x). Any claim about pipelining should survive substituting the words "more instructions per second" for "faster".

"The cache is transparent, so software cannot do much about it." Transparent means correct, not performance-neutral. Blocking changed the matmul's memory traffic from \(O(n^3)\) to \(O(n^3/b)\) without changing a single arithmetic operation, an asymptotic improvement available purely in software, and Hong and Kung's bound says it is the best possible. Most performance engineering on CPUs is exactly this, restructuring computation so the fixed hardware sees locality.

"The hardware runs my program in the order I wrote it." The compiler reorders, the out-of-order core reorders, and the memory system makes some reorderings visible to other cores. The store-buffer litmus test yields the "impossible" result on every x86 machine, and weaker models permit far more. Single-threaded code cannot tell (that is the point of precise state), but any cross-thread communication outside the atomics discipline is a bug in waiting, typically discovered when x86 code first runs on Arm.

"GPUs are fast because they have thousands of cores." The deeper truth is that GPUs are fast because they stopped paying for latency. They carry no large per-thread caches, predictors, or windows, memory latency is hidden by warp oversubscription, and the saved transistors are spent on arithmetic and bandwidth. A single GPU thread is far slower than a CPU thread. The machine wins only when tens of thousands of threads exist. "Thousands of slow threads, scheduled for free" explains real GPU performance behavior, while "thousands of cores" predicts none of it.

"Maximize occupancy to maximize GPU performance." Occupancy is a means (enough warps to satisfy Little's law), not an end. Past the point where latency is covered, more resident warps shrink each thread's register budget, forcing spills. The fastest matmul kernels run at modest occupancy with enormous per-thread register tiles, as Problem 5's trade-off showed. Chase coverage of stalls, not the occupancy percentage.

"More TFLOPS means faster machine learning." Peak flops bound only kernels right of the ridge point. On the measured H100 the balance is 243.5 flops/byte. Every elementwise op in a model sits below intensity 1 and is bound by the 3 TB/s memory system at under one percent of peak, and naive attention threw away a measured 23.6x by materializing its score matrix. For most workloads bandwidth and kernel structure, not peak arithmetic, decide the wall clock, which is why HBM capacity and bandwidth, not TFLOPS, are the scarce commodities in the accelerator market.

Self-check

References

1. Hennessy, J. and Patterson, D. Computer Architecture: A Quantitative Approach, 6th ed., Morgan Kaufmann, 2017.
2. Patterson, D. and Hennessy, J. Computer Organization and Design: The Hardware/Software Interface, RISC-V Edition, 2nd ed., Morgan Kaufmann, 2020.
3. Harris, S. and Harris, D. Digital Design and Computer Architecture, RISC-V Edition, Morgan Kaufmann, 2021.
4. Shen, J. P. and Lipasti, M. Modern Processor Design: Fundamentals of Superscalar Processors, Waveland Press, 2013 (orig. 2005).
5. Nagarajan, V., Sorin, D., Hill, M., and Wood, D. A Primer on Memory Consistency and Cache Coherence, 2nd ed., Morgan & Claypool Synthesis Lectures, 2020.
6. Sze, V., Chen, Y.-H., Yang, T.-J., and Emer, J. "Efficient Processing of Deep Neural Networks: A Tutorial and Survey," Proceedings of the IEEE, 2017. arXiv:1703.09039
7. Tomasulo, R. M. "An Efficient Algorithm for Exploiting Multiple Arithmetic Units," IBM Journal of Research and Development 11(1), 1967. doi:10.1147/rd.111.0025
8. Smith, J. E. "A Study of Branch Prediction Strategies," ISCA, 1981.
9. Yeh, T.-Y. and Patt, Y. "Two-Level Adaptive Training Branch Prediction," MICRO-24, 1991.
10. Seznec, A. and Michaud, P. "A Case for (Partially) TAgged GEometric History Length Branch Prediction," Journal of Instruction-Level Parallelism, vol. 8, 2006.
11. Hill, M. and Smith, A. J. "Evaluating Associativity in CPU Caches," IEEE Transactions on Computers 38(12), 1989.
12. Hong, J.-W. and Kung, H. T. "I/O Complexity: The Red-Blue Pebble Game," STOC, 1981.
13. Lamport, L. "How to Make a Multiprocessor Computer That Correctly Executes Multiprocess Programs," IEEE Transactions on Computers C-28(9), 1979.
14. Adve, S. and Gharachorloo, K. "Shared Memory Consistency Models: A Tutorial," IEEE Computer 29(12), 1996.
15. Sewell, P., Sarkar, S., Owens, S., Zappa Nardelli, F., and Myreen, M. "x86-TSO: A Rigorous and Usable Programmer's Model for x86 Multiprocessors," CACM 53(7), 2010. doi:10.1145/1785414.1785443
16. Tullsen, D., Eggers, S., and Levy, H. "Simultaneous Multithreading: Maximizing On-Chip Parallelism," ISCA, 1995.
17. Dennard, R. et al. "Design of Ion-Implanted MOSFETs with Very Small Physical Dimensions," IEEE Journal of Solid-State Circuits 9(5), 1974.
18. Esmaeilzadeh, H., Blem, E., St. Amant, R., Sankaralingam, K., and Burger, D. "Dark Silicon and the End of Multicore Scaling," ISCA, 2011.
19. Horowitz, M. "Computing's Energy Problem (and What We Can Do About It)," ISSCC, 2014.
20. Kung, H. T. "Why Systolic Architectures?" IEEE Computer 15(1), 1982.
21. Jouppi, N. et al. "In-Datacenter Performance Analysis of a Tensor Processing Unit," ISCA, 2017. arXiv:1704.04760
22. Chen, Y.-H., Emer, J., and Sze, V. "Eyeriss: A Spatial Architecture for Energy-Efficient Dataflow for Convolutional Neural Networks," ISCA, 2016.
23. Williams, S., Waterman, A., and Patterson, D. "Roofline: An Insightful Visual Performance Model for Multicore Architectures," CACM 52(4), 2009. doi:10.1145/1498765.1498785
24. Genc, H. et al. "Gemmini: Enabling Systematic Deep-Learning Architecture Evaluation via Full-Stack Integration," DAC, 2021. arXiv:1911.09925
25. Waterman, A. and Asanović, K. (eds.) The RISC-V Instruction Set Manual. github.com/riscv/riscv-isa-manual
26. NVIDIA. NVIDIA H100 Tensor Core GPU Architecture (Hopper whitepaper), 2022. AMD, CDNA 3 Architecture whitepaper, 2023.
27. Hooker, S. "The Hardware Lottery," CACM 64(12), 2021. arXiv:2009.06489