DeepSeek DeepGEMM 2026年4月重磅更新：Mega MoE融合算子、FP4精度与极致性能优化

一、背景：为什么DeepGEMM的每次更新都牵动整个AI Infra圈

2026年4月16日，DeepSeek旗下高性能算子库DeepGEMM发布了代号为"Public Release 26/04"的重大更新，一口气带来了Mega MoE、FP8xFP4 GEMM、FP4 Indexer、PDL（Programmatic Dependent Launch）、JIT编译加速等多项关键特性。GitHub上这条PR迅速收获了136个👍和35个🚀，再次印证了DeepGEMM在AI Infra领域无可替代的技术地位。

如果你还不了解DeepGEMM是什么，可以先用一句话概括：它是目前开源世界里对NVIDIA GPU GEMM优化做得最彻底、最Clean的库，没有之一。 H800上实测1550 TFLOPS的峰值性能、轻量级JIT架构、借鉴CUTLASS但大幅简化的代码结构——这些特性让DeepGEMM既是顶级大模型训练的底层支柱，也是GPU内核优化学习者的最佳教科书。

而这一次更新的意义，远不止"又多了一些kernel"这么简单。Mega MoE代表了MoE计算范式的一次范式跃迁——从多kernel流水线到单mega-kernel的工程化跨越；FP4精度支持则意味着AI Infra正式进入Ultra-low precision时代；而PDL和JIT加速，则是DeepSeek在工程细节上对极致性能追求的又一次体现。

本文将深入解析这次更新的每一项核心技术，从原理到代码，从架构到调优，逐一拆解。

二、技术全景：这次更新到底带来了什么

先上一张全景图，然后我们逐项深入：

三、Mega MoE：从"流水线工厂"到"一体化产线"

3.1 传统MoE的计算瓶颈在哪里

要理解Mega MoE的价值，首先得搞清楚传统MoE（Mixture of Experts）的计算流程是怎么跑的。

一个标准的MoE层，在前向传播时的计算路径是这样的：

Token → Router → Dispatch → Expert 1..N (各自独立计算) → Combine → 输出
                      ↓
              Linear1 (FP8)
                      ↓
              SwiGLU激活
                      ↓
              Linear2 (FP8)

看起来线性流畅，但工程实现上，每一步都对应着独立的CUDA kernel调用：

1. Dispatch Kernel：将token分发到各个Expert的显存区域，涉及跨SM的数据搬运
2. Linear1 Kernel：FP8 GEMM，计算量巨大，但需要独立启动
3. SwiGLU Kernel：非线性激活，涉及额外内存读写
4. Linear2 Kernel：又一次FP8 GEMM，与Linear1结构相同但权重不同
5. Combine Kernel：将各Expert输出汇总回主序列

问题来了——在这些独立的kernel之间，存在大量的GPU空闲等待时间。更致命的是，在多Expert Parallelism（EP）场景下，token需要在不同GPU之间通过NVLink传输，而传统的多kernel方案里，通信和计算是串行的：算完才能传，传完才能算，形成典型的"气泡"（bubble）。

3.2 Mega MoE的破局思路

Mega MoE的核心思想一句话概括：把这条流水线焊死成一个mega-kernel，让NVLink通信和Tensor Core计算同时进行。

具体来说，它做了两件事：

第一件事：操作融合（Operation Fusion）

将原本分散的5个步骤——dispatch、Linear1（FP8xFP4）、SwiGLU、Linear2（FP8xFP4）、combine——融合进一个mega-kernel。这意味着：

• 消除了5次kernel启动开销（每次约5-10μs）
• 消除了中间结果的全局内存写入/读取（每次约50-200ns per element）
• 极大减少了显存占用峰值

第二件事：计算-通信重叠（Computation-Communication Overlap）

这是最关键也最hardcore的部分。在EP（Expert Parallelism）多GPU场景下：

• 当GPU A正在对本地Expert执行Tensor Core计算时，GPU B也在同时进行计算
• 与此同时，NVLink正在并行传输需要跨GPU交换的token数据
• Mega MoE通过**对称内存（Symmetric Memory）**机制，确保通信和计算访问的是同一块物理内存区域，从而实现真正的流水线并行

从代码层面看，这种重叠的调度极为复杂。DeepSeek在PR中说明Mega MoE仅支持FP8 × FP4 MoE，且要求PyTorch >= 2.9，这暗示了PyTorch 2.9引入了一些对融合算子调度至关重要的新特性（如CUDA Graph的改进、异步执行的优化等）。

3.3 Mega MoE代码实战

以下是基于官方测试代码的Mega MoE使用示例：

import torch
import deep_gemm

# 准备输入数据
num_tokens = 8192
num_experts = 8
hidden = 7168
intermediate_hidden = 18432
num_topk = 2

# 权重（FP4格式，需要UE8M0 scaling factor）
l1_weights = ...  # [num_experts, intermediate_hidden, hidden]
l2_weights = ...  # [num_experts, hidden, intermediate_hidden]

# 获取对称内存缓冲区
# 对称内存是多EP通信的关键：所有GPU访问同一地址
buffer = deep_gemm.get_symm_buffer_for_mega_moe(
    group,                    # torch.distributed.ProcessGroup
    num_experts,
    num_max_tokens_per_rank,  # 最大token数，用于预分配
    num_topk,
    hidden,
    intermediate_hidden
)

# 权重格式转换：FP4 + UE8M0 SF
transformed_l1, transformed_l2 = deep_gemm.transform_weights_for_mega_moe(
    l1_weights, l2_weights
)

# 输入数据填充到缓冲区
buffer.x[:num_tokens].copy_(x_fp8)
buffer.x_sf[:num_tokens].copy_(x_sf)        # 输入缩放因子
buffer.topk_idx[:num_tokens].copy_(topk_idx) # 每个token选中的Expert ID
buffer.topk_weights[:num_tokens].copy_(topk_weights)  # 路由权重

# 一行代码触发整个Mega MoE流水线
y = torch.empty((num_tokens, hidden), dtype=torch.bfloat16, device='cuda')
deep_gemm.fp8_fp4_mega_moe(y, transformed_l1, transformed_l2, buffer)

整个Mega MoE调用，从外部看只是一行，但实际上内部发生的是：

Thread Block 1: [计算] Linear1_GEMM    [通信] NVLink_Recv(token_A→B)
Thread Block 2: [通信] NVLink_Send(token_B→A)  [计算] SwiGLU
Thread Block 3: [计算] Linear2_GEMM    [通信] NVLink_Recv(token_C→B)
                 ...
(所有操作在同一个kernel launch中并行交错执行)

3.4 为什么Mega MoE是MoE工程化的里程碑

回顾MoE优化的历史，我们能看到一条清晰的演进路径：

阶段1到阶段2解决了"慢"的问题（减少launch开销），阶段2到阶段3解决了"等"的问题（消除气泡）。这基本上是MoE计算优化的"最终形态"了——如果还有更进一步，那可能是将Router本身也融合进mega-kernel，但这涉及到动态控制流，在GPU上实现难度极高。

四、FP8xFP4 GEMM：Ultra-low Precision的最后一块拼图

4.1 精度格式的基础知识

在深入FP8xFP4之前，先简要梳理一下AI训练/推理中常用的精度格式：

FP8有两种格式：

• E4M3：1个符号位 + 4个指数位 + 3个尾数位，动态范围窄但精度高，适合值域范围可控的激活
• E5M2：1个符号位 + 5个指数位 + 2个尾数位，动态范围宽但精度低，适合梯度等值域宽广的场景

FP4只有E2M1：1个符号位 + 2个指数位 + 1个尾数位，动态范围仅约±6，精度极低，但存储空间只有FP8的一半。

4.2 FP8xFP4 GEMM的工程挑战

为什么DeepGEMM要做"FP8权重 × FP4激活"这种组合，而不是更常见的FP8×FP8？

答案是：极大模型推理时的显存瓶颈。

在超大规模MoE模型中，激活值（activation，即计算过程中的中间结果）的显存占用有时甚至超过权重本身。以DeepSeek V3这样的超大MoE模型为例：

• 权重：可以通过专家分割分布到多张GPU，每张卡负担可控
• 激活值：每个token都要经过所有参与计算的Expert，产生大量中间结果，必须驻留在单卡显存中

因此，用FP4来存储/计算激活值，能将这部分显存开销再砍一半——这对于需要在有限显存下跑大模型的场景意义重大。

但FP4的工程挑战极为严峻：

挑战一：量化误差累积

E2M1的动态范围只有约[-6, +6]，而激活值通常远超这个范围。需要精心设计的per-token或per-channel缩放因子（scaling factor）来将激活值映射到FP4的可表示范围。DeepGEMM使用UE8M0格式来存储这些缩放因子——4个UE8M0打包进一个torch.int，既保证了精度，又节省了存储。

挑战二：混合精度数值稳定性

FP8×FP4的计算结果在累加过程中会涉及不同的数值范围。需要精心设计累加精度（比如在中间结果用FP32累加），以及缩放因子的传播策略。DeepGEMM借鉴了CUTLASS的fine-grained scaling策略，但做了大量简化。

挑战三：硬件指令支持

NVIDIA Hopper（SM90）和Blackwell（SM100）架构的Tensor Core原生支持FP8计算，但FP4需要更复杂的软件层面模拟。DeepGEMM通过精细的Triton或CUDA实现，在SM90/SM100上实现了FP8输入、FP4权重的高效计算。

4.3 FP8xFP4 GEMM的调用方式

import deep_gemm

# A: FP8格式的激活值 (shape: [M, K])
A_fp8 = torch.randint(0, 255, (1024, 4096), dtype=torch.uint8, device='cuda')
A_sf = torch.rand(1024, 128, device='cuda')  # per-token/per-block scaling factor

# B: FP4格式的权重 (shape: [N, K])，存储为4bit打包
B_fp4_packed = ...  # [N, K // 2], uint8存储两个FP4
B_sf = torch.rand(128, 4096 // 128, device='cuda')  # UE8M0打包的缩放因子

# 执行FP8 × FP4 GEMM
# 输出为BF16
D = deep_gemm.fp8_fp4_gemm(A_fp8, A_sf, B_fp4_packed, B_sf)

4.4 FP4在LLM中的适用边界

FP4是精度损失最大的量化格式，并非所有场景都适用。根据DeepSeek的技术积累和业界经验，FP4的适用场景主要有：

1. 训练良好的大模型的推理阶段：经过大量训练数据校准，激活值分布相对稳定
2. MoE的中间激活：专家内部计算，残差连接少，激活值范围可控
3. 权重量化后的推理：FP4权重 + FP8激活的组合已被验证可行

不适用的场景：模型初期收敛阶段、激活值分布极不稳定的场景、生成式任务的早期token（分布偏异）。

五、FP4 Indexer：让Multi-Token Prediction走得更远

5.1 什么是MQA logits与Indexing

在DeepSeek V3.2及后续版本中引入的Lightning Indexer中，一个关键组件是MQA（Multi-Query Attention）的logits计算。这个logits用于判断token之间的相关性权重，直接影响KV缓存的调度效率。

FP4 Indexer（内部代号可能是对MTP——Multi-Token Prediction的支持扩展）将这部分计算扩展到FP4精度，核心价值在于：

• 支持更大的MTP窗口：FP4带来的显存节省，可以支撑更长的Multi-Token Prediction序列
• 更快的logits计算：FP4的计算密度比FP8更高（同样Tensor Core运算量下，数据量更小）

# FP4 Indexer调用示例（参考官方test代码）
q = torch.randint(0, 255, (seq_len, num_heads, head_dim), 
                  dtype=torch.uint8, device='cuda')  # FP8 query
kv = torch.randint(0, 255, (seq_len_kv, head_dim), 
                   dtype=torch.uint8, device='cuda')  # FP8 key-value

# FP4 logits: [seq_len, num_heads, seq_len_kv]
# 对应每个query token到每个KV token的attention logit
logits = deep_gemm.fp8_fp4_mqa_logits(
    q, kv,
    weights,      # [seq_len, num_heads] 用于加权的权重
    cu_seq_len_k_start,  # KV序列的起始边界
    cu_seq_len_k_end,    # KV序列的结束边界
    clean_logits=True    # 未填充位置设为-inf
)

5.2 为什么Indexing值得单独做优化

Indexing（索引）是LLM推理中的高频操作。每次KV缓存查询，都需要根据当前token的query向量，在所有历史KV token中计算相关性。这个操作的核心是：

logit[i, j] = ReLU(softmax(Q[i] @ K[j])) * weight[i]

其中Q、K是FP8，但计算结果（logits）如果用更高精度存储会占用大量显存。FP4 Indexer将这个结果也量化到FP4，使得：

• MQA的KV Cache可以做得更大
• 多轮对话的上下文窗口可以更长
• 推理时的显存峰值进一步降低

六、PDL（Programmatic Dependent Launch）：消灭kernel依赖的最后一滴CPU开销

6.1 传统kernel依赖管理的代价

在CUDA中，当kernel B依赖kernel A的结果时，传统做法是：

# 伪代码
result_A = kernel_A<<<blocks, threads>>>(args_A)  # async, 不等待
cudaStreamSynchronize(stream)  # CPU侧显式等待A完成
result_B = kernel_B<<<blocks, threads>>>(args_B)  # 现在安全地启动B

问题在于：cudaStreamSynchronize是CPU-GPU同步点。每次同步都需要CPU和GPU之间的一次握手，在高度串行的kernel链路上，这笔开销不可忽视。一个典型的MoE kernel链可能有5-10次这样的同步，累计延迟可达数百微秒。

6.2 PDL的解决方案

PDL（Programmatic Dependent Launch）将依赖管理完全移到了GPU侧：

• 由第一个kernel在完成前主动触发下一个kernel的启动
• CPU只需要发起整个链的第一个kernel
• GPU内部的硬件调度器/软件依赖机制确保正确的执行顺序

这意味着CPU和GPU之间只需要一次握手，而不是N次。DeepSeek PR #304中引入PDL支持，配合Mega MoE一起使用，可以最大化融合算子的效率收益。

启用方式极为简单：

import deep_gemm
deep_gemm.set_pdl(True)  # 启用PDL

七、JIT编译加速：让优化不耽误开发迭代

7.1 传统JIT的问题

DeepGEMM的核心设计哲学之一是JIT编译（Just-In-Time Compilation）——所有kernel在运行时编译，不需要预编译。但这带来了一个问题：首次编译时间。

深度学习训练通常需要反复调试、修改shape、重跑benchmark。如果每次都重新编译kernel，几分钟的编译时间会严重影响开发效率。

7.2 NVRTC加速方案

DeepGEMM从2025年5月就引入了**NVRTC（NVIDIA Runtime Compilation）**支持，本次更新进一步优化了NVRTC的集成效率：

# 启用NVRTC加速，编译速度提升约10倍
export DG_JIT_USE_NVRTC=1

# 可选：设置JIT缓存目录，复用编译结果
export DG_JIT_CACHE_DIR=/path/to/persistent/cache

NVRTC的原理是用NVIDIA提供的运行时编译器（而非完整的NVCC）来编译PTX代码。虽然编译速度更快，但DeepSeek也诚实指出：某些场景下NVRTC编译的kernel可能比NVCC慢1-5%。因此默认为关闭（DG_JIT_USE_NVRTC=0），由用户根据场景选择。

7.3 JIT编译的环境变量体系

DeepGEMM提供了一套完整的环境变量来控制JIT行为：

这套体系让用户可以在开发阶段（快速编译）和生产阶段（极致性能）之间自由切换。

八、GEMM动态优化：MoE GEMM的隐藏性能红利

8.1 Dynamic Swap A/B

在传统的GEMM实现中，A矩阵（activation）和B矩阵（权重）的驻留位置（寄存器/Shared Memory/GMEM）是固定的。但MoE场景下，不同Expert的形状和活跃程度差异很大——有的Expert处理1000个token，有的只处理10个。

DeepGEMM PR #304引入的Dynamic Swap A/B机制，会根据实时的shape和occupancy数据，动态决定A和B的内存布局和交换策略，从而在各种MoE GEMM场景下都能获得接近最优的performance。

这本质上是一个运行时自适应的tile调度器——根据硬件利用率反馈，实时调整数据搬运策略。

8.2 GEMM Heuristics重构

DeepSeek同时对GEMM的启发式配置选择器（heuristics）做了全面重构。原来的heuristics是针对特定shape hard-coded的参数，新的版本则：

• 引入了成本模型来评估不同配置的预期性能
• 支持在线学习：根据实际benchmark数据动态调整配置权重
• 为DeepEPv2 MoE GEMM Layout专门优化了调度策略

九、实战：如何用DeepGEMM构建高性能MoE推理

9.1 环境准备

# 基础环境
# - NVIDIA SM90 (Hopper) 或 SM100 (Blackwell) GPU
# - CUDA 12.3+ (SM90) / CUDA 12.9+ (SM100)
# - Python 3.8+
# - PyTorch 2.9+ (Mega MoE必需)
# - CUTLASS 4.0+

git clone --recursive git@github.com:deepseek-ai/DeepGEMM.git
cd DeepGEMM
./develop.sh    # 编译C++ JIT模块
./install.sh    # 安装Python包

9.2 完整的MoE Forward Pipeline

以下是一个整合了所有新特性的MoE推理示例：

import torch
import torch.distributed as dist
import deep_gemm

# ============ 配置 ============
NUM_EXPERTS = 8
HIDDEN = 7168
INTERMEDIATE = 18432
SEQ_LEN = 4096
TOPK = 2
NUM_GPUS = 8  # EP并行度

# ============ 权重准备（示例，实际从模型加载）============
l1_weights = torch.randn(
    NUM_EXPERTS, INTERMEDIATE, HIDDEN,
    dtype=torch.float16, device='cuda'
).to(torch.float8_e4m3fn)  # FP8

l2_weights = torch.randn(
    NUM_EXPERTS, HIDDEN, INTERMEDIATE,
    dtype=torch.float16, device='cuda'
).to(torch.float8_e4m3fn)  # FP8

# ============ 输入准备 ============
x = torch.randn(
    SEQ_LEN, HIDDEN,
    dtype=torch.bfloat16, device='cuda'
)
x_fp8, x_sf = deep_gemm.cast_to_fp8(x)  # FP8激活 + 缩放因子

# ============ Router ============
# 简化Router实现
router = torch.randn(SEQ_LEN, NUM_EXPERTS, device='cuda')
topk_weights, topk_idx = torch.topk(router, TOPK, dim=-1)
topk_weights = torch.softmax(topk_weights, dim=-1)

# ============ Mega MoE Forward ============
# 初始化EP通信组
group = dist.new_group([i for i in range(NUM_GPUS)])

# 获取对称缓冲区
buffer = deep_gemm.get_symm_buffer_for_mega_moe(
    group, NUM_EXPERTS, SEQ_LEN // NUM_GPUS, TOPK, HIDDEN, INTERMEDIATE
)

# 填充缓冲区
buffer.x[:SEQ_LEN].copy_(x_fp8)
buffer.x_sf[:SEQ_LEN].copy_(x_sf)
buffer.topk_idx[:SEQ_LEN].copy_(topk_idx.to(torch.int32))
buffer.topk_weights[:SEQ_LEN].copy_(topk_weights)

# 权重转换（FP8 → FP8xFP4 Mega MoE格式）
transformed_l1, transformed_l2 = deep_gemm.transform_weights_for_mega_moe(
    l1_weights, l2_weights
)

# 启用优化
deep_gemm.set_pdl(True)
deep_gemm.set_num_sms(132)  # 使用132个SM（H800共132个SM）

# 执行Mega MoE
output = torch.empty((SEQ_LEN, HIDDEN), dtype=torch.bfloat16, device='cuda')
deep_gemm.fp8_fp4_mega_moe(output, transformed_l1, transformed_l2, buffer)

print(f"Mega MoE output shape: {output.shape}")
print(f"Output mean: {output.mean().item():.4f}, std: {output.std().item():.4f}")

9.3 Benchmark脚本

import torch
import time
import deep_gemm

def benchmark_mega_moe(shape, num_runs=100, warmup=10):
    """Benchmark Mega MoE性能"""
    # ... 初始化代码 ...
    
    # Warmup
    for _ in range(warmup):
        deep_gemm.fp8_fp4_mega_moe(output, l1, l2, buffer)
    torch.cuda.synchronize()
    
    # Benchmark
    times = []
    for _ in range(num_runs):
        start = time.perf_counter()
        deep_gemm.fp8_fp4_mega_moe(output, l1, l2, buffer)
        torch.cuda.synchronize()
        elapsed = time.perf_counter() - start
        times.append(elapsed * 1000)  # ms
    
    mean = sum(times) / len(times)
    p50 = sorted(times)[len(times)//2]
    p99 = sorted(times)[int(len(times)*0.99)]
    
    print(f"Mean: {mean:.2f}ms | P50: {p50:.2f}ms | P99: {p99:.2f}ms")
    
    # 计算TFLOPS
    # Mega MoE的计算量：2 * M * K * N (两次GEMM)
    flops = 2 * SEQ_LEN * HIDDEN * INTERMEDIATE * 2
    tflops = flops / (mean * 1e-3) / 1e12
    print(f"Effective TFLOPS: {tflops:.1f}")

十、底层架构解析：DeepGEMM为什么能又快又Clean

10.1 借鉴CUTLASS但不依赖它

DeepGEMM的设计哲学在README中有明确表述："借鉴了CUTLASS和CuTe的思想，但避免了对其模板或代数的重度依赖。"

这背后有深刻的技术洞察。CUTLASS是NVIDIA官方维护的GEMM优化库，提供了极为精细的模板抽象——TileIterator、ThreadMap、MmaAtom……这套体系功能强大，但学习曲线极其陡峭，而且编译时间感人（整个CUTLASS编译一次需要数十分钟）。

DeepGEMM只保留了GEMM优化中最核心的几个概念：

// DeepGEMM核心模块结构
csrc/
  apis/
    gemm.hpp      // GEMM API定义
    attention.hpp // Attention kernel API
    einsum.hpp   // 爱因斯坦求和API
  jit_kernels/
    sm90/         // Hopper架构实现
    sm100/        // Blackwell架构实现
  core/           // 公共底层
deep_gemm/
  include/        // Python绑定的C++头文件
tests/            // 完整测试套件

代码量控制在几千行级别（对比CUTLASS的数万行），但性能不输甚至超越。这种"Less is More"的设计理念，是DeepGEMM最值得学习的工程哲学。

10.2 SM90 vs SM100的差异

DeepGEMM同时支持Hopper（SM90）和Blackwell（SM100）架构，但两者有重要差异：

UE8M0是一种-packed格式，将4个FP32值打包进一个32位整数，这在SM100上由Tensor Core直接支持，效率远超SM90的软件实现。

十一、深度技术细节：那些让性能飞跃的关键trick

11.1 Fine-grained Scaling

DeepGEMM的GEMM核心使用了per-block或per-segment的fine-grained scaling：

# 传统方法：整个矩阵用一个缩放因子
sf = max(abs(A)) / max_val_fp8
A_scaled = A / sf

# DeepGEMM方法：每个block用独立缩放因子
sf = per_block_max(abs(A_block)) / max_val_fp8  # [num_blocks]
A_scaled_block = A_block / sf_view  # 广播

这种精细的缩放策略能更充分地利用FP8的动态范围，减少量化误差。在极大规模矩阵中，这种差异可以累积出显著的精度收益。

11.2 对称内存与NVLink协同

Mega MoE中对称内存（Symmetric Memory）的设计值得单独拿出来讲。

在多GPU EP场景下，传统的做法是：

1. GPU A准备要发给B的数据
2. GPU A通过NVLink发送
3. GPU B接收并存入本地内存
4. GPU B执行计算

这意味着NVLink传输和GPU计算必须串行——因为数据的目的地不同。

对称内存的做法是：A和B事先协商好一块"虚拟共享内存"，A写入的位置恰好是B读取的位置，NVLink硬件可以在A写入的同时让B开始读取，物理上对齐传输和计算的时序。这需要NVLink硬件和软件栈的深度协同。

DeepSeek的PR注释中提到"Overlapping NVLink communication and tensor core computation"，正是这种对称内存机制的实现。

11.3 TMA（Tensor Memory Access）的精细使用

Hopper架构引入的TMA（Tensor Memory Access）是DeepGEMM性能的关键之一。TMA允许线程块以 collective 的方式执行 coalesced 的全局内存到共享内存的搬运，相比传统的 __ldg + 手动 shared memory 填充，有几个优势：

• 硬件级别的2D寻址：自动处理非对齐和跨步访问
• 自动warp级同步：TMA操作在warp内部是隐式同步的
• 支持异步执行：TMA可以与计算overlap

DeepGEMM的kernel大量使用TMA进行A、B矩阵的加载，配合Double Buffering（双缓冲）策略，使得数据加载时间完全被计算时间覆盖。

十二、性能数据与对标

12.1 历史性能数据回顾

根据DeepGEMM的Release History，核心性能指标如下：

1550 TFLOPS是什么概念？H800的理论峰值是989 TFLOPS（FP8稠密）或1979 TFLOPS（FP8稀疏）。DeepGEMM在稠密FP8下跑出了接近稀疏峰值的性能，这说明其在利用率和tile调度上做到了极致。

12.2 Mega MoE的性能预期

DeepSeek官方表示Mega MoE的详细性能数据将在后续发布。但从技术分析可以推断：

• 计算-通信重叠理论上可将端到端MoE延迟降低30-50%（取决于MoE的E/P比例）
• 5-in-1融合可消除约50μs的kernel launch开销
• 动态Swap A/B在非均衡负载的MoE场景下可获得**10-20%**额外加速

十三、对AI Infra生态的影响

13.1 开源库的角色

DeepGEMM不是DeepSeek的"独门绝技"，而是整个开源AI Infra生态的公共基础设施。它的价值体现在几个层面：

对框架开发者：可以用DeepGEMM替代手写的low-performance GEMM，快速获得高性能算子
对大模型研究者：可以用DeepGEMM作为benchmark基线，评估新架构的计算效率
对学习者：DeepGEMM的Clean代码是GPU内核优化的最佳教科书

13.2 与DeepSeek V4的关联

新浪财经报道，DeepSeek V4将首发支持FP4精度，并且会深度适配NVIDIA SM100/Blackwell GPU。结合这次DeepGEMM的更新来看，V4很可能会是：

• 大量使用FP8×FP4混合精度
• MoE层全面采用Mega MoE架构
• DeepEPv2作为分布式通信基础

这意味着DeepGEMM不只是"工具库"，而是DeepSeek大模型技术栈的核心基础设施。

十四、总结与展望

这次更新的核心收获

回顾DeepGEMM 2026年4月更新，有三个最重要的技术演进：

1. Mega MoE代表着MoE计算从"多kernel流水线"到"一体化融合"的范式跃迁。 计算与通信的完全重叠，是分布式训练/推理优化的终极目标，DeepSeek率先在开源领域实现了工程化落地。

2. FP4进入实用化阶段，标志着Ultra-low Precision在LLM中的适用边界正在快速扩展。 从FP32到FP16到FP8再到FP4，每一次精度压缩都带来了可观的效率收益，但也伴随着越来越复杂的工程挑战。DeepGEMM证明了FP8×FP4是可行的，这为未来的FP2精度探索铺平了道路。

3. 工程细节决定成败。 PDL、Faster JIT、Dynamic Swap A/B……这些看起来不起眼的优化，累积起来就是质的飞跃。DeepSeek在底层工程上的极致追求，是其技术护城河的核心组成部分。

对开发者的建议

如果你正在做LLM训练/推理优化：

1. 优先使用DeepGEMM而非手写GEMM：除非你有超越DeepGEMM的内核优化能力，否则没有重复造轮子的必要
2. 关注PyTorch 2.9+的新特性：Mega MoE需要PyTorch 2.9，PyTorch生态的演进与DeepGEMM的性能紧密绑定
3. 建立性能基准：用DeepGEMM建立baseline，持续追踪不同shape下的性能变化
4. 学习它的代码：DeepGEMM是GPU内核优化最好的教科书，值得逐行研读

展望未来

DeepGEMM的演进路径已经清晰可见：

• FP2精度支持（如果未来硬件支持的话）
• 更大的Mega MoE融合范围（将Router也融入mega-kernel？）
• 对更多硬件架构的支持（AMD ROCm？国产GPU？）

无论如何演进，DeepGEMM"Clean + High Performance"的设计哲学不会改变。在AI Infra这个越来越复杂的领域里，它像一股清流，证明了简单和极致可以并存。

参考资料

1. DeepSeek DeepGEMM GitHub: https://github.com/deepseek-ai/DeepGEMM
2. PR #304 - Mega MoE, FP4 Indexer: https://github.com/deepseek-ai/DeepGEMM/pull/304
3. DeepSeek V4消息: https://finance.sina.com.cn/tech/roll/2026-04-17/doc-inhuszrw2758980.shtml
4. DeepSeek悄悄更新 Mega MoE: https://www.toutiao.com/article/7629560298604610058/
5. DeepGEMM性能优化秘籍: https://blog.csdn.net/gitblog_00475/article/details/152251300
6. NVIDIA CUTLASS: https://github.com/NVIDIA/cutlass
7. NVIDIA CuTe: https://github.com/NVIDIA/cutlass/tree/main/include/cute
8. DeepEP (Expert Parallelism通信库): https://github.com/deepseek-ai/DeepEP