DeepEP 完全拆解：DeepSeek 开源通信库如何把 MoE 推理的 dispatch/combine 压到硬件极限

专家并行（EP）：把专家摊到多张 GPU 上#

先复习一下 MoE 层的推理过程。一个 MoE 层包含一个路由器（router/gating network）和一组专家 FFN。对每个输入 token，路由器计算它对每个专家的评分，选出 top-k 个专家，token 的隐藏状态被送到这些专家里各自过一遍 FFN，最后按路由权重加权求和。用符号表示：设 token 数为 $B$，专家数为 $N$，每个 token 选 $K$ 个专家（top-k），隐藏维度为 $H$。

在单张 GPU 上，所有 $N$ 个专家都在这张卡里，路由只是计算，不需要通信。但模型大到一定程度（比如 DeepSeek-V3 的 256 个专家、671B 参数），一张 GPU 装不下全部专家，就必须把专家分布到多张 GPU 上——这就是专家并行（EP）。EP 下，任意一个 token 的 top-k 专家很可能分布在不同的 GPU 上，token 数据就必须在 GPU 之间搬运。

EP 组里的每张 GPU（一个 rank）要做两件事：

• dispatch（分发）：把本 rank 上的 token 按路由结果发送到持有对应专家的远程 rank；
• combine（合并）：专家计算完后，把结果按路由权重加权求和，送回 token 原本所在的 rank。

vLLM MoE 实践指南（AMD ROCm 博客的 vLLM MoE Playbook）把整个 MoE 层的数据流画得很清楚：注意力层每个 GPU 各自处理自己的请求，到专家层时通过 AllToAll（dispatch）把 token 重分布到持有专家的 GPU，专家算完后再 AllToAll（combine）送回：

实际部署中 MoE 层还常与张量并行（TP）混合使用：注意力层和 dense 层按 TP 切分（每卡算 1/8，层间 AllReduce 同步），专家层则按 EP 摊到各卡——这样每个 token 的 KV 计算走 TP、专家计算走 EP，两种并行模式的通信模式完全不同（AllReduce 是”大家算一部分、结果汇总”，all-to-all 是”数据按路由重分布”）：

通信量有多大：一个公式算出账单#

每层每张卡要搬多少数据？dispatch 阶段，每个 token 要发给 $K$ 个专家，每个专家的输入是隐藏向量 $x \in \mathbb{R}^H$。假设每个 rank 上平均有 $b$ 个 token，数据类型占 $s$ 字节（BF16 为 2，FP8 为 1），那么每个 rank 单层需要发出的逻辑数据量约为：

combine 阶段数据量相同（每个专家输出也是 $H$ 维向量，且要附带路由权重）。以 DeepSeek-V3 生产配置为例：解码阶段每个 rank 约 128 个 token、$K=8$、$H=7168$、FP8 分发（$s=1$）：

7.3 MB 看起来不大，但注意这是每一层的通信量，而一个 MoE 大模型有几十上百个 MoE 层，每个 token 每层都要搬一次。更重要的是：这 7.3 MB 必须在极短时间内搬完。如果预算 80 微秒，需要的带宽就是约 95 GB/s——而一张 H800 连的单张 InfiniBand CX7 网卡只有 400 Gb/s（约 50 GB/s）。也就是说，通信几乎是物理上不可能的，除非把架构设计到极致，或者想办法减少真正跨网络的流量。这个矛盾就是 DeepEP 存在的全部理由。

带宽不对称：NVLink 与 RDMA 之间的鸿沟#

MoE 通信横跨两个带宽截然不同的域：

NVLink 比 RDMA 快 3 倍以上。把数据从 NVLink 域搬到 RDMA 域（或反过来）的速度，由慢的那一端决定——这就像从宽 160 米的路拐进宽 50 米的路，车流不可能超过 50 米路的容量。所以任何好的 MoE 通信设计都必须回答同一个问题：如何让数据在昂贵的 RDMA 链路上少跑，在便宜的 NVLink 链路上多跑？

DeepSeek-V3 的 group-limited gating 让问题更特殊#

DeepSeek-V3 的 MoE 路由（DeepSeek-V3 论文）采用分组受限门控（group-limited gating）：256 个专家被分成 128 组，每组 2 个专家；路由器先选 top-4 组，再在这 4 组里选出 top-8 个专家。部署时，同一组的两个专家会被尽量放在同一节点上。这样一来，每个 token 的 8 个专家都集中在少数几个节点——跨节点的数据量小但集中，节点内的转发需求大。

这个特性直接催生了 DeepEP 的第一个设计点：常规 kernel 专门针对”非对称域带宽转发”（asymmetric-domain bandwidth forwarding）优化，也就是数据从 NVLink 域转发到 RDMA 域、或从 RDMA 域转发回 NVLink 域的场景。比如某个节点收到跨节点来的 token 后，还要通过 NVLink 转发给节点内持有不同专家的其他 GPU。

朴素 all-to-all 的三个问题#

在 DeepEP 之前，MoE 的跨节点通信一般用 NCCL 的 AllToAll，它有三个明显的浪费：

1. 重复发送：一个 token 被路由到同一节点上的多个专家（group-limited gating 下这很常见），朴素实现会对每个专家各发一份跨节点数据，同一份数据在 RDMA 上反复传输；
2. 不利用 NVLink：节点内转发没有跟跨节点传输联动，慢速链路成了瓶颈；
3. 同步开销大：collective 通信需要全局同步点，rank 之间互相等待，长尾效应严重——通信耗时取决于数据量最大的那个连接。

DeepEP 的每个设计几乎都是针对这三条逐个击破的。

非对称带宽转发与两阶段通信#

常规 kernel 的完整数据流可以概括为两阶段：节点内聚合，节点间单次传输。

以 dispatch 为例。设一个节点上有 4 张 GPU（rank 0-3），每个 token 的 top-8 专家可能落在本节点或其他节点。DeepEP 的做法是：

• 阶段一（节点内）：每个 rank 通过 NVLink 把需要发给同一远程节点的 token 数据汇总（聚合/去重）到该节点的一个”出口” rank；
• 阶段二（节点间）：出口 rank 通过 RDMA 把聚合后的数据一次发给目标节点；
• 目标节点收到后，再通过 NVLink 把数据转发给节点内持有对应专家的各个 GPU。

去重的效果立竿见影：一个 token 被路由到同一节点上的多个专家时，跨节点只传一份数据，到达后再复制。combine 方向反过来做分层归约（hierarchical reduction）：同一 token 在多个专家上的输出，先在节点内通过 NVLink 加权求和归约成一份，再把这一份结果通过 RDMA 送回原 rank——跨节点流量同样被大幅压缩。

这就是”非对称域带宽转发”的全部含义：让慢速的 RDMA 只承载”每个 token 每节点一份”的最小流量，让快速的 NVLink 承担聚合、复制与转发。带宽不对称问题被结构性地规避了。

SM 数量控制：把通信 kernel 塞进少量 SM#

常规 kernel 支持通过 Buffer.set_num_sms(N) 显式控制参与通信的 SM 数量。为什么需要这个？因为通信 kernel 和计算 kernel 共用 GPU 的 SM，如果通信 kernel 独占所有 SM，计算就没法同时进行；如果只给它一小部分 SM，它就有足够多的空闲周期，可以和计算并行推进。DeepSeek-V3 训练中，H800 上最多约 20 个 SM 被分配给通信相关操作（README 示例代码默认 set_num_sms(24)），剩余 110+ 个 SM 专心做 GEMM。这就是”用 SM 数换重叠度”的思路：通信慢一点没关系，只要它在计算的阴影里完成，总时间就是两者取 max 而不是相加。

隐式的 CPU 等待：dispatch 里的一个微妙设计#

常规 kernel 的 dispatch 有个细节：调用方一开始不知道本 rank 会收到多少 token（取决于全局路由结果）。所以 dispatch 内部有一个隐式的 CPU 等待：GPU 端先跑一个计数内核，算出每个 rank 的接收 token 数，把信号发回 CPU，CPU 等这个信号到达后再分配接收缓冲、启动真正的数据搬运。官方文档用下面这张图描述这个流程：

这个”CPU 等 GPU 信号”的设计意味着常规 kernel 不兼容 CUDA graph（除非在节点内场景指定 num_worst_tokens，用最坏情况预留缓冲）。这正是低延迟 kernel 存在的另一个理由——解码阶段必须兼容 CUDA graph。

性能数据：逼近 NVLink 与 RDMA 的极限#

官方在 H800（NVLink 约 160 GB/s）+ CX7 400 Gb/s 网卡（约 50 GB/s）上、按 DeepSeek-V3 预训练配置（4096 token/批、hidden 7168、top-4 组、top-8 专家、FP8 分发、BF16 合并）测得：

节点内 153–158 GB/s 意味着 NVLink 利用率超过 95%（按官方标注的 160 GB/s 峰值算）。节点间最有趣的是 EP32 反而比 EP16 高（58 vs 43 GB/s）——EP 规模增大时，每个 rank 的发送目标变多，聚合效应更强，跨节点流量分布更均匀，直到 EP64 才因更多 rank 之间的同步开销回落。这也解释了生产环境把 EP 组规模取在 32-64 这个区间：组太小聚合效应不足，组太大同步开销会吃掉收益。

为什么解码阶段需要完全不同的设计#

解码（decode）阶段和 prefill 有本质区别：每个 rank 每轮只有几十到两百多个 token（KV 缓存注意力逐个生成），对延迟极度敏感——每多 100 微秒，用户感知的 TTFT/TPOT 就多 100 微秒。这时候：

• 高吞吐 kernel 的 NVLink 转发路径多一跳，延迟高；
• 占 SM 的通信 kernel 会挤占计算资源，拖慢 GEMM；
• CPU 参与同步会引入不可控的抖动。

所以低延迟 kernel 做了三个激进选择：纯 RDMA（不做 NVLink 转发）、SM 只负责发指令然后立即退出、用 hook 实现计算-通信重叠。

纯 RDMA：为延迟砍掉转发路径#

低延迟 kernel 完全走 RDMA，节点内转发被去掉。代价是节点内的流量也走网卡（所以官方要求 QP 数量必须等于本地专家数，每个专家对应一个 Queue Pair 以保证并发度），以及缓冲内存消耗巨大（官方建议 num_max_dispatch_tokens_per_rank 小于 256，否则显存扛不住）。这是典型的”用显存和带宽换延迟”。

hook 机制：RDMA 在后台跑，GPU 一个 SM 都不占#

这是 DeepEP 最精妙的设计。调用 low_latency_dispatch(..., return_recv_hook=True) 后，kernel 做的事只是：SM 把 RDMA 工作请求（WQE）提交出去，然后立刻退出。真正的数据传输由网卡异步完成，完全不占用 SM。提交动作本身只花微秒级时间，而传输时间（几十到几百微秒）全部躲进了计算阶段的阴影里。

等待接收数据的时机由调用方通过 hook 控制：hook() 被调用时才真正消费接收缓冲。官方文档用下面这张图说明双微批（double-batch）重叠的用法——每个微批的注意力/分发/专家计算/合并四段交错，RDMA 流量在后台持续流动：

这正是 DeepSeek-V3/R1 推理系统”双批重叠”机制的底层实现（系统层面还叠加了 DualPipe 的跨层流水线）。vLLM 的 EP 解码路径、SGLang 的 DeepEP 后端都用这个 hook 接口。由于 kernel 提交后立即返回、没有 CPU 等待 GPU 信号，低延迟模式兼容 CUDA graph（replay 时需恢复 buffer 状态）。

为什么 combine 比 dispatch 慢？#

先看数据。官方在 H800 + CX7 上、按生产配置（128 token/批、hidden 7168、top-8 专家、FP8 分发、BF16 合并）测得：

几个值得解释的点：

1. combine 延迟约是 dispatch 的 1.5 倍。combine 除了传输，还要做加权求和（乘以 top-k 权重再累加），并且跨节点的归约涉及 RDMA 原子操作（fetch-and-add 类），原子操作在 InfiniBand 上的开销远高于普通 put/get。dispatch 是纯数据搬运，combine 是”搬运 + 归约”，天然更重。
2. 表格里的”RDMA 带宽”是逻辑带宽口径。V2 README 的注释明确说明：“这些结果是逻辑带宽，例如 EP 8×2 场景下的 90 GB/s 实际包含了本 rank 内的流量”。低延迟模式中，被路由到本 rank 本地专家的 token 不走网卡，但仍计入数据量；所以 EP8 的 98 GB/s 和 127 GB/s 都不能直接跟单网卡 50 GB/s 峰值比较。真实网卡流量低于表格数字，这是评估任何这类 benchmark 都要注意的口径。
3. EP 规模越大延迟越高（77 µs → 194 µs），因为数据要发往更多 rank、每个 rank 的同步与原子操作变多；EP256 时延迟与 EP128 基本持平，说明此时已经完全是带宽受限（39-40 GB/s 接近网卡上限）。
4. 对照前面算的账：EP8 场景下 7.3 MB 逻辑数据在 77 µs 内完成，等效约 95 GB/s——与官方 98 GB/s 吻合，说明 dispatch 是纯带宽受限的，kernel 本身的开销已经被压到接近零。

统一接口与解析式资源计算#

V1 里高吞吐和低延迟是两套 API（Buffer + 手动指定 SM 数），V2 统一成单个 ElasticBuffer 接口，SM 数与 QP 数不再需要手动调优，而是用 get_theoretical_num_sms() 按带宽模型解析式计算——输入节点内/节点间带宽（sm_read_gbs、rdma_gbs、nvlink_gbs 等）与是否重叠的标志，直接算出最优 SM 数。官方 README 给出示例：

1
_buffer = ElasticBuffer(
2
    group,
3
    num_max_tokens_per_rank=num_max_tokens_per_rank,
4
    hidden=hidden,
5
    num_topk=num_topk,
6
    use_fp8_dispatch=use_fp8_dispatch,
7
)
8
# V2 解析式计算最优 SM 数，无需再手动调优
9
_num_comm_sms = _buffer.get_theoretical_num_sms(num_experts, num_topk)

V2 的性能与代价#

相比 V1：峰值性能最高提升 1.3 倍，SM 占用最多节省 4 倍——V3 风格训练场景的通信 SM 从 24 个降到 4-6 个，性能持平或更好。EP 规模上限从几十扩展到 EP2048。

V2 还新增了三个”0 SM”实验特性：Engram（远程内存访问）、PP（流水线并行通信）、CP（上下文并行，用 Copy Engine）。代价也很明确：缓冲内存消耗比 V1 大；V1 的 0-SM 纯 RDMA 低延迟模式在 V2 中不再支持（低延迟路径并入 Elastic kernel，占用少量 SM）。另外社区反馈的已知问题：ld.global.nc 这类激进 PTX 仍在使用（可用环境变量关闭），部分场景下逻辑带宽口径依旧。

集成与社区分支#

• 推理框架：vLLM 与 SGLang 都以 DeepEP 作为 EP 后端；Megatron-LM 也提供接入。
• 零拷贝分支：腾讯网络平台部的 PR 去掉了 PyTorch tensor 与通信 buffer 之间的拷贝，显著降低常规 kernel 的 SM 占用（zero-copy PR）。
• hybrid-ep 分支：用 TMA 指令实现最小 SM 占用、支持更大 NVLink 域、单 batch 细粒度重叠，还支持无 NVLink 环境的 PCIe kernel 与 NVFP4 数据类型。
• AntGroup-opt 分支：Normal-SMFree（把 SM 从 RDMA 路径中彻底摘除）、LL-SBO（用信号机制把 Down GEMM 与 combine 发送重叠）、LL-Layered（跨节点低延迟算子用 rail 优化转发）。
• MORI：AMD 团队基于 ROCm 的 MORI 后端让 DeepEP 在 AMD GPU 上跑通（低延迟模式）。

腾讯 TRMT 的三层优化#

腾讯云在其技术博客（腾讯云社区《万字详解大模型推理加速核心原理》）中披露了对 DeepEP 的深度改造：控制面用 IBGDA 内存语义 API 把计算与通信封装为统一 kernel，消除 kernel 启动开销，控制面延迟从 3 µs 降到 0.5 µs；数据面通过 PTX 指令精细控制数据在 HBM 与 L1/L2 之间的搬移，避免通信流量冲刷掉模型权重与 KV 缓存（bypass L2）；传输层构建全互联架构，用动态 QP 端口分配让双端口网卡同时跑满带宽，弥补 DeepEP 只考虑单网卡的不足。整体通信算子耗时减少约 60%。

局限#

• 硬件绑定：V1 依赖 NVSHMEM + IBGDA，且要求 mlx5 网卡驱动，可移植性差（社区论文对比中，主机代理（host-proxy）方案在 EFA 等网络上更通用）；
• 精度与确定性：节点内 BF16 预归约带来精度损失与非确定性，对要求逐位复现的训练场景是隐患；
• RoCE 兼容：官方只在 InfiniBand 上充分测试，RoCE 是”理论上兼容”；
• V2 的取舍：缓冲内存更大、0-SM RDMA 低延迟模式被移除。