A Deep Dive Into The Flash MLA FP8 Decoding Kernel on Hopper

Read route. Cloned the repository and read the document, the companion 2025-04-22 deep-dive, the README, and the kernel source in csrc/sm90/decode/sparse_fp8/. Repository HEAD at read time: 15f13e5030374295491c5ce31b02d7e63a7772c6, 2026-07-27.

Evidence grade. Every performance figure below is DeepSeek measuring DeepSeek's own kernel and is labelled inline as such. No independent reproduction was found, and none is possible from this seat — see Standing. The mechanism claims, unlike the performance claims, were checked directly against the shipped source and are marked VERIFIED IN SOURCE where they were.

Abstract (verbatim)

The source is an engineering blog post and carries no formal abstract. Its opening two paragraphs, verbatim and in full, stand in its place:

"With the release of DeepSeek-V3.2, we have doubled the context length of our models from 64K tokens to 128K tokens. This puts significant pressure on GPU memory (a single request with 128K tokens requires a KVCache of size $576 \times 2 \times 62 \times 128 \times 1024 = 8.72\ \mathrm{GiB}$), which can lead to out-of-memory (OOM) errors or under-utilized GPUs due to small batch sizes. To address this, we introduced FP8 KVCache for DeepSeek-V3.2.

However, writing a high-performance decoding kernel is challenging due to the need for dequantization and its sparse memory access patterns. In this blog, we share the story behind our new FP8 sparse decoding kernel for Hopper GPUs. We will first explain our FP8 KVCache format, then provide a theoretical analysis of clock cycles, and finally detail the techniques used in our new kernel."

Arithmetic check on the one number in that passage: 576 x 2 x 62 x 128 x 1024 = 9,361,686,528 bytes = 8.719 GiB. The stated 8.72 GiB is correct.

Key Contributions

1. A named bottleneck that is neither compute nor memory bandwidth. The kernel is dequantization-bound: the CUDA cores converting FP8 to BF16 cannot feed the tensor cores fast enough. This is stated with a cycle budget rather than asserted (see Methodology), and it is the contribution that matters most for this rung — it is a bottleneck that exists only at the kernel layer and is invisible from either neighbouring rung.

2. "Crossover" — halving dequantization work without halving the work done. Because MLA decode is Multi-Query Attention, every query head in a token attends to the same key head. Two CTAs handling different query-head halves therefore need identical dequantized K/V. Launching them as a Hopper cluster of 2 lets each dequantize half and hand its half to its partner through Distributed Shared Memory, so each CTA does half the conversion work and still ends up with the full block. VERIFIED IN SOURCE.

3. A working use of Hopper Distributed Shared Memory in a production inference kernel. The exchange uses st.async into the peer CTA's shared memory with a cluster transaction barrier for synchronisation, rather than a round trip through L2 or global memory. VERIFIED IN SOURCE — see Methodology for the exact PTX.

4. A measured before/after on unchanged hardware, algorithm and numerics. 250 -> 410 TFLOPS. The isolation is what gives the number its value: nothing about the model, the precision, or the GPU changed.

Methodology

The FP8 KV cache format

MLA decode behaves like MQA: 128 query heads, 1 key head, head_dim_k = 576, head_dim_v = 512. The first 512 elements of each token's KV entry are quantized tile-wise (tile size 1 x 128) to float8_e4m3, producing 512 FP8 values plus 4 float32 scale factors. The remaining 64 elements — the RoPE part — are left in bfloat16, the source stating they "are sensitive to precision loss".

Per token that is 512 + 16 + 128 = 656 bytes, against 576 x 2 = 1,152 bytes for a bf16 cache: a 43.1% reduction per token (derived here from the source's own field layout; the source does not state the percentage). The README's separate description of the wire format gives the same 656-byte breakdown, so the two documents agree.

Inside the kernel the FP8 values are dequantized back to bfloat16 and concatenated with the untouched RoPE values; both GEMMs (QK and score-V) then run in bfloat16 with float32 accumulate. The cache is compressed; the math is not. That distinction is what creates the bottleneck the rest of the document is about.

The cycle budget that names the bottleneck

The source works in cycles per SM, which is the right unit for a kernel-layer argument.

• Tensor core budget. Each SM does 4,096 MMA FLOPs per clock, derived by the source as 989 TFlops / 1830 MHz / 132 SMs on H800. Check: 989e12 / 1.83e9 / 132 = 4,094.2 — rounded to 4,096, consistent.
• MMA cost. With 64 query heads per CTA: 64 x (576 + 512) x 2 / 4096 ≈ 34 cycles per K/V token. Check: exactly 34.0.
• Dequantization cost. H800 cannot cast float8_e4m3 to bfloat16 directly, so each value takes four steps — FP8 to half, half to float32, float32 to bfloat16, then multiply by the float32 scale. Costed against NVIDIA's published native arithmetic throughputs: (1/64 + 1/64 + 1/16 + 1/256) x 512 ≈ 50 cycles per token. Check: exactly 50.0.

50 > 34, so the tensor cores idle waiting on format conversion. Every step of this derivation reproduces exactly. The argument is sound on its own terms; its weakest joint is that the dequantization figure is a lower bound assembled from documented instruction throughputs ("we need at least"), not a profiled measurement, so the real gap is at best this wide.

Crossover, and how it is implemented

CTAs are launched in clusters of 2, each responsible for 64 of the 128 query heads of the same query token. Each CTA loads half the quantized K/V (128-bit __ldg), dequantizes its half, writes the result to its own shared memory, and simultaneously st.asyncs the same data into its partner's shared memory. A cluster transaction barrier synchronises the exchange; afterwards both CTAs hold the full dequantized block.

Checked against csrc/sm90/decode/sparse_fp8/splitkv_mla.cuh at HEAD:

• #include <cutlass/cluster_launch.hpp>; static_assert(CLUSTER_SIZE == 1 || CLUSTER_SIZE == 2).
• const int idx_in_cluster = CLUSTER_SIZE == 1 ? 0 : head_block_idx % 2; — the split.
• static constexpr int NUM_TOKENS_PER_THREAD = CLUSTER_SIZE == 1 ? 2 : 1; — the literal halving of per-thread dequantization work, which is the whole claim in one line.
• get_peer_addr(...) for both the peer shared-memory buffer and the peer barrier (bar_k_remote_ready) — Distributed Shared Memory addressing.
• st_async_128b(...) emits st.async.weak.shared::cluster.mbarrier::complete_tx::bytes.v2.s64 — a 2 x 64-bit (128-bit) asynchronous store into cluster shared memory with mbarrier transaction completion. The blog's "wide load / st.async" description is exact, not a simplification.

A CLUSTER_SIZE == 1 path is retained throughout, consistent with the "without the crossover technique" baseline the source compares against — though note the source does not state that the 250 TFLOPS figure was produced by this exact code path, and that link is inference on my part rather than something the document establishes.

Companion primary source

The compute-bound framing rests on the earlier deep-dive, docs/20250422-new-kernel-deep-dive.md (committed 2025-04-22, LaTeX fix 2025-04-23), which supplies the roofline. Compute-to-memory ratio for MLA decode is ≈ 2·h_q·s_q. On H800 SXM5 the document gives peak bandwidth 3.35 TB/s and peak 990 TFlops, dropping to ≈ 865 TFlops practical under throttling to ~1600 MHz. It concludes the kernel is compute-bound when h_q·s_q ≥ (1/2)(865/3.35) = 128. Check: that expression evaluates to 129.1, not 128 — a rounding in the source, immaterial to the conclusion since DeepSeek runs decode without tensor parallel, fixing h_q at 128.

Results

All figures are the authors' own benchmarks of their own kernels. [Evidence: vendor self-measurement — DeepSeek measuring DeepSeek. No independent reproduction located. Not reproducible from this seat: no Hopper or Blackwell hardware.]

The headline, and what it isolates

Benchmark configuration, stated: batch_size=128, num_heads=128, s_q=2, topk=2048, H800 SXM5. Derived: 1.64x over the no-crossover baseline, 1.84x at the larger topk, and the crossover kernel reaches 64.1% of the dense bf16 kernel's peak.

The 1.64x is the number worth keeping. Same GPU, same attention algorithm, same bf16 MMA numerics, same FP8 cache format — the only change is which CTA does which half of the conversion and how the halves are shared. Whatever that multiple is, the kernel layer owns all of it.

The source also reports the crossover kernel's execution time matching the dense kernel's at a sequence length of roughly 3,000, with the advantage widening beyond that.

The repository's own published figures

From README.md, with per-line as-of dates taken from git blame rather than the file's mtime:

From the 2025-04-22 companion: the dense decode kernel reaches "up to 80% Tensor Core utilization (of the throttled theoretical peak) and 3 TB/s memory bandwidth", against a previous 3000 GB/s and 580 TFlops, while running ~2% slower than the prior ping-pong version in memory-bound settings.

Two consistency notes on those figures, both mine rather than the source's:

• The 80% and the 660 do not describe the same measurement. 660 / 865 = 76.3% of throttled peak. "Up to 80%" is a best-case utilization; the headline throughput corresponds to a slightly lower one. Not a contradiction, but they should not be quoted as the same claim.
• 3000 GB/s against a 3.35 TB/s peak is 89.6% of stated peak bandwidth, using the peak figure the same authors give in the companion document. On the memory-bound side of the roofline this kernel has already taken most of what the machine has.

Freshness

The repository is live — HEAD 2026-07-27, 15f13e50 — but the README's performance lines have not been touched since 2025-10-01, and the deep-dive since 2025-09-30. Ten and a half months of kernel work (including SM100 additions) sit behind numbers that have not been restated. Treat every figure above as an as-of-October-2025 measurement of a codebase that has since moved, not as the kernel's current standing.

Limitations

Stated by the source.

• The 410 TFLOPS falls short of the 640 TFLOPS dense bf16 peak. The source's own explanation: it is a sparse kernel with a small topk (2048), so prologue and epilogue overhead is proportionally larger. Raising topk to 32768 recovers it to ~460 — which concedes that the headline number is partly an artefact of the chosen configuration.
• The advantage over dense decoding only materialises past a sequence length of ~3,000.
• The cycle analysis is explicitly a floor ("we need at least ... cycles").

Not stated, and load-bearing.

• Single configuration, single GPU, no error bars. One shape on one part. No variance, no run count, no methodology for how throughput was computed from wall time.
• No accuracy evaluation whatsoever. The document argues the FP8 format preserves accuracy structurally — fine-grained 1x128 tiles, RoPE left unquantized — but reports no perplexity, no benchmark score, nothing. A KV-cache quantization change without a quality number is a half-reported result; the reader is asked to take the accuracy leg on the design's word.
• Vendor-measured throughout. DeepSeek benchmarking DeepSeek's kernel, powering DeepSeek's model, in support of DeepSeek's serving economics. The mechanism is checkable and I checked it; the numbers are not independently checkable.
• Hopper-bound by construction. Distributed Shared Memory is an SM90 feature. The README's own B200 sparse decode figure — 350 TFLOPS, below the 410 on the older H800 — is the visible cost: the trick does not carry to the newer architecture as written. That is a portability result the document does not draw attention to, and it is the most interesting thing in the README.
• The FP8-to-BF16 conversion cost is a property of one chip generation. The entire bottleneck exists because H800 lacks a direct float8_e4m3 -> bfloat16 cast. Silicon that adds one deletes the problem and the contribution together.

Why it matters for this rung

The binding constraint, named with evidence

For this rung the constraint is not FLOPs and not bandwidth — it is whether the parts of the SM that are not the tensor core can keep the tensor core fed. This source names that constraint in the hardest available currency: 50 cycles of format conversion against 34 cycles of matrix math, per KV token, on a specific part, with the derivation reproducing exactly. The tensor cores were idle roughly a third of the time and no roofline drawn at the level above would have shown it: by the FLOPs/byte test in the companion document the kernel is comfortably compute-bound, and it still was not compute-limited. That gap between the two is the compilers rung's whole subject matter.

The number this rung owns

1.64x on decode throughput (250 -> 410 TFLOPS, H800 SXM5, as of 2025-09-29), from a kernel-layer change alone.

It qualifies because everything else is held fixed: same silicon, same attention algorithm, same bf16 MMA numerics, same FP8 cache format, same model. No other rung can claim any part of it. Decode throughput is the denominator of output-token cost on the serving side, so a 1.64x here is a direct, if partial, move on the price of a token — and therefore on the price of a finished task, which is where this KB is ultimately pointed.

Two honest qualifications on that number. It is the vendor's own measurement at a single configuration. And it is stale by ten and a half months against a repository that has kept moving.

What the neighbouring rungs would claim, and why it sits here

hardware would claim it, and has the better long-run argument: the bottleneck exists only because H800 has no direct FP8-to-BF16 cast. Silicon with that instruction erases both the problem and the fix. But that is a claim on the next part, not this one. On hardware that shipped, the gain came from software, and the fallback path is still in the tree.

serving would claim it, since FP8 KV cache and sparse attention are KV-budget moves and DSA is a serving-layer algorithm. But serving owns the decision to store fewer bytes; this document is about what happens after that decision, when those bytes have to become numbers again fast enough to matter. The crossover changes no bytes stored and no tokens attended — it reallocates conversion work between two CTAs. That is squarely "the layer between a model graph and the silicon".

models has no claim. Nothing here touches architecture or weights.

There is also a portability finding that belongs to this rung and nowhere else: the same kernel scores 350 TFLOPS on B200 against 410 on H800. A newer, faster part running an older part's kernel more slowly is the cleanest possible statement that the compiler and kernel layer, not the silicon, is what is being measured — and it is the shape of the "portability across accelerators" problem this rung exists to track.

Standing

read it. Connor has not run these kernels and cannot: they require H800 SXM5 or B200 hardware that is not in reach. Nothing here was reproduced.

What was actually done, so the standing is not overstated: the repository was cloned and read at a named commit; every arithmetic step in the source was recomputed and reproduces; the crossover mechanism was verified against the shipped CUDA and down to the PTX instruction; per-figure as-of dates were recovered with git blame rather than assumed; and three things the source does not say — the 43.1% per-token saving, the 76.3%-vs-80% utilization discrepancy, and the B200-slower-than-H800 portability result — were derived here.

That is a genuine close-read of a primary artefact and it is still read it. Per the teaching doctrine this caps anything built on this source at L2 (Mechanism): it can explain how a kernel-layer bottleneck is found and removed, and it cannot make a first-person performance claim. Any piece that reaches for L3 on this material is seat inflation.