AI Chip Architectures

At the 2018 International Symposium on Computer Architecture, John Hennessy and David Patterson delivered their Turing Lecture 'A New Golden Age for Computer Architecture,' noting that single-threaded CPU performance growth had fallen from 52% per year in the 1980s to 3% by 2018, and advocating for domain-specific architectures. Their example, Google's TPU v1, achieved 29× the throughput of a CPU on neural-network inference with 80× better energy efficiency. The prediction of a 'Cambrian explosion' of novel architectures has materialized, with NVIDIA leading via GPUs, AMD following with 6 GW commitments from OpenAI and Meta, Google TPUs serving Anthropic with up to a million chips, Amazon Trainium running over a million chips for Anthropic, and Cerebras now serving OpenAI inference, while Groq's LPU was folded into NVIDIA via a $20B acquihire.

At the 2018 International Symposium on Computer Architecture , and John Hennessy https://en.wikipedia.org/wiki/John L. Hennessy delivered their David Patterson https://en.wikipedia.org/wiki/David Patterson computer scientist Turing Lecture: . "A New Golden Age for Computer Architecture" https://dl.acm.org/doi/10.1145/3282307 In the 1980s, when Hennessy and did their Turing Award-winning research, Patterson single-threaded CPU performance grew 52% a year. By 2018, with the end of and Moore's Law https://en.wikipedia.org/wiki/Moore%27s law , the rate was 3%. Dennard Scaling https://en.wikipedia.org/wiki/Dennard scaling There was a need for domain-specific architectures DSAs . Their worked example was Google's , already in production: 29× the throughput of a CPU on neural-network inference, at 80× better energy efficiency. The closing prediction: TPU v1 https://en.wikipedia.org/wiki/Tensor Processing Unit "the next decade will see a Cambrian explosion of novel computer architectures." This prediction came true. Today, we now have dozens of architectures in serious development. GPUs , , TPUs , LPUs , NPUs , DPUs , ASICs , wafer-scale engines , reconfigurable dataflow , neuromorphic , photonic . Particularly, these architectures focus on compute for analog . AI The architectures that have won real deployment so far: GPUs NVIDIA, AMD , TPU, Trainium , the systolic-array accelerators , and the Cerebras Wafer-Scale Engine . Groq LPU NVIDIA is the clear frontrunner; follows, with 6 GW commitments from both AMD OpenAI https://openai.com/index/openai-amd-strategic-partnership/ and Meta https://www.amd.com/en/newsroom/press-releases/2026-2-24-amd-and-meta-announce-expanded-strategic-partnersh.html . train Gemini and will TPUs serve Anthropic with up to a million chips https://www.anthropic.com/news/expanding-our-use-of-google-cloud-tpus-and-services ; Anthropic also runs Claude on over a million https://techcrunch.com/2026/03/22/an-exclusive-tour-of-amazons-trainium-lab-the-chip-thats-won-over-anthropic-openai-even-apple/ . chips Trainium Cerebras now serves OpenAI inference https://openai.com/index/cerebras-partnership/ ; the was Groq LPU folded into NVIDIA via a $20B acquihire https://www.datacenterdynamics.com/en/news/nvidia-builds-out-lpu-chip-team-following-20bn-groq-acquihire-announcement-rumored-for-gtc/ . This post aims to survey these varying approaches - their philosophy , architecture , scaling methods scale-up and scale-out , and software stack how you program the chip . The Problem AI compute is dominated by matrix multiplication . A transformer is a sequence of matmuls: , Q/K/V projection , attention , output projection - interleaved with element-wise ops: normalisation, activation, residual adds. FFN Training a frontier model performs multiply-accumulate operations matmuls are a sequence of multiply-accumulates . The shape of those matmuls depends on the workload. pushes a batch of sequences forward through every layer, backpropagates the loss, and updates the weights, with thousands of tokens flowing through the same weight matrix at once. Training is the prompt-ingestion phase of inference: the full input sequence projected through the model in a single pass, before the first output token has been produced. Both training & prefill stack many tokens against the same weight matrix, so each layer's math is a large Prefill multiply GEMM , with high matrix-matrix arithmetic intensity compute-bound . is autoregressive: the model emits one token at a time, each conditioned on every token before it, and token Decode N+1 cannot begin until token N has been produced. Only one token gets projected per step, so every matmul becomes a product GEMV . Producing one token requires a full pass over every weight in the model, plus a full read of the matrix-vector KV Cache for attention. Arithmetic intensity drops by orders of magnitude versus prefill . Inference systems recover some of that intensity by batching tokens to promote those GEMVs back to GEMMs: stacks many users' continuous batching decode steps, stacks K drafted tokens per request and verifies them in one pass, and speculative decoding folds the same trick inside the model itself. This achieves higher utilisation of the matmul units, and pushes up the Ops/B. For continuous batching, each user's request still reads its own multi-token prediction KV Cache , so long-context decode shifts from weight-bandwidth-bound to KV-bandwidth-bound. The architecture problem here is moving the numbers to where the matmuls happens fast enough. This is known as the : compute has scaled exponentially, memory bandwidth has not. memory wall Each architecture proposes a different strategy for winning the data-movement game. Understanding a chip reduces to four questions: where does data live, how does it move to the compute units, what do the compute units look like, and how do chips talk to each other at scale . NVIDIA GPU The NVIDIA GPU is a massively parallel processor . The philosophy is that a programmable chip with thousands of threads, orchestrated by a host CPU and exposed through , is the right machine to run CUDA https://docs.nvidia.com/cuda/cuda-c-programming-guide/ workloads. Each generation adds acceleration primitives onto programmable parallelisable without changing the programming model. The same chip trains transformers, serves inference, renders graphics, and runs scientific simulation Streaming Multiprocessors . accelerated computing Genealogy unified shaders and the SIMT execution model . warp schedulers , IEEE-754 FP64 . Architecture An NVIDIA GPU is a group of throughput-oriented cores, a deep memory hierarchy to keep them fed, + just enough scheduling silicon to keep thousands of threads in flight . The cores are , replicated 100+ times per package: 80 on Streaming Multiprocessors V100 , 108 on A100 , 132 on H100 , 148 on B200 , 160 on B300 , 224 on Rubin . Inside every SM sits the same recipe: four , each with its own SM Sub-Partitions , warp scheduler , 16k×32-bit dispatch unit , scalar register file lanes, a CUDA Core for transcendentals, and a private port into the SM's Special Function Unit . The four partitions share an Tensor Cores L1/shared-memory block, and the TMA . Threads are grouped into of 32 that execute in warps SIMT lock-step; dozens of resident warps per partition let the scheduler hide memory/arithmetic stalls by switching between them. Compute are the original compute throughput, and for AI they still own everything that isn't a matmul: activations, residual adds, normalization, address arithmetic. But, a transformer block is ~99% matmul FLOPs, so the overwhelming compute throughput comes from the CUDA Cores . Tensor Cores These cores execute fused matrix multiply-accumulate on small matrix tiles, The full matmul is broken into output tiles: to produce one output tile, a kernel walks the shared inner dimension , drawing from a row-strip of the left input matrix and from a column-strip of the right, and folds each partial product into a running accumulator. holds the partial sum so far, is the updated value carried into the next step. After the inner loop completes, is one finished tile of the full output matrix; the whole matmul is built from many of these tile MMAs . Tile shapes are written M × N × K , is the output tile size, and is how much of the inner dimension the instruction contracts over in one fire; the rest of the matmul's axis is walked by the kernel's inner loop. The accumulator is sticky across that loop: each MMA's output becomes the next MMA's input , so the equation is really in place: successive instructions fold their partial products into the same storage until the K-axis is fully walked. V100 's first-gen unit 8 per SM ran a warp-level 16×16×16 FP16 MMA . A100 's 3rd-gen unit added TF32 , BF16 , FP64 matmul, and 2:4 structured sparsity . H100 's 4th-gen unit added native FP8 and pulled the abstraction up from a warp to a : 128 cooperating threads firing an asynchronous warp group at 64×256×16 shape that runs in the background while the issuing warps load the next tile. wgmma B200 's 5th-gen unit went further still: a of 256×256×16 with operands split across a pair of SMs, native two-SM MMA FP4 , and a dedicated 256 KB scratchpad per SM that holds accumulator tiles instead of bleeding into the register file. Tensor Memory TMEM Rubin 's 6th-gen unit extends FP4 throughput, adds native FP6, and pairs with a 3rd-gen Transformer Engine that does adaptive NVFP4 micro-block scaling in hardware, keeping the per-tile quantization metadata on the Tensor Core path, rather than through the CUDA Cores . What stays constant across all six generations is that the matmul lives inside the thread/warp hierarchy , but the number of threads it takes to issue one has shrunk, and the issue itself has decoupled from execution. Volta 's mma.sync is warp- collective and synchronous: all 32 threads in a warp execute it together, each lane holding register fragments of A, B, and the accumulator D , and the warp blocks until it completes. Hopper 's widens the issuer to a wgmma.mma async warp-group of 128 threads, moves B into a shared-memory descriptor A becomes optional: either registers or a descriptor, kernel's choice , and returns immediately : the matmul runs in the background while the warp-group queues the next tile, with completion tracked via wgmma.commit group / wgmma.wait group . Blackwell 's tcgen05.mma completes the migration: A joins B in shared-memory descriptors or A comes from TMEM directly , and the accumulator D lands in TMEM rather than the register file. With every operand off the lanes, there is no per-thread state for an issue to coordinate, so a single thread fires the instruction and returns immediately , with completion signalled by an mbarrier the consumer warp waits on. The rest of the warp, and the issuing thread itself, is free for other work in the meantime. A scales the same model across two SMs: one thread on each SM in a paired cluster issues coordinated CTA -pair variant MMAs that share operands across the pair, composing the 256×256×16 two-SM tile under the same async/ mbarrier completion, just promoted to a cluster-level barrier so the pair stays in step.The matmul has grown bigger and lighter on the issuing threads at the same time: an instruction that started as 32 lanes acting in lockstep is now closer to a single descriptor-driven command , dispatched from inside the warp model but no longer executed by it. That decoupling is what makes transformer attention kernels efficient on a GPU. The warp can run softmax, apply a mask, or pre-load the next tile while the matmul is in flight; the overlap of matmul and the surrounding element-wise work is the structure of every modern attention kernel FlashAttention-3 , FA4 , and it depends on the matrix instruction not blocking the warp. Memory The on-chip hierarchy is hardware-managed caches at every level, with software hints layered on top . Off-chip is : 32 GB HBM HBM2 on V100, 80 GB HBM3 on H100, 192 GB HBM3e on B200, 288 GB on B300, 288 GB HBM4 on Rubin. A chip-level sits between HBM and the SMs: 6 MB on V100, 40 MB on A100, 50 MB on H100, 60 MB on B200 split into two 30 MB banks across the L2 Cache two-die package , with locality-aware residency controls so that hot tiles can be pinned to the near die . Inside each SM, 256 KB of unified is partitioned at kernel launch between hardware-managed L1 and a programmer-controlled scratchpad. The register file is another ~256 KB per SM, sliced four ways across the partitions. L1/SMEM Blackwell adds a fifth tier: , 256 KB per SM dedicated to TMEM MMA accumulators and addressed only by the Tensor Core, pulling the operand-residency pressure out of the general register file. Movement between tiers has been progressively decoupled from the warp. Pre-Ampere, loading a tile was synchronous: each thread issued its own global load, the warp blocked until every fragment landed in registers, and a second pass copied them to shared memory; every tile burned warp lanes on address arithmetic and on the wait. Ampere introduced : per-thread async copies HBM → SMEM that bypass registers entirely, with the warp committing groups of in-flight copies and waiting only when the consumer needs the data. cp.async Hopper replaced that with the , a dedicated DMA engine: one thread submits a multi-dimensional tile descriptor base address, leading dimension, swizzle , the engine handles all the address arithmetic and writes into shared memory, and completion is signalled by an TMA . The whole warp is freed from load issue and address math; the kernel just queues descriptors. TMA also supports mbarrier cluster-level multicast : one HBM read fans out to every SM in a thread-block cluster , turning what used to be N separate loads into one. Blackwell extends TMA again: direct loads into TMEM , so accumulator tiles stream in without staging through SMEM. The trajectory is one less thing the warp has to do per tile, generation after generation. Warp Specialisation The Hopper-era programming idiom is : inside one block, some warps act as warp specialisation that issue back-to-back producers TMA loads; others act as that fire consumers on freshly-arrived tiles. Synchronisation between them is no longer the old SM-wide wgmma syncthreads barrier; it is memory barriers in shared memory and asynchronous transaction barriers attached to TMA completions, allowing fine-grained producer/consumer handshakes at warp granularity rather than block granularity. The pattern that has become the reference for every modern attention kernel mbarrier FlashAttention-3 , ping-pong GEMMs, the Blackwell CUTLASS https://github.com/NVIDIA/cutlass FA4 kernel is the same recipe: a TMA-driven producer pipeline feeds a wgmma consumer pipeline through shared memory and TMEM, with mbarrier handshakes and Hopper+ tying multiple SMs into one cooperative compute unit so that the two-SM MMA of Blackwell composes naturally on top. thread-block clusters Numerics FP32 was the historical default; Volta brought with FP32 accumulate and the FP16 loss-scaling tricks that made it trainable; Ampere added FP32 range, FP16 mantissa, drop-in for FP32 matmul , TF32 , and 2:4 BF16 that doubles effective throughput on pruned weights. Hopper introduced native structured sparsity in both FP8 E4M3 and E5M2 , paired with the which auto-scales activations layer-by-layer to keep them inside FP8 dynamic range. Blackwell halved precision again with Transformer Engine and shipped FP4 block-level shared exponents that recover most of the accuracy lost at FP4 , together with a 2nd-gen Transformer Engine that retargets the auto-scaling pipeline to FP4. Rubin's 3rd-gen Transformer Engine adds microscaling MX formats NVIDIA's tightened FP4 variant and native NVFP4 with more aggressive sparsity. The chip layout itself is now part of the numerics story: B100/B200/B300 are FP6 stitched by a ~10 TB/s two reticle-limit dies link and presented to software as one logical GPU, with 8 HBM stacks on the package; Rubin extends the chiplet recipe to dual-die at ~336 B transistors with 8 HBM4 stacks. Every generation buys roughly 2× per-watt throughput by cutting bits in half and restoring accuracy with a finer-grained scaling scheme, and increasingly, by bonding more silicon into the package. NV-HBI Bets The workload is a moving target attention variants, novel model architectures , so keep every block programmable and let the developer write Bet 1: Programmability. CUDA https://docs.nvidia.com/cuda/cuda-c-programming-guide/ . Even the specialised units are exposed through that model rather than as fixed-function blocks.Latency is unpredictable and data-dependent, so hide it not with a static schedule but with massive thread overcommit, up to 64 resident warps per SM, with the hardware Bet 2: Hide Latency with Massive Multithreading. warp scheduler picking a ready warp every cycle.The matrix unit is the overwhelming compute throughput, but it must live behind the same warp/thread abstraction that everything else uses, so wrap it in Bet 3: Warp-wrapped Matmul. mma.sync → wgmma → tcgen05.mma - rather than expose it as a fixed-function pipe. This enables a single kernel to fuse matmul, softmax, and element-wise ops in one pass.Make the memory hierarchy Bet 4: Async Memory Hierarchy. explicit and programmer-managed rather than implicit and compiler-scheduled . Keep the L2 cache , but expose SMEM and TMEM as named scratchpads, and layer async machinery on top: TMA for bulk copies, TMEM for the matmul accumulator,for the producer/consumer handshake. The hierarchy is mbarrier software-pipelined inside a programmable kernel, not statically scheduled by a compiler against a known-latency scratchpad.Every transistor spent on a warp scheduler, register-file, or coherent cache is a transistor not spent on a Bet 5: Amortised SIMT Tax. MAC ; accept the tax, and pay it down two ways: a Tensor Core now big enough that the SIMT machinery is amortised across a much larger MAC count, and units like TMEM trading away some general-purpose flexibility for MAC density. Scaling There are two regimes for scaling: and scale-up . scale-out AI infrastructure uses both: bandwidth-hungry collectives tensor parallelism , MoE expert routing stay inside the scale-up domain; data parallelism and pipeline parallelism cross the scale-out fabric. Scale-up The scale-up stack is plus NVLink . NVSwitch NVLink implements a between GPUs, so a load or store on one GPU can target another GPU's cache-coherent fabric HBM with the hardware handling address translation and coherence. But NVLink by itself is point-to-point: one link connects exactly two chips. NVSwitch is a dedicated chip that every GPU connects to, routing traffic so every GPU can simultaneously communicate with every other at full crossbar NVLink bandwidth, non-blocking and all-to-all . Together they defined the 8-GPU baseboard, pairing eight HGX H100 SXM modules with x86 hosts AMD EPYC or Intel Xeon over PCIe Gen5 . Hopper also shipped a Grace -paired form: the bonded one GH200 Grace Hopper Superchip Grace ARM CPU to one H100 over at 900 GB/s, eliminating the NVLink-C2C PCIe host-device hop. Modules scaled up into pairs and rack-level GH200 NVL2 . Blackwell makes the pairing the default. The GH200 NVL32 module fuses one GB200 Grace with two B200s over NVLink-C2C , and stitches 36 of them into a single liquid-cooled scale-up domain: 72 GPUs, 36 NVL72 Grace CPUs, 13.5 TB of HBM and 17 TB of LPDDR5X as one flat, coherent address space. Rubin steps this in two. ships in 2026 as a Rubin-generation refresh inside the same NVL144 Oberon -class rack: 72 Rubin packages, badged as 144 GPUs under NVIDIA's new die-counting convention, with HBM4 and NVLink 6 doubling per-package bandwidth. The actual rack-scale jump is Rubin Ultra in 2027: packs 144 four-die Rubin Ultra packages into the new NVL576 chassis for 576 GPU dies in one coherent domain. Kyber That density is held together by . NVL72's NVLink fabric runs over 5,184 cables blind-mated through a backplane ~2 miles of cabling per rack, no passive copper in-cable retimers , the SerDes living on the GPU and switch ASICs themselves , carrying ~130 TB/s of all-to-all bandwidth across the 72 GPUs. NVIDIA estimates the copper choice saves roughly 20 kW per rack against an optical equivalent that would have needed pluggable transceivers on every link. Copper is what makes rack as one GPU economically practical: at sub-2-metre runs it still wins on power, cost, and signal integrity per dollar; beyond that, the bits have to go on glass. NVL144 stays inside Oberon and copper continues to work because the package count 72 is unchanged from NVL72; the cabling doesn't have to lengthen, just transmit faster on Gen 6 SerDes. Rubin Ultra 's NVL576 holds the same copper line by reshaping the rack: the new Kyber form factor is roughly twice the height of Oberon and packs all 576 GPU dies into one enclosure, sized specifically so every NVLink path stays within passive-copper reach even at 144 four-die packages and tens of thousands of cables. Scale-out The scale-out stack comes from their acquisition of Mellanox . Unlike NVLink , scale-out fabrics are not coherent : nodes keep separate address spaces, and data crosses only via explicit initiated by software, typically wrapped in RDMA collectives like NCCL all-reduce or all-to-all . The reference cluster is the : eight NVL72 racks stitched together over DGX SuperPOD Quantum-X800 InfiniBand yield 576 Blackwell GPUs under a single scheduler, and training clusters scale further by tiling SuperPODs. Rubin SuperPODs in 2026 keep the same 8-rack pattern with NVL144 yielding 1,152 GPUs per SuperPOD instead of 576 . Rubin Ultra in 2027 scales the recipe up an order of magnitude: Kyber racks of 576 GPU dies each, stitched together over Quantum-X Photonics CPO , putting thousands of GPUs under one scheduler. Every GPU has its own ConnectX NIC into that fabric. Blackwell nodes run ConnectX-8 at 800 Gbps per GPU, an order of magnitude less bandwidth than per-GPU NVLink , and latencies climb from nanoseconds to microseconds. Rubin moves to ConnectX-9 at 1.6 Tbps per GPU, doubling the per-GPU scale-out bandwidth as the per-rack scale-up domain grows from 72 to 576 GPUs. Alongside each NIC sits a BlueField DPU , adding ARM cores and accelerators to offload storage, networking, and security from the host CPU. For customers who prefer Ethernet to InfiniBand , is a lossless-Ethernet alternative tuned for AI traffic. Spectrum-X The crossover from copper to glass happens at the rack boundary. Inside the NVL72 the spine is copper; once a link has to cross racks at 800 Gbps it is optical . Passive copper DAC tops out at roughly 1.5–2 metres at 200 G/lane, well short of cross-rack reach, so today's SuperPOD spine rides over OSFP-RHS pluggable transceivers , each module carrying its own laser, modulator, photodetector, and DSP. A SuperPOD spine fanning out to thousands of GPUs is, in optical terms, tens of thousands of pluggables drawing tens of kilowatts on transceiver lasers alone. With Rubin, that optical layer collapses into the switch ASIC . Quantum-X Photonics InfiniBand and Spectrum-X Photonics Ethernet replace the pluggables with : lasers, modulators, and photodetectors bonded onto the switch package via TSMC COUPE. NVIDIA claims ~4× fewer lasers and ~3.5× lower link power than the OSFP-pluggable equivalent. The chiplet logic that turned the GPU into a two-die package and stacked HBM next to it is now showing up at the network layer: vertical integration of compute, memory, co-packaged optics and photonics on one substrate. recently opened the scale-up fabric itself: third-party CPUs and NVLink Fusion XPUs can now join NVLink domains, letting hyperscalers build semi-custom racks around NVIDIA's interconnect without designing their own coherent fabric from scratch. Software is the natural programming model for a CUDA https://docs.nvidia.com/cuda/cuda-c-programming-guide/ processor. You write a kernel one piece of code executed once per thread and launch it across thousands of threads organised into blocks and warps; the programmer decides what they share, when they synchronise, and which piece of the problem each one handles. That is why the abstraction has barely changed in eighteen years, and why every CUDA kernel written since 2007 would still compile and run on Blackwell. massively parallel That continuity is both the moat and the constraint. Each new generation introduces new hardware Tensor Cores , TMA , TMEM onto the same kernel-and-warps model, exposed as intrinsics in and PTX https://docs.nvidia.com/cuda/parallel-thread-execution/ : SASS https://docs.nvidia.com/cuda/cuda-binary-utilities/ mma.sync , wgmma.mma async , and so on. NVIDIA cannot radically rethink the SM because too much code depends on it; in return, every investment in CUDA software compounds across generations.On top of PTX sits a stack constructed over two decades. and cuBLAS https://docs.nvidia.com/cuda/cublas/ for math and DNN primitives; cuDNN https://developer.nvidia.com/cudnn , encoding decades of GEMM expertise in templated C++; CUTLASS https://github.com/NVIDIA/cutlass for paged attention, in-flight batching, and speculative decoding; framework bindings through TensorRT-LLM https://github.com/NVIDIA/TensorRT-LLM , PyTorch https://pytorch.org/ , and Triton https://triton-lang.org/ . JAX https://github.com/jax-ml/jax , one of the most important algorithmic rewrites in modern AI, tiles attention to avoid materialising the matrix. Its four generations FA1 through FA4 have each been hand-optimised for the latest NVIDIA silicon FA3 for Hopper's async pipelines, FA4 for Blackwell , with ports to other hardware trailing by months or years. FlashAttention https://arxiv.org/abs/2205.14135 Most of this stack is written by people NVIDIA does not pay. The moat is not CUDA itself; it is two decades of third-party kernels, libraries, and tooling, and the millions of developers who have learned the API along the way. NVIDIA also ships human expertise alongside the silicon. They embed dozens of their own engineers inside frontier labs and hyperscaler teams, writing kernels for each new model architecture and tuning them to each new silicon generation. Whatever a lab wants to train next month tends to run well on NVIDIA much faster than other platforms. Switching off NVIDIA is therefore not just rewriting the kernels and libraries. It is re-training the mental models of an entire engineering workforce, and losing the NVIDIA engineers who today sit inside the building. Google TPU The is a TPU https://en.wikipedia.org/wiki/Tensor Processing Unit . The philosophy is, rather than a programmable chip that can run any massively-parallel workload, focus on a single primitive dense matrix-multiplication on a large matrix multiplication machine systolic array https://en.wikipedia.org/wiki/Systolic array and let the compiler plan every cycle and every byte of memory ahead of time. No hardware scheduler, no cache, no threads/warps. Each generation grows the XLA https://openxla.org/xla pod , with thousands of chips wired through the interconnect into one coherent machine. A TPU has no ambition to render graphics or run scientific simulation; it exists to train and serve Google's workloads search, translation, recommendation, Gemini more efficiently per watt than any general-purpose alternative. ICI Genealogy optical circuit switches Palomar ; SparseCores ; both BF16 & ; 4,096-chip pods. INT8 Architecture A TPU chip is a matmul engine wrapped in just enough silicon to keep it fed . The unit of compute is the : flagship chips from TensorCore v2 onward carry two per package; efficiency-tuned chips v4i , v5e , v6e carry one. Inside every TensorCore sits the same five-component recipe: one or more for matrix math, a MXUs for VPU element-wise math , a that runs the show, an Scalar Unit for cross-lane reductions, and an attached XLU , plus accumulator queues feeding and draining the MXU. From Transpose/Permute Unit v4 onward each chip also carries dedicated dataflow engines outside the TensorCore 4 per chip on SparseCore v4 , v5p , and Ironwood ; 2 per chip on Trillium , explicitly carved out to absorb the embedding-lookup workload the systolic array was the wrong shape for. Every block sits on a single VLIW issue plane driven by a that fills all eight functional slots of a 322-bit bundle every cycle. There is no instruction cache miss, no Core Sequencer warp scheduler , no out-of-order engine, no branch predictor : the compiler is the scheduler, and the silicon area saved is spent on more MACs . TensorCore The is the systolic array. MXU v1 shipped one 256×256 INT8 inference array; v2 was the first training-capable TPU and introduced 128×128 cells doing BF16 multiply with FP32 accumulate INT8 came back to the MXU at v4 onwards at equivalent throughput . Cell counts per TensorCore grew from there: 1 MXU on v2 → 2 on v3 → 4 on v4 / v5e / v5p . Trillium went back to 256×256 65,536 multiply-accumulate cells per array per cycle , and Ironwood , 8t , and 8i all kept the 256×256 shape. To compute , matrix B's values are pre-loaded one weight per cell: dataflow, the choice that distinguishes TPUs from weight-stationary output-stationary arrays elsewhere. Activations enter from the left edge, propagate one column per cycle, multiply against the resident weight at every cell, and partial sums flow downward into accumulator queues at the bottom. Once data enters the array no memory access occurs: each weight is reused for every activation that passes through, each activation is reused 128 or 256 times across the row. Data reuse is wired into the silicon, not arbitrated by a cache. The dominant cost in computing is not the multiplication itself a few picojoules but reading and writing memory at 100–1000× more energy per access; the systolic array deletes that cost by construction. The trade-off is : a 128×128 matmul on a 256×256 array wastes 75% of the silicon, so underfill XLA https://openxla.org/xla tiles , pads , and schedules dimensions to multiples of 128 or 256 on v6e+ and the model code is written with those quanta in mind. The is the second-fiddle compute engine but is in many ways the more interesting microarchitectural object: every TPU is a 2D vector machine, not a 1D SIMD machine. The VPU's register file holds 2D VPU . On VREGs v4 / v5p the shape is 8, 128 : 128 wide, 8 lanes deep, 32 v4 or 64 v5p registers per core, with 4 independent floating-point ALUs per lane, sublane . The lane axis matches the systolic array's input width, so the lane count presumably widened to 256 alongside the MXU on sublanes Trillium and Ironwood ; Google has not published post-v5p VPU dimensions. The sublane axis lets the VPU stream tiles through the MXU at one matmul per X clocks where X is the sublane dimension . Most of the speedup in modern TPU programs comes from : quantisation, layernorm, softmax, activation, and bias-add all run on the VPU in the same cycles the MXU is running the matmul behind them. Cross-lane reductions the awkward case for any 2D vector ISA are handled by the VPU/MXU overlap : slow, expensive, and a known compiler hot spot. Layout transforms that misalign with the 2D shape are absorbed by the dedicated XLU , sparing a round-trip through memory. Transpose/Permute Unit The is the smallest block and arguably the most consequential: a single-threaded, dual-issue integer ALU with 32 32-bit registers and 4 KiB of Scalar Unit for control state, paired with an Imem holding the program. It is the only block that does instruction fetch; every cycle it pulls a 322-bit VLIW bundle, executes its own two scalar slots locally address arithmetic, loop counters, branches, sync-register checks , and dispatches the remaining six slots to the rest of the chip: 2 vector ALU VPU , 2 vector load/store HBM↔VMEM SMEM DMA , 2 matrix push/pop the MXU queue . Synchronisation between blocks is explicit: track when MXU and VPU pipelines are busy, and the compiler inserts barrier checks rather than the hardware tracking dependencies. The Scalar Unit is what makes the rest of the TensorCore look like fixed-function dataflow: every cycle, one place decides what eight things happen, and there is no dynamic sync flags reorder buffer to undo a bad decision. Memory The on-chip memory hierarchy is the same idea as the compute side: there are no caches, every level is software-managed . Off-chip is 16 GB on v2/v5e, 32 GB on v3/v4/v6e, 95 GB on v5p, 192 GB on Ironwood, 216–288 GB on the v8 generation , and on-chip is a hand-stacked tier of explicitly-addressable scratchpads. Closest to compute is HBM , the vector scratchpad feeding both the VPU and the MXU input queues, sized 32 MiB on v4, 128 MiB on v5e, and stretched to 384 MiB on the inference-tuned VMEM v8i precisely to hold an entire KV cache on chip. Above it sits , introduced with CMEM v4 at 128 MiB: a slower, larger SRAM staging area between HBM and VMEM that absorbs fused-op intermediates. The Scalar Unit has its own ~10 MiB for control state on v4 and a tiny scalar register file. Every tensor in the program is pinned to one tier at compile time; XLA's SMEM buffer-assignment pass schedules DMAs across tiers so that data arrives just before the cycle that consumes it. The hardware does no prefetching, no eviction, no coherence ; when the compiler gets it right, the array never stalls; when it gets it wrong, there is no fallback path. SparseCore The block outside the TensorCore that breaks the systolic mould is , introduced with SparseCore v4 . Recommender and ranking models live on embedding lookups billions of indices into vast tables , and the access pattern is the inverse of dense matmul: irregular , indirect , all-to-all . A 256×256 systolic array is exactly the wrong shape. SparseCore is a with 16 compute tiles and dedicated dataflow processor scratchpads, sitting alongside the TensorCore and absorbing SPMEM scatter , gather , and segmented-reduce primitives plus the data-dependent all-to-all traffic that sharded embedding tables generate. This achieves 5–7× speedup on embedding -heavy models for ~5% of die area and power. v4 shipped 4 SparseCores per chip, v5p kept that count, Trillium dropped to 2, and Ironwood went back to 4 2 per chiplet on its dual-die layout . The v8i Zebrafish inference chip removes SparseCore entirely and replaces it with a on the I/O chiplet: different problem collective CAE Collectives Acceleration Engine reductions during autoregressive decode , same idea carve a small accelerator off the main core to absorb a workload the systolic array is the wrong shape for . Numerics TPU v1 was INT8 -only inference; v2 switched this for as the canonical training format: same dynamic range as BF16 FP32 , half the memory, no loss-scaling tricks. v4 reintroduced native INT8 support. Ironwood then added native support both E4M3 and E5M2 for ~2× the throughput of BF16 in the same area. v8 adds native FP8 plus FP4 inside the MXU itself, which deletes the VPU dequant overhead that Ironwood still paid. block-scale multiplication is hardware-supported on every modern TensorCore: rounding decisions made by the lower mantissa bits acting as a probability, which preserves the expected value of low-precision accumulations across long training runs and is one of the small details that lets BF16/FP8 close the accuracy gap to FP32. Stochastic rounding At the chip boundary sit the ports themselves 4 ports on the ICI 2D-torus chips v2/v3/v5e/v6e, 6 on the 3D-torus flagships v4/v5p/v7/8t , and the DCN NIC for scale-out. From a chip-level perspective the ICI ports look like just another set of DMA engines the Core Sequencer can target inside a VLIW bundle: a remote-tensor send is the same instruction class as a VMEM-to-HBM transfer, and the compiler treats collectives as part of the same overall schedule it builds for compute and local memory. Bets Matmul dominates the workload, so spend the silicon on a systolic array. Bet 1: Systolic array. Compute is cheap and memory is expensive, so reuse data in the wires of the array and replace caches with software-managed scratchpads. Bet 2: Software scratchpads. The workload is statically predictable, so move scheduling into the compiler: VLIW issue, no Bet 3: Compiler scheduling. speculation , no out-of-order, no dynamic scheduler .Power matters more than peak, so delete every transistor that does not Bet 4: MAC-only silicon. multiply-add : every cache tag, every branch predictor, every reorder buffer.The dense matmul array is the wrong shape for some real workloads Bet 5: Dedicated off-array engines. embeddings , collectives , so carve out small dedicated engines SparseCore, CAE rather than warp the main core to fit them. Scaling The TPU's scale-up story is the inverse of NVIDIA's. Where NVLink + NVSwitch make every other GPU's HBM look like local memory a hardware-managed coherent address space , Google's ICI is message-passing . There is no remote-load semantics , no cache coherence , no crossbar . Every multi-chip operation is an explicit collective compiled by . The scale-up domain is tied together not by a switch fabric but by a XLA https://openxla.org/xla chips wired directly to their neighbours with torus edge wrap and stitched at the rack boundary by . optical circuit switches 2D or 3D torus over ICI . XLA emits SPMD collectives that tightly choreograph thousands of TPUs as one program. No coherence, but huge bisection bandwidth at low latency. Scale-up ICI links come straight out of the TPU die: high-speed serial lanes , direct-attach copper inside a 64-chip cube a 4×4×4 arrangement that lives in one liquid-cooled rack , optical between cubes. Per-chip aggregate ICI bandwidth has scaled from ~250 GB/s on v2 to 1.2 TB/s bidirectional on , and Ironwood that on 2× . Topology alternates by generation: v8t 2D torus on the efficiency-tuned chips v2 , v3 , v5e , v6e , 3D torus on the flagships v4 , v5p , v7 , v8t . The piece with no NVIDIA analogue is the : a Palomar OCS that sits between cubes. Tiny mirrors physically rotate to map any input fibre to any output. A 3D-MEMS optical circuit switch v4 superpod uses 48 Palomar switches to wire 64 cubes 4,096 chips into one 3D torus; v5p and Ironwood scale the same scheme up. Reconfiguration is millisecond-class, not nanosecond, but that's fine, because OCS is : pick a topology at job start, run it for a week, then reconfigure for the next workload. Three problems collapse into one component: topology reconfiguration per workload circuit-switched twisted tori give up to 70% better bisection , sub-pod slicing on demand, and when a chip dies, the fault tolerance OCS optically swaps in a spare cube and the run continues without losing the ICI domain . This makes the the unit of scale-up: equivalent in role to NVIDIA's NVL72, two orders of magnitude bigger. superpod v4 was 4,096 chips; v5p , 8,960; is 9,216 chips arranged as 144 cubes of 64, presenting Ironwood TPU v7 as one coherent 1.77 PB of HBM ~68 PB/s and 42.5 ExaFLOPS FP8 ICI domain. stretches this to TPU 8t Sunfish . 9,600 chips, 2 PB of HBM ~62 PB/s , and 121 ExaFLOPS FP4 has TPU 8i Zebrafish . 8i replaces torus with a new hierarchical 1,024 chips, ~295 TB of HBM 8.8 PB/s , and ~10 ExaFLOPS FP4 high-radix topology called 4-chip ring → 8-board group → up to 36 groups linked by Boardfly OCS , cutting all-to-all latency in half. This is designed for MoE inference. A 3D torus excels when collectives are nearest-neighbour ring all-reduce uses every link every cycle , but MoE expert routing is the opposite pattern, all-to-all : every chip ships unique fragments to every other, and round-trip latency is bounded by the longest-hop pair. A 1,024-chip 3D torus has a 16-hop diameter; Boardfly 's ring → group → OCS hierarchy compresses that to 7. Scale-out Through TPU v7 , scale-out ran over a single fabric: , Jupiter via all-optical at the spine since 2022 , the same 3D-MEMS family as Apollo OCS Palomar , scaled across the building. Google uses the same primitive optical circuit switching at every layer from rack to datacenter spine; that is the architectural signature nobody else has. Jupiter today carries 13 Pb/s of bisection per building. With , scale-out split into two fabrics. East-west TPU-to-TPU traffic moved to TPU 8t , a dedicated accelerator fabric; Virgo Jupiter retained the role: storage access, general compute, and inter-site scaling. north-south Virgo is a topology built on flat , two-layer , non-blocking switches: every TPU is at most two switches from any other. One high-radix Virgo cluster links 134,000+ TPU 8t s at 47 Pb/s of bisection 4× the per-chip bandwidth and 40% lower unloaded latency than the prior DCN generation , with multi-planar fault isolation and sub-millisecond telemetry that lets the scheduler kill stragglers before they wreck a step. The architectural payoff is that each layer can now evolve independently: scale-up, east-west scale-out, and front-end can iterate on different cadences without rewiring the others. Per-chip scale-out bandwidth is on the order of 100 Gbps on , and Ironwood that on 4× , but still two orders of magnitude less than per-chip v8t ICI . This bandwidth gap dictates partitioning: tensor parallelism and MoE expert routing stay inside ICI ; data parallelism and pipeline parallelism cross the scale-out fabric. Google's framework, plumbed into Multislice , lets a single XLA https://openxla.org/xla SPMD program span multiple slices in different pods ; the compiler emits hierarchical collectives ring all-reduce inside each slice, higher-level reduce across . The structure is exactly the trick for hiding the ICI / DCN bandwidth gap: as much work as possible stays inside the slice over fast ICI, leaving only the cross-slice residual to pay the slow-fabric cost. Above this sits . Where Pathways NCCL + Slurm + Megatron -style schedulers drive SPMD from many controllers, Pathways drives the entire job from client and one virtualises multiple "islands" pods with their own ICI domains connected over DCN . It does gang scheduling , elastic training when a slice fails, OCS reshapes the topology and Pathways resumes from the last checkpoint on the new shape , and cross-region orchestration . was the first frontier model trained across multiple datacenters; Gemini Ultra Pathways stitches them into one synchronous SPMD job. The philosophy: the compiler is the scheduler, the torus is the topology, and the optical switch is the universal reconfigurable substrate , at every layer from rack to datacenter. Software The TPU stack is compiler-driven where CUDA is . On a GPU, the developer writes the kernel and the kernel-driven framework strings kernels together; the compiler 's job is mostly local . On a TPU, the developer writes a numerical program in and JAX https://github.com/jax-ml/jax is responsible for everything below it: which operations XLA https://openxla.org/xla fuse , where each tensor lives , how it is laid out across the 2D vector registers, when DMAs from HBM to VMEM issue, how the 322-bit VLIW bundles are scheduled , how the program shards across thousands of chips. There is no hardware fallback: no warp scheduler , no cache, no out-of-order engine to paper over a bad schedule . The compiler is the system. The trade-off is the central one of the architecture: . XLA gets closer to the theoretical ceiling without hand-tuning, but closing the remaining gap is harderThe compilation path is . JAX → JAXpr → StableHLO → HLO → LLO → VLIW bundles traces a Python function into a JAX typed functional IR JAXpr under , lowers it to jit the StableHLO OpenXLA -standardised, versioned op-set of ~100 statically-shaped primitives that all front-ends now emit , which XLA ingests as and runs through its pass pipeline: HLO collapse operation fusion pointwise + reduction + matmul into one kernel so intermediates never hit HBM , decide the layout assignment 2D tiling of every tensor so it streams into the MXU without a transpose : substantially harder than on 1D SIMD machines because both the registers and the systolic inputs are 2D , every tensor pinned to either VMEM, CMEM, or HBM with buffer assignment overlap windows pre-computed , , and finally a VLIW SPMD partitioning scheduler that fills all eight slots of every bundle. HLO lowers to Low-Level Optimizer , the TPU-specific LLO IR , and LLO emits the final VLIW stream. A well-compiled program overlaps MXU systolic execution, VPU element-wise math , and HBM↔VMEM DMA in the same bundle every cycle. Multi-chip execution is : one program, sharded data, hierarchical SPMD collectives , emitted by now being replaced by GSPMD , an Shardy MLIR -native successor that lands as the default in early 2026 . The user expresses sharding declaratively with + Mesh annotations on a few key tensors; the PartitionSpec compiler propagates shardings through the rest of the graph and inserts all-reduces , all-gathers , and reduce-scatters where the layout changes. When the compiler picks a wrong collective, drops the user into shard map per-device code with explicit local shapes and explicit manual SPMD collectives , composable inside jit so a single kernel can be hand-partitioned without giving up auto-partitioning everywhere else. This is the inverse of the PyTorch idiom: and FSDP wrap the model in a runtime that issues DeepSpeed collectives at module boundaries; GSPMD/Shardy partitions the whole graph as a compiler problem. is the escape hatch: JAX's kernel-writing language, broadly the TPU equivalent of Pallas on GPUs. Pallas kernels are written in JAX-flavoured Python, lowered through Triton the Mosaic MLIR -based TPU backend to LLO, and embedded back into HLO as a custom op. It exists because XLA cannot always synthesise the optimum for novel attention variants, fused MoE dispatch, or anything that demands manual VMEM tiling and DMA scheduling: a optimisation, where the win is in the FlashAttention-class schedule and not the algebra. targets H100/Blackwell with the same front-end, so a kernel author can write once and lower to either substrate. The library tier above this is uniformly JAX-native: Pallas:Mosaic-GPU for modules, Flax NNX for optimisers, Optax for asynchronous distributed checkpointing, Orbax for input pipelines, Grain for post-training/RL, Tunix for quantisation. Google's reference training stacks Qwix for LLMs including DeepSeek-V3-class MoE, and MaxText for Flux, Wan 2.1 sit at the top, in pure JAX; MaxDiffusion sits beneath, exposed to the user as Pathways , so a single Python client can drive a job across thousands of chips and several pod-islands without giving up the JAX programming model. pathwaysutils The PyTorch path is real but second-class. uses a torch xla mechanism: every PyTorch op records into an HLO LazyTensor graph that compiles on the next barrier, with the compiled artifact cached by graph-shape hash. PyTorch/XLA 2.x added , GSPMD-style sharding annotations through an XLA backend, a torch.compile integration , and PyTorch/XLA 2.7 C++11-ABI builds with materially faster tracing. The gap to JAX is real JAX's primitives map more cleanly to JAX bridge StableHLO , and complex parallelism strategies are better-covered , which is why powered by the vLLM TPU plugin announced at Cloud Next 2025 lowers tpu-inference model, JAX-defined or PyTorch-defined, through a every . unified JAX→XLA path , announced April 2026, is Google's response: a native PyTorch experience with eager mode, TorchTPU torch.distributed , and torch.compile over XLA, on track to replace torch xla.Compared to CUDA, the TPU ecosystem is centralised, not sprawling . Almost everything below the framework XLA, JAX, Flax, Optax, Pallas, MaxText, Pathways, Shardy, Mosaic is open-sourced by Google itself, evolving in lockstep with the silicon. There are far fewer third-party kernels than CUDA's decades of accumulation; the moat is thinner where the workload looks weird, deeper where the workload looks like Gemini. The recent language is the explicit framing: chip, ICI fabric, OCS, XLA, Pathways, Pallas, MaxText, vLLM, and Pathways co-released as one product, with v8t/v8i continuing the same model under a single tpu-inference lowering path. Ironwood v7 "codesigned AI stack" and Triton narrow the gap on the NVIDIA side kernel-driven and compiler-driven are converging , but the philosophical poles are still real: torch.compile on TPU the compiler is the only interface that matters; on GPU the compiler is one of several. AMD GPU The are built on a different bet from NVIDIA: where NVIDIA each generation expands what each SM can AMD Instinct https://www.amd.com/en/products/accelerators/instinct.html GPUs do , AMD has held the conservative since Compute Unit GCN 2012 and reinvested into the package: matched or beat the contemporary NVIDIA flagship on HBM capacity every generation since 2021; the first datacenter GPU CDNA 3 ; the first coherent 3D-stacked CPU+GPU APU MI300A ; and an open ecosystem ROCm , HIP , OCP MX, UALink . Genealogy XCDs hybrid-bonded onto IODs via TSV ; FP8 ; Infinity Cache ; coherent CPU+GPU APU on MI300A; powered El Capitan . Architecture | AMD | NVIDIA | |---|---| | Compute Unit CU | Streaming Multiprocessor SM | | SIMD | SM Sub-Partition | | SIMD Lane | CUDA Core FP32 ALU | | Wavefront wave64 | Warp warp32 | | Matrix Core | Tensor Core | | MFMA | mma.sync / wgmma / tcgen05.mma | | VGPR / SGPR | Register File | | LDS Local Data Share | SMEM Shared Memory | | Infinity Fabric | NVLink | Where NVIDIA's architectural ambition lives inside each SM new tensor primitives, new async machinery, new operand stores each generation , AMD's lives between the CUs , in how many of them can be bonded into a single coherent package. The CU itself is conservative: four 16-lane SIMDs , one shared scalar unit, a 64 KB Local Data Share , an L1 vector cache, a per-SIMD VGPR file with a CU-shared SGPR pool, and since CDNA 1 a Matrix Core running MFMA . The shape hasn't meaningfully changed since GCN in 2012; what scales is the count 120 CUs on MI100 , 220 on MI250X , 304 on MI300X , 256 on MI355X and the packaging that bonds them. A wavefront of 64 threads streams across the 16 SIMD lanes over 4 cycles, with many wavefronts resident per SIMD that the scheduler switches between to hide stalls. There's nothing exotic in here; what's interesting about CDNA is everything outside the CU. Compute Inside the CU, the SIMDs and the Matrix Core run side by side. The four SIMDs handle everything element-wise: activations, normalization, residuals, address arithmetic. The Matrix Core handles the matmul. The split is the same as NVIDIA's CUDA Cores / Tensor Cores split, but the matrix abstraction has evolved on a very different curve. NVIDIA's Tensor Core climbed the thread hierarchy: a 32-thread warp on Volta , a 128-thread warp-group on Hopper , a single thread plus an optional two-SM cluster on Blackwell . AMD's Matrix Core stayed put. Every generation of MFMA from MI100 in 2020 through MI355X in 2025 is wavefront-scoped: one wave64 issues a single matrix op V MFMA , the four SIMDs cooperate to drive it, and operands come from the wavefront's register file: A and B from VGPRs , C and D usually from the dedicated AGPR file. The instruction got faster and the format set widened, but the issuer and the scope did not. The one feeder-side concession came with CDNA 4: a dedicated MFMA transpose-load from LDS that hands operands to the Matrix Core already in the layout it wants, small in spirit to NVIDIA's TMA, but the matrix op itself stayed wave-issued. The throughput numbers tell the format story directly. CDNA 1 launched in 2020 with FP32 / FP16 / BF16 / INT8 at 256 / 1024 / 512 / 1024 FLOPs per CU per cycle, with native BF16 support alongside A100 . CDNA 2 doubled the FP64 path to a full-rate matrix at 256 FLOPs/CU/cycle: uniquely AMD, the bet that put MI250X into Frontier . CDNA 3 reached parity with H100 on FP8 at 4,096 FLOPs E4M3 + E5M2 , added 2:4 structured sparsity , and added a TF32 -equivalent path that runs FP32 matmul at the FP64-matrix rate by truncating mantissas. CDNA 4 doubled again to FP4 at 16,384 FLOPs and FP6 with OCP MX block-scaling , and added mixable A/B precision in one MFMA: FP8 × FP4, for example. The same generation halved per-CU FP64 throughput, the first AMD chip to trade HPC density for AI density rather than ship both. The wavefront-scope decision shows up in two costs. Divergence. A half-empty wave64 wastes 32 lanes where a half-empty warp32 wastes 16. For workloads with mostly-uniform control flow this is a small price; for irregular workloads it hurts. Overlap. NVIDIA's asynchronous, descriptor-driven matmul decouples issue from execution: the issuing thread fires the instruction and moves on; the Tensor Core runs in the background; the warp can run softmax, apply a mask, or pre-load the next tile while the previous matmul is still in flight. AMD's wavefront-collective MFMA gives the wave no equivalent: the same wave that issued the matmul can't simultaneously do meaningful vector work while it's pending. Overlap is possible across separate wavefronts, but has to be staged in software with explicit wavefront barriers, which is more fragile and consumes more wave slots and registers. How much this matters depends on the workload. Pure dense GEMM DGEMM, the inner loop of large-batch training has nothing useful to do during the matmul; both engines saturate; async buys little. These are exactly the workloads where AMD has historically led at exascale HPC Frontier on MI250X, El Capitan on MI300A . Transformer attention FlashAttention-3 , FA4 interleaves matmul with softmax, masking, and KV-cache reads, and the async overlap is the whole structure of those kernels. AMD has to recreate the pipeline by hand, which lags NVIDIA's hardware-level support. sit in the same camp: address-irregular work that wants to run alongside the matmul. MoE dispatch, paged attention, speculative decode NVIDIA's matrix-instruction abstraction has moved further across generations warp → warp-group → single-thread async + cluster , and AMD hasn't followed. Memory AMD's memory hierarchy has fewer general-purpose tiers than NVIDIA's, with one giant cache that NVIDIA does not have at all. From the CU outward: a 64 KB LDS scratchpad software-managed, 32-bank, AMD's analog of NVIDIA's SMEM , a vector L1 16 KB on early CDNA, 32 KB from MI300X onward , a per- XCD L2 of a few MB. The L2 isn't coherent across XCDs, though; coherence happens one tier above L2. That tier is the Infinity Cache : 256 MB on MI300X, distributed across the four IODs , 16-way set-associative, ~12 TB/s measured, more than twice MI300X's 5.3 TB/s of HBM3 . It originated on RDNA gaming GPUs to compensate for narrow GDDR buses; AMD reused the IP for AI on CDNA 3, where attention KV reuse and weight reuse fit a large LLC unusually well. NVIDIA bet on bigger HBM bandwidth instead 8 TB/s on B200 , scaling with HBM4 on Rubin , and AMD bet on the cache. Off-chip, the HBM capacity grows aggressively: 32 → 64 → 128 → 192 → 256 → 288 GB across MI100 / MI210 / MI250X / MI300X / MI325X / MI350X , matching or exceeding the contemporary NVIDIA flagship in every generation from 2021 onward. The bet is that inference workloads are increasingly capacity-bound, and that the chip with more memory wins. Numerics The format trajectory tracks the precision-halving pattern that everyone in AI silicon shares: FP32 → FP16 → FP8 → FP4, restoring accuracy each step with finer-grained scaling. The AMD-specific axis is openness . CDNA 4's FP4 and FP6 use : the same numeric format as OCP MX block-scale multiplication Blackwell 's MXFP4 and TPU v8's MXU, but specified by an open consortium AMD, NVIDIA, Intel, Meta, Microsoft, Qualcomm, ARM that AMD helped found, rather than by any single vendor. The format that ships in MI355X is identical to what ships in B200 and TPU v8. The CDNA 4 inflection deserves its own line: per-CU FP64 throughput halved. MI300X served training, HPC, and inference together; MI355X is an AI chip first. The full-rate FP64-matrix bet that powered Frontier hasn't been killed, but it's no longer carrying the weight. Chiplets The packaging is where CDNA stops looking like NVIDIA and starts being something else. CDNA 1's MI100 was monolithic 7 nm. CDNA 2's MI250X was AMD's first multi-chip GPU: two Aldebaran GCDs side-by-side on a 2.5D EFB organic substrate, joined by 4 in-package Infinity Fabric links at 400 GB/s aggregate, but presented to software as two separate GPUs. CDNA 3 is the move that changed everything. Eight TSMC N5, ~115 mm² each are stacked in 3D via XCDs hybrid bonding sub-micron-pitch TSMC SoIC , no microbumps onto four TSVs TSMC N6 below. The IODs carry the I/O dies Infinity Cache , the HBM3 PHYs, the Infinity Fabric links, and PCIe Gen 5 ; each IOD hosts two XCDs above and two HBM stacks beside. The four IODs are stitched by at 4.8 TB/s bisection, so the 153-billion-transistor package looks like one GPU to the kernel: cache and address space unified at the IOD layer. NVIDIA stayed monolithic through Infinity Fabric AP H100 and only went to two reticle-limit dies on B200 via 2.5D CoWoS-L . AMD got to 3D stacking a generation earlier, at smaller per-die area: different bets on the same packaging frontier. The pushed the bet further. Replace 2 of the 8 XCDs with three Zen 4 MI300A APU , leave HBM and Infinity Cache and the IODs intact, and let CPU and GPU share one physical address space backed by HBM3 with hardware coherence. There is no host-device copy. There is no pinned memory. There is no PCIe in the path. Zen 4 cores and CDNA 3 XCDs read from the same pages. NVIDIA's CCDs Grace-Hopper bridges two packages over NVLink-C2C ; MI300A is one . 11,039 nodes of 4× MI300A is the deployment that justified it. El Capitan On CDNA 4's MI355X , eight XCDs are still 3D-stacked via SoIC onto base dies below, but the XCDs move to TSMC N3P with 32 active CUs apiece 256 total, vs 304 on MI300X ; the per-XCD count dropped to free area for bigger Matrix Cores and a 160 KB LDS . The four MI300X IODs collapse to two, each twice as wide on TSMC N6, hosting four XCDs above and four HBM3E stacks beside. Each IOD now carries its own 128 MB slice of the 256 MB Infinity Cache , half the HBM PHYs, its share of the Infinity Fabric links, and PCIe Gen 5 . Infinity Fabric AP between the two IODs runs at 5.5 TB/s bisection ~15% above CDNA 3 , and the eight stacks shift to 12-Hi HBM3E for 288 GB at 8 TB/s, 50% more capacity than MI300X on the same pin count. The package totals 185 billion transistors and still presents as one GPU to the kernel. Bets HPC and AI are the same bet Bet 1: HPC then AI. until they aren't : ship full-rate FP64 matrix from CDNA 2 through CDNA 3, then bifurcate at CDNA 4 once inference economics decisively favour low precision.Match or beat the contemporary NVIDIA flagship on HBM capacity every generation since 2021, and add a 256 MB last-level Bet 2: Memory capacity. Infinity Cache that absorbs the reuse H100 must hit HBM for.3D-stack compute on cache and I/O before NVIDIA does: TSMC Bet 3: Early 3D-stacking. SoIC hybrid-bonded XCDs on IODs in 2023, while NVIDIA stayed monolithic until 2025.The Bet 4: Coherent CPU+GPU. MI300A APU is the most chiplet-aggressive product ever shipped, and the El Capitan deployment is the proof. Bet 5: Open scale-up fabric. UALink and OCP MX over NVLink and proprietary FP4. Scaling The memory bet has a scaling consequence: when 8 MI300X chips hold 1.5 TB of HBM and 8 MI350X chips hold 2.3 TB, you can fit a 405B-parameter model in FP8 inside a single 8-GPU box weights, KV cache, and headroom for longer contexts and bigger batches , where the same model on 8× H100 640 GB requires careful sharding. For inference workloads through 2024–2025, AMD's scale-up didn't need to match NVL72 at the rack to be competitive at the box. For training at the frontier, it did, and AMD didn't have an answer until 2026. InfiniBand . Pensando NICs Pollara 400 , Vulcano 800 implement the Ultra Ethernet Consortium 's UET RDMA transport; Broadcom Tomahawk 6 supplies the switch ASIC and CPO . Scale-up Through MI355X, AMD's scale-up means an 8-GPU OAM platform over Infinity Fabric . Each MI300X has 7 IF links one to every peer in the box at 128 GB/s bidirectional, giving 896 GB/s of per-GPU mesh bandwidth in a fully-connected all-to-all topology. MI350X bumps each link to 153.6 GB/s ~1,075 GB/s per GPU but keeps the 8-GPU shape. The platform conforms to OCP's UBB 2.0: the same mechanical socket as an NVIDIA HGX baseboard, so server vendors can ship AMD or NVIDIA on the same chassis without redesigning the system. What AMD didn't ship through MI355X was a rack-scale equivalent of NVL72. Customers running larger models on MI300X clusters scaled across multiple 8-GPU boxes via Ethernet, paying scale-out latency for what NVIDIA users could keep inside scale-up. This was the gap that mattered for training, and the gap that is built to close. Helios Helios is AMD's first rack-scale scale-up domain, shipping in 2H 2026 alongside MI455X . 72 GPUs per rack, ~31 TB HBM4 , 1.4 PB/s aggregate HBM bandwidth, 2.9 ExaFLOPS FP4 / 1.4 ExaFLOPS FP8, 260 TB/s of scale-up bandwidth, 43 TB/s of scale-out. The form factor is Open Rack Wide ORW Meta's 2025 OCP submission, double-wide and liquid-cooled , not an AMD-proprietary chassis. Building on Meta's reference design rather than designing a rack from scratch is a deliberate AMD bet: any hyperscaler standardised on ORW can deploy Helios without bespoke datacenter facilities work. The fabric is : Ultra Accelerator Link, an open consortium standard AMD helped found alongside Apple, AWS, Cisco, Google, HPE, Intel, Meta, Microsoft, and Synopsys. UALink 200G 1.0 April 2025 defines a 200 GT/s lane and 800 Gbps per direction, with switched topologies scaling to 1,024 accelerators per pod. The promise is a cache-coherent interconnect comparable to NVLink but unowned: any vendor can build a UALink switch, any accelerator can talk UALink, the standard belongs to the consortium rather than to the strongest seller. UALink The catch: native UALink switching silicon won't ship in volume until 2027 . Astera Labs' Scorpio, plus competing parts from Auradine, Enfabrica, and Xconn, are all targeting late-2026 / 2027 deployment. Helios at launch uses Infinity Fabric tunnelled over standard Ethernet as a stopgap, preserving the programming model while waiting for native UALink fabric. Native UALink switching arrives with MI500 in 2027. At launch, Helios is closer to a fast Ethernet-tunnelled coherent cluster than to NVL72's true cache-coherent NVLink domain: a real concession on the timeline, paid in exchange for hitting 2H 2026 with a competitive product. UALoE Scale-out AMD does not ship InfiniBand . The whole scale-out stack is Ethernet, anchored on a different open standard: the . Ultra Ethernet Consortium UEC UEC 1.0 released June 2025 defines : a new RDMA transport over standard Ethernet, with packet spraying, SACK-based selective retransmit, and modern congestion control. UET is not RoCEv2 which encapsulates InfiniBand transport in Ethernet frames ; it's a clean redesign of RDMA semantics for scale-out AI fabrics. AMD is a founding member alongside Broadcom, Cisco, Meta, and Microsoft. Same play as UALink: own the standard, not the implementation. Ultra Ethernet Transport UET The NIC is , the networking startup AMD acquired in 2022. Pensando is the current AI NIC: 400 GbE, P4-programmable, UEC-ready, PCIe Gen 5, paired with MI300X / MI355X. Pollara 400 ships in 2026 alongside MI455X: UEC 1.0 compliant, PCIe Gen 6, native UALink interfaces, 8× the per-GPU scale-out bandwidth of Pollara. Vulcano 800 is the front-end DPU 16× Arm Neoverse-N1, dual 400 GbE for storage / SDN / firewall, equivalent to NVIDIA's Salina 400 BlueField , distinct from the AI back-end NIC. The switch silicon, though, isn't AMD's. Helios's 43 TB/s scale-out fabric runs through : a 102.4 Tbps Ethernet switch ASIC with co-packaged optics "Davisson" . AMD has no in-house Broadcom Tomahawk 6 CPO and no in-house switch ASIC; the optical layer is partner silicon. NVIDIA owns its entire stack: InfiniBand, Spectrum-X Ethernet, ConnectX, BlueField, Quantum-X Photonics CPO, all in-house. AMD owns one tier NIC + DPU via Pensando and bets that open standards plus best-of-breed partner silicon will outpace vertical integration. The industry has moved AMD's way. Dell'Oro reports Ethernet handled more than twice the AI scale-out fabric volume of InfiniBand in 2025; AWS, Microsoft, Meta, Oracle, and xAI have all standardised on Ethernet for their AMD-based AI clusters. The remaining question isn't whether Ethernet can match InfiniBand on RDMA semantics UEC closes that gap but whether Helios can close the rack-scale gap with NVL72 fast enough to win frontier training workloads that today default to NVIDIA. Software is the open-source counterpoint to ROCm https://rocm.docs.amd.com/ . Where NVIDIA's stack is proprietary and vertically integrated cuBLAS, cuDNN, TensorRT-LLM ship as binary blobs maintained by NVIDIA alone , ROCm is GitHub-native and bets on open standards PyTorch, Triton, vLLM, OCP MX rather than a walled-garden library set. The software gap with NVIDIA is real, but AMD's strategy is to close it through the open community rather than build a parallel CUDA stack from scratch. CUDA https://docs.nvidia.com/cuda/cuda-c-programming-guide/ The bottom of the stack is , AMD's CUDA-compatible C++ runtime. HIP translates CUDA source to HIP automatically. Bulk HPC code HACC, Laghos, QMCPack ports at 80–95% out of the box: the CORAL-2 number. Modern AI kernels port worse: anything that reaches for Hopper- or Blackwell-specific primitives hipify TMA descriptors, , wgmma tcgen05.mma has no clean ROCm analog and has to be rewritten by hand.Above HIP sits a library tier structured to mirror NVIDIA's, one-to-one by name: for cuBLAS; rocBLAS https://github.com/ROCm/rocBLAS for cuBLASLt; hipBLASLt https://github.com/ROCm/hipBLASLt for cuDNN; MIOpen https://github.com/ROCm/MIOpen for NCCL; RCCL https://github.com/ROCm/rccl and its modern Composable Kernel ck-tile DSL for CUTLASS; rocprofv3 / rocprof-sys / rocprof-compute for the Nsight family. There is no first-party analog of TensorRT-LLM, though. AMD's answer is to back as the open-source serving engine and ship AMD-specific operators vLLM https://github.com/vllm-project/vllm that plug into it; the dedicated ROCm CI for vLLM took test-pass rate from 37% to 93% across early 2026. AITER The PyTorch path is first-class. Eager-mode PyTorch has run on ROCm since 2018; torch.compile lowers through Triton, and Triton's ROCm backend with AOTriton for ahead-of-time math kernels is upstream. There is no XLA-style intermediate IR; ROCm compiles direct to HIP / Triton / CK. As Triton becomes the default kernel path in PyTorch, much of the porting cost evaporates: a kernel that runs through torch.compile works on both CUDA and ROCm without source change. This is the architectural bet beneath AMD's open strategy: Triton's Python DSL becomes the cross-vendor lingua franca that sidesteps the need for a CUDA-equivalent kernel ecosystem. is the load-bearing case. FlashAttention is production on MI300X via Composable Kernel; PyTorch defaults to CK or AOTriton on ROCm. FA2 Hopper-tuned is partially supported via AITER + CK, but Dao-AILab's canonical implementation remains CUDA-only. FA3 Blackwell, March 2026 has no ROCm port at all. FA4 , Hazy Research's MI355X port of ThunderKittens November 2025 , claims forward-pass parity with hand-tuned AITER in ~500 lines. The pattern: open-source academic kernels close the AMD tail months after NVIDIA's, not years. HipKittens https://hazyresearch.stanford.edu/blog/2025-11-09-hk Production deployment has validated the strategy. Microsoft Azure's ND MI300X v5 instances went GA in May 2024; OpenAI runs GPT inference on them. Meta ships Llama 3 / Llama 4 inference on MI300X via the Grand Teton platform. Oracle OCI's went GA in September 2024, with MI355X following in 2026. These are real serving fleets at hyperscaler scale, not pilots. BM.GPU.MI300X.8 The honest gap is still real. Independent benchmarks Phoronix, March 2026 put ROCm 7.2 at 10–25% slower than equivalent CUDA on standard PyTorch / vLLM / SGLang workloads, at equivalent precision on equivalent silicon. ROCm 7 reached feature parity but not perf parity . The FlashAttention-4 tail research code that exploits Blackwell's newest primitives is where NVIDIA's moat remains most durable; it has no clean ROCm analog and waits for a hand-written AITER kernel or HipKittens-class community port. NVIDIA ships engineers inside frontier labs; AMD ships kernels through GitHub. The strategies converge on common workloads Llama inference, attention, dense transformer training but the long tail of novel research code still costs MI300X / MI355X deployments engineering time NVIDIA users don't pay. Cerebras WSE builds the Cerebras https://www.cerebras.ai/ . The philosophy: the largest chip ever shipped memory wall is a consequence of cutting the wafer. A fab prints dozens of dies onto 300 mm of silicon and saws them apart; the industry then spends its most exotic engineering HBM , NVLink , CoWoS , 5,184 copper cables per rack wiring the pieces back together at a small fraction of on-die bandwidth. Cerebras skips the saw. The is one piece of silicon: 84 Wafer-Scale Engine reticle fields , 46,225 mm², 900,000 dataflow cores, and every byte of on-chip memory in SRAM one cycle from a compute unit. Genealogy Architecture A GPU is a hierarchy: threads inside warps inside SMs, dies inside packages inside racks, each boundary with its own bandwidth, its own latency, its own programming construct; every accelerator built from dies inherits some version of it. The WSE is a flat plane : 900,000 identical cores tiled edge-to-edge in a 2D mesh, with no shared cache, no global memory, and no boundary of any kind between one core and the other 899,999. Each core is tiny, ~38,000 µm² on WSE-2 , roughly half SRAM and half logic, peaking at 30 mW: 48 kB of local SRAM, sixteen general-purpose registers, a six-stage pipeline, a 4-wide FP16 FMAC SIMD 8-wide on WSE-3 , and a five-port router into the fabric. Execution is : a core sits idle until a dataflow arrives, control bits in the wavelet select which handler task fires, and eight hardware wavelet switch cycle-by-cycle as tensor operands arrive and drain. No warps, no microthreads warp schedulers , no caches to miss, no reorder buffer: the arrival of data is the schedule . The Wafer A stepper exposes a wafer one reticle at a time, ~850 mm² per shot, which is why every conventional chip lives under that ceiling and why B200 became two dies the moment NVIDIA pressed against it . Cerebras prints the same ~550 mm² die 84 times in a 12×7 grid, like any other customer of TSMC, and then, in a process co-developed with TSMC, lays extra high-level metal across the <1 mm scribe lines where the saw would normally run. The mesh crosses each seam on a source-synchronous parallel interface 2,880 GB/s per die on WSE-3 , and the entire inter-die layer costs ~97 W. To software the seams do not exist: one uniform mesh, one chip. Wafer-scale has been tried before and it failed on yield: a single defect in a monolithic wafer-computer kills the whole wafer, which is what buried the idea in the 1980s. Cerebras's answer is granularity. A defect on an H100 disables an entire ~6 mm² SM; the same defect on a WSE disables one 0.05 mm² core. WSE-3 fabricates ~970,000 cores and ships 900,000: the ~7% spare pool, plus redundant fabric links, lets the hardware remap around every defect and restore a full logical mesh. The Core The unusual part of the core is not the datapath; it is what an instruction is . Alongside the sixteen general-purpose registers sit 44 , each holding a tensor descriptor: data-structure registers DSRs base address , extent , and stride , up to four dimensions. Instructions name their operands by DSR, so a single FMAC instruction says multiply the arriving stream against this resident tensor and accumulate into that one , and the hardware streams elements for as long as the tensor lasts. There is no software loop around the multiply and no instruction fetch per element; the loop lives in the descriptor. NVIDIA spent five Tensor Core generations walking the matmul toward a single descriptor-driven command ; on a WSE core, a tensor instruction has no other form. Sequencing is the fabric's job. A color is a statically routed virtual channel with a handler task bound to it at compile time, so sending a wavelet on a color is invoking code on the destination core: the 16 control bits are the call, the 16 data bits the argument. The task scheduler holds the in-flight tensor operations on the core's eight microthreads and switches among them every cycle by operand availability. It is the same stall-hiding job a warp scheduler does with 64 resident warps, done with eight contexts, because the latency being hidden is a busy SRAM bank or a neighbour hop, not an HBM round trip. The 48 kB of local SRAM is organised for the datapath rather than for locality: eight single-ported 6 kB banks deliver two 64-bit reads and one 64-bit write every cycle, exactly two 4-element FP16 operands in and one result out, the width of the WSE-2 FMAC. A 256-byte software-managed cache 512 B on WSE-3 keeps the hottest values beside the pipeline. This is the machine's thesis in miniature: per core, memory bandwidth and compute are matched exactly, and the wafer inherits that balance 900,000 times over. Compute There is no matrix unit on the wafer. NVIDIA, Google, and AMD all concentrate their FLOPs in a dedicated matmul engine Tensor Core , MXU , Matrix Core and differ mainly in how that engine is fed; Cerebras assembles matmul out of the fabric. A GEMM runs as a wafer-wide choreography: each arriving weight is broadcast along a row of cores holding activations, every core fires a multiply-accumulate against its resident slice an AXPY per weight , and partial sums reduce across the mesh. The data reuse a Tensor Core gets from a register tile and an MXU gets from its wiring, the WSE gets from geometry: activations never move, so the only operand in flight is the one being multiplied. The FLOPs ledger needs care, because the number Cerebras prints is not the number to compare. WSE-3's headline 125 PFLOPS is sparse FP16 : it assumes the hardware's roughly 8× zero-skipping payoff on ideally sparse tensors. Dense is roughly derived: 900,000 cores × 8-wide FMAC × 1.1 GHz; Cerebras publishes no official dense figure . That is real compute, but it is not the point: per watt, dense FLOPs on the wafer lose to every contemporary GPU. The wafer was never a FLOPs machine. It is a 15.8 PFLOPS FP16 , and the FLOPs exist to keep up with the SRAM. bandwidth machine Zero-skipping is where dataflow earns its keep. Because computation is triggered by arriving data, a zero never triggers anything: zeros are filtered at the sender , and the receiving core never sees them and never spends the cycle. This is unstructured, element-granular sparsity, the general case that NVIDIA's 2:4 structured sparsity only samples. It is also, so far, an unexercised option. Cerebras's own sparse-pretraining results SPDF https://arxiv.org/abs/2303.10464 : 75% sparsity at 1.3B parameters; a follow-up at 6.7B are vendor-authored and sub-7B, and no flagship customer model has been disclosed as sparse-trained: Jais 2 , the biggest run on the hardware, is dense. The only silicon that can harvest unstructured sparsity has yet to ship a headline model that uses it. Memory The hierarchy is one tier: 44 GB of SRAM in 48 kB slices inside the cores, and nothing else on the wafer . No HBM, no L2, no eviction policy; every byte is one cycle from an FMAC. The quoted bandwidth is 21 PB/s, and the number deserves its flag: it is the sum of 900,000 local SRAM ports, an on-wafer aggregate, not a point-to-point link, and not comparable to an HBM figure. The honest comparison is bytes per FLOP: the wafer can feed ~1.3 bytes per dense FP16 FLOP, where a B200 gets ~0.002 from HBM. On that axis every GPU and TPU is starved; the WSE is the only machine in balance. Decode , the phase that is a pure bandwidth problem one full read of the weights per token , is the phase the wafer turns out to be shaped for. The other side of the tier is the edge of it. The wafer's connection to everything else is 12×100 GbE: 1.2 Tb/s , barely more than the single ConnectX-8 NIC attached to one Blackwell GPU. Between on-wafer SRAM and off-wafer Ethernet sit . NVIDIA's hierarchy descends gradually, each tier a few times slower than the last; the WSE has two tiers with a cliff between them. The wafer is an island, and the island's superpower and its cage are the same fact. five orders of magnitude And the island is not growing. SRAM density has effectively stopped scaling on leading nodes: WSE-3 carries just 10% more SRAM than WSE-2 despite a full node shrink and a 54% jump in transistor count. Logic keeps shrinking; the six-transistor SRAM cell does not. The architecture's scarcest resource is the one thing the next process node no longer buys. Weight Streaming Training on the wafer inverts the flow everyone else takes for granted: on a GPU or TPU, weights are resident and activations stream through; on a WSE, activations are resident and weights stream through . Master weights live in , a DRAM-and-flash appliance beside the cluster. Layer by layer, weights stream across the wafer, trigger multiply-accumulates against the activations pinned in SRAM, and leave; gradients stream back out on the backward pass, and the optimizer step runs inside MemoryX on CPUs a weight update is O parameters of element-wise work with no reuse, so CPU-class compute keeps pace . The wafer never stores weights, "not even temporarily" MemoryX Cerebras's phrase https://www.kisacoresearch.com/sites/default/files/documents/cs weight streaming white paper - cerebras.pdf . Model size is bounded by MemoryX, not by the 44 GB; the 44 GB bounds activations and batch. What this buys is the programming model. One wafer holds a full layer's activations, so there is no tensor parallelism , no pipeline parallelism , no FSDP sharding: a 70B model is written as a single-device program, and multi-system scaling is pure through data parallelism , a broadcast/reduce tree that fans one weight stream out to N wafers and sums their gradients on the way home. The parallelism-strategy spreadsheet that dominates GPU training simply has no Cerebras page. SwarmX What it costs is scale, in the market's own revealed preference. The spec sheet says 2,048 CS-3s; the largest cluster ever disclosed is 64 Condor Galaxy 3 . The largest from-scratch model ever disclosed on the platform is Jais 2 at 70B parameters and 2.6T tokens , trained by anchor customer G42 with Cerebras engineers embedded. Nothing above 70B, from anyone, in the seven years since CS-1. And utilisation MFU , the number GPU labs publish as a matter of course at 35–45%, has never been disclosed for any Cerebras run. Numerics The numerics fit in a sentence: FP16 and BF16 with FP32 accumulate , plus from WSE-3 a 16-wide 8-bit integer path that the Hot Chips disclosure labels fixed-point. No FP8, no FP4, no microscaling. While every other vendor halves precision each generation and buys the accuracy back with block scaling, Cerebras still computes in 16-bit and markets it as a quality differentiator "the original 16-bit weights" . The tension is obvious: SRAM capacity is the architecture's scarcest resource, and 8-bit weights would halve the number of wafers a model needs. Whether 16-bit-only is numerical conviction or a datapath roadmap gap is an open question; no primary Cerebras source shows floating-point 8 anywhere on the wafer. Bets The die boundary is the tax the rest of the industry pays: SerDes, interposers, HBM stacks, cables, switches. Stitch 84 reticle fields in metal and the highest-bandwidth boundary in rival systems does not exist at all. Bet 1: Don't cut the wafer. Trade capacity for bandwidth at the steepest ratio in the industry: 44 GB at an on-wafer aggregate 21 PB/s. Balance the machine instead of hiding imbalance behind a hierarchy. Bet 2: SRAM is the only memory. 900,000 tiny cores triggered by arriving wavelets, with matmul assembled from broadcast, FMAC, and mesh reduction: skipping a zero is free rather than a special mode. Bet 3: Dataflow cores, no matrix unit. Weight streaming decouples model size MemoryX from wafer memory 44 GB and collapses cluster scaling to pure data parallelism. Bet 4: Weights move, activations stay. The wafer re-reads an entire model per token faster than anything built on HBM; price that speed as a premium product instead of competing on cost per token. Bet 5: Sell latency, not throughput. Scaling Scale-up and scale-out mean something different here. NVIDIA's scale-up problem make 72 packages behave like one device is solved on the WSE by lithography: the coherent domain ships from the fab in one piece. What remains is everything past the wafer's edge, and no other machine hits its edge as hard or as early. colors , native broadcast, 214 Pbit/s aggregate fabric bandwidth. Fixed at 46,225 mm² by the size of a 300 mm wafer. Scale-up The wafer's internal fabric has no SerDes , no cables, no transceivers, and no marginal cost per link: routing is compiled, each hop is one cycle, and a broadcast is a native fabric primitive rather than a switch feature. Where NVL72 spends 5,184 copper cables and a tray of NVSwitch ASICs to give 72 GPUs 130 TB/s of all-to-all, the WSE's equivalent domain is a single lithographic object. The catch is that the domain size is a constant. NVIDIA's scale-up domain grows every generation NVL72 to NVL576 across three years ; the wafer has been 46,225 mm² since 2019 and will stay there. 300 mm is the largest wafer the industry runs the 450 mm transition died a decade ago , so Cerebras's scale-up roadmap is whatever the next node yields in density: there is no more area to be had. Scale-out Training scale-out is SwarmX , and it only does one thing: replicate. Broadcast the weight stream to N wafers, reduce their gradients on the return path; batch grows with system count, model size does not. The claimed ceiling of 2,048 systems "256 exaFLOPS", sparse has never been built; 64 has. Inference abandons weight streaming entirely; the arithmetic is fatal. Streaming a 70B model's 140 GB from MemoryX for every decoded token over a ~150 GB/s pipe would cost roughly a second per token. So inference parks the weights in SRAM and shards the model across wafers at layer boundaries: Llama 70B on "as few as four" CS-3s, pipeline-parallel over Ethernet, each additional wafer contributing 44 GB of weight-plus- KV capacity and 23 kW of load. The speeds are real, and independently verified. measured 1,850 tokens/s on Llama 3.1 8B and 446 on 70B at the August 2024 launch, 969 on Llama 405B 240 ms to first token , and 2,522 on Llama 4 Maverick in 2025, ~2.4× the best published Blackwell number of the time. Vendor-quoted peaks run higher 2,100 on 70B with Artificial Analysis speculative decoding ; 3,000 on GPT-OSS-120B, where the live independent measurement sits nearer 2,000 . No GPU provider comes close on per-user decode speed. The economics are the sharp edge. Forty-four GB per wafer means a frontier-scale model consumes fleets: SemiAnalysis https://newsletter.semianalysis.com/p/cerebras-faster-tokens-please estimates ~24 CS-3s for a 1.6T-parameter-class model that fits in a handful of GPU racks, each system an analyst-estimated ~$450k bill of materials selling at a list price around $2–3M never officially disclosed . During decode the wafer's enormous FLOPs mostly idle; Cerebras has declined to disclose batch sizes and has never published per-system throughput. Per-token API pricing runs roughly 3–5× GPU-based providers for the same open models, and Llama 405B was quietly dropped from the API, which SemiAnalysis reads as serving economics that didn't clear. Fixed SRAM also prices context: KV cache lives in the same 44 GB as weights, so long contexts steal capacity and force more systems per replica; the API caps at 131K tokens while frontier providers serve 256K–1M. MoE is served Qwen3-235B at ~1,500 tokens/s, vendor-quoted but is the format's worst case: a huge parameter footprint touched a few experts at a time, held in the most expensive memory. The market has priced this honestly. Mistral's Le Chat ~1,100 tokens/s , Perplexity Sonar, and Meta's Llama API all pay for the latency; in January 2026 OpenAI signed for 750 MW of CS-3 capacity through 2028 , reported above $10B https://www.cnbc.com/2026/01/14/cerebras-scores-openai-deal-worth-over-10-billion.html at signing and since grown past $20B https://finance.yahoo.com/technology/ai/articles/cerebras-systems-openai-tout-20b-040208708.html , the largest endorsement wafer-scale has ever received. The first flagship to ship on that capacity is , launched July 2026 at a quoted 750 tokens/s. GPT-5.6 Sol https://openai.com/index/gpt-5-6/ Software The stack is compiler-driven like the TPU's, but through a much narrower aperture: the Cerebras compiler is a kernel matcher , not a general code generator. cerebras.pytorch traces the training step through lazy tensors into Torch-MLIR and a graph IR, then matches subgraphs against a library of hand-written kernels, falling back to slower auto-generated ones for ops with no match. The documented constraints https://training-api.cerebras.ai/en/rel-2.4.0/wsc/tutorials/cstorch-limitations.html are stark by GPU standards: static graphs only, no dynamic shapes, no data-dependent control flow, no eager tensor access mid-step, and a PyTorch version pinned behind upstream. The best independent practitioner account SURF https://servicedesk.surf.nl/wiki/spaces/WIKI/pages/112592526/Evaluation+Cerebras+CS-2 , the Dutch national compute centre reports unsupported layer types and no 1:1 porting path for standard PyTorch code. And there is no kernel escape hatch. CUDA's answer to a novel attention variant is write a kernel ; the TPU's is Pallas ; ROCm's is Triton . The Cerebras ML stack has no user kernel path at all: when the matcher misses badly, the fix is a Cerebras engineer. A separate SDK language, , exposes the raw machine tasks, wavelets, colors and has produced striking HPC results a CSL TotalEnergies stencil code https://arxiv.org/abs/2204.03775 at ~228× an A100, a Gordon Bell finalist on 48 CS-2s , but it is a separate world, unconnected to the PyTorch flow. Every flagship model on the platform Jais, BTLM, Med42 was co-developed with embedded Cerebras staff. There is a strange immunity in this. FlashAttention , the defining kernel lineage of the GPU era, is a scheme for tiling attention through a memory hierarchy, and the WSE has no hierarchy to tile against: the optimisation class that costs AMD years of porting lag simply does not apply. But the immunity and the poverty are the same fact. The third-party kernel ecosystem that compounds on CUDA has no surface to attach to here; every kernel improvement in the platform's history has one author. Where does that leave the wafer? Owning a real niche, honestly won: batch-one decode speed, independently verified, paid for by customers who price latency above cost. Around the niche, hard walls: 3–5× per-token pricing, a 70B training ceiling seven years in, revenue still ~86% concentrated in two Abu-Dhabi-linked customers in 2025 per the S-1 filings around its May 2026 IPO , and a scarcest resource, SRAM density, that stopped scaling just as models kept growing. Hennessy and Patterson promised a Cambrian explosion; the WSE is its most extreme body plan, the one that decided the memory wall was a packaging choice and spent 46,225 mm² of silicon refusing to make it. AWS Trainium , the team behind AWS's Annapurna Labs cards and Nitro CPUs, built Graviton as a Trainium . The compute core takes the TPU's proven playbook a 128×128 fast-follower weight-stationary systolic array , software-managed scratchpads, whole-program compilation down to sharing Google's compiler outright. The scale-out fabric is the XLA https://openxla.org/xla -offloaded network that already carries the rest of AWS. What is genuinely Amazon's is narrow and deliberate: dedicated collective-communication silicon bolted onto the borrowed core, and the vertical integration to price a chip that only has to beat NVIDIA Nitro inside AWS . Genealogy FP8 acceleration, 96 GB HBM3 ; the 64-chip UltraServer . Powers Project Rainier . MXFP8/MXFP4 ; the NeuronSwitch all-to-all fabric replaces the torus. 144-chip UltraServer. Architecture The other captive-silicon story belongs to Google, and Trainium is best read as the TPU's thesis rebuilt inside a different cloud. The bets underneath are the same a systolic array fed from software-managed SRAM , scheduled ahead of time by a compiler, with no caches and no thread scheduler , but the unit is assembled differently. A Trainium chip carries a small number of NeuronCores 2 on Trn1 , 8 on Trn2 and Trn3 , and each NeuronCore is not one monolithic matmul engine but a : a cluster of decoupled, specialised engines the 128×128 systolic array , a Tensor Engine for reductions, a Vector Engine for pointwise math, and a programmable Scalar Engine of eight 512-bit vector processors for whatever fits none of the other three. Around them sit the data-movers: 128 GPSIMD Engine , a DMA engines that sequences transfers, and from Trn2 dedicated Sync Engine for collectives. There are no warps and no wavefronts; the engines run as a statically-scheduled dataflow pipeline, and the load-bearing design decisions are about what surrounds the systolic array, not the array itself. CC-Cores Compute The owns the matmul FLOPs; the other three engines own everything else. It is a 128×128 grid of processing elements 16,384 Tensor Engine MACs run weight-stationary : one operand tile is loaded into the array and held in place LoadStationary , the other streams through it MultiplyMoving , and partial sums land in , a small accumulator SRAM the engine can read-add-write so a contraction longer than 128 folds into place along the axis. This is the same tile PSUM MMA at the heart of every matmul accelerator; but where NVIDIA wraps it in the warp hierarchy and Google issues it from a VLIW bundle, Trainium exposes it as a pair of explicit instructions against a named scratchpad. The array is physically fixed at 128×128 across all three generations; what changes is how many products it packs per cell. Trn1 's NeuronCore-v2 ran BF16 /FP16 with FP32 accumulate and offered FP8 only at the BF16 rate no speedup . Trn2 's v3 double-pumps FP8 to present an effective 256×128 array, the first Trainium with a real 2× on 8-bit. Trn3 's v4 packs microscaling operands to present an effective 512×128 at 4× the BF16 rate. The count of physical multiply-add cells never moves; the datapath just feeds them narrower numbers. The other three engines are what keep the array busy. The handles cross-element reductions layernorm, softmax, pooling ; the Vector Engine handles one-in-one-out pointwise ops activations, GELU ; the Scalar Engine , eight fully-programmable vector processors running C, absorbs anything that maps to none of them. A well-compiled step overlaps all four: the Tensor Engine grinds a matmul while the Vector Engine runs the previous tile's softmax and the DMA engines stage the next, the same producer/consumer overlap that makes TPU and GPU attention kernels efficient, expressed here as separate physical engines rather than separate warps or VLIW slots. The design pays off when a layer decomposes cleanly onto the four engine types, which transformers largely do. It pays a tax at the edges: an operator that fits none of the specialised engines falls to the programmable GPSIMD Engine GPSIMD path, slower, and the part of the machine most likely to bottleneck a novel architecture. It is Trainium's version of the long-tail cost every non-GPU accelerator carries. Memory The memory hierarchy is the compute philosophy applied to storage: three tiers, all software-managed, no hardware cache anywhere . AWS's own documentation draws the contrast, noting that unlike a CPU or GPU the NeuronCore has no cache and that "all memory movement is explicit in the program itself." Off-chip is 32 GB on Trn1, 96 GB HBM HBM3 on Trn2, 144 GB HBM3e on Trn3 . On-chip, closest to the engines, is the : the main scratchpad, roughly 20× HBM bandwidth, organised in 128 partitions and sized per NeuronCore at 24 MiB v2 , 28 MiB v3 , 32 MiB v4 . Between the array and SBUF sits State Buffer SBUF , a 2 MiB accumulator dedicated to matmul outputs. Data moves HBM → SBUF → Tensor Engine → PSUM → SBUF, every hop issued by the compiler; nothing is prefetched or evicted by hardware. PSUM This is exactly Google's VMEM bet, an explicit scratchpad the compiler must schedule perfectly with no cache to paper over a mistake, and the opposite of NVIDIA's hardware-managed L2 and L1 . Trainium inherits both the ceiling and the fragility that come with it: when the schedule is right the engines never stall, and when it is wrong there is no fallback path. The design runs a generous HBM budget against modest peak FLOPs, so per unit of compute Trainium carries more memory than a comparable NVIDIA part. On absolute capacity, though, it trails: Trn2's 96 GB sits below the H200 and B200 , and Trn3's 144 GB 2025 sits below the 192 GB B200 and 288 GB B300 it ships against. So the lever AWS actually pulls when it argues the economics of serving a large model is not memory leadership but price : cost per unit of compute and HBM, on silicon it builds and rents itself. Numerics Trainium tracks the same precision-halving curve as everyone else FP32 → BF16 → FP8 → FP4 , with two Trainium-specific wrinkles. The first is : rather than fix configurable FP8 E4M3 and E5M2 like Hopper, the Tensor Engine takes an adjustable exponent bias and supports E5M2, E4M3, and E3M4, letting the compiler trade range for precision per tensor. The second is that Trn3 's FP4 buys no extra throughput : OCP MXFP4 operands are up-converted to MXFP8 before they reach the array, so FP4 runs at the FP8 rate and saves only memory and bandwidth, not compute. Both generations lean on the industry's accuracy-recovery tricks: microscaling block exponents from Trn3, and hardware on every generation. The one figure to distrust is the sparse peak: AWS headlines a 4× FP8 number that its own architecture pages put at 2× over dense FP8 the 4× is relative to dense BF16 , so the marketed acceleration and the datapath do not quite agree. stochastic rounding Collectives in Silicon The block with no clean analogue on a GPU is the . Distributed training and inference spend a large fraction of their wall-clock in collective-communication core collectives : every gradient step is an all-reduce , every MoE layer an all-to-all . On a GPU those collectives run as NCCL kernels on the same SMs doing the math, so communication and compute contend for the same silicon and the overlap has to be won in software. Trainium carves the function out into dedicated hardware: 20 per Trn2 chip, wired straight to the CC-Cores ports, executing all-reduce, all-gather, reduce-scatter, and all-to-all while the Tensor and Vector engines keep running. It is the same move Google made with NeuronLink SparseCore and Cerebras made with its off-core zero filter: find a workload the main engine is the wrong shape for, and spend a little area on a purpose-built block beside it rather than steal cycles from the core. Communication becomes something the chip does concurrently , not something it pauses to do. Bets Annapurna designs chip, server, rack, Bet 1: The cloud is the product, the chip is a component. Nitro network, and cloud API as one stack, so Trainium only has to win on price-performance inside AWS, never on a merchant-silicon spec sheet.A 128×128 Bet 2: Borrow the compute thesis, don't reinvent it. weight-stationary array, software-managed SBUF / PSUM scratchpads, and whole-program compilation are the TPU's bets, reused down to sharing Google's OpenXLA . The effort saved goes into the network and the rack.Dedicated Bet 3: Collectives belong in silicon. CC-Cores overlap all-reduce and all-to-all with compute in hardware, instead of running them as kernels that steal FLOPs from the matmul units.Scale-out is Bet 4: Reuse the cloud's own network. EFA with the SRD transport: the same Nitro -offloaded, packet-sprayed RDMA that already runs the rest of AWS. No InfiniBand .Trn1 and Trn2 copied the TPU's Bet 5: Move the topology to the workload. torus ; Trn3's NeuronSwitch replaces it with a switched all-to-all fabric as MoE traffic outgrew nearest-neighbour. Honestly, this is following the playbook: first Google's, now NVIDIA's. Scaling Trainium's scaling inherits its split from the rest of AWS: a tightly-coupled domain for the chips that must act as one, and the cloud's general-purpose NeuronLink fabric for everything beyond it. The scale-up domain is not cache-coherent shared memory the way EFA NVLink is; AWS markets the UltraServer as a pooled multi-terabyte memory, but underneath it is message-passing over point-to-point links, closer in spirit to the TPU's ICI than to an NVSwitch crossbar. NeuronLink binds chips into one UltraServer . Through Trn2 the topology is a torus 16 chips per instance in a 4×4 2D torus, 64 per UltraServer in a 4×4×4 3D torus ; Trn3 replaces it with the NeuronSwitch all-to-all fabric. Message-passing, not coherent load/store. Elastic Fabric Adapter over Ethernet, offloaded to Nitro . The SRD transport sprays each flow across many paths and delivers reliably but out-of-order; UltraClusters reach hundreds of thousands of chips over the 10p10u fabric. Scale-up NeuronLink is Trainium's chip-to-chip fabric, the role NVLink plays for NVIDIA and ICI for the TPU. Through Trn2 it wires chips into a , exactly the TPU's choice: a single torus instance is 16 chips in a 4×4 2D torus at ~1.28 TB/s per chip, and the trn2 joins four instances into 64 chips on a 4×4×4 3D torus, presenting 83 dense Trn2 UltraServer FP8 PetaFLOPS and ~6 TB of HBM as one scale-up domain. The third torus axis is deliberately thin the inter-instance ring runs at ~256 GB/s per chip against 1.28 TB/s inside an instance , which is the torus's characteristic trade: cheap wiring and huge nearest-neighbour bandwidth, at the cost of many hops across the diameter. AWS positions the 64-chip UltraServer against NVIDIA's 72-GPU NVL72 ; the aggregate compute is in the same league, but a torus is not a crossbar , and the two behave very differently on traffic that is not nearest-neighbour. That trade is why Trn3 abandons the torus. is a switched NeuronSwitch-v1 fabric that roughly doubles inter-chip bandwidth and, more importantly, flattens the diameter so any chip reaches any other in one switched hop. The Trn3 UltraServer scales to 144 chips for 362 dense FP8 PetaFLOPS and 20.7 TB of all-to-all HBM3e . The motivation is the one that also pushed Google toward high-radix topologies for MoE inference: expert routing is all-to-all, the worst case for a torus, and a switch turns the longest-hop pair into a single crossing. Trainium's interconnect roadmap is a compressed re-run of the industry's: adopt the torus while the workload is nearest-neighbour, switch to a crossbar when it is not. Scale-out Scale-out is not bespoke; it is the same fabric AWS already runs. Every Trainium instance carries an Elastic Fabric Adapter NIC into the datacenter network 3.2 Tbps per Trn2 instance , and the transport is , offloaded to the SRD Scalable Reliable Datagram cards rather than run on the accelerator. SRD is AWS's clean-sheet answer to Nitro RDMA : instead of the single ordered flow of RoCE or InfiniBand , it sprays each message across up to 64 parallel paths and delivers reliably but out-of-order, pushing reassembly up to the collective library and sidestepping the head-of-line blocking a single congested path would cause. It is the transport AWS built for its cloud generally, repurposed for the accelerator fabric. At the top of the hierarchy is the , stitched together by the UltraCluster network AWS's shorthand for ~10 petabits/s of bandwidth at under 10 microseconds of latency across a datacenter and scaling to hundreds of thousands of chips. The proof point is 10p10u : roughly half a million Trainium2 chips across multiple US datacenters, brought online for Project Rainier Anthropic in late 2025; by early 2026 Claude was running on over a million chips, the largest commitment any external lab has made to a non-NVIDIA training platform. It exists because the economics close end to end. AWS claims Trainium2 delivers 30–40% better price-performance than its Hopper -class GPU instances an AWS figure, measured against last-generation NVIDIA rather than Blackwell , and because Amazon owns every layer from the Nitro card to the API, that margin is Amazon's to set. Software Trainium's software makes the borrowing explicit: the is a Neuron SDK https://awsdocs-neuron.readthedocs-hosted.com/ . The Neuron compiler compiler-first stack built on the same OpenXLA foundation as the TPU neuronx-cc ingests XLA HLO graphs and lowers them to a binary that the Neuron runtime loads onto the NeuronCores; the front-end IR is Google's, and Google's own OpenXLA announcements list Trainium as a first-class NEFF PJRT device alongside the TPU. runs PyTorch through torch-neuronx PyTorch/XLA 's LazyTensor tracing record ops, compile the graph at a step boundary , and lowers JAX through jax-neuronx StableHLO . On the spectrum from kernel-driven CUDA at one pole to whole-program XLA at the other, Trainium sits almost on top of the TPU: the compiler is the system, and it is largely the same compiler. Where it diverges is the escape hatch. XLA alone cannot always synthesise the optimum for a novel attention variant or a fused MoE dispatch, so Neuron ships , a Python, tile-level kernel language that exposes the four engines and the NKI Neuron Kernel Interface SBUF / PSUM scratchpads directly. It is Trainium's or its Pallas : the same idea of a tile DSL that drops beneath the whole-program compiler when a kernel's win is in the Triton schedule , not the algebra. Below it, a maps collective-communication library all-reduce and all-to-all onto the CC-Cores and the NeuronLink topology the NCCL analogue , and provides the sharded-training layer. NeuronX Distributed The gap to CUDA and even to the TPU's stack is maturity, not design. NKI, the JAX path, and the distributed library were all still in beta through late 2024; a ported model runs only on AWS, with no cross-vendor fallback; and the vLLM backend trails the upstream project. The clearest tell is how the anchor tenant works: Anthropic does not simply target Trainium through PyTorch, it embeds with Annapurna, writes its own low-level NKI kernels, and upstreams fixes into the Neuron stack. Trainium is production-viable at the frontier, but at the frontier it is co-engineered, not turnkey: the compiler is inherited and excellent, but the surrounding ecosystem is young. Groq LPU The is a Groq https://groq.com/ LPU machine. Every other chip spends silicon tolerating uncertainty: caches to hide memory latency, schedulers to fill stalls, arbiters to resolve contention it cannot predict. The LPU deletes all of it. Strip out every deterministic component no cache, no branch predictor, no arbiter, no reorder buffer, not even an on-chip crossbar and hand the entire scheduling problem to the compiler, which places every instruction and every byte on an exact cycle. What is left is a chip whose latency is known before it runs. Where the reactive moved scheduling into the compiler but kept TPU HBM and a dynamic network, Groq removed the last sources of nondeterminism: memory is all SRAM , and the network is scheduled too, so hundreds of chips run as one clock-exact program. Genealogy Jonathan Ross , who started Google's TPU as a 20% project, leaves to build a deterministic inference chip. ISCA 2022: software-scheduled networking extends the deterministic schedule across thousands of chips via a compiled Dragonfly . SF4X ; it never shipped a reported failed tapeout . Language Processing Unit ; the company pivots from selling cards to selling tokens, on record decode speeds. non-exclusive license to the LPU technology and hires Ross and much of the team. GTC 2026 as a latency co-processor beside Rubin NVL72, via Attention-FFN disaggregation . Architecture The rest of the field is built from a replicated core : tile one SM , TensorCore , CU , or dataflow core across the die and farm work out to the copies. The LPU is built the other way. It takes a single conventional core and : instruction control, the vector ALUs, the matrix units, the memory, and the network each become a pulls it apart , a full-height column of identical hardware, and the columns stand side by side across the die. Homogeneous down each slice, heterogeneous across the chip. Data does not sit in a register file waiting to be issued onto a unit; it functional slice horizontally through the slices like parts down an assembly line, East and West, one register hop per cycle, while streams VLIW instructions issue Northward from the control slices to meet it. Nothing in the datapath reacts: the compiler knows where every operand is on every cycle, and the hardware just turns the clock. The streaming is the identity: this design launched as the , and carried that name until the 2024 rebrand to Tensor Streaming Processor TSP Language Processing Unit . The vertical axis is SIMD width. The chip is 320 lanes tall, organised as 20 of 16 lanes each a 21st is a spare, fused out for yield and invisible to software , and every slice acts on all 320 lanes at once. The horizontal axis is time. There are 64 logical superlanes per lane, 32 flowing East and 32 West, and on every tick each stream advances one slice in its direction until it is consumed or falls off the edge of the die. A slice reads operands off the passing streams, computes, and writes results back onto streams bound for the next slice. The die is mirrored into two hemispheres around a central vector unit, so a value produced once can be consumed by slices on either side. stream registers Compute The LPU keeps the same division of labour as everything else, matrix work on dedicated units and the rest on a vector engine, but arranges both as slices in the stream. The matrix path is the : four independent 320×320 multiply-accumulate planes two per hemisphere , 409,600 multipliers in all, taking INT8 or FP16 operands into INT32 or FP32 accumulators. Weights install across a plane all of them in under 40 cycles , then activations stream through and products accumulate. At 900 MHz that is roughly MXM , and, unusually, the number carries no sparsity asterisk: the TSP refuses to skip zeros at all, because a data-dependent skip would make execution time data-dependent, and determinism is the one property it will not trade. 750 INT8 TOPS and 188 FP16 TFLOPS The vector path is the in the centre of the die: 16 ALUs per lane arranged as a 4×4 mesh, 5,120 32-bit ALUs, running activations, normalisation, quantisation, and residual adds. Because compute is VXM rather than issued to a shared unit, an operand can march through a chain of VXM ALUs and straight into an MXM plane on consecutive cycles without touching memory: the operator fusion a GPU kernel builds by hand is here just the physical order of the slices. A third slice type, the spatial , handles the movement the straight-line stream cannot express: lane shifts, a 320-lane permute, transposes, and the chip-to-chip links all live here, so rearranging data across lanes is a first-class operation rather than a round-trip through SRAM. SXM Memory There is no HBM, no DRAM, and no cache. On-chip is the slices: 230 MB of SRAM in 88 slices 44 per hemisphere , every byte a single cycle from a compute slice, ~80 TB/s aggregate. That is the whole hierarchy: one tier, flat, software-addressed, with none of the eviction, prefetch, or coherence machinery that would introduce a variable-latency access. MEM The consequence is the defining constraint of the architecture. 230 MB does not hold a model. Llama-2 70B in FP16 is 140 GB, so it has to be sharded across hundreds of chips , its weights spread over the aggregate SRAM of a whole rack or more: the deployed configuration was ~576 LPUs. Where a GPU parks the model in HBM on a handful of packages and streams tokens past it, the LPU spreads the model in SRAM across a cluster and streams tokens through the cluster. The chip count is set by capacity, not compute: the weights have to fit. It is the same trade Cerebras makes SRAM only, no HBM , reached from the opposite direction: Cerebras keeps one enormous die and gives up capacity per wafer; Groq keeps a normal-sized die and gives up ever fitting a model on one. Numerics The numerics are the road not taken. Every other vendor here has been halving precision each generation, FP16 to FP8 to FP4 with block scaling to buy the accuracy back. The TSP stayed at FP16 and INT8 with FP32 accumulate and never shipped FP8 or FP4 in silicon. Its one numeric idea is : a 320-element dot product fused into a single rounding step with FP32 accumulation, so an FP16 multiplier array lands close to FP32 accuracy on the reduction Groq reports ~0.05% max error against an FP32 baseline . TruePoint Whether 16-bit was conviction or a datapath that never got its low-precision refresh is hard to separate from the fact that the second-generation chip never shipped. SRAM capacity is the architecture's scarcest resource, and 8-bit weights would halve the chips a model needs; a machine this capacity-bound had every reason to want FP8 and did not get it on silicon. It is the same open question that hangs over Cerebras's 16-bit-only datapath, and the same tension: the vendor most starved for capacity computing at the widest precision. Determinism Every other accelerator hides latency; the LPU exposes it. The ISA carries the execution latency of each instruction, the datapaths are fixed-latency by construction, and so the compiler computes ahead of time the exact cycle on which every result appears. Nothing in the hardware can disturb that schedule: no cache to miss, no arbiter to stall on, no branch to mispredict, no speculation to unwind. Groq's own measurement is the proof: 24,240 runs of BERT-Large returned inside a ~75 µs band, and the compiler's predicted latency sat within 2% of measured. This is the TPU's instinct move scheduling into the compiler, delete the hardware that second-guesses it taken one step further. The TPU compiler schedules a chip; the LPU compiler schedules a system , because the determinism holds across the network too. And it is the exact inverse of Cerebras, whose cores are , firing whenever an operand happens to arrive: the WSE reacts to data, the LPU is timed to it. Both machines delete the scheduler; one replaces it with arrival, the other with a clock. dataflow Bets Delete every reactive component caches, arbiters, predictors, reorder buffers and let the compiler own every cycle. Bet 1: Determinism over tolerance. Disaggregate the core into slices and stream operands through them, so fusion is the floorplan and data reuse lives in the wires, not a register-file dance. Bet 2: Spatial functional slices. No HBM, at any capacity cost. Trade the ability to hold a model on-chip for single-cycle, fixed-latency access, accept models must span hundreds of chips. Bet 3: SRAM is the only memory. Make the chips their own routers and compile the communication cycle-by-cycle, so a thousand-chip cluster is one deterministic program with no switches and no congestion. Bet 4: Schedule the network too. Optimise for tokens per second per user at batch 1, the regime GPUs are worst at, and price that speed as the product rather than competing on cost per token. Bet 5: Sell latency, not throughput. Scaling Scaling an LPU is unlike anything else here, because there is no separate scale-up fabric to build: the chip is already a switch. Each LPU carries up to 16 chip-to-chip links 11 exposed on the card and acts simultaneously as a compute endpoint and a router. Wire the chips directly to each other and the cluster is a RealScale : no glueless multiprocessor NICs , no switch ASICs , no top-of-rack switch. And because determinism holds across those links, the entire cluster runs on one compile-time schedule. Dragonfly of nodes: 9 per rack 72 chips, one node a hot spare , scaling to a spec'd 10,440 chips, every hop still on a compiled, deterministic schedule. Scale-up The node is 8 LPUs, fully connected: 7 of each chip's links wire it to the other seven, so every chip in the node is one hop from every other. The remaining four links on each chip 32 across the node bundle into what the ISCA paper calls a 32-port virtual router, the node's uplink into the larger fabric. There is no baseboard switch and no coherent address space; a remote operand is not loaded, it is scheduled to arrive, injected by the source chip on a cycle the compiler chose and consumed by the destination on the cycle it lands. Scale-out Beyond the node, nodes wire into a : 9 nodes make a 72-chip rack the ninth a hot spare, so 64 active , and the topology scales to a specified 10,440 chips with any two under six hops apart. The fabric is Dragonfly : routing and flow control move to compile time, and the paper's framing is blunt, software-scheduled scheduled, not routed . There is no back-pressure and no dynamic arbitration, because the compiler has already proven the receiver is ready; links carry forward error correction instead of retransmission, because a retry would perturb the schedule. Keeping a rack of independently-clocked chips in lockstep is its own problem: the links are , and the fabric maintains a global consensus time with plesiochronous exchanged every 256 cycles over a spanning tree, with periodic deskew instructions stalling each chip back into alignment. The payoff Groq reports is that an 8-way Hardware-Aligned Counters all-reduce matches an A100 / NVSwitch node on large tensors and beats it on small ones, where a scheduled fabric pays none of the handshake latency a dynamic one does. The cost is written into the physics of the memory bet. A model replica is not a box, it is a rack or eight : Llama-2 70B on ~576 chips carried, by one analysis, 144 host CPUs and 144 TB of host RAM alongside the LPUs, against two CPUs for an 8-GPU server. The wafer under each chip is cheap 14 nm GlobalFoundries, reportedly under $6k, against ~$16k for an H100-class part , but you need hundreds of them, and during decode most of their enormous compute sits idle while the SRAM does the work. put it plainly: the LPU wins the bill of materials per token when you optimise for latency, and loses to GPUs by roughly an order of magnitude on throughput per dollar once you batch. The architecture is not competing on cost. It is competing on speed. SemiAnalysis https://newsletter.semianalysis.com/p/groq-inference-tokenomics-speed-but Software The programming model is the purest expression of the compiler is the machine . There are no kernels . You hand the Groq compiler a model from PyTorch , TensorFlow, or ONNX ; it lowers to a small tensor op set and statically schedules every instruction, every stream, and every chip-to-chip transfer. Nobody writes a or hand-tunes a tile, because there is no dynamic hardware to hand-tune against. Groq's demonstration was bringing up LLaMA in four days with a team of under ten, against the months of hand-kernel work the same model took to tune on a GPU. The stack around the compiler a profiler, a runtime, the wgmma GroqFlow bring-up path is small and closed, and GroqFlow was archived in 2025 as the company stopped selling cards and started selling tokens.That pivot is the tell about what the architecture is for. The LPU is inference-only by construction Ross's framing is that training is a local game and inference a global one , and it is unbeaten at a single thing: single-user decode latency. Independent measurement backs the claim, with clocking Groq among the fastest token-per-second providers on open models. It is badly matched to the rest: a model that will not fit in a rack of SRAM, a workload that wants big batches for throughput-per-dollar, or dynamic control flow a static schedule cannot express. Artificial Analysis https://artificialanalysis.ai/providers/groq MoE is served, but its data-dependent expert routing sits awkwardly against a compiler that wants to know everything in advance, and Groq has published little on how it reconciles the two. The epilogue is that the buyer of all this was NVIDIA. In December 2025 NVIDIA took a to the LPU technology and hired Ross and much of the team. It was not an acquisition: no products, customer contracts, or equity changed hands, per NVIDIA's own 10-K, though the roughly $13B paid at closing led the press to call it one. At non-exclusive license GTC 2026 the technology reappeared as the , a rack of 256 SRAM-only inference chips sitting beside NVIDIA Groq 3 LPU Rubin NVL72 and splitting the transformer between them: the GPUs run attention , the LPUs run the feed-forward and MoE layers, with Dynamo orchestrating the hand-off. The most deterministic architecture in AI ended up as a latency co-processor inside the most programmable one. GroqCloud, meanwhile, still serves tokens on the original 14 nm silicon. Comparison All arithmetic figures are peak values at the stated precision; entries are dense unless the vendor does not publish the basis. Memory bandwidth is the native tier shown: HBM for GPUs, TPUs, and Trainium; aggregate on-chip SRAM for Cerebras and Groq. Those numbers are not directly comparable. Scale-up bandwidth follows each vendor's convention and can mean per-chip aggregate, rack aggregate, or true bisection. Per-chip | Company | Year | Chip | Accelerator memory | Memory BW | Flagship dense FLOPs | TDP | Scale-up BW | |---|---|---|---|---|---|---|---| | 2023 | H100 SXM5 | 80 GB HBM3 | 3.4 TB/s | 1.98 PetaFLOPS FP8 | 700 W | 900 GB/s | | | 2024 | H200 SXM | 141 GB HBM3e | 4.8 TB/s | 1.98 PetaFLOPS FP8 | 700 W | 900 GB/s | | | 2024 | B200 | 192 GB HBM3e | 8 TB/s | 4.5 PetaFLOPS FP8 / 9 PetaFLOPS FP4 | 1,000 W | 1.8 TB/s | | | 2025 | B300 | 288 GB HBM3e | 8 TB/s | 7.5 PetaFLOPS FP8 / 15 PetaFLOPS FP4 | 1,400 W | 1.8 TB/s | | | 2026 | Rubin | 288 GB HBM4 | ~13 TB/s | ~17 PetaFLOPS FP8 / ~50 PetaFLOPS FP4 | ~1,500 W | 3.6 TB/s | | | 2027 | Rubin Ultra | 1 TB HBM4e | ~32 TB/s | ~33 PetaFLOPS FP8 / ~100 PetaFLOPS FP4 | ~1,800 W | 3.6 TB/s | | | 2023 | TPU v5p | 95 GB HBM2e | 2.8 TB/s | 0.46 PetaFLOPS BF16 | n/d | 1.2 TB/s | | | 2025 | TPU Ironwood v7 | 192 GB HBM3e | 7.4 TB/s | 4.6 PetaFLOPS FP8 | n/d | 1.2 TB/s | | | 2026 | TPU v8t Sunfish | 216 GB HBM3e | 6.5 TB/s | 12.6 PetaFLOPS FP4 | n/d | n/d | | | 2023 | MI300X | 192 GB HBM3 | 5.3 TB/s | 2.6 PetaFLOPS FP8 | 750 W | 896 GB/s | | | 2024 | MI325X | 256 GB HBM3e | 6.0 TB/s | 2.6 PetaFLOPS FP8 | 1,000 W | 896 GB/s | | | 2025 | MI355X | 288 GB HBM3e | 8 TB/s | 10 PetaFLOPS FP8 / 20 PetaFLOPS FP4 | 1,400 W | 1,075 GB/s | | | 2026 | MI455X | TBD | TBD | ~40 PetaFLOPS FP4 | TBD | n/d | | | 2021 | WSE-2 | 40 GB SRAM on-wafer | 20 PB/s aggregate | 7.5 PetaFLOPS FP16 | 23 kW system | domain = the wafer | | | 2024 | WSE-3 | 44 GB SRAM on-wafer | 21 PB/s aggregate | ~15.8 PetaFLOPS FP16 | 23 kW system | domain = the wafer | | | 2022 | Trainium1 | 32 GB HBM2e | 820 GB/s | 0.19 PetaFLOPS BF16/FP8 | n/d | n/d | | | 2024 | Trainium2 | 96 GB HBM3 | 2.9 TB/s | 1.3 PetaFLOPS FP8 | ~500 W | 1.28 TB/s | | | 2025 | Trainium3 | 144 GB HBM3e | 4.9 TB/s | 2.5 PetaFLOPS FP8 | n/d | n/d | | | 2020 | GroqChip 1st-gen TSP/LPU | 230 MB SRAM | 80 TB/s on-chip aggregate | 0.188 PetaFLOPS FP16 | 215 W | 330 GB/s 11-link card | | | 2026 | NVIDIA Groq 3 LP30 | 500 MB SRAM | 150 TB/s on-chip aggregate | ~1.2 PetaFLOPS FP8 | n/d | 2.5 TB/s | Per-rack / pod | Company | Year | System | Chips | Aggregate dense FLOPs | Accelerator memory total | Scale-up fabric BW | | |---| What this shows Per-chip FP8 has converged. B200 4.5 PF , Ironwood 4.6 PF , and MI355X 10 PF sit within ~2× of each other. The per-chip arms race is close; the rack and pod are where the architectures diverge. HBM capacity is AMD's persistent win. 192 → 256 → 288 GB across 2023–2025 has matched or beaten NVIDIA every generation. NVIDIA caught up at 288 GB only with B300 late 2025 ; Rubin Ultra retakes the lead at 1 TB / package in 2026. Rack-scale scale-up is NVIDIA's win until 2026. GB200 / GB300 NVL72 was the only coherent rack-scale domain shipping in 2024–2025; AMD scaled up at the box and didn't reach rack scale until Helios. The TPU sidesteps the question: its torus is the rack and the cluster at once. TPU pods dwarf any NVIDIA rack in chip count. Ironwood pod = 9,216 chips for 42.5 ExaFLOPS FP8; NVL576 = 576 GPUs for ~5 ExaFLOPS FP8. The TPU's flat-rate-per-chip × massive-pod recipe yields more aggregate compute per system, at the cost of per-chip bandwidth. Power per chip is rising fast. 700 W Hopper → 1,000 W Blackwell, MI325X → 1,400 W B300, MI355X → ~1,800 W Rubin Ultra, analyst . Liquid cooling becomes mandatory above ~1,000 W; air cooling effectively ends with Hopper. Scale-out NIC bandwidth doubles each NVIDIA generation. 400 Gbps CX-7, Hopper → 800 Gbps CX-8, Blackwell → 1.6 Tbps CX-9, Rubin . AMD lags one generation Pollara 400 → Vulcano 800 , reflecting Pensando's smaller install base and later integration. Cerebras breaks the table's axes. No HBM at all: 44 GB of on-wafer SRAM at an aggregate 21 PB/s, ~1.3 bytes per dense FLOP where the GPU rows sit near 0.002. The cost is visible in the same row: less total memory than a single H200, dense FLOPs per watt behind every contemporary GPU, and an empty scale-up column because the coherent domain is the wafer itself. Trainium competes on economics, not the spec sheet. Per-chip it trails Trn2's 1.3 PF FP8 is roughly a quarter of MI355X , but the Trn2 UltraServer reached 64-chip rack-scale scale-up in 2024 alongside NVL72, as a message-passing torus rather than a coherent crossbar, and Trn3 pivots to the switched NeuronSwitch fabric. AWS owns every layer from the Nitro card to the API, and one anchor tenant Anthropic, over a million Trainium2 chips validates it at frontier scale. Groq trades capacity for SRAM bandwidth, then scales the memory pool with chip count. The first GroqRack exposes only 14 GB across 64 active chips; Groq 3 LPX grows that to 128 GB across 256 chips at 40 PB/s aggregate SRAM bandwidth. Its 12 TB DDR5 tier and pairing with Rubin show that the LPU complements, rather than replaces, a large-memory GPU rack.