NVIDIA GPU

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 Streaming Multiprocessors, replicated 100+ times per package: 80 on V100, 108 on A100, 132 on H100, 148 on B200, 160 on B300, 224 on Rubin. Inside every SM sits the same recipe: four SM Sub-Partitions, each with its own warp scheduler, dispatch unit, 16k×32-bit register file, scalar CUDA Core lanes, a Special Function Unit for transcendentals, and a private port into the SM's Tensor Cores. The four partitions share an L1/shared-memory block, and the TMA. Threads are grouped into warps of 32 that execute in SIMT lock-step; dozens of resident warps per partition let the scheduler hide memory/arithmetic stalls by switching between them.NVIDIA GPU由一组面向吞吐量的核心、深层内存层次结构以及足够的调度硅片组成，以保持数千个线程同时运行。核心是流式多处理器（SM），在每个封装中复制100多次：V100上有80个，A100上有108个，H100上有132个，B200上有148个，B300上有160个，Rubin上有224个。每个SM内部的配置相同：四个SM子分区，每个分区拥有独立的线程束调度器、分发单元、16k×32位寄存器堆、标量CUDA核心通道、用于超越函数计算的特殊功能单元（SFU），以及通往SM Tensor Core的私有端口。四个分区共享一个L1/共享内存块和TMA。线程被分组为32个一组的线程束（Warp），以SIMT锁步方式执行；每个分区内有数十个驻留线程束，调度器通过在它们之间切换来隐藏内存/算术延迟。

Compute计算

CUDA Cores 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 Tensor Cores.CUDA核心是最初的计算吞吐量来源，对于AI而言，它们依然负责处理非矩阵乘法的所有任务：激活、残差相加、归一化、地址计算。但是，Transformer块中约99%的浮点运算（FLOPs）属于矩阵乘法，因此绝大部分计算吞吐量来自Tensor Core。

These cores execute fused matrix multiply-accumulate on small matrix tiles, $D = A \cdot B + C$ The full matmul is broken into output tiles: to produce one output tile, a kernel walks the shared inner dimension $K$, drawing $A$ from a row-strip of the left input matrix and $B$ from a column-strip of the right, and folds each partial product into a running accumulator. $C$ holds the partial sum so far, $D$ is the updated value carried into the next step. After the inner loop completes, $D$ is one finished tile of the full output matrix; the whole matmul is built from many of these tile MMAs.这些核心在小矩阵瓦片（Tile）上执行融合矩阵乘加指令，即D=A⋅B+C。完整的矩阵乘法被分解为输出瓦片：为了生成一个输出瓦片，内核会遍历共享的内部维度K，从左侧输入矩阵的行带中提取A，从右侧输入矩阵的列带中提取B，并将每个部分乘积累加到运行中的累加器中。C保存了迄今为止的部分和，D是更新后的值，被带入下一步。内部循环完成后，D就是一个完整的输出矩阵瓦片；整个矩阵乘法由许多这样的瓦片乘加运算构建而成。

Tile shapes are written M × N × K, $M \times N$ is the output tile size, and $K$ is how much of the inner dimension the instruction contracts over in one fire; the rest of the matmul's $K$ axis is walked by the kernel's inner loop. The accumulator is sticky across that loop: each MMA's output $D$ becomes the next MMA's input $C$, so the equation is really $C \leftarrow A \cdot B + C$ in place: successive instructions fold their partial products into the same storage until the K-axis is fully walked.瓦片形状记为M×N×K，M×N是输出瓦片大小，K是指令在一次激发中收缩的内部维度大小；矩阵乘法K轴的其余部分由内核的内部循环遍历。累加器在循环中保持不变：每次MMA的输出D成为下一次MMA的输入C，因此方程实际上是原地进行的C←A⋅B+C：后续指令将它们的部分乘积累加到相同的存储中，直到K轴完全遍历完成。

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 warp group: 128 cooperating threads firing an asynchronous wgmma at 64×256×16 shape that runs in the background while the issuing warps load the next tile. B200's 5th-gen unit went further still: a two-SM MMA of 256×256×16 with operands split across a pair of SMs, native FP4, and a dedicated 256 KB Tensor Memory (TMEM) scratchpad per SM that holds accumulator tiles instead of bleeding into the register file. 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.V100的第一代单元（每个SM有8个）运行线程束级的16×16×16 FP16 MMA。A100的第三代单元增加了TF32、BF16、FP64矩阵乘法以及2:4结构化稀疏。H100的第四代单元增加了原生FP8支持，并将抽象层从线程束提升到线程束组（Warp Group）：128个协同线程以64×256×16的形状触发异步wgmma，在后台运行，同时发出指令的线程束加载下一个瓦片。B200的第五代单元更进一步：采用256×256×16的双SM MMA，操作数跨两个SM拆分，支持原生FP4，每个SM配备专用的256 KB张量内存（TMEM）暂存器，用于存放累加器瓦片，避免溢出到寄存器堆。Rubin的第六代单元扩展了FP4吞吐量，增加了原生FP6支持，并与第三代Transformer引擎配对，该引擎在硬件中执行自适应NVFP4微块缩放，将每瓦片的量化元数据保持在Tensor Core路径上，而不是通过CUDA核心。

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 wgmma.mma_async widens the issuer to a 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.这六代产品中保持不变的是：矩阵乘法位于线程/线程束层次结构内，但发出指令所需的线程数已减少，且指令的发出与执行已解耦。Volta的mma.sync是线程束集体且同步的：线程束中的所有32个线程一起执行，每个通道持有A、B和累加器D的寄存器片段，且线程束会阻塞直到完成。Hopper的wgmma.mma_async将发出者扩展为128个线程的线程束组，将B移入共享内存描述符（A变为可选：可以是寄存器或描述符，由内核选择），并立即返回：矩阵乘法在后台运行，同时线程束组排队下一个瓦片，通过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 CTA-pair variant scales the same model across two SMs: one thread on each SM in a paired cluster issues coordinated 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.Blackwell的tcgen05.mma完成了这一迁移：A与B一起进入共享内存描述符（或A直接来自TMEM），累加器D落在TMEM而不是寄存器堆中。由于每个操作数都离开了通道，发出指令时无需协调每个线程的状态，因此单个线程即可触发指令并立即返回，通过消费者线程束等待的mbarrier信号来标记完成。同时，线程束的其余部分以及发出指令的线程本身可以自由执行其他工作。CTA对变体在两个SM之间扩展了相同的模型：配对集群中每个SM上的一个线程发出协调的MMA，在对之间共享操作数，在相同的异步/mbarrier完成机制下构成256×256×16的双SM瓦片，只需提升到集群级屏障，使两者保持同步。

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.矩阵乘法变得越来越大，同时对发出指令的线程负担越来越轻：最初作为32个通道锁步执行的指令，现在更接近于单个由描述符驱动的命令，从线程束模型内部发出，但不再由其执行。

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.这种解耦使得Transformer注意力内核在GPU上非常高效。线程束可以在矩阵乘法运行期间执行Softmax、应用掩码或预加载下一个瓦片；矩阵乘法与周围逐元素运算的重叠是现代注意力内核（FlashAttention-3, FA4）的结构基础，这取决于矩阵指令不会阻塞线程束。

Memory内存

The on-chip hierarchy is hardware-managed caches at every level, with software hints layered on top. Off-chip is HBM: 32 GB 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 L2 Cache 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 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 L1/SMEM 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.片上层次结构是各级硬件管理的缓存，之上叠加了软件提示。片外是HBM：V100配备32 GB HBM2，H100配备80 GB HBM3，B200配备192 GB HBM3e，B300配备288 GB，Rubin配备288 GB HBM4。芯片级L2缓存位于HBM和SM之间：V100为6 MB，A100为40 MB，H100为50 MB，B200为60 MB（拆分为双芯封装中的两个30 MB存储体，具有位置感知驻留控制，使热瓦片可以固定在附近的芯片上）。在每个SM内部，256 KB的统一L1/共享内存（SMEM）在内核启动时在硬件管理的L1和程序员控制的暂存器之间进行分区。寄存器堆是每个SM约256 KB的另一个存储层，跨分区分为四部分。

Blackwell adds a fifth tier: TMEM, 256 KB per SM dedicated to MMA accumulators and addressed only by the Tensor Core, pulling the operand-residency pressure out of the general register file.Blackwell增加了第五层：TMEM，每个SM 256 KB，专门用于MMA累加器，仅由Tensor Core寻址，将操作数驻留压力从通用寄存器堆中释放出来。

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 cp.async: 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. Hopper replaced that with the TMA, 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 mbarrier. The whole warp is freed from load issue and address math; the kernel just queues descriptors. TMA also supports 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.各层之间的数据移动已逐渐与线程束解耦。Ampere之前，加载瓦片是同步的：每个线程发出自己的全局加载，线程束阻塞直到每个片段进入寄存器，并通过第二遍复制到共享内存；每个瓦片都在地址计算和等待上消耗了线程束通道。Ampere引入了cp.async：每线程异步复制HBM→SMEM，完全绕过寄存器，线程束提交正在进行的复制组，仅在消费者需要数据时才等待。Hopper用TMA（专用DMA引擎）取代了它：一个线程提交多维瓦片描述符（基地址、前导维度、重排），引擎处理所有地址计算并写入共享内存，完成情况由mbarrier发出信号。整个线程束从加载发出和地址计算中解放出来；内核只需排队描述符。TMA还支持集群级多播：一次HBM读取扇出到线程块集群中的每个SM，将曾经的N次独立加载转变为一次。Blackwell再次扩展了TMA：直接加载到TMEM，因此累加器瓦片无需通过SMEM暂存即可直接流入。每一代产品都在减少线程束处理每个瓦片所需的工作量。

Warp Specialisation线程束专业化

The Hopper-era programming idiom is warp specialisation: inside one block, some warps act as producers that issue back-to-back TMA loads; others act as consumers that fire wgmma on freshly-arrived tiles. Synchronisation between them is no longer the old SM-wide __syncthreads() barrier; it is mbarrier (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 (FlashAttention-3, CUTLASS ping-pong GEMMs, the Blackwell 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 thread-block clusters (Hopper+) tying multiple SMs into one cooperative compute unit so that the two-SM MMA of Blackwell composes naturally on top.Hopper时代的编程习惯是线程束专业化（Warp Specialisation）：在一个块内，一些线程束充当生产者，发出背靠背的TMA加载；另一些充当消费者，在刚到达的瓦片上触发wgmma。它们之间的同步不再是旧的SM级__syncthreads()屏障；而是mbarrier（共享内存中的内存屏障）和附加到TMA完成情况的异步事务屏障，允许在线程束粒度而非块粒度进行细粒度的生产者/消费者握手。成为现代注意力内核（FlashAttention-3, CUTLASS乒乓GEMM, Blackwell FA4内核）参考的模式是相同的配方：TMA驱动的生产者流水线通过共享内存和TMEM为wgmma消费者流水线提供数据，配合mbarrier握手和线程块集群（Hopper+），将多个SM连接成一个协同计算单元，使Blackwell的双SM MMA自然地在之上组合。

Numerics数值计算

FP32 was the historical default; Volta brought FP16 with FP32 accumulate and the loss-scaling tricks that made it trainable; Ampere added TF32 (FP32 range, FP16 mantissa, drop-in for FP32 matmul), BF16, and 2:4 structured sparsity that doubles effective throughput on pruned weights. Hopper introduced native FP8 in both E4M3 and E5M2, paired with the Transformer Engine which auto-scales activations layer-by-layer to keep them inside FP8 dynamic range. Blackwell halved precision again with FP4 and shipped microscaling MX formats (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 NVFP4 (NVIDIA's tightened FP4 variant) and native FP6 with more aggressive sparsity. The chip layout itself is now part of the numerics story: B100/B200/B300 are two reticle-limit dies stitched by a ~10 TB/s NV-HBI 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.FP32是历史默认值；Volta带来了支持FP32累加的FP16以及使其可训练的损失缩放技巧；Ampere增加了TF32（FP32范围，FP16尾数，FP32矩阵乘法的直接替代品）、BF16以及在修剪后的权重上使有效吞吐量翻倍的2:4结构化稀疏。Hopper引入了原生FP8（E4M3和E5M2），并配合Transformer引擎，该引擎逐层自动缩放激活，使其保持在FP8动态范围内。Blackwell再次将精度减半，引入FP4并发布微缩放MX格式（块级共享指数，恢复了FP4损失的大部分精度），以及第二代Transformer引擎，将自动缩放流水线重定向至FP4。Rubin的第三代Transformer引擎增加了NVFP4（NVIDIA的强化FP4变体）和具有更激进稀疏性的原生FP6。芯片布局本身现在也是数值计算故事的一部分：B100/B200/B300是两块通过约10 TB/s NV-HBI链路缝合的光刻极限芯片，在软件层面呈现为单个逻辑GPU，封装上有8个HBM堆栈；Rubin将小芯片配方扩展到双芯，晶体管数量约3360亿，带有8个HBM4堆栈。每一代产品通过减半位数并利用更细粒度的缩放方案恢复精度，以及越来越多地将更多硅片结合到封装中，实现了约2倍的每瓦吞吐量增长。

Bets押注

• Bet 1: Programmability. The workload is a moving target (attention variants, novel model architectures), so keep every block programmable and let the developer write CUDA. Even the specialised units are exposed through that model rather than as fixed-function blocks.押注1：可编程性。工作负载是一个移动的目标（注意力变体、新型模型架构），因此保持每个块的可编程性，并让开发者编写CUDA。即使是专用单元也通过该模型暴露，而不是作为固定功能块。
• Bet 2: Hide Latency with Massive Multithreading. 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 warp scheduler picking a ready warp every cycle.押注2：通过大规模多线程隐藏延迟。延迟是不可预测且依赖于数据的，因此不通过静态调度来隐藏它，而是通过大规模线程超额订阅（每个SM最多64个驻留线程束），由硬件线程束调度器在每个周期选择一个就绪的线程束。
• Bet 3: Warp-wrapped Matmul. 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 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.押注3：线程束包装的矩阵乘法。矩阵单元是压倒性的计算吞吐量来源，但它必须位于其他所有功能使用的相同线程束/线程抽象之后，因此将其包装在mma.sync → wgmma → tcgen05.mma中，而不是将其暴露为固定功能管道。这使得单个内核可以在一次传递中融合矩阵乘法、Softmax和逐元素运算。
• Bet 4: Async Memory Hierarchy. Make the 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, mbarrier for the producer/consumer handshake. The hierarchy is software-pipelined inside a programmable kernel, not statically scheduled by a compiler against a known-latency scratchpad.押注4：异步内存层次结构。使内存层次结构显式且由程序员管理，而不是隐式且由编译器调度。保留L2缓存，但将SMEM和TMEM暴露为命名暂存器，并在其上叠加异步机制：用于批量复制的TMA，用于矩阵乘法累加器的TMEM，用于生产者/消费者握手的mbarrier。层次结构在可编程内核内部进行软件流水线处理，而不是由编译器针对已知延迟的暂存器进行静态调度。
• Bet 5: Amortised SIMT Tax. Every transistor spent on a warp scheduler, register-file, or coherent cache is a transistor not spent on a 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.押注5：摊销SIMT税。花费在线程束调度器、寄存器堆或一致性缓存上的每一个晶体管都是没有花费在MAC（乘加单元）上的晶体管；接受这种税收，并以两种方式偿还：Tensor Core现在足够大，使得SIMT机制在更大的MAC计数上被摊销；以及TMEM等单元以牺牲部分通用灵活性为代价换取MAC密度。

Scaling扩展

There are two regimes for scaling: scale-up and scale-out.扩展有两种模式：纵向扩展（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.AI基础设施同时使用两者：带宽敏感的集合通信（张量并行、MoE专家路由）留在纵向扩展域内；数据并行和流水线并行则跨越横向扩展结构。

Scale-up

The scale-up stack is NVLink plus NVSwitch. NVLink implements a cache-coherent fabric between GPUs, so a load or store on one GPU can target another GPU's 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 crossbar chip that every GPU connects to, routing traffic so every GPU can simultaneously communicate with every other at full NVLink bandwidth, non-blocking and all-to-all. 纵向扩展栈是NVLink加NVSwitch。NVLink在GPU之间实现了缓存一致性结构，因此一个GPU上的加载或存储可以以硬件处理地址转换和一致性的方式寻址另一个GPU的HBM。但NVLink本身是点对点的：一个链路仅连接两个芯片。NVSwitch是一个专用的交叉开关芯片，每个GPU都连接到它，路由流量，使每个GPU都能以全NVLink带宽同时与所有其他GPU通信，无阻塞且全对全。

Together they defined the HGX 8-GPU baseboard, pairing eight H100 SXM modules with x86 hosts (AMD EPYC or Intel Xeon) over PCIe Gen5. Hopper also shipped a Grace-paired form: the GH200 Grace Hopper Superchip bonded one Grace ARM CPU to one H100 over NVLink-C2C at 900 GB/s, eliminating the PCIe host-device hop. Modules scaled up into GH200 NVL2 pairs and rack-level GH200 NVL32. Blackwell makes the pairing the default. The GB200 module fuses one Grace with two B200s over NVLink-C2C, and NVL72 stitches 36 of them into a single liquid-cooled scale-up domain: 72 GPUs, 36 Grace CPUs, 13.5 TB of HBM and 17 TB of LPDDR5X as one flat, coherent address space. Rubin steps this in two. NVL144 ships in 2026 as a Rubin-generation refresh inside the same 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: NVL576 packs 144 four-die Rubin Ultra packages into the new Kyber chassis for 576 GPU dies in one coherent domain.它们共同定义了HGX 8-GPU基板，通过PCIe Gen5将8个H100 SXM模块与x86主机（AMD EPYC或Intel Xeon）配对。Hopper还发布了Grace配对形式：GH200 Grace Hopper超级芯片通过NVLink-C2C以900 GB/s的速度将一个Grace ARM CPU绑定到一个H100，消除了PCIe主机-设备跳跃。模块扩展为GH200 NVL2对和机架级GH200 NVL32。Blackwell使这种配对成为默认。GB200模块通过NVLink-C2C将一个Grace与两个B200融合，NVL72将36个模块缝合成一个液冷纵向扩展域：72个GPU、36个Grace CPU、13.5 TB HBM和17 TB LPDDR5X作为一个平坦、一致的地址空间。Rubin将其翻倍。NVL144于2026年在同一Oberon级机架内作为Rubin代更新发布：72个Rubin封装，在NVIDIA新的芯片计数惯例下标记为144个GPU，HBM4和NVLink 6使每封装带宽翻倍。真正的机架级飞跃是2027年的Rubin Ultra：NVL576将144个四芯Rubin Ultra封装打包到新的Kyber机架中，在单个一致域内实现576个GPU芯。

That density is held together by passive copper. NVL72's NVLink fabric runs over 5,184 cables blind-mated through a backplane (~2 miles of cabling per rack, no 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.这种密度通过无源铜缆连接。NVL72的NVLink结构通过背板盲插的5,184根电缆运行（每个机架约2英里长的布线，电缆内无重定时器，SerDes本身位于GPU和交换机ASIC上），在72个GPU之间承载约130 TB/s的全对全带宽。NVIDIA估计，与需要每个链路都有可插拔收发器的光学等效方案相比，选择铜缆每个机架可节省约20 kW功率。铜缆使得“机架即GPU”在经济上变得可行：在2米以内的短距离传输中，它在功耗、成本和每美元信号完整性方面仍占优势；超过这个距离，信号就必须通过光纤传输。

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.NVL144保持在Oberon机架内，且铜缆继续有效，因为封装数量（72）与NVL72相同；布线不需要延长，只需在Gen 6 SerDes上以更快的速度传输。Rubin Ultra的NVL576通过重塑机架保持了同样的铜缆方案：新的Kyber外形尺寸约为Oberon的两倍高，并将所有576个GPU芯封装在一个外壳中，其尺寸专门设计，使每个NVLink路径即使在144个四芯封装和数万根电缆的情况下，也保持在无源铜缆的覆盖范围内。

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 RDMA initiated by software, typically wrapped in NCCL collectives like all-reduce or all-to-all. The reference cluster is the DGX SuperPOD: eight NVL72 racks stitched together over 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.横向扩展栈来自他们对Mellanox的收购。与NVLink不同，横向扩展结构不是一致性的：节点保持独立的地址空间，数据仅通过软件发起的显式RDMA跨越，通常包装在NCCL集合通信（如all-reduce或all-to-all）中。参考集群是DGX SuperPOD：通过Quantum-X800 InfiniBand缝合在一起的八个NVL72机架，在单个调度器下产生576个Blackwell GPU，训练集群通过平铺SuperPOD进一步扩展。2026年的Rubin SuperPOD保持了相同的8机架模式，使用NVL144（每个SuperPOD产生1,152个GPU，而不是576个）。2027年的Rubin Ultra将此配方扩大了一个数量级：每个机架包含576个GPU芯的Kyber机架，通过Quantum-X Photonics CPO缝合在一起，将数千个GPU置于单个调度器之下。

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, Spectrum-X is a lossless-Ethernet alternative tuned for AI traffic.每个GPU都有自己的ConnectX NIC接入该结构。Blackwell节点运行ConnectX-8，每个GPU 800 Gbps，带宽比每GPU NVLink低一个数量级，延迟从纳秒级攀升至微秒级。Rubin转向每个GPU 1.6 Tbps的ConnectX-9，随着每机架纵向扩展域从72个GPU增长到576个，每GPU的横向扩展带宽翻倍。每个NIC旁边都有一个BlueField DPU，增加ARM核心和加速器，以从主机CPU卸载存储、网络和安全任务。对于喜欢以太网而非InfiniBand的客户，Spectrum-X是一种针对AI流量优化的无损以太网替代方案。

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.铜缆到光纤的交叉点发生在机架边界。在NVL72内部，主干是铜缆；一旦链路需要在800 Gbps下跨越机架，它就是光学的。无源铜缆DAC在200 G/通道下大约在1.5–2米达到极限，远未达到跨机架的距离，因此今天的SuperPOD主干运行在OSFP-RHS可插拔收发器上，每个模块都带有自己的激光器、调制器、光电探测器和DSP。从光学角度看，扇出到数千个GPU的SuperPOD主干是数万个可插拔模块，仅收发器激光器就消耗数万千瓦功率。

With Rubin, that optical layer collapses into the switch ASIC. Quantum-X Photonics (InfiniBand) and Spectrum-X Photonics (Ethernet) replace the pluggables with co-packaged optics: 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, and photonics on one substrate.随着Rubin的发布，该光学层被压缩进交换机ASIC中。Quantum-X Photonics（InfiniBand）和Spectrum-X Photonics（以太网）通过共封装光学器件取代了可插拔模块：激光器、调制器和光电探测器通过TSMC COUPE键合到交换机封装上。NVIDIA声称，与OSFP可插拔等效方案相比，激光器数量减少了约4倍，链路功耗降低了约3.5倍。将GPU变成双芯封装并将其与HBM堆叠在一起的小芯片逻辑，现在出现在网络层：计算、内存和光子学在单个基板上的垂直集成。

NVLink Fusion recently opened the scale-up fabric itself: third-party CPUs and XPUs can now join NVLink domains, letting hyperscalers build semi-custom racks around NVIDIA's interconnect without designing their own coherent fabric from scratch.NVLink Fusion最近开放了纵向扩展结构本身：第三方CPU和XPU现在可以加入NVLink域，让超大规模数据中心运营商无需从头设计自己的相干结构，即可围绕NVIDIA的互连构建半定制机架。

Software软件

CUDA is the natural programming model for a massively parallel 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.CUDA是大规模并行处理器的自然编程模型。你编写一个内核（一段代码，每个线程执行一次），并将其启动到组织成块和线程束的数千个线程中；程序员决定它们共享什么、何时同步以及每个线程处理问题的哪一部分。这就是为什么该抽象在十八年里几乎没有改变，也是为什么自2007年以来编写的每个CUDA内核在Blackwell上仍然可以编译和运行的原因。

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 PTX and SASS: 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.这种连续性既是护城河也是约束。每一代新产品都在相同的内核和线程束模型上引入新硬件（Tensor Core、TMA、TMEM），以PTX和SASS中的内在函数暴露：mma.sync, wgmma.mma_async等。NVIDIA无法彻底重新思考SM，因为太多代码依赖于它；作为回报，在CUDA软件上的每一项投资都在几代产品中产生复利。

On top of PTX sits a stack constructed over two decades. cuBLAS and cuDNN for math and DNN primitives; CUTLASS, encoding decades of GEMM expertise in templated C++; TensorRT-LLM for paged attention, in-flight batching, and speculative decoding; framework bindings through PyTorch, Triton, and JAX.在PTX之上是一个构建了二十年的栈。用于数学和DNN原语的cuBLAS和cuDNN；编码了数十年GEMM专业知识的模板化C++库CUTLASS；用于分页注意力、动态批处理和投机解码的TensorRT-LLM；通过PyTorch、Triton和JAX提供的框架绑定。

FlashAttention, one of the most important algorithmic rewrites in modern AI, tiles attention to avoid materialising the $O(N^2)$ 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是现代AI中最重要算法重写之一，它对注意力机制进行分块以避免实现O(N^2)矩阵。它的四代产品（FA1到FA4）每一代都针对最新的NVIDIA芯片进行了手工优化（FA3针对Hopper的异步流水线，FA4针对Blackwell），而对其他硬件的移植则滞后数月或数年。

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不支付工资的人编写的。护城河不是CUDA本身；而是二十年的第三方内核、库和工具，以及一路上学习该API的数百万开发者。

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.NVIDIA还在芯片之外提供人力资源。他们将数十名自己的工程师嵌入到前沿实验室和超大规模数据中心团队中，为每种新模型架构编写内核，并针对每一代新芯片进行调优。实验室下个月想要训练的任何模型，往往都能在NVIDIA平台上比其他平台运行得更好。因此，放弃NVIDIA不仅仅是重写内核和库的问题。这是在重新训练整个工程团队的思维模式，并失去今天坐在办公楼里的NVIDIA工程师。

Google TPUGoogle TPU

Architecture

A TPU chip is a matmul engine wrapped in just enough silicon to keep it fed. The unit of compute is the TensorCore: flagship chips from 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 MXUs for matrix math, a VPU for element-wise math, a Scalar Unit that runs the show, an XLU for cross-lane reductions, and an attached Transpose/Permute Unit, plus accumulator queues feeding and draining the MXU. From v4 onward each chip also carries dedicated SparseCore dataflow engines outside the TensorCore (4 per chip on 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 Core Sequencer that fills all eight functional slots of a 322-bit bundle every cycle. There is no instruction cache miss, no 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.TPU 芯片本质上是一个矩阵乘法引擎，外围封装了维持其运转所需的最小硅片面积。计算的基本单元是 TensorCore：从 v2 开始，旗舰芯片每个封装内包含两个 TensorCore；效率优化型芯片（v4i, v5e, v6e）则包含一个。每个 TensorCore 内部都遵循相同的五组件设计：一个或多个用于矩阵运算的 MXU、一个用于逐元素运算的 VPU、一个负责统筹的标量单元（Scalar Unit）、一个用于跨通道归约（cross-lane reductions）的 XLU，以及一个附加的转置/排列单元（Transpose/Permute Unit），再加上为 MXU 提供数据输入和输出的累加器队列。从 v4 开始，每个芯片在 TensorCore 之外还配备了专用的 SparseCore 数据流引擎（v4, v5p 和 Ironwood 上每芯片 4 个；Trillium 上每芯片 2 个），专门用于处理收缩阵列（systolic array）结构不匹配的嵌入查找（embedding-lookup）工作负载。每个模块都位于由核心定序器（Core Sequencer）驱动的单个 VLIW 发射平面上，该定序器每个周期都会填充 322 位指令包中的所有八个功能槽位。这里没有指令缓存缺失、没有乱序调度器、没有乱序执行引擎，也没有分支预测器：编译器即是调度器，节省下来的硅片面积全部用于增加 MAC（乘加运算单元）。

TensorCore

The MXU is the systolic array. 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.MXU 即是收缩阵列。v1 发布时搭载了一个 256×256 的 INT8 推理阵列；v2 是首款具备训练能力的 TPU，引入了 128×128 的单元，执行 BF16 乘法并进行 FP32 累加（从 v4 开始，INT8 回归 MXU，吞吐量相当）。每个 TensorCore 的单元数量随之增长：v2 为 1 个 MXU → v3 为 2 个 → v4/v5e/v5p 为 4 个。Trillium 回归到 256×256（每个阵列每周期 65,536 个乘加单元），而 Ironwood、8t 和 8i 都保持了 256×256 的规格。

To compute $C = A \times B$, matrix B's values are pre-loaded one weight per cell: weight-stationary dataflow, the choice that distinguishes TPUs from 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 underfill: a 128×128 matmul on a 256×256 array wastes 75% of the silicon, so 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.为了计算 C=A×B，矩阵 B 的值被预先加载，每个单元存储一个权重：这种“权重静止”（weight-stationary）数据流是 TPU 区别于其他输出静止阵列的关键选择。激活值从左侧进入，每周期传播一列，与每个单元中驻留的权重相乘，部分和向下流入底部的累加器队列。一旦数据进入阵列，便无需内存访问：每个权重被重复用于通过它的每一个激活值，每个激活值在行内被重复使用 128（或 256）次。数据重用直接硬连线在硅片中，而非由缓存仲裁。计算中的主要成本并非乘法本身（仅需几个皮焦耳），而是内存读写，其能耗比乘法高出 100–1000 倍；收缩阵列通过其结构设计消除了这一成本。其代价是“填充不足”（underfill）：在 256×256 的阵列上执行 128×128 的矩阵乘法会浪费 75% 的硅片，因此 XLA 会将维度平铺、填充并调度为 128（v6e+ 上为 256）的倍数，模型代码编写时也会考虑到这些量子化单位。

The VPU 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 VREGs. On v4/v5p the shape is (8, 128): 128 lanes wide, 8 sublanes 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 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 VPU/MXU overlap: 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 XLU: slow, expensive, and a known compiler hot spot. Layout transforms that misalign with the 2D shape are absorbed by the dedicated Transpose/Permute Unit, sparing a round-trip through memory.VPU 是辅助计算引擎，但在微架构层面却更为有趣：每个 TPU 都是 2D 向量机，而非 1D SIMD 机。VPU 的寄存器文件包含 2D 向量寄存器（VREG）。在 v4/v5p 上，其形状为 (8, 128)：128 通道宽，8 子通道深，每个核心 32（v4）或 64（v5p）个寄存器，每个（通道，子通道）配有 4 个独立的浮点 ALU。通道轴与收缩阵列的输入宽度匹配，因此在 Trillium 和 Ironwood 上，通道数推测随 MXU 一起扩展到了 256；Google 尚未公布 v5p 之后的 VPU 维度。子通道轴允许 VPU 以每 X 个时钟周期一个矩阵乘法的速度通过 MXU 流式传输瓦片（tile）（其中 X 为子通道维度）。现代 TPU 程序的大部分加速来自 VPU/MXU 的重叠执行：量化、层归一化（layernorm）、Softmax、激活和偏置加法（bias-add）都在 VPU 上运行，而 MXU 则在后台执行矩阵乘法。跨通道归约（对于任何 2D 向量 ISA 来说都是棘手的情况）由 XLU 处理：速度慢、成本高，且是已知的编译器热点。与 2D 形状不对齐的布局转换由专用的转置/排列单元处理，避免了往返内存的开销。

The Scalar Unit 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 SMEM 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 DMA), 2 matrix (push/pop the MXU queue). Synchronisation between blocks is explicit: sync flags 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 reorder buffer to undo a bad decision.标量单元是最小但可以说是最重要的模块：一个单线程、双发射的整数 ALU，配有 32 个 32 位寄存器和 4 KiB 用于控制状态的 SMEM，并与存放程序的 Imem 配对。它是唯一执行指令获取的模块；每个周期它提取一个 322 位的 VLIW 指令包，在本地执行其自身的两个标量槽（地址计算、循环计数器、分支、同步寄存器检查），并将剩余的六个槽位分发给芯片的其他部分：2 个向量 ALU（VPU）、2 个向量加载/存储（HBM↔VMEM DMA）、2 个矩阵（推入/弹出 MXU 队列）。模块间的同步是显式的：同步标志跟踪 MXU 和 VPU 流水线何时繁忙，编译器插入屏障检查，而不是由硬件跟踪依赖关系。标量单元使 TensorCore 的其余部分看起来像固定功能的数据流：每个周期，一个地方决定发生哪八件事，没有动态重排序缓冲区来撤销错误的决策。

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 HBM (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 VMEM, 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 v8i precisely to hold an entire KV cache on chip. Above it sits CMEM, introduced with 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 SMEM (~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 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.片上内存层次结构与计算侧的设计理念相同：没有缓存，每一层都由软件管理。片外是 HBM（v2/v5e 为 16 GB，v3/v4/v6e 为 32 GB，v5p 为 95 GB，Ironwood 为 192 GB，v8 代为 216–288 GB），片上则是分层的手动堆叠式、可显式寻址的暂存存储器（scratchpad）。最靠近计算的是 VMEM，即为 VPU 和 MXU 输入队列提供数据的向量暂存存储器，v4 上大小为 32 MiB，v5e 上为 128 MiB，而在专为推理优化的 v8i 上扩展到了 384 MiB，精确地为了将整个 KV 缓存保持在片上。其上是 CMEM，于 v4 引入，大小为 128 MiB：这是 HBM 和 VMEM 之间一个更大、更慢的 SRAM 暂存区，用于吸收融合算子的中间结果。标量单元有自己的 SMEM（v4 上约 10 MiB 用于控制状态）和一个微小的标量寄存器文件。程序中的每个张量在编译时都被固定到某一层；XLA 的缓冲区分配通道跨层调度 DMA，以便数据在消耗它的周期之前恰好到达。硬件不做预取、不做驱逐、不做一致性；当编译器正确调度时，阵列永不停止；当调度错误时，没有回退路径。

SparseCoreSparseCore

The block outside the TensorCore that breaks the systolic mould is SparseCore, introduced with 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 dataflow processor with 16 compute tiles and dedicated SPMEM scratchpads, sitting alongside the TensorCore and absorbing 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 CAE (Collectives Acceleration Engine) on the I/O chiplet: different problem (collective 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).TensorCore 之外打破收缩阵列模式的模块是 SparseCore，于 v4 引入。推荐和排序模型依赖于嵌入查找（对庞大表进行数十亿次索引），其访问模式与密集矩阵乘法相反：不规则、间接、全对全（all-to-all）。256×256 的收缩阵列形状完全不匹配。SparseCore 是一个带有 16 个计算瓦片和专用 SPMEM 暂存存储器的数据流处理器，与 TensorCore 并列，吸收了散布（scatter）、聚集（gather）和分段归约（segmented-reduce）原语，以及分片嵌入表产生的与数据相关的全对全流量。这在嵌入密集型模型上实现了 5–7 倍的加速，仅消耗约 5% 的芯片面积和功耗。v4 每芯片搭载 4 个 SparseCore，v5p 保持了该数量，Trillium 降至 2 个，Ironwood 回归到 4 个（双芯片布局中每个小芯片 2 个）。v8i（Zebrafish）推理芯片完全移除了 SparseCore，并在 I/O 小芯片上代之以 CAE（集体加速引擎）：问题不同（自回归解码过程中的集体归约），理念相同（从主核心中剥离一个小加速器以吸收收缩阵列形状不匹配的工作负载）。

Numerics

TPU v1 was INT8-only inference; v2 switched this for BF16 as the canonical training format: same dynamic range as FP32, half the memory, no loss-scaling tricks. v4 reintroduced native INT8 support. Ironwood then added native FP8 support (both E4M3 and E5M2) for ~2× the throughput of BF16 in the same area. v8 adds native FP4 plus block-scale multiplication inside the MXU itself, which deletes the VPU dequant overhead that Ironwood still paid. Stochastic rounding 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.TPU v1 仅支持 INT8 推理；v2 将其切换为 BF16 作为标准训练格式：与 FP32 具有相同的动态范围，内存占用减半，且无需损失缩放技巧。v4 重新引入了原生 INT8 支持。Ironwood 随后增加了原生 FP8 支持（E4M3 和 E5M2），在相同面积下吞吐量约为 BF16 的 2 倍。v8 增加了原生 FP4 以及 MXU 内部的块级缩放乘法，这消除了 Ironwood 仍需支付的 VPU 反量化开销。随机舍入（stochastic rounding）在所有现代 TensorCore 上均有硬件支持：舍入决策由作为概率的低位尾数位决定，这在长训练运行中保持了低精度累加的期望值，也是让 BF16/FP8 缩小与 FP32 精度差距的细节之一。

At the chip boundary sit the ICI ports themselves (4 ports on the 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.芯片边界处是 ICI 端口本身（2D 环面芯片 v2/v3/v5e/v6e 上有 4 个端口，3D 环面旗舰芯片 v4/v5p/v7/8t 上有 6 个），以及用于横向扩展的 DCN NIC。从芯片级视角看，ICI 端口看起来就像 Core Sequencer 可以在 VLIW 指令包中定位的另一组 DMA 引擎：远程张量发送与 VMEM 到 HBM 的传输属于同一指令类，编译器将集体操作视为其为计算和本地内存构建的整体调度的一部分。

Bets

• Bet 1: Systolic array. Matmul dominates the workload, so spend the silicon on a systolic array.赌注 1：收缩阵列。矩阵乘法占据了工作负载的主导地位，因此将硅片资源投入到收缩阵列中。
• Bet 2: Software scratchpads. Compute is cheap and memory is expensive, so reuse data in the wires of the array and replace caches with software-managed scratchpads.赌注 2：软件暂存存储器。计算廉价而内存昂贵，因此在阵列的连线中重用数据，并用软件管理的暂存存储器取代缓存。
• Bet 3: Compiler scheduling. The workload is statically predictable, so move scheduling into the compiler: VLIW issue, no speculation, no out-of-order, no dynamic scheduler.赌注 3：编译器调度。工作负载是静态可预测的，因此将调度转移到编译器：VLIW 发射，无推测，无乱序，无动态调度器。
• Bet 4: MAC-only silicon. Power matters more than peak, so delete every transistor that does not multiply-add: every cache tag, every branch predictor, every reorder buffer.赌注 4：仅含 MAC 的硅片。功耗比峰值性能更重要，因此删除每一个不执行乘加运算的晶体管：每一个缓存标签、每一个分支预测器、每一个重排序缓冲区。
• Bet 5: Dedicated off-array engines. The dense matmul array is the wrong shape for some real workloads (embeddings, collectives), so carve out small dedicated engines (SparseCore, CAE) rather than warp the main core to fit them.赌注 5：专用的阵列外引擎。密集矩阵乘法阵列对于某些实际工作负载（嵌入、集体操作）来说形状不对，因此剥离出小型专用引擎（SparseCore, CAE），而不是为了适配它们而扭曲主核心。

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 XLA. The scale-up domain is tied together not by a switch fabric but by a torus (chips wired directly to their neighbours with edge wrap) and stitched at the rack boundary by optical circuit switches.TPU 的扩展故事与 NVIDIA 正好相反。NVIDIA 通过 NVLink + NVSwitch 使每个 GPU 的 HBM 看起来像本地内存（硬件管理的统一地址空间），而 Google 的 ICI 是消息传递机制。没有远程加载语义，没有缓存一致性，没有交叉开关。每个多芯片操作都是由 XLA 编译的显式集体操作。扩展域不是通过交换结构连接的，而是通过环面（芯片直接与邻居相连，边缘环绕）连接，并在机架边界通过光路交换机缝合在一起。

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 Ironwood, and 2× that on v8t. Topology alternates by generation: 2D torus on the efficiency-tuned chips (v2, v3, v5e, v6e), 3D torus on the flagships (v4, v5p, v7, v8t).ICI 链路直接从 TPU 裸片引出：高速串行通道，64 芯片立方体（一个液冷机架中的 4×4×4 排列）内使用直接连接铜缆，立方体间使用光缆。单芯片聚合 ICI 带宽已从 v2 的约 250 GB/s 扩展到 Ironwood 的 1.2 TB/s 双向带宽，v8t 上则翻倍。拓扑结构随代际变化：效率优化型芯片（v2, v3, v5e, v6e）采用 2D 环面，旗舰芯片（v4, v5p, v7, v8t）采用 3D 环面。

The piece with no NVIDIA analogue is the Palomar OCS: a 3D-MEMS optical circuit switch that sits between cubes. Tiny mirrors physically rotate to map any input fibre to any output. A 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 circuit-switched: 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 (twisted tori give up to 70% better bisection), sub-pod slicing on demand, and fault tolerance (when a chip dies, the OCS optically swaps in a spare cube and the run continues without losing the ICI domain).没有 NVIDIA 对应物的部分是 Palomar OCS：位于立方体之间的 3D-MEMS 光路交换机。微小的镜面物理旋转以将任何输入光纤映射到任何输出光纤。v4 超级集群使用 48 个 Palomar 交换机将 64 个立方体（4,096 个芯片）连接成一个 3D 环面；v5p 和 Ironwood 采用了相同的扩展方案。重配置是毫秒级的，而非纳秒级，但这没关系，因为 OCS 是电路交换的：在作业开始时选择拓扑，运行一周，然后为下一个工作负载重新配置。三个问题合并为一个组件：每个工作负载的拓扑重配置（扭曲环面可提供高达 70% 的更好对分带宽）、按需子集群切片，以及容错（当芯片损坏时，OCS 从光学上换入备用立方体，运行无需丢失 ICI 域即可继续）。

This makes the superpod the unit of scale-up: equivalent in role to NVIDIA's NVL72, two orders of magnitude bigger. v4 was 4,096 chips; v5p, 8,960; Ironwood (TPU v7) is 9,216 chips arranged as 144 cubes of 64, presenting 1.77 PB of HBM (~68 PB/s) and 42.5 ExaFLOPS FP8 as one coherent ICI domain.这使得超级集群成为扩展单元：在角色上等同于 NVIDIA 的 NVL72，但规模大两个数量级。v4 为 4,096 个芯片；v5p 为 8,960 个；Ironwood（TPU v7）为 9,216 个芯片，排列为 144 个 64 芯片立方体，作为一个统一的 ICI 域提供 1.77 PB 的 HBM（约 68 PB/s）和 42.5 ExaFLOPS 的 FP8 算力。

TPU 8t (Sunfish) stretches this to 9,600 chips, 2 PB of HBM (~62 PB/s), and 121 ExaFLOPS FP4. TPU 8i (Zebrafish) has 1,024 chips, ~295 TB of HBM (8.8 PB/s), and ~10 ExaFLOPS FP4. 8i replaces torus with a new hierarchical high-radix topology called Boardfly (4-chip ring → 8-board group → up to 36 groups linked by 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.TPU 8t（Sunfish）将其扩展到 9,600 个芯片、2 PB 的 HBM（约 62 PB/s）和 121 ExaFLOPS 的 FP4 算力。TPU 8i（Zebrafish）拥有 1,024 个芯片、约 295 TB 的 HBM（8.8 PB/s）和约 10 ExaFLOPS 的 FP4 算力。8i 用一种称为 Boardfly 的新型分层高基数拓扑取代了环面（4 芯片环 → 8 板组 → 多达 36 个由 OCS 链接的组），将全对全延迟减半。这是为 MoE 推理设计的。当集体操作是最近邻时，3D 环面表现出色（环形全归约每周期使用每条链路），但 MoE 专家路由是相反的模式，即全对全：每个芯片向其他所有芯片发送唯一片段，往返延迟受最长跳对限制。1,024 芯片的 3D 环面直径为 16 跳；Boardfly 的环 → 组 → OCS 分层结构将其压缩为 7。

Scale-out

Through TPU v7, scale-out ran over a single fabric: Jupiter, all-optical at the spine since 2022 via Apollo OCS, the same 3D-MEMS family as 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.截至 TPU v7，横向扩展在单一结构上运行：Jupiter，自 2022 年起通过 Apollo OCS（与 Palomar 相同的 3D-MEMS 系列，在整个建筑内扩展）在主干网实现全光化。Google 在从机架到数据中心主干网的每一层都使用相同的原语（光路交换）；这是其他人没有的架构特征。Jupiter 目前每栋建筑承载 13 Pb/s 的对分带宽。

With TPU 8t, scale-out split into two fabrics. East-west TPU-to-TPU traffic moved to Virgo, a dedicated accelerator fabric; Jupiter retained the north-south role: storage access, general compute, and inter-site scaling. Virgo is a flat, two-layer, non-blocking topology built on high-radix switches: every TPU is at most two switches from any other. One Virgo cluster links 134,000+ TPU 8ts 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.随着 TPU 8t 的推出，横向扩展分为两个结构。东西向 TPU 到 TPU 的流量转移到 Virgo（专用加速器结构）；Jupiter 保留了南北向角色：存储访问、通用计算和跨站点扩展。Virgo 是一种基于高基数交换机的扁平、双层、无阻塞拓扑：每个 TPU 到任何其他 TPU 最多经过两个交换机。一个 Virgo 集群以 47 Pb/s 的对分带宽连接 134,000+ 个 TPU 8t（单芯片带宽是前一代 DCN 的 4 倍，空载延迟降低 40%），具有多平面故障隔离和亚毫秒级遥测，允许调度器在落后者破坏步骤之前将其杀死。架构上的回报是每一层现在都可以独立演进：扩展、东西向横向扩展和前端可以在不同的节奏下迭代，而无需重新布线其他层。

Per-chip scale-out bandwidth is on the order of 100 Gbps on Ironwood, and 4× that on v8t, but still two orders of magnitude less than per-chip 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.单芯片横向扩展带宽在 Ironwood 上约为 100 Gbps，在 v8t 上是该值的 4 倍，但仍比单芯片 ICI 低两个数量级。这种带宽差距决定了分区：张量并行和 MoE 专家路由保留在 ICI 内部；数据并行和流水线并行跨越横向扩展结构。

Google's Multislice framework, plumbed into XLA, lets a single 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.Google 的 Multislice 框架植入 XLA，允许单个 SPMD 程序跨越不同集群中的多个切片；编译器发出分层集体操作（切片内环形全归约，跨切片更高级别的归约）。这种结构正是隐藏 ICI/DCN 带宽差距的诀窍：尽可能多的工作保留在切片内部通过快速 ICI 完成，仅将跨切片残余部分支付慢速结构成本。

Above this sits Pathways. Where NCCL + Slurm + Megatron-style schedulers drive SPMD from many controllers, Pathways drives the entire job from one client and 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. Gemini Ultra was the first frontier model trained across multiple datacenters; Pathways stitches them into one synchronous SPMD job.在此之上是 Pathways。当 NCCL + Slurm + Megatron 风格的调度器从多个控制器驱动 SPMD 时，Pathways 从一个客户端驱动整个作业，并虚拟化连接在 DCN 上的多个“岛屿”（具有各自 ICI 域的集群）。它执行组调度、弹性训练（当切片失败时，OCS 重塑拓扑，Pathways 从新形状上的最后一个检查点恢复）和跨区域编排。Gemini Ultra 是第一个跨多个数据中心训练的前沿模型；Pathways 将它们缝合为一个同步的 SPMD 作业。

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 kernel-driven. On a GPU, the developer writes the kernel and the framework strings kernels together; the compiler's job is mostly local. On a TPU, the developer writes a numerical program in JAX and XLA is responsible for everything below it: which operations 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 harder.TPU 堆栈是编译器驱动的，而 CUDA 是内核驱动的。在 GPU 上，开发者编写内核，框架将内核串联起来；编译器的主要工作是局部的。在 TPU 上，开发者在 JAX 中编写数值程序，XLA 负责其下的一切：哪些操作融合，每个张量驻留在哪里，它如何跨 2D 向量寄存器布局，何时发出从 HBM 到 VMEM 的 DMA，如何调度 322 位 VLIW 指令包，程序如何跨数千个芯片分片。没有硬件回退：没有乱序调度器、没有缓存、没有乱序执行引擎来掩盖糟糕的调度。编译器即系统。权衡是架构的核心：XLA 在无需手动调整的情况下更接近理论上限，但缩小剩余差距更难。

The compilation path is JAX → JAXpr → StableHLO → HLO → LLO → VLIW bundles. JAX traces a Python function into a typed functional IR (JAXpr) under jit, lowers it to StableHLO (the OpenXLA-standardised, versioned op-set of ~100 statically-shaped primitives that all front-ends now emit), which XLA ingests as HLO and runs through its pass pipeline: operation fusion (collapse pointwise + reduction + matmul into one kernel so intermediates never hit HBM), layout assignment (decide the 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), buffer assignment (every tensor pinned to either VMEM, CMEM, or HBM with overlap windows pre-computed), SPMD partitioning, and finally a VLIW scheduler that fills all eight slots of every bundle. HLO lowers to LLO (Low-Level Optimizer), the TPU-specific 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.编译路径为 JAX → JAXpr → StableHLO → HLO → LLO → VLIW 指令包。JAX 在 jit 下将 Python 函数跟踪为类型化的函数式 IR（JAXpr），将其降低为 StableHLO（OpenXLA 标准化的、版本化的算子集，包含约 100 个所有前端现在都发出的静态形状原语），XLA 将其作为 HLO 摄入并运行通过其通道流水线：操作融合（将点对点 + 归约 + 矩阵乘法折叠为一个内核，使中间结果永远不会触及 HBM）、布局分配（决定每个张量的 2D 平铺，以便它无需转置即可流式传输到 MXU：这比在 1D SIMD 机器上难得多，因为寄存器和收缩输入都是 2D 的）、缓冲区分配（每个张量在编译时固定到 VMEM、CMEM 或 HBM，并预先计算重叠窗口）、SPMD 分区，最后是 VLIW 调度器，填充每个指令包的所有八个槽位。HLO 降低到 LLO（低级优化器，TPU 特定的 IR），LLO 发出最终的 VLIW 流。一个编译良好的程序在每个周期都在同一个指令包中重叠 MXU 收缩执行、VPU 逐元素数学和 HBM↔VMEM DMA。

Multi-chip execution is SPMD: one program, sharded data, hierarchical collectives, emitted by GSPMD (now being replaced by Shardy, an MLIR-native successor that lands as the default in early 2026). The user expresses sharding declaratively with Mesh + PartitionSpec annotations on a few key tensors; the 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, shard_map drops the user into manual SPMD (per-device code with explicit local shapes and explicit 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: FSDP and DeepSpeed wrap the model in a runtime that issues collectives at module boundaries; GSPMD/Shardy partitions the whole graph as a compiler problem.多芯片执行是 SPMD：一个程序，分片数据，分层集体操作，由 GSPMD 发出（现在正被 Shardy 取代，这是一个 MLIR 原生的后继者，将于 2026 年初成为默认设置）。用户通过几个关键张量上的 Mesh + PartitionSpec 注释声明式地表达分片；编译器通过图的其余部分传播分片，并在布局变化处插入全归约、全聚集和归约散布。当编译器选择错误的集体操作时，shard_map 会将用户降级为手动 SPMD（带有显式本地形状和显式集体操作的每设备代码），可在 jit 中组合，因此单个内核可以手动分区，而无需放弃其他地方的自动分区。这与 PyTorch 的惯用法相反：FSDP 和 DeepSpeed 将模型包装在一个在模块边界发出集体操作的运行时中；GSPMD/Shardy 将整个图的分区作为一个编译器问题。

Pallas is the escape hatch: JAX's kernel-writing language, broadly the TPU equivalent of Triton on GPUs. Pallas kernels are written in JAX-flavoured Python, lowered through Mosaic (the 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 FlashAttention-class optimisation, where the win is in the schedule and not the algebra. Pallas:Mosaic-GPU 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: Flax NNX for modules, Optax for optimisers, Orbax for asynchronous distributed checkpointing, Grain for input pipelines, Tunix for post-training/RL, Qwix for quantisation. Google's reference training stacks (MaxText for LLMs including DeepSeek-V3-class MoE, and MaxDiffusion for Flux, Wan 2.1) sit at the top, in pure JAX; Pathways sits beneath, exposed to the user as pathwaysutils, so a single Python client can drive a job across thousands of chips and several pod-islands without giving up the JAX programming model.Pallas 是逃生舱：JAX 的内核编写语言，大致相当于 GPU 上的 Triton。Pallas 内核以 JAX 风格的 Python 编写，通过 Mosaic（基于 MLIR 的 TPU 后端）降低到 LLO，并作为自定义算子嵌入回 HLO。它的存在是因为 XLA 不能总是为新颖的注意力变体、融合的 MoE 分发或任何需要手动 VMEM 平铺和 DMA 调度的工作负载合成最优解：FlashAttention 类的优化，其优势在于调度而非代数。Pallas:Mosaic-GPU 以相同的前端针对 H100/Blackwell，因此内核作者可以编写一次并降低到任一基板。其上的库层统一为 JAX 原生：Flax NNX 用于模块，Optax 用于优化器，Orbax 用于异步分布式检查点，Grain 用于输入流水线，Tunix 用于训练后/RL，Qwix 用于量化。Google 的参考训练堆栈（MaxText 用于 LLM，包括 DeepSeek-V3 级 MoE；MaxDiffusion 用于 Flux, Wan 2.1）位于顶部，使用纯 JAX；Pathways 位于底层，作为 pathwaysutils 向用户公开，因此单个 Python 客户端可以跨数千个芯片和多个集群岛驱动作业，而无需放弃 JAX 编程模型。

The PyTorch path is real but second-class. torch_xla uses a LazyTensor mechanism: every PyTorch op records into an HLO 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, torch.compile integration through an XLA backend, a JAX bridge, 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 StableHLO, and complex parallelism strategies are better-covered), which is why vLLM TPU (powered by the tpu-inference plugin announced at Cloud Next 2025) lowers every model, JAX-defined or PyTorch-defined, through a unified JAX→XLA path. TorchTPU, announced April 2026, is Google's response: a native PyTorch experience with eager mode, torch.distributed, and torch.compile over XLA, on track to replace torch_xla.PyTorch 路径是真实的但属于二等公民。torch_xla 使用 LazyTensor 机制：每个 PyTorch 算子记录到一个在下一个屏障处编译的 HLO 图中，编译后的工件由图形状哈希缓存。PyTorch/XLA 2.x 增加了 GSPMD 风格的分片注释、通过 XLA 后端的 torch.compile 集成、JAX 桥接，以及（PyTorch/XLA 2.7）具有明显更快跟踪速度的 C++11-ABI 构建。与 JAX 的差距是真实的（JAX 的原语能更清晰地映射到 StableHLO，且复杂的并行策略覆盖得更好），这就是为什么 vLLM TPU（由 Cloud Next 2025 上宣布的 tpu-inference 插件驱动）通过统一的 JAX→XLA 路径降低每个模型（无论是 JAX 定义还是 PyTorch 定义）。2026 年 4 月宣布的 TorchTPU 是 Google 的回应：一种通过 XLA 具有 eager 模式、torch.distributed 和 torch.compile 的原生 PyTorch 体验，有望取代 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 Ironwood (v7) "codesigned AI stack" 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. Triton and torch.compile narrow the gap on the NVIDIA side (kernel-driven and compiler-driven are converging), but the philosophical poles are still real: on TPU the compiler is the only interface that matters; on GPU the compiler is one of several.与 CUDA 相比，TPU 生态系统是集中的，而非蔓延的。框架之下的几乎所有内容（XLA, JAX, Flax, Optax, Pallas, MaxText, Pathways, Shardy, Mosaic）都是由 Google 自己开源的，并与硅片同步演进。第三方内核比 CUDA 数十年的积累要少得多；当工作负载看起来很奇怪时，护城河较薄；当工作负载看起来像 Gemini 时，护城河更深。最近的 Ironwood (v7) “协同设计的 AI 堆栈”语言是显式的框架：芯片、ICI 结构、OCS、XLA、Pathways、Pallas、MaxText、vLLM 和 Pathways 作为一个产品共同发布，v8t/v8i 在单一 tpu-inference 降低路径下继续相同的模式。Triton 和 torch.compile 缩小了 NVIDIA 侧的差距（内核驱动和编译器驱动正在融合），但哲学上的极点仍然存在：在 TPU 上，编译器是唯一重要的接口；在 GPU 上，编译器是多个接口之一。

AMD GPU

Architecture

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.NVIDIA 的架构野心存在于每个 SM 内部（每一代都有新的张量原语、新的异步机制、新的操作数存储），而 AMD 的野心存在于 CU 之间，在于有多少 CU 可以键合到单个一致性封装中。CU 本身是保守的：四个 16 通道 SIMD，一个共享标量单元，64 KB 本地数据共享（LDS），一个 L1 向量缓存，每个 SIMD 的 VGPR 文件以及 CU 共享的 SGPR 池，以及（自 CDNA 1 以来）运行 MFMA 的矩阵核心。自 2012 年 GCN 以来，其形状没有发生实质性变化；扩展的是数量（MI100 上 120 个 CU，MI250X 上 220 个，MI300X 上 304 个，MI355X 上 256 个）以及键合它们的封装。64 个线程的 wavefront 在 4 个周期内跨越 16 个 SIMD 通道流式传输，每个 SIMD 中驻留许多 wavefront，调度器在它们之间切换以隐藏停顿。这里没有什么异国情调；CDNA 有趣之处在于 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.在 CU 内部，SIMD 和矩阵核心并排运行。四个 SIMD 处理所有逐元素操作：激活、归一化、残差、地址计算。矩阵核心处理矩阵乘法。这种拆分与 NVIDIA 的 CUDA 核心/张量核心拆分相同，但矩阵抽象在非常不同的曲线上演进。

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.NVIDIA 的张量核心爬升了线程层次结构：Volta 上的 32 线程 warp，Hopper 上的 128 线程 warp-group，Blackwell 上单个线程加可选的双 SM 集群。AMD 的矩阵核心保持原位。每一代 MFMA（从 2020 年的 MI100 到 2025 年的 MI355X）都是 wavefront 作用域的：一个 wave64 发出一个矩阵算子（V_MFMA_*），四个 SIMD 协作驱动它，操作数来自 wavefront 的寄存器文件：A 和 B 来自 VGPR，C 和 D 通常来自专用的 AGPR 文件。指令变快了，格式集也变宽了，但发行者和作用域没有变。馈送侧的唯一让步出现在 CDNA 4：来自 LDS 的专用 MFMA 转置加载，将操作数以其想要的布局交给矩阵核心，精神上类似于 NVIDIA 的 TMA，但矩阵算子本身仍然是 wave 发出的。

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.吞吐量数字直接讲述了格式故事。CDNA 1 于 2020 年发布，FP32 / FP16 / BF16 / INT8 分别为 256 / 1024 / 512 / 1024 FLOPs/CU/周期，与 A100 一起提供原生 BF16 支持。CDNA 2 将 FP64 路径加倍至 256 FLOPs/CU/周期的全速率矩阵：AMD 独有，该赌注将 MI250X 推向了 Frontier。CDNA 3 在 FP8 上以 4,096 FLOPs (E4M3 + E5M2) 与 H100 持平，增加了 2:4 结构化稀疏性，并增加了 TF32 等效路径，通过截断尾数以 FP64 矩阵速率运行 FP32 矩阵乘法。CDNA 4 再次加倍至 FP4 16,384 FLOPs 和带有 OCP MX 块缩放的 FP6，并在一个 MFMA 中增加了可混合的 A/B 精度：例如 FP8 × FP4。同一代将每 CU FP64 吞吐量减半，这是 AMD 第一款为了 AI 密度而牺牲 HPC 密度的芯片，而非两者兼顾。

The wavefront-scope decision shows up in two costs.wavefront 作用域决策表现为两个成本。

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.发散。半空的 wave64 浪费 32 个通道，而半空的 warp32 浪费 16 个。对于具有大部分统一控制流的工作负载，这是很小的代价；对于不规则工作负载，它会造成损害。

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.重叠。NVIDIA 的异步、描述符驱动的矩阵乘法将发布与执行解耦：发布线程触发指令并继续；张量核心在后台运行；warp 可以在前一个矩阵乘法仍在飞行时运行 Softmax、应用掩码或预加载下一个瓦片。AMD 的 wavefront 集体 MFMA 没有给 wave 提供等效功能：发出矩阵乘法的同一个 wave 在挂起时不能同时执行有意义的向量工作。重叠在单独的 wavefront 之间是可能的，但必须在软件中通过显式的 wavefront 屏障进行分段，这更脆弱且消耗更多的 wave 槽位和寄存器。

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. MoE dispatch, paged attention, speculative decode sit in the same camp: address-irregular work that wants to run alongside the matmul.这有多重要取决于工作负载。纯密集 GEMM（DGEMM，大规模批处理训练的内部循环）在矩阵乘法期间没有有用的事情可做；两个引擎都饱和；异步购买很少。这正是 AMD 历史上在百亿亿次 HPC 领域领先的工作负载（Frontier 上的 MI250X，El Capitan 上的 MI300A）。Transformer 注意力（FlashAttention-3, FA4）将矩阵乘法与 Softmax、掩码和 KV 缓存读取交织在一起，异步重叠是这些内核的整个结构。AMD 必须手动重新创建流水线，这落后于 NVIDIA 的硬件级支持。MoE 分发、分页注意力、推测解码处于同一阵营：地址不规则的工作负载希望与矩阵乘法一起运行。

NVIDIA's matrix-instruction abstraction has moved further across generations (warp → warp-group → single-thread async + cluster), and AMD hasn't followed.NVIDIA 的矩阵指令抽象在几代中走得更远（warp → warp-group → 单线程异步 + 集群），而 AMD 没有跟进。

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.AMD 的内存层次结构比 NVIDIA 拥有更少的通用层级，且拥有一个 NVIDIA 完全没有的巨大缓存。从 CU 向外：64 KB LDS 暂存存储器（软件管理，32 库，AMD 类似 NVIDIA 的 SMEM），向量 L1（早期 CDNA 上 16 KB，MI300X 之后 32 KB），每 XCD 几 MB 的 L2。然而，L2 在 XCD 之间不是一致的；一致性发生在 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.该层是 Infinity Cache：MI300X 上 256 MB，分布在四个 IOD 上，16 路组相联，测量约为 12 TB/s，是 MI300X 5.3 TB/s HBM3 的两倍多。它起源于 RDNA 游戏 GPU，以补偿窄 GDDR 总线；AMD 在 CDNA 3 上为 AI 重用了该 IP，其中注意力 KV 重用和权重重用非常适合大型 LLC。NVIDIA 转而押注更大的 HBM 带宽（B200 上 8 TB/s，随 Rubin 上的 HBM4 扩展），而 AMD 押注于缓存。

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.片外，HBM 容量激进增长：MI100 / MI210 / MI250X / MI300X / MI325X / MI350X 分别为 32 → 64 → 128 → 192 → 256 → 288 GB，自 2021 年以来每一代都匹配或超过当代的 NVIDIA 旗舰。赌注是推理工作负载越来越受容量限制，而拥有更多内存的芯片会获胜。

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 OCP MX block-scale multiplication: the same numeric format as 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.格式轨迹追踪了 AI 硅片中每个人都共享的精度减半模式：FP32 → FP16 → FP8 → FP4，每一步都以更细粒度的缩放恢复精度。AMD 特定的轴是开放性。CDNA 4 的 FP4 和 FP6 使用 OCP MX 块缩放乘法：与 Blackwell 的 MXFP4 和 TPU v8 的 MXU 相同的数字格式，但由 AMD 帮助创立的开放联盟（AMD, NVIDIA, Intel, Meta, Microsoft, Qualcomm, ARM）指定，而非任何单一供应商。MI355X 中发布的格式与 B200 和 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.CDNA 4 的拐点值得单独列出：每 CU FP64 吞吐量减半。MI300X 同时服务于训练、HPC 和推理；MI355X 首先是 AI 芯片。驱动 Frontier 的全速率 FP64 矩阵赌注并没有被杀死，但它不再承载重担。

Chiplets小芯片（Chiplets）

The packaging is where CDNA stops looking like NVIDIA and starts being something else.封装是 CDNA 停止看起来像 NVIDIA 并开始成为其他东西的地方。

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 1 的 MI100 是单片 7 nm。CDNA 2 的 MI250X 是 AMD 的第一款多芯片 GPU：两个 Aldebaran GCD 在 2.5D EFB 有机基板上并排，由 4 个封装内 Infinity Fabric 链路连接，聚合带宽 400 GB/s，但向软件呈现为两个独立的 GPU。

CDNA 3 is the move that changed everything. Eight XCDs (TSMC N5, ~115 mm² each) are stacked in 3D via TSMC SoIC hybrid bonding (sub-micron-pitch TSVs, no microbumps) onto four I/O dies (TSMC N6) below. The IODs carry the 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 Infinity Fabric AP 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 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.CDNA 3 是改变一切的举措。8 个 XCD（TSMC N5，每个约 115 mm²）通过 TSMC SoIC 混合键合（亚微米间距 TSV，无微凸块）在 3D 中堆叠在下方的 4 个 I/O 裸片（TSMC N6）上。IOD 承载 Infinity Cache、HBM3 PHY、Infinity Fabric 链路和 PCIe Gen 5；每个 IOD 在上方托管 2 个 XCD，并在侧面托管 2 个 HBM 堆栈。4 个 IOD 由 Infinity Fabric AP 以 4.8 TB/s 的对分带宽缝合，因此 1530 亿晶体管的封装对内核来说看起来像一个 GPU：缓存和地址空间在 IOD 层统一。NVIDIA 在 H100 之前保持单片，仅通过 2.5D CoWoS-L 在 B200 上采用了两个标线限制裸片。AMD 早一代实现了 3D 堆叠，且每个裸片面积更小：在相同的封装前沿有不同的赌注。

The MI300A APU pushed the bet further. Replace 2 of the 8 XCDs with three Zen 4 CCDs, 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 Grace-Hopper bridges two packages over NVLink-C2C; MI300A is one. El Capitan (11,039 nodes of 4× MI300A) is the deployment that justified it.MI300A APU 将赌注推得更远。用 3 个 Zen 4 CCD 替换 8 个 XCD 中的 2 个，保持 HBM、Infinity Cache 和 IOD 不变，并让 CPU 和 GPU 共享由 HBM3 支持并具有硬件一致性的单一物理地址空间。路径中没有主机设备复制。没有固定内存。路径中没有 PCIe。Zen 4 核心和 CDNA 3 XCD 从相同的页面读取。NVIDIA 的 Grace-Hopper 通过 NVLink-C2C 桥接两个封装；MI300A 是一个。El Capitan（11,039 个 4× MI300A 节点）是证明其合理性的部署。

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.在 CDNA 4 的 MI355X 上，8 个 XCD 仍然通过 SoIC 3D 堆叠在下方的基底裸片上，但 XCD 移动到 TSMC N3P，每个拥有 32 个活动 CU（总共 256 个，而 MI300X 为 304 个；每个 XCD 的数量减少以释放面积用于更大的矩阵核心和 160 KB LDS）。4 个 MI300X IOD 折叠为 2 个，每个在 TSMC N6 上宽两倍，上方托管 4 个 XCD，侧面托管 4 个 HBM3E 堆栈。每个 IOD 现在承载其 256 MB Infinity Cache 的 128 MB 切片、一半的 HBM PHY、其 Infinity Fabric 链路份额和 PCIe Gen 5。两个 IOD 之间的 Infinity Fabric AP 以 5.5 TB/s 的对分带宽运行（比 CDNA 3 高约 15%），8 个堆栈转移到 12-Hi HBM3E，在相同的引脚数量下以 8 TB/s 提供 288 GB，容量比 MI300X 多 50%。该封装总计 1850 亿晶体管，并且仍然向内核呈现为一个 GPU。

Bets

• Bet 1: HPC then AI. HPC and AI are the same bet 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.押注 1：先 HPC 后 AI。在 HPC 和 AI 出现分歧之前，它们是同一项押注：从 CDNA 2 到 CDNA 3 交付全速率 FP64 矩阵，然后在推理经济性决定性地倾向于低精度时，于 CDNA 4 进行分叉。
• Bet 2: Memory capacity. Match or beat the contemporary NVIDIA flagship on HBM capacity every generation since 2021, and add a 256 MB last-level Infinity Cache that absorbs the reuse H100 must hit HBM for.押注 2：内存容量。自 2021 年以来，每一代产品在 HBM 容量上都匹配或超越当时的 NVIDIA 旗舰产品，并增加了 256 MB 的末级 Infinity Cache，以吸收 H100 必须访问 HBM 才能完成的重用数据。
• Bet 3: Early 3D-stacking. 3D-stack compute on cache and I/O before NVIDIA does: TSMC SoIC hybrid-bonded XCDs on IODs in 2023, while NVIDIA stayed monolithic until 2025.押注 3：早期 3D 堆叠。在 NVIDIA 之前实现缓存和 I/O 的 3D 堆叠：2023 年在 IOD 上采用 TSMC SoIC 混合键合的 XCD，而 NVIDIA 直到 2025 年仍保持单片设计。
• Bet 4: Coherent CPU+GPU. The MI300A APU is the most chiplet-aggressive product ever shipped, and the El Capitan deployment is the proof.押注 4：相干 CPU+GPU。MI300A APU 是有史以来最激进的小芯片产品，El Capitan 超级计算机的部署就是证明。
• Bet 5: Open scale-up fabric. UALink and OCP MX over NVLink and proprietary FP4.押注 5：开放式扩展互联架构。用 UALink 和 OCP MX 取代 NVLink 和专有的 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.内存押注带来了扩展性后果：当 8 片 MI300X 芯片拥有 1.5 TB HBM，而 8 片 MI350X 拥有 2.3 TB 时，你可以在单个 8-GPU 盒子内放入一个 FP8 格式的 405B 参数模型（包括权重、KV 缓存以及用于更长上下文和更大批次的余量），而在 8× H100（640 GB）上运行同一模型则需要仔细的分片。对于 2024–2025 年的推理工作负载，AMD 的扩展架构无需在机架层面匹配 NVL72 即可在盒子层面保持竞争力。但在前沿训练方面，它必须做到这一点，而 AMD 直到 2026 年才给出答案。

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.在 MI355X 之前，AMD 的扩展是指基于 Infinity Fabric 的 8-GPU OAM 平台。每片 MI300X 拥有 7 条 IF 链路（连接盒内每一个对等点），双向带宽为 128 GB/s，在全连接的全互联拓扑中提供每个 GPU 896 GB/s 的网状带宽。MI350X 将每条链路提升至 153.6 GB/s（每个 GPU 约 1,075 GB/s），但保持 8-GPU 的形态。该平台符合 OCP 的 UBB 2.0 标准：与 NVIDIA HGX 基板使用相同的机械插槽，因此服务器供应商无需重新设计系统即可在同一机箱上提供 AMD 或 NVIDIA 产品。

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 Helios is built to close.AMD 在 MI355X 之前没有交付机架级的 NVL72 等效产品。在 MI300X 集群上运行大型模型的客户通过以太网跨多个 8-GPU 盒子进行扩展，不得不承担扩展延迟，而 NVIDIA 用户则可以将数据保留在扩展互联内部。这就是训练领域中关键的差距，也是 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.Helios 是 AMD 首个机架级扩展域，于 2026 年下半年随 MI455X 一起发布。每个机架 72 个 GPU，约 31 TB HBM4，1.4 PB/s 聚合 HBM 带宽，2.9 ExaFLOPS FP4 / 1.4 ExaFLOPS FP8，260 TB/s 的扩展互联带宽，43 TB/s 的外部扩展带宽。其外形采用 Open Rack Wide (ORW)（Meta 2025 年提交给 OCP 的标准，双宽且液冷），而非 AMD 专有机箱。基于 Meta 的参考设计而非从零开始设计机架是 AMD 的刻意押注：任何标准化 ORW 的超大规模数据中心无需进行定制化的设施改造即可部署 Helios。

The fabric is UALink: 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（Ultra Accelerator Link），这是一个 AMD 与 Apple、AWS、Cisco、Google、HPE、Intel、Meta、Microsoft 和 Synopsys 共同创立的开放联盟标准。UALink 200G 1.0（2025 年 4 月）定义了 200 GT/s 的通道和每个方向 800 Gbps 的带宽，通过交换拓扑可扩展至每个 Pod 1,024 个加速器。其承诺是提供一种可与 NVLink 相媲美但不受单一厂商控制的缓存一致性互联：任何供应商都可以制造 UALink 交换机，任何加速器都可以使用 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 UALoE (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.问题在于：原生 UALink 交换芯片直到 2027 年才能大规模出货。Astera Labs 的 Scorpio 以及来自 Auradine、Enfabrica 和 Xconn 的竞争产品都瞄准 2026 年底/2027 年的部署。Helios 在发布时使用 UALoE（通过标准以太网隧道传输 Infinity Fabric）作为权宜之计，在等待原生 UALink 架构的同时保留编程模型。原生 UALink 交换功能将于 2027 年随 MI500 一起推出。发布时，Helios 更像是一个快速的以太网隧道相干集群，而非 NVL72 那种真正的缓存一致性 NVLink 域：这是一个在时间节点上的妥协，目的是为了在 2026 年下半年以具备竞争力的产品入市。

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).AMD 不提供 InfiniBand。整个外部扩展堆栈基于以太网，锚定在另一个开放标准上：超以太网联盟（UEC）。

UEC 1.0 (released June 2025) defines Ultra Ethernet Transport (UET): 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.UEC 1.0（2025 年 6 月发布）定义了超以太网传输（UET）：一种基于标准以太网的新型 RDMA 传输协议，具有数据包喷洒（packet spraying）、基于 SACK 的选择性重传和现代拥塞控制功能。UET 不是 RoCEv2（后者将 InfiniBand 传输封装在以太网帧中）；它是针对大规模 AI 架构对 RDMA 语义的彻底重新设计。AMD 是与 Broadcom、Cisco、Meta 和 Microsoft 一起创始的成员。与 UALink 的策略相同：拥有标准，而非拥有实现。

The NIC is Pensando, the networking startup AMD acquired in 2022. Pollara 400 is the current AI NIC: 400 GbE, P4-programmable, UEC-ready, PCIe Gen 5, paired with MI300X / MI355X. Vulcano 800 ships in 2026 alongside MI455X: UEC 1.0 compliant, PCIe Gen 6, native UALink interfaces, 8× the per-GPU scale-out bandwidth of Pollara. Salina 400 is the front-end DPU (16× Arm Neoverse-N1, dual 400 GbE) for storage / SDN / firewall, equivalent to NVIDIA's BlueField, distinct from the AI back-end NIC.网卡是 Pensando，即 AMD 在 2022 年收购的网络初创公司。Pollara 400 是当前的 AI 网卡：400 GbE，P4 可编程，支持 UEC，PCIe Gen 5，与 MI300X / MI355X 配套。Vulcano 800 将于 2026 年随 MI455X 一起出货：兼容 UEC 1.0，PCIe Gen 6，原生 UALink 接口，外部扩展带宽是 Pollara 的 8 倍。Salina 400 是用于存储/SDN/防火墙的前端 DPU（16× Arm Neoverse-N1，双 400 GbE），相当于 NVIDIA 的 BlueField，与 AI 后端网卡不同。

The switch silicon, though, isn't AMD's. Helios's 43 TB/s scale-out fabric runs through Broadcom Tomahawk 6: a 102.4 Tbps Ethernet switch ASIC with co-packaged optics ("Davisson"). AMD has no in-house 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.不过，交换机芯片并非 AMD 自有。Helios 的 43 TB/s 外部扩展架构运行在 Broadcom Tomahawk 6 上：这是一款 102.4 Tbps 以太网交换机 ASIC，带有共封装光学器件（“Davisson”）。AMD 没有内部 CPO，也没有内部交换机 ASIC；光学层是合作伙伴的硅片。NVIDIA 拥有其整个堆栈：InfiniBand、Spectrum-X 以太网、ConnectX、BlueField、Quantum-X 光子 CPO，全部自研。AMD 拥有其中一层（通过 Pensando 实现的网卡+DPU），并押注开放标准加顶尖的合作伙伴芯片将胜过垂直整合。

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.行业已经转向了 AMD 的方向。Dell'Oro 报告显示，2025 年以太网处理的 AI 外部扩展架构容量是 InfiniBand 的两倍多；AWS、Microsoft、Meta、Oracle 和 xAI 都已将其基于 AMD 的 AI 集群标准化为以太网。剩下的问题不在于以太网能否在 RDMA 语义上匹配 InfiniBand（UEC 弥补了这一差距），而在于 Helios 能否足够快地缩小与 NVL72 的机架级差距，从而赢得目前默认使用 NVIDIA 的前沿训练工作负载。

Software

ROCm is the open-source counterpoint to CUDA. 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.ROCm 是 CUDA 的开源对应方案。NVIDIA 的堆栈是专有且垂直整合的（cuBLAS、cuDNN、TensorRT-LLM 作为仅由 NVIDIA 维护的二进制块发布），而 ROCm 是原生于 GitHub 的，并押注于开放标准（PyTorch、Triton、vLLM、OCP MX）而非封闭的库集。与 NVIDIA 的软件差距是客观存在的，但 AMD 的策略是通过开放社区来弥补，而不是从零开始构建一个平行的 CUDA 堆栈。

The bottom of the stack is HIP, AMD's CUDA-compatible C++ runtime. hipify 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 (TMA descriptors, wgmma, tcgen05.mma) has no clean ROCm analog and has to be rewritten by hand.堆栈底层是 HIP，这是 AMD 的 CUDA 兼容 C++ 运行时。hipify 会自动将 CUDA 源码转换为 HIP。大量的 HPC 代码（HACC、Laghos、QMCPack）开箱即用，移植率达到 80–95%：这就是 CORAL-2 的数据。现代 AI 内核的移植效果较差：任何涉及 Hopper 或 Blackwell 特定原语（TMA 描述符、wgmma、tcgen05.mma）的代码都没有干净的 ROCm 模拟，必须手动重写。

Above HIP sits a library tier structured to mirror NVIDIA's, one-to-one by name: rocBLAS for cuBLAS; hipBLASLt for cuBLASLt; MIOpen for cuDNN; RCCL for NCCL; Composable Kernel (and its modern 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 vLLM as the open-source serving engine and ship AMD-specific operators (AITER) that plug into it; the dedicated ROCm CI for vLLM took test-pass rate from 37% to 93% across early 2026.在 HIP 之上是按照 NVIDIA 命名方式一一对应的库层：rocBLAS 对标 cuBLAS；hipBLASLt 对标 cuBLASLt；MIOpen 对标 cuDNN；RCCL 对标 NCCL；Composable Kernel（及其现代 ck-tile DSL）对标 CUTLASS；rocprofv3 / rocprof-sys / rocprof-compute 对标 Nsight 系列。不过，AMD 没有 TensorRT-LLM 的第一方模拟。AMD 的做法是支持 vLLM 作为开源服务引擎，并提供插入其中的 AMD 特定算子（AITER）；vLLM 的专用 ROCm CI 在 2026 年初将测试通过率从 37% 提高到了 93%。

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.PyTorch 路径是一等公民。Eager 模式的 PyTorch 自 2018 年起就在 ROCm 上运行；torch.compile 通过 Triton 进行降级编译，而 Triton 的 ROCm 后端（带有用于预编译数学内核的 AOTriton）已进入上游。这里没有 XLA 风格的中间 IR；ROCm 直接编译为 HIP / Triton / CK。随着 Triton 成为 PyTorch 中的默认内核路径，大部分移植成本消失了：一个通过 torch.compile 运行的内核可以在 CUDA 和 ROCm 上运行而无需更改源码。这就是 AMD 开放策略背后的架构押注：Triton 的 Python DSL 成为跨供应商的通用语言，从而规避了对 CUDA 等效内核生态系统的需求。

FlashAttention is the load-bearing case. FA2 is production on MI300X via Composable Kernel; PyTorch defaults to CK or AOTriton on ROCm. FA3 (Hopper-tuned) is partially supported via AITER + CK, but Dao-AILab's canonical implementation remains CUDA-only. FA4 (Blackwell, March 2026) has no ROCm port at all. HipKittens, 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.FlashAttention 是典型的承重案例。FA2 在 MI300X 上通过 Composable Kernel 实现生产级应用；PyTorch 在 ROCm 上默认使用 CK 或 AOTriton。FA3（针对 Hopper 优化）通过 AITER + CK 获得部分支持，但 Dao-AILab 的规范实现仍然仅限 CUDA。FA4（Blackwell，2026 年 3 月）完全没有 ROCm 移植。HipKittens（Hazy Research 在 2025 年 11 月推出的 ThunderKittens 的 MI355X 移植版）声称用约 500 行代码实现了与手动调优的 AITER 同等的前向传递性能。模式是：开源学术内核在 NVIDIA 发布后几个月内填补 AMD 的短板，而不是几年。

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 BM.GPU.MI300X.8 went GA in September 2024, with MI355X following in 2026. These are real serving fleets at hyperscaler scale, not pilots.生产部署验证了该策略。Microsoft Azure 的 ND MI300X v5 实例于 2024 年 5 月正式发布；OpenAI 在其上运行 GPT 推理。Meta 通过 Grand Teton 平台在 MI300X 上发布 Llama 3 / Llama 4 推理。Oracle OCI 的 BM.GPU.MI300X.8 于 2024 年 9 月正式发布，MI355X 紧随其后于 2026 年推出。这些都是超大规模数据中心级别的真实服务集群，而非试点项目。

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.客观差距依然存在。独立基准测试（Phoronix，2026 年 3 月）显示，在标准 PyTorch / vLLM / SGLang 工作负载下，在同等硅片和同等精度上，ROCm 7.2 比等效的 CUDA 慢 10–25%。ROCm 7 实现了功能对等，但未实现性能对等。FlashAttention-4 的短板（利用 Blackwell 最新原语的研究代码）是 NVIDIA 护城河最稳固的地方；它没有干净的 ROCm 模拟，只能等待手写的 AITER 内核或 HipKittens 类社区移植。NVIDIA 将工程师派驻到前沿实验室；AMD 通过 GitHub 发布内核。这些策略在通用工作负载（Llama 推理、注意力机制、稠密 Transformer 训练）上趋于一致，但大量新型研究代码仍然需要 MI300X / MI355X 部署投入 NVIDIA 用户无需支付的工程时间。

Cerebras WSECerebras WSE

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 dataflow: a core sits idle until a wavelet arrives, control bits in the wavelet select which handler task fires, and eight hardware microthreads switch cycle-by-cycle as tensor operands arrive and drain. No warps, no warp schedulers, no caches to miss, no reorder buffer: the arrival of data is the schedule.GPU 是一种层级结构：SM 内部有 Warp，Warp 内部有线程；封装内部有裸片，机架内部有封装，每个边界都有自己的带宽、延迟和编程结构；任何由裸片构建的加速器都继承了这种结构。WSE 是一个平坦的平面：90 万个相同的核心在 2D 网格中边缘对齐，没有共享缓存，没有全局内存，核心之间没有任何边界。每个核心都很小，在 WSE-2 上约为 38,000 平方微米，大约一半是 SRAM，一半是逻辑，峰值功耗 30 毫瓦：48 KB 本地 SRAM，16 个通用寄存器，6 级流水线，4 宽 FP16 FMAC SIMD（WSE-3 上为 8 宽），以及一个连接到架构的五端口路由器。执行是数据流驱动的：核心在小波（wavelet）到达前处于闲置状态，小波中的控制位选择触发的处理程序任务，当张量操作数到达并消耗时，8 个硬件微线程逐周期切换。没有 Warp，没有 Warp 调度器，没有缓存缺失，没有重排序缓冲区：数据的到达即是调度。

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.步进机一次曝光一个掩模场，每次拍摄约 850 平方毫米，这就是为什么每个传统芯片都生活在这个天花板之下（也是为什么 NVIDIA 在 B200 达到极限时将其变为双裸片的原因）。Cerebras 像 TSMC 的任何其他客户一样，在 12×7 的网格中打印 84 次相同的约 550 平方毫米裸片，然后通过与 TSMC 共同开发的工艺，在通常用于锯切的 <1 毫米划线处铺设额外的顶级金属。网格在每个接缝处通过源同步并行接口（WSE-3 上每个裸片 2,880 GB/s）交叉，整个裸片间层成本约为 97 瓦。对软件而言，接缝不存在：一个统一的网格，一个芯片。

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.晶圆级技术曾被尝试过，但因良率而失败：单片晶圆计算机上的一个缺陷会摧毁整片晶圆，这就是 20 世纪 80 年代该想法被埋没的原因。Cerebras 的答案是颗粒度。H100 上的一个缺陷会禁用整个约 6 平方毫米的 SM；WSE 上的同一个缺陷仅禁用一个 0.05 平方毫米的核心。WSE-3 制造约 97 万个核心并出货 90 万个：约 7% 的备用池加上冗余架构链路，让硬件能够绕过每个缺陷并恢复完整的逻辑网格。

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 data-structure registers (DSRs), each holding a tensor descriptor: 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.核心不同寻常之处不在于数据路径，而在于指令的定义。除了 16 个通用寄存器外，还有 44 个数据结构寄存器（DSR），每个寄存器保存一个张量描述符：基地址、范围和步长，最多支持四个维度。指令通过 DSR 命名操作数，因此单个 FMAC 指令表示将到达的数据流与驻留张量相乘并累加到另一个张量中，硬件会根据张量的持续时间流式传输元素。乘法周围没有软件循环，每个元素也没有指令获取；循环存在于描述符中。NVIDIA 花了五代 Tensor Core 的时间将矩阵乘法推向单一描述符驱动的命令；在 WSE 核心上，张量指令没有其他形式。

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.排序是架构的工作。颜色（color）是一种静态路由的虚拟通道，在编译时绑定了处理程序任务，因此在颜色上发送小波就是调用目标核心上的代码：16 个控制位是调用，16 个数据位是参数。任务调度器将飞行中的张量操作保留在核心的 8 个微线程上，并根据操作数的可用性在每个周期进行切换。这与 Warp 调度器处理 64 个驻留 Warp 的隐藏延迟工作相同，只不过是用 8 个上下文完成的，因为隐藏的延迟是繁忙的 SRAM 存储体或邻居跳转，而不是 HBM 往返。

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.48 KB 的本地 SRAM 是为数据路径而非局部性组织的：8 个单端口 6 KB 存储体每周期提供两次 64 位读取和一次 64 位写入，正好对应两个 4 元素 FP16 操作数输入和一个结果输出，即 WSE-2 FMAC 的宽度。一个 256 字节的软件管理缓存（WSE-3 上为 512 B）将最热的数据保留在流水线旁边。这是机器微缩版的论点：每个核心的内存带宽和计算能力完全匹配，晶圆将这种平衡继承了 90 万次。

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.晶圆上没有矩阵单元。NVIDIA、Google 和 AMD 都将 FLOPs 集中在专用的矩阵乘法引擎（Tensor Core、MXU、Matrix Core）中，区别仅在于如何为该引擎供电；Cerebras 则通过架构组装矩阵乘法。GEMM 作为晶圆级的编排运行：每个到达的权重沿持有激活值的核心行广播，每个核心对驻留切片进行乘加运算（每个权重一个 AXPY），部分和在网格中进行归约。Tensor Core 从寄存器瓦片中获得的数据重用，以及 MXU 从其布线中获得的数据重用，WSE 通过几何结构获得：激活值从不移动，因此唯一在飞行中的操作数就是正在被乘的那个。

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 15.8 PFLOPS FP16 (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 bandwidth machine, and the FLOPs exist to keep up with the SRAM.FLOPs 账本需要谨慎，因为 Cerebras 打印的数字不是可比较的数字。WSE-3 的 125 PFLOPS 标题数据是稀疏 FP16：它假设硬件在理想稀疏张量上具有约 8 倍的跳过零收益。稠密计算约为 15.8 PFLOPS FP16（推导：90 万核心 × 8 宽 FMAC × 1.1 GHz；Cerebras 未发布官方稠密数字）。这是真实的计算能力，但重点不在此：按每瓦特计算，晶圆上的稠密 FLOPs 输给了每一款现代 GPU。晶圆从来不是 FLOPs 机器。它是一台带宽机器，FLOPs 的存在是为了跟上 SRAM。

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: 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.跳过零（Zero-skipping）是数据流发挥作用的地方。由于计算是由到达的数据触发的，零不会触发任何操作：零在发送者处被过滤，接收核心永远看不到它们，也不会消耗周期。这是非结构化的、元素颗粒度的稀疏性，是 NVIDIA 的 2:4 结构化稀疏性仅采样的通用情况。这也是目前尚未行使的选项。Cerebras 自己的稀疏预训练结果（SPDF：13 亿参数下 75% 的稀疏性；67 亿参数的后续研究）是供应商编写的，且小于 70 亿参数，目前尚未披露任何旗舰客户模型采用稀疏训练：硬件上最大的运行模型 Jais 2 是稠密的。唯一能够利用非结构化稀疏性的硅片尚未发布使用它的旗舰模型。

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.层级只有一层：核心内 48 KB 切片中的 44 GB SRAM，晶圆上没有其他东西。没有 HBM，没有 L2，没有驱逐策略；每个字节距离 FMAC 仅一个周期。引用的带宽是 21 PB/s，这个数字名副其实：它是 90 万个本地 SRAM 端口的总和，是晶圆上的聚合值，而非点对点链路，不能与 HBM 数字相比。诚实的比较是每 FLOP 的字节数：晶圆可以为每个稠密 FP16 FLOP 提供约 1.3 字节，而 B200 从 HBM 获得约 0.002 字节。在这个轴上，每个 GPU 和 TPU 都处于饥饿状态；WSE 是唯一保持平衡的机器。解码阶段（每个 Token 对权重的完整读取，纯粹的带宽问题）正是晶圆所擅长的。

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 five orders of magnitude. 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.层级的另一面是它的边缘。晶圆与外界的连接是 12×100 GbE：1.2 Tb/s，几乎只比连接到单个 Blackwell GPU 的单个 ConnectX-8 网卡多一点。在片上 SRAM 和片外以太网之间存在五个数量级的差距。NVIDIA 的层级逐渐下降，每一层比前一层慢几倍；WSE 有两层，中间有悬崖。晶圆是一座孤岛，孤岛的超能力和它的牢笼是同一事实。

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.而且孤岛并没有在增长。SRAM 密度在领先节点上实际上已停止扩展：尽管工艺节点完全收缩且晶体管数量跃升 54%，但 WSE-3 的 SRAM 仅比 WSE-2 多 10%。逻辑部分在不断缩小；六晶体管 SRAM 单元却没有。该架构最稀缺的资源正是下一个工艺节点不再能买到的东西。

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 MemoryX, 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" (Cerebras's phrase). Model size is bounded by MemoryX, not by the 44 GB; the 44 GB bounds activations and batch.在晶圆上训练颠倒了其他人认为理所当然的流程：在 GPU 或 TPU 上，权重是驻留的，激活值流式传输；在 WSE 上，激活值是驻留的，权重流式传输。主权重位于集群旁边的 DRAM 和闪存设备 MemoryX 中。层层权重流过晶圆，触发针对驻留在 SRAM 中的激活值的乘加运算，然后离开；梯度在反向传播中流回，优化器步骤在 MemoryX 的 CPU 上运行（权重更新是 O(参数) 的元素级工作，没有重用，因此 CPU 级计算可以跟上）。晶圆从不存储权重，“甚至不是暂时存储”（Cerebras 的话）。模型大小受 MemoryX 限制，而非 44 GB；44 GB 限制了激活值和批次。

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 data parallelism through SwarmX, 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.这换来的是编程模型。一个晶圆拥有完整层的激活值，因此没有张量并行、流水线并行或 FSDP 分片：70B 模型被编写为单设备程序，多系统扩展是通过 SwarmX 实现的纯数据并行，SwarmX 是一个广播/归约树，将一个权重流扇出到 N 个晶圆，并在返回途中汇总它们的梯度。GPU 训练中占主导地位的并行策略电子表格在 Cerebras 这里根本没有页面。

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.代价是市场所揭示的偏好中的规模。规格表上写着 2,048 个 CS-3；有史以来披露的最大集群是 64 个（Condor Galaxy 3）。自 CS-1 发布七年来，平台上披露的最大的从零开始训练的模型是 Jais 2，参数量 70B，Token 数 2.6T，由锚定客户 G42 在 Cerebras 工程师嵌入的情况下训练。自 CS-1 以来，七年间没有任何人训练过超过 70B 的模型。而且 GPU 实验室通常以 35–45% 发布的使用率（MFU），从未在任何 Cerebras 运行中披露过。

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.数值计算可以用一句话概括：FP16 和 BF16 加 FP32 累加，加上（从 WSE-3 开始）一个 16 宽的 8 位整数路径，Hot Chips 披露将其标记为定点数。没有 FP8，没有 FP4，没有微缩放。当其他所有供应商每一代都将精度减半并通过块缩放买回精度时，Cerebras 仍然以 16 位计算并将其作为质量差异化因素推销（“原始 16 位权重”）。张力是显而易见的：SRAM 容量是该架构最稀缺的资源，8 位权重将使模型所需的晶圆数量减半。16 位唯一性是数值上的坚定信念还是数据路径路线图的缺失，这是一个悬而未决的问题；没有主要的 Cerebras 来源显示晶圆上的任何地方有浮点 8。

Bets

• Bet 1: Don't cut the wafer. 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.押注 1：不要切割晶圆。裸片边界是行业其他部分支付的税：SerDes、中介层、HBM 堆栈、电缆、交换机。用金属缝合 84 个掩模场，竞争对手系统中最高带宽的边界根本不存在。
• Bet 2: SRAM is the only memory. 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.押注 2：SRAM 是唯一的内存。以业内最陡峭的比例用容量换带宽：44 GB，晶圆上聚合 21 PB/s。平衡机器，而不是将不平衡隐藏在层级之后。
• Bet 3: Dataflow cores, no matrix unit. 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.押注 3：数据流核心，没有矩阵单元。90 万个由到达的小波触发的微型核心，矩阵乘法由广播、FMAC 和网格归约组装而成：跳过零是免费的，而不是特殊模式。
• Bet 4: Weights move, activations stay. Weight streaming decouples model size (MemoryX) from wafer memory (44 GB) and collapses cluster scaling to pure data parallelism.押注 4：权重移动，激活值保持。权重流式传输将模型大小（MemoryX）与晶圆内存（44 GB）解耦，并将集群扩展折叠为纯数据并行。
• Bet 5: Sell latency, not throughput. 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.押注 5：销售延迟，而非吞吐量。晶圆重新读取整个模型每个 Token 的速度比任何基于 HBM 的机器都快；将这种速度定价为高端产品，而不是在每个 Token 的成本上竞争。

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.扩展互联（Scale-up）和外部扩展（Scale-out）在这里有不同的含义。NVIDIA 的扩展互联问题（使 72 个封装表现得像一个设备）在 WSE 上通过光刻解决：相干域从晶圆厂出来时就是一块。剩下的就是晶圆边缘之外的一切，没有其他机器能像它那样硬、那样早地触及其边缘。

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.晶圆的内部架构没有 SerDes、电缆、收发器，也没有每条链路的边际成本：路由是编译的，每个跳转是一个周期，广播是原生架构原语而非交换机功能。当 NVL72 花费 5,184 根铜缆和一托盘 NVSwitch ASIC 为 72 个 GPU 提供 130 TB/s 的全互联带宽时，WSE 的等效域是一个单一的光刻对象。问题在于域大小是一个常数。NVIDIA 的扩展互联域每一代都在增长（三年内从 NVL72 到 NVL576）；晶圆自 2019 年以来一直是 46,225 平方毫米，并将保持不变。300 毫米是行业运行的最大晶圆（450 毫米的过渡十年前就死了），因此 Cerebras 的扩展路线图就是下一个节点在密度上所能提供的：没有更多的面积可得。

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.训练外部扩展是 SwarmX，它只做一件事：复制。将权重流广播到 N 个晶圆，在返回路径上归约它们的梯度；批次随系统数量增长，模型大小则不增长。声称的 2,048 个系统的上限（“256 exaFLOPS”，稀疏）从未被构建过；64 个系统是有的。

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.推理完全放弃了权重流式传输；算术上是致命的。在约 150 GB/s 的管道上为每个解码的 Token 从 MemoryX 流式传输 70B 模型的 140 GB 数据，每个 Token 大约需要一秒。因此，推理将权重驻留在 SRAM 中，并在层边界处跨晶圆分片模型：Llama 70B 在“少至四个”CS-3 上，通过以太网流水线并行，每个额外的晶圆贡献 44 GB 的权重加 KV 容量和 23 kW 的负载。

The speeds are real, and independently verified. Artificial Analysis 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 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.速度是真实的，并经过独立验证。Artificial Analysis 在 2024 年 8 月发布时测得 Llama 3.1 8B 为 1,850 tokens/s，70B 为 446 tokens/s，Llama 405B 为 969 tokens/s（首个 Token 延迟 240 毫秒），2025 年 Llama 4 Maverick 为 2,522 tokens/s，约为当时最佳 Blackwell 发布数据的 2.4 倍。供应商引用的峰值更高（70B 配合推测解码为 2,100；GPT-OSS-120B 为 3,000，而实时独立测量接近 2,000）。没有 GPU 提供商在每用户解码速度上能接近这一点。

The economics are the sharp edge. Forty-four GB per wafer means a frontier-scale model consumes fleets: SemiAnalysis 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.经济性是尖锐的边缘。每个晶圆 44 GB 意味着前沿规模的模型消耗集群：SemiAnalysis 估计 1.6T 参数级模型需要约 24 个 CS-3，而该模型仅需少量 GPU 机架即可容纳，每个系统估计物料清单约为 45 万美元，标价约为 200–300 万美元（从未正式披露）。在解码过程中，晶圆巨大的 FLOPs 大多处于闲置状态；Cerebras 拒绝披露批次大小，也从未发布过每个系统的吞吐量。每个 Token 的 API 定价对于相同的开源模型来说，大约是基于 GPU 提供商的 3–5 倍，Llama 405B 被悄悄从 API 中移除，SemiAnalysis 将其解读为服务经济性不达标。固定的 SRAM 也为上下文定价：KV 缓存与权重驻留在相同的 44 GB 中，因此长上下文会窃取容量并迫使每个副本需要更多系统；API 上限为 131K Token，而前沿提供商服务 256K–1M。MoE 得到支持（Qwen3-235B 约为 1,500 tokens/s，供应商报价），但这是该格式的最坏情况：巨大的参数足迹在每次只触及少数专家，且保存在最昂贵的内存中。

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 at signing and since grown past $20B, the largest endorsement wafer-scale has ever received. The first flagship to ship on that capacity is GPT-5.6 Sol, launched July 2026 at a quoted 750 tokens/s.市场已经诚实地为这一点定价。Mistral 的 Le Chat（约 1,100 tokens/s）、Perplexity Sonar 和 Meta 的 Llama API 都为延迟买单；2026 年 1 月，OpenAI 签署了到 2028 年 750 MW 的 CS-3 容量，据报道签署时价值超过 100 亿美元，此后增长到 200 亿美元以上，这是晶圆级技术获得的最高背书。在该容量上出货的首个旗舰是 GPT-5.6 Sol，于 2026 年 7 月推出，报价为 750 tokens/s。

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 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, the Dutch national compute centre) reports unsupported layer types and no 1:1 porting path for standard PyTorch code.堆栈像 TPU 一样是编译器驱动的，但通过一个窄得多的孔径：Cerebras 编译器是一个内核匹配器，而不是通用的代码生成器。cerebras.pytorch 通过惰性张量将训练步骤追踪到 Torch-MLIR 和图形 IR 中，然后将子图与手写内核库进行匹配，对于没有匹配的算子回退到较慢的自动生成内核。按 GPU 标准，文档化的约束是严苛的：仅限静态图，没有动态形状，没有数据依赖的控制流，步骤中间没有急切的张量访问，PyTorch 版本固定在上游之后。最好的独立从业者账户（SURF，荷兰国家计算中心）报告了不支持的层类型，且没有标准 PyTorch 代码的 1:1 移植路径。

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, CSL, exposes the raw machine (tasks, wavelets, colors) and has produced striking HPC results (a TotalEnergies stencil code 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.而且没有内核逃生舱口。CUDA 对新型注意力变体的回答是编写内核；TPU 是 Pallas；ROCm 是 Triton。Cerebras ML 堆栈根本没有用户内核路径：当匹配器严重失误时，修复者是 Cerebras 工程师。一种单独的 SDK 语言 CSL 暴露了原始机器（任务、小波、颜色），并产生了惊人的 HPC 结果（TotalEnergies 的模板代码约为 A100 的 228 倍，在 48 个 CS-2 上入围 Gordon Bell 奖），但这是一个独立的世界，与 PyTorch 流程无关。平台上的每个旗舰模型（Jais, BTLM, Med42）都是与嵌入的 Cerebras 员工共同开发的。

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.这其中有一种奇怪的免疫力。FlashAttention，GPU 时代最具代表性的内核谱系，是一种通过内存层级平铺注意力机制的方案，而 WSE 没有层级可供平铺：导致 AMD 多年移植滞后的优化类问题在这里根本不适用。但免疫力和贫困是同一事实。在 CUDA 上复合的第三方内核生态系统在这里没有附着表面；平台历史上每一次内核改进都只有一位作者。

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.晶圆的未来在哪里？拥有一个真实的、诚实赢得的利基市场：批次为 1 的解码速度，经过独立验证，由将延迟定价高于成本的客户买单。在利基市场周围，是坚硬的墙壁：每个 Token 的定价是 3–5 倍，70B 的训练上限维持了七年，2025 年收入仍约 86% 集中在两个与阿布扎比相关的客户身上（根据 2026 年 5 月 IPO 前后的 S-1 文件），以及最稀缺的资源 SRAM 密度，在模型不断增长时停止了扩展。Hennessy 和 Patterson 预言了一场寒武纪大爆发；WSE 是其最极端的身体计划，它决定内存墙是一个封装选择，并花费 46,225 平方毫米的硅片拒绝制造它。

AWS TrainiumAWS Trainium

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 cluster of decoupled, specialised engines: a Tensor Engine (the 128×128 systolic array), a Vector Engine for reductions, a Scalar Engine for pointwise math, and a programmable GPSIMD Engine of eight 512-bit vector processors for whatever fits none of the other three. Around them sit the data-movers: 128 DMA engines, a Sync Engine that sequences transfers, and (from Trn2) dedicated CC-Cores 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.另一个自研芯片的故事属于谷歌，Trainium最好被理解为TPU理念在不同云环境下的重构。其底层的赌注是一样的（由软件管理的SRAM提供数据的脉动阵列，由编译器提前调度，没有缓存，也没有线程调度器），但单元的组装方式不同。一个Trainium芯片携带少量NeuronCore（Trn1上有2个，Trn2和Trn3上有8个），每个NeuronCore并非单一的矩阵乘法引擎，而是由解耦的专用引擎组成的集群：张量引擎（128×128脉动阵列）、用于规约运算的向量引擎、用于点对点数学运算的标量引擎，以及一个由八个512位向量处理器组成的可编程GPSIMD引擎，用于处理不适合前三种引擎的任务。在它们周围是数据搬运单元：128个DMA引擎、一个用于序列化传输的同步引擎，以及（从Trn2开始）用于集合通信的专用CC-Core。这里没有线程束（warps）或波前（wavefronts）；引擎作为静态调度的流数据管道运行，负载设计的决策重点在于脉动阵列周围的组件，而非阵列本身。

Compute

The Tensor Engine owns the matmul FLOPs; the other three engines own everything else. It is a 128×128 grid of processing elements (16,384 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 PSUM, a small accumulator SRAM the engine can read-add-write so a contraction longer than 128 folds into place along the $K$ axis. This is the same $D = A \cdot B + C$ tile 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.张量引擎负责矩阵乘法浮点运算，其余三个引擎负责其他所有工作。它是一个128×128的处理单元网格（16,384个MAC），以权重静态（weight-stationary）方式运行：一个操作数分块被加载到阵列中并保持不动（LoadStationary），另一个操作数流经它（MultiplyMoving），部分和存入PSUM（一个小型的累加器SRAM），引擎可以对其进行读-加-写操作，从而使长于128的收缩运算能够沿着K轴折叠完成。这是每个矩阵乘法加速器核心中相同的D=A⋅B+C分块矩阵乘法（MMA）；但NVIDIA将其包装在线程束层级中，谷歌通过VLIW指令包发出，而Trainium将其作为针对命名暂存区的成对显式指令暴露出来。

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.该阵列在所有三代产品中物理上固定为128×128；改变的是每个单元打包的乘积数量。Trn1的NeuronCore-v2以FP32累加运行BF16/FP16，FP8仅能以BF16速率运行（无加速）。Trn2的v3通过双倍泵送FP8，实现了等效的256×128阵列，这是首款在8位运算上实现真正2倍加速的Trainium芯片。Trn3的v4通过打包微缩放操作数，实现了等效的512×128阵列，达到BF16速率的4倍。物理乘加单元的数量从未改变；数据通路只是向它们输送了更窄的数字。

The other three engines are what keep the array busy. The Vector Engine handles cross-element reductions (layernorm, softmax, pooling); the Scalar Engine handles one-in-one-out pointwise ops (activations, GELU); the GPSIMD 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 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.其他三个引擎负责保持阵列的繁忙状态。向量引擎处理跨元素规约（层归一化、Softmax、池化）；标量引擎处理一对一的点对点运算（激活函数、GELU）；GPSIMD引擎包含八个完全可编程的向量处理器，运行C代码，吸收所有不属于前述引擎的任务。一个编译良好的步骤会重叠所有四个引擎的工作：张量引擎进行矩阵乘法，同时向量引擎运行前一个分块的Softmax，DMA引擎调度下一个分块，这种生产者/消费者重叠使得TPU和GPU的注意力机制内核高效，在这里表现为独立的物理引擎，而不是独立的线程束或VLIW槽位。当层结构能清晰地分解到这四种引擎类型上时，该设计就能发挥出优势，Transformer模型大多如此。它在边缘情况会付出代价：如果一个算子不适合任何专用引擎，就会落入可编程的GPSIMD路径，速度较慢，这也是该架构中最容易成为瓶颈的部分。这是Trainium版本的“长尾成本”，每个非GPU加速器都必须承担。

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 HBM (32 GB on Trn1, 96 GB HBM3 on Trn2, 144 GB HBM3e on Trn3). On-chip, closest to the engines, is the State Buffer (SBUF): 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 PSUM, 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.内存层级结构是应用于存储的计算哲学：三层架构，全部由软件管理，没有任何硬件缓存。AWS自己的文档指出了这种对比，指出与CPU或GPU不同，NeuronCore没有缓存，并且“所有内存移动都在程序本身中显式完成”。片外是HBM（Trn1为32 GB，Trn2为96 GB HBM3，Trn3为144 GB HBM3e）。片上最靠近引擎的是状态缓冲区（SBUF）：主要暂存区，带宽约为HBM的20倍，组织成128个分区，每个NeuronCore的大小为24 MiB（v2）、28 MiB（v3）、32 MiB（v4）。阵列和SBUF之间是PSUM，一个专用于矩阵乘法输出的2 MiB累加器。数据流向为HBM → SBUF → 张量引擎 → PSUM → SBUF，每一跳都由编译器发出指令；没有任何数据由硬件预取或驱逐。

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.这正是谷歌VMEM的赌注，即一个编译器必须完美调度且没有缓存来掩盖错误的显式暂存区，与NVIDIA硬件管理的L2和L1缓存截然相反。Trainium继承了这种设计带来的上限和脆弱性：当调度正确时，引擎永不卡顿；当调度错误时，没有回退路径。该设计在适度的峰值浮点运算能力下提供了慷慨的HBM预算，因此与同类NVIDIA部件相比，Trainium每单位计算携带的内存更多。但在绝对容量上，它有所落后：Trn2的96 GB低于H200和B200，Trn3的144 GB（2025年）低于其竞争对手192 GB的B200和288 GB的B300。因此，AWS在论证服务大型模型的经济性时，真正使用的杠杆不是内存领先，而是价格：在自己构建和租赁的硅片上，每单位计算和HBM的成本。

Numerics

Trainium tracks the same precision-halving curve as everyone else (FP32 → BF16 → FP8 → FP4), with two Trainium-specific wrinkles. The first is configurable FP8: rather than fix 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 stochastic rounding 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.Trainium遵循与其他所有人相同的精度减半曲线（FP32 → BF16 → FP8 → FP4），并有两个Trainium特有的改进。首先是可配置的FP8：张量引擎不像Hopper那样固定E4M3和E5M2，而是采用可调的指数偏置，并支持E5M2、E4M3和E3M4，允许编译器为每个张量在范围和精度之间进行权衡。其次，Trn3的FP4不提供额外的吞吐量：OCP MXFP4操作数在到达阵列前被上转换为MXFP8，因此FP4以FP8速率运行，仅节省内存和带宽，不节省计算。两代产品都借鉴了业界的精度恢复技巧：Trn3的微缩放块指数，以及每一代硬件上的随机舍入。唯一不可信的数据是稀疏峰值：AWS宣称的4倍FP8数字在自己的架构页面中被标为密集FP8的2倍（4倍是相对于密集BF16而言），因此营销宣传的加速与数据通路并不完全吻合。

Collectives in Silicon硅片上的集合通信

The block with no clean analogue on a GPU is the collective-communication core. Distributed training and inference spend a large fraction of their wall-clock in 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 CC-Cores per Trn2 chip, wired straight to the NeuronLink 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 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.GPU上没有直接对应物的模块是集合通信核心。分布式训练和推理在集合通信上花费了大量的挂钟时间：每个梯度步骤都是一次All-Reduce，每个MoE层都是一次All-to-All。在GPU上，这些集合通信作为NCCL内核在进行数学运算的同一SM上运行，因此通信和计算争夺相同的硅片资源，重叠必须通过软件实现。Trainium将该功能剥离到专用硬件中：每个Trn2芯片有20个CC-Core，直接连接到NeuronLink端口，在张量和向量引擎持续运行时执行All-Reduce、All-Gather、Reduce-Scatter和All-to-All。这与谷歌在SparseCore上的动作以及Cerebras在片外零过滤器上的动作相同：找到主引擎形状不匹配的工作负载，并在旁边花费少量面积构建专用模块，而不是从核心中窃取周期。通信成为芯片并发执行的任务，而不是暂停计算去执行的任务。

Bets

• Bet 1: The cloud is the product, the chip is a component. Annapurna designs chip, server, rack, 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.赌注1：云是产品，芯片是组件。Annapurna将芯片、服务器、机架、Nitro网络和云API作为一个堆栈进行设计，因此Trainium只需要在AWS内部实现性价比胜出，而无需在商用硅片规格表上获胜。
• Bet 2: Borrow the compute thesis, don't reinvent it. A 128×128 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.赌注2：借鉴计算论文，不要重新发明。128×128权重静态阵列、软件管理的SBUF/PSUM暂存区和全程序编译是TPU的赌注，甚至共享了谷歌的OpenXLA。节省下来的精力投入到网络和机架设计中。
• Bet 3: Collectives belong in silicon. Dedicated 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.赌注3：集合通信属于硅片。专用CC-Core在硬件中将All-Reduce和All-to-All与计算重叠，而不是作为从矩阵乘法单元窃取浮点运算能力的内核运行。
• Bet 4: Reuse the cloud's own network. Scale-out is EFA with the SRD transport: the same Nitro-offloaded, packet-sprayed RDMA that already runs the rest of AWS. No InfiniBand.赌注4：复用云自身的网络。横向扩展采用带有SRD传输的EFA：即已经在AWS其余部分运行的、Nitro卸载的、数据包喷洒式RDMA。没有InfiniBand。
• Bet 5: Move the topology to the workload. Trn1 and Trn2 copied the TPU's 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.赌注5：将拓扑结构迁移到工作负载。Trn1和Trn2复制了TPU的环面结构；Trn3的NeuronSwitch用交换式全对全结构取代了它，因为MoE流量增长超过了最近邻通信。老实说，这遵循了行业剧本：先是谷歌的，现在是NVIDIA的。

Scaling

Trainium's scaling inherits its split from the rest of AWS: a tightly-coupled NeuronLink domain for the chips that must act as one, and the cloud's general-purpose EFA fabric for everything beyond it. The scale-up domain is not cache-coherent shared memory the way 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.Trainium的扩展继承了其与AWS其余部分的分割：为必须协同工作的芯片提供紧密耦合的NeuronLink域，为其他所有内容提供云的通用EFA结构。横向扩展域不是NVLink那样的缓存一致性共享内存；AWS将UltraServer作为池化多TB内存进行营销，但在底层它是基于点对点链路的消息传递，在精神上更接近TPU的ICI，而不是NVSwitch交叉开关。

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 torus, exactly the TPU's choice: a single trn2 instance is 16 chips in a 4×4 2D torus at ~1.28 TB/s per chip, and the Trn2 UltraServer joins four instances into 64 chips on a 4×4×4 3D torus, presenting 83 dense 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.NeuronLink是Trainium的芯片间结构，扮演着NVIDIA的NVLink和TPU的ICI的角色。通过Trn2，它将芯片连成环面，这正是TPU的选择：单个Trn2实例是4×4二维环面中的16个芯片，每个芯片约1.28 TB/s，Trn2 UltraServer将四个实例连接成4×4×4三维环面中的64个芯片，作为一个扩展域提供83个密集FP8 PetaFLOPS和约6 TB HBM。第三个环面轴刻意做得很窄（实例间环路每个芯片约256 GB/s，而实例内部为1.28 TB/s），这是环面结构的典型权衡：布线便宜且拥有巨大的最近邻带宽，代价是跨越直径需要多次跳转。AWS将64芯片的UltraServer定位为NVIDIA 72-GPU NVL72的竞争对手；总计算能力处于同一量级，但环面不是交叉开关，两者在非最近邻流量上的表现大相径庭。

That trade is why Trn3 abandons the torus. NeuronSwitch-v1 is a switched all-to-all 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 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.这种权衡是Trn3放弃环面的原因。NeuronSwitch-v1是一种交换式全对全结构，带宽大约是芯片间带宽的两倍，更重要的是，它扁平化了直径，使任何芯片都能在一次交换跳转内到达任何其他芯片。Trn3 UltraServer扩展到144个芯片，提供362个密集FP8 PetaFLOPS和20.7 TB HBM3e。其动机与推动谷歌为MoE推理采用高基数拓扑的动机相同：专家路由是全对全的，这对环面来说是最坏的情况，而交换机将最长跳转路径变成了一次跨越。Trainium的互联路线图是行业路线图的压缩重演：在工作负载为最近邻时采用环面，在不是时切换到交叉开关。

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 SRD (Scalable Reliable Datagram), offloaded to the Nitro cards rather than run on the accelerator. SRD is AWS's clean-sheet answer to 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.横向扩展不是定制的；它是AWS已经在运行的相同结构。每个Trainium实例携带一个进入数据中心网络的EFA网卡（每个Trn2实例3.2 Tbps），传输协议是SRD（可扩展可靠数据报），卸载到Nitro卡而不是在加速器上运行。SRD是AWS对RDMA的全新回答：它不是RoCE或InfiniBand的单一有序流，而是将每条消息喷洒到多达64条并行路径上，并可靠但乱序地交付，将重组推向集合通信库，规避了单一拥塞路径会导致的队头阻塞。这是AWS为其云构建的传输协议，被重新用于加速器结构。

At the top of the hierarchy is the UltraCluster, stitched together by the 10p10u 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 Project Rainier: roughly half a million Trainium2 chips across multiple US datacenters, brought online for 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.层级结构的最顶端是UltraCluster，由10p10u网络（AWS的速记，指跨数据中心在10微秒延迟下约10 petabits/s的带宽）缝合在一起，扩展到数十万个芯片。证明点是Rainier项目：2025年底为Anthropic上线的跨多个美国数据中心的约50万个Trainium2芯片；到2026年初，Claude运行在超过100万个芯片上，这是任何外部实验室对非NVIDIA训练平台做出的最大承诺。它的存在是因为经济性在端到端上能够闭环。AWS声称Trainium2提供了比其Hopper级GPU实例高30-40%的性价比（AWS数据，对比上一代NVIDIA而非Blackwell），并且因为亚马逊拥有从Nitro卡到API的每一层，该利润空间由亚马逊设定。

Software

Trainium's software makes the borrowing explicit: the Neuron SDK is a compiler-first stack built on the same OpenXLA foundation as the TPU. The Neuron compiler (neuronx-cc) ingests XLA HLO graphs and lowers them to a NEFF 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 PJRT device alongside the TPU. torch-neuronx runs PyTorch through PyTorch/XLA's LazyTensor tracing (record ops, compile the graph at a step boundary), and jax-neuronx lowers JAX through 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.Trainium的软件明确了这种借鉴：Neuron SDK是一个以编译器为先的堆栈，建立在与TPU相同的OpenXLA基础之上。Neuron编译器（neuronx-cc）摄取XLA HLO图并将其降低为NEFF二进制文件，由Neuron运行时加载到NeuronCore上；前端IR是谷歌的，谷歌自己的OpenXLA公告将Trainium列为与TPU并列的一等PJRT设备。torch-neuronx通过PyTorch/XLA的LazyTensor追踪（记录操作，在步骤边界编译图）运行PyTorch，jax-neuronx通过StableHLO降低JAX。在从内核驱动的CUDA到全程序XLA的谱系中，Trainium几乎位于TPU之上：编译器就是系统，而且在很大程度上是同一个编译器。

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 NKI (Neuron Kernel Interface), a Python, tile-level kernel language that exposes the four engines and the SBUF/PSUM scratchpads directly. It is Trainium's Pallas (or its Triton): the same idea of a tile DSL that drops beneath the whole-program compiler when a kernel's win is in the schedule, not the algebra. Below it, a collective-communication library maps all-reduce and all-to-all onto the CC-Cores and the NeuronLink topology (the NCCL analogue), and NeuronX Distributed provides the sharded-training layer.它与TPU的区别在于“逃生舱”。仅靠XLA并不总能为新型注意力变体或融合MoE调度合成最优方案，因此Neuron发布了NKI（Neuron Kernel Interface），这是一种Python分块级内核语言，直接暴露四个引擎和SBUF/PSUM暂存区。它是Trainium的Pallas（或其Triton）：相同的分块DSL理念，当内核的优势在于调度而非代数时，它会深入到全程序编译器之下。在其下方，集合通信库将All-Reduce和All-to-All映射到CC-Core和NeuronLink拓扑（NCCL的类似物），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.与CUDA（甚至TPU堆栈）的差距在于成熟度，而非设计。NKI、JAX路径和分布式库在2024年底仍处于测试阶段；移植的模型只能在AWS上运行，没有跨供应商回退；vLLM后端落后于上游项目。最明显的标志是锚定租户的工作方式：Anthropic不仅仅通过PyTorch定位Trainium，它还与Annapurna深度合作，编写自己的底层NKI内核，并将修复程序上游合并到Neuron堆栈中。Trainium在最前沿是生产可用的，但在最前沿它是协同工程的，而不是交钥匙的：编译器是继承的且非常出色，但周围的生态系统还很年轻。

Groq LPUGroq LPU

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 pulls it apart: instruction control, the vector ALUs, the matrix units, the memory, and the network each become a functional slice, 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 streams horizontally through the slices like parts down an assembly line, East and West, one register hop per cycle, while 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 Tensor Streaming Processor (TSP), and carried that name until the 2024 rebrand to Language Processing Unit.该领域的其他产品都是基于复制核心构建的：在一个晶圆上平铺SM、TensorCore、CU或数据流核心，并将工作分发给副本。LPU的构建方式恰恰相反。它取一个传统核心并将其拆解：指令控制、向量ALU、矩阵单元、内存和网络各自成为一个功能切片，即一列全高的相同硬件，这些列并排站在晶圆上。切片内同构，芯片间异构。数据不会坐在寄存器文件中等待被发出到单元；它像流水线上的零件一样水平流经切片，向东或向西，每个周期跳一个寄存器，同时VLIW指令从控制切片向北发出以与数据汇合。数据通路中没有任何东西是反应式的：编译器知道每个操作数在每个周期位于何处，硬件只是转动时钟。这种流式传输就是其身份：该设计最初作为张量流处理器（TSP）发布，直到2024年更名为语言处理单元。

The vertical axis is SIMD width. The chip is 320 lanes tall, organised as 20 superlanes 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 stream registers 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.垂直轴是SIMD宽度。芯片有320个通道高，组织为20个16通道的超级通道（第21个是备用，为良率熔断，对软件不可见），每个切片同时作用于所有320个通道。水平轴是时间。每个通道有64个逻辑流寄存器，32个向东流，32个向西流，每个周期每个流向其方向推进一个切片，直到被消耗或从晶圆边缘掉落。切片从经过的流中读取操作数，进行计算，并将结果写回绑定到下一个切片的流中。晶圆在中央向量单元周围镜像为两个半球，因此产生一次的值可以被两侧的切片消耗。

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 MXM: 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 750 INT8 TOPS and 188 FP16 TFLOPS, 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.LPU保持了与其他架构相同的分工，矩阵工作在专用单元上，其余工作在向量引擎上，但都将它们排列为流中的切片。矩阵路径是MXM：四个独立的320×320乘加平面（每半球两个），总共409,600个乘法器，接收INT8或FP16操作数到INT32或FP32累加器。权重安装在平面上（全部在40个周期内），然后激活值流过，乘积进行累加。在900 MHz下，这大约是750 INT8 TOPS和188 FP16 TFLOPS，而且不同寻常的是，该数字没有稀疏性星号：TSP根本拒绝跳过零，因为依赖数据的跳过会使执行时间依赖于数据，而确定性是它绝不会交易的属性。

The vector path is the VXM 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 spatial 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 SXM, 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.向量路径是晶圆中心的VXM：每个通道16个ALU，排列为4×4网格，5,120个32位ALU，运行激活、归一化、量化和残差相加。由于计算是空间的而不是发出到共享单元的，操作数可以在连续周期内穿过VXM ALU链并直接进入MXM平面，而无需触及内存：GPU内核手动构建的算子融合在这里只是切片的物理顺序。第三种切片类型SXM处理直线流无法表达的移动：通道移位、320通道置换、转置和芯片间链路都驻留在这里，因此跨通道重排数据是一等操作，而不是通过SRAM的往返。

Memory

There is no HBM, no DRAM, and no cache. On-chip is the MEM 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.没有HBM，没有DRAM，也没有缓存。片上是MEM切片：88个切片中的230 MB SRAM（每半球44个），每个字节距离计算切片仅一个周期，总带宽约80 TB/s。这就是整个层级结构：单层、扁平、软件寻址，没有任何会引入变量延迟访问的驱逐、预取或一致性机制。

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.其结果是该架构的决定性约束。230 MB装不下模型。FP16下的Llama-2 70B为140 GB，因此必须分片到数百个芯片上，权重分散在整个机架或更多的SRAM中：部署配置约为576个LPU。GPU将模型停在少数封装的HBM中并流过Token，而LPU将模型分散在集群的SRAM中并流过Token。芯片数量由容量决定，而非计算：权重必须装得下。这是Cerebras采取的相同权衡（仅SRAM，无HBM），从相反方向达成：Cerebras保留一个巨大的晶圆并放弃每晶圆容量；Groq保留一个常规尺寸的晶圆并放弃在单个晶圆上装下模型。

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 TruePoint: 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).数值是未走之路。这里其他所有供应商每一代都在将精度减半，从FP16到FP8到FP4，并使用块缩放来换回精度。TSP停留在FP16和INT8，带有FP32累加，从未在硅片上出货FP8或FP4。它唯一的一个数值理念是TruePoint：一个320元素点积融合到一个带有FP32累加的舍入步骤中，因此FP16乘法器阵列在规约上接近FP32精度（Groq报告相对于FP32基准的最大误差约为0.05%）。

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.16位是信念还是从未获得低精度更新的数据通路，很难与第二代芯片从未出货的事实分开。SRAM容量是架构最稀缺的资源，8位权重会将模型所需的芯片减半；一个如此受容量限制的机器有充分理由想要FP8，但没有在硅片上实现。这是悬在Cerebras仅16位数据通路上的同一个开放问题，也是同样的张力：最渴求容量的供应商在最宽的精度上进行计算。

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.其他所有加速器都隐藏延迟；LPU暴露它。ISA携带每条指令的执行延迟，数据通路通过构造是固定延迟的，因此编译器提前计算出每个结果出现的精确周期。硬件中没有任何东西可以干扰该调度：没有缓存缺失，没有仲裁器停顿，没有分支预测错误，没有推测需要撤销。Groq自己的测量是证明：24,240次BERT-Large运行返回在约75 µs的区间内，编译器的预测延迟与测量值相差在2%以内。

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 dataflow, 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.这是TPU的本能（将调度移入编译器，删除猜疑它的硬件）更进一步。TPU编译器调度一个芯片；LPU编译器调度一个系统，因为确定性也跨越了网络。它与Cerebras正好相反，Cerebras的核心是数据流，只要操作数到达就会触发：WSE对数据做出反应，LPU对时钟做出反应。两台机器都删除了调度器；一个用到达替换它，另一个用时钟替换它。

Bets

• Bet 1: Determinism over tolerance. Delete every reactive component (caches, arbiters, predictors, reorder buffers) and let the compiler own every cycle.赌注1：确定性高于容错。删除所有反应式组件（缓存、仲裁器、预测器、重排序缓冲区），让编译器拥有每一个周期。
• Bet 2: Spatial functional slices. 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.赌注2：空间功能切片。将核心拆解为切片并流过操作数，因此融合是平面布局，数据重用存在于导线中，而不是寄存器文件的舞蹈中。
• Bet 3: SRAM is the only memory. 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.赌注3：SRAM是唯一的内存。没有HBM，不计容量成本。用在芯片上容纳模型的能力，换取单周期、固定延迟访问，接受模型必须跨越数百个芯片。
• Bet 4: Schedule the network too. 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.赌注4：调度网络。使芯片成为自己的路由器，并按周期编译通信，因此千芯片集群是一个没有交换机、没有拥塞的确定性程序。
• Bet 5: Sell latency, not throughput. 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.赌注5：销售延迟，而非吞吐量。在Batch 1（GPU最差的机制）下优化每用户每秒Token数，并将该速度作为产品定价，而不是竞争每Token成本。

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 RealScale 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 glueless multiprocessor: no 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.扩展LPU与这里其他任何东西都不同，因为没有单独的扩展结构需要构建：芯片本身就是一个交换机。每个LPU携带多达16个芯片间RealScale链路（卡上暴露11个），并同时作为计算端点和路由器。将芯片直接连在一起，集群就是一个无胶多处理器：没有网卡，没有交换机ASIC，没有架顶式交换机。由于确定性跨越了这些链路，整个集群运行在一个编译时调度表上。

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.节点是8个LPU，完全连接：每个芯片的7个链路将其连接到其他七个，因此节点中的每个芯片距离其他芯片都只有一跳。每个芯片上剩余的四个链路（节点总共32个）捆绑成ISCA论文所称的32端口虚拟路由器，即节点到更大结构的上传链路。没有底板交换机和一致性地址空间；远程操作数不是加载的，而是被调度到达的，由源芯片在编译器选择的周期注入，并在目的地到达的周期被消耗。

Scale-out

Beyond the node, nodes wire into a Dragonfly: 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 software-scheduled: routing and flow control move to compile time, and the paper's framing is blunt, 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 plesiochronous, and the fabric maintains a global consensus time with Hardware-Aligned Counters 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 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.节点之外，节点连入Dragonfly：9个节点组成72芯片机架（第九个是热备，所以64个活跃），拓扑扩展到指定的10,440个芯片，任意两个距离小于六跳。结构是软件调度的：路由和流控制移动到编译时，论文的框架很直白，调度而非路由。没有背压和动态仲裁，因为编译器已经证明接收者准备好了；链路携带前向纠错而不是重传，因为重试会扰乱调度。保持一架独立时钟芯片同步是它自己的问题：链路是准同步的，结构通过每256个周期在生成树上交换的硬件对齐计数器维持全局共识时间，并带有周期性的去偏斜指令，将每个芯片停顿回对齐状态。Groq报告的收益是，8路All-Reduce在大张量上匹配A100/NVSwitch节点，在小张量上击败它，调度结构不需要动态结构那样的握手延迟。

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. SemiAnalysis 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.代价写在内存赌注的物理学中。模型副本不是一个盒子，而是一个机架（或八个）：根据一项分析，约576个芯片上的Llama-2 70B携带了144个主机CPU和144 TB主机RAM，而8-GPU服务器仅需两个CPU。每个芯片下的晶圆很便宜（14纳米GlobalFoundries，据报道低于6千美元，对比H100级部件约1.6万美元），但你需要数百个，在解码期间，它们巨大的计算能力大部分处于闲置状态，而SRAM在工作。SemiAnalysis直言不讳：当你优化延迟时，LPU在每Token物料清单上胜出；一旦你分批（batch），在每美元吞吐量上它比GPU输掉一个数量级。该架构不在成本上竞争。它在速度上竞争。

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 wgmma 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 GroqFlow bring-up path) is small and closed, and GroqFlow was archived in 2025 as the company stopped selling cards and started selling tokens.编程模型是“编译器就是机器”的最纯粹表达。没有内核。你将PyTorch、TensorFlow或ONNX中的模型交给Groq编译器；它降低到一小组张量操作并静态调度每一条指令、每一个流和每一个芯片间传输。没有人写wgmma或手动调整分块，因为没有动态硬件可以调整。Groq的演示是在四天内由不到十人的团队带起LLaMA，而同样的模型在GPU上调整需要数月的手动内核工作。编译器周围的堆栈（分析器、运行时、GroqFlow启动路径）很小且封闭，GroqFlow在2025年随着公司停止卖卡开始卖Token而被归档。

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 Artificial Analysis 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. 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.这次转型是该架构用途的信号。LPU通过构造仅用于推理（Ross的框架是训练是本地游戏，推理是全球游戏），并且在单件事上未被击败：单用户解码延迟。独立测量支持这一说法，Artificial Analysis将Groq列为开放模型上最快的每秒Token提供商之一。它与其余部分匹配得很差：装不进SRAM机架的模型，想要大分批以获得每美元吞吐量的工作负载，或静态调度无法表达的动态控制流。MoE得到服务，但其依赖数据的专家路由与希望提前知道一切的编译器不兼容，Groq很少发布关于它如何调和两者的信息。

The epilogue is that the buyer of all this was NVIDIA. In December 2025 NVIDIA took a non-exclusive license 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 GTC 2026 the technology reappeared as the NVIDIA Groq 3 LPU, a rack of 256 SRAM-only inference chips sitting beside 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.尾声是这一切的买家是NVIDIA。2025年12月，NVIDIA获得了LPU技术的非独家许可，并聘请了Ross和大部分团队。这不是收购：根据NVIDIA自己的10-K，没有产品、客户合同或股权易手，尽管收盘时支付的约130亿美元导致媒体称其为收购。在GTC 2026上，该技术作为NVIDIA Groq 3 LPU重现，这是一个由256个纯SRAM推理芯片组成的机架，坐在Rubin NVL72旁边，在它们之间拆分Transformer：GPU运行注意力，LPU运行前馈和MoE层，由Dynamo编排移交。AI中最确定性的架构最终成为了最可编程架构内部的延迟协处理器。与此同时，GroqCloud仍在原始14纳米硅片上提供Token服务。