DeepSeek V4 源码剖析

第13章 DeepGEMM：FP4 / FP8 矩阵乘法内核

作者 杨艺韬 · 9,477 字 · 发布于 · 更新于

好的 GEMM kernel 在模型代码里是隐形的，但它决定了一个想法能不能变成产品。

先说一件容易搞混的事：V4 开源的 inference/kernel.py 里，fp8_gemm / fp4_gemm 是用 TileLang 现写的参考实现（@tilelang.jit 装饰、几十行 DSL），不是在调 DeepGEMM。DeepGEMM 是另一个独立仓库，是生产路径上真正跑的那套 kernel 库。本章讲的是后者。

13.1 引子：为什么 cuBLAS 不够用

NVIDIA 的 cuBLAS 是矩阵乘法的事实标准——FP16 / BF16 / FP8 都有官方实现。为什么 V4 不直接用 cuBLAS，要自己写一个 DeepGEMM？

技术报告第 3.3 节给了最直接的一条：传统 cuBLAS 做不到 batch invariance（同一个 token 的输出不随它和谁拼在一个 batch 里而变，逐位一致），所以 V4 把矩阵乘法"端到端换成了 DeepGEMM"，并且为此在多数场景放弃了 split-k。除此之外，还有四个工程上的理由：

理由一：scale 配方要与模型逐一对齐

cuBLAS 的 FP8 GEMM 早期只有 per-tensor 这类简单 scale；Hopper 上的 per-128 与 128×128 分块 scale 要到 CUDA 12.9（2025 年）才进 cuBLASLt。DeepGEMM 在 2025-02 开源时就是照 DeepSeek 的"激活 1×128、权重 128×128"细粒度 scale 写的，到了 V4 的 SM100 路径又换成打包 UE8M0 scale——scale 的粒度、格式与内存布局都跟着模型的量化配方走。

理由二：FP8 激活 × FP4 权重的组合

V4 的 routed expert 是"FP8 激活 × FP4 权重、权重 scale 沿 K 每 32 个元素一个"这种非对称组合。cuBLASLt 自 CUDA 12.8 起在 Blackwell 上支持块缩放的 FP4 / FP8，但 V4 要的是这套组合再叠加下面的 grouped / 融合形态。DeepGEMM 里对应的是一整个 fp8_fp4_gemm_{nt,nn,tn,tt} 系列加 m_grouped_fp8_fp4_gemm_*。顺带一提，DeepGEMM 并没有和 cuBLAS 划清界限——它自己就 wrap 了一层 cublaslt_gemm_{nt,nn,tn,tt}（csrc/jit_kernels/impls/smxx_cublaslt.hpp），tests/test_fp8_fp4.py 就拿它当计时对照。

理由三：MoE 的稀疏 dispatch

V4 的 MoE 是"每个 token 选 6 个 expert"——每个 expert 看到的 token 数不同（per-batch 不规则）。cuBLAS 的 grouped 接口（CUDA 12.5 起，实验性）只覆盖 FP16 / BF16 输入；DeepGEMM 则专门做了"只按 M 轴分组"的 FP8 / FP8×FP4 grouped GEMM，还有 decode 阶段用的 masked 布局（§13.6）。

理由四：V4 的特化优化空间

为 DeepSeek 模型定制的 kernel 库可以把整段计算吃进一个 kernel。DeepGEMM 里最典型的两个例子：fp8_fp4_mega_moe 把 EP dispatch → linear1 (FP8×FP4) → SwiGLU → linear2 (FP8×FP4) → EP combine 融成一个 mega-kernel，让 NVLink 通信和 tensor core 计算重叠；tf32_hc_prenorm_gemm（csrc/apis/hyperconnection.hpp）则把第 10 章 hc_pre 里那句 F.linear(x, hc_fn) * rsqrt 连同平方和一起算了——它的签名是 (a: bf16, b: fp32, d: fp32, sqr_sum: fp32, num_splits=None)，sqr_sum 是每行的平方和（tests/test_hyperconnection.py 的参考值是 a.float().square().sum(-1)），rsqrt 就从它算。num_splits 对应技术报告第 3.3 节讲的那个难点：mHC 这个矩阵乘输出维度只有 24，batch 很小时只能用 split-k，而朴素的 split-k 不确定；报告的解法是"每个 split 的部分和单独输出，再由后续 kernel 做确定性归约"。源码正是这样——传了 num_splits 时，d 与 sqr_sum 多出一维 [num_splits, ...]，每个 split 各写各的，测试里再 d.sum(0) 合起来。cuBLAS 是通用库，不会为这种上下游形态开专门的口子。

DeepGEMM 仓库（github.com/deepseek-ai/DeepGEMM）2025-02 随 V3 那一代开源，2026-04 的更新（PR #304）才加入 V4 用到的 FP8×FP4 GEMM、Mega MoE 与 FP4 Indexer。

13.2 DeepGEMM 仓库结构

截至本章基准的 891d57b（2026-04-24，与 V4 同日）：

DeepGEMM/
├── csrc/                                  # host 侧：Python binding + JIT 调度
│   ├── python_api.cpp                     # 唯一的 PyTorch binding 入口
│   ├── apis/                              # 对外 API 分组
│   │   ├── gemm.hpp                       # fp8 / fp8_fp4 / bf16 GEMM
│   │   ├── mega.hpp                       # Mega MoE
│   │   ├── attention.hpp                  # MQA logits（lightning indexer 打分）
│   │   ├── hyperconnection.hpp            # tf32_hc_prenorm_gemm
│   │   └── einsum.hpp / layout.hpp / runtime.hpp
│   ├── jit/                               # 运行时编译：compiler / cache / kernel_runtime
│   └── jit_kernels/
│       ├── heuristics/                    # 按 shape 选 block size 的启发式（sm90.hpp / sm100.hpp / config.hpp）
│       └── impls/                         # 每条路径一个 .hpp：sm90_fp8_gemm_1d1d、sm100_fp8_fp4_gemm_1d1d、
│                                          #   sm100_fp8_fp4_mega_moe、smxx_cublaslt、sm90_tf32_hc_prenorm_gemm …
├── deep_gemm/
│   ├── include/deep_gemm/                 # device 侧：真正的 kernel
│   │   ├── impls/*.cuh                    # 约 7.9K 行，一条路径一个文件
│   │   ├── mma/{sm90,sm100}.cuh           # WGMMA / tcgen05 封装
│   │   ├── ptx/{tma,wgmma,tcgen05,ld_st}.cuh
│   │   └── epilogue/ scheduler/ layout/ comm/ common/
│   ├── legacy/                            # A100 的 Triton 老 kernel
│   └── mega/                              # Mega MoE 的 Python 侧
├── third-party/tilelang_ops/              # 少量 TileLang 算子
└── tests/                                 # test_fp8_fp4 / test_mega_moe / test_attention / test_hyperconnection …

体量：device 侧 deep_gemm/include/deep_gemm/impls/*.cuh 约 7.9K 行，整个 C++/CUDA 侧（含 host 调度与 PTX 封装）约 2.1 万行，Python 侧（deep_gemm/ 包加 tests/）约 3.2K 行。

仓库的设计哲学，README 自己写得很清楚：借鉴 CUTLASS / CuTe 的概念，但不重度依赖它们的模板与代数——每条 (架构 × 精度) 路径一个独立文件，核心 kernel 函数数量刻意保持得少。还有一条关键工程决策：所有 kernel 都在运行时 JIT 编译，安装时完全不编译 CUDA。

13.3 fp8_gemm 的接口与算法

DeepGEMM 的命名约定是 D = C + A @ B，后缀两个字母表示 A / B 的转置状态。FP8 那一族对外暴露的是 fp8_gemm_nt / nn / tn / tt——fp8_gemm_nt 算的是 D = C + A @ B.T。SM90 只支持 NT（A 行主序、B 列主序），SM100 四种布局都支持。

有一个细节值得先说破：在这一版里 fp8_gemm_* 根本不是独立实现，而是 fp8_fp4_gemm_* 的别名——csrc/apis/gemm.hpp:659 那几行 m.attr("fp8_gemm_nt") = m.attr("fp8_fp4_gemm_nt"); 写得明明白白（m_grouped_ 那三个同理）。真正的分叉在 fp8_fp4_gemm_nt 内部按架构与 scale dtype 做（同文件 :86 起）：arch_major == 9 且 scale 是 FP32 → 走 sm90_fp8_gemm_1d1d / 1d2d；arch_major == 10 且 scale 是 torch::kInt（打包 UE8M0）→ 走 sm100_fp8_fp4_gemm_1d1d；其余直接 DG_HOST_UNREACHABLE。注意这里判断的是入口转换之后的 scale：分叉前那句 layout::transform_sf_pair_into_required_layout(...)（同文件 :83）会先把 scale 整理成 kernel 要的布局，在 SM100 上还会把 FP32 scale 自动打包成 UE8M0（除非传 disable_ue8m0_cast）。所以"FP8 GEMM"和"FP8×FP4 GEMM"在 API 层是同一个入口，权重到底是 FP8 还是 FP4 由张量 dtype 决定。

两条必须记住的约束（README 明确写了）：

  1. LHS 的 scale 必须是 TMA 对齐且转置过的布局——库提供 transform_sf_into_required_layout / get_mn_major_tma_aligned_tensor 帮你转（这一版的 GEMM 入口也会自动调用上面那句转换）
  2. scale 的数据格式两代不同：SM90 要 FP32 scale；SM100 要打包过的 UE8M0——4 个 UE8M0 塞进一个 torch.int。SM100 上若传 FP32 scale，入口打包时只取 FP32 的 8 位指数（deep_gemm/include/deep_gemm/impls/smxx_layout.cuh:102 的 >> 23），尾数被直接丢掉——scale 本身不是 2 的幂时，结果会静默偏差而不报错（按源码推断，本机无 GPU 未实跑；CPU 上模拟这一步，0.75 会变成 0.5、3 会变成 2）

另外，输入转置与 FP8 casting 不由这个库负责，README 直说了让用户自己实现或融进前一个 kernel。

算法核心：

对每个输出 tile (m_tile, n_tile):
    accumulator = 0  (FP32)
    对每个 k_tile:
        # 1. 从 global → SMEM 加载 x[m_tile, k_tile] 和 w[n_tile, k_tile] (FP8)
        # 2. 加载对应的 scale (SM90: FP32；SM100: 打包 UE8M0)
        # 3. MMA 指令做 FP8 × FP8 → FP32 (SM90: WGMMA；SM100: tcgen05)
        accumulator += scale_x * scale_w * (x_block @ w_block.T)
    # 4. 把 accumulator (FP32) 转 BF16 写回
    output[m_tile, n_tile] = bfloat16(accumulator)

关键工程点：

点 1：scale 在哪一步乘上去，两代不同

SM90 上，每个 K 块做完一次 WGMMA，kernel 从 SMEM 读出 FP32 的 scale_a、scale_b，在 CUDA core 上把 scale_a * scale_b 乘到这一块的部分和上，再加进 FP32 累加器（deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d2d.cuh:328 起）——就是上面伪代码那一行。SM100 上则不用软件乘：UE8M0 scale 被拷进 TMEM，由 tcgen05.mma ... kind::mxf8f6f4.block_scale 指令在 MMA 内部施加（deep_gemm/include/deep_gemm/ptx/tcgen05.cuh:53）。ue8m0 是纯指数，所以硬件做这一步很便宜。

scale 张量本身很小：按字节算，128×128 分块的 FP8 权重，scale 是权重本体的 1/16384；每 32 个元素一个 scale 的 FP4 权重是 1/16；1×128 的激活是 1/128。

点 2：tile 大小由启发式选，不是写死的

csrc/jit_kernels/heuristics/ 下面有 sm90.hpp / sm100.hpp / config.hpp，按 (M, N, K) 与架构挑 block size。库还给了 deep_gemm.set_block_size_multiple_of(...) 让调用方约束块大小（BLOCK_M / BLOCK_N）是某个数的倍数（csrc/jit_kernels/heuristics/sm90.hpp:40、csrc/jit_kernels/heuristics/sm100.hpp:52 起按它筛候选）——可以用来把 GEMM tile 对齐到量化 block（V4 的 FP8 权重是 128×128），让每个 tile 干净地对应一组 scale。

点 3：异步加载 + 计算重叠

H100 / B200 的 cp.async.bulk + TMA 让全局内存到 SMEM 的拷贝异步进行——计算前一个 tile 的同时拷贝下一个 tile 的数据。SMEM 被切成多级（stage），级数由启发式按 SMEM 容量算出；SM90 那边还要求至少 3 级、小 tile 至少 4 级，注释写的理由就是"藏住 TMA 延迟"（csrc/jit_kernels/heuristics/sm90.hpp:105-107）。

13.4 fp8_gemm 的异步流水线图

把 §13.3 的算法用图表达——展示异步加载与计算重叠：

flowchart LR
  subgraph Cycle1["Cycle 1"]
    A1["cp.async.bulk<br/>load tile 0 → SMEM"]
  end
  subgraph Cycle2["Cycle 2"]
    A2["cp.async.bulk<br/>load tile 1"]
    W1["wait tile 0"]
  end
  subgraph Cycle3["Cycle 3"]
    A3["cp.async.bulk<br/>load tile 2"]
    W2["wait tile 1"]
    C1["WGMMA compute<br/>tile 0 → FP32 acc"]
  end
  subgraph Cycle4["Cycle 4"]
    A4["load tile 3"]
    W3["wait tile 2"]
    C2["compute tile 1"]
  end
  subgraph Output["...output stage"]
    Acc["FP32 累加器"]
    Cast["→ BF16 输出"]
  end
  Cycle1 --> Cycle2 --> Cycle3 --> Cycle4 --> Output
  C1 -.scale_x × scale_w.-> Acc
  C2 -.scale_x × scale_w.-> Acc
  Acc --> Cast

图里每个 cycle 同时做 3 件事：加载下一个 tile + 等上一个 tile + 算上上个 tile，画的是 3 级的情形。实际级数由启发式按 SMEM 容量定（§13.3 点 3），目的一样：让 SMEM 持续被使用、TensorCore 持续有活干。

如果某个 stage 落后（如 SMEM 不够大、cp.async 延迟过长），整条流水线会"气泡化"——TFlops 大幅下降。这是 DeepGEMM 调优时最重要的诊断点。

13.5 fp4_gemm 的特殊处理

先说清楚 DeepGEMM 里 FP4 那条路的形态：它叫 fp8_fp4_gemm，不是"纯 FP4 GEMM"——激活是 FP8、权重才是 FP4，正好对应第 12 章 §12.7 里 linear() 那句"FP4 weight 也要先把 x 量化到 FP8"。deep_gemm/include/deep_gemm/impls/ 下带 FP4 的 GEMM 只有 sm100_fp8_fp4_gemm_1d1d.cuh 和 sm100_fp8_fp4_mega_moe.cuh（另有 sm100_fp4_mqa_logits.cuh / sm100_fp4_paged_mqa_logits.cuh 给 FP4 Indexer 打分用）——FP4 这条路是 SM100 独有的：host 侧 csrc/utils/layout.hpp:48 直接断言 FP4 操作数只能出现在 arch_major == 10 上。SM90 那边的 GEMM 只有 sm90_fp8_gemm_1d1d / 1d2d 和 sm90_bf16_gemm。

V4 自带的 TileLang 参考实现（inference/kernel.py 的 fp4_gemm_kernel）走的是"先转换再算"的路子，它的 docstring 把策略写得很直白（下面前两行是意译，第三行是原句）：

# 载入 [block_N, sub_K] (sub_K=32) 的 FP4 子块，
# 经 FP32 中转 cast 成 FP8，再做 FP8×FP8 GEMM
B_shared[i, j] = T.Cast(FP8, T.Cast(FP32, B_fp4_shared[i, j]))

这就是"没有原生 FP4 MMA 时怎么办"的标准答案：FP4 只用来省 global memory 带宽与容量，进 tensor core 之前抬回 FP8。DeepGEMM 在 SM100 上则能用 tcgen05 那一代的指令直接吃 FP8 与 FP4 混合的操作数（deep_gemm/include/deep_gemm/ptx/tcgen05.cuh 的 SM100_MMA_MXF8F6F4_SS，发的是 kind::mxf8f6f4.block_scale）。不过技术报告第 1 节也说明，现有硬件上 FP4×FP8 的峰值算力与 FP8×FP8 相同，只是理论上未来硬件可以再高效 1/3——眼下 FP4 权重的收益主要是显存与访存。

至于"B200 比 H100 快多少"——README 里没有这个数字，DeepGEMM 的公开性能数据只有下面 §13.10 引的那一条。

13.6 grouped GEMM：MoE 的关键 kernel

V4 的 MoE forward 用循环逐 expert 跑：

for i in range(self.experts_start_idx, self.experts_end_idx):
    if counts[i] == 0:
        continue
    expert = self.experts[i]
    idx, top = torch.where(indices == i)
    y[idx] += expert(x[idx], weights[idx, top, None])

这个循环在小 batch 下效率很低——每个 expert 只算 1-2 个 token，GEMM 的 launch overhead 主导计算。

DeepGEMM 提供 grouped GEMM——但它的分组方式与 CUTLASS 的传统 grouped GEMM 不同，README 特意强调了这点：只在 M 轴分组，N 和 K 必须固定。理由正是 MoE 的形状特点——同一层里所有 expert 的 [dim, inter_dim] 是一样的，变的只有"每个 expert 分到多少 token"。

两种布局，对应两个部署阶段：

contiguous 布局（训练前向 / 推理 prefill）：把所有 expert 的 token 按 expert 顺序拼成一个大张量，每段的起点必须对齐到 GEMM 的 M block（用 get_mk_alignment_for_contiguous_layout() 查）：

deep_gemm.m_grouped_fp8_gemm_nt_contiguous(...)       # FP8 × FP8
deep_gemm.m_grouped_fp8_fp4_gemm_nt_contiguous(...)   # FP8 激活 × FP4 权重（对应 V4 routed expert 的精度配方）

masked 布局（推理 decode）：开了 CUDA graph 之后，CPU 根本不知道每个 expert 收到几个 token——这时传一个 mask，kernel 只算有效部分：

deep_gemm.m_grouped_fp8_gemm_nt_masked(...)
deep_gemm.m_grouped_fp8_fp4_gemm_nt_masked(...)

README 直接点名了它的搭档：masked 版的示例用法（README 原文 example usage）就是把 DeepEP 低延迟 kernel 的输出喂进来（第 16 章）。

还有一条给反向用的：k_grouped_fp8_gemm_tn_contiguous，M 和 N 固定、按 K 轴分组，服务 MoE 的 weight 梯度。

再往上一层是 Mega MoE：fp8_fp4_mega_moe 把 EP dispatch、linear1、SwiGLU、linear2、EP combine 全部融进一个 mega-kernel，用对称内存（symmetric memory）让 NVLink 通信与 tensor core 计算重叠。它要求多进程启动，用法是先 get_symm_buffer_for_mega_moe(...) 拿 buffer、再 transform_weights_for_mega_moe(...) 把 FP4 权重（带 UE8M0 scale）转成需要的排布。完整例子在 tests/test_mega_moe.py。

这一族 kernel 把 N 个独立 GEMM 合并成一次 launch——大幅降低 launch overhead，对小 batch MoE 至关重要。

inference/model.py 里那段 Python 循环是参考实现——写出来是让人读懂语义的，不是给生产用的。生产路径要么走 m_grouped_* 系列，要么直接上 Mega MoE。

13.7 grouped GEMM 在 MoE 中的调度图

把 m_grouped_* 系列在 V4 MoE forward 中的位置画清楚：

flowchart TB
  Input["x: [B*S, 7168] BF16"] --> Gate
  Gate --> TopK["topk indices: [B*S, 6]"]
  TopK --> Sort["DeepEP dispatch 后<br/>按 expert 重排成连续布局"]
  Sort --> Grouped["x_grouped: [total_tokens, 7168]<br/>+ grouped_layout: [total_tokens]<br/>（每行所属 expert）"]
  Grouped --> GG["m_grouped_fp8_fp4_gemm_nt_contiguous<br/>(decode 用 _masked)<br/>1 次 launch 跑 N 个 expert"]
  GG --> ExpertOuts["expert outputs"]
  ExpertOuts --> Combine["乘路由权重后<br/>DeepEP combine 求和"]
  Combine --> Y["y: [B*S, 7168]"]
  classDef gemm fill:#312e81,stroke:#a78bfa
  class GG gemm

关键点：m_grouped 系列一次 launch 处理所有 expert 的 GEMM——不需要 N 次 kernel launch。这是 MoE 在小 batch 下不被 launch overhead 拖死的根本工程优化。图里 dispatch → GEMM → combine 这三步，Mega MoE 干脆把它们融成了一个 kernel。

13.8 与 cutlass / cuBLAS / Triton 的对比

把 DeepGEMM 与同类方案对比（cuBLAS 一行按 CUDA 12.5 / 12.8 / 12.9 的官方 release notes）：

方案 FP8 FP4 per-block scale grouped GEMM 主要用户
cuBLAS ✅ ✅（CUDA 12.8 起，Blackwell 块缩放） CUDA 12.9 起支持 Hopper 128×128 等分块 仅 FP16/BF16（12.5 起，实验性） 通用
cutlass ✅ 部分 灵活 ✅ 框架开发者
Triton ✅ ✅ (DSL) 灵活 ✅ 算子工程师
DeepGEMM ✅ ✅（FP8 激活 × FP4 权重，SM100） ✅ UE8M0（SM100 打包版）/ FP32（SM90） ✅（只分 M 轴 / K 轴） DeepSeek V3/V4

DeepGEMM 的定位是"为 DeepSeek V3/V4 量身定制的 GEMM"——不追求通用，但在自己的特化场景下追求极致性能。

如果你想在自己的项目里用 V4 那套 FP8 激活 × FP4 权重 + UE8M0 的配方，DeepGEMM 是最省事的选择——cutlass / Triton 都要你自己把 scale 布局、grouped 语义、epilogue 组装一遍。代价是它只覆盖 SM90 / SM100（A100 只剩 deep_gemm/legacy/ 里的 Triton 老 kernel，而且只有 BF16 的 grouped GEMM）。

13.9 编译与部署

DeepGEMM 的环境要求（README「Requirements」原文口径）：

  • GPU：NVIDIA SM90 或 SM100 架构——A100 (SM80) 不在支持之列
  • Python 3.8+
  • 编译器：支持 C++20
  • CUDA Toolkit：SM90 要 12.3 以上（README 强烈建议 12.9+ 以获得最佳性能）；SM100 要 12.9 以上
  • PyTorch 2.1+（Mega MoE 的对称内存要 2.9+）
  • CUTLASS 4.0+ 与 {fmt}——都可以用 git submodule 拉

安装：

# submodule 必须一起 clone
git clone --recursive git@github.com:deepseek-ai/DeepGEMM.git
cd DeepGEMM

# 开发模式：链好头文件并编译 CPP JIT 模块
./develop.sh

# 或者直接安装
./install.sh

注意这里没有编译 CUDA kernel 这一步——所有 kernel 都是运行时 JIT 出来的，install.sh 只构建那个轻量的 C++ JIT 模块。这也意味着第一次跑某个 shape 会有编译延迟（下面 §13.14 的调优点 4）。

跑测试（仓库里只有 tests/，没有单独的 benchmark/ 目录，性能与正确性都在同一批脚本里）：

python tests/test_fp8_fp4.py
python tests/test_mega_moe.py       # 需要多进程启动
python tests/test_hyperconnection.py

13.10 性能数字：README 里到底写了什么

这一节要克制。DeepGEMM 的 README 给出的唯一一条整体性能数字是 2025-04-18 那条 news：

DeepGEMM now achieves up to 1550 TFLOPS on H800

其余的表述都是定性的——README 说它的性能"匹配或超过专家手调的库"（matches or exceeds expert-tuned libraries across various matrix shapes），但没有给 cuBLAS 的对比数字，也没有给 H100 vs B200 的倍数。

所以本节能负责任地说的只有：

  • H800 上 FP8 GEMM 的峰值达到过 1550 TFLOPS（2025-04 的优化批次，见 PR #74 / #78 / #81 / #86）
  • 2026-04-16 那批更新（PR #304）带来了 Mega MoE、FP8×FP4 GEMM、FP4 Indexer、PDL 与更快的 JIT；Mega MoE 的 benchmark 在 PR #316
  • 顺带澄清一条容易记串的：「2025-04-22 腾讯网络平台部贡献、最多 30% 提升」那条 news（PR #130）是 DeepEP README 里的，不在 DeepGEMM 的 news 列表里——DeepGEMM 的 news 只有 2026.04.16 / 2025.09.28 / 2025.07.20 / 2025.05.14 / 2025.05.07 / 2025.04.18 这六条

README 之外，唯一的一手数字在技术报告第 3.1 节：Mega MoE 背后那套细粒度 EP 方案，相对强的非融合基线，一般推理负载加速 1.50～1.73 倍，RL rollout、高速 agent 服务这类延迟敏感场景最高 1.96 倍（报告称在 NVIDIA GPU 与华为昇腾 NPU 上都验证过）。除此之外，任何"DeepGEMM 比 cuBLAS 快 X%"的说法都没有一手出处——要数字就自己在目标硬件上跑 tests/，那才是对你的 shape 有意义的数。

13.11 动手实验：跑 DeepGEMM benchmark

git clone --recursive git@github.com:deepseek-ai/DeepGEMM.git
cd DeepGEMM
./develop.sh

# 跑 V4 典型 GEMM 形状
python -c "
import torch, time
import deep_gemm

M, N, K = 4096, 7168, 7168

# NT 布局：A 行主序 [M,K]，B 列主序 [N,K]，算 D = A @ B.T
a = torch.randn(M, K, dtype=torch.bfloat16, device='cuda').to(torch.float8_e4m3fn)
b = torch.randn(N, K, dtype=torch.bfloat16, device='cuda').to(torch.float8_e4m3fn)
d = torch.empty(M, N, dtype=torch.bfloat16, device='cuda')

# scale 的布局/dtype 两代不同：SM90 要 FP32，SM100 要打包过的 UE8M0。
# 别自己拼，用库里的 transform_sf_into_required_layout / get_mn_major_tma_aligned_tensor。
# 完整可跑的构造见 tests/generators.py 与 tests/test_fp8_fp4.py。

# warmup 一次，让 JIT 把这个 shape 的 kernel 编出来并进缓存
# deep_gemm.fp8_gemm_nt((a, a_sf), (b, b_sf), d)
# torch.cuda.synchronize()
"

照抄这段之前先读 tests/test_fp8_fp4.py 和 tests/generators.py——scale 的构造是这套 API 最容易写错的地方，README 专门强调了 LHS scale 必须是 TMA 对齐且转置的布局，而 SM90 / SM100 要的 scale dtype 还不一样。库里的 deep_gemm.testing.bench / bench_kineto 提供了现成的计时封装，比自己写 time.perf_counter() 靠谱（bench 默认先跑 5 次 warmup，JIT 编译的首次开销落在 warmup 里；它还会先用 256 MB 数据冲掉 L2）。

13.12 DeepGEMM 设计哲学的几个潜规则

DeepGEMM 的代码读起来与 cutlass、Triton 这些通用 GEMM 库非常不同。差异不在算法本身，而在工程哲学。把它的几条潜规则归纳出来：

潜规则一：拒绝模板抽象

README 的原话是"借鉴 CUTLASS 与 CuTe 的概念，但避免重度依赖它们的模板与代数"。落到目录上就是 deep_gemm/include/deep_gemm/impls/ 下一条路径一个 .cuh——打开 sm90_fp8_gemm_1d1d.cuh 看到的就是 SM90 那条路的具体优化，不需要在模板里"展开演算"。

（注意"不用模板"是相对 CUTLASS 而言的：DeepGEMM 的 kernel 仍然是 C++ 模板函数，只是层数浅、参数少，而且靠 JIT 在运行时把部分 shape 特化成编译期常量（由各 API 的 compiled_dims 参数决定，GEMM 默认是 "nk"，M 不参与）——这比在编译期展开一棵模板树更直接。）

潜规则二：与上下游 op 主动融合

cuBLAS 把 GEMM 当作"独立黑盒"——输入张量、输出张量、不关心上下游。DeepGEMM 则在值得的地方把 GEMM 和上下游揉在一起。

但融合的边界要看清：DeepGEMM 并不替你做 act_quant。README 明确写着"输入转置或 FP8 casting 这类操作必须由用户单独处理，请自行实现或融进前面的 kernel"——它接收的就是"已量化的 FP8 + scale"。真正的融合发生在别处：fp8_fp4_mega_moe 把 dispatch / linear1 / SwiGLU / linear2 / combine 融进一个 kernel，tf32_hc_prenorm_gemm 把 HC 的投影与平方和融在一起。这类融合省下的是算子边界的 global memory 往返；公开的收益数字只有 Mega MoE 那一组（技术报告第 3.1 节，见 §13.10）。

潜规则三：tile size 由启发式定，不由调用方定

cutlass / Triton 让你手选 tile size。DeepGEMM 把这件事收进 csrc/jit_kernels/heuristics/——sm90.hpp / sm100.hpp / config.hpp 按 shape 和架构自动挑，调用方只能通过 set_block_size_multiple_of / set_num_sms / set_tc_util 这几个旋钮施加约束，不能直接指定。这削弱了灵活性，但保证了"DeepGEMM 跑出的数字就是这套启发式认为的最优"。想看它挑了什么，设 DG_PRINT_CONFIGS=1。

潜规则四：训练与推理共用一套库

这条要特别澄清，因为很容易想当然地以为"公开的是推理版"。README 的 2025.05.14 news 写得很清楚：DeepGEMM 提供 dense 与 MoE 的 weight gradient kernel。API 层面也能看到——k_grouped_fp8_gemm_tn_contiguous 这个"按 K 轴分组"的接口，就是专给 MoE 权重反向用的（前向那边分的是 M 轴）。所以这不是"训推分离"，而是同一套 kernel 同时服务两端。

这四条潜规则让 DeepGEMM 不能直接替代 cuBLAS（它自己也 wrap 了一层 cuBLASLt 当对照），但作为 V4 这种特化场景的 kernel 引擎，它的取舍完全合理。如果你的项目想从 DeepGEMM 借鉴某个优化，要做好"无法直接复用、必须重写到自己代码里"的心理准备。

13.13 DeepGEMM 与 V4 训练栈的串接

DeepGEMM 不只是给推理用。技术报告第 3.3 节说 V4 为了 batch invariance 把 cuBLAS 端到端换成了 DeepGEMM，目标是"训练与推理之间逐位可复现"。具体串接点：

前向 GEMM：按技术报告第 3.4 节，FP4 QAT 只在后训练阶段引入，训练 step 里 expert 权重先量化到 FP4、再无损反量化回 FP8，走的是现有的 FP8 训练框架（即 FP8×FP8 GEMM）；只有推理与 RL rollout 才直接用真 FP4 权重（对应 DeepGEMM 的 FP8×FP4 路径）。两边数值一致靠的是"FP4→FP8 反量化无损"，而不是训练与推理跑同一个 FP4 kernel。

反向 GEMM：backward 的 weight grad 计算与 forward GEMM 的形状不同——这正是 k_grouped_fp8_gemm_tn_contiguous / k_grouped_bf16_gemm_tn_contiguous 存在的理由（M、N 固定，按 K 轴分组）。README 的 2025.05.14 news 把这类 kernel 称作 "weight gradient kernels for dense and MoE backward"。

BF16 那一族：bf16_gemm_* 与 m_grouped_bf16_gemm_* 也在库里——凡是不适合走低精度的计算（比如需要保精度的中间步骤）都可以落到这条路上。

以上都能从 DeepGEMM 的 API、README 与技术报告核实。但"V4 的训练脚本具体怎么调这些 API"无法核实——V4 的训练代码不公开。本节只说库这一侧提供了什么，不替 V4 团队编排训练流程。

13.14 DeepGEMM 性能调优经验

如果你部署 V4 后发现 DeepGEMM 没跑到预期吞吐，常见原因：

调优点 1：GPU 时钟与温度

低精度 GEMM 是标准的功耗墙场景——tensor core 满载时最先撞上的往往是 power limit 而不是算力。技术报告第 3.1 节也点了这一条：极致融合让计算、访存、网络同时高负载，功耗降频成了关键的性能限制因素。如果 nvidia-smi 看到 SW Power Cap / HW Thermal Slowdown 之类的 throttle reason，先解决散热与供电再调软件。这一项没有通用的下降百分比，以你自己 nvidia-smi -q -d PERFORMANCE 的读数为准。

调优点 2：CUDA Stream 配置

DeepGEMM 的 kernel 发在 PyTorch 的当前 stream 上（at::cuda::getCurrentCUDAStream()，csrc/jit/kernel_runtime.hpp:141），和同一 stream 上的其他算子串行执行。与它有数据依赖的算子本来就得排队；只有确实无依赖的工作（比如通信、另一路计算），才值得放到另一条 stream 上与 GEMM 并行。

调优点 3：TMA 对齐与 scale 布局

比 bank conflict 更常见的坑是 scale 张量。README 反复强调 LHS 的 scale 必须是 TMA 对齐且转置的布局，库为此提供了 get_tma_aligned_size / get_mn_major_tma_aligned_tensor / get_mn_major_tma_aligned_packed_ue8m0_tensor / transform_sf_into_required_layout 一整套工具（这一版的 GEMM 入口也会自动转换，形状不对会被 check_sf_layout 断言拦下）。真正危险的是数值：SM100 上 FP32 scale 被打包成 UE8M0 时只留指数（§13.3），scale 不是 2 的幂就会静默偏差——错法是"结果不对"而不是"报错"。

调优点 4：JIT 预热与缓存

DeepGEMM 的所有 kernel 都是运行时 JIT 出来的，第一次遇到参与特化的某组维度（GEMM 默认是 N、K）会现编译——生产部署前必须用真实形状做一遍 warmup，否则第一个用户请求要等编译。相关旋钮：DG_JIT_CACHE_DIR 指定缓存目录（默认 $HOME/.deep_gemm，容器里记得挂持久卷，否则每次重启白编一遍）；DG_JIT_USE_NVRTC=1 换 NVRTC 编译器换取更快的编译速度（README 提示某些场景可能损失性能）；DG_JIT_PRINT_LOAD_TIME=1 打印 kernel 加载耗时；deep_gemm.set_ignore_compile_dims(True) 让所有维度都不参与 JIT 的特化（逐个 API 调整则用 compiled_dims 参数），减少要编译的 kernel 变体数量。

13.15 DeepGEMM 与 vLLM 中现有 GEMM 路径的并存

vLLM 已经有自己的 GEMM 实现——通常用 cuBLAS 或 CUTLASS。V4 集成 DeepGEMM 时不能简单"全部替换"——必须让两套 GEMM 路径并存。

为什么不能全部替换：vLLM 支持多个模型（Llama、Mistral、Qwen 等）、多种硬件与权重格式；DeepGEMM 只支持 SM90 / SM100，接口也是照 DeepSeek 的量化配方开的（FP8 / FP8×FP4 + 特定 scale 布局）。它虽然也有 bf16_gemm_*，但覆盖面远不足以替掉 vLLM 现有的 cuBLAS / CUTLASS 路径。

vLLM 已经这么做了，而且不用等 V4——0.8.5 里就有两处现成的接入点，可以直接对着读：

  • vllm-0.8.5/vllm/model_executor/layers/quantization/fp8.py:41 用 importlib.util.find_spec("deep_gemm") 探测库是否存在，:447 起的 Fp8MoEMethod.__init__ 再叠条件——环境变量 VLLM_USE_DEEP_GEMM 打开、has_deep_gemm 为真、current_platform.is_cuda() 且 has_device_capability(90)——都满足才把 self.allow_deep_gemm 置真，否则打日志降级
  • vllm-0.8.5/vllm/model_executor/layers/fused_moe/deep_gemm_moe.py:22 的 _valid_deep_gemm() 还要再按形状放行一次：N、K 都得能被 dg.get_m_alignment_for_contiguous_layout() 整除、M 不能小于它，N <= 512 直接拒（注释写的理由是 permute/unpermute 算子还不够好），张量必须 contiguous，expert_map 不为空也拒

这套"环境变量 + 算力 + 形状"的三重门是现成的模板：dispatch 发生在运行期（同一个 vLLM 实例可能加载不同模型），检查开销相对一次 grouped GEMM 可以忽略。V4 要接进来，dispatch 机制不必重新设计，但工作量不止在 _valid_deep_gemm 补一条 FP4 分支：0.8.5 调的还是 DeepGEMM 旧版接口——dg.m_grouped_gemm_fp8_fp8_bf16_nt_contiguous、dg.get_col_major_tma_aligned_tensor（vllm-0.8.5/vllm/model_executor/layers/fused_moe/deep_gemm_moe.py:219、:271），这两个名字在 891d57b 里已经不存在（DeepGEMM 2025-07 的 PR #112 重构时改名，现在对应 m_grouped_fp8_gemm_nt_contiguous 一族；get_m_alignment_for_contiguous_layout 倒还留着别名，见 deep_gemm/utils/layout.py:20）。所以 0.8.5 连本章基准的 DeepGEMM 都要先适配新接口。

与 PagedAttention 的协调：DeepGEMM 处理 attention 内部的 Q/K/V/O linear，但 PagedAttention 处理 attention 计算本身。两者完全解耦——一个负责"线性投影"、一个负责"注意力计算"。集成时不需要让它们共享代码。

第 19 章会展开部署的工程接缝；这里先强调 DeepGEMM 与 vLLM 现有 GEMM 的并存——这是部署 V4 后其他模型也要继续跑的工程现实。

13.16 DeepGEMM 的"开发者工效" 哲学

读完 DeepGEMM 仓库后，会注意到它的另一个工程哲学——对开发者的工效极致优化。具体体现：

工效点 1：每个 impl 文件可独立读懂

DeepGEMM 不像 cutlass 那样模板嵌套——deep_gemm/include/deep_gemm/impls/ 下每个 .cuh 可以独立打开、独立读完。多数在几百行，最长的 sm100_fp8_fp4_mega_moe.cuh 约 1380 行。README 自己就把"作为学习 NVIDIA GPU kernel 优化技术的干净资料"写进了定位。这种"可上手性"对开源项目极重要。

工效点 2：测试驱动开发

DeepGEMM 的 tests/ 按 kernel 族分文件，形状表里直接放着 DeepSeek 模型的实际形状（比如 tests/generators.py 的 (24576, 1536)、(7168, 16384)，tests/test_hyperconnection.py 的 n=24, k=28672），FP8 / FP8×FP4 / BF16 各种精度组合与布局都会枚举一遍。

工效点 3：正确性与性能测在同一批脚本里

与 FlashMLA 那种 tests/ + benchmark/ 分家的组织不同，DeepGEMM 没有单独的 benchmark 目录——tests/test_fp8_fp4.py 之类的脚本既跑数值比对（calc_diff 断言），也用 deep_gemm.testing 里的 bench_kineto 计时，还顺手测一遍 cuBLASLt 作对照。好处是"改完一个 kernel，正确性与吞吐一次看完"；代价是跑一遍测试比纯功能测试慢。

工效点 4：接口跟着模型形状走

注意这里不是说 inference/model.py 调了 DeepGEMM——它调的是同目录的 TileLang 参考实现（本章开头已经说过）。真正的协同体现在接口形状上：grouped GEMM 只分 M 轴（因为同层 expert 形状一样）；tf32_hc_prenorm_gemm 的签名里直接带 sqr_sum（因为 hc_pre 要 rsqrt）；csrc/apis/attention.hpp 里的 fp8_gemm_nt_skip_head_mid 按 head_splits=(left, mid, right) 把每个 head 的输出写成"left｜空出 mid｜right"的排布，中间那段不写（下标映射见 deep_gemm/include/deep_gemm/epilogue/transform.cuh:20）。这些 API 都是照着 DeepSeek 自家模型的形状开的口子，不是通用 GEMM 库会有的东西。

工效点 5：安装时不编译 CUDA

DeepGEMM 把编译整个挪到了运行时——install.sh 只构建那个轻量的 C++ JIT 模块，kernel 在第一次用到时才编。这让"改代码 → 测试"的循环里没有"等整个 kernel 库重编"这一步；代价是首次调用某个 shape 要等，所以有 DG_JIT_CACHE_DIR 与 DG_JIT_USE_NVRTC 这些旋钮。README 的 2025.07.20 那条 news 说这次重构的目标就是"低 CPU 开销的 JIT CPP 模块"。

这 5 个工效点是 DeepGEMM 在 V3 时代积累、V4 时代成熟的工程美学。结果上看，V4 这次的配套更新（PR #304，README 记为 2026.04.16；基准 891d57b 与 V4 同日）是和模型一起到位的。

13.17 DeepGEMM 在"未来硬件"上的演进路径

DeepGEMM 当前覆盖 SM90 / SM100，但 NVIDIA 还有更多硬件路线图。把 DeepGEMM 在未来硬件上的演进做个推测。

下一代数据中心 GPU（纯推测）：

预期改进：

  • TensorCore 对 FP4 / FP6 的进一步原生加速
  • 更大 SMEM（256-512 KB / SM）
  • 改进的 TMA 指令

每上一代新架构，DeepGEMM 就要新增一套 smXXX_* 的 impl 与 heuristics——从仓库里 SM90 与 SM100 两套并存的结构看，这个增量是可预期的工作量。

再下一代训练 GPU（推测 2027 年）：

预期改进：

  • 更高的 NVLink 带宽（H100 的 NVLink 4 已经是单卡双向 900 GB/s，下一代还会往上走）
  • 更大显存

DeepGEMM 在新架构上的主要工作是"给 heuristics 补一套新的 block size 表 + 适配新的 MMA 指令"——算法骨架不变。

国产 GPU（如华为昇腾 / 寒武纪）：

DeepGEMM 的 kernel 深度绑定 NVIDIA 的 TMA / WGMMA / tcgen05 与 PTX，换架构等于重写。V4 的 README 里没有关于国产芯片适配的表述；技术报告第 3.1 节倒是提到，细粒度 EP 方案在 NVIDIA GPU 与华为昇腾 NPU 上都做了验证，但开源的只是 CUDA 版 mega-kernel（即 DeepGEMM 里的 Mega MoE）。所以昇腾上的实现如果存在，也是另一份代码库的事，本节不替它做承诺。

消费级 GPU（5090 / 6090）：

消费级 GPU 没有 NVLink + 显存有限——DeepGEMM 在这上面的优化优先级低（RTX 5090 这一代是 SM120，本来就不在 DeepGEMM 声明的 SM90 / SM100 之列）。社区可能会有"消费级版本"的 DeepGEMM fork，但大概率不会由 DeepSeek 团队维护（推测）。

理解这些演进路径让你做"硬件投资规划"。但要提醒一句：以上关于下一代硬件的内容全是推测——DeepGEMM 当前只声明支持 SM90 与 SM100，NVIDIA 未发布产品的规格不在本专栏可核实的范围内。

13.18 DeepGEMM 工程师速记

版本与依赖（README「Requirements」口径）：

  • GPU：SM90 或 SM100（A100 只剩 deep_gemm/legacy/ 里的 Triton 老 kernel）
  • CUDA：SM90 需 12.3+（建议 12.9+），SM100 需 12.9+
  • 编译器支持 C++20；PyTorch 2.1+（Mega MoE 要 2.9+）；CUTLASS 4.0+ 与 {fmt}（submodule）
  • 安装时不编译 CUDA——kernel 全部运行时 JIT，首次跑某个 shape 要等编译

核心 API（deep_gemm/__init__.py 导出）：

  • fp8_gemm_{nt,nn,tn,tt}：FP8 e4m3 GEMM（SM90 仅 NT）
  • fp8_fp4_gemm_{nt,nn,tn,tt}：FP8 激活 × FP4 权重（SM100）
  • bf16_gemm_{nt,nn,tn,tt}：BF16 GEMM
  • m_grouped_{fp8,fp8_fp4,bf16}_gemm_nt_{contiguous,masked}：MoE 前向 / prefill 用 contiguous，decode + CUDA graph 用 masked
  • k_grouped_{fp8,bf16}_gemm_tn_contiguous：MoE 权重反向
  • fp8_fp4_mega_moe + get_symm_buffer_for_mega_moe + transform_weights_for_mega_moe：融合 MoE
  • fp8_fp4_mqa_logits / fp8_fp4_paged_mqa_logits：lightning indexer 的打分 kernel（V3.2 时代的 fp8_mqa_logits / fp8_paged_mqa_logits 仍作为 legacy 导出）
  • tf32_hc_prenorm_gemm：HyperConnection 的 pre-norm 投影
  • cublaslt_gemm_{nt,nn,tn,tt}：cuBLASLt 对照路径
  • 运行时旋钮：set_num_sms / set_tc_util / set_pdl / set_block_size_multiple_of / set_ignore_compile_dims

性能基线：README 只给了一条——H800 上 FP8 GEMM 达到过 1550 TFLOPS。其他数字请自测。

与 vLLM / SGLang 的集成：

  • 按 weight dtype 分派：FP4 / FP8 走 DeepGEMM，BF16 dense 模型继续走 cuBLAS / cutlass
  • V4 适配落地前，别假设任何具体的集成接口

调试工具：

  • python tests/test_fp8_fp4.py / tests/test_mega_moe.py 等——正确性与性能都在这里
  • deep_gemm.testing.bench / bench_kineto 提供带 warmup 的计时封装
  • DG_PRINT_CONFIGS=1 看每个 shape 选中的 config；DG_JIT_DEBUG=1 看 JIT 细节
  • DG_JIT_DUMP_SASS=1 / DG_JIT_DUMP_PTX=1 导出汇编；DG_JIT_PTXAS_CHECK=1 断言 kernel 没用到 local memory
  • nsys profile 看 kernel 时间分布

常见问题：

  • 安装/加载失败：先查 CUDA 与 C++20 编译器版本，再看 submodule 是不是没 clone
  • 首次调用巨慢：JIT 在编译，配好 DG_JIT_CACHE_DIR 并做 warmup
  • 结果不对：九成是 scale 布局——用 transform_sf_into_required_layout 而不是自己拼

13.19 DeepGEMM 的"异步流水线" 优化

DeepGEMM 接近硬件极限的关键是异步流水线——把数据加载与计算重叠，让 GPU 不空闲。把这条流水线展开。

传统同步 GEMM 的问题：

Cycle 1: 从 global memory 加载 tile 0 到 SMEM
Cycle 2: GPU 等待加载完成
Cycle 3: TensorCore 计算 tile 0
Cycle 4: 加载 tile 1
Cycle 5: 等待
Cycle 6: 计算 tile 1

按这张示意图，6 个 cycle 里只有 2 个在算——TensorCore 大部分时间在等内存。

DeepGEMM 的异步流水线：

Cycle 1: 启动 cp.async.bulk 加载 tile 0
Cycle 2: 启动加载 tile 1, 同时 wait tile 0
Cycle 3: 计算 tile 0, 同时启动加载 tile 2, wait tile 1
Cycle 4: 计算 tile 1, 同时启动加载 tile 3, wait tile 2
...

加载与计算重叠——稳态下 TensorCore 几乎不必停下来等数据。实际能跑到多少利用率取决于 shape 与 stage 数，README 没有给通用数字。

实现关键：

  • cp.async.bulk：H100 / B200 的异步拷贝指令，不阻塞 SM
  • TMA descriptor：预定义的拷贝模板，启动开销极低
  • multi-stage SMEM：SMEM 切成多份，正在用的 + 正在加载的 + 备用
  • fence + wait：精确控制同步点，避免 race

与 sparse_attn 的串行依赖：

DeepGEMM 处理 attention 的 Q/K/V/O linear，sparse_attn 处理 attention 计算本身。两者是串行的——linear 必须先完成才能开始 attention，这条依赖链是单层延迟的一个硬约束。

优化方向：

理论上 linear 与 attention 可以做"算子融合"——把 Q 投影直接接进 attention kernel。这在 891d57b 里还没有，是否会成为未来 DeepGEMM 的方向属于推测。

理解流水线让你 debug 性能问题——如果某个 GEMM 跑得慢，可能是流水线被打断（如 SMEM 不够、依赖未满足）。

13.20 延伸阅读

  • DeepGEMM 仓库 README 与 tests/：本章主要参考（基准 commit 891d57b，2026-04-24）
  • cutlass / CuTe 文档：DeepGEMM 借鉴了它们的概念，理解它们有助于读懂 mma/ 与 ptx/ 下的封装
  • NVIDIA Hopper Architecture Whitepaper：WGMMA / TMA 指令细节
  • 本专栏第 12 章：FP4 / FP8 / ue8m0 的格式细节
  • 本专栏第 14 章：QAT 训练时的 act_quant 与 GEMM 配合

13.21 本章小结

  • DeepGEMM 是为 V3/V4 量身定制的 kernel 库——不只是 GEMM，还包含 Mega MoE、MQA 打分与 HyperConnection
  • 关键设计：按架构区分的细粒度 scale（SM90 用 FP32、SM100 用打包 UE8M0）+ 只分 M/K 轴的 grouped GEMM + 全量运行时 JIT
  • SM90 走 WGMMA 做 FP8 GEMM；没有原生低精度 MMA 时，FP4 只用来省带宽，进 tensor core 前抬回 FP8
  • FP8×FP4 那条路（fp8_fp4_gemm_*）只在 SM100 上有；SM90 侧只有 FP8 与 BF16
  • README 给出的唯一整体性能数字是 H800 上 FP8 GEMM 达到过 1550 TFLOPS，没有与 cuBLAS 的对比数字；技术报告另给出 Mega MoE 细粒度 EP 相对非融合基线 1.50～1.73 倍（延迟敏感场景最高 1.96 倍）
  • V4 选 DeepGEMM 而不是 cuBLAS，技术报告给的一手理由是 batch invariance（第 3.3 节）
  • 集成到 vLLM 等推理引擎是 V4 部署的必经之路

第 14 章：QAT 与 act_quant——V4 训练时的"假量化"全链路。