MLA 最佳算子全景调研

Decode / Prefill / Sparse — 各平台最强 MLA Kernel 横评 (2025-2026)

目录

1. MLA 算子生态总览 2. Decode 算子排行榜 (Hopper / H100 / H800) 3. Decode 算子排行榜 (Blackwell / B200 / B300) 4. Prefill 算子排行榜 5. Sparse MLA (DSA) 算子 6. AMD ROCm 平台 MLA 算子 7. 各框架默认算子选择策略 8. 总结与推荐

1. MLA 算子生态总览

当前 MLA 算子生态已经相当丰富，主要分为 Decode（memory-bound, Absorbed 路径）和 Prefill（compute-bound, Non-Absorbed 路径）两大类：

MLA 算子生态图谱 Decode 算子 (Absorbed / MQA) Memory-bound, 直接用 c_t + k_rope FlashMLA DeepSeek | Hopper/Blackwell | OSS 660 TFLOPS ThunderMLA Stanford Hazy | Hopper | OSS +20-35% FlashMLA-ETAP H20 优化 | 长序列 | OSS 2.78x @64K FlashInfer MLA MLSys Best Paper | 通用 | OSS FP8 支持 CutlassMLA NVIDIA | Blackwell SM100 | 闭源 CUTLASS 3.9 TRTLLM MLA NVIDIA TensorRT-LLM | Blackwell 生产级 Prefill 算子 (Non-Absorbed / MHA) Compute-bound, 需要解压 c_t → full K/V FlashAttention-4 (FA4) NVIDIA | Blackwell 1605 TFLOPS FlashAttention-3 (FA3) Tri Dao | Hopper 标准基线 FlashInfer Prefill 通用 | FP8 | MLA 适配 灵活 Sparse MLA (DSA) 算子 DeepSeek V3.2 Sparse Attention FlashMLA Sparse Prefill 640 / Decode 410 TFLOPS H800

2. Decode 算子排行榜 Hopper / H100 / H800

排名 算子名称 来源 峰值性能 vs FlashMLA 核心特点
🥇 ThunderMLA Stanford Hazy Research 183 TFLOPS, 1520 GB/s
(H100 SXM)		 +20~35% 完全融合的 "megakernel"，消除多次 kernel launch 开销；单 kernel 完成 absorption + attention + output projection
🥈 FlashMLA DeepSeek (开源) 660 TFLOPS, 3000 GB/s
(H800 SXM5)		 基线 Paged KV Cache 支持，变长序列优化，BF16/FP8 支持；Hopper 标杆实现
🥉 FlashMLA-ETAP 社区 (H20 优化) 89 TFLOPS
(H20, BS=16, 64K seq)		 2.78x @64K (H20) 转置注意力流水线，减少 padding 开销；专为 H20 (低带宽) 设计，长序列场景大幅领先
4 FlashInfer MLA FlashInfer Team 接近 FlashMLA ~1x FP8 KV Cache 原生支持，CUDA Graph 兼容，Paged + Ragged 灵活布局；MLSys 2025 最佳论文
5 TileLang MLA TileLang (开源 DSL) 接近 FlashMLA ~1x 用户可自定义 tile 策略，与 FlashMLA 性能持平，显著优于 Triton 实现
Hopper 冠军：ThunderMLA — 在 H100 SXM 上，ThunderMLA 以 183 TFLOPS / 1520 GB/s 的成绩超越 FlashMLA (144 TFLOPS / 1199 GB/s) 约 27%，是目前 Hopper 平台已知最快的 MLA decode kernel。其核心创新是完全融合的 megakernel 设计，避免了 FlashMLA 多次 kernel launch 的开销。
注意 H800 vs H100：FlashMLA 在 H800 SXM5 上报告的 660 TFLOPS / 3000 GB/s 与 ThunderMLA 在 H100 SXM 上的 183 TFLOPS 不能直接比较——H800 的数字是在更大 batch size 下 compute-bound 区域的峰值，而 ThunderMLA 的数字是 representative 实际负载下的表现。相同条件下 ThunderMLA 更快 20-35%。

3. Decode 算子排行榜 Blackwell / B200 / B300

排名 算子名称 来源 峰值性能 核心特点
🥇 FlashMLA (Blackwell) DeepSeek ~1450 TFLOPS (B200) Blackwell 原生 tcgen05.mma 支持，FP8/FP4 KV Cache，Paged Attention
🥈 CutlassMLA NVIDIA (CUTLASS 3.9) 接近 FlashMLA SM100 原生，CUTLASS 模板库，weight-absorbed decoding；与 Fused Reduction 配合
🥉 TRTLLM MLA NVIDIA TensorRT-LLM 生产级优化 端到端推理引擎集成，自动 kernel 选择，FP4 MoE + MLA 联合优化
4 FlashInfer MLA FlashInfer 默认后端 Blackwell 上 vLLM/SGLang 的默认 MLA decode 后端，FP8 支持
Blackwell 冠军：FlashMLA (Blackwell) — 在 B200 上峰值可达 ~1450 TFLOPS，是 Hopper 的 2x+ 提升。NVIDIA 的 CutlassMLA 和 TRTLLM MLA 也提供了接近水平的生产级实现。Blackwell 的 tcgen05.mma 指令 (128×256×K) 相比 Hopper 的 WGMMA 有显著优势。

4. Prefill 算子排行榜

Prefill 阶段 MLA 走 Non-Absorbed (MHA) 路径，本质上使用标准的 FlashAttention 变体，只是 head_dim 不同：

排名 算子名称 平台 峰值性能 核心特点
🥇 FlashAttention-4 (FA4) Blackwell 1605 TFLOPS (B200)
硬件利用率 71%		 Blackwell 原生，tcgen05.mma + Tensor Memory，MLA prefill 最优选择
🥈 SageAttention3 (FP4) Blackwell 1038 TOPS (RTX 5090)
5x vs FA on 5090		 FP4 量化注意力，利用 Blackwell FP4 Tensor Core；精度 tradeoff
🥉 FlashAttention-3 (FA3) Hopper 标准基线 WGMMA + TMA 异步流水线，SGLang chunked prefix cache 唯一支持的后端
4 FlashInfer Prefill Ampere+ 接近 FA3 灵活的 Ragged/Paged 布局，FP8 Q/K，cascade attention 支持
5 FlashMLA Sparse Prefill Hopper 640 TFLOPS (H800) Token-level sparse attention for DSA，仅 V3.2+ 模型
Prefill 冠军：FA4 (Blackwell) — 1605 TFLOPS，71% 硬件利用率。FA3 (Hopper) 仍然是 Hopper 平台的标准选择，且是 SGLang chunked prefix cache 的唯一后端。

5. Sparse MLA (DSA) 算子

DeepSeek V3.2 引入了 DeepSeek Sparse Attention (DSA)，在 token 级别做稀疏注意力，需要专门的 sparse kernel：

算子阶段峰值性能 (H800)峰值性能 (B200)说明
FlashMLA Sparse Prefill Prefill 640 TFLOPS 开发中 token-level 稀疏 prefill，FP8 KV Cache 支持
FlashMLA Sparse Decode Decode 410 TFLOPS ~350 TFLOPS token-level 稀疏 decode，与 DSA selection mask 配合
FlashInfer Sparse MLA Decode 支持中 支持中 FP8 KV cache + CUDA Graph，vLLM 默认 FP8 后端
Sparse MLA 现状：目前仅 DeepSeek V3.2+ 模型使用 DSA 稀疏注意力。FlashMLA Sparse 在 H800 上 prefill 达到 640 TFLOPS，decode 达到 410 TFLOPS。Blackwell 上的 sparse kernel 优化仍在进行中（B200 decode 约 350 TFLOPS）。

6. AMD ROCm 平台 MLA 算子 AMD

算子GPU性能核心特点
AITER MLA MI300X / MI325X / MI355X 17x vs non-AITER baseline
1.2-1.6x vs Triton MLA		 AMD 官方 AI Tensor Engine，手工优化汇编级 kernel，vLLM 推荐默认后端
FlashInfer MLA (ROCm) MI300X+ 与 AITER 持平 FlashInfer 0.2 ROCm 移植，FP8 prefill 支持
TileLang MLA MI300X 接近 AITER DSL 生成，自动调优，与 AITER 手工汇编达到 performance parity
Triton MLA MI300X+ 基线 Triton 后端，通用但较慢
AMD 冠军：AITER MLA — AMD 官方的 AI Tensor Engine for ROCm，在 MI300X 上相比无优化基线提速 17x，比 Triton MLA 快 1.2-1.6x。vLLM 在 ROCm 上推荐 ROCM_AITER_MLA 作为默认后端。

7. 各框架默认算子选择策略

vLLM 的自动选择

GPUPrefill 默认Decode 默认FP8 KV Cache
Hopper (H100/H800)FA3 / FlashInferFlashMLAFlashInfer Sparse MLA
Blackwell (B200/B300)FA4CutlassMLA / FlashInferFlashInfer MLA
AMD MI300XAITER FAAITER MLAAITER MLA

SGLang 的选择

GPUPrefill 默认Decode 默认Chunked Prefix Cache
HopperFA3FlashMLA / FlashInfer / CutlassMLA仅 FA3
BlackwellFA3/FA4TRTLLM MLA / FlashMLA仅 FA3
AmpereFlashInferFlashInfer MLA不支持

8. 总结与推荐

按场景推荐

场景推荐算子原因
Hopper Decode (极致性能) ThunderMLA 比 FlashMLA 快 20-35%，完全融合的 megakernel
Hopper Decode (生产稳定) FlashMLA DeepSeek 官方，Paged KV Cache，生态最成熟
Hopper Decode (H20 长序列) FlashMLA-ETAP H20 上 64K 序列比 FlashMLA 快 2.78x
Blackwell Decode FlashMLA Blackwell ~1450 TFLOPS on B200，tcgen05 原生
Blackwell Prefill FA4 1605 TFLOPS, 71% 硬件利用率
Hopper Prefill FA3 标准选择，SGLang chunked prefix cache 唯一后端
Sparse MLA (DSA) FlashMLA Sparse 目前唯一成熟的 DSA sparse kernel
AMD MI300X AITER MLA AMD 官方优化，17x 加速，vLLM 默认
FP8 KV Cache (通用) FlashInfer MLA 最广泛的 FP8 支持，跨平台兼容
2026 年 Q1 最强 MLA 算子总结：

Decode (Hopper)：ThunderMLA (Stanford Hazy) — 183 TFLOPS / 1520 GB/s on H100，比 FlashMLA 快 20-35%
Decode (Blackwell)：FlashMLA Blackwell — ~1450 TFLOPS on B200
Decode (H20 长序列)：FlashMLA-ETAP — 2.78x speedup @64K
Prefill (Blackwell)：FlashAttention-4 — 1605 TFLOPS on B200
Prefill (Hopper)：FlashAttention-3 — 最成熟，支持 chunked prefix cache
Sparse (DSA)：FlashMLA Sparse — Prefill 640 / Decode 410 TFLOPS on H800
AMD：AITER MLA — MI300X 默认最优
MLA 最佳算子全景调研 | 2025-2026 | 基于公开 benchmark 数据整理