Technical Analysis & Systems Synthesis by DistributedApps.ai | Synthesizing frontier research from DeepSeek, Moonshot AI, Zhipu AI, Alibaba, NVIDIA, and Google DeepMind into production engineering blueprints.
This is Chapter 5 of the 10-part series The Physics & Engineering of Frontier LLM Inference. I announced the full curriculum, including every later chapter, in Announcing the 10-Part Substack Series.
Read the earlier parts first: Chapter 1: The Physics of LLM Inference, Chapter 2: The KV Cache Frontier, Chapter 3: Next-Gen Speculative Decoding, Chapter 4: Extreme Quantization.
Modern deep learning accelerators possess terrifying theoretical compute capacity. An NVIDIA Hopper H100 SXM5 GPU delivers 989.5 TFLOPS of dense FP16/BF16 matrix arithmetic and nearly 2,000 TFLOPS of FP8 Tensor Core throughput. The Blackwell B200 raises this ceiling to 2.25 PFLOPS of FP16 and 4.5 PFLOPS of FP8 compute. Yet, if you execute textbook scaled dot-product attention in PyTorch on an H100 across a 16,000-token sequence, your observed hardware utilization will rarely exceed 15% Model FLOPs Utilization (MFU)—wasting over 800 TFLOPS of raw compute capacity per GPU.
Why does this compute collapse occur?
The answer lies in the physics of the memory hierarchy. Attention is not bounded by how fast Tensor Cores can execute matrix multiply-accumulate operations; it is bounded by the speed at which operands can be moved across the silicon. On an H100 GPU, the High-Bandwidth Memory (HBM3) bus delivers 3.35 TB/s of bandwidth with a memory latency of 400 to 800 clock cycles. In stark contrast, the on-chip Static Random-Access Memory (SRAM)—distributed across 132 Streaming Multiprocessors (SMs) as 228 KB of Shared Memory and L1 cache per SM—delivers an aggregate bandwidth exceeding 33 TB/s with a latency of just 15 to 30 clock cycles.
When standard attention materializes the intermediate N × N attention score matrix in global HBM, it forces the GPU to transfer O(N²) bytes over the narrow 3.35 TB/s bus. Because the arithmetic intensity of memory-bound elementwise operators (softmax scaling, masking, exponentiation, and normalization) is strictly below 1.0 FLOP/byte—against an H100 hardware saturation roofline of 295.4 FLOPs/byte—the execution pipeline stalls completely, starving the Tensor Cores.
Unlocking 70% to 80%+ MFU across frontier LLM architectures requires complete algorithm-hardware co-design: restructuring attention algorithms so that intermediate matrices never leave ultra-fast on-chip SRAM, and orchestrating asynchronous hardware engines to overlap memory movement directly with Tensor Core arithmetic.
========================================================================================
THE 2026 HARDWARE-AWARE ATTENTION SUITE
========================================================================================
1. ONLINE SOFTMAX TILING (Milakov-Gimelshein)
Eliminates O(N²) HBM materialization. Computes exact running row-maxima and
normalizers in a single pass entirely inside on-chip SRAM / VMEM registers.
IO Complexity compressed from O(N²) down to O(N² d² M⁻¹ + Nd).
2. FLASHATTENTION-3 (Hopper TMA & Warp Specialization)
Exploits Hopper H100 / Blackwell B200 Tensor Memory Accelerators (TMA) for
zero-register global-to-shared transfers. Decouples Producer warps from Consumer
warpgroups. Leverages asynchronous WGMMA and FP8 block scaling for 75-80% MFU.
3. FLASHDECODING & FLASHDECODING++ (Split-K Sequence Parallelism)
Solves the single-query (Q_len = 1) decode occupancy collapse. Partitions the
KV cache across sequence splits (Split-K), saturating all 132 SMs and reducing
128k decode latency from 27.2 ms to 3.40 ms (8.0x speedup).
4. GOOGLE TPU PALLAS KERNELS (TPU v5p & Trillium v6e)
Bypasses high-level XLA graph compilers with imperative JAX Pallas kernels.
Controls 128×128 systolic Matrix Multiply Units (MXUs) and Vector Processing Units
(VPUs) via double-buffered DMA BlockSpecs and ping-pong VMEM memory spaces.
========================================================================================
In this exhaustive 6,300+ word series, we tear open the silicon, mathematics, and low-level kernel code powering modern attention acceleration across NVIDIA Hopper/Blackwell GPUs and Google Cloud TPUs:
**Memory Hierarchy Physics & Attention IO Complexity:**Mathematical derivation of operational intensity, hardware saturation thresholds, and exact HBM traffic equations.**Milakov-Gimelshein Online Normalization:**Inductive mathematical proof of single-pass running softmax rescaling, register accumulator recurrence relations, and backward-pass activation recomputation.**FlashAttention-3 Microarchitecture on Hopper:**Deep dive into hardware TMA descriptors (cp.async.bulk.tensor
), Producer-Consumer Warp Specialization, asynchronous WGMMA matrix multiply instructions, interleaved GEMM overlapping, and FP8 E4M3 dynamic tile scaling. Includes a production C++/CUDA/PTX kernel listing.**FlashDecoding & FlashDecoding++:**The mechanics of sequence-parallel Split-K decomposition, partial log-sum-exp workspace management, hierarchical reduction trees, and asynchronous softmax estimation. Includes a production OpenAI Triton kernel listing.**Google TPU Pallas Optimization:**DeepMind's Pallas framework in JAX, 128×128 systolic MXU tile alignment, double-buffered Direct Memory Access (DMA) pipelines, and VMEM scratchpad management on TPU v5p and Trillium (v6e). Includes a complete JAX/Pallas production kernel listing.**Empirical Benchmarks & Decision Matrix:**Real-world throughput (TFLOPS), MFU (%), and latency across H100, H200, B200, TPU v5p, and Trillium across context lengths from 4k to 1,048,576 tokens (1M context), along with an engineering selection flowchart.**2026+ Frontiers:**Microscaling data formats (MXFP4, NVFP4), State-Space hybrid kernel fusion (Mamba-2 / SSD + Attention), and hardware-assisted sparse verification engines.**Academic Citations:**Complete primary lab references and system papers.
Below is the complete architectural roadmap covered in this chapter:
**The Arithmetic Intensity Crisis of Scaled Dot-Product Attention:**SRAM block tiling and the Milakov-Gimelshein online softmax normalization theorem.**FlashAttention-3 Microarchitecture:**Tensor Memory Accelerator (TMA) asynchronous global memory loads and hardware warp specialization on Hopper/Blackwell.**FlashDecoding & FlashDecoding++ Split-K Reduction Trees:**Partitioning KV sequences across thread blocks to restore GPU occupancy during autoregressive decoding.**Google TPU v5p & Trillium v6e Pallas Kernel Pipelines:**Matrix Multiply Units (MXUs), Vector Memory (VMEM) double-buffering, and asynchronous DMA execution.**Attention Kernel Benchmark Suite:**Multi-accelerator comparison measuring Model FLOPs Utilization (MFU), latency, and context scaling.
To access the complete technical treatise, production code repositories, Triton/CUDA kernels, and infrastructure sizing templates for this chapter, upgrade to a paid subscription today.
Special Offer: Get 50% OFF the annual subscription to the 2026 Foundation Model Inference Series using the link below:
