How AI Accelerator Architecture Actually Works

The instruction that is actually thousands of instructions

A framework issues one call — matmul, or a fused kernel underneath it — and a researcher reading the trace sees a single line. Nothing about that line is single. Underneath it, an accelerator is moving thousands of operand values into a grid of arithmetic cells, holding some of them still, streaming others past, and combining partial results in a strict cycle-by-cycle order that has to be exactly right or the answer is wrong. The earlier piece in this series established that an AI accelerator is best understood as a memory system with arithmetic attached, and organised that argument around the roofline model: peak arithmetic rate, achievable bandwidth, and the ridge point at which one stops being the limit and the other starts ^[7]. That framing answers whether a kernel can reach a chip’s peak rate. It does not answer the question this article is about: what physically happens, cell by cell and cycle by cycle, when it does.

That mechanism has a name older than the transformer, older than the GPU, older than the term “AI accelerator” itself. It is the systolic array, and almost every matrix engine shipping today — Google’s Tensor Processing Unit, the tensor cores inside an NVIDIA GPU, the mobile inference blocks inside a phone’s neural processor — is a variation on an idea published in 1982 for a completely different generation of hardware problems.

Why pump data through a grid at all

H. T. Kung’s original argument was about cost, not about neural networks, which did not yet exist as a hardware target. Special-purpose chips of the era were expensive to design and built in small volumes, and the bottleneck was rarely the arithmetic — it was getting data in and out of memory fast enough to keep custom logic busy. Kung’s proposed fix was to arrange simple processing cells in a regular array and pump data through the array from multiple directions at once, so that each value fetched from memory is used by many cells before it is discarded, rather than fetched once per use ^[1]. The name comes from that pumping: data moves through the array the way blood is pumped through tissue, in a rhythm, reaching every cell without a separate trip back to the source each time.

Google’s first Tensor Processing Unit is the cleanest large-scale instance of that idea built for neural-network arithmetic. Its Matrix Multiply Unit is a 256×256 grid of 8-bit multiply-accumulate cells — 65,536 of them — and the paper’s authors are explicit about why: “as reading a large SRAM uses much more power than arithmetic, the matrix unit uses systolic execution to save energy by reducing reads and writes” of the on-chip buffer ^[2]. The mechanism they describe is worth reading at face value rather than paraphrasing away, because it is the concrete answer to “how does a matmul move through the chip”: weights are preloaded into the array from the top, activation data streams in from the left, and “a given 256-element multiply-accumulate operation moves through the matrix as a diagonal wavefront” — control and data pipelined so that, from software’s point of view, 256 inputs appear to be read at once and 256 accumulators appear to update instantly, even though the underlying event is a wave crossing the grid one diagonal at a time ^[2].

That wavefront is the whole mechanism. Every cell does one thing: it holds a value, multiplies it by whatever arrives from one neighbour, adds whatever arrives from another neighbour, and passes results onward next cycle. No cell needs to know what a “convolution” or a “transformer layer” is. The complexity of a neural network arithmetic graph gets flattened, at this level, into nothing more than a sequence of values entering a grid at a fixed cadence.

Four places to park the data

A weight held in place while activations stream past it is one choice, not the only one, and the choice has a name: it is a weight-stationary dataflow. Chen, Emer, and Sze’s Eyeriss paper formalised the vocabulary that the field now uses to describe the alternatives, and it is worth stating their definitions precisely rather than loosely, because the differences are the entire content of an accelerator’s scheduling strategy. Under weight stationary, “each PE holds a single weight in the [register file] at a time,” and results are passed to a neighbour or to the on-chip buffer — “the PE array operates as a systolic array with little local control,” which is exactly the TPU mechanism above ^[3]. Under output stationary, each cell instead accumulates one output value across the entire reduction, so the partial sum never leaves its cell but both operands must be re-fetched and streamed past it for every step. A third option the paper calls no local reuse dispenses with per-cell storage almost entirely, keeping cells as bare arithmetic and pushing all reuse into a larger shared buffer.

None of the three, the authors argue, is uniformly good, because each only reduces one category of data movement — weight traffic, output traffic, or neither — while leaving the others unmanaged. Their proposed fourth option, row stationary, tries to reduce all three simultaneously by mapping a one-dimensional row of a convolution’s multiply-accumulate operations onto each cell: filter weights reused horizontally across a row of cells, feature-map values reused diagonally, and partial sums accumulated vertically down a column ^[3]. Measured under matched hardware area and parallelism constraints against the other three dataflows on AlexNet’s actual layer shapes, row stationary came out 1.4 to 2.5 times more energy-efficient in the convolutional layers and at least 1.3 times more efficient in the fully-connected layers at batch sizes above 16 — a result the paper backs with a fabricated test chip, not simulation alone ^[3].

The reason those multiples are worth taking seriously rather than treating as a rounding difference is what the same paper reports about the cost of moving a byte in the first place. Normalised to a register-file access, the paper’s own energy comparison puts a fetch from the on-chip global buffer at roughly six times the cost, a hop between neighbouring cells at roughly twice the cost, and a trip out to off-chip DRAM at roughly two hundred times the cost of the same access held locally ^[3]. Once that ratio is fixed, dataflow selection is not a stylistic choice between four roughly equivalent options. It is the decision that determines whether a given operand’s hundred-times-cheaper local trip is used, or its two-hundred-times-more-expensive trip to memory is paid instead.

Formalising that trade-off well enough to search it automatically is its own research problem. Kwon and colleagues’ MAESTRO framework represents a dataflow as a small set of compiler-friendly directives rather than a hand-drawn diagram, and uses that representation to infer which forms of reuse a given mapping actually exploits. Run as a design-space search tool rather than a hand analysis, it explored 480 million candidate accelerator designs and identified 2.5 million that were valid, at an average search rate of about 170,000 designs evaluated per second ^[4]. The scale of that number is the point: with four named dataflow families and a real chip’s array dimensions, buffer sizes, and bandwidth all interacting, “which dataflow is best” stopped being answerable by intuition well before accelerator design became a mainstream discipline.

Tiling a real matmul onto a finite grid

An array is a fixed size. A real matrix multiply almost never is. Reconciling the two is called tiling, and its cost is not a rounding error — it is large enough that Google published a specific worked example against their own hardware to make the point.

The TPU’s matrix unit holds one 64 KiB tile of weights at a time, plus a second tile buffered in the background so that the 256 cycles it takes to shift a new tile into place can be hidden behind computation already underway ^[2]. When a matrix does not divide evenly into 256×256 tiles, the leftover has to be padded and processed as a partial tile anyway, and the paper reports exactly what that costs on a real workload: an LSTM layer that needs a 600×600 matrix multiply takes nine tiling steps on the 256×256 array, for 18 microseconds of total tiling time. A larger, 512×512 array would need only four steps — but each step takes four times as long, for 32 microseconds overall, because three of those four steps are mostly padding wasted on a matrix that does not fill the larger grid. The authors describe the effect as “analogous to internal fragmentation of large pages, only worse since it’s in two dimensions” ^[2]. A bigger array is not a strictly better array; it is a better array for the matrices that happen to fill it.

That single case study is why current-generation matrix units are not all built to one size. Google’s own architecture documentation describes the matrix multiply unit inside a TPU’s TensorCore as a 128×128 systolic array of multiply-accumulators on most generations through TPU v5p — 16,384 cells performing 16,384 multiply-accumulate operations per cycle — expanding to a 256×256 array, quadrupling per-cycle operations, starting with the Trillium (v6e) generation ^[5]. That is a vendor’s own account of its own roadmap and should be read as such, but the mechanism it documents is independent of any one vendor’s numbers: “the TPU loads the parameters from HBM into the Matrix Multiplication Unit… as each multiplication is executed, the result is passed to the next multiply-accumulator… No memory access is required during the matrix multiplication process” ^[5]. That last sentence is the entire promise of the systolic design stated as plainly as it can be stated: once a tile is loaded, the array runs for as long as data keeps streaming through it, without a single trip back to memory.

The memory hierarchy that makes this work has three tiers with three very different costs, and the ratios from the previous section apply directly to them: a handful of registers inside each cell holding the stationary operand and a running partial sum; an on-chip scratchpad — the TPU calls it a Unified Buffer, sized at 24 MiB on the original chip, large enough to hold both an active tile and its double-buffered successor ^[2]; and off-chip high-bandwidth memory beyond that, holding the full weight and activation tensors a tile at a time is drawn from. Every tiling decision is a decision about how much of a real matrix problem can be kept inside the cheap tiers before the expensive tier has to be touched again.

Inside a GPU tensor core: the same mechanism, one instruction at a time

A GPU does not build one enormous physical grid the way a TPU does. Instead, it builds a small, fixed-function matrix unit — a tensor core — inside each streaming multiprocessor, and invokes it repeatedly through an instruction rather than by continuously streaming data across a large array. NVIDIA’s Parallel Thread Execution instruction set documents the mma instruction family in exactly this form: instructions named by the tile shape they consume and produce, such as mma.m16n8k16, where the first two numbers describe the output tile’s rows and columns and the third describes the reduction dimension being summed over in that one instruction ^[9]. Executing that instruction is a cooperative act, not a single-thread operation: the 32 threads of a warp each hold a fragment of the input matrices in their own registers, and the instruction coordinates the whole warp to compute the tile’s output together ^[9]. It is a miniature version of the same idea Kung described in 1982 — many small compute elements combining locally held operands in a fixed pattern — implemented in a warp’s register file instead of a physical wire grid, and re-invoked by software once per tile rather than streamed continuously by hardware.

Two consequences follow from doing it this way, and both show up as measurable, independently verified behaviour rather than vendor claims taken on faith. First, the interface matters as much as the silicon underneath it: Sun and colleagues’ microbenchmarking study of real tensor cores found that the legacy, easier-to-use wmma programming interface can silently miss the fast hardware path altogether for some instruction shapes — they report that the mma.m8n8k4 PTX instruction, when issued through the pattern their benchmarks exercised, compiled down to a set of ordinary floating-point-unit instructions roughly ten times slower than the dedicated tensor-core path, rather than the fused hardware operation a developer would expect ^[11]. The array being fast is not the same statement as the instruction stream reaching it. Second, the same study documents that current tensor-core generations expose a fine-grained, hardware-enforced 50% structured sparsity mode directly in the instruction set — a claim about the instruction interface’s own hardware, verified by the paper’s own microbenchmarks rather than taken from a datasheet ^[11].

NVIDIA’s own architecture documentation for Hopper describes the further step of decoupling the instruction from the data movement that feeds it: alongside synchronous mma instructions, Hopper adds asynchronous warpgroup-level wgmma instructions that let a group of four warps issue one larger matrix operation while a separate hardware unit, the Tensor Memory Accelerator, moves the operand tiles between global and shared memory in the background — reported by the vendor as delivering twice the per-streaming-multiprocessor matrix throughput of the prior Ampere generation at equivalent numeric precision, and four times at the newly added 8-bit floating-point precision ^[8]. That is a disclosed comparison against the vendor’s own prior product and should be read as exactly that, not as a claim about any competitor. What it documents independent of the specific multiples is a design direction: moving the “wait for data to arrive” problem out of the instruction stream and into dedicated hardware, which is precisely the problem a physically streaming systolic array never has in the first place because the array itself is the data-movement schedule.

Utilization: why the datasheet number is a ceiling nobody touches

The roofline model’s ridge point, $I^{*} = P_{\max} / \beta$↗, says whether a kernel’s arithmetic intensity is high enough that memory bandwidth is not the binding constraint ^[7]. It says nothing about whether the array can actually deliver $P_{\max}$↗ once that bar is cleared, because $P_{\max}$↗ itself assumes every cell is doing useful work on every cycle — and a real systolic pipeline spends some of its cycles filling and draining rather than computing.

Consider an idealised $r \times c$↗ array holding one operand stationary, with a second operand streamed through in a sequence of length $L$↗ — for instance, $L$↗ activation vectors passed through a weight-stationary array of $r$↗ input rows and $c$↗ output columns. Before the first result emerges, the data has to propagate across the array’s diagonal, costing roughly $r + c - 2$↗ cycles; after the last input enters, the same number of cycles is needed to drain the pipeline. Total latency is therefore approximately

while the useful work performed is $r \cdot c \cdot L$↗ multiply-accumulates — one per cell, once per streamed step, in steady state. Dividing useful work by the total cell-cycles available, $r \cdot c \cdot T$↗, gives an idealised utilization

As the streamed sequence $L$↗ grows large relative to the array’s dimensions, $\eta$↗ approaches one and the fill-and-drain overhead becomes negligible. But when $L$↗ is comparable to $r$↗ and $c$↗ — a small batch, a short sequence, a matrix dimension that barely exceeds the array’s own size — the overhead is not a rounding error, it is a large fraction of every pass through the array. This is a simplified model, not a specification of any one vendor’s pipeline, but it is the same effect Google’s own engineers measured directly rather than modelled: in their reported case study of one convolutional workload, the TPU spent less than half its cycles performing matrix operations at all, and on the cycles it did spend computing, only about half of the 65,536 available multiply-accumulate cells “held useful weights because some layers… have shallow feature depths” — with roughly 35% of all cycles lost simply waiting for a new weight tile to load ^[2]. Independently, the SCALE-Sim simulator was built specifically because the research community lacked tooling to see this kind of effect at all, and its authors report using it to show, across vision, speech, text, and game-playing workloads, that memory bandwidth, dataflow choice, and an array’s aspect ratio each materially change realised runtime and energy for kernels that share the same nominal peak throughput ^[6]. A vendor’s peak-TOPS figure describes the array. It does not describe what any particular workload, at any particular batch size, will actually draw from it.

Precision formats change the geometry, not just the numerics

Reduced-precision arithmetic is usually explained as a numerics question — how much mantissa a model can tolerate losing. Inside the array, it is at least as much a geometry question: a narrower number occupies less area and less register width, so the same physical grid of cells can be built denser, or the same die area can be repurposed to raise throughput at the narrower format specifically, rather than uniformly across all formats. NVIDIA’s own comparison for Hopper’s fourth-generation tensor cores states this as a discrete jump rather than a smooth curve: the vendor reports double the matrix throughput of the prior generation at shared precisions, and four times the throughput specifically at the newly supported 8-bit floating-point format, which is a statement about hardware built to specifically exploit the smaller operand, not only about tolerable numerical error ^[8].

Structured sparsity is the same geometric argument taken further: instead of shrinking every value, it removes a fixed, hardware-known fraction of them and redesigns the array to skip the resulting zeros rather than waste cycles multiplying by them. Liu, Whatmough, and Mattina’s Systolic Tensor Array work generalises the conventional scalar processing cell into a small “tensor” cell handling several operands at once, and extends that design with a structured block-sparse format the authors call density-bound block. Against a conventional systolic array baseline at matched throughput, the plain tensor-cell redesign alone reduces circuit area by up to 2.08 times and power by up to 1.36 times; adding the structured-sparse extension on top, for models specifically trained to fit the sparse block format, reaches a 3.14-times area reduction and a 1.97-times power reduction at the same throughput, while remaining backward compatible with ordinary dense models ^[12]. Sun and colleagues’ independent microbenchmarking of shipping NVIDIA hardware confirms the same idea reached production silicon rather than staying a research proposal: current tensor-core generations expose a fine-grained 50% structured-sparsity mode directly as an instruction-level feature, measured on real chips rather than read off a specification sheet ^[11].

The mechanism-level statement worth holding onto is narrow but important: a precision or sparsity format is not merely a tolerance the model has to survive. It is a decision about how many arithmetic cells fit in a given area, or how many of a fixed grid’s cells are allowed to do nothing on a given cycle — decisions made in the array’s physical design, long before a single weight is trained.

Workload co-design: choosing a dataflow per layer, not per chip

Nothing above implies that one dataflow is correct for a whole accelerator. A convolution with many output channels and a small spatial footprint has a very different reuse pattern from a wide, shallow fully-connected layer or from the long, thin matrix multiplies inside attention — and the Eyeriss energy comparison already showed that no single one of weight-, output-, or row-stationary dominates every shape and every batch size uniformly ^[3]. Parashar and colleagues’ Timeloop tool takes this seriously as a search problem rather than a design-time constant: it represents an accelerator’s architecture and a workload’s loop structure in a common form, then evaluates candidate mappings — including which loops run spatially across the array and which run temporally over cycles — to project runtime and energy for each one, making it possible to compare architectures and mappings on equal terms rather than by hand ^[10]. MAESTRO’s 480-million-design search exists for the same reason: at real chip scale, with real buffer sizes and real bandwidth limits, no engineer is evaluating that space by intuition alone ^[4].

The practical split this produces across today’s shipping hardware is a genuine architectural fork, not a matter of one approach being more advanced. A TPU’s systolic array commits to one physical dataflow in silicon at design time and relies on the compiler to tile and schedule workloads onto that fixed structure as efficiently as it can ^[2]. A GPU’s tensor cores, invoked instruction by instruction across independently scheduled warps, keep the mapping decision open at every kernel launch, trading some of the TPU’s zero-overhead streaming efficiency for the flexibility to run a completely different tile shape and precision on the next instruction issued a moment later ^[9]. Neither design has made the other obsolete, and the field’s own tooling — Timeloop, MAESTRO, SCALE-Sim — exists precisely because which choice wins depends on the workload, the batch size, and the layer shape in front of it, not on a fixed ranking between architectures ^{[10] [6]}.

Reading a scheduling or utilization claim

A short discipline follows from all of the above, and it is the practical content of this article.

Ask what “peak” assumes. A quoted peak throughput figure assumes full array occupancy on every cycle. Fill, drain, and tile-boundary padding are not edge cases; they are the normal operating condition for any workload whose relevant dimension is not large relative to the array.

Ask which dataflow a comparison used, not just which chip. Two accelerators with identical peak arithmetic can differ by more than double in realised energy or runtime on the same workload purely from a weight-stationary versus row-stationary choice, holding hardware area fixed ^[3].

Separate the array’s capability from the instruction interface that reaches it. A tensor core’s hardware sparsity or its fastest matrix-multiply path is only used if the specific instruction shape and API a program issues actually maps onto it; measured, not assumed, behaviour is what independent microbenchmarking is for ^[11].

Treat a precision or sparsity throughput multiple as a hardware disclosure, not a universal ratio. It describes cells built for that specific format; it says nothing about a workload that has not been validated to tolerate the format in the first place.

Predictions, with the observations that would falsify them

These are forecasts, clearly separated from the sourced analysis above. Horizon: 15 August 2029.

One. Hardware-enforced structured sparsity, not just narrower numeric formats, will become a standard feature of newly announced matrix engines across vendors, because it multiplies effective throughput within a fixed cell budget rather than only within a fixed precision. Disconfirmed if a majority of major matrix-engine architectures announced in 2029 still ship dense-only arithmetic arrays with no hardware path for skipping structured zeros.

Two. Automated dataflow and tile-mapping search — the Timeloop and MAESTRO approach — will move from a research and design-time tool into runtime or compile-time software that selects a mapping per layer shape on already-shipped hardware, rather than remaining a chip-design exercise performed once before tape-out. Disconfirmed if production inference and training compiler stacks in 2029 still apply one fixed dataflow mapping to every layer of a model regardless of its shape.

Three. The two-dimensional fragmentation effect Google documented for its own 512×512 matrix-unit experiment will keep pushing new architectures toward configurable, foldable, or multiple-array-size designs rather than ever-larger single monolithic grids. Disconfirmed if the dominant new matrix-engine architectures of 2029 have grown their single systolic array’s physical dimensions substantially without adding tiling, folding, or reconfiguration flexibility to match a wider range of workload shapes.

None of these requires a discontinuity in the underlying physics. They are extrapolations of mechanisms already visible and independently documented in current hardware and its own published measurements.

What to take away

A matrix multiply, at the level where the arithmetic actually happens, is a wavefront crossing a grid of identical cells, timed to the cycle, with a decision already made about which operand stays parked in a cell’s own register and which one streams past it. That decision — weight-stationary, output-stationary, row-stationary — is not a footnote; it is worth a factor of two or more in energy for the same silicon area, because moving a byte off-chip can cost two hundred times what using it locally costs. Fitting a real matrix onto a fixed-size array is a tiling problem with a documented, measured penalty when the fit is bad, whether the mismatch shows up as a TPU’s own reported page-fragmentation analogy or as a GPU instruction that quietly misses its fast hardware path. Reduced precision and structured sparsity are not just numerical concessions; they are decisions about how many of those cells physically exist and how many of them are allowed to sit idle on a given cycle. None of it is visible in a peak-TOPS number, and all of it is knowable — because, unusually for a hardware topic this deep, the chip designers themselves have published exactly how it works and exactly where it falls short.