DiegoVallejo

The Anatomy of an AI GPU: From SIMT to Silicon Photonics

hardwareaiarchitectureengineering

Every conversation about AI infrastructure eventually collapses into a spec-sheet shouting match: PFLOPS, HBM gigabytes, interconnect bandwidth. Those numbers are real, but they hide the actual engineering story. A modern AI GPU is a stack of seven distinct constraint systems, and each layer only makes sense in light of the one below it.

This post walks the whole stack, from thread scheduling to photons, comparing how NVIDIA and AMD attack each layer. No benchmarks, no brand loyalty. Just the physics and the architecture. And if you have never touched hardware, do not worry: each section opens with the plain-software version of the idea before diving into the silicon.

An NVIDIA H100 accelerator board

An NVIDIA H100 accelerator. Everything in this post lives somewhere on a board like this one. Photo: Geekerwan, Wikimedia Commons (CC BY-SA 4.0).

I. Silicon Logic: SMs, CUs, and the SIMT Execution Model

If you come from software, here is the one-paragraph version. A CPU core is a brilliant soloist: it runs one thread very fast and switches tasks constantly. A GPU is a marching band: thousands of simple workers that all play the same note at the same time. That trade is what makes GPUs absurdly good at math on big arrays and clumsy at branchy, unpredictable code.

The fundamental building block of a GPU is not a core in the CPU sense. It is a throughput machine built around SIMT: Single Instruction, Multiple Threads. One instruction is fetched once and executed in lockstep across a group of threads, amortizing the enormous cost of fetch and decode over dozens of data lanes. In marching-band terms: the conductor reads the sheet music once, and 32 musicians play it simultaneously.

NVIDIA calls its block the Streaming Multiprocessor (SM); AMD calls its equivalent the Compute Unit (CU). The philosophy is shared, the grouping is not:

  • NVIDIA Warps: a rigid group of 32 threads. The warp is the atomic unit of scheduling. Divergence inside a warp (an if where half the threads take each branch) serializes both paths, and every lane pays for both.
  • AMD Wavefronts: historically Wave64, 64 threads per wavefront, which amortizes instruction issue even further but doubles the divergence penalty. RDNA and later CDNA generations migrated toward Wave32 precisely because narrower groups synchronize more cheaply and waste fewer lanes on divergent control flow.
SIMT DIVERGENCE: WARP (32) VS WAVEFRONT (64)NVIDIA Warp · 32 lanes Branch A · 16 lanes active Branch B · 16 lanes active (A idles) 2 serialized passes · 50% utilizationAMD Wave64 · 64 lanes Branch A · 32 of 64 lanes Branch B · 32 of 64 lanes Same 50%, but twice the lanes stalled per passWhy Wave32 won: narrower groups = cheaper sync, less divergence wasteWave32Wave32Both vendors converged on ~32-lane scheduling granularity for control-heavy code.

The more interesting story is that the classic warp-synchronous model is quietly dissolving. In the newest tensor pipelines, a single thread can issue a Matrix Multiply-Accumulate operation (NVIDIA's tcgen05.mma family is the visible example) that executes asynchronously on dedicated hardware. The warp no longer marches in lockstep through the math; it launches work and gets out of the way. For tensor workloads this effectively sidesteps the traditional branch divergence penalty, because the expensive operation is no longer tied to per-lane control flow at all.

Mental model: the SM is becoming less like a platoon of lockstep soldiers and more like a dispatch desk. Threads file work orders; specialized silicon executes them out of band.

II. Specialized Silicon: Tensor Cores vs. Matrix Core Engines

Plain-software version first: a tensor core is to matrix multiplication what a GPU is to graphics. Instead of computing a matrix multiply as thousands of individual multiply-and-add instructions, the hardware has a dedicated circuit that eats whole tiles of the matrices in one gulp. If your code calls matmul, this is the silicon that actually runs it.

General-purpose FP32 and FP64 units are, for AI purposes, legacy hardware. Nearly all training and inference FLOPS now flow through dedicated matrix engines, and the frontier of those engines is precision reduction. Precision here just means how many bits you spend per number: FP32 is a 32-bit float like a JavaScript number, FP4 squeezes a value into 4 bits, the range of a single hex digit. Fewer bits per number means more numbers per second through the same wires.

Die shot of the NVIDIA GA100 GPU showing the grid of streaming multiprocessors

A die shot of NVIDIA's GA100. The repeating tiles are SMs; each contains tensor cores. Note how much of the chip is copy-pasted structure. Photo: Fritzchens Fritz, Wikimedia Commons (CC0).

The arithmetic justification is simple. For a matrix multiply of dimensions by , cost scales with the reduction dimension . Halving operand width doubles how many elements fit per register, per wire, and per multiplier array:

  • NVIDIA: the Rubin (R100) generation ships 6th-generation Tensor Cores with native 4-bit NVFP4 paths. Native support matters: the hardware doubles throughput per clock along the K dimension instead of emulating narrow types on wider units. Blackwell made FP4 viable; Rubin makes it the default assumption.
  • AMD: CDNA 4 and the upcoming UDNA line answer with Matrix Cores posting dense FP4/FP6 PFLOPS and hardware support for microscaling (MX) formats, where small blocks of values share an exponent. MX formats are a pragmatic compromise: FP4 dynamic range is brutal, and per-block scaling recovers accuracy without giving back the bandwidth win. This matters disproportionately for mixture-of-experts models, where enormous weight tensors are touched sparsely and bandwidth per expert dominates.

The subtext of this entire section: precision is now a system-design variable, not a numerical afterthought. The question is no longer "does FP4 lose accuracy" but "where in the network can I spend accuracy to buy bandwidth."

III. Advanced Packaging: Breaking the Reticle Limit

For the non-hardware readers: chips are printed, literally, by projecting patterns of light through a mask onto silicon. The projector has a maximum image size, called the reticle, of roughly . You cannot print a bigger monolithic chip no matter how much you pay, the same way you cannot print an A2 poster on an A4 printer. Both vendors hit that wall generations ago, and both answered the same way you would: print multiple pages and glue them together seamlessly. That gluing is called advanced packaging, and it is now core competitive technology:

  • NVIDIA Rubin joins two reticle-sized compute dies with NV-HBI, a die-to-die link fast and wide enough that software sees one logical GPU.
  • AMD MI350/MI400 went chiplet-first earlier: multiple XCD compute dies around IOD I/O dies, a philosophy inherited from EPYC.

The glue is TSMC's CoWoS family, and the S/L distinction is worth knowing:

  • CoWoS-S: everything sits on a single monolithic silicon interposer. Excellent wiring density, but the interposer itself is reticle-limited, which caps total package size.
  • CoWoS-L: an organic substrate with small embedded local silicon bridges only where dies actually meet. Density where you need it, cheap substrate where you do not, and a much larger total package canvas.
COWOS-S VS COWOS-LCoWoS-SHBMGPU diemonolithic silicon interposer (reticle-limited)package substrateCoWoS-LHBMdie 0die 1local silicon bridges only at die seamsorganic substrate (large canvas)S: max wiring density, capped package size. L: density only where dies meet,enabling near-monolithic latency across hundreds of billions of transistors + HBM.Simplified cross-sections, not to scale.
AMD Fiji package showing the GPU die surrounded by HBM stacks on a silicon interposer

AMD's Fiji, the 2015 ancestor of this whole approach: a GPU die and four HBM stacks sharing one silicon interposer. Modern CoWoS packages are this idea scaled up tenfold. Photo: Wikimedia Commons (CC BY-SA 4.0).

The punchline: "one GPU" is now a legal fiction maintained by packaging engineers. The die boundary moved from the product level to the interconnect level, and the reticle limit stopped being a ceiling and became a floor plan.

IV. Execution Behavior: Prefill vs. Decode Dynamics

Here is the section that explains most real-world LLM performance mysteries, and you have felt it yourself: paste a huge prompt into a chatbot and the first token takes a while, then the rest stream out at a steady rhythm. That is not one system being slow; it is two completely different workloads wearing one trench coat:

  • Prefill (prompt processing): the entire prompt is processed at once as large matrix-matrix multiplications (GEMM). Arithmetic intensity is high, tensor cores saturate, and the GPU is compute-bound. This is the phase spec sheets are written for.
  • Decode (token generation): one token at a time. Each step is a vector-matrix multiply (GEMV) that must stream the entire set of model weights out of HBM to produce a single token. Arithmetic intensity collapses and the GPU becomes memory-bandwidth-bound. Your million-dollar tensor cores idle while HBM does all the work.

A useful back-of-envelope: decode throughput is bounded by

which is why every serious inference optimization attacks either the numerator (HBM4, more stacks) or the denominator (quantization, MoE sparsity).

THE ROOFLINE: PREFILL VS DECODEArithmetic intensity (FLOPs / byte)Perfbandwidth roof (HBM)compute roof (tensor cores)Prefill (GEMM)Decode (GEMV)tensor cores idle, HBM saturatedSame GPU, same model: two different bottlenecks depending on the phase.

Then there is the KV cache, the third tenant of memory. Every generated token appends keys and values that all future tokens must read. Long contexts and large batches inflate it until it competes with the weights themselves for SRAM and HBM residency. The hardware responses are direct:

  • Native FP8/FP4 KV cache quantization, cutting cache footprint 2x to 4x with hardware-supported formats rather than software repacking.
  • NVIDIA's Tensor Memory Accelerator (TMA) updates, including inline descriptor modification, so address arithmetic and tensor metadata handling stop consuming instruction issue slots that decode kernels cannot spare.

If you remember one thing: prefill buys FLOPS, decode buys bandwidth. A chip can win one phase and lose the other, and "tokens per second" quietly averages over both.

V. System Constraints: Power Delivery and Thermal Limits

Modern accelerators are not gate-limited; they are electron-limited. The interesting engineering has moved from logic design to power delivery and heat removal. For scale: a flagship AI GPU draws more power than a household microwave, continuously, from an area smaller than a credit card, and a rack holds dozens of them.

Water cooling infrastructure at the Swiss National Supercomputing Centre

Water cooling plant at the Swiss National Supercomputing Centre (CSCS). The plumbing is part of the computer now. Photo: CSCS, Wikimedia Commons (CC BY-SA 3.0).

The Power Delivery Network (PDN). Flagship accelerator sockets have crossed . At roughly one volt of core supply, that is over a thousand amps entering a package the size of your palm. Three consequences:

  1. IR drop: at these currents, resistive losses across the package and die metal ( ) eat meaningful voltage before it reaches a transistor. Margins that were noise at 300 W are failure modes at 1200 W.
  2. Transients: tensor cores switch from idle to full-power in nanoseconds. The resulting voltage spikes can undershoot logic thresholds unless decoupling sits physically close to the load.
  3. On-package VRMs: the fix for both is moving the final voltage conversion stage onto the package itself, delivering higher voltage (lower current) across the board and converting at the last possible millimeter.

Thermal management. Heat is not uniform. A dense tensor core array under sustained GEMM load forms a localized hot spot that can throttle the whole die while the average temperature still looks fine. Air cooling cannot extract kilowatt-class, concentrated flux; this is why rack-scale designs (NVIDIA's Vera Rubin NVL72 class systems, AMD's Helios) mandate direct-to-chip liquid cooling as a baseline, not a premium option. The coolant loop is now part of the computer.

Field note: when a datacenter GPU underperforms its spec sheet, the cause is more often volts and degrees than FLOPS. The silicon is usually capable; the environment is the constraint.

VI. The Memory Wall and Silicon Photonics

The software analogy here is the one every backend dev knows: your database query is rarely slow because the CPU cannot compute; it is slow because the data is not where the compute is. GPUs have exactly that problem, at nanosecond scale. Compute grew faster than memory for forty years, and the bill is due. Three fronts:

Simplified cross-section diagram of a High Bandwidth Memory stack next to a processor die

HBM in cross-section: DRAM dies stacked vertically and wired straight down into the package, millimeters from the GPU. Distance is bandwidth. Diagram: Wikimedia Commons (CC BY-SA 4.0).

On-package memory: HBM4. The next stacks push past per stack toward aggregate package bandwidths of . Capacity strategies diverge: AMD's MI400 line leans into a massive footprint, up to via 12-high stacks, betting that fitting bigger models (or more KV cache) per package wins inference economics. NVIDIA's Rubin-class parts ship around of HBM4, betting the balance point lies closer to bandwidth per die and rack-level disaggregation.

Co-Packaged Optics (CPO). Copper SerDes at PAM4 is near the end of the line: at those rates, a copper trace a few tens of centimeters long loses most of its signal, and the DSP burned recovering it costs real watts per lane. The answer is moving the optical engine onto the GPU substrate itself, so electrical signals travel millimeters and photons handle everything beyond the package. Fewer retimers, less signal conditioning, and a large cut in interconnect power, which at cluster scale is measured in megawatts.

Scale-out fabrics. NVIDIA extends its proprietary vertical stack with NVLink 6 at of fabric bandwidth per GPU. AMD upgrades Infinity Fabric and, more strategically, backs UALink, an open standard for accelerator fabrics. It is the classic topology of the industry repeating itself: a vertically integrated proprietary fabric versus an open coalition, this time at exascale.

VII. Software Abstraction: The Final Constraint

Everything above is potential energy. A GPU executes what a compiler and runtime manage to map onto it, and this is the layer where the two ecosystems differ most. If sections I through VI felt foreign, this one will not: it is the same story as every platform war you have lived through, and the moat is the same as always. Not the hardware. The developer experience.

CUDA's advantage is not the language; it is twenty years of co-design. Every architectural feature in sections I through VI (async MMA, TMA descriptors, FP4 paths, multi-die address spaces) ships with day-one compiler support, profiling visibility, and a library ecosystem (cuBLAS, cuDNN, CUTLASS, Triton backends) already tuned for it. Fine-grained control over the hardware is exposed, documented, and battle-tested because the same company builds both sides.

ROCm's trajectory is real but asymmetric. The hardware, as this whole post argues, is competitive and in some dimensions (memory capacity, open fabrics) ahead. The engineering effort now is a frictionless abstraction layer: HIP as a porting path, day-one framework support in PyTorch, Triton as a vendor-neutral kernel language. The gap is not correctness anymore; it is the long tail of kernels, profilers, and tribal knowledge that make performance predictable rather than merely possible.

The honest assessment: hardware parity is necessary but not sufficient. The spec sheet tells you the ceiling; the software stack decides how close to the ceiling you actually live.

Closing: Seven Layers, One Budget

Walk the stack once more, bottom-up: photons feed HBM, HBM feeds matrix engines, matrix engines are scheduled by SIMT front-ends, all of it inside a power and thermal envelope, glued by packaging, and finally made usable (or not) by a compiler.

The deepest shift is that these layers no longer improve independently. FP4 exists because HBM bandwidth is precious. CoWoS-L exists because the reticle is fixed. Liquid cooling exists because tensor cores are dense. CPO exists because copper gave up. Modern accelerator design is not about making any one layer faster; it is about spending one shared budget of joules, bytes, and square millimeters across all seven layers at once.

That is also why the spec-sheet shouting match misses the point. The interesting question is never "how many PFLOPS." It is: which layer is your workload actually bound by, and who engineered that layer best?