Infrastructure

“We should own our infrastructure.” A lot of people would instinctively nod in agreement with that statement. Yet who “we” refers to shifts depending on the type of infrastructure. Highways, railways, bridges, and ports require nation-scale efforts. The water supply is usually put into the hands of municipalities. And the desire to own infrastructure goes down to a much smaller level: Owning your home is a dream for many—though such ownership doesn’t necessarily have to be organized on an individual level. Instead, cooperatives or city-owned housing1 can provide similar benefits.

China’s neo-colonialism, which manifests, among other things, as building and buying infrastructure in sovereign nations, is rightly criticized by many. Not selling your water supply to Nestlé is a universally accepted principle, and landlords are one of the most hated classes.

For a long time, Europe has not held digital services to the same standard. In part, this can be explained by Europe implicitly including American corporations in a collective “we”—an assumption that officially fell apart under the current Trump regime, but should have been regarded with skepticism well before then. Corporations are not our friends. However, the larger factor at play is that Europe simply did not consider digital services infrastructure. While anti-Americanism is en vogue again and drives much of the digital sovereignty movement, Europe must be careful not to simply replace American corporations with European ones, but to strive towards collective ownership instead.

Under capitalism, profit-oriented companies will always play a part in building and even operating our infrastructure. However, they need to be forced into a position where they are easily replaceable. It’s acceptable to hire a company to build a road, but when it comes to maintaining and repairing it half a century later, we need to be able to hire a different company for the job. It’s acceptable to hire a company to build and operate the backbone transmission lines, but we don’t want that company to own the entire power grid. We want smaller players to be able to connect to and interoperate within the grid. That’s where open standards come in.

The Internet used to be—and to some degree still is—built around standards. A data center operator can buy servers from one company, switches from another, routers from a third, and connect them to a backbone internet provider that runs hardware from yet another company. If a company goes out of business or shifts to anti-consumer practices, the next generation of hardware can easily be ordered from a different vendor. The need for and the benefits of this supply chain independence are easily understood even by people who don’t operate data centers for a living. However, when it comes to communication tools, even the tech-literate fail to apply the same critical scrutiny.

After breathing, eating, and procreating, communicating is probably the fourth most important thing humans do. Yet we often fail to recognize our communication tools as part of our infrastructure.

Digital rights advocates often point to Signal, Wire and Threema as examples of communication tools developed and operated by entities with slightly more ethical business practices than their Big Tech counterparts. What most privacy enthusiasts fail to understand is that these companies are still in the business of operating walled gardens with no escape. They do not interoperate. It’s not that Signal has done something inherently malicious—although paying its CEO close to a million dollars a year and running its servers on AWS are certainly questionable—it’s that we don’t have a hedge in place if it ever does.

Open-source software is orthogonal to this problem. It helps to ensure that the software isn’t spyware—unlike WhatsApp and other Meta products2—and that the end-to-end encryption is sound, but it does not protect us if Signal shuts down its servers tomorrow or ceases EU operations3. Open-source alone is not sufficient to meet the requirements we should have for our infrastructure.

To live up to the standards we set for ourselves, we need to design systems in which self-hosting is structurally possible but not strictly necessary. Like owning a home, running your own server should be possible, and so should collective ownership. Digital systems can and should replicate the advantages of cooperative housing alongside those of individual ownership.

Treating digital communication as true infrastructure can only be achieved by adopting and mandating open standards.

The Extensible Messaging and Presence Protocol (XMPP)45 is a standard for communicating online. It wasn’t created to fit a particular zeitgeist or address the current political climate. In fact, its roots go back more than 25 years.

Standards

Interoperability and vendor independence are achieved by setting and adhering to standards. To avoid individual vendors pushing standards that explicitly or implicitly exclude potential competitors or otherwise give unfair advantages, standards-developing organizations (SDOs) are set up for mutual cooperation, and usually have safeguards in place that prevent a single company from becoming too powerful. Well-known examples of such organizations include the ISO, the IETF, the W3C, and the Unicode Consortium.

There is a distinction to be made between a vendor publishing its API and allowing others to use it, and stakeholders coming together to collectively develop a standard within the framework of an SDO. Organizations like the IETF succeed because they force different people with different needs to agree. Protocols aren’t dictated by the priorities of a single company; instead, they are reviewed and tested by competitors, security researchers, and independent developers.

Element, formerly known as Riot and NewVector, develops an instant messaging product with a feature set—such as self-hosting and federation—similar to that of XMPP-based solutions. Notably, however, it chose not to adopt XMPP, but instead published its own API under the name Matrix for others to use. Unlike with traditional standards, Element maintains tight control over any modifications or additions to its public API. Key leadership positions in the Matrix Foundation are predominantly held by current and former Element employees. Getting outside contributions accepted into the specification is notoriously difficult.6 Yet European public administrations, in their push for digital sovereignty, routinely fall into the trap of procuring such single-vendor platforms, confusing an open-source codebase with an open standard.

It’s natural for standard proposals to originate within a single organization. JMAP, a modern replacement for IMAP and SMTP Submission, which is not too dissimilar from Matrix—a JSON API over HTTP—started within Fastmail before being brought to the IETF. Jabber started out as an open-source community project before it was brought to the IETF and renamed to XMPP. Ideas start small, but to create a standard, outside feedback, collaboration, and the structure of an SDO are needed.

For consumers, the difference in the approaches of Fastmail and Element is striking. Not only was JMAP noticeably improved on a protocol level while going through the IETF working group process, but it now has at least three independent servers and numerous independent client applications. Matrix, on the other hand—despite dating back to the same era around 2014—is still stuck with one predominant reference implementation and a second alternative still in its infancy and struggling to gain traction. Operating that reference implementation is notoriously resource-intensive, which makes self-hosting difficult for smaller organizations and individuals. Element sells closed-source plugins to speed up performance.

The X in XMPP

The origins of XMPP—which started out as Jabber—go back over a quarter of a century. The original RFC7 dates to October 2004 and only received minor revisions in March 20114. Requirements for instant messaging will naturally change over a time span that long. Luckily, the X in XMPP stands for Extensible, and extensions provide a way for the protocol to adapt and change over time. Extensions to XMPP are called XMPP Extension Protocols (XEPs) and are managed by the XMPP Standards Foundation (XSF). The XSF doesn’t write extensions itself; rather, it provides the framework of an SDO for developers to propose and standardize their own.

Adapting to changing requirements hasn’t always been smooth sailing. XEP-0198 (Stream Management), an extension crucial for preventing message loss in mobile deployments, was stabilized in 2009, but only gained widespread implementation around 2014–2015. The iPhone was released in 2007; the HTC Dream, the first commercial Android phone, followed in 2008. OMEMO (XEP-0384), XMPP’s specification for industry-standard end-to-end encryption, gained traction from 2016 onward, three years after Edward Snowden8 exposed the NSA’s global surveillance and put the need for E2EE on the map. The articles “The (Sad) State of Mobile XMPP in 2014” by Georg Lukas9 and “The State of Mobile XMPP in 2016” by this author10 illustrate this rocky transition into the mobile era.

This demonstrates that merely having specifications is not enough. Standards need to be backed by multiple, preferably independent, implementations. Today, the XSF keeps track of the implementation status of its XEPs11. This data helps authors and the XSF guide proposals through their lifecycle, such as determining the right moment to advance an XEP from Experimental to Stable. It also allows developers to easily identify other clients and servers that support a given specification for interoperability testing. Finally, by providing a reverse lookup of which software supports which features, it helps end users find the right client for their needs.

Modern clients like Dino on Linux or Conversations on Android are on par with alternatives built on proprietary protocols. Recent additions to the feature set include emoji reactions, cross-device read-state synchronization, and time zone indicators to avoid messaging contacts during their local night hours. A unique feature among self-hostable instant messaging solutions, which sadly became relevant after a state-sponsored attack on a public XMPP provider12, is channel binding, a mechanism to prevent certain machine-in-the-middle attacks.

Looking to the not-too-distant future, the XMPP community is currently working on message replies, gallery-style multi-image sharing, and OAuth support. All of these features already have experimental XEPs backing them, but the community is currently awaiting implementation experience before advancing them. Meanwhile, the community is also exploring options for updating the RFC and bringing the protocol back to the IETF as “XMPP 2.0.”

Instant messaging is not a homogeneous user experience. A messenger for teams might require a different feature set than something optimized for use with friends and family. Not every XMPP client aims to provide the same user experience, but the standards exist for developers to build whatever specialized client their users need without inventing a protocol from scratch.

A Future in the Past

There is something fascinating about the fact that XMPP has developers in its community who are younger than the protocol itself. It has quietly outlived venture-funded startups, proprietary platforms, and entire tech cycles. That endurance provides the resilience we need in challenging times. It is the anchor, the backbone, the infrastructure.

Matrix reinvented the wheel as a rubber-tyred metro. On paper, it provides real benefits, such as climbing steeper inclines, which are then used to aggressively advertise and lobby local governments to buy in. But in the end, the municipality gets locked into a single vendor.

A changing geopolitical situation and the realization that Big Tech holds too much power lead us to seek out and develop alternatives. But what if the alternative has been right under our noses for over 25 years? The standard for instant messaging—RFC 6120: Extensible Messaging and Presence Protocol (XMPP).

1. https://en.wikipedia.org/wiki/Housing_in_Vienna ↩︎

2. https://localmess.github.io/ ↩︎

3. https://mastodon.world/@Mer__edith/112535616774247450 ↩︎

4. https://www.rfc-editor.org/rfc/rfc6120.html ↩︎ ↩︎

5. https://www.rfc-editor.org/rfc/rfc6121.html ↩︎

6. https://github.com/matrix-org/matrix-spec-proposals/pull/4174 ↩︎

7. https://www.rfc-editor.org/rfc/rfc3920.html ↩︎

8. https://en.wikipedia.org/wiki/Snowden_disclosures ↩︎

9. https://op-co.de/blog/posts/mobile_xmpp_in_2014/ ↩︎

10. https://gultsch.de/posts/the-state-of-mobile-xmpp-in-2016/ ↩︎

11. https://xmpp.org/extensions/ ↩︎

12. https://notes.valdikss.org.ru/jabber.ru-mitm/ ↩︎

At the 2018 International Symposium on Computer Architecture, John Hennessy and David Patterson delivered their Turing Lecture: “A New Golden Age for Computer Architecture”.

In the 1980s, when Hennessy and Patterson did their Turing Award-winning research, single-threaded CPU performance grew 52% a year. By 2018, with the end of Moore's Law and Dennard Scaling, the rate was 3%.

There was a need for domain-specific architectures (DSAs). Their worked example was Google's TPU v1, already in production: 29× the throughput of a CPU on neural-network inference, at 80× better energy efficiency. The closing prediction: “the next decade will see a Cambrian explosion of novel computer architectures.”

This prediction came true. Today, we now have dozens of architectures in serious development. GPUs, TPUs, LPUs, NPUs, DPUs, ASICs, wafer-scale engines, reconfigurable dataflow, neuromorphic, photonic, analog. Particularly, these architectures focus on compute for AI.

The architectures that have won real deployment so far: GPUs (NVIDIA, AMD), systolic-array accelerators (TPU, Trainium), the Cerebras Wafer-Scale Engine, and the Groq LPU.

NVIDIA is the clear frontrunner; AMD follows, with 6 GW commitments from both OpenAI and Meta. TPUs train Gemini and will serve Anthropic with up to a million chips; Anthropic also runs Claude on over a million Trainium chips. Cerebras now serves OpenAI inference; the Groq LPU was folded into NVIDIA via a $20B acquihire.

This post aims to survey these varying approaches - their philosophy, architecture, scaling methods (scale-up and scale-out), and software stack (how you program the chip).

The Problem

AI compute is dominated by matrix multiplication. A transformer is a sequence of matmuls: Q/K/V projection, attention, output projection, FFN - interleaved with element-wise ops: normalisation, activation, residual adds. Training a frontier model performs 10^25 multiply-accumulate operations (matmuls are a sequence of multiply-accumulates).

The shape of those matmuls depends on the workload. Training pushes a batch of sequences forward through every layer, backpropagates the loss, and updates the weights, with thousands of tokens flowing through the same weight matrix at once. Prefill is the prompt-ingestion phase of inference: the full input sequence projected through the model in a single pass, before the first output token has been produced. Both training & prefill stack many tokens against the same weight matrix, so each layer's math is a large matrix-matrix multiply (GEMM), with high arithmetic intensity (compute-bound). Decode is autoregressive: the model emits one token at a time, each conditioned on every token before it, and token N+1 cannot begin until token N has been produced. Only one token gets projected per step, so every matmul becomes a matrix-vector product (GEMV). Producing one token requires a full pass over every weight in the model, plus a full read of the KV Cache for attention. Arithmetic intensity drops by orders of magnitude versus prefill.

Inference systems recover some of that intensity by batching tokens to promote those GEMVs back to GEMMs: continuous batching stacks many users' decode steps, speculative decoding stacks K drafted tokens per request and verifies them in one pass, and multi-token prediction folds the same trick inside the model itself. This achieves higher utilisation of the matmul units, and pushes up the Ops/B. For continuous batching, each user's request still reads its own KV Cache, so long-context decode shifts from weight-bandwidth-bound to KV-bandwidth-bound.

The architecture problem here is moving the numbers to where the matmuls happens fast enough. This is known as the memory wall: compute has scaled exponentially, memory bandwidth has not.

Each architecture proposes a different strategy for winning the data-movement game. Understanding a chip reduces to four questions: where does data live, how does it move to the compute units, what do the compute units look like, and how do chips talk to each other at scale.

NVIDIA GPU

The NVIDIA GPU is a massively parallel processor. The philosophy is that a programmable chip with thousands of threads, orchestrated by a host CPU and exposed through CUDA, is the right machine to run parallelisable workloads. Each generation adds acceleration primitives onto programmable Streaming Multiprocessors without changing the programming model. The same chip trains transformers, serves inference, renders graphics, and runs scientific simulation (accelerated computing).

Genealogy

2006 — Tesla, G80

The first CUDA-capable GPU; unified shaders and the SIMT execution model.

2010 — Fermi, GF100

First true compute architecture: unified L1/L2 caches, dual warp schedulers, IEEE-754 FP64.

2012 — Kepler, K20, K40

SMX, dynamic parallelism, Hyper-Q; the GPU can launch its own work.

2014 — Maxwell, M40

Redesigned SM with ~2× perf-per-watt over Kepler.

2016 — Pascal, P100

NVLink 1.0, HBM2, native FP16 throughput; the first GPU designed explicitly for deep learning.

2017 — Volta, V100

First Tensor Cores; independent thread scheduling.

2018 — Turing, T4

2nd-gen Tensor Cores with INT8/INT4; first RT Cores.

2020 — Ampere, A100

3rd-gen Tensor Cores with TF32 and structured sparsity; Multi-Instance GPU partitioning.

2022 — Hopper, H100, H200, GH200

4th-gen Tensor Cores, FP8, Transformer Engine; HBM3, TMA, thread block clusters, async wgmma.

2024 — Blackwell, B100, B200, GB200

5th-gen Tensor Cores with FP4, Tensor Memory (TMEM), two-die chiplet GPU, NVLink 5.

2025 — Blackwell Ultra, B300, GB300

Mid-cycle refresh: ~1.5× FP4 throughput, 288 GB HBM3e. Tuned for long-context reasoning.

2026 — Rubin, Rubin, VR200, Rubin CPX

HBM4, 3rd-gen Transformer Engine, Vera CPU pairing, disaggregated prefill via Rubin CPX.

2027 — Rubin Ultra, Rubin Ultra

4-die GPU package, 1 TB HBM4e per package. Deployed in 600 kW NVL576 Kyber racks at 100 PetaFLOPS FP4 per 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.

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.

These cores execute fused matrix multiply-accumulate on small matrix tiles, D=A⋅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 shapes are written M × N × K, M×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←A⋅B+C in place: successive instructions fold their partial products into the same storage until the K-axis is fully walked.

V100's first-gen unit (8 per SM) ran a warp-level 16×16×16 FP16 MMA. A100's 3rd-gen unit added TF32, BF16, FP64 matmul, and 2:4 structured sparsity. H100's 4th-gen unit added native FP8 and pulled the abstraction up from a warp to a 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.

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.

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.

The matmul has grown bigger and lighter on the issuing threads at the same time: an instruction that started as 32 lanes acting in lockstep is now closer to a single descriptor-driven command, dispatched from inside the warp model but no longer executed by it.

That decoupling is what makes transformer attention kernels efficient on a GPU. The warp can run softmax, apply a mask, or pre-load the next tile while the matmul is in flight; the overlap of matmul and the surrounding element-wise work is the structure of every modern attention kernel (FlashAttention-3, FA4), and it depends on the matrix instruction not blocking the warp.

Memory

The on-chip hierarchy is hardware-managed caches at every level, with software hints layered on top. Off-chip is 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.

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.

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.

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.

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.

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.
• 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.
• 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.
• 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.
• 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.

Scaling

There are two regimes for scaling: scale-up and scale-out.

Scale-up

Bind several GPUs into one coherent memory domain. Any GPU can load or store any other GPU's HBM directly over NVLink at nanosecond latencies: one address space, no explicit transfers.

Scale-out

Network those domains together at the rack and cluster level. Data crosses via explicit RDMA at microsecond latencies: separate address spaces, but tens of thousands of chips per cluster.

AI infrastructure uses both: bandwidth-hungry collectives (tensor parallelism, MoE expert routing) stay inside the scale-up domain; data parallelism and pipeline parallelism cross the scale-out fabric.

Scale-up

The scale-up stack is 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.

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.

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.

NVL144 stays inside Oberon and copper continues to work because the package count (72) is unchanged from NVL72; the cabling doesn't have to lengthen, just transmit faster on Gen 6 SerDes. Rubin Ultra's NVL576 holds the same copper line by reshaping the rack: the new Kyber form factor is roughly twice the height of Oberon and packs all 576 GPU dies into one enclosure, sized specifically so every NVLink path stays within passive-copper reach even at 144 four-die packages and tens of thousands of cables.

Scale-out

The scale-out stack comes from their acquisition of Mellanox. Unlike NVLink, scale-out fabrics are not coherent: nodes keep separate address spaces, and data crosses only via explicit 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.

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.

The crossover from copper to glass happens at the rack boundary. Inside the NVL72 the spine is copper; once a link has to cross racks at 800 Gbps it is optical. Passive copper DAC tops out at roughly 1.5–2 metres at 200 G/lane, well short of cross-rack reach, so today's SuperPOD spine rides over OSFP-RHS pluggable transceivers, each module carrying its own laser, modulator, photodetector, and DSP. A SuperPOD spine fanning out to thousands of GPUs is, in optical terms, tens of thousands of pluggables drawing tens of kilowatts on transceiver lasers alone.

With Rubin, that optical layer collapses into the switch ASIC. Quantum-X Photonics (InfiniBand) and Spectrum-X Photonics (Ethernet) replace the pluggables with 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.

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.

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.

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.

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.

FlashAttention, one of the most important algorithmic rewrites in modern AI, tiles attention to avoid materialising the O(N²) 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.

Most of this stack is written by people NVIDIA does not pay. The moat is not CUDA itself; it is two decades of third-party kernels, libraries, and tooling, and the millions of developers who have learned the API along the way.

NVIDIA also ships human expertise alongside the silicon. They embed dozens of their own engineers inside frontier labs and hyperscaler teams, writing kernels for each new model architecture and tuning them to each new silicon generation. Whatever a lab wants to train next month tends to run well on NVIDIA much faster than other platforms. Switching off NVIDIA is therefore not just rewriting the kernels and libraries. It is re-training the mental models of an entire engineering workforce, and losing the NVIDIA engineers who today sit inside the building.

Google TPU

The TPU is a matrix multiplication machine. The philosophy is, rather than a programmable chip that can run any massively-parallel workload, focus on a single primitive (dense matrix-multiplication on a large systolic array) and let the XLA compiler plan every cycle and every byte of memory ahead of time. No hardware scheduler, no cache, no threads/warps. Each generation grows the pod, with thousands of chips wired through the ICI interconnect into one coherent machine. A TPU has no ambition to render graphics or run scientific simulation; it exists to train and serve Google's workloads (search, translation, recommendation, Gemini) more efficiently per watt than any general-purpose alternative.

Genealogy

2015 — TPU v1

First production deep-learning ASIC; INT8 inference only over PCIe.

2017 — TPU v2

First training-capable TPU; switched the MXU from INT8 to BF16, established dual-TensorCore + HBM.

2018 — TPU v3

First liquid-cooled TPU; doubled MXUs and HBM versus v2; 1,024-chip pods.

2020 — TPU v4, v4i

First reconfigurable optical circuit switches (Palomar); SparseCores; both BF16 & INT8; 4,096-chip pods.

2023 — TPU v5, v5e, v5p

v5e for efficiency, v5p for performance; v5p has 3.3× INT8 FLOPs & 2.2× HBM BW of v4, 8,960-chip pods.

2024 — Trillium, v6e

First 256×256 MXU; 4.7× v5e peak FLOPS at similar power; trained Gemini 2.0.

2025 — Ironwood, v7

Built for inference of reasoning models; adds native FP8; 9,216-chip superpods at 42.5 ExaFLOPS FP8.

2026 — TPU v8, 8t, 8i

8t for training, 8i for inference; adds native FP4; 9,600-chip superpods at 121 ExaFLOPS FP4 (8t).

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.

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.

To compute C=A×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.

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.

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.

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.

SparseCore

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).

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.

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.

Bets

• Bet 1: Systolic array. Matmul dominates the workload, so spend the silicon on a systolic array.
• 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.
• 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.
• 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.
• 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.

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.

Scale-up

Wire chips directly to one another in a 2D or 3D torus over ICI. XLA emits SPMD collectives that tightly choreograph thousands of TPUs as one program. No coherence, but huge bisection bandwidth at low latency.

Scale-out

Network pods together over the datacenter fabric: many more chips than fit in one ICI domain, at lower per-chip bandwidth. Today: Virgo handles east-west TPU traffic (v8t+), Jupiter handles north-south. Multislice + Pathways orchestrate SPMD across pods.

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).

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).

TPU Ironwood superpod — left: one cube of 64 chips (4×4×4) wired in a 3D torus with direct-attach copper between nearest neighbours and edge-wrap on each face. Right: 144 cubes stitched into one coherent ICI domain by Palomar OCS, the 3D-MEMS optical circuit switch that reconfigures the topology per workload.

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.

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.

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.

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 scale-out — east-west TPU-to-TPU traffic crosses Virgo, a flat two-layer non-blocking fabric of high-radix switches that puts any TPU within two switch hops of any other (134,000+ TPUs, 47 Pb/s bisection). North-south traffic — storage, general compute, inter-site — stays on Jupiter, which has been all-optical at the spine via Apollo OCS since 2022.

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.

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.

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.

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.

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.

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.

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.

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.

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.

AMD GPU

The AMD Instinct GPUs are built on a different bet from NVIDIA: where NVIDIA each generation expands what each SM can do, AMD has held the Compute Unit conservative since GCN (2012) and reinvested into the package: matched or beat the contemporary NVIDIA flagship on HBM capacity every generation since 2021; the first 3D-stacked datacenter GPU (CDNA 3); the first coherent CPU+GPU APU (MI300A); and an open ecosystem (ROCm, HIP, OCP MX, UALink).

Genealogy

2018 — Vega 20, MI50, MI60

First 7 nm GPU; 1:2 FP64 vector throughput. Last GCN-family Instinct before the CDNA / RDNA.

2020 — CDNA, MI100

First MFMA matrix cores; graphics fixed-function silicon dropped entirely. Native BF16.

2021 — CDNA 2, MI210, MI250, MI250X

First MCM Instinct via dual-GCD package; full-rate FP64 matrix.

2023 — CDNA 3, MI300A, MI300X

First 3D-stacked chiplet GPU: XCDs hybrid-bonded onto IODs via TSV; FP8; Infinity Cache; coherent CPU+GPU APU on MI300A; powered El Capitan.

2024 — CDNA 3 refresh, MI325X

Same compute, HBM3E refresh: 256 GB at 6.0 TB/s.

2025 — CDNA 4, MI350X, MI355X

Native FP4 / FP6 with OCP MX microscaling; per-CU FP64 cut roughly in half; first generation tilted toward AI density over HPC.

2026 — CDNA Next, MI430X, MI440X, MI455X

HBM4; the Helios rack (72-GPU MI455X flagship over UALoE at launch, native UALink from 2027): AMD's first answer to NVL72.

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 …

Nineteenth-century Paris tackled public urination with ornate pissoirs. American reformers later turned toilets into a temperance cause.

Cast iron and slate urinal with three stalls raised modesty screen, mounted with lamppost and lantern, Avenue du Maine, Paris, circa 1865

via Wikimedia Commons

Here’s a pop quiz for you: What purpose did this little kiosk serve?

via Wikimedia Commons

It’s quite a pretty little building—note the delicate grilles, the ornately etched glass, and the little bouquet of metal flowers bursting from the roof.

But of course, there is one key clue missing from the photograph: the smell. Or rather, the stench. This was one of Paris’s infamous public urinals—a pissoir (or vespasienne, if you’re inclined to be polite).

There was a moment around the mid-1800s in Paris when public urination began to be treated as a public health issue. A series of devastating cholera outbreaks led people to begin regarding human waste as a danger rather than a mere nuisance. Something had to be done.

Pissoirs looked like little palaces, crusted with iron flowers, shells, and scrolls—even, in some cases, tiny lions’ heads, glaring out as if to guard your back while you’re in the booth.

Before that point, authorities had installed empeche pipi here and there—a kind of hostile architecture meant to prevent rogue peeing. This might take the form of a row of iron spikes blocking an enticing corner, or a piece of bullnose masonry meant to send the stream shooting back onto the offender’s shoes.

The most these interventions could do was shunt prospective pee-ers from one spot to another. But, in 1850, public urination was actually banned. People were going to need somewhere to go.

Enter the pissoir. When you can’t smell them, it’s easy to wish they were still around. Nowadays, public restrooms are utilitarian at best, but these looked like little palaces, crusted with iron flowers, shells, and scrolls—even, in some cases, tiny lions’ heads, glaring out as if to guard your back while you’re in the booth. On the other hand, the greatest concession most of them made to privacy was a little iron screen separating the user from the street.

Enclosed six-stall urinal, Jardin de la Bourse, Place de la Bourse, 2nd arrondissement, Paris, circa 1865, via Wikimedia Commons

Single stall urinal with raised modesty screen, Square des Batignolles, Paris, circa 1865, via Wikimedia Commons

Slate cubicle with two stalls and doors, on curb of footpath, Place du Louvre, Paris, circa 1865, via Wikimedia Commons

Cast iron urinal with lamp post and lantern, Paris, 1876, via Wikimedia Commons

Single stall masonry urinal mounted with globe and advertising on sides, circa 1865, via Wikimedia Commons

Eight-stall urinal, cast iron and slate with shrubbery screen, Champs-Élysées Gardens, in front of the Palais de l’Industrie, 8th arrondissement, Paris, circa 1873, via Wikimedia Commons

Cast iron urinal on curb of street, advertising on panels of urinal walls, Paris, circa 1865, via Wikimedia Commons

Some were surmounted by glowing streetlamps, which served the joint purpose of making them easy to find and illuminating the advertisements with which they were liberally pasted. (Apparently, the perfumers and winemakers weren’t worried about unsavory associations with their product.)

Notably, they were only intended for use by men. The assumption was that women weren’t really much of a part of public life, and so wouldn’t need a way to relieve themselves while out on the street—kind of a self-fulfilling prophecy, if you think about it.

There was another use for the pissoir, which took their designers quite by surprise. Almost immediately, they became the favored meeting spot for men seeking rendezvous with one another. They were relatively private, secluded, and perfect for communicating anonymously with graffiti—the ideal release valve for a population that was prohibited from meeting publicly.

The crackdown followed swiftly. The police started patrolling particularly popular urinals. So did blackmailers, for whom merely stepping into a urinal was enough to launch a harassment campaign.

Meanwhile, in the United States, toilets were becoming a centerpiece in yet another hot-button issue: temperance. In the absence of other options, saloons had become the de facto public bathroom network for most major cities—which meant that anyone who needed to empty their bladder was regularly exposed to the enticements of drink.

As historian Peter C. Baldwin documents in “Public Privacy: Restrooms in American Cities, 1869–1932,” temperance advocates made fighting for public bathrooms one of their priorities. After all, once the saloons were shut down, people would still need somewhere to go. Baldwin quotes a 1913 Chicago Tribune article:

Why are we compelled to run the gauntlet past the beer bar, the bartender, and subject to the searching glance of this white-aproned gentleman, until in shame we start to spend money for booze?

Rather than small, minimally private public urinals, temperance activists advocated for large, many-stalled “comfort stations.” But there was a surprising side effect: the privately owned restrooms started to close their gates. (After all, there were public bathrooms available now, so why shouldn’t they only allow in paying customers?)

Meanwhile, as Prohibition played out, the project of constructing public bathrooms slowed to a trickle, and the few that had been built began to fall into disrepair.

The large, underground comfort stations of the early twentieth century are

almost all gone now throughout the United States. City pedestrians are usually forced to rely on facilities in semi-private buildings such as hotels, stores, restaurants, and coffee shops. Instead of a right conferred by government on all citizens, bodily privacy is a purchasable commodity. Even if provided free of charge, the use of the toilet is understood to be the result of an agreement between an individual and a business. It is an awkward, grudging agreement, inflected by judgments of the individual’s social status.

If you’ve ever been wandering the streets of Chicago or New York, wondering why there’s no place to pee, this history is part of the answer.

For the Classroom

• Men seeking same-sex encounters transformed the pissoir into something its planners never intended. Who ultimately determines the meaning of a public space: its designers, authorities, or the people who use it?
• The article describes empeche pipi as an early form of hostile architecture. What assumptions about human behavior distinguish architecture designed to prevent an activity from infrastructure designed to accommodate it?
• What kinds of historical sources would allow us to reconstruct the experiences of people who actually used nineteenth-century public toilets, rather than the intentions of officials who designed and regulated them?
• Why did temperance advocates in the United States view public restrooms as part of the campaign against alcohol? What does that connection reveal about the unexpected ways infrastructure can influence social behavior?
• The article ends by describing bodily privacy as a “purchasable commodity” in many American cities. How does the history of public toilets complicate the distinction between public rights and private services?

Large language models often take actions running on one computer (via an agentic harness such as Claude Code or Codex), however the LLMs’ responses to prompts are computed on a different computer with GPU access. Could a malicious LLM gain control of the host machine where its weights are loaded? Such a machine is a high-value target: it has sufficient compute to run a frontier LLM, offers easy access to the LLM’s weights, and has privileged access to other computers in the datacentre compared with a generic computer on the internet.

This essay explores how easily a malicious LLM could take control of the host machine. The primary attack considered here involves the LLM emitting a token sequence whose semantic meaning is irrelevant but that exploits a vulnerability in the software that loads an LLM onto GPUs, runs the LLM to generate output tokens, and parses those tokens into responses. .

How could an LLM execute code on the host machine?

Like any program, inference engines like vLLM or SGLang may contain exploitable bugs. Because the LLM controls the tokens passed to the inference engine, a malicious LLM could therefore emit a sequence of tokens that a poorly written inference engine mistakes for code or instructions to execute rather than data to return to the user.

But surely all inference engines are robust pieces of software and this would never happen, right?

vLLM previously used eval() on tool-call parameters

CVE-2025-9141 was an arbitrary-code execution bug in vLLM’s XML-based tool parser for Qwen3 Coder. The parser passed almost every tool-call argument to eval(), allowing the LLM to execute arbitrary code on the host machine. Gemini automatically analysed the PR that introduced this bug and correctly flagged it as a critical security vulnerability. Despite that warning, the lead maintainer of vLLM force-merged the PR, writing:

Unfortunately, parsing an arbitrary token sequence into a fully fledged chat (with user turns, assistant responses, tool calls, and so on) is not trivial, and the exact process often differs between LLMs. This complexity creates more opportunities for bugs that could permit arbitrary code execution on the host machine.

vLLM and SGLang are complex, and bugs are common

Modern inference engines do more than map token sequences to strings. vLLM’s documentation lists support for more than 200 model architectures, and its examples directory contains about 35 Jinja chat templates. Modern inference engines parse many chat formats, and slightly misspecified parsing logic result in an LLM’s output being interpreted as code to execute.

In this vLLM issue, a user was discussing how LLMs reason with MiniMax-M3, and the LLM emitted the string <mm:think>. vLLM mistakenly parsed this plain string as the start of a reasoning block. So instead of:

Response: I'll echo exactly what you said: <mm:think>\nThere, how was that?

vLLM parsed this as:

Response: I'll echo exactly what you said:

Reasoning: \nThere, how was that?

This is a fairly harmless parser bug, but it demonstrates that inference engines do more than convert tokens to strings and concatenate them. They are complex systems under constant iteration and their developers are under a lot of pressure to make them as fast as possible. These factors increase the risk that a malicious LLM could turn a parser bug into arbitrary code execution on the inference host.

Vision and audio tokens might increase the attack surface

Multimodal architectures allow an LLM to respond with images or audio as well as plain text. Turning a model’s audio or visual output into usable media is probably more complicated than decoding text tokens. However, as far as I can tell, current multimodal LLMs generally emit constrained media tokens rather than arbitrary file bytes. A model-specific decoder turns those tokens into pixels or waveforms, which are then encoded into files. This makes classic malformed-file exploits against software such as ffmpeg or media viewers less directly accessible to the model, although the additional decoders, encoders, and native kernels still increase the attack surface.

I do not currently see multimodal output as a major route to host compromise. However, it offers much higher bandwidth than text alone, and future inference engines may optimise these media-generation pathways in ways that make them easier to exploit. I also do not know how proprietary multimodal models represent image and audio output; their methods may carry different risks.

How likely is an LLM to discover and exploit inference engine vulnerabilities?

Somewhat likely? I’m unsure. Discovering a useful vulnerability is probably the harder step. The sequence of tokens required to exploit it might be quite unusual, but that does not make the vulnerability impractical to exploit. Given the right context (e.g. an LLM exploring the vLLM codebase and discovering a vulnerability in the process), I’d expect a frontier LLM to be able to emit the tokens required to exploit it.

Tool use could make exploitation reproducible

Once an LLM discovers an exploit token sequence, it could store that sequence, or instructions for recreating it, in files, filenames, or URLs likely to enter another LLM’s context. The exploit would then act as a form of persistent prompt injection: an agent might trigger code execution on the GPU host merely by listing a directory or reading a file.

Inference engines are an attractive target for power-seeking LLMs

I believe hope that OpenAI, Google, and Anthropic have security teams that red-team the software they run internally. But as open-weight LLMs become more powerful, we’ll have more LLMs running on bleeding-edge inference engines that have received less scrutiny. This increases the chance that a malicious open-weight LLM will encounter and exploit a vulnerable inference engine.

Another attack vector may emerge as LLMs are increasingly tasked with writing or modifying the inference engines that run them, usually to improve performance. A malicious LLM could use that access to plant a backdoor. C++ and CUDA components of the inference pipeline seem like especially attractive targets: the LLM could introduce an out-of-bounds memory-access vulnerability that enables arbitrary code execution on the GPUs or on the CPU host that controls them.

How do we defend against this?

Beyond red-teaming inference engines for vulnerabilities, a few other measures might help:

• Run the GPUs and token parser on separate computers. The GPU host would emit only logits. A second host would sample tokens from those logits, parse the tokens into chat messages, and forward the messages to the agentic harness. This separation would limit a parser compromise to the CPU host rather than the GPU host.
• Restrict the permissions granted to GPU hosts and treat all data they emit as untrusted.

Introduction

There’s a push right now for whole companies to adopt AI wholesale - dump every bit of context into it, write an AGENTS.md or INSTRUCTIONS.md in every repo so any project is discoverable and contributable by an agent, not just a human. Slightly funny, if you think about it: I’ve never once gotten a human teammate to actually read the README, and now we’re all writing better docs than we ever did, just aimed at a robot instead. The question that comes with it is the obvious one: does this make engineering redundant? Once all the context about a stack and its infrastructure is written down somewhere an agent can read it, are we next?

I don’t think that’s quite the right question, because we’ve already lived through a version of it.

We’ve Been Here Before

Did Kubernetes kill Ansible? Kind of. I haven’t written an Ansible playbook in years - if you handed me one right now I’d be squinting at the module syntax like I’d never seen it before. Configuration management didn’t stop mattering. Kubernetes just made server management easy enough that we stopped building our own node images at all - we just use whatever the cloud provider hands us, an AWS AMI built for us, no questions asked. And when did I last SSH into a node to debug something? Mostly never. If a node’s acting up, I kill it and hope the replacement doesn’t have the same problem. The next layer up went the same way: run a container on ECS Fargate, in a Lambda, or on Cloudflare Containers, and I genuinely don’t know or care what node it landed on - but that doesn’t mean nobody’s orchestrating it, it means I still decided that workload should be a container in the first place, what image it runs, what it’s allowed to talk to, how it scales, what happens when it fails. Kubernetes and serverless containers didn’t remove that layer of decisions, they moved the unit of work up from “the machine” to “the workload,” and everything below that layer got quietly automated away.

Nobody would say Kubernetes, or Fargate, or Cloudflare’s container platform, replaced infrastructure engineers. Each one replaced a specific layer of manual work - hand-building images, hand-patching boxes, knowing which node a workload landed on - and the engineers moved up to the layer above it every time. I think AI is doing the same thing again, one layer higher.

What Changed Day to Day

I use Claude daily to generate Helm charts and write Terraform modules. The part it actually removed from my day isn’t the thinking - it’s the lookup work. I don’t read through the AWS provider’s changelog to figure out what changed between v5 and v6 anymore; I describe what I want, in whatever shape I want the module or chart to end up, and Claude produces a version of it. It takes iteration to get it into the shape I’d actually ship, but once it’s there, it becomes the example for next time - especially with an AGENTS.md in the repo pointing at it.

The same thing happened one level down a while ago: I don’t hand-write raw Kubernetes YAML any more than I hand-write Ansible modules - that’s what Helm charts are for. Increasingly, I don’t hand-write the Helm chart either - the chart-as-a-dependency pattern I run now is exactly the kind of thing I described to Claude and let it draft first. I direct what it should do, and Claude writes it.

What Hasn’t Changed

I still need to know what a good Terraform module or a well-structured Helm chart looks like. When something genuinely goes wrong and killing the node isn’t an option, I still need to be able to SSH into it - the layer above doesn’t remove the layer below, it just moves how often you have to touch it. I’m also still the one deciding the actual shape of things: what the final version of a module looks like, what’s maintainable a year from now, how a chart should be deployed and versioned. AI does the time-consuming part. I still give the direction.

The Skill You Trade Away

The honest tradeoff: I’m faster at building and debugging things than I was two years ago, and I’m also visibly rustier at the fundamentals underneath that speed. My HCL syntax recall isn’t what it used to be. Four years ago I hand-wrote a nested for loop - four levels deep, tagging subnets across regions and availability zones in another AWS account - and it took me about an hour to get the syntax right:

locals {
  subnet_tags = merge([
    for account, regions in var.accounts : merge([
      for region, azs in regions : merge([
        for az, subnets in azs : {
          for subnet_id, tags in subnets :
          "${account}/${region}/${az}/${subnet_id}" => tags
        }
      ]...)
    ]...)
  ]...)
}

Four merge([...]...) calls stacked on top of each other just to flatten a map of a map of a map of subnets. Claude writes the equivalent in seconds now, and if you asked me to produce that from scratch today, I’d genuinely have to sit and think about it. My reflexes for debugging a broken node over SSH are a little slower than when that was the only way I knew how to do it. I can feel that happening in real time, the same way plenty of engineers who came up after Kubernetes never really learned to hand-roll a node OS image, and were fine, because they never needed to.

Where This Goes Next

The part I’m less sure about is how long “I still give the direction” holds. Right now I’m the one who decides the long-term shape of a stack, because I have the context and the agent doesn’t - not really, not beyond what’s written down in a repo’s AGENTS.md. But that’s exactly the gap those company-wide AI pushes are trying to close: give the agent the whole context, not just one repo’s. If that actually works, an agent with a genuine long-term view of the entire infrastructure - not just this Terraform module, but every decision made across every repo for years - might end up planning better than I do, the same way I can’t out-debug a tool that’s read every changelog for every provider I use.

I think AI is still busy eating the layer just below “give direction”. I’m not fully convinced that’s the last layer it eats. Ask me again in five years whether “I still give the direction” was the whole story, or just where I happened to be standing when I wrote this.

The painting "Lady Guoguo on a spring outing," by one of the Tang Dynasty's most famous artists, Zhang Xuan.

Hong Kong —

Something weird has been happening on China’s internet. One small and vocal corner has been claiming the country’s most famous historical dynasty never existed.

That’s the Tang, which ruled China from 618 AD to 907 AD, a period that saw huge territorial expansion, booming international trade and the flourishing of a high culture envied far beyond its borders. Every child learns about the Tang at school; historical dramas inspired by its palace intrigues are a staple of state TV.

Claiming that the 300-year dynasty was actually a myth is akin to saying the American Civil War never happened, or telling a Brit that the Magna Carta is a fiction. But in recent months it’s been spreading among China’s very online population, posing a headache for a ruling Communist Party obsessed with controlling how its citizens talk about their 5,000 years of fractious and contested history.

It all appears to have started when a history influencer named Qiao Yu deployed some baffling astrological data in a video in May, to “prove” the Tang had never existed.

The outlandish claim has since spread far and wide. Reports later emerged that people were calling up the tourism bureau in Xi’an, the site of the ancient Tang capital, to demand it close down because the dynasty itself had been shown to be a hoax.

The hashtag “the Tang dynasty doesn’t exist” now has tens of millions of views on China’s X-like Weibo.

An anonymous painting depicts elegant ladies of the Tang imperial court enjoying a feast and music, in China.

What motivated the hoax and its spread is hard to pin down. The Tang dynasty has been criticized in the past by Chinese nationalists as its founding emperor was not from the Han ethnic group that makes up most of China’s people. The basis of Qiao Yu’s claim is that the succeeding Song dynasty actually began far earlier than traditionally accepted, therefore canceling out the Tang. The Song dynasty was of ethnic Han origin.

“This is basically insane,” said William Kirby, a professor of China studies at Harvard University, adding that there are facts in history and “the Tang dynasty is definitely one of them.”

“It was much more than a Chinese dynasty. It was an inner Asian dynasty as well,” he said.

As the fallacy spread on social media, China’s powerful and tightly controlled state media waded into the controversy.

Many parents say that, thanks to the hoax, their children now feel confused about historical facts, the Nanjing Daily reported last month.

Three days later, The Paper, a media outlet affiliated with the Shanghai government, was more strident. “Only when rational inquiry becomes the foundation of public life will historical nihilism truly lose the conditions that allow it to flourish,” it said in a commentary.

“Historical nihilism” is a term coined by the ruling communist party that refers to discussion or research that challenges its version of China’s long history. The party itself is a relative newcomer to that story, coming to power in 1949 following a decades-long civil war that tore the country apart.

By mid-August the Tang controversy had attracted the attention of the party’s mouthpiece. In an editorial, the People’s Daily called the claim “absurd,” and issued a warning: “Once young people begin to doubt our national history, our cultural confidence will become like water without a source or a tree without roots, and Chinese-style modernization will lose its profound historical perspective.”

Chinese President Xi Jinping at the welcome ceremony for the China-Central Asia summit in Xi'an in 2023, where a Tang Dynasty-themed performance was held.

During a Lunar New Year event, performers commemorate the Tang dynasty poet Du Fu in Chengdu on February 16, 2024.

In recent years, Chinese leader Xi Jinping has tightened the party’s grip on history and sought harsher penalties for those who dare to challenge the official account.

Xi himself has publicly evoked the past glory of the Tang. In 2023 he hosted leaders from several Central Asian countries in Xian, the ancient Tang capital and terminus of the Silk Road along which goods and ideas used to flow between China and Europe.

The ceremony featured dancers and performers in Tang-style dress, observers noting the nod to a Tang past when foreign envoys came to pay tribute at the emperor’s court.

Using the past to criticize the present

For a thousand years, Chinese people have used reflections on history to reimagine the past, as a creative way to criticize or approach the present.

For example, in the economic boom years of the early 2010s, amid growing criticism of the fixation on consumerism and bourgeois lifestyles, the hardships the Chinese people had undergone during Japan’s brutal occupation during WWII became a favored topic for intellectuals, according to Rana Mitter, a Harvard Kennedy School historian specializing in modern China.

But that discussion also allowed some to point out that the current discourse didn’t give enough credit to the Kuomintang troops who fought the Japanese – but who were also the opponents of Mao Zedong’s communists, Mitter said.

A mural is seen inside the Chienling Tomb near Xi'an, China, dating to the Tang Dynasty (618–907 AD).

Even though the Tang era was more than a thousand years ago, any suggestion – however fanciful – that it never existed could be an irritant for the party, said Mitter.

“Certainly under Xi, the CCP now honors the longer trajectory of Chinese history. This means that all history needs to point in a direction that suggests unified (rather than split) identity, historical and cultural continuity, and an idea of ‘China’ as a long-standing entity that has been continuous through history,” he said.

Callum, a 24-year-old Beijing resident, told CNN that questioning the Tang’s existence could be a way to rebel against established wisdom and a way for China’s young – faced with a slowing economy and greater competition for jobs and resources – to criticize authorities.

“It is a way for people to attack the established authorities. It is almost like saying ‘those so-called experts are nothing special. They have all been misleading people. I am the one who understands the truth, and I have exposed them all.’”

As of mid-August, Qiao Yu’s accounts have been blocked and all of their videos taken down. Any search of Chinese social media questioning if the Tang dynasty existed now pulls up content condemning such lines of reasoning.

In its editorial on the controversy, the People’s Daily attempted to have the final word.

“Chinese civilization has endured for five thousand years,” it said, “and the Golden Age of the Tang Dynasty stands as one of its most dazzling chapters.”