The Softmax Wall and the 2:4 Sparse Attention Trap

TL;DR#

For the past decade, optimizing Transformer attention was treated as a synonym for accelerating matrix multiplication. Every hardware generation followed the same playbook: drop the numerical precision of the Tensor Cores or Matrix Cores—from FP32 to FP16/BF16, then to FP8, and now to 4-bit MXFP4 and NVFP4—to double the matrix FLOPs per clock.

That playbook just hit a physical wall of its own making. By scaling matrix throughput $2\times$ to $4\times$ per generation while leaving Special Function Units (SFUs), vector ALUs, and shared-memory bandwidth behind, chipmakers inverted the attention roofline.

In FP8 and FP4, attention is no longer bottlenecked by matrix multiplication. It is bottlenecked by the transcendental and vector math of softmax and the on-chip memory traffic required to shuttle intermediate scores between execution domains:

• On AMD’s Instinct MI355X (gfx950, CDNA 4) and Instinct MI455X (gfx1250, CDNA 5): On MI355X, dropping from BF16 to FP8 doubles matrix throughput, so the two attention GEMMs ($\mathbf{Q}\mathbf{K}^\top$ and $\mathbf{P}\mathbf{V}$) for a $32 \times 128$ tile take just 1,024 cycles. The exact FP8 softmax between them, however, runs on the vector ALU at unchanged speed and takes 896 vector cycles—leaving AMD’s default AITER kernel with its matrix pipe idle 31.7% of the time. To break that bottleneck in ZYFA (released this week), Zyphra had to bypass the hardware exponential instruction altogether with “Lean P”—an IEEE/FP8 bit-pattern trick that computes $2^x$ in e4m3 using one floating-point multiply-add and one integer conversion, cutting softmax from 896 to 384 cycles and lifting steady-state matrix utilization from 60% to 86%. Even then, preserving dynamic range in FP8 probabilities flips more operand bits in the multiplier trees, triggering a switching-power clock drop from 2,108 MHz to 1,628 MHz on real Mistral-7B activations. In CDNA 5 (gfx1250, MI455X), AMD overhauled the compute core—switching to Wave32 WMMA, expanding addressable VGPRs $4\times$ to 1,024 per wave, adding VOPD3 dual-issue and hardware block-scaled conversions (V_CVT_SCALEF32), and doubling transcendental throughput ($2\times$ vs. MI355X). Yet because CDNA 5 simultaneously scales OCP MXFP8 and MXFP4 matrix throughput by $4\times$ over MI355X (reaching 20.13 PFLOPS in FP8/MXFP8 and 40.26 PFLOPS in MXFP4 against 5.03 PFLOPS in BF16), low-bit matrix engines on MI455X outpace the doubled transcendental units by an even wider margin.
• On NVIDIA’s Blackwell (sm_100), Blackwell Ultra (sm_103), and Vera Rubin (sm_107): From Hopper H100 to Blackwell B200, BF16 Tensor Core throughput per SM doubled ($4{,}096 \to 8{,}192\text{ ops/clk/SM}$), while the multifunction exponential unit (MUFU.EX2) stayed frozen at $16\text{ ops/clk/SM}$ and shared-memory read bandwidth stayed frozen at $128\text{ B/clk/SM}$. For a $128 \times 128 \times 128$ tile on B200, MUFU.EX2 takes 1,024 cycles—already tying BF16 matmuls ($1{,}024\text{ cycles}$), taking $2\times$ longer than FP8 matmuls ($512\text{ cycles}$), and taking $4\times$ longer than NVFP4 matmuls ($256\text{ cycles}$). FlashAttention-4 had to route 10%–25% of its exponentials onto general-purpose FP32 FMA pipelines via degree-3 Sollya polynomial emulation and skip ~90% of online softmax rescaling steps. NVIDIA responded by doubling SFU throughput in Blackwell Ultra (B300) to $32\text{ ops/clk/SM}$. Vera Rubin (sm_107) did not double it again for FP32. NVIDIA’s own Rubin spec table lists EX2 at 32 ops/clk/SM for FP32 and 64 for FP16, so the second doubling only exists for half-precision exponentials. Rubin meanwhile doubled per-SM FP8 Tensor Core throughput. With an FP32 exponential, a Rubin FP8 tile is $2\times$ SFU-bound and an NVFP4 tile $4\times$, worse than B300. Rubin only gets back to a zero-slack FP8 tie if the kernel moves the exponential itself to FP16. NVIDIA’s most direct hardware answer, the Rubin CPX context GPU with “3× exponent operations,” was pulled from the roadmap at GTC 2026.

Against this inverted roofline, a second hardware pitch has resurfaced: accelerating dynamic attention with fine-grained 2:4 Sparse Tensor Cores—most prominently in NVIDIA’s Vera Rubin (sm_107), which adds the PTX 9.4 instruction tcgen05.ld.red.spcompress to prune 2 of every 4 attention logits while loading from Tensor Memory (TMEM) so that $\mathbf{P}\mathbf{V}$ can run on tcgen05.mma.sp, alongside 4:2 structured-sparse matrix instructions (V_SWMMAC) in AMD’s CDNA 5 (gfx1250). The headline numbers are still sold with sparsity. Rubin’s 50 PF NVFP4 “inference” figure is Transformer Engine compute with a new, undisclosed “adaptive compression” scheme that SemiAnalysis says is neither 2:4 nor Blackwell’s pairwise 4:8; the dense training figure is 35 PF. AMD’s MI455X datasheet lists 2:4 peaks only for FP16, BF16 and INT8, at 10.07 PF, which is half of the same chip’s 20.13 PF of dense FP8.

Let me state my position up front: applying 2:4 (or FP4 pairwise 4:8) hardware sparse matrix units to attention is an architectural trap.

1. It accelerates the shrinking part of the loop while dumping repack and online-pruning work onto the bottleneck. On Blackwell and Rubin’s tcgen05 architecture, Tensor Cores write $\mathbf{S} = \mathbf{Q}\mathbf{K}^\top$ into TMEM and require operand A ($\mathbf{P}$) for $\mathbf{P}\mathbf{V}$ to come from TMEM or Shared Memory (SMEM)—never from thread registers. Yet Rubin’s tcgen05.ld.red.spcompress writes its compressed FP32 values (cdata) and 2-bit metadata (mdata) into general-purpose registers. Software must still execute sparse softmax in registers, downcast dtypes, repack $\mathbf{P}$ into the MMA A-layout, repack and swizzle the 2-bit metadata, store both back to SMEM/TMEM (tcgen05.st), and wait on a completion fence before tcgen05.mma.sp can issue. (AMD’s CDNA 5 avoids that TMEM round-trip because V_SWMMAC keeps operands and indices in VGPRs, but CDNA 5 lacks a hardware spcompress instruction, leaves unscaled FP8 at the same $K=128$ tile depth as dense V_WMMA, and disables structured sparsity entirely on block-scaled V_WMMA_SCALE* instructions.) Across Ada, Hopper, and Blackwell, online pruning and compression of $\mathbf{P}$ consumes 70% to 87% of total runtime, turning a $1.49\times\text{--}1.77\times$ standalone FP16 sparse-GEMM speedup into a net $0.38\times\text{--}0.51\times$ slowdown in FP16 (and trailing dense NVFP4 at $0.93\times\text{--}0.96\times$ in the $4\times$ tier).
2. At equal bits and equal Tensor Core tiers, dropping one dense precision step crushes 2:4 on both speed and accuracy. Even before counting metadata or 2:4 structural constraints, analytical distribution models and convex SDP lower bounds (Kuzmin et al., 2023) show that uniform quantization dominates unstructured pruning at every practical bit-width above 3 bits. Once 2:4’s 2-bit index per kept value ($1\text{ bit}$ per original element) is counted, 2:4 on $2m$-bit values costs the exact same memory as dense $(m+1)$-bit quantization ($\text{2:4 + FP16} = 9\text{ bits} > \text{FP8}$). Quantization error energy shrinks $4\times$ with every added bit, whereas 2:4 deletes a fixed fraction of tensor energy ($e_{2:4} = 13.2\%$ on Gaussian data) that never shrinks with bit-width—and Hadamard rotations ($\mathbf{Q}, \mathbf{K}$) make 2:4 twice as bad by spreading energy evenly across each 4-tuple. Across 221,408 measured configurations (Qin, 2026), dense FP8 and dense NVFP4 beat 2:4 at the $2\times$ and $4\times$ compute tiers on every operand.
3. Perplexity hides severe degradation in fact recall and multi-step reasoning, single-sided pruning caps speedup at $1.29\times$, and Blackwell/Rubin’s pairwise 4:8 FP4 constraint collapses model quality. Evaluating pruned models by WikiText-2 perplexity is a mirage: parametric fact recall breaks at just 20%–30% sparsity (Jin et al., 2024; Jaiswal et al., 2024), pruning small weights monotonically and irreversibly impairs harder tasks by destroying the pre-trained loss basin (Yin et al., 2024), and on Qwen3-1.7B, 2:4 sparsity collapses reasoning accuracy from 54.0% to 15.3% (Shao et al., 2026). Trying to rescue 2:4 accuracy by relaxing to 6:8 sparsity (SlideSparse) forces a $1.5\times$ activation duplication tax that caps theoretical speedup at $1.33\times$. Inside attention, pruning $\mathbf{P}$ alone leaves $\mathbf{Q}\mathbf{K}^\top$ dense ($1.29\times$ Amdahl ceiling), pruning $\mathbf{Q}$ or $\mathbf{K}$ wrecks early layers and long-context retrieval, and at the NVFP4 floor hardware forces pairwise 4:8 sparsity—blowing up one-shot Llama-3.1-8B perplexity from $6.98$ to $108.83$.

You can see the verdict in the production serving codebases. In 2026, vLLM and LLM Compressor deleted their 2:4 sparsity kernels for lack of adoption. Where fine-grained sparsity genuinely pays off—such as $>99\%$ unstructured activation sparsity in ReLU-gated FFNs via Sakana AI and NVIDIA’s Tile-wise ELLPACK (TwELL) kernels on CUDA Cores, or static weights combined with low-rank adapters below the 4-bit floor—it works by skipping $>99\%$ of column/row fetches or operating offline on static weights, neither of which applies to 2:4 attention. In attention, vLLM v0.33.0, SGLang v0.5.23, and FlashInfer v0.7.3—across NVIDIA Hopper/Blackwell/Rubin, AMD MI355X/MI455X, and Google TPUv7 SparseCores—converged on coarse token-level sparse attention (DeepSeek Sparse Attention / DSA, QSA, MSA): a lightweight ReLU indexer selects the top $k = 512\text{ to }2{,}048$ tokens out of $1\text{M}$, skipping 99.8% of the $\mathbf{Q}\mathbf{K}^\top$ matmuls, 99.8% of the softmax exponentials, and 99.8% of the KV HBM reads before the main attention loop even starts, and then executes the surviving tokens as dense tiles at peak low-precision Tensor Core speed.

1. Why did FP8 and FP4 make the softmax the bottleneck?#

To see why softmax suddenly became the governor of long-context prefill and training, start from the two-stage arithmetic structure of self-attention. Given query, key, and value matrices $\mathbf{Q}, \mathbf{K}, \mathbf{V} \in \mathbb{R}^{L \times d}$ for an attention head of dimension $d$, tiled into blocks of $M$ queries by $N$ keys, each iteration of the inner FlashAttention loop executes two matrix multiply-accumulates (MMAs) separated by a row-wise online softmax:

Count the operations required for a single $M \times N$ tile with head dimension $d$:

1. Matrix multiply-accumulate work ($\mathbf{Q}\mathbf{K}^\top$ and $\mathbf{P}\mathbf{V}$): Each of the two GEMMs performs $2MNd$ multiply-adds, for a total of $$W_{\text{MMA}} = 4MNd \;\;\text{FLOPs}$$ Crucially, $W_{\text{MMA}}$ runs on dedicated matrix engines (NVIDIA Tensor Cores, AMD Matrix Cores) whose per-cycle throughput $\Theta_{\text{MMA}}(b)$ doubles (or quadruples) every time you halve the operand bit-width $b$—from 16-bit (BF16/FP16) to 8-bit (FP8 E4M3/MXFP8) to 4-bit (NVFP4/MXFP4).
2. Non-matmul softmax and rescaling work ($\mathbf{S} \mapsto \mathbf{P}$ and $\mathbf{O}$ correction): Between the two GEMMs, every one of the $M \times N$ score elements in $\mathbf{S}$ must undergo a row-max reduction, a scaled subtraction, a base-2 transcendental exponentiation $$p_{ij} = \exp\!\left(\frac{s_{ij} - m_i}{\sqrt{d}}\right) = 2^{\left(s_{ij} - m_i\right) \cdot \left(\log_2 e / \sqrt{d}\right)}, \qquad W_{\text{exp}} = MN \;\;\text{exponentials}$$ a row-sum reduction $\ell_i = \sum_j p_{ij}$, a down-conversion to the MMA input format (BF16, FP8, or FP4), and—whenever a block’s row maximum increases—a vector rescaling of the $M \times d$ running output accumulator $\mathbf{O}$.

Look at the fundamental scaling asymmetry between $W_{\text{MMA}}$ and $W_{\text{exp}}$:

• The ratio of matrix FLOPs to transcendental exponentials is fixed solely by the head dimension $d$: $$\frac{W_{\text{MMA}}}{W_{\text{exp}}} = \frac{4MNd}{MN} = 4d$$ For the standard head dimension $d = 128$ used across Llama, Qwen, DeepSeek, Kimi, and GLM, every 1 exponential evaluation is paired with exactly $4 \times 128 = 512$ matrix FLOPs (and for $d = 64$, as in Meta’s 816B Muse Spark Open, only 256 matrix FLOPs per exponential).
• Lowering GEMM precision to FP8 or FP4 shrinks the cycle time of those 512 matrix FLOPs by $2\times$ to $8\times$, but does not reduce $W_{\text{exp}} = MN$ by a single operation. Because attention logits $\mathbf{S} = \mathbf{Q}\mathbf{K}^\top$ must be accumulated and exponentiated in FP32 to avoid catastrophic overflow or underflow before normalization, the softmax stage still has to evaluate $M \times N$ exponentials in FP32 on Special Function Units (NVIDIA’s MUFU.EX2) or vector ALUs (AMD’s VALU v_exp_f32), and still has to move $\mathbf{S}$, $\mathbf{P}$, and $\mathbf{O}$ across on-chip memories (TMEM, SMEM/LDS, and the register file).

For an attention kernel to remain compute-bound on its matrix engines without stalling on softmax, the hardware’s per-SM/per-WGP matrix throughput $\Theta_{\text{MMA}}(b)$ and exponential throughput $\Theta_{\text{exp}}$ must satisfy:

At $d = 128$, the moment a chip’s matrix throughput exceeds $512\times$ its exponential throughput ($\Theta_{\text{MMA}} / \Theta_{\text{exp}} > 512$), the exponential unit takes more clock cycles per tile than both matrix multiplications combined. At $d = 64$, that crossover happens at $\Theta_{\text{MMA}} / \Theta_{\text{exp}} > 256$.

1.1 How did the cycle budget change from Hopper (sm_90) to Blackwell (sm_100)?#

Now plug the actual per-SM hardware throughputs of NVIDIA’s Hopper H100 (sm_90) and Blackwell B200 / GB200 (sm_100) into that inequality, following the roofline analysis in FlashAttention-4 (Zadouri, Hoehnerbach, Shah, Liu, Thakkar, & Dao, arXiv:2603.05451, §2.2 & §3.1):

1. On Hopper H100 (sm_90):

   • BF16 Tensor Cores: $4{,}096\text{ FLOPs/clk/SM}$ ($\sim 1\text{ PFLOPS}$ dense across 132 SMs).
   • FP8 Tensor Cores: $8{,}192\text{ FLOPs/clk/SM}$ ($\sim 2\text{ PFLOPS}$ dense).
   • Exponential unit (MUFU.EX2): $16\text{ ops/clk/SM}$.
   • Shared Memory (SMEM) read bandwidth: $128\text{ bytes/clk/SM}$.
   • Check the ratio at $d = 128$: in BF16, $\Theta_{\text{MMA}} / \Theta_{\text{exp}} = 4{,}096 / 16 = 256 \le 512$. For a $128 \times 128 \times 128$ tile, the two BF16 MMAs take $T_{\text{MMA}} = \frac{4 \cdot 128^3}{4096} = \mathbf{2{,}048\text{ cycles}}$, while the $128^2 = 16{,}384$ exponentials take $T_{\text{exp}} = \frac{128^2}{16} = \mathbf{1{,}024\text{ cycles}}$. Because matrix math took $2\times$ longer than MUFU.EX2, FlashAttention-3 could hide the entire BF16 softmax behind the Tensor Cores using a two-warpgroup ping-pong schedule. In FP8 on Hopper, however, $\Theta_{\text{MMA}} / \Theta_{\text{exp}} = 8{,}192 / 16 = 512$: $T_{\text{MMA}}$ dropped to $1{,}024\text{ cycles}$, already tying $T_{\text{exp}} = 1{,}024\text{ cycles}$.

2. On Blackwell B200 / GB200 (sm_100): As documented in zartbot’s November 2025 microarchitectural analysis (Inside Nvidia GPU: Discussing Blackwell’s Limitations and Predicting Rubin’s Microarchitecture) and the FlashAttention-4 roofline model, NVIDIA aggressively expanded the 5th-generation Tensor Core (tcgen05.mma) and added 256 KB of Tensor Memory (TMEM) plus DSMEM per SM—consuming so much die area that the SM count per die was trimmed to 80 (148 enabled across two dies on B200, 160 on GB200) while the SFU (MUFU.EX2) and SMEM bandwidth were left unscaled:

   • BF16 Tensor Cores: Doubled to $8{,}192\text{ FLOPs/clk/SM}$ ($2.25\text{ PFLOPS} / (1.85\text{ GHz} \times 148\text{ SMs})$).
   • FP8 Tensor Cores: Doubled to $16{,}384\text{ FLOPs/clk/SM}$ ($4.5\text{ PFLOPS}$ on B200, $5.0\text{ PFLOPS}$ on GB200).
   • NVFP4 Tensor Cores: Quadrupled to $32{,}768\text{ FLOPs/clk/SM}$ ($9.0\text{--}10.0\text{ PFLOPS}$ dense).
   • Exponential unit (MUFU.EX2): Still $16\text{ ops/clk/SM}$ ($1.0\times$ scaling from Hopper).
   • Shared Memory (SMEM) read bandwidth: Still $128\text{ bytes/clk/SM}$ ($1.0\times$ scaling from Hopper).

Compute the exact cycle budget per SM on B200 (sm_100) for a $128 \times 128 \times 128$ forward tile ($M = N = d = 128$):

• MMA compute cycles: $$T_{\text{MMA}}(\text{BF16}) = \frac{4 \cdot 128^3}{8192} = \mathbf{1{,}024\text{ cycles}}, \quad T_{\text{MMA}}(\text{FP8}) = \frac{4 \cdot 128^3}{16384} = \mathbf{512\text{ cycles}}, \quad T_{\text{MMA}}(\text{NVFP4}) = \frac{4 \cdot 128^3}{32768} = \mathbf{256\text{ cycles}}$$
• SMEM read cycles (BF16 operands): $\mathbf{Q}\mathbf{K}^\top$ is a shared-shared (SS) MMA reading $128 \times 128$ of $\mathbf{Q}$ and $128 \times 128$ of $\mathbf{K}^\top$ from SMEM; $\mathbf{P}\mathbf{V}$ is a tensor-shared (TS) MMA reading $\mathbf{P}$ directly from TMEM and $128 \times 128$ of $\mathbf{V}$ from SMEM. At 2 bytes/element and $128\text{ B/clk}$: $$T_{\text{SMEM}}(\text{fwd, BF16}) = \frac{2 \times (128 d + 128 d + 128 N)}{128} = \frac{3 M N d}{8192} = \mathbf{768\text{ cycles}}$$
• Exponential unit (MUFU.EX2) cycles: $$T_{\text{exp}} = \frac{M N}{16} = \frac{128^2}{16} = \mathbf{1{,}024\text{ cycles}} \quad\text{(independent of GEMM precision)}$$

And in the backward pass on B200 (which chains 5 MMAs: $\mathbf{S}^\top = \mathbf{K}\mathbf{Q}^\top$, $\mathbf{dP}^\top = \mathbf{V}\mathbf{dO}^\top$, $\mathbf{dV} = \mathbf{P}^\top\mathbf{dO}$, $\mathbf{dQ} = \alpha\,\mathbf{dS}\mathbf{K}$, and $\mathbf{dK} = \alpha\,\mathbf{dS}^\top\mathbf{Q}$), 5 BF16 MMAs take $T_{\text{MMA}} = \frac{10 \cdot 128^3}{8192} = \mathbf{2{,}560\text{ cycles}}$, while 1-CTA SMEM traffic (reading 8 BF16 MMA operands from SMEM = $2{,}048\text{ cycles}$, writing BF16 $\mathbf{dS}$ to SMEM = $256\text{ cycles}$, and writing + TMA-reading FP32 $\mathbf{dQ}$ = $1{,}024\text{ cycles}$) takes $T_{\text{SMEM}} = \mathbf{3{,}328\text{ cycles}}$—exceeding Tensor Core compute time by 30% (and requiring 2-CTA cooperative MMA with $384\text{ cycles}$ of DSMEM $\mathbf{dS}$ exchange to bring total SMEM traffic down to $2{,}688\text{ cycles}$).

Look at what happened to the forward pass on B200:

• Even in BF16, $\Theta_{\text{MMA}} / \Theta_{\text{exp}} = 8{,}192 / 16 = 512 = 4d$: the SFU ($1{,}024\text{ cycles}$) is already tied 1:1 with the Tensor Cores ($1{,}024\text{ cycles}$), before counting a single cycle of row-max reduction, row-sum reduction, scale subtraction, or TMEM load/store.
• In FP8, $\Theta_{\text{MMA}} / \Theta_{\text{exp}} = 1{,}024 = 8d$: the SFU ($1{,}024\text{ cycles}$) takes $2\times$ longer than both Tensor Core GEMMs combined ($512\text{ cycles}$).
• In NVFP4, $\Theta_{\text{MMA}} / \Theta_{\text{exp}} = 2{,}048 = 16d$: the SFU ($1{,}024\text{ cycles}$) takes $4\times$ longer than both Tensor Core GEMMs combined ($256\text{ cycles}$).

1.2 Did B300 (sm_103) and Rubin (sm_107) fix the exponential unit?#

At this point, a hardware defender will raise the natural objection: “NVIDIA recognized the B200 SFU mistake. Blackwell Ultra (B300 / GB300, sm_103) doubled SFU throughput from 16 to 32 ops/clk/SM, and Vera Rubin (sm_107) doubles it again.”

The first half is true. The second half is true only for half precision. Check both against NVIDIA’s own specifications and benchmarks.

In Blackwell Ultra (B300 / GB300, sm_103), NVIDIA doubled MUFU.EX2 throughput from 16 to 32 ops/clk/SM. In NVIDIA’s mufu_ex2_bench microbenchmark (Making Softmax More Efficient with NVIDIA Blackwell Ultra), isolating raw MUFU.EX2 loops across the full GPU confirms the $2\times$ jump:

• GB200 (sm_100): $2{,}454\text{ Gop/s}$ (BF16x2, $4{,}908\text{ GFLOPS}$), $4{,}938\text{ Gop/s}$ (BF16), and $4{,}943\text{ Gop/s}$ (FP32).
• GB300 (sm_103): $4{,}996\text{ Gop/s}$ (BF16x2, $9{,}992\text{ GFLOPS}$), $9{,}738\text{ Gop/s}$ (BF16), and $10{,}024\text{ Gop/s}$ (FP32).

On DeepSeek-V3 forward propagation (FPROP) in FP8, that single change—doubling MUFU.EX2 without changing FP8 Tensor Core FLOPs—boosted attention FPROP throughput by ~35%. That 35% jump is the clearest possible proof of how badly B200 was SFU-starved in FP8.

Now look at what 32 MUFU.EX2 ops/clk/SM on B300 actually buys you against the precision ladder:

• For a $128 \times 128 \times 128$ tile on B300 (sm_103), $T_{\text{exp}} = \frac{128^2}{32} = \mathbf{512\text{ cycles}}$.
• In BF16 ($T_{\text{MMA}} = 1{,}024\text{ cycles}$), B300 merely restores the 2:1 MMA-to-SFU cushion that Hopper H100 already had in 2022.
• In FP8 ($T_{\text{MMA}} = 512\text{ cycles}$), $T_{\text{exp}} = 512\text{ cycles}$ is a zero-margin 512-vs.-512 tie: the SFU must stay 100% utilized on every single clock cycle just to keep the FP8 Tensor Cores from stalling, leaving zero slack for row-max, row-sum, or tcgen05.ld/tcgen05.st latency.
• In NVFP4—where B300 sm_103 also widened dense Tensor Core reduction depth from $K=64$ to $K=96$, pushing dense NVFP4 from $10\text{ PFLOPS}$ to $15\text{ PFLOPS}$ ($\sim 49{,}152\text{ ops/clk/SM}$ at $K=96$, or $32{,}768$ at $K=64$)—$T_{\text{MMA}}$ is $171\text{ to }256\text{ cycles}$ against $T_{\text{exp}} = \mathbf{512\text{ cycles}}$. Even on B300, NVFP4 Tensor Cores still outrun the doubled SFU by $2\times$ to $3\times$.

Now Vera Rubin (sm_107). Start with what Rubin did to the Tensor Cores:

• By moving PCIe, NVLink 6, and NVLink-C2C off the compute dies onto separate I/O dies connected via NV-HBI, Rubin expands the SM count by $1.4\times$ (from 160 SMs on B300 to 224 SMs on Rubin, with 896 Tensor Cores) and raises clock frequency by $\sim 25\%$.
• Rubin doubles the per-cycle $K$-dimension throughput of the Tensor Core for FP8 (and expands NVFP4 from $K=96$ to $K=128$). Across the GPU, dense BF16 rises from $2.25\text{ PFLOPS}$ (B300) to $4.0\text{ PFLOPS}$ ($1.78\times$), dense FP8 jumps from $5.0\text{ PFLOPS}$ (B300) to $17.5\text{ PFLOPS}$ ($3.5\times$ chip-wide, or $2.0\times$ per Tensor Core per clock; NVIDIA’s rack page gives 1.26 EF FP8 training for 72 GPUs), and NVFP4 reaches $35\text{ PFLOPS}$ dense training / $50\text{ PFLOPS}$ inference.

Then read the SFU row of NVIDIA’s own Rubin compute table (Inside the NVIDIA Rubin Platform: Six New Chips, One AI Supercomputer). It lists “Softmax acceleration (SFU EX2 Ops/Clk/SM for FP32 | FP16)” as 16 for Blackwell and 32 | 64 for Rubin. B300 already runs FP32 MUFU.EX2 at 32 per clock per SM (the GB300 mufu_ex2_bench numbers above). So Rubin’s FP32 exponential rate per SM is the same as B300’s. The new doubling is the FP16 EX2 path.

Do the per-SM per-clock arithmetic for a $128 \times 128 \times 128$ tile on Rubin. Per-SM FP8 throughput doubles to $32{,}768\text{ ops/clk/SM}$, so $T_{\text{MMA}}(\text{FP8}) = 256\text{ cycles}$. NVFP4 reaches $65{,}536\text{ ops/clk/SM}$, so $T_{\text{MMA}}(\text{NVFP4}) = 128\text{ cycles}$. Against that:

• With an FP32 exponential (32 per clock): $T_{\text{exp}} = 128^2 / 32 = \mathbf{512\text{ cycles}}$. FP8 is now $2\times$ SFU-bound and NVFP4 $4\times$. That is worse than B300, where FP8 was a 512-vs-512 tie. Rubin went back to B200’s ratios.
• With an FP16 exponential (64 per clock): $T_{\text{exp}} = \mathbf{256\text{ cycles}}$. FP8 returns to a zero-slack $256\text{-vs.-}256$ tie and NVFP4 stays $2\times$ SFU-bound.

So the second doubling helps only kernels that move the exponential itself to half precision. That is a numerics decision, not a free speedup. The scaled logit and its exponential then carry FP16’s 10-bit mantissa and narrower range instead of FP32’s, before $\mathbf{P}$ is rounded again to BF16, FP8 or FP4 for the $\mathbf{P}\mathbf{V}$ MMA. It is the same kind of trade ZYFA’s Lean P makes in software on AMD (§2.2): precision that the destination format will throw away anyway is spent on throughput instead. On Rubin, NVIDIA built that trade into the SFU and left the FP32 path where B300 had it.

1.2.1 What about the vector units, Rubin CPX, and Rubin Ultra?

The software escape hatch is narrowing too. FA4’s way around the SFU (§2.1) moves 10%–25% of the exponentials onto FP32 FMA pipes. NVIDIA’s same Rubin blog publishes the FP32 vector rates: 67 TFLOPS (Hopper), 80 (Blackwell), 130 (Rubin). Against dense FP8 matrix throughput of roughly 2 PF (H100), 4.5 PF (B200) and 17.5 PF (Rubin), the ratio of FP8 matrix FLOPs to FP32 vector FLOPs goes from about 30:1 to 56:1 to 135:1. From Blackwell to Rubin the FMA pipes grew $1.6\times$ while FP8 matrix grew about $3.9\times$. The idle FMA capacity that FA4 borrows shrinks per matrix FLOP each generation, so polynomial emulation buys less on Rubin than it did on B200.

Rubin CPX was NVIDIA’s dedicated answer, and it was pulled. In September 2025 NVIDIA announced Rubin CPX, a monolithic GDDR7 GPU for the context (prefill) phase with up to 30 PF of NVFP4 and “3x faster attention capabilities compared with NVIDIA GB300 NVL72 systems.” Third-party notes from the launch describe the mechanism as “3x exponent operations” in “attention acceleration cores.” That is a chip built around the softmax wall, for exactly the long-context prefill this article is about. It never shipped. CPX was absent from the six-chip Rubin launch at CES 2026 and from the GTC 2026 roadmap. Ian Buck told reporters “we’ve pulled CPX. It’s still a good idea,” and NVIDIA said prefill work goes back to the Rubin GPU, with Groq 3 LPX racks taking the decode-specialist slot. A revived CPX with 168 GB of HBM4 was reported in August 2026 by one analyst (Ming-Chi Kuo) and has not been confirmed. As shipped, the Rubin generation’s context processing runs on a GPU whose FP32 EX2 rate per SM equals B300’s.

Rubin Ultra changes package size, not the per-SM ratio, as far as anything is published. NVIDIA’s only official Rubin Ultra spec is from March 2025: four reticle-sized dies, 100 PF FP4 and 1 TB of HBM4E per package, for 2H 2027 in Kyber racks. In June 2026 SemiAnalysis reported the quad-die design was cut to two dies; NVIDIA has not confirmed that or published a replacement. Either way, NVIDIA has published no SM-level numbers for Rubin Ultra: no SFU rate, no per-SM Tensor Core rate. Adding dies multiplies matrix units and SFUs together. The tile ledger above only moves if the SM itself changes, so there is nothing yet to suggest Rubin Ultra fixes it.

A note on names, since sources mix them:

• Rubin is the shipping GPU (two dies, 224 SMs, 288 GB HBM4). SemiAnalysis calls it R200; “R100” appears in some early coverage but is not an NVIDIA product name.
• VR200 NVL144 was the GTC 2025 name for the Rubin rack, counting dies. NVIDIA went back to counting packages, so the same rack ships as Vera Rubin NVL72.
• Rubin Ultra (R300 in some trackers) pairs with Kyber. “NVL576” meant one Kyber rack counted by dies in 2025; in the GTC 2026 roadmap, Kyber NVL144 is one rack and Vera Rubin Ultra NVL576 is eight racks.
• Feynman (2028) has no published performance numbers.

1.3 Does AMD hit the same wall on CDNA 4 (gfx950) and CDNA 5 (gfx1250)?#

This asymmetric scaling wall is not an artifact of NVIDIA’s SM layout. It hits AMD’s Instinct line on both CDNA 4 (gfx950, MI355X) and the newly disclosed CDNA 5 (gfx1250, MI455X / MI450)—first through CDNA 4’s static wave-scheduling constraints, and then through CDNA 5’s $4\times$ microscaling matrix jump.

1.3.1 Why does AITER leave the MI355X (gfx950) matrix pipe idle 31.7% of the time?

On MI355X (gfx950), each Compute Unit (CU) contains four Wave64 SIMD/Matrix Core quadrants. Moving from BF16 to FP8 e4m3 doubles the throughput of the Matrix Fused Multiply-Add (MFMA) units. However, the softmax between $\mathbf{Q}\mathbf{K}^\top$ and $\mathbf{P}\mathbf{V}$ executes on the Vector ALU (VALU), whose throughput does not change when the GEMMs switch from BF16 to FP8.

In Zyphra’s ZYFA technical report (Zyphra FlashAttention on AMD, released this week for MI355X on ROCm 7.2.3), the authors traced the exact cycle breakdown for a single SIMD processing a tile of $M = 32$ queries against $N = 128$ keys at head dimension $d = 128$:

• FP8 Matrix Core work ($\mathbf{Q}\mathbf{K}^\top$ + $\mathbf{P}\mathbf{V}$): $1{,}024\text{ cycles}$.
• Exact FP8 Softmax work on the Vector ALU (VALU): $896\text{ vector cycles}$.

Look at those two numbers: $896\text{ vector cycles}$ of softmax against $1{,}024\text{ matrix cycles}$ of FP8 GEMM ($87.5\%$ of the entire matrix budget).

On paper, a naive roofline model would say: $\max(1{,}024,\; 896) = 1{,}024\text{ cycles}$, so if you overlap VALU softmax for tile $k$ with MFMA matrix multiplication for tile $k+1$, the 896 vector cycles should hide completely behind the 1,024 matrix cycles.

Why doesn’t that happen in silicon on MI355X? Because unlike NVIDIA’s tcgen05 (which runs asynchronously in a separate TMEM domain and signals mbarrier completion in hardware), MI355X does not track matrix instruction completion in hardware. On CDNA 4, both MFMA instructions and VALU instructions are issued by the same wave schedulers and share the vector register file (split into 256 architectural VGPRs and 256 accumulation VGPRs), with no hardware scoreboard for asynchronous MFMA completion. Consequently, overlapping the 896 cycles of VALU softmax with the 1,024 cycles of MFMA requires static instruction-order interleaving across waves sharing a SIMD.

Because $\mathbf{P}_k = \operatorname{softmax}(\mathbf{S}_k)$ has a hard data dependency on the completion of the $\mathbf{Q}\mathbf{K}_k^\top$ MFMA, and the $\mathbf{P}_k\mathbf{V}_k$ MFMA has a hard data dependency on the completion of the softmax VALU instructions (plus cross-lane row-max and row-sum reductions and VGPR-to-LDS staging), $1{,}024 - 896 = 128\text{ cycles}$ of slack is far too small to absorb pipeline bubbles:

• In AITER (AMD’s default production kernel library), Zyphra’s thread traces show that the MI355X matrix pipe sits completely idle for 31.7% of cycles, capping 32K-token FP8 attention at $1{,}944\text{ TFLOP/s}$ (39% of the MI355X FP8 datasheet peak) on $d=128$ and just $1{,}237\text{ TFLOP/s}$ on DeepSeek’s MLA shape ($192/128$).
• Even when Zyphra implemented a hand-scheduled two-wave ping-pong schedule (where two waves share each SIMD, one issuing MFMA instructions while the other computes softmax probabilities on the VALU, and then swapping), running exact FP8 softmax (896 cycles) only kept the matrix pipe busy for 60% of steady-state cycles.

1.3.2 What changes on CDNA 5 (gfx1250, Instinct MI455X)?

How does AMD’s next-generation CDNA 5 architecture on the Instinct MI455X / MI450 (gfx1250)—detailed in AMD’s Introducing AMD CDNA 5 Architecture whitepaper and the preliminary CDNA5 Instruction Set Architecture Reference Guide (July 2026)—change this equation?

In CDNA 5, AMD made the most sweeping microarchitectural shift since the original GCN/CDNA split, converging several RDNA 4 (gfx12) execution concepts with datacenter matrix and memory engines across 8 TSMC N2 Accelerator Complex Dies (XCDs):

1. From Wave64 CUs and MFMA to Wave32 WGPs and WMMA:
   • The 32 Compute Units per XCD of MI355X are replaced by 32 Work-Group Processors (WGPs) per XCD (2 Shader Engines $\times$ 16 WGPs per XCD = 256 active WGPs chip-wide, running at up to 2,400 MHz; Whitepaper pp. 6, 21).
   • Each WGP (“Double Compute Unit”) contains four 32-lane SIMD32 units and four scalar ALUs, and operates exclusively in Wave32 mode (dropping Wave64 entirely; ISA Ch. 1, p. 3), supporting up to 64 resident waves per WGP (16 per SIMD32; Whitepaper p. 10) and up to 32 workgroups per WGP.
   • CDNA’s old Wave64 MFMA instructions are replaced by Wave32 Wave Matrix Multiply-Accumulate (WMMA) instructions (V_WMMA_* and block-scaled V_WMMA_SCALE*; ISA §7.12, Table 43), operating on $16 \times 16$ output tiles (plus $32 \times 16$ for FP4).
2. Transcendental and Vector Upgrades ($2\times$ Throughput + Native BF16 & TANH):
   • Just as NVIDIA doubled MUFU.EX2 in B300, AMD’s CDNA 5 whitepaper (p. 7) explicitly states that existing transcendental operations have twice ($2\times$) the throughput of the same instructions on the Instinct MI355X, specifically targeting activation and softmax bottlenecks. (The preliminary CDNA 5 ISA Reference Guide does not publish a per-cycle instruction-rate table for transcendental opcodes, though LLVM gfx1250 scheduling definitions also show FP32 transcendental latency dropping from 9 cycles on RDNA 4 to 8 cycles on gfx1250.)
   • The ISA also adds native BF16 transcendental instructions (V_EXP_BF16, V_LOG_BF16, V_RCP_BF16, V_RSQ_BF16, V_SQRT_BF16, V_SIN_BF16, V_COS_BF16), hardware hyperbolic tangent (V_TANH_F32, V_TANH_F16, V_TANH_BF16, directly accelerating attention logit soft-capping), and single-lane pseudo-scalar transcendental ALU ops (V_S_EXP_F32, V_S_EXP_F16, V_S_LOG_F32, V_S_LOG_F16, V_S_RCP_F32, V_S_RCP_F16, V_S_RSQ_F32, V_S_RSQ_F16, V_S_SQRT_F32, V_S_SQRT_F16; ISA §7.3 & §7.10).
   • For general vector math, CDNA 5 retains packed 64-bit FP32 operations (V_PK_FMA_F32, V_PK_ADD_F32, V_PK_MUL_F32; ISA §7.7) and adds RDNA-style dual-issue VALU instructions (VOPD 64-bit 2-source and a new VOPD3 96-bit 3-source encoding; ISA §7.8), delivering 315 TFLOPS of peak vector/matrix FP32 and vector FP16 throughput across the GPU (Whitepaper p. 21).
3. The $4\times$ Matrix Scaling Leap in FP8/MXFP8 and MXFP4: Now compare that $2\times$ transcendental speedup against what CDNA 5 did to the matrix engines per clock (Whitepaper pp. 4, 7, 21 & ISA §7.12, Table 43):
   • Matrix FP16 / BF16 (V_WMMA_F32_16X16X32_{F16,BF16}, $K=32$): Doubled $2\times$ per clock over MI355X to 5.03 PFLOPS ($8{,}192\text{ FLOPs/clk/WGP}$, or $2{,}048\text{ FLOPs/clk/SIMD32}$ at $2{,}400\text{ MHz}$).
   • OCP MXFP8, FP8, and MXFP6 (V_WMMA_F32_16X16X128_{FP8,BF8} and V_WMMA_SCALE{,16}_F32_16X16X128_F8F6F4, $K=128$): Quadrupled $4\times$ per clock over MI355X to 20.13 PFLOPS ($32{,}768\text{ FLOPs/clk/WGP}$, or $8{,}192\text{ FLOPs/clk/SIMD32}$). Notice that on MI455X, FP8/MXFP8 matrix peak ($20.13\text{ PF}$) is $4\times$ the BF16 matrix peak ($5.03\text{ PF}$)—not $2\times$—because CDNA 5 expanded the FP8/F6 WMMA reduction depth to $K = 128$ ($4\times$ BF16’s $K = 32$).
   • OCP MXFP4 (V_WMMA_SCALE{,16}_F32_32X16X128_F4, $32 \times 16 \times 128$): Quadrupled $4\times$ per clock over MI355X to 40.26 PFLOPS ($65{,}536\text{ FLOPs/clk/WGP}$, or $16{,}384\text{ FLOPs/clk/SIMD32}$—an $8\times$ jump over MI455X’s own BF16 matrix rate).

Look at what that $4\times$-vs.-$2\times$ scaling asymmetry does to the $32 \times 128$ SIMD tile ($d=128$, $W_{\text{MMA}} = 4 \times 32 \times 128 \times 128 = 2{,}097{,}152\text{ FLOPs}$) when moving from MI355X (gfx950) to MI455X (gfx1250):

• In BF16, both matrix throughput and transcendental throughput doubled ($2\times$) per WGP/CU over MI355X, keeping the BF16 matrix-to-softmax ratio unchanged.
• In FP8 / MXFP8, however, per-WGP matrix throughput scaled $4\times$ ($1{,}024\text{ matrix cycles}$ on MI355X $\to$ $256\text{ matrix cycles}$ on MI455X at $8{,}192\text{ FLOPs/clk/SIMD32}$), while only transcendental throughput scaled $2\times$. ZYFA’s 896-cycle exact softmax on MI355X also contains non-transcendental VALU work (row max, row sum, scaling, FP8 conversion), and AMD does not publish per-cycle rates for either, so the same softmax on MI455X lands somewhere between $448$ cycles (if every one of the 896 cycles were transcendental) and $896$ cycles (if none were). Either end sits above the $256$-cycle FP8 matrix budget.
• And in MXFP4 ($16{,}384\text{ FLOPs/clk/SIMD32}$), both GEMMs for a $32 \times 128$ tile finish in $128\text{ matrix cycles}$.

The same squeeze shows up in AMD’s datasheet totals. MI455X lists 315 TFLOPS of vector FP32 against 20.13 PFLOPS of dense FP8 matrix, a 64:1 ratio of matrix to vector FLOPs. On NVIDIA’s published numbers, Rubin sits at about 135:1 (§1.2.1). AMD left more vector headroom per matrix FLOP than NVIDIA did, but the software exponential tricks of §2 still run on a vector side that grew more slowly than the matrix side.

1.3.3 What about the rest of the MI400 family (MI450, MI430X, MI440X)?

AMD’s naming moved during 2026, so here is what each name refers to and whether it changes the argument:

• MI450 is the earlier label for the whole generation. AMD’s Q2 2026 results still say “MI450 Series” for its Helios deals, and this week’s SGLang and vLLM notes say “MI450” for gfx1250 hardware. At the July 23, 2026 launch AMD defined the MI400 series as two GPUs, MI455X and MI430X.
• MI455X is the Helios rack GPU and the only MI400 part with a full public datasheet. Everything in §1.3.2 refers to it. The datasheet also shows how far CDNA 5 was split by precision: MI455X lists 5 TFLOPS of FP64, matrix or vector.
• MI430X is the HPC and sovereign-AI part: 288 TFLOPS of hardware FP64, 432 GB of HBM4, availability in 2027, and systems announced for ORNL’s Discovery and two European machines. AMD has not published its low-precision matrix rates or its transcendental rate. It does not change the attention ledger, and nothing published lets you compute one for it.
• MI440X was announced at CES 2026 as an eight-GPU enterprise server built from a low-precision CDNA 5 subset, like MI455X. It was not in the July launch definition, and AMD has published no compute, memory or availability figures for it.

So for attention on AMD’s 2026–27 parts, the MI455X numbers above are the only ones there are. The CES framing does support one point: AMD is now building different CDNA 5 dies or configurations per precision envelope (FP64 for MI430X, FP4–BF16 for MI440X and MI455X). That is the same choice NVIDIA made in Rubin’s FP64 vector rate, which fell from 40 TFLOPS on Blackwell to 33. Vendors are trading high-precision units for low-precision matrix area. On the AI parts, exponentials got at most $2\times$ out of that trade (on MI455X, and only for FP16 on Rubin), while FP8 matrix got $4\times$ (MI455X over MI355X) or $3.5\times$ (Rubin over B300).

Put the cycle budgets across NVIDIA and AMD generations side by side for a standard attention tile ($d = 128$, normalized to a $128 \times 128$ tile on an NVIDIA SM / AMD 4-SIMD WGP, and shown per-SIMD for AMD’s $32 \times 128$ tile):

*Note: B300 (sm_103) supports dense $K=96$ for NVFP4 ($1.5\times$ B200’s $K=64$), and Rubin (sm_107) expands NVFP4 to $K=128$ ($2\times$ B200’s $K=64$) while doubling FP8 $K$ depth. Rubin’s EX2 rates are NVIDIA’s published 32 (FP32) and 64 (FP16) ops/clk/SM. $^\dagger$For MI455X (gfx1250), per-WGP matrix rates are derived directly from AMD’s published peaks ($5.03\text{ PF}$ BF16, $20.13\text{ PF}$ FP8/MXFP8, $40.26\text{ PF}$ MXFP4 at $2{,}400\text{ MHz}$ across 256 WGPs); exact per-cycle transcendental instruction rates are not published in the preliminary CDNA 5 ISA guide, so $448\text{--}896$ brackets MI355X’s 896-cycle exact softmax under the whitepaper’s stated $2\times$ transcendental throughput: $448$ if all of those cycles were transcendental, $896$ if none were.

Conclusion: Across both NVIDIA (sm_100 $\to$ sm_107) and AMD (gfx950 $\to$ gfx1250), the transition from BF16 to FP8 and FP4 broke the foundational assumption of early FlashAttention kernels—that non-matmul softmax work is cheap enough to hide effortlessly behind slow matrix multiplications. Once matrix units take 128 to 1,024 cycles per tile while exact softmax takes 256 to 1,024 cycles, hardware SFU/transcendental doublings alone only tread water. They fall further behind when FP8/FP4 matrix depth quadruples (CDNA 5) or when the doubling covers only FP16 exponentials (Rubin). To keep the matrix engines fed, software had to rewrite how exponentials and online normalization are computed.

2. How does software get around the exponential unit, and what does it cost?#

When hardware transcendental units cannot keep up with low-precision matrix engines, kernel engineers have only two mathematical degrees of freedom:

1. Replace hardware exp2 calls with software bit-manipulation and polynomial FMA instructions that execute on underutilized ALU/FMA pipes; and
2. Eliminate redundant vector rescaling and reduction operations in the online softmax loop.

Both FlashAttention-4 (on NVIDIA Blackwell B200) and Zyphra’s ZYFA (on AMD MI355X) pull these exact levers—and AMD’s CDNA 5 (gfx1250) ISA bakes several of those software workarounds directly into silicon.

2.1 How does FlashAttention-4 do it on Blackwell?#

Look at how FlashAttention-4 (Zadouri et al., 2026) attacks the $1{,}024\text{-cycle}$ MUFU.EX2 wall on Blackwell B200 (sm_100).

First, FA4 partitions each 512-thread CTA into four specialized 128-thread warpgroups (WG0–WG3 in the CuTe-DSL / TIRx reference implementation) and uses PTX setmaxnreg instructions to dynamically redistribute the SM’s 65,536 32-bit registers to the warps that actually hold the softmax rows:

• WG3 (Orchestration: Warp 0 = tcgen05 MMA, Warp 1 = TMA Load, Warp 2 = TMA Store): Because tcgen05.mma is launched by a single elected thread (elect_sync) and writes its accumulator directly to Tensor Memory (TMEM) rather than registers, WG3 sheds registers via setmaxnreg.dec.sync.aligned.u32(48) (48 registers/thread).
• WG2 (Output Correction & Epilogue): Sheds registers via setmaxnreg.dec.sync.aligned.u32(64) (64 registers/thread).
• WG0 and WG1 (Ping-Pong Softmax for Q Stages 0 and 1): Acquire the freed registers via setmaxnreg.inc.sync.aligned.u32(200) (200 registers/thread, totaling $128 \times (200 + 200 + 64 + 48) = 65{,}536$ registers). That 200-register budget allows each of the 128 threads in a softmax warpgroup to load and hold one entire 128-element FP32 row of $\mathbf{S}$ in private registers (s_chunk_buf), eliminating inter-thread warp shuffles during the row-max and row-sum reductions.

Even with zero warp-shuffle overhead and WG0/WG1 alternating on two query tiles ($Q^H$ and $Q^L$), the hardware MUFU.EX2 pipe is still capped at $16\text{ ops/clk/SM}$. While MUFU.EX2 is 100% saturated during softmax, the SM’s general-purpose FP32 FMA pipelines sit mostly idle.

FA4 exploits those idle FMA pipelines by implementing a hybrid hardware/software exponential:

1. Cody–Waite Range Reduction ($2^x = 2^{\lfloor x \rfloor} \cdot 2^{x_{\text{frac}}}$): For any scaled logit $x = (s_{ij} - m_i)\log_2 e / \sqrt{d}$ (clamped at $-127$ to prevent underflow), decompose $x$ into its integer floor $n = \lfloor x \rfloor \in \mathbb{Z}$ and fractional part $x_{\text{frac}} = x - \lfloor x \rfloor \in [0, 1)$. In hardware, FA4 extracts $\lfloor x \rfloor$ in two FP32 instructions using the magic-number round-down trick: add $2^{23} + 2^{22} = 12{,}582{,}912.0$ to force the fractional bits out of the 23-bit IEEE-754 mantissa, then subtract it back in round-toward-$-\infty$ mode.
2. Degree-3 Sollya Minimax Polynomial on Packed f32x2 FMA Units: On the unit interval $x_{\text{frac}} \in [0, 1)$, FA4 approximates $2^{x_{\text{frac}}}$ with a degree-3 polynomial whose coefficients are optimized via the Sollya package to minimize maximum relative error: $$2^{x_{\text{frac}}} \approx p_0 + x_{\text{frac}}\bigl(p_1 + x_{\text{frac}}(p_2 + p_3 x_{\text{frac}})\bigr), \qquad \begin{cases} p_0 = 1.0, & p_1 = 0.69514614, \\ p_2 = 0.22756439, & p_3 = 0.07711909 \end{cases}$$ Using Horner’s rule, evaluating this cubic polynomial on pairs of logits takes just three packed fma.rn.ftz.f32x2 instructions:

   // Horner evaluation of (((p3 * r) + p2) * r + p1) * r + 1.0 on packed f32x2
   fma.rn.ftz.f32x2  l10, l9, l6, l5;
   fma.rn.ftz.f32x2  l10, l10, l9, l4;
   fma.rn.ftz.f32x2  l10, l10, l9, l3;

3. IEEE-754 Exponent Bit-Shift Reconstruction: Because $2^{x_{\text{frac}}} \in [1, 2)$, its IEEE-754 FP32 representation already has biased exponent $127$ and holds the fractional mantissa bits in bits $[0..22]$. Multiplying by $2^{\lfloor x \rfloor}$ is therefore a single integer shift-and-add: shift the integer $\lfloor x \rfloor$ left by 23 bits into the FP32 exponent field ([23..30]) and add the bit representation of $2^{x_{\text{frac}}}$.

Why doesn’t FA4 run all 128 exponentials per row through FMA emulation? Because polynomial emulation requires extra temporary registers for intermediate values and polynomial coefficients, consumes more register-file read/write bandwidth, and has higher instruction latency than a single MUFU.EX2. Emulating 100% of the row would spill registers past the 200-register ceiling. Instead, FA4 tunes the split (emu_pairs): 75%–90% of each row executes on hardware MUFU.EX2, while 10%–25% executes in parallel on the FMA units, expanding effective exponential throughput just enough to break the MUFU bottleneck without a single register spill.

And what about numerical accuracy? Look at the error measurements across 4 million random inputs in $[0, 1)$ against an FP64 oracle (Table 2 of FlashAttention-4):

At the raw FP32 level, the 3-FMA cubic polynomial has a max relative error of $8.77 \times 10^{-5}$. Once cast to BF16 (or FP8) to form the operand $\mathbf{P}$ for $\mathbf{P}\mathbf{V}$, however, BF16’s 7-bit mantissa rounding error ($3.89 \times 10^{-3}$) is $44\times$ larger than the polynomial approximation error. After rounding to BF16, the degree-3 FMA polynomial matches hardware MUFU.EX2 to within 1 BF16 ULP on 99% of inputs and achieves the exact same mean relative error ($1.41 \times 10^{-3}$).

FA4 pairs this with a second non-matmul optimization: Conditional Online Softmax Rescaling. In standard FlashAttention-2 and -3, whenever KV block $j$ produces a candidate row maximum $r_c = \max(r_{\text{old}}, \operatorname{rowmax}(\mathbf{S}_j))$ strictly greater than the previous reference $r_{\text{old}}$, the kernel immediately rescales the running output accumulator $\mathbf{O}_{j-1}$ by $a_{\text{scale}} = 2^{(r_{\text{old}} - r_c)\alpha}$. On Blackwell, where $\mathbf{O}$ lives in TMEM (O_region), rescaling $\mathbf{O}_{j-1}$ is expensive: the Correction warpgroup (WG2) must read the $128 \times 128$ FP32 tile out of TMEM into registers (tcgen05.ld), multiply every element by $a_{\text{scale}}$, and write it back to TMEM (tcgen05.st) before the next $\mathbf{P}\mathbf{V}$ MMA can accumulate into $\mathbf{O}$.

FA4 notices that FP32 has a massive exponent range ($2^{-126}$ to $2^{127}$), so subtracting $r_{\text{old}}$ instead of $r_c$ does not threaten overflow unless the new maximum jumps by a large margin. Setting a threshold $\tau = \log_2(256) = 8.0$ (allowing unnormalized block weights up to $2^8 = 256$), FA4 modifies the online update rule to:

Each warp in WG2 checks its 32 rows simultaneously using vote.sync.any.pred (any_needs_rescale). Whenever all 32 rows stay within the $\tau = 8.0$ slack window ($a_{\text{scale}} = 1.0$), the warp skips the entire TMEM -> registers -> multiply -> TMEM correction pass and immediately signals the p_o_rescale barrier so the $\mathbf{P}\mathbf{V}$ MMA can proceed. Because the final epilogue normalizes $\mathbf{O}_{\text{final}}$ by the true accumulated denominator $\ell_{\text{final}}$ on the same reference scale, skipping intermediate rescales preserves exact mathematical equivalence while cutting vector rescaling operations by ~10×. Combined with its CuTe-DSL async pipeline, FA4 reaches 1,613 TFLOP/s in BF16 on B200 (71% hardware utilization, $1.1\times\text{--}1.3\times$ faster than cuDNN 9.13 and $2.1\times\text{--}2.7\times$ faster than Triton).

2.2 How does ZYFA do it on MI355X (gfx950), and what changes on CDNA 5 (gfx1250)?#

On AMD’s Instinct MI355X (gfx950), the problem is even sharper: in FP8 attention, the target probability matrix $\mathbf{P}$ consumed by the $\mathbf{P}\mathbf{V}$ MFMA instruction is not 16-bit BF16; it is 8-bit FP8 e4m3.

In AMD’s default AITER kernel, computing $\mathbf{P}$ follows the textbook path: evaluate the FP32 base-2 exponential on the vector ALU (v_exp_f32), scale it, and convert the FP32 result down to FP8 e4m3. Across a $32 \times 128$ tile, that exact FP8 softmax takes 896 vector cycles against 1,024 matrix cycles, starving the matrix pipe.

In ZYFA, Zyphra asked a first-principles question: if the destination of $2^x$ is an 8-bit e4m3 float with only 3 mantissa bits, why compute a 23-bit FP32 transcendental on the SFU/VALU only to throw away 20 bits of mantissa a cycle later?

Look at the binary layout of an unsigned FP8 e4m3 number:

• Bit 7 is the sign bit ($0$ for non-negative probabilities $p_{ij} > 0$).
• Bits $[6..3]$ are the 4-bit biased exponent $E \in \{0, \dots, 15\}$ (bias $b = 7$).
• Bits $[2..0]$ are the 3-bit mantissa $M \in \{0, \dots, 7\}$, representing the significand $1 + M/8$.

Now read that 8-bit e4m3 byte directly as an unsigned integer $I \in \{0, \dots, 127\}$:

Because an e4m3 float has 3 mantissa bits ($2^3 = 8$ quantization steps per octave), its integer code $I$ grows by exactly 8 every time the floating-point value doubles ($x \mapsto x + 1$). And on the fractional interval $x_{\text{frac}} \in [0, 1)$, the piecewise-linear secant $1 + x_{\text{frac}}$ already approximates $2^{x_{\text{frac}}}$ to within a maximum error of $1 - \frac{1 + \ln(\ln 2)}{\ln 2} \approx 0.086$, which is smaller than a single e4m3 mantissa step ($1/8 = 0.125$).

This observation—a modern FP8 incarnation of Schraudolph’s (1999) IEEE-754 bit-level exponential—gives ZYFA its core mechanism, Lean P:

1. One Multiply-Add + One Integer Conversion: Instead of calling the hardware exponential unit and then casting FP32 to FP8, Lean P computes the integer bit pattern of the e4m3 probability directly from the scaled logit $x$ using a single floating-point multiply-add ($8x + \text{const}$) and a single float-to-integer conversion. The low byte of that integer is the valid e4m3 floating-point probability.
2. 57% Reduction in Softmax Cycles: Eliminating the transcendental instruction and separate float-to-FP8 conversion slashes softmax vector work per $32 \times 128$ block from 896 vector cycles down to 384 vector cycles ($2.33\times$ fewer VALU cycles).
3. Two-Wave Ping-Pong Schedule: With softmax compressed to 384 vector cycles against 1,024 matrix cycles, ZYFA pairs two waves on each SIMD in a ping-pong schedule: while Wave 0 issues MFMA matrix instructions, Wave 1 computes Lean P probabilities on the VALU, and then they swap. Because 384 vector cycles now fit comfortably inside the 1,024-cycle MFMA window, steady-state matrix-pipe utilization jumps from 60% (with exact P) to 86% (with Lean P).
4. Power-of-2 Block Scaling and Rounded Row Sum via MFMA: To prevent small probabilities from underflowing e4m3’s narrow dynamic range and to eliminate quantization bias in the denominator, ZYFA scales each 64-key block of probabilities by a power of 2 before the Lean P rounding step—a factor that CDNA 4’s block-scaled MFMA instruction removes in hardware at zero extra cycle cost—and normalizes by the sum of the rounded FP8 probabilities (computed by issuing 1 extra matrix instruction per block rather than burning VALU cycles on a vector reduction).

Look at what Lean P + block scaling + rounded row sum achieves on MI355X:

• Throughput: On 32K-token Gaussian benchmarks, ZYFA reaches 2,672 TFLOP/s in FP8 (53% of MI355X’s FP8 datasheet peak, vs. AITER’s 1,944 TFLOP/s / 39%—a 1.37× gain) at $d=128$, and on DeepSeek’s MLA shape ($192/128$) it reaches 2,464 TFLOP/s vs. AITER’s 1,237 TFLOP/s (1.99×).
• Accuracy (62%–78% Lower Error than Exact Hardware Exp Without Block Scaling): Against an FP64 reference, even though Lean P uses a 1-FMA bit-pattern approximation for $2^x$, ZYFA’s power-of-2 block scaling and rounded row sum cut total FP8 attention error to $0.22\times$ AITER’s on Mistral-7B (78% lower error) and $0.38\times$ on Qwen3-8B (62% lower error). Across the full Mistral-7B model, KL divergence from FP64 drops from $5.9 \times 10^{-3}$ (AITER) to $3.4 \times 10^{-3}$ (ZYFA), with bitwise-identical outputs across 7,200 consecutive kernel launches.

How does CDNA 5 (gfx1250, MI455X) change the AITER / ZYFA pipeline?

When porting AITER, ZYFA, or FlyDSL attention kernels from MI355X (gfx950) to CDNA 5 (gfx1250, MI450 / MI455X)—where initial gfx1250 enablement for gpt-oss-120b, DeepSeek-R1, DeepSeek-V4-Flash, MiniMax-M3, and FlyDSL / AITER tgemm landed in ROCm 10.0 / 10.1, vLLM v0.33.0 (#58761, #58800), and SGLang v0.5.23 (#41988, #42699, #44671, #45304, #45365, #45758, #46366) this week—six ISA changes directly alter the attention cycle ledger:

1. 4× Addressable Unified VGPRs (1,024 VGPRs/Wave via S_SET_VGPR_MSB): On CDNA 4, a wave was restricted to 256 architectural VGPRs plus 256 MFMA-only AccVGPRs. On CDNA 5 (ISA §3.3.2.1 & §3.3.2.3; Coelacanth’s Dream), the split AccVGPR file is eliminated and a single Wave32 can allocate up to 1,024 unified 32-bit VGPRs (in blocks of 16). To address registers beyond V255 without breaking 8-bit register fields in standard encodings, CDNA 5 introduces S_SET_VGPR_MSB (and MODE bits DST_VGPR_MSB, SRC0_VGPR_MSB, SRC1_VGPR_MSB, SRC2_VGPR_MSB), which supplies the upper 2 bits ([9:8]) of each operand’s 10-bit VGPR address. (While RDNA 4’s dynamic VGPR allocation instruction is a no-op on gfx1250, static allocation of up to 1,024 VGPRs per wave allows an attention wave to keep large $\mathbf{S}$, $\mathbf{P}$, and $\mathbf{O}$ subtiles resident in registers without AccVGPR-to-VGPR copies or LDS spills.)
2. Hardware VALU/WMMA Co-Execution, S_DELAY_ALU, and 16 Named Barriers per Workgroup: On MI355X, overlapping VALU softmax with MFMA required brittle static instruction-order interleaving across waves. On CDNA 5 (ISA §5.7.2.1 & §5.8), 16-bit and smaller XDL WMMA instructions are tracked in S_DELAY_ALU as TRANS operations (DEP codes 5–7), allowing normal VALU instructions and WMMA instructions to co-execute. By default, the arbiter stalls a wave for the multi-cycle duration of its WMMA so peer waves on the same SIMD can co-execute VALU softmax instructions; setting SCHED_MODE[2] (DISABLE_XDL_ARB_STALL) instead allows a single wave to issue multiple WMMA instructions back-to-back and immediately issue independent VALU instructions while the matrix ops execute. Combined with 16 hardware named barriers per workgroup (S_BARRIER_INIT, S_BARRIER_JOIN, S_BARRIER_SIGNAL, S_BARRIER_WAIT, allocated from 64 per WGP; ISA §5.6.2) and LDS split barriers (DS_ATOMIC_ASYNC_BARRIER_ARRIVE_B64; ISA §11.2.2), CDNA 5 supports true producer/consumer wave specialization.
3. Dual-Issue VOPD / VOPD3 and Packed V_PK_FMA_F32: Because MI455X quadruples FP8/MXFP8 and MXFP4 matrix throughput ($4\times$) while doubling transcendental throughput ($2\times$), software exponential paths like ZYFA’s Lean P (1 FMA + 1 int cast) or FA4’s Horner cubic polynomial become even more essential on CDNA 5. CDNA 5 accelerates those non-transcendental softmax instructions with both 64-bit packed FP32 math (V_PK_FMA_F32, V_PK_MUL_F32, V_PK_ADD_F32) and dual-issue VALU encodings (VOPD and the new 96-bit 3-source VOPD3, which removes VOPD’s even/odd destination bank restriction; ISA §7.8.1).
4. Pseudo-Scalar Transcendentals (V_S_EXP_F32, V_S_RCP_F32) and Native V_TANH_* / V_EXP_BF16: In Wave32 attention, when a row or sub-tile maximum $m_i$ and denominator $\ell_i$ are reduced across lanes, computing the per-row rescaling factor $a_{\text{scale}} = 2^{(r_{\text{old}} - r_c)\alpha}$ or final reciprocal $1/\ell_i$ previously burned a full 32-lane vector transcendental instruction. CDNA 5’s pseudo-scalar transcendental ops (V_S_EXP_F32, V_S_EXP_F16, V_S_RCP_F32, V_S_RCP_F16; ISA §7.10) read an SGPR/constant and write the result directly to an SGPR in a single lane, while V_TANH_{F32,F16,BF16} and V_EXP_{F16,BF16} accelerate half-precision activation and soft-capping steps.
5. Hardware Block-Scaled Conversions (V_CVT_SCALEF32_PK*) + 4-DWORD Block-Scaled V_WMMA_SCALE*: ZYFA’s biggest accuracy win on MI355X came from computing a per-block power-of-2 scale in software before casting $\mathbf{P}$ to FP8. CDNA 5 turns block-scaled conversion into a native VALU instruction family: V_CVT_SCALEF32_PK8_{FP8,BF8,FP4}_{F32,F16,BF16} and V_CVT_SCALEF32_PK16_{FP6,BF6}_{F32,F16,BF16} (plus stochastic-rounding variants V_CVT_SCALEF32_SR_*; ISA §7.3 & §7.6.3), which feed directly into CDNA 5’s 4-DWORD VOP3PX2 block-scaled matrix instructions (V_WMMA_SCALE_F32_16X16X128_F8F6F4 for 32-element blocks and V_WMMA_SCALE16_F32_16X16X128_F8F6F4 for 16-element blocks, supporting E8M0, E5M3, and fractional E4M3 scales; ISA §7.12.6).
6. Tensor Data Mover (TDM), Async Global-to-LDS Loads, Cluster Multicast, and Transpose Loads (DS_LOAD_TR* / GLOBAL_LOAD_TR*): Each WGP on MI455X includes a dedicated Tensor Data Mover (TDM) (TENSOR_LOAD_TO_LDS, TENSOR_STORE_FROM_LDS, tracked by TENSORcnt; ISA §10.11) that copies up to 5D tiles between HBM4 and the 320 KB LDS without using VGPRs, optionally inserting bank-conflict-free LDS padding (pad_enable), broadcasting tiles across up to 16 WGPs in a Shader Engine cluster (workgroup_mask), or executing indexed row gathers (up to 16 16-bit or 8 32-bit row indices per descriptor). Once in LDS (or directly from global memory), DS_LOAD_TR16_B128, DS_LOAD_TR8_B64, DS_LOAD_TR6_B96, and DS_LOAD_TR4_B64 (and GLOBAL_LOAD_TR*; ISA §10.9 & §11.2.4) load and transpose 16-bit, 8-bit, 6-bit, or 4-bit matrix tiles into VGPRs in hardware, eliminating the VALU/LDS transpose instructions that previously competed with softmax.

2.3 Why did the faster kernel run at a lower clock on MI355X?#

When Zyphra measured ZYFA and AITER on captured real-model activations at sustained clocks (16K context on Mistral-7B, Qwen3-8B, and DeepSeek-V2-Lite), they uncovered one of the most important hardware-software interactions in modern low-precision inference.

While ZYFA achieved a 1.99× speedup on synthetic Gaussian inputs (and 1.76× on DeepSeek-V2-Lite’s real MLA activations), its wall-clock speedup on real Mistral-7B and Qwen3-8B activations was 1.16× and 1.17×—significantly smaller than the cycle-count reduction predicted. Why?

Because on the Mistral-7B layer, AITER ran at a sustained GPU clock of 2,108 MHz, while ZYFA caused the MI355X to throttle its clock down to 1,628 MHz—a 480 MHz (22.8%) frequency drop.

Why did a faster, more accurate kernel run at a 22.8% lower clock frequency on the exact same GPU and the exact same layer activations?

Look at the physics of dynamic CMOS power dissipation inside a matrix multiplier array:

where $\alpha_{\text{toggle}} \in [0, 1]$ is the bit-toggling activity factor—the fraction of logic gates and operand bits that flip between $0$ and $1$ on each clock cycle.

1. Why AITER runs at 2,108 MHz: In real language-model attention (unlike synthetic zero-mean Gaussian tensors), attention distributions are heavy-tailed: a few keys have large probabilities, and thousands of keys in the tail have very small probabilities. Without ZYFA’s per-64-key power-of-2 block scaling, AITER lets those small tail probabilities underflow to exact zero (0x00) or subnormals when cast to FP8 e4m3. When the $\mathbf{P}\mathbf{V}$ MFMA array streams thousands of 0x00 bytes into its multiplier trees, the operand bits do not flip ($\alpha_{\text{toggle}} \to 0$). Dynamic switching power drops, and MI355X’s power management firmware boosts $f_{\text{clk}}$ to 2,108 MHz.
2. Why ZYFA drops to 1,628 MHz (and still wins): ZYFA’s per-64-key power-of-2 block scaling rescales each block upward before rounding so that small tail probabilities survive as non-zero e4m3 bit patterns (which is why ZYFA’s error is 62%–78% lower). Feeding dense, non-zero mantissa and exponent bits into the Matrix Cores on every cycle maximizes bit-toggling activity $\alpha_{\text{toggle}}$ at 86% matrix-pipe utilization, driving the chip into its TDP/current limit and forcing DVFS to pull the clock down to 1,628 MHz. As Zyphra’s authors emphasize, this is why low-precision attention kernels must always be benchmarked on real model activations at sustained clocks rather than synthetic Gaussian inputs in short bursts.

Even after absorbing that 22.8% clock drop, ZYFA’s cycle reduction is so large that it still finishes 1.16×–1.76× faster in wall-clock time and delivers 21%–22% higher energy efficiency (1.94 and 1.91 TFLOP/J on Mistral-7B and Qwen3-8B, vs. AITER’s 1.60 and 1.56 TFLOP/J).

Conclusion: On both Blackwell and MI355X (and even more so under CDNA 5’s $4\times$ low-bit matrix scaling), the only way software reaches 53%–71%+ of peak attention throughput is by treating the non-matmul softmax path—every MUFU.EX2 instruction, every VALU cycle, every TMEM load/store, and every bit-flip in the probability tile—as the primary scarce resource. Now watch what happens when a hardware architect tries to bolt 2:4 Sparse Tensor Cores onto that exact bottleneck.

3. Can 2:4 sparse Tensor Cores speed up attention?#

Once you understand that modern attention is bottlenecked by softmax and the on-chip data path between $\mathbf{Q}\mathbf{K}^\top$ and $\mathbf{P}\mathbf{V}$, examine the hardware pitch for 2:4 Sparse Tensor Core Attention—most recently in NVIDIA’s Vera Rubin (sm_107) architecture disclosures (Inside NVIDIA Rubin GPU Architecture and PTX ISA 9.4) and in AMD’s 4:2 structured-sparse V_SWMMAC instructions on CDNA 5 (gfx1250).

Here is the hardware pitch in its strongest form:

1. In NVIDIA’s Vera Rubin (sm_107) architecture preview and PTX ISA 9.4: The first attention GEMM ($\mathbf{S} = \mathbf{Q}\mathbf{K}^\top$) runs densely on Tensor Cores, writing the dense FP32 score tile $\mathbf{S}$ into Tensor Memory (TMEM). Next, instead of loading all $M \times N$ scores into registers with a standard tcgen05.ld, Rubin introduces a new compound PTX 9.4 instruction:

   tcgen05.ld.red.spcompress.sync.aligned.32x32b.x32.b32.rowop::max  {cdata, mdata}, redval, [tmem_addr];

   While reading the dense FP32 scores out of TMEM, tcgen05.ld.red.spcompress performs 2:4 selection in hardware: from every group of 4 consecutive FP32 logits, it retains the 2 largest values (cdata), generates 2-bit-per-element positional metadata (mdata), and simultaneously computes the fragment’s local row maximum (redval via .red with rowop::max). Because only 50% of the logits survive into cdata, the softmax warps only have to evaluate ex2.approx((x - max) * log2(e)) and normalize on half as many values (cutting SFU work in half), and the second attention GEMM ($\mathbf{O} = \mathbf{P}_{\text{sp}}\mathbf{V}$) can then execute on Sparse Tensor Cores (tcgen05.mma.sp) at up to $2\times$ the dense MMA instruction rate.
2. The Broader 2:4 Hardware Promise: Since Ampere (sm_80) in 2020 (Mishra et al., arXiv:2104.08378), NVIDIA’s keynote slides and datasheets have quoted “2× the math throughput of dense matrix units” whenever an operand satisfies the 2:4 pattern (at least 2 zeros in every contiguous 4-element group along the GEMM reduction dimension $K$, encoded with 2-bit positional indices).

Even on static linear weights—where pruning and metadata compression happen offline before deployment—practitioners have learned to treat the “2× throughput” datasheet figure with skepticism. As documented across NVIDIA’s own benchmarks, SemiAnalysis’s Hopper Tensor Core analysis, and Spheron’s production benchmarks (Myth: NVIDIA’s 2:4 Sparsity Doubles Your Inference Throughput), real-world 2:4 sparse GEMMs on static weights typically land between $1.1\times$ and $1.4\times$ over dense equivalents (e.g., $\sim 1.2\times$ on A100 TensorRT 8.0, $1.1\times\text{--}1.2\times$ single-stream on A100 Sparse Llama, and $1.0\times$—zero measured speedup—on an NVIDIA L4 serving a verified 2:4-pruned Qwen3-4B-Instruct across concurrency 1–128), due to sparse kernel maturity gaps relative to cuBLAS, TDP power ceilings, and memory-bandwidth bottlenecks at low batch sizes. SemiAnalysis put it bluntly in its Rubin CPX analysis: sparsity “has yet to actually deliver the promised benefits – falling well short of the 2x pickup it promises.”

The 2026–27 datasheets still lead with sparse numbers, and the two vendors now do it differently:

• NVIDIA Rubin introduces a third sparsity scheme. NVIDIA’s Rubin table gives 50 PF NVFP4 “inference”, footnoted as “Transformer Engine compute,” next to 35 PF NVFP4 “training”, footnoted as “dense compute.” NVIDIA attributes the gap to “new hardware-accelerated adaptive compression designed to boost NVFP4 performance while preserving accuracy.” SemiAnalysis reads the same ratio as sparse to dense (it puts Rubin at 50 PF sparse vs 33.3 PF dense, and CPX at 30 vs 20, a 3:2 ratio) and says the scheme “is unlike the 2:4 structured sparsity used in Hopper and Ampere, and it isn’t like Blackwell’s 4:8 pairwise structured sparsity either.” NVIDIA has not published what pattern it accepts, whether the compression happens offline on weights or online on activations, or whether it applies to the attention GEMMs at all. Until it does, the honest reading is that Rubin’s dense NVFP4 rate is 35 PF, and the rest of this article’s tests apply as soon as the mechanism is public. If it compresses static weights offline, §4’s equal-bit and capability arguments apply. If it compresses activations such as $\mathbf{P}$ online, §3.3’s preprocessing cost and §3.4’s Amdahl ceiling apply.
• AMD publishes 2:4 peaks only where they cannot beat its own dense FP8. The MI455X datasheet lists “with structured sparsity” peaks for exactly three formats, Matrix FP16, BF16 and INT8, each 10,066 TFLOPS, double their 5,033 dense. It lists no sparse figure for FP8, MXFP8, MXFP6 or MXFP4. On the same chip, dense FP8 is 20,133 TFLOPS. So the fastest sparse number AMD will put on paper for MI455X is half of its dense FP8 number. AMD’s MI400 page describes the feature the same way: “INT8, and FP16/BF16 sparsity support for AI.” The ISA does define FP8 V_SWMMAC instructions (§3.2), but AMD has not published a rate for them.

When you move from static weights to dynamic attention, where the sparse operand $\mathbf{P}$ (or $\mathbf{Q}, \mathbf{K}, \mathbf{V}$) must be pruned, normalized, quantized, and packed into hardware metadata format on the fly inside every attention tile, the hardware pitch collides with four independent mechanisms.

3.1 Where do the operands live on tcgen05, and why does that force a 7-step repack?#

First, trace where Tensor Core and Matrix Core operands actually live across NVIDIA and AMD architectures:

On older register-resident NVIDIA architectures like Ada Lovelace (sm_89), mma.sp.sync.aligned.m16n8k32 is a warp-synchronous instruction whose inputs $\mathbf{A}$, $\mathbf{B}$, metadata $\mathbf{E}$, and accumulator $\mathbf{D}$ all live in the thread register file (RMEM). That register-resident symmetry is what academic Ada kernels such as DynSpAttn (Yu et al., ACM ICS 2026) relied on: when $\mathbf{S} = \mathbf{Q}\mathbf{K}^\top$ finishes in thread registers, each thread can prune 2 of 4 scores in-register, run exp2 in-register, pack two FP16 values and 2-bit indices into registers, and pass those exact registers directly into the next mma.sp.sync instruction for $\mathbf{P}_{\text{sp}}\mathbf{V}$ without writing $\mathbf{P}_{\text{sp}}$ or its metadata to shared memory.

That register-resident data path no longer exists on datacenter Blackwell (sm_100/sm_103) or Vera Rubin (sm_107). To scale tile sizes up to $128 \times 128$ and $128 \times 256$ without exhausting the 65,536-register file, tcgen05 completely severed the Tensor Core from the CUDA Core register file:

• tcgen05.mma writes the dense score accumulator $\mathbf{S} = \mathbf{Q}\mathbf{K}^\top$ into Tensor Memory (TMEM).
• CUDA Cores (and SFUs) cannot operate on TMEM directly; softmax warps must load $\mathbf{S}$ out of TMEM into general-purpose thread registers (RMEM) via tcgen05.ld.
• Conversely, tcgen05.mma and tcgen05.mma.sp cannot read operand A ($\mathbf{P}$) or sparse metadata from general-purpose thread registers. Operand A must reside in TMEM or Shared Memory (SMEM), and for tcgen05.mma.sp, the 2-bit-per-element metadata operand must also reside at a TMEM address.

Now trace what actually happens on Vera Rubin (sm_107) when you use PTX 9.4’s tcgen05.ld.red.spcompress instruction (zartbot, Nvidia Rubin Architecture Preview, §1.4):

• tcgen05.ld.red.spcompress always operates on FP32 (f32) data in TMEM, performs 2:4 selection, emits 2-bit (b2) metadata, and supports vector lengths from .x4 to .x128.
• Where does it put the output? It places the compressed FP32 logits (cdata) and the 2-bit metadata (mdata) into general-purpose thread registers (RMEM).
• Even inside the register file, ld.spcompress does not cut register pressure by 50%: metadata registers add overhead for shorter vector widths, while long vector widths (.x128) inflate peak register live ranges inside the exact softmax warpgroups (WG0/WG1) that are already pinned against their 200-register/thread ceiling.
• Furthermore, rowop::max only returns the local maximum redval of the loaded fragment—it is not a complete softmax instruction.

Because cdata and mdata land in thread registers while tcgen05.mma.sp requires sparse operand A in SMEM or TMEM and its metadata in TMEM, the software pipeline between $\mathbf{Q}\mathbf{K}^\top$ and $\mathbf{P}\mathbf{V}$ is forced to execute a 7-step register-to-TMEM/SMEM repack round-trip on every single tile:

[TMEM: dense f32 scores S]
   │  (tcgen05.ld.red.spcompress.rowop::max)
   ▼
[RMEM: compressed f32 logits (cdata) + b2 metadata (mdata) + local redval]
   │
   ├─► Step 1: Cross-fragment / cross-warp row-max reduction + online softmax
   │           (ex2.approx((x - max) * log2(e)), row sum, and O rescaling sync)
   ├─► Step 2: Dtype conversion in RMEM (f32 probabilities -> bf16 / fp8 / nvfp4 + block scales)
   ├─► Step 3: MMA A-layout repacking in RMEM (swizzle compressed non-zeros for tcgen05.mma.sp)
   ├─► Step 4: Metadata-layout repacking in RMEM (pack & swizzle b2 mdata into TMEM metadata format)
   ├─► Step 5: Shared-memory or TMEM store (tcgen05.st for P_sp + tcgen05.st for TMEM metadata)
   ├─► Step 6: Asynchronous store completion wait (tcgen05.wait::st.sync.aligned + mbarrier arrive)
   └─► Step 7: Issue tcgen05.mma.sp (reads P_sp from TMEM/SMEM and metadata from TMEM)

Compare that 7-step sparse obstacle course against what dense FlashAttention-4 does on the exact same TMEM hardware (§3.1.2 of the FA4 paper and its CuTe-DSL / TIRx reference implementation):

• In dense FA4, the 128-Row $\times$ 512-Column TMEM allocation (tmem) has zero metadata region. S0 occupies physical columns [0, 128), and P0 is simply an in-place float16 view (P_region) aliasing the upper 64 physical columns [64, 128) of S0.
• The softmax warpgroup loads $\mathbf{S}$, computes $\mathbf{P}$, casts to 16-bit/8-bit, and issues contiguous tcgen05.st.sync.aligned.32x32b.x16.b32 stores straight into [64, 128)—immediately kicking off the first segment (K_SPLIT = 96) of the dense $\mathbf{P}\mathbf{V}$ MMA after just three 32-column chunks are written, with zero metadata repacking, zero metadata TMEM stores, and zero TMEM column fragmentation.

In short, on tcgen05, ld.spcompress saves half the ex2.approx calls in Step 1 only to pile Steps 3, 4, and 5 (register layout shuffling, metadata bit-packing, and extra TMEM metadata writes) right back onto the softmax warpgroups’ CUDA Core ALU and TMEM-store pipelines.

3.2 Does AMD CDNA 5 (gfx1250) avoid that round-trip?#

Look at the fourth row of the operand table above: AMD’s CDNA 5 (gfx1250, Instinct MI455X).

Unlike NVIDIA’s tcgen05, AMD did not move matrix accumulators into a separate Tensor Memory address space on CDNA 5. Instead, CDNA 5 expanded each Wave32’s addressable vector register file to 1,024 unified VGPRs and kept all matrix operands in VGPRs. For structured sparsity, CDNA 5 provides the V_SWMMAC instruction family (ISA §7.12, Tables 43–44 & §7.12.3–7.12.5):

• V_SWMMAC_F32_16X16X64_{F16,BF16}, V_SWMMAC_F16_16X16X64_F16, V_SWMMAC_BF16_16X16X64_BF16, and V_SWMMAC_BF16F32_16X16X64_BF16 for 16-bit data (where a packed $16 \times 32$ Matrix A in 8 VGPRs represents a 4:2 sparse $16 \times 64$ matrix multiplied by a dense $64 \times 16$ Matrix B in 16 VGPRs);
• V_SWMMAC_{F32,F16}_16X16X128_{FP8,BF8}_{FP8,BF8} and V_SWMMAC_I32_16X16X128_IU8 for 8-bit data (where a packed $16 \times 64$ Matrix A in 8 VGPRs represents a 4:2 sparse $16 \times 128$ matrix multiplied by a dense $128 \times 16$ Matrix B in 16 VGPRs).
• In every V_SWMMAC instruction, the packed A-matrix (SRC0), dense B-matrix (SRC1), 2-bit sparse index data (SRC2, fitting in 1 VGPR where each lane holds sixteen 2-bit codes with idx0 < idx1), and accumulator C/D (VDST) all reside in VGPRs.

Let’s state the architectural consequence plainly: because CDNA 5 keeps V_WMMA and V_SWMMAC operands and 2-bit index metadata in VGPRs, the tcgen05 register-to-TMEM/SMEM repack round-trip (§3.1) does not apply to AMD’s sparse path on CDNA 5. On gfx1250, the dense $\mathbf{S} = \mathbf{Q}\mathbf{K}^\top$ output of V_WMMA is already in VGPRs, and V_SWMMAC can consume the pruned $\mathbf{P}_{\text{sp}}$ and its index register SRC2 directly from VGPRs without writing them to LDS or TMEM.

Does that rescue 4:2 sparse attention on CDNA 5? No—because three CDNA 5 hardware constraints and three general algorithmic walls immediately take its place:

1. CDNA 5 has no hardware spcompress instruction—so 4:2 selection and index packing run on the VALU: Unlike Vera Rubin’s tcgen05.ld.red.spcompress, the CDNA 5 ISA has no hardware instruction that selects the top 2 of 4 FP32 logits and emits 2-bit indices. On gfx1250, software must compare every 4 consecutive FP32 scores in VGPRs, select the top 2 while enforcing the hardware’s ordering rule idx0 < idx1 (ISA §7.12.3, p. 101), pack the surviving values into SRC0, and bit-pack sixteen 2-bit index nibbles per lane into the SRC2 VGPR—all using VALU instructions on the exact vector pipeline that softmax is already competing for.
2. In unscaled FP8, CDNA 5’s dense V_WMMA already matches V_SWMMAC’s $16 \times 16 \times 128$ tile shape: Look at ISA Table 43 (§7.12, p. 95) and the RAW data-hazard table (§7.12.1, p. 98). For 16-bit F16/BF16, V_SWMMAC doubles the $K$ dimension per instruction from $K = 32$ (V_WMMA_F32_16X16X32_BF16, requiring 5 V_NOP / co-exec hazard slots) to $K = 64$ (V_SWMMAC_F32_16X16X64_BF16, 3 V_NOP slots). For 8-bit FP8/BF8, however, CDNA 5 already provides a dense $K = 128$ instruction (V_WMMA_F32_16X16X128_FP8_FP8, 4 V_NOP slots) alongside the $K = 64$ instruction (V_WMMA_F32_16X16X64_FP8_FP8, 2 V_NOP slots), while sparse V_SWMMAC_F32_16X16X128_FP8_FP8 (3 V_NOP slots) is capped at the exact same $16 \times 16 \times 128$ expanded shape (as Chips and Cheese also noted from the gfx1250 LLVM target definitions, keeping $K$ the same as the wide dense instruction). The preliminary CDNA 5 ISA guide publishes no rate for V_SWMMAC. The MI455X datasheet gives one only for the 16-bit and INT8 forms: 10,066 TFLOPS, against 20,133 TFLOPS for dense FP8 V_WMMA.
3. CDNA 5’s block-scaled F8F6F4 and F4 WMMA instructions do not support structured sparsity at all: Look at ISA §7.12.6 (p. 103) for V_WMMA_SCALE_F32_16X16X128_F8F6F4, V_WMMA_SCALE16_F32_16X16X128_F8F6F4, V_WMMA_SCALE_F32_32X16X128_F4, and V_WMMA_SCALE16_F32_32X16X128_F4:

   “These matrix operations do not supported structured sparsity.” Those block-scaled instructions are the exact instructions that deliver MI455X’s headline 20.13 PFLOPS MXFP8 and 40.26 PFLOPS MXFP4 throughput and implement ZYFA-style hardware block scaling (eliminating 62%–78% of FP8 attention quantization error). On CDNA 5, choosing 4:2 V_SWMMAC means forfeiting both MXFP4 ($40.26\text{ PF}$) and hardware block scaling on FP8.

And beyond those ISA specifics, CDNA 5 faces the exact same online pruning overhead (§3.3), single-sided Amdahl ceiling (§3.4), and equal-bit / equal-tier precision and capability losses (§4) as NVIDIA GPUs.

3.3 How much does pruning and packing P cost in practice?#

How much does online pruning, repacking, and metadata compression of the attention probability matrix $\mathbf{P}$ actually cost in wall-clock time on real GPUs?

In “Is 2:4 Sparsity Still Worth It? Semi-Structured Sparsity versus Quantization in LLM Inference” (Qin, October 2026, arXiv:2610.28947, §3.5, §4.2, and Appendix E, Tables 13–14), Xiuyi Qin measured this exact operation across three GPU generations—Ada Lovelace (L40S), Hopper (H100), and Blackwell (RTX PRO 6000, sm_120)—using CUPTI hardware kernel timers.

To give the 2:4 side the fastest possible online compression path, Qin did not rely on cuSPARSELt’s built-in cusparseLtSpMMAPrune + cusparseLtSpMMACompress API (which takes 15,086 µs on Blackwell and 47,354 µs on H100 for a $65{,}536 \times 16{,}384$ $\mathbf{P}$ matrix) or PyTorch’s semi-structured conversion (3,330 µs on Blackwell, 4,796 µs on H100). Instead, Qin used a hand-written fused Triton prune-and-compress kernel (reimplementing HieraSparse’s kernel bit-for-bit, plus fused FP8 quantization) that reads $\mathbf{P}$ once, selects the top 2 of every 4 entries, and writes the compressed values and 2-bit metadata indices once—running at 63% to 102% of the GPU’s raw memory-copy bandwidth on H100 and Blackwell (1,892 to 2,277 µs on Blackwell, 1,530 to 1,728 µs on H100; Table 13).