Garden | DeepEP

MoE 提到：一层 MoE 要搬两次 token 隐藏态，dispatch 一次、combine 一次，量级随 $d \times k$ 增长。EP 一旦跨出单机，这两次 all-to-all 就落在 IB 上，未优化时能吃掉六成训练时间。

DeepEP 是 DeepSeek 开源的专为 MoE 和专家并行（Expert Parallelism, EP）设计的通信库¹。提供了一系列优化的通信 Kernel，实现了以下能力：

高度优化的 all-to-all 通信
- 分层协作：节点内走 NVLink + NVSwitch，节点间走 RDMA
- 提供 Dispatch 和 Combine 两个通信原语
两套针对不同场景的 kernel：
- 高吞吐 kernel 面向训练与推理 prefill。节点内 NVLink + 节点间 RDMA 通信。
- 低时延 kernel 面向 Inference Decoding。使用纯 RDMA 通信来最小化时延。
原生 FP8 dispatch，通信量相比 BF16 减半
显式的 SM 数量控制，让通信 kernel 能和计算 kernel 抢得可控，从而做 overlap

为什么 NCCL 的 all-to-all 不够用

在 DeepEP 之前，MoE 框架普遍基于 NCCL 的 p2p send/recv 实现 all-to-all 通信，但是带宽利用率并不高，参考 Alltoall > NCCL all2all 的问题。

第一，带宽是不对称的，而 all-to-all 这个原语假装它是对称的。 H800 上 NVLink 提供约 160 GB/s，CX 7 IB 网卡约 50 GB/s，差 3.2 倍（H800 的 NVLink 本身已经为出口合规从 900 GB/s 砍下来过一刀²）。All-to-all 的语义里没有「目标 rank 在不在本机」这回事，一个 token 要发给 8 张远端 GPU，就是 8 份独立的 IB 流量；哪怕这 8 张卡在同一台机器上，NVLink 那 3.2 倍的余量也一点用不上。

第二，长度是数据相关的。 all-to-all 要求预先知道收发长度，而 MoE 的收发长度取决于 router 当前这一步吐出什么。标准做法是两阶段：先用一次小 all-to-all 换计数，再搬数据。第一阶段的结果必须回到 host 才能分配输出 tensor，这就插入了一次强制的 device-host 同步。

第三，通信要和计算抢 SM。 GPU 上的通信 kernel 不是免费的旁路，它占 SM、占 L2、占带宽。NCCL 用多少 SM、什么时候用，调用方基本无法控制；而 MoE 训练恰恰需要把 all-to-all 藏进 GEMM 的影子里，这要求通信侧的资源占用是一个可以调的旋钮。

DeepEP 对这三件事各给了一个答案：

分层转发。这是 DeepEP 性能优势的主要来源
把 layout 计算做成一等公民
把 SM 数量做成参数

Normal 模式下 DeepEP API 使用

1. 创建 Buffer

1
2
3


from deep_ep import Buffer

buffer = Buffer(group, num_nvl_bytes, num_rdma_bytes)

DeepEP 不像 NCCL 那样每次调用临时分配显存，而是一次性预分配一大块通信 buffer 反复用。

num_nvl_bytes 是机内 NVLink 通信区
num_rdma_bytes 是跨机 RDMA 通信区

1
2
3
4
5
6
7


num_nvl_bytes = max(
    config.get_nvl_buffer_size_hint(hidden_bytes, group.size()), num_nvl_bytes
)

num_rdma_bytes = max(
    config.get_rdma_buffer_size_hint(hidden_bytes, group.size()), num_rdma_bytes
)

Buffer 不是「一块缓存」，而是「一块所有 rank 都能直接读写的共享显存」。

2. 计算 dispatch layout

get_dispatch_layout 完全在 GPU 本地计算，输入只有 router 出来的 topk_idx（形状 [num_tokens, topk]）。

注意 get_dispatch_layout 是一个独立的 API，而不是塞进 dispatch 里 ——这样你可以提前算、和别的计算 overlap

1
2
3
4
5


(num_tokens_per_rank,      # [num_ranks] 的数组, 我要发给每个 rank 多少 token
 num_tokens_per_rdma_rank, # [num_rdma_ranks] 的数组, 我要发给每台机器多少 token
 num_tokens_per_expert,    # [num_experts] 的数组, 我要发给每个 expert 多少 token
 is_token_in_rank,         # [num_tokens, num_ranks] 的 bool 矩阵, 标记 token 是否从某个 rank 发送
 event) = buffer.get_dispatch_layout(topk_idx, num_experts)

3. Dispatch Forward

1
2
3
4
5
6
7


recv_x, recv_topk_idx, recv_topk_weights, num_recv_tokens_per_expert_list, handle, event = \
    buffer.dispatch(
        x,                          # [num_tokens, hidden]
        topk_idx=..., topk_weights=...,
        num_tokens_per_rank=..., num_tokens_per_rdma_rank=...,
        is_token_in_rank=..., num_tokens_per_expert=...,
    )

DeepEP 常规 kernel 的 CPU/GPU 时序：CPU launch notify 后必须等 GPU 回传 tensor 尺寸才能分配显存，随后 launch dispatch；combine 直接复用 dispatch 留下的 layout 信息，不必再等一次

notifiy_dispatch

CPU 自旋等待

dispatch kernel

4. Combine Forward

注意这里只传 handle，不用再传任何 layout 信息。

这是 API 设计上最值得注意的一点：handle 是 dispatch 返回的一个不透明对象，里面封装了「这次通信的完整路由表」。combine 就是拿着它把 dispatch 反着走一遍。

1

combined_x, _, event = buffer.combine(y, handle)

5. Dispatch Backward

因为 dispatch 和 combine 互为逆运算，反向传播根本不用写新 kernel——dispatch 的 backward 就是 combine，combine 的 backward 就是 dispatch。

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11


def dispatch_backward(grad_recv_x: torch.Tensor, grad_recv_topk_weights: torch.Tensor, handle: Tuple) -> \
        Tuple[torch.Tensor, torch.Tensor, EventOverlap]:
    global _buffer

    # The backward process of MoE dispatch is actually a combine
    # For more advanced usages, please refer to the docs of the `combine` function
    combined_grad_x, combined_grad_recv_topk_weights, event = \
        _buffer.combine(grad_recv_x, handle, topk_weights=grad_recv_topk_weights, async_finish=True)

    # For event management, please refer to the docs of the `EventOverlap` class
    return combined_grad_x, combined_grad_recv_topk_weights, event

6. Combine Backward

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13


def combine_backward(grad_combined_x: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],
                     handle: Tuple, previous_event: Optional[EventOverlap] = None) -> \
        Tuple[Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]], EventOverlap]:
    global _buffer

    # The backward process of MoE combine is actually a dispatch
    # For more advanced usages, please refer to the docs of the `dispatch` function
    grad_x, _, _, _, _, event = _buffer.dispatch(grad_combined_x, handle=handle, async_finish=True,
                                                 previous_event=previous_event,
                                                 allocate_on_comm_stream=previous_event is not None)

    # For event management, please refer to the docs of the `EventOverlap` class
    return grad_x, event

Low Latency 模式下 DeepEP API 使用

1. 创建 Buffer

1
2
3
4
5
6
7


num_rdma_bytes = Buffer.get_low_latency_rdma_size_hint(
    num_max_dispatch_tokens_per_rank, hidden, group.size(), num_experts)

assert num_experts % group.size() == 0
_buffer = Buffer(group, 0, num_rdma_bytes,           # 注意 nvl_bytes = 0
                 low_latency_mode=True,
                 num_qps_per_rank=num_experts // group.size())

num_nvl_bytes 传 0——纯 RDMA，连机内也走网卡不走 NVLink。
QP 数必须等于本地 expert 数（官方标了 must）。每个本地 expert 独占一条 RDMA 队列，这样不同 expert 的数据互不阻塞。这是「为低延迟而设计」的典型体现。
显存消耗远大于 normal 模式，所以官方建议 num_max_dispatch_tokens_per_rank（即 decode batch size）小于 256。它是按最坏情况静态分配的——这正是它能兼容 CUDA Graph 的原因（没有动态 size，也就不需要 CPU 同步）。

不对称带宽：把 IB 流量除掉一个 fan-out

考虑 EP=64，8 台机器每台 8 卡，256 个 routed expert 每卡 4 个，每个 token 选 8 个 expert。

先问一个纯组合数的问题：一个 token 平均会命中多少张不同的卡、多少台不同的机器？把 $E$ 个 expert 均分成 $G$ 组，某一组完全没被命中意味着 $k$ 个选择全落在这组之外，于是命中组数的期望是

$$ \mathbb{E}[\text{命中组数}] = G\left(1 - \frac{\binom{E - E/G}{k}}{\binom{E}{k}}\right) $$

代进去（$E=256$、$k=8$）：

分组方式	$G$	每组 expert 数	命中组数期望
按卡分（EP=64）	64	4	7.68
按机器分（8 机）	8	32	5.30

8 个选择几乎总是落在 8 张不同的卡上（7.68），但这 7.68 张卡只分布在 5.30 台机器上。这个差值就是可以省下来的 IB 流量。

朴素实现按卡发，每张远端卡一份独立的 IB 拷贝。分层转发按机器发：一个 token 对每台目标机器只过 IB 一次，落地后由那台机器的接收 GPU 通过 NVLink 转发给机内所有目标卡。把本机那部分剔掉（本机 8 张卡平均被命中 0.96 张，本机被命中的概率 0.66），真正过网的量是：

实现方式	每个 token 的 IB 拷贝数
朴素按卡发（每张远端卡一份）	$7.68 - 0.96 = 6.72$
分层转发（每台远端机器一份）	$5.30 - 0.66 = 4.63$

IB 流量降到 0.69 倍，也就是 1.45 倍加速——前提是新增的 NVLink 那一跳能被完全藏住。

这个比例还能继续往下压——只要愿意约束路由。DeepSeek-V3 的 node-limited gating 规定每个 token 最多发往 $M=4$ 个节点，于是 IB 侧的拷贝数被硬压到 4 以内（上面的组合数公式此时不再适用，因为路由不再是自由的），而 NVLink 侧要多承担一些 fan-out。这笔交换划不划算，取决于 NVLink 还有多少余量，而余量恰好就是那个 3.2 倍：

机内 fan-out 的免费额度 = NVLink 带宽 / IB 带宽。 只要平均每台机器展开的 expert 数不超过 3.2，NVLink 那一跳就能完全藏在 IB 那一跳背后，端到端时间仍由 IB 决定。

V3 论文把这句话说成了一个很漂亮的结论：4 个节点 × 3.2 expert/节点 = 12.8，所以「虽然实际只选 8 个 routed expert，在相同通信成本下这个数字最多可以放大到 13 个」³。DeepEP 的常规 kernel 就是把这句话变成代码的那部分——它必须让 IB 与 NVLink 两跳真正流水起来，而不是先收完再转发，否则 3.2 这个额度立刻退化成串行的加法。

顺带说明一个常见误读：这里的两跳不是「先发到某台机器的 rank 0 再散开」。落点是目标节点上与发送方节点内序号相同的那张卡（nvl_rank 相同），所以 8 张卡的 IB 流量天然分散在 8 张网卡上，不会在某一张卡上形成热点。这也解释了 DeepEP 高吞吐模式为什么和 NVSHMEM 的 NVSHMEM_HCA_PE_MAPPING 冲突——它把 NVSHMEM 意义上的 PE 定义成了 rdma rank，同一节点内所有 PE 的 mype_node 都是 0，NIC 映射会全部选到同一张网卡上，细节见 NVSHMEM 那篇。

两套 kernel：吞吐与时延是两个问题

同一个 all-to-all，训练和 decode 想要的东西正好相反。DeepEP V1 干脆写了两套。

	常规 kernel（normal）	低时延 kernel（low-latency）
目标场景	训练、推理 prefill	推理 decode
优化目标	打满带宽	压低单次延迟
数据路径	IB 到对端同号卡，再 NVLink 转发（两跳）	纯 RDMA 直达目标卡，机内也不用 NVLink
NVSHMEM 传输层	IBRC（CPU 代理线程）	IBGDA（GPU 直接发起）
收发长度	先 `notify_dispatch` 换计数，按实际长度分配	不换计数，按最大容量静态分配
显存开销	小（环形队列复用）	大（每 expert 预留 `max_tokens × num_ranks`）
CUDA Graph	不兼容（有 CPU 等待）	兼容
overlap 方式	独立 comm stream + event	recv hook，不占 SM
典型数字（H800）	EP8 内节点 153 GB/s；EP32 跨节点 58 GB/s	EP8 dispatch 77 μs；EP256 dispatch 194 μs

两列的差异全部来自同一个取舍：是否愿意为了省一次同步而预留最坏情况的显存。常规 kernel 选择省显存，代价是 CPU 必须等 GPU 告诉它要收多少；低时延 kernel 选择省同步，代价是 buffer 按 num_max_dispatch_tokens_per_rank × num_ranks 开——这也是官方建议 decode batch size 控制在 256 以内的原因。

顺着这个取舍还能看懂时延表里一件反直觉的事：combine 的延迟大约是 dispatch 的两倍（EP256 时 360 μs vs 194 μs）。combine 搬的是 BF16 而 dispatch 搬的是 FP8，量翻一倍；而且 combine 要做加权求和，收到的每一份都得读出来累加，不是单纯的拷贝。

常规 kernel：channel 与 warp specialization

先算 layout，再搬数据

dispatch 之前要先调 get_dispatch_layout，它拿本地的 topk_idx 算出四张表：num_tokens_per_rank、num_tokens_per_rdma_rank、num_tokens_per_expert、is_token_in_rank。这一步在 GPU 上做，纯本地，不通信。

真正的通信从 notify_dispatch 开始——它是那次「只搬计数」的小 all-to-all 的 DeepEP 版本，但换到的东西比 output_splits 多得多：

rdma_channel_prefix_matrix (num_rdma_ranks, num_channels)：每个 channel 发往每个节点的 token 数前缀和
gbl_channel_prefix_matrix (num_ranks, num_channels)：每个 channel 发往每张卡的 token 数前缀和
recv_rdma_rank_prefix_sum / recv_gbl_rank_prefix_sum：本节点、本卡要收多少
moe_recv_counter 与 moe_recv_expert_counter：总接收量、以及每个本地 expert 的接收量

前缀和的粒度是「channel × 目标」而不是「目标」，因为下面每个 channel 是独立流水的，它必须知道自己那一段数据在对端缓冲区里的确切偏移，否则多个 channel 会互相踩。

moe_recv_counter 这个值必须回到 host——PyTorch 要用它分配 recv_x。这就是那次躲不掉的 CPU 等待，也是常规 kernel 不能进 CUDA Graph 的唯一原因：

图里两个细节值得留意。一是 Notify tensor size ASAP——notify kernel 里计数一算完就立刻回写，不等其他工作结束，把 CPU 的等待窗口压到最短。二是右下那条 Reuse layout information：combine 完全不需要再 notify 一次，它调的是 cached_notify，把 dispatch 阶段算好的 layout 反着用。同样地，dispatch 传入 handle 参数时也会走 cached 路径——反向传播里 combine 的 backward 就是一次 dispatch，layout 与 forward 完全一致。

一个 channel 两个 SM，五种 warp 角色

DeepEP 把本卡的 token 切成 num_channels 段，每段由一个 channel 独立发送，而 num_channels = num_sms / 2。所以 set_num_sms(24) 的含义是 12 个并行的发送流水线；V3 论文里那句「20 个 SM 划成 10 个通信通道」就是 num_sms = 20。

每个 channel 的两个 SM 分工不同（sm_id % 2 决定），每个 SM 16 个 warp，角色在 kernel 入口一次性静态分配好：

1
2
3
4
5


enum class WarpRole { kRDMASender, kRDMASenderCoordinator,
                      kRDMAAndNVLForwarder, kForwarderCoordinator, kNVLReceivers };
// 每 SM 16 warp = kNumDispatchRDMASenderWarps(7) + 1 + LEGACY_NUM_MAX_NVL_PEERS(8)
const auto num_channels = num_sms / 2, channel_id = sm_id / 2;
const bool is_forwarder = sm_id % 2 == 0;

偶号 SM：8 个 RDMAAndNVLForwarder（一个 warp 盯一张机内卡）+ 1 个 ForwarderCoordinator
奇号 SM：7 个 RDMASender + 1 个 RDMASenderCoordinator + 8 个 NVLReceivers

数据在这些角色之间的流向，就是前面那两跳：

flowchart LR
    X["本卡 token 第 c 段"] --> S["RDMASender x7<br/>写 RDMA send_buffer"]
    S --> Q1(("rdma_channel<br/>环形队列"))
    Q1 -->|"nvshmem put"| F["RDMAAndNVLForwarder x8<br/>按机内目标卡分流"]
    F --> Q2(("nvl_channel<br/>IPC 环形队列"))
    Q2 --> R["NVLReceivers x8<br/>写 recv_x"]
    SC["RDMASenderCoordinator"] -.->|"聚合 head/tail"| S
    FC["ForwarderCoordinator"] -.->|"回推 rdma_channel_head"| F

两个队列的实现机制不同，这是理解代码的关键：rdma_channel 的收发缓冲区分别在两台机器上，头尾指针必须靠 NVSHMEM 显式同步；nvl_channel 的数据和指针都放在 CUDA IPC 共享内存里，机内所有卡直接读写，不需要额外同步。

Coordinator 这两个角色最容易被忽略，但它们解释了「为什么不是 8 个 warp 各自管好自己就行」。以 ForwarderCoordinator 为例：远端的 rdma_channel_head 只有在全部 8 个 forwarder 都确认转发完某段数据之后才能推进，否则会把还没转走的数据覆盖掉。这个「取 8 个 warp 的最小进度」的归约本身需要一个 warp 来做，让它顺带负责回写远端指针，比让 8 个 warp 互相 poll 便宜得多。RDMASenderCoordinator 同理，它还额外承担攒批的职责：只有当某个 channel 里待发的 token 超过 num_max_rdma_chunked_send_tokens 才真正发一次 RDMA write，避免小消息把消息率打满。

combine 基本是把这张图反过来跑：NVLSender → NVLAndRDMAForwarder → RDMAReceiver，只是多了两次 reduce。dispatch 时一个 token 从 (dst_rdma_rank, src_nvl_rank) 扇出到多个 (dst_rdma_rank, dst_nvl_rank)，combine 时这些副本要先在机内合并，再跨机合并——所以 forwarder 的 warp 数量明显多于 dispatch 侧（kNumCombineForwarderWarps = 24），因为它干的不只是拷贝。dispatch 阶段留下的 combined_nvl_head、combined_rdma_head 两张表就是给这两次 reduce 用的：它们记录每个 token 当初被扇到了哪些卡，reduce 才知道该等谁。

低时延 kernel：省掉一次同步，和一个不占 SM 的 hook

decode 阶段 batch 只有一百来个 token，通信量小到延迟完全由固定开销主导。这时候两跳转发不再是优势——多一跳就是多一次等待。所以低时延 kernel 干脆放弃 NVLink，即使目标卡就在本机也走 RDMA 直达；对应地它把所有 GPU 放进同一个 NVSHMEM team，而不再是「同号卡组成一个 team」。

传输层也换了。常规模式用 IBRC，请求由 CPU 上的代理线程发起；低时延模式启用 IBGDA，GPU SM 直接和 NIC 交互。初始化时那几个环境变量就是干这件事的：

1
2
3
4
5


os.environ['NVSHMEM_DISABLE_P2P'] = '1'
os.environ['NVSHMEM_IB_ENABLE_IBGDA'] = '1'
os.environ['NVSHMEM_IBGDA_NIC_HANDLER'] = 'gpu'
os.environ['NVSHMEM_IBGDA_NUM_RC_PER_PE'] = f'{num_qps_per_rank}'
os.environ['NVSHMEM_QP_DEPTH'] = '1024'

IBGDA 对小消息的收益接近腰斩延迟，原理和 IBRC 的对比见 NVSHMEM 那篇⁴。这里只强调一处 DeepEP 特有的约定：num_qps_per_rank 必须等于本地 expert 数，因为 kernel 里一个 warp group 对应一个 expert，让每个 expert 独占一条 QP 才能避免 doorbell 争抢。NVSHMEM_QP_DEPTH = 1024 也是同一个思路——把队列开到一定不会满，kernel 里就可以省掉检查 WQ 空槽的分支。

没有 notify_dispatch 的代价直接写在返回值的 shape 上：

packed_recv_x: [num_local_experts, num_max_dispatch_tokens_per_rank * num_ranks, hidden]
packed_recv_count: [num_local_experts]

第二维是最坏情况——假设所有 rank 的所有 token 都涌向同一个 expert。实际有效的行数由 packed_recv_count 给出，其余是垃圾。用显存换掉一次 CPU 同步，换来的就是 CUDA Graph 兼容（这对 decode 太重要了，几百微秒的 kernel launch 开销本来就和通信同量级）。scaling factor 那一维被存成 column-major，是为了后续 GEMM 能直接用 TMA 读。

recv hook：把「等数据」这件事从 SM 上摘下来

低时延 kernel 拆成 LOW_LATENCY_SEND_PHASE 和 LOW_LATENCY_RECV_PHASE 两段，return_recv_hook=True 时只跑发送段，接收段延迟到用户调 hook() 时才执行。

这个设计解决的是一个很具体的浪费。传统的通信-计算 overlap 靠两条 stream：通信 kernel 在 stream 1 上跑，计算 kernel 在 stream 0 上跑，两者抢 SM。但 RDMA 传输期间通信 kernel 其实什么也没干，它只是在自旋等数据到达——占着 SM 空转。

两种 overlap 方式对比：上半为传统双 stream 方案，dispatch/combine 各自占用 SM；下半为 recv hook 方案，RDMA 在后台传输，单条 stream 上只有计算 kernel，dispatch/combine 的 issue 与 receive 被拆开插在计算之间

下半张图里只有一条 stream，四个格子全是计算。Dispatch 0 issue 之后 RDMA 请求已经发出去了，SM 立刻转去算 Attention 1；等算完再调 Dispatch 0 的 hook 去收。整个传输过程「零 SM 占用」——严格说是零 SM 等待，发起请求那一下当然还是要 SM 的。代价是需要把模型前向手工切成两个 micro-batch 并交错编排，且四段（attention / dispatch / MoE / combine）的耗时不会自动对齐，得按 workload 调。

SM 数量是一种要分配的资源

前面反复出现 num_sms，值得把它单独拎出来说，因为这是 DeepEP 区别于通用集合通信库最本质的一点：它把「这次通信允许占用多少 SM」交给调用方。

V1 是 Buffer.set_num_sms(24) 这样一个静态设置，官方建议跑一遍测试拿 auto-tune 的结果。为什么需要调：给多了，通信更快但计算被挤；给少了，通信成为瓶颈。最优值同时取决于网络带宽、hidden size、top-k 和模型本身的计算强度——这些量在不同集群上都不一样，所以只能测。

DeepSeek-V3 训练时给的是 20 个 SM，这个数字要和 DualPipe 放在一起看：DualPipe 把前向与反向的计算/通信阶段交错编排，all-to-all 藏在 GEMM 的影子里，而 20/132 这个占用比就是「影子」的宽度。V3 论文在硬件建议那节把这件事说得更直白：SM 现在被迫干四种活——在 IB 与 NVLink 域之间转发、在 RDMA buffer 与输入输出 buffer 之间搬数据、给 combine 做 reduce、管理细粒度的显存布局——他们希望未来有硬件把这些卸载掉³。

为了少占 SM，V1 还用了一处 undefined-behavior 的 PTX：用只读指令 ld.global.nc.L1::no_allocate.L2::256B 去读 volatile 数据。.nc 走 non-coherent cache 本来不该用于读别人正在写的数据，但 Hopper 上 non-coherent cache 与 L1 统一，加上 .L1::no_allocate 保证 L1 里不留脏数据，实测正确且明显更快。这种事只有自己写 kernel 才敢做，也是不用 NCCL 换来的自由度之一。平台不同挂了的话，DISABLE_AGGRESSIVE_PTX_INSTRS=1 可以关掉。

V2：从 NVSHMEM 换到 NCCL Gin

V2 是一次彻底重构，几个变化都指向同一个方向：用更少的 SM 支撑更大的规模。

后端换了。 NVSHMEM 换成 NCCL 的 Gin backend，理由是后者 header-only、更轻，而且能复用已有的 NCCL communicator——不必再为 EP 单独 bootstrap 一套通信域。所有 kernel 改为 JIT 编译，安装时不再需要 nvcc。

两套 API 合成一套。 高吞吐与低时延统一到 ElasticBuffer，V1 那对 dispatch / low_latency_dispatch 不再是两个函数，而是同一个函数的两种模式：

hybrid mode：分层的 RDMA + NVLink 通信，也就是 V1 常规 kernel 那两跳，对多平面 / 多轨网络更友好
direct mode：直达，也就是 V1 低时延 kernel 的路径

术语也跟着换了：nvl_rank / rdma_rank 变成更中立的 scale-up rank / scale-out rank，因为「机内是不是 NVLink」这个假设已经不再成立（PCIe 环境、更大的 NVLink 域都在支持范围内）。

SM 数量从测出来变成算出来。 这是 V2 最漂亮的一处。get_theoretical_num_sms() 不再依赖 auto-tune，而是直接建模：先用前面那个组合数公式算出期望命中的 scale-out rank 数与总 rank 数，据此推出每个 token 平摊的 HBM 读写量和 RDMA / NVLink 流量，取比值找出被哪一侧卡住，最后用「瓶颈带宽 ÷ 单 SM 的 HBM 带宽」得到需要几个 SM：

1
2
3
4
5
6
7


if bounded_traffic > 0:
    num_sms = max(
        bounded_gbs / bounded_traffic * sm_read  / sm_read_gbs,   # 默认 200 GB/s
        bounded_gbs / bounded_traffic * sm_write / sm_write_gbs,  # 默认 50 GB/s
    )
num_sms = align(max(4, math.ceil(num_sms * 1.25)), 2)
num_sms = num_sms if self.prefer_overlap_with_compute else max(num_sms, 64)

留意 prefer_overlap_with_compute 这个开关，它把「要不要给计算让路」变成了显式选项：想 overlap 就用算出来的最小值，不想 overlap（比如纯通信 benchmark）就至少给 64 个。1.25 是安全余量，最低 4 个、且必须是偶数——偶数这个约束还是来自「一个 channel 两个 SM」。

效果是显著的：V3 那种配置下 SM 从 24 降到 4–6，性能不降反升；整体相比 V1 峰值提升最多 1.3 倍，SM 省最多 4 倍。V2 的实测数据：

架构	拓扑	dispatch 瓶颈带宽	combine 瓶颈带宽	SM 数
SM90	EP 8×2	90 GB/s (RDMA)	81 GB/s (RDMA)	12
SM90	EP 8×4	61 GB/s (RDMA)	61 GB/s (RDMA)	6
SM100	EP 8（纯 NVLink）	726 GB/s	740 GB/s	64
SM100	EP 8（纯 NVLink）	643 GB/s	675 GB/s	24

测试配置是 8K tokens/batch、hidden 7168、top-8、FP8 dispatch + BF16 combine。最后两行是同一配置下「最大性能」与「最少 SM」两个点，说明 SM 从 64 砍到 24 只损失 11% 带宽——这正是 prefer_overlap_with_compute 想让你做的那个交换。

规模上限抬到 EP2048。 配合的是把「0 SM」这件事推到更多场景：0 SM Engram（远端 KV cache 拉取，走 RDMA）、0 SM PP（流水线并行的 send/recv，走 RDMA）、0 SM CP（上下文并行，走 Copy Engine）。名字里的 DeepEP 已经从 Expert Parallel 变成了 DeepEveryParallel。

代价也要记住： buffer 占用比 V1 更大；低时延 EP 不再支持 0 SM RDMA；Engram / PP / CP 都还是 experimental。V1 文档移到了 docs/legacy.md，代码在 csrc/kernels/legacy/，短期内不会删。

用起来会踩到的地方

负载不均衡不是 DeepEP 能解决的。 kernel 打满带宽只保证「搬得快」，不保证「各卡搬得一样多」。热点 expert 所在的卡收到 15 行而邻居只收到 9 行，快的那张卡照样干等——这是 MoE 那篇里说的 straggler，得靠 EPLB 这类 expert 重排、或者负载均衡机制在算法侧解决。V2 roadmap 里那句「用 EP replay 处理不均衡以缩小中间 buffer」说明官方也在往这个方向走。

网络侧有几个必须配的开关。 IB 的 virtual lane 要给 EP 流量单独划一条，避免和别的流量互相干扰（V1 用 NVSHMEM_IB_SL，V2 用 sl_idx 参数或 EP_OVERRIDE_RDMA_SL）。adaptive routing 这条要按 kernel 分开看：V1 的常规 kernel 要求 AR 关闭，而代码里检测不到 AR 状态，配错了不会报错、只会表现为结果不对；支持 AR 的只有低时延 kernel，对它的建议是重载开、轻载关。V2 统一成了「一律建议开」。congestion control 两个版本都建议关掉，因为它伤峰值带宽。

buffer 要复用。 Buffer / ElasticBuffer 的构造是重操作（要交换 IPC handle、建连接），标准写法是在框架初始化时建一次，之后靠官方给的 size hint 函数判断够不够、不够才重建，两个版本的示例代码都是这个套路。低时延模式还有个额外限制：内部只有两块 buffer 轮转，所以同一时刻不能持有超过两个 low-latency kernel 的结果 tensor。

FP8 是非对称的。 dispatch 用 FP8（[num_tokens, hidden] 的 e4m3 加 [num_tokens, hidden // 128] 的 scale），combine 用 BF16。方向不同精度不同，是因为 combine 要做累加，量化误差会被放大；细节可参考 FP8 混合精度。

生态：分叉比主干还热闹

DeepEP 的 README 底部那一长串分支和 fork 本身就是一份「all-to-all 还能怎么优化」的清单，挑几个有代表性的：

Zero-copy（腾讯网络平台部）：去掉 PyTorch tensor 与通信 buffer 之间的拷贝，常规 kernel 的 SM 占用大幅下降
Hybrid-EP：用 TMA 指令进一步压 SM，支持更大 NVLink 域、PCIe 环境和 NVFP4
Normal-SMFree（蚂蚁网络平台部）：把 comm kernel 执行与 NIC 传输解耦，从 RDMA 路径上彻底摘掉 SM
LL-SBO / LL-Layered（蚂蚁）：前者让 down GEMM 与 combine send 重叠，后者用轨道优化的转发降低跨节点 LL 延迟
Mori-EP / uccl-ep：AMD ROCm 与异构 GPU/NIC（EFA、Broadcom）支持
DeepXTrace（蚂蚁）：定位慢 rank 的诊断工具——大 EP 下「哪张卡拖后腿」本身就是个难题

顺着看还有一批不走 DeepEP 路线的同类工作： COMET 做的是细粒度的计算-通信融合， Triton-Distributed 想把这类 kernel 写成可移植的 Triton， Ring-EP 与 Bumi-EP 则换了通信模式。

参考

NVSHMEM：初始化流程与核心机制——DeepEP V1 的通信底座，IBRC / IBGDA 对比
all-to-all——原语语义、两阶段动态 splits
MoE 与 DeepSeek-V3 ——为什么会有这两次 all-to-all，以及 node-limited routing 的来历
MoonEP
UltraEP

其他参考：

DeepEP: an efficient expert-parallel communication library, https://github.com/deepseek-ai/DeepEP ↩︎
Insights into DeepSeek-V3: Scaling Challenges and Reflections on Hardware for AI Architectures, ISCA 2025, https://arxiv.org/abs/2505.09343 ↩︎
DeepSeek-V3 Technical Report, https://arxiv.org/abs/2412.19437 ↩︎ ↩︎
NVIDIA OpenSHMEM Library (NVSHMEM) Documentation, https://docs.nvidia.com/nvshmem/api/index.html ↩︎