研究报告 · AI INFRASTRUCTURE

大模型训推算子与 Infra 研究报告

基于 2024–2026 年主流模型架构的算子需求、训推差异与"倒挂"分析

版本v1.1
日期2026 年 10 月
样本12 个主流模型系列
资料截止2026-10

00执行摘要

本报告梳理了 2024 年底至 2026 年 10 月公开技术细节的 12 个主流大模型系列，从架构出发分析训练与推理两侧的算子需求，归纳出以下结论：

  1. 架构已收敛到"超稀疏 MoE + 高效注意力 + 原生低精度"。 样本中的 MoE 模型激活参数比例大多在 3%–6% 之间，路由专家数从 128 增长到 512。注意力分化为 KV 压缩、稀疏选择、线性混合、局部窗口四条路线，且经常组合使用。
  2. 同一个模型在训练和推理中对应两套不同的算子。 MLA 的 decode 走权重吸收路径，线性注意力的 decode 走递推路径，MoE 的 decode 是小批量、访存受限的 GEMM。这些路径在训练中根本不存在，必须为推理单独实现。
  3. 硬件决定了瓶颈在哪里。 无论哪代显卡，一个权重矩阵每步约需 300 个 token 才能喂满算力；超稀疏 MoE 因此需要上万 token 的全局 batch 和上百卡专家并行（DeepSeek 的 decode 部署用了 144 卡）。GPU 代际之间 kernel 互不兼容（Hopper、Blackwell 数据中心卡、Blackwell 消费级卡各需一套），FP8 方言不同（OCP 与 FNUZ）会让数值差 2 倍。详见第 3 章。
  4. "倒挂"体现在四个方面： 推理侧算子变体数量远多于训练侧（阶段 × 精度 × 缓存 × 并行的组合）；新架构先在训练栈落地，但能否上线取决于推理算子是否就绪；推理 Infra 的成熟度正在反向约束架构选择；推理集群混合多代硬件，算子还要按硬件代际重复实现。
  5. 头部团队用"开源推理算子"来解决 day-0 问题。 DeepSeek（FlashMLA、DeepGEMM）、Moonshot（KDA kernel）、Qwen（依托 flash-linear-attention）都在模型发布时同步提供推理 kernel，vLLM、SGLang 才能当天支持。
  6. 训推数值一致性正在成为一等需求。 RL 后训练中 rollout 与训练的 logprob 偏差会让 on-policy 退化为 off-policy。DeepSeek-V4 报告中提出了 batch-invariant 的确定性 kernel 库，实现训练与推理逐 bit 一致。
  7. 建议：把推理算子成本纳入架构评审；训推共用 kernel 源码（统一 DSL）；用 QAT 让训练与部署使用同一低精度格式；把训推一致性纳入持续集成。详见第 12 章。

01研究范围与方法

样本

  • 12 个模型系列，时间跨度 2024-12 至 2026-06（技术报告发布时间）
  • 覆盖 DeepSeek、Moonshot（Kimi）、阿里（Qwen）、智谱（GLM）、MiniMax、OpenAI（gpt-oss）、Meta（Llama）、Google（Gemma）
  • 入选标准：公开了足以分析算子需求的架构细节（技术报告、模型配置或推理引擎实现）

资料来源与限定

  • 一手资料优先：技术报告、模型卡 / config、vLLM 与 SGLang 官方博客、开源 kernel 仓库
  • 闭源模型（GPT、Gemini、Claude 等系列）未公开架构细节，不做逐项分析
  • 参数以官方口径为准；不同来源有出入处已注明。文中性能数字均引自原文，未经本报告复现

02主流模型架构全景

模型发布总 / 激活参数注意力MoE 配置低精度与其他上下文
DeepSeek-V32024-12671B / 37BMLA256 路由 + 1 共享，激活 8FP8 训练（细粒度块缩放）、MTP128K [1]
DeepSeek-V3.22025-09 / 12671B / 37BMLA + DSA 稀疏注意力同 V3Indexer 走 FP8128K [3][4]
DeepSeek-V4
Pro / Flash		2026-061.6T / 49B
284B / 13B		CSA + HCA 混合（压缩 + 稀疏）DeepSeekMoE，与 V3 基本一致路由专家 FP4 QAT、mHC、Muon、MTP1M [5]
Kimi K22025-071.04T / 32BMLA（64 头）384 路由 + 1 共享，激活 8MuonClip 优化器128K [8]
Kimi Linear2025-1048B / 3BKDA : MLA = 3 : 1（MLA 层不用位置编码）MoE—1M [9]
Qwen3-Next2025-0980B / 3BGated DeltaNet : 门控注意力 = 3 : 1高稀疏 MoEMTP262K [10][11]
Qwen3.5（旗舰）2026-02397B / 17BGated DeltaNet : 门控注意力 = 3 : 1；全注意力层 GQA 32/2512 路由，激活 10 + 1 共享原生多模态262K [12]
GLM-52026-02744B / 40BMLA + DSA256 路由 + 1 共享，激活 8（sigmoid 打分）提供 FP8 权重200K [13][14]
MiniMax-M22025-10230B / 10B全注意力（放弃了上一代的线性混合）MoE—— [16]
gpt-oss-120b2025-08117B / 5.1BGQA 64/8；128 窗口滑窗与全注意力 1 : 1 交替；注意力 sink128 专家，激活 4专家权重原生 MXFP4128K [18][19]
Llama 4
Maverick / Scout		2025-04400B / 17B
109B / 17B		iRoPE：3 层 8K 分块注意力 + 1 层 NoPE 全局128 / 16 专家—Scout 10M [20]
Gemma 3 27B2025-0327B（稠密）1024 滑窗 : 全局 = 5 : 1；GQA稠密—128K [21]

表 2-1 主流模型架构要点。GLM-5 的层数在技术报告（80）与 config（78）之间不一致，激活参数按官方 40B 计。DeepSeek-V4 总参数部分第三方来源写作 1.65T，此处按技术报告。

图 2-1　MoE 模型的激活参数比例

激活参数 ÷ 总参数，数值越低表示越稀疏。数据来自表 2-1。

除 Llama 4 Scout（16 专家）外，所有 MoE 模型的激活比例都落在 3%–6.5% 之间。万亿级模型（Kimi K2、DeepSeek-V4-Pro）最稀疏，约 3%。

2.1 六个共性观察

① MoE 全面普及且更稀疏

除 Gemma 3 外全部是 MoE。专家数从 128（gpt-oss、Llama 4 Maverick）到 256（DeepSeek-V3、GLM-5）、384（Kimi K2）、512（Qwen3.5），每 token 激活 4–10 个专家。

② 注意力不再有统一答案

KV 压缩（MLA）、稀疏选择（DSA、CSA）、线性混合（Gated DeltaNet、KDA）、局部窗口（滑窗、分块）四条路线并存，并经常混用：如 MLA + DSA、KDA + MLA、滑窗 + 全局。

③ 低精度从推理前移到训练

DeepSeek-V3 用 FP8 训练；DeepSeek-V4 对路由专家做 FP4 量化感知训练；gpt-oss 直接以 MXFP4 发布专家权重。训练格式与部署格式开始统一。

④ 长上下文成为标配

DeepSeek-V4 和 Kimi Linear 支持 1M，Llama 4 Scout 宣称 10M。注意力和 KV cache 取代 FFN，成为推理成本的主要来源，这也是②的根本驱动力。

⑤ 训练侧出现新的"非标"组件

Muon / MuonClip 优化器（Kimi K2、DeepSeek-V4）、mHC 超连接（DeepSeek-V4）、QK-Clip 等，都需要训练框架提供新的算子或分布式实现。

⑥ 闭源与开源的信息差

闭源旗舰不公开架构，但其服务面对的是同类问题。开源模型的技术报告是目前了解算子需求最可靠的窗口。

03硬件基础：显卡如何决定算子

算子的形态一半由模型架构决定，另一半由硬件决定。本章先列出主流 AI 加速卡的关键参数，再用屋顶线模型（roofline）说明：为什么同一个模型在训练、prefill、decode 中的瓶颈完全不同，以及硬件代际、互联、数值格式如何直接改变算子实现。

3.1 主流加速卡关键参数

型号架构BF16FP8FP4显存带宽互联（单卡）FP8 脊点
A100 SXMAmpere（sm_80）312——80 GB2.0 TB/sNVLink3 600 GB/s—
H100 SXMHopper（sm_90）9891,979—80 GB3.35 TB/sNVLink4 900 GB/s591
H800 SXMHopper（sm_90）9891,979—80 GB3.35 TB/sNVLink 400 GB/s，实测单向约 160 GB/s [2]591
H200Hopper（sm_90）9891,979—141 GB4.8 TB/sNVLink4 900 GB/s412
B200Blackwell（sm_100）2,2504,5009,000180–192 GB7.7–8 TB/sNVLink5 1.8 TB/s≈563
B300 / GB300Blackwell Ultra（sm_103）≈2,250–2,500≈4,500–5,00013,500–15,000288 GB8 TB/sNVLink5 1.8 TB/s≈563–625
RubinRubin（224 SM）—17,50035,000（训练口径）288 GB HBM419.2–22 TB/sNVLink6 3.6 TB/s≈800–910
MI300XCDNA3（gfx942）1,3072,615—192 GB5.3 TB/sInfinity Fabric 8 卡互联493
MI325XCDNA3（gfx942）1,3072,615—256 GB6.0 TB/sInfinity Fabric 8 卡互联436
MI355XCDNA4（gfx950）2,5165,03310,066288 GB8 TB/sInfinity Fabric 8 卡互联629

表 3-1 主流 AI 加速卡参数。算力单位为 TFLOPS，均为稠密口径（不含 2:4 稀疏，厂商宣传值常为稀疏口径的 2 倍）。FP8 脊点 = FP8 稠密算力 ÷ HBM 带宽，单位 FLOP/Byte。B300 与 Rubin 的部分参数在不同来源间存在出入（Rubin 带宽 19.2 vs 22 TB/s、NVLink 3 vs 3.6 TB/s），此处给出区间 [29][30][31][33][41]。

两条规律 ① 算力长期比带宽增长得快。A100→H100，BF16 算力 ×3.2，带宽只 ×1.6；H100→Rubin，FP8 算力约 ×8.8，带宽约 ×6.6。脊点越来越高，访存受限的算子（主要是 decode）相对越来越"吃亏"。
② 精度每降一档，算力翻倍。BF16 → FP8 → FP4 依次 ×2。低精度格式不只是省显存，而是直接决定能用到多少算力，这是 FP8 训练、FP4 QAT 兴起的硬件动因。

3.2 屋顶线：为什么 prefill 吃算力、decode 吃带宽

脊点（达到峰值算力所需的算术强度）= 峰值算力 ÷ 显存带宽 线性层 GEMM：M 个 token × 权重矩阵 K×N，每个权重占 b 字节 计算量 = 2 · M · K · N 访存量 ≈ K · N · b （M 较小时，读权重占主导） 算术强度 ≈ 2M / b FLOP/Byte H100，BF16（b=2）： M ≥ 295 才能喂满算力 H100，FP8 （b=1）： 2M ≥ 591 → M ≥ 296 B200，FP4 （b=0.5）：4M ≥ 9000/8 → M ≥ 281

结论：无论哪一代硬件、哪种精度，一个权重矩阵每次大约要同时处理 300 个 token 才能把算力用满。

3.3 MoE 把 decode 的问题放大了几十倍

MoE 中每个专家只收到 k/E 比例的 token。要让每个专家的 GEMM 也达到约 300 个 token，全局每步需要的 token 数就要乘以 E/k：

模型路由专家数 E每 token 激活 k每个专家分到的比例喂满所有专家所需的全局 batch（≈ 300 × E / k）
DeepSeek-V3 / GLM-525681/32≈ 9,600 token / 步
gpt-oss-120b12841/32≈ 9,600 token / 步
Kimi K238481/48≈ 14,400 token / 步
Qwen3.5-397B512101/51≈ 15,400 token / 步
Llama 4 Maverick12811/128≈ 38,400 token / 步

表 3-2 超稀疏 MoE 的 decode 批量需求（理想均匀路由下的估算，不含共享专家）。

单机 8 卡很难凑出上万的并发 decode token，所以超稀疏 MoE 的高效推理必须依赖大规模专家并行（EP）：把专家分散到上百张卡，所有卡上的请求合在一起喂给专家。DeepSeek 公开的 V3/R1 线上部署正是如此 [27]：

Prefill 部署单元

  • 4 节点 32 卡，路由专家 EP32，MLA 与共享专家 DP32
  • 每卡 9 个路由专家 + 1 个共享专家
  • 两个 micro-batch 交替执行，用一个的计算掩盖另一个的通信

Decode 部署单元

  • 18 节点 144 卡，路由专家 EP144，MLA 与共享专家 DP144
  • 每卡只放 2 个路由专家 + 1 个共享专家，权重占用小，显存留给 KV cache
  • 注意力层拆成两步，配合 5 级流水线掩盖通信
  • 额外部署 32 个冗余专家做负载均衡（256 + 32 = 288 = 144 × 2）

3.4 注意力 decode 的算术强度

Decode 时注意力要读完整的 KV cache。它的算术强度取决于一份缓存被多少个 query 头复用：

GQA：每个 KV 头被 g 个 query 头共享（g = query 头数 ÷ KV 头数） 算术强度 ≈ g FLOP/Byte（BF16） MLA（权重吸收后）：所有 h 个头共享同一份 576 维 latent（512 + 64） 每个缓存 token 的计算量 ≈ h × (576 + 512) × 2，访存 = 576 × 2 字节 算术强度 ≈ 0.94 h → 128 头时 ≈ 242，64 头时 ≈ 121 使用 MTP（每步 2 个 query token）时再 ×2

图 3-1　Decode 注意力的算术强度（FLOP/Byte，BF16 KV）

竖线为 H100 / H800 的 BF16 脊点。低于脊点为访存受限，接近或高于脊点才可能计算受限。

普通 GQA 的 decode 注意力远低于脊点，只能拼带宽；MLA 因为 128 个头共享一份缓存，算术强度接近脊点，配合 MTP 后可以进入计算受限区。

3.5 KV cache 与显存容量

长上下文下，限制并发的往往不是权重，而是 KV cache。并发上不去，3.2 节说的 batch 就凑不够。下图是各模型在长上下文中每多 1 个 token 新增的缓存量：

图 3-2　长上下文下每 token 新增 KV cache（KB，BF16）

只统计随序列长度增长的部分；滑窗、分块、线性注意力层的固定开销见下表。1 KB = 1024 字节。
模型随长度增长的部分每条序列的固定开销
Llama 3.1 70B（参照）80 层 × 2 × 8 KV 头 × 128 × 2 B = 320 KB / token无
GLM-578 层 × 576 × 2 B ≈ 87.8 KB / token另有 DSA indexer key 缓存（FP8，较小）
DeepSeek-V3 / Kimi K261 层 × 576 × 2 B ≈ 68.6 KB / tokenV3.2 另有 indexer 缓存
Llama 4 Maverick12 个全局层 × 4 KB = 48 KB / token [39]36 个分块层最多缓存 8192 token：约 1.2 GB
gpt-oss-120b18 个全注意力层 × 2 × 8 × 64 × 2 B = 36 KB / token18 个滑窗层 × 128 token：约 4.7 MB
Qwen3.5-397B15 个全注意力层 × 2 × 2 KV 头 × 256 × 2 B = 30 KB / token [38]45 个 GDN 层的状态：64 头 × 128 × 128 × 4 B（FP32）≈ 180 MB，相当于约 6,100 token 的 KV

表 3-3 KV cache 计算明细。DeepSeek-V4 未列出：其报告称 1M 上下文下 KV 仅为 V3.2 的 10%（Pro）/ 7%（Flash）[5]。

算一笔账：DeepSeek-V3 的 FP8 权重约 700 GB
  • 8 × H100（640 GB）：权重都放不下，必须多机部署。
  • 8 × H200（1,128 GB）：剩约 430 GB，可存约 610 万 token 的 BF16 KV，相当于约 47 条 128K 序列。
  • 8 × B200（约 1,440 GB）：剩约 740 GB，可存约 1,050 万 token，相当于约 80 条 128K 序列。
以上未扣除激活、通信缓冲与运行时开销，实际可用更少。作为对比，gpt-oss-120b 的专家权重是 MXFP4，单张 80 GB 显卡就能放下 [18]。

3.6 GPU 代际：每一代都要重写核心 kernel

架构矩阵指令与关键特性对算子的影响代表 kernel
Ampere（sm_80）warp 级 mma.sync、cp.async通用基线FlashAttention-2
Hopper（sm_90）warpgroup 异步 wgmma、TMA 张量搬运、线程块集群、FP8需要 warp 专用化、生产者-消费者流水线；FP8 累加精度受限（见 5.5 节）FlashAttention-3、FlashMLA、DeepGEMM
Blackwell 数据中心（sm_100 / sm_103）tcgen05 MMA、Tensor Memory（每 SM 256 KB）、2-CTA MMA、原生块缩放 MX / NVFP4累加器放在 TMEM，softmax 要在 TMEM 和寄存器间来回搬运；Hopper 的 wgmma kernel 必须重写FlashAttention-4（CuTe DSL）
Blackwell 消费级 / 工作站（sm_120 / sm_121）扩展的 mma.sync；没有 TMEM、tcgen05、wgmma；每 block 共享内存 99 KB同叫"Blackwell"，却不能运行 sm_100 的 kernel，写法实际回到 Ampere 风格走 FA2 路径
AMD CDNA3（gfx942）MFMA 矩阵指令，FP8 为 FNUZ 方言独立的 kernel 栈（ROCm、AITER）；FP8 语义与 NVIDIA 不同—
AMD CDNA4（gfx950）MFMA，OCP FP8，原生 MXFP8 / MXFP6 / MXFP4格式与 OCP 标准对齐—

表 3-4 GPU 架构代际与 kernel 兼容性。sm_90 用 wgmma、sm_100 用 tcgen05、sm_120 用 mma.sync，三者之间没有 kernel 兼容性；FlashAttention-4 只在 sm_100 / sm_103 上启用 [37]。

训练集群通常是单一代际的最新硬件；推理集群出于成本考虑，往往混合多代硬件（包括退役的训练卡）。因此推理侧的 kernel 矩阵还要再乘以硬件代际数，这是 7.4 节将讨论的又一层"倒挂"。

3.7 互联：scale-up 与 scale-out

配置Scale-up（节点内 / 超节点）Scale-out（跨节点）对并行策略的含义
H800 集群（DeepSeek-V3 训练）NVLink 实测约 160 GB/s400 Gbps IB，约 50 GB/s，有效约 40 GB/s节点内外带宽差约 4 倍 [2]
HGX H100 / H2008 卡，NVLink 900 GB/s通常每卡 400 Gbps 网卡TP / EP 一般限制在 8 卡内
GB200 / GB300 NVL7272 卡单一 NVLink 域，每卡 1.8 TB/s，全域 130 TB/s机柜间网络可在一个 NVLink 域内做 72 卡宽 EP [30]
Rubin NVL72每卡 3.6 TB/s，全域 260 TB/s机柜间网络单域带宽再翻倍 [29]

表 3-5 典型互联配置。

硬件约束如何塑造了 DeepSeek-V3 的训练方案 [1][2]

硬件约束如何塑造了推理通信算子

DeepEP 提供两套 kernel [40]：

高吞吐模式

节点内走 NVLink、节点间走 RDMA，并做 NVLink 与 RDMA 之间的转发。用于训练和 prefill 这类大批量场景。

低延迟模式

纯 RDMA，基于 hook 的通信计算重叠，不占用 SM。用于 decode 这类小批量、延迟敏感的场景。

同样一个"专家并行 All-to-All"，训练和推理需要两套完全不同的实现。NVL72 这类大 NVLink 域出现后，EP 可以全部走 NVLink，最优的通信 kernel 又会变一次。

3.8 数值格式与硬件支持

硬件FP8 方言块缩放（MX 等）原生支持FP4
NVIDIA HopperOCP E4M3FN / E5M2否，需在 CUDA core 上软件处理 scale 和 FP32 提升否
NVIDIA Blackwell（数据中心）OCPMXFP8 / MXFP6 / MXFP4、NVFP4MXFP4、NVFP4
AMD MI300X / MI325X（gfx942）FNUZ：指数偏置 8、最大值 ±240、无 inf、无 −0否否
AMD MI350X / MI355X（gfx950）OCP：指数偏置 7、最大值 ±448MXFP8 / MXFP6 / MXFP4MXFP4

表 3-6 硬件的低精度格式支持 [32]。FNUZ 与 OCP 的 E4M3 位布局相同，但指数偏置差 1，同一个字节按错误的方言解读，数值正好差 2 倍 [34]。

案例：DeepSeek-V4-Flash 移植到 MI300X [34][35][36] 社区开发者把 V4-Flash 部署到单张 MI300X 时，遇到的问题几乎覆盖了本报告讨论的每一类：
  • FP8 方言：vLLM 的很多 FP8 路径只区分 E4M3 / E5M2，不区分 FNUZ / OCP。V4 的 KV 压缩器、量化写缓存、滑窗 K 缓存、Lightning Indexer 缓存写入都必须改用平台对应的 FP8 类型，否则 scale 和缓存字节对不上。
  • 缺少注意力快路径：没有调优 kernel 的形状会回退到通用 Triton 实现，速度慢数倍。
  • MoE 路由错误：MXFP4 在该硬件上回退到模拟后端时，专家 mask 形状判断错误，token 被送到了错误的专家。
  • 训练侧引入的算子：mHC 的 TileLang 实现依赖 PDL（programmatic dependent launch），在 MI350X 上不受支持；vLLM 官方 recipe 一度把各 AMD 型号都标为不支持 V4。
  • 结果：经过一系列修复和 kernel 调优，单卡 MI300X 单流 decode 约 168 tok/s。
这个案例说明：训练团队为某一硬件写的算子（这里是 mHC），会直接决定模型能否在其他硬件上部署。

04架构趋势与算子需求映射

每一项架构选择都会落到具体算子上。本章把第 2 章的趋势逐项展开为训练侧和推理侧的算子需求。

4.1 超稀疏 MoE

代表：DeepSeek-V3 / V4、Kimi K2、Qwen3.5、GLM-5、gpt-oss、Llama 4
核心算子
  • 路由打分与 Top-k（softmax 或 sigmoid，含无辅助损失的负载均衡偏置）
  • Token 重排（permute / unpermute）
  • Grouped GEMM
  • 专家并行 All-to-All（dispatch / combine）
训练侧要求
  • 大 batch 下的 Grouped GEMM 前向 + 反向
  • All-to-All 与计算重叠（DeepSeek-V3 的 DualPipe）[1]
  • FP8 分发激活以节省通信
推理侧要求
  • Decode 时每个专家只分到少量 token，GEMM 变成访存受限
  • 低延迟 All-to-All（区别于训练的高吞吐模式）
  • 专家权重 FP8 / FP4 存储与反量化
  • DeepSeek-V4 把计算、通信、访存融合进单个 MoE kernel [5]

4.2 KV 压缩注意力（MLA 与少 KV 头）

代表：DeepSeek-V3 / V3.2、Kimi K2、GLM-5、Kimi Linear 的全注意力层；Qwen3.5 全注意力层仅 2 个 KV 头
核心算子
  • 低秩 KV 投影与 latent cache
  • 非对称 head 维度注意力（如 qk 192 / v 128）
  • 解耦 RoPE
训练侧要求
  • 展开为多头形式做前向 + 反向
  • 需要注意力 kernel 支持 qk 与 v 维度不同
推理侧要求
  • Prefill 通常仍用展开形式（计算密集）
  • Decode 用权重吸收，等价为对 latent cache 的 MQA，需要专门 kernel（如 FlashMLA）
  • Paged latent cache、FP8 KV
  • Kimi K2 把注意力头从 128 降到 64，主要就是为了推理成本 [8]

4.3 稀疏注意力（DSA / CSA / HCA）

代表：DeepSeek-V3.2、DeepSeek-V4、GLM-5
核心算子
  • Lightning Indexer：FP8 打分，复杂度仍为 O(L²) 但计算量远小于主注意力
  • 长序列 Top-k 选择
  • 按索引 gather 的稀疏注意力
  • V4：KV 压缩器 + 滑窗 + 稀疏选择的组合 [5]
训练侧要求
  • V3.2 先做 indexer 预热：保留稠密注意力、冻结其余参数，再切换到稀疏训练 [3]
  • V4 随序列长度分阶段（4K→16K→64K→1M）逐步开启稀疏
推理侧要求
  • Indexer 单独的 key + scale 缓存 [7]
  • Prefill 与 decode 需要分别处理 indexer [6]
  • vLLM 对该模型只支持 block size 64，因为 indexer cache 和 FlashMLA 都按此定制 [6]

4.4 线性 / 混合注意力

代表：Qwen3-Next、Qwen3.5（Gated DeltaNet）、Kimi Linear（KDA）；反例：MiniMax-M2 回到全注意力
核心算子
  • Chunkwise 并行扫描（delta rule + 门控）
  • 短卷积（causal conv1d）
  • 与全注意力层交替（3 : 1）
训练侧要求
  • Chunkwise 并行形式的前向 + 反向
  • Kimi Linear 用特化的 DPLR 转移矩阵，减少计算量 [9]
  • MiniMax 指出很多线性架构在训练时也是访存受限 [16]
推理侧要求
  • Decode 走递推形式，每步更新固定大小的状态
  • 状态缓存与 KV cache 混合管理（vLLM 的 hybrid KV cache manager）[11]
  • 前缀缓存、投机解码的状态回滚、低精度状态存储都还不成熟 [16]

4.5 局部注意力与注意力 sink

代表：gpt-oss（128 滑窗 + sink）、Gemma 3（1024 滑窗 5 : 1）、Llama 4（8K 分块 + NoPE 全局）
核心算子
  • 带滑窗 / 分块 mask 的注意力
  • 可学习 sink：softmax 分母上额外的加性项，不对应真实 token [18]
  • 注意力温度缩放（Llama 4）
训练侧要求
  • 注意力 kernel 支持各类 mask 与 sink 的反向
  • mask 模式通常在已有 kernel 上扩展即可
推理侧要求
  • 不同层使用不同长度的 KV cache（环形缓冲 vs 全量）
  • sink 需要改 kernel：vLLM 团队为 gpt-oss 专门让 FlashAttention 3 支持了 sink [18]

4.6 原生低精度（FP8 / FP4）

代表：DeepSeek-V3（FP8 训练）、DeepSeek-V4（FP4 QAT）、gpt-oss（MXFP4 发布）、GLM-5（FP8 权重）
核心算子
  • 块缩放 GEMM（激活 1×128、权重 128×128）
  • 在线量化 / 反量化
  • FP4 × FP8 混合精度 GEMM
训练侧要求
  • 部分硬件 FP8 矩阵单元内部只保留约 14 位精度，DeepSeek 每 4 次 WGMMA 把部分和提升到 FP32 [1]
  • 反向需要沿另一维度重新量化
  • embedding、输出头、门控、归一化、注意力保持高精度 [1]
推理侧要求
  • MXFP4：两个 FP4 打包为 1 字节，在 matmul kernel 内即时解包 [18]
  • 不支持的硬件上回退到 BF16 [18]
  • V4 中 FP4 可无损反量化到 FP8，复用既有 FP8 流水线 [5]
  • KV cache 量化（推理独有）

4.7 多 token 预测（MTP）与新训练组件

代表：DeepSeek-V3 / V4、Qwen3-Next（MTP）；Kimi K2、DeepSeek-V4（Muon 系优化器）；DeepSeek-V4（mHC）
核心算子
  • MTP 模块（训练目标 / 推理草稿头）
  • Muon 的 Newton–Schulz 正交化迭代
  • QK-Clip、mHC 等结构性操作
训练侧要求
  • Muon 需要对完整矩阵做运算，与 ZeRO / FSDP 的参数分片存在张力
  • MuonClip 解决了 attention logit 爆炸问题 [8]
推理侧要求
  • MTP 作为投机解码的草稿来源，需要验证 kernel 和树形注意力
  • 与稀疏 / 线性注意力组合时的状态管理更复杂
  • mHC 增加的残差流小算子需要融合

05关键算子拆解

本章用公式和流程，把五个最关键的算子拆开来讲，说明它们在训练和推理中具体怎么算、为什么需要不同的 kernel。

5.1 MLA：训练和 decode 是两种计算

缓存内容（每层每 token）： c_t = W_DKV · h_t 512 维 latent k_t^R = RoPE(W_KR · h_t) 64 维，所有头共享 训练 / prefill（展开形式，DeepSeek-V3 共 128 头）： k_t,i = [ W_UK,i · c_t ; k_t^R ] 每头 192 维（128 + 64） v_t,i = W_UV,i · c_t 每头 128 维 → 标准多头注意力，qk 与 v 的维度不同 decode（权重吸收形式）： q̃_i = W_UK,iᵀ · q_i^C 把上投影吸收进 query，得到 512 维 score_i(t) = q̃_i · c_t + q_i^R · k_t^R o_i = W_UV,i · Σ_t p_i(t) · c_t 先对 latent 加权求和，再上投影 → 等价于 128 个 query 头共享同一个 576 维 key / 512 维 value 的 MQA

5.2 DSA：indexer + top-k + 稀疏注意力

1. Lightning Indexer 打分（FP8，少量 indexer 头，ReLU 代替 softmax）： I_t,s = Σ_j w_t,j · ReLU( q_t,j^I · k_s^I ) 2. 为每个 query 选出得分最高的 k 个历史 token（V3.2 中 k = 2048）： S_t = TopK_s( I_t,s ) 3. 只在选中的 token 上做注意力（基于 MLA 的 MQA 形式）： o_t = Attention( q_t , { c_s : s ∈ S_t } ) 复杂度：indexer O(L²) 但计算很轻；主注意力从 O(L²) 降到 O(L·k)
步骤训练推理对应 kernel
Indexer 打分预热阶段只训练 indexer，主注意力保持稠密 [3]Prefill 与 decode 分开处理；decode 读 paged indexer 缓存DeepGEMM fp8_mqa_logits（含 paged 版本）[6]
Top-k稀疏阶段开启每个 query 在最长 128K 个分数中选 2048 个专用 top-k kernel
稀疏注意力前向 + 反向按索引 gather latent，block size 固定为 64FlashMLA 稀疏 kernel [4]
缓存无额外的 indexer key + scale 缓存 [7]推理引擎内存池

表 5-1 DSA 各步骤在训练与推理中的实现。DeepSeek-V4 在此基础上增加了 KV 压缩（每 m 个 token 压成一条）和滑窗，并与重度压缩的 HCA 层交替 [5]。

5.3 Gated DeltaNet / KDA：并行训练、递推推理

状态更新（每个头维护一个 d_k × d_v 的状态矩阵 S）： S_t = S_{t−1} · α_t · ( I − β_t k_t k_tᵀ ) + β_t v_t k_tᵀ o_t = S_t · q_t α_t：遗忘门（Gated DeltaNet 为每头一个标量；KDA 改为逐通道的对角矩阵） β_t：写入强度（delta rule） 训练 / prefill：按 64 个 token 分块，块内用 WY 表示转成矩阵乘，块间传递状态 decode：每个 token 做一次 O(d_k · d_v) 的状态更新，读写整个状态矩阵

5.4 MoE 的一次前向

① 路由打分 s = sigmoid(W_r · h)（DeepSeek 系：选择时加负载均衡偏置，加权时不加） ② Top-k 每 token 选 k 个专家；训练时限制最多跨 4 个节点 ③ Permute 按目标专家重排 token ④ Dispatch All-to-All 把 token 发到专家所在的卡（激活以 FP8 发送） ⑤ Grouped GEMM 每个专家：gate / up 投影 → SwiGLU → down 投影 ⑥ Combine All-to-All 把结果发回（BF16） ⑦ Unpermute 恢复顺序，按路由权重加权求和，再加上共享专家的输出
环节训练推理 Prefill推理 Decode
④⑥ 通信与另一 micro-batch 的计算重叠（DualPipe）；反向时 dispatch / combine 互换DeepEP 高吞吐模式 + 双 micro-batch 重叠DeepEP 低延迟模式 + 5 级流水线 [27]
⑤ GEMM每专家数千 token，计算受限，需反向每专家 token 多，与训练相近每专家 token 少，访存受限；专家权重常用 FP8 / FP4
负载均衡无辅助损失的偏置调整冗余专家 + prefill 专用均衡器冗余专家 + decode 专用均衡器

表 5-2 MoE 各环节在三个场景中的差异。

看似微小的结构改动也会产生新的 kernel 需求。例如 gpt-oss 的 SwiGLU 带有截断（limit = 7.0）和残差项 [18]，标准的融合 SwiGLU kernel 不能直接用。

5.5 FP8 细粒度 GEMM（DeepSeek-V3 方案）

量化粒度： 激活：每 1 × 128 个元素一个 scale（每个 token 沿 K 维每 128 个元素一组） 权重：每 128 × 128 块一个 scale scale 在线计算（当前块的最大绝对值），不用历史统计的延迟量化 计算流程（Hopper）： Tensor Core 执行 FP8 MMA，沿 K 维每累加 128 个元素（4 条 WGMMA） → 把部分和搬到 CUDA core 的 FP32 寄存器，乘上激活与权重两侧的 scale 后累加 反向： dgrad / wgrad 需要沿另一个维度分组 → 激活要转置后重新量化

06同一模型，两套算子

把第 4 章按"训练 / 推理 prefill / 推理 decode"三个场景重新排列，可以看到：大部分算子族在 decode 阶段都有训练中不存在的专属实现路径。

算子族训练推理 Prefill推理 Decode差异程度
MLA展开多头，前向 + 反向展开多头权重吸收后的 MQA，读 latent cache新路径
稀疏注意力先稠密预热，再稀疏训练Indexer 打分 + Top-k + 稀疏注意力Paged indexer cache + 稀疏 gather新路径
线性注意力Chunkwise 并行Chunkwise 并行逐步递推 + 状态缓存新路径
MoE大 batch Grouped GEMM + 高吞吐 All-to-All大 batch，与训练相近小 batch、访存受限 GEMM + 低延迟 All-to-All形状不同
低精度 GEMMFP8 块缩放（含反向转置量化）W8A8 / FP8 / FP4Weight-only 或 FP4 反量化 GEMM格式可能不同
局部注意力 / sink带 mask 的注意力带 mask 的注意力分层不同长度 KV cache + sink缓存不同
MTP额外损失项—投机解码：草稿 + 验证新路径
KV / 状态缓存不存在写入Paged 读写、量化、前缀复用推理独有
采样不存在（RL rollout 除外）首 token 采样Top-p / Top-k、约束解码推理独有
Muon / mHC优化器与残差结构mHC 前向mHC 前向训练为主

表 6-1 主要算子族在三个场景下的实现形态。"新路径"指该场景需要训练中不存在的独立 kernel。

含义 训练栈跑通了一个新模型，只说明"推理 prefill"大致可行。"推理 decode"这一列（真正决定服务成本的部分）往往需要从零实现。

07训推算子的"倒挂"

本报告所说的"倒挂"，指训练侧与推理侧在算子上的投入、时序和决定权，与它们的实际复杂度和成本占比方向相反。具体有三个层面。

7.1 数量倒挂：推理侧的变体组合爆炸

训练中一个算子通常只需要"前向 + 反向"两个实现，形状也相对固定。推理侧同一个算子要乘上多个维度：

执行阶段Prefill / Decode / Chunked prefill / 投机验证
×
精度BF16 / FP8 / FP4；KV cache 精度另算
×
缓存Paged block 大小、混合缓存、前缀复用
×
并行TP / EP / DP 注意力 / PD 分离
×
Batch 形状从 1 到数千 token，各自最优 tile 不同

图 7-1 推理算子的变体维度（示意）。实际组合会被裁剪，但量级远超训练侧。

7.2 时序倒挂：训练先行，推理定生死

架构设计与训练研究团队在训练栈上实现新算子（前向 + 反向），只需满足训练形状
模型发布社区和客户立即需要能跑、跑得快的推理实现
推理算子补齐Decode 路径、缓存、低精度、投机解码等全部要在推理引擎中实现

头部团队的做法是在发布时同步开源推理 kernel，把第三步提前到发布当天：

模型同步提供的推理算子引擎支持
DeepSeek-V3.2DeepGEMM（indexer logits，含 paged 版本）、FlashMLA（稀疏注意力）、TileLang 参考实现 [4]vLLM、SGLang 发布当天支持 [6][7]
Qwen3-Next依托社区 flash-linear-attention 的 Triton kernelvLLM 发布当天支持，并新增混合缓存管理 [11]
Kimi Linear开源 KDA kernel 与 vLLM 实现 [9]vLLM
gpt-ossFlashInfer 的 MXFP4 MoE kernel；vLLM 改造 FlashAttention 3 支持 sink [18][19]vLLM、Transformers；不支持 MXFP4 的硬件回退 BF16

表 7-1 Day-0 推理支持的典型做法。

7.3 决定权倒挂：推理 Infra 反向约束架构

过去是"架构定了，Infra 去适配"。现在越来越多的案例显示，推理 Infra 的成熟度在反过来决定架构：

MiniMax-M2 回到全注意力

MiniMax 在小规模实验中看到线性混合注意力与全注意力相当，但大规模下多跳推理等能力明显下降。他们同时指出线性 / 稀疏注意力的 Infra 不如全注意力成熟：训练访存受限、对低精度存储敏感、缺少原生前缀缓存、与投机解码的集成不明确 [16][17]。

Kimi K2 减少注意力头

在 DeepSeek-V3 架构基础上，K2 把注意力头从 128 降到 64，以降低推理开销 [8]。

稀疏注意力向既有 kernel 靠拢

DeepSeek-V4 用 FP4 可无损转 FP8 的设计复用现有 FP8 流水线 [5]。另据报道，MiniMax 后续的稀疏注意力方案也以复用 FlashAttention 类 kernel 为设计目标 [17]。

7.4 硬件倒挂：训练用一代卡，推理要跑很多代卡

训练集群追求最新硬件，且通常是同构的；推理集群出于成本考虑，会混合多代甚至多厂商的硬件。由于各代架构之间 kernel 不兼容（表 3-4），同一个新算子在推理侧要为每种硬件各实现一遍：

小结 倒挂的本质是：推理侧的复杂度和成本占比更高，但新算子的定义权和最早的实现都在训练侧。能把两侧拉齐的团队，才能在发布新架构的同时把它变成可用的服务。

08训推数值一致性

当训练和推理使用不同的算子实现时，数值偏差不可避免。在离线评测中这通常只是噪声，但在以下场景中会变成实质问题：

RL 后训练

Rollout 由推理引擎生成，训练框架重算 logprob。即使在同一硬件上，两者也有差异，会让 on-policy 训练悄悄变成 off-policy；MoE 和量化 rollout 下更严重 [23][24]。

在线蒸馏

教师模型由推理引擎服务、学生在训练栈中学习，同样依赖两侧概率分布的一致。

可复现性

Batch 大小随负载变化，kernel 的归约策略随之改变，导致相同输入得到不同输出 [22]。

8.1 偏差来源与对策

来源说明对策
Batch 不变性不同 batch 大小下 kernel 选择不同的并行与归约策略（如 split-K）Batch-invariant kernel [22]
实现路径不同训练用展开 MLA、推理 decode 用吸收 MLA；线性注意力 chunkwise vs 递推训推共用 kernel 源码，或在训练中复用推理路径
精度格式训练 BF16 / FP8，部署 FP4 / INT4；KV cache 量化QAT，让训练与部署使用同一格式（DeepSeek-V4、gpt-oss）
MoE 路由跳变微小偏差让 token 选中不同专家Routing replay：训练复用 rollout 的路由结果 [24]
残余偏差系统层面无法完全消除截断 / 掩码重要性采样（TIS / MIS）等算法修正 [23]

表 8-1 训推数值偏差的主要来源与对策。

业界最新进展 DeepSeek-V4 技术报告描述了一个 batch-invariant 的确定性 kernel 库，实现了训练与推理之间的逐 bit 可复现 [5]。这意味着"训推一致性"已经从算法侧的补偿手段，上升为算子库的设计目标。需要注意，batch-invariant 方案通常有性能代价，且只在同一硬件与软件栈内成立。

09开源算子生态地图

项目主要维护方覆盖算子训练推理
FlashAttentionDao-AILab稠密 / 变长 / 滑窗注意力✓✓
FlashMLADeepSeekMLA decode、DSA 稀疏注意力—✓
DeepGEMMDeepSeekFP8 块缩放 GEMM、Grouped GEMM、indexer logits✓✓
DeepEPDeepSeek专家并行 All-to-All（高吞吐 / 低延迟两种模式）✓✓
flash-linear-attention社区Gated DeltaNet、KDA 等线性注意力（Triton）✓✓
FlashInfer社区推理注意力、采样、MoE（含 MXFP4）—✓
TritonOpenAI 与社区Kernel DSL，多硬件后端✓✓
TileLang社区Kernel DSL；DeepSeek-V3.2 参考实现与 V4 使用 [4][5]✓✓
batch-invariant-opsThinking MachinesBatch 不变的 matmul、归一化、注意力 [22]—✓

表 9-1 主要开源算子项目。"✓"表示该项目被用于相应场景。

两个趋势值得关注：

10算子成熟度评估

综合以上资料，对主要算子族在各场景的成熟度做定性评估。本表为作者基于公开资料的判断，不是测量结果。

算子族训练推理 Prefill推理 Decode跨硬件可移植性主要短板
稠密 / GQA 注意力成熟成熟成熟可用长上下文成本高
滑窗 / 分块 / sink成熟成熟可用可用分层缓存管理、sink 需改 kernel
MLA成熟成熟成熟早期Decode kernel 高度依赖特定硬件
DSA / CSA 稀疏注意力可用可用可用早期Block size 等约束多，实现集中在少数团队
线性 / 混合注意力可用可用可用可用前缀缓存、投机解码、低精度状态
MoE GEMM + EP 通信成熟成熟可用早期低延迟 All-to-All 依赖特定互联
FP8 块缩放 GEMM成熟成熟成熟可用格式语义与累加精度各家不同
FP4（MXFP4 等）早期可用可用早期硬件原生支持面窄，私有格式不互通
MTP 投机解码成熟—可用可用与稀疏 / 线性注意力组合复杂

表 10-1 算子成熟度定性评估。成熟：有多个生产级实现；可用：有可用实现但有明显约束；早期：实现稀少或性能差距大。

11跨平台部署的附加挑战

当训练和推理（或推理的不同阶段）运行在不同的硬件和软件栈上时，上述问题会进一步放大：

这部分的详细分析（包括逐层栈对比、量化格式对比表、PD 分离部署和 KV 传输估算）见配套文档 《训推分离的异构栈》。

12结论与建议

12.1 结论

主流模型的架构创新几乎都集中在注意力、MoE 和低精度这三块，而它们恰好是推理成本的主体。新架构带来的算子需求在推理侧被成倍放大，且大部分 decode 路径在训练中不存在。算子与 Infra 能力正在从"支撑模型"变成"决定模型"。

12.2 建议

面向模型研发

  • 在架构评审中加入"推理算子成本"一项：新算子是否需要 decode 专属路径、是否兼容 paged cache、前缀缓存和投机解码
  • 架构选型时参考 MiniMax-M2 的经验：小规模实验之外，还要评估 Infra 成熟度
  • 优先选择能复用现有高性能 kernel 的设计

面向训练 Infra

  • 用 Triton / TileLang 等 DSL 编写新算子，使同一源码可服务推理
  • 训练阶段采用 QAT，让训练格式等于部署格式
  • 为 Muon 等新优化器准备分布式实现方案

面向推理 Infra

  • 与模型团队同步开发 decode 路径，目标是发布即支持
  • 建立 kernel 变体预算，控制阶段 × 精度 × 缓存 × 并行的组合数
  • 为混合注意力模型提前设计统一的状态 / KV 缓存管理

面向硬件选型

  • 按"每百万 token 成本"而不是单卡价格比较，同时考虑算力、带宽、显存容量和 kernel 成熟度
  • 超稀疏 MoE 推理优先看 scale-up 域大小和跨节点带宽，它决定能否做大规模 EP
  • 长上下文服务优先看 HBM 容量，KV cache 决定并发上限
  • 引入新硬件前，核对 FP8 方言、块缩放格式支持，以及目标模型关键算子（MLA、稀疏 / 线性注意力、mHC 等）在该架构上的实现状态

面向推理部署

  • MLA 类模型的注意力优先用数据并行，避免张量并行切分头数导致算术强度下降
  • 根据表 3-2 估算 decode 所需的全局 batch，据此规划 EP 规模
  • 为混合注意力模型预留状态快照空间，以支持前缀缓存和投机解码

面向一致性与质量

  • 把训推 logprob 差异、MoE 路由一致率纳入 CI 指标
  • RL 场景默认启用 TIS / MIS 等修正，并监控训推 KL
  • 评估 batch-invariant kernel 的性能代价，在 RL rollout 等关键链路优先采用

附录术语表与参考文献

A. 术语表

MoE
Mixture of Experts，混合专家；每个 token 只激活部分专家
MLA
Multi-head Latent Attention；把 KV 压缩为低秩潜向量以减少缓存
DSA
DeepSeek Sparse Attention；用 lightning indexer 为每个 query 选出 top-k 个 token 参与注意力
CSA / HCA
DeepSeek-V4 的压缩稀疏注意力 / 重度压缩注意力
Gated DeltaNet / KDA
基于 delta rule 的门控线性注意力；KDA 是 Kimi 提出的通道级门控变体
GQA / MQA
分组查询注意力 / 多查询注意力，多个 query 头共享 KV 头
MTP
Multi-Token Prediction；训练时额外预测后续 token，推理时可用作投机解码草稿
QAT
量化感知训练，在训练中模拟低精度行为
EP / TP
专家并行 / 张量并行
Prefill / Decode
推理的两个阶段：一次性处理 prompt / 逐 token 生成
Batch-invariant
算子输出不随 batch 大小变化的性质，是推理可复现的前提
TIS / MIS
截断 / 掩码重要性采样，用于修正 RL 中的训推偏差
算术强度
每读写 1 字节数据所做的浮点运算次数（FLOP/Byte）
屋顶线 / 脊点
Roofline 模型；脊点 = 峰值算力 ÷ 带宽，算术强度低于脊点的算子受带宽限制
Scale-up / Scale-out
节点内或超节点内的高速互联（NVLink 等）/ 跨节点网络（InfiniBand、RoCE 等）
WGMMA / TMA
Hopper 的 warpgroup 异步矩阵乘指令 / 张量内存搬运单元
tcgen05 / TMEM
Blackwell 数据中心卡的张量核心指令 / 存放累加器的 Tensor Memory
FNUZ
AMD MI300 系列的 FP8 变体：无 inf、无负零，指数偏置比 OCP 标准大 1
PDL
Programmatic Dependent Launch，NVIDIA 让相邻 kernel 提前启动、重叠执行的机制
DP 注意力
注意力层按请求做数据并行，而不是按头做张量并行

B. 参考文献