DeepEP 整体架构分析
本文档是对 DeepEP(DeepEveryParallel)通信库的独立架构分析,聚焦其作为 MoE Data Movement Runtime 的系统设计。
分析基于 DeepEP v2.1.0 源码,覆盖 Python API → C++ binding → JIT 编译 → CUDA kernel 全链路。
1. 项目定位:DeepEP 是什么
DeepEP 是一个面向 MoE(Mixture of Experts)的数据搬运运行时(Data Movement Runtime),核心定位为:
- 高性能 Expert Parallelism(EP)all-to-all 通信库:提供 dispatch(token 路由到对应 expert 所在 rank)与 combine(expert 输出 reduce 回原 rank)两个基本原语
- 低精度支持:原生 FP8 dispatch + BF16 combine
- 极低 SM 占用:通信 kernel 仅占用极少 SM(4-64 个),将计算资源留给 GEMM
- 实验性扩展:PP send/recv、Engram(RDMA 远程内存访问)、CP、AGRS
DeepEP 的设计哲学是通信 kernel 不应与计算争抢 SM。这与传统集合通信库(NCCL)的设计目标有本质区别——NCCL 追求通用性与全吞吐,DeepEP 追求"刚好够用"的通信能力 + 最小 SM 开销。
flowchart LR
subgraph 目标["DeepEP 设计目标"]
A["高吞吐 all-to-all<br/>逼近硬件带宽极限"]
B["低延迟<br/>RDMA 直连"]
C["极低 SM 占用<br/>4-64 SMs for comm"]
D["FP8/BF16 混合精度"]
end
subgraph 场景["应用场景"]
E["MoE Training"]
F["MoE Inference Prefill"]
G["MoE Inference Decode"]
H["实验: PP/CP/Engram"]
end
A --> E
B --> G
C --> E
D --> F
2. 架构分层总览
DeepEP 采用经典的四层架构,从上到下依次为:
flowchart TB
subgraph L1["Layer 1: Python API"]
P1["ElasticBuffer / Buffer"]
P2["EPHandle / EventOverlap"]
P3["Utils: envs, event, math, comm"]
end
subgraph L2["Layer 2: C++ Binding (pybind11)"]
B1["python_api.cpp"]
B2["jit::register_apis"]
B3["elastic::register_apis"]
B4["legacy::register_apis"]
end
subgraph L3["Layer 3: JIT Compilation"]
J1["Compiler / NVCCCompiler"]
J2["KernelRuntimeCache"]
J3["KernelRuntime (cubin load)"]
J4["IncludeParser / DeviceRuntime"]
end
subgraph L4["Layer 4: CUDA Kernels + Backend"]
K1["elastic/dispatch.hpp · combine.hpp"]
K2["elastic/engram.hpp · pp_send_recv.hpp"]
K3["legacy/internode.cu · intranode.cu"]
K4["NCCL Gin Backend / NVSHMEM Backend"]
end
L1 --> L2 --> L3 --> L4
2.1 Layer 1:Python API 层
入口文件 deep_ep/__init__.py 在 import 时完成三件关键事情:
- NCCL 版本校验(
check_nccl_so):读取/proc/self/maps检查运行时加载的libnccl.so是否唯一,并与链接时的 NCCL 二进制对比,防止 PyTorch 自带 NCCL 与系统 NCCL 冲突 - JIT 运行时初始化(
init_jit):调用_C.init_jit(library_root, cuda_home, nccl_root)设置编译器路径 - 导出符号:
ElasticBuffer、Buffer、EPHandle、EventOverlap、Config、topk_idx_t
# __init__.py 核心初始化流程
check_nccl_so() # 1. 确保 NCCL 版本一致
init_jit() # 2. 初始化 JIT 编译器路径
from .buffers.legacy import Buffer # V1
from .buffers.elastic import ElasticBuffer, EPHandle # V2
from .utils.event import EventOverlap
2.2 Layer 2:C++ Binding 层
csrc/python_api.cpp 是唯一的 pybind11 模块入口,注册四个类别的 API:
| 注册函数 | 命名空间 | 功能 |
|---|---|---|
is_sm90_compiled |
全局 | 是否编译 SM90 特性 |
register_apis(m) |
jit |
init_jit |
register_apis(m) |
elastic |
ElasticBuffer 全套方法 |
register_apis(m) |
legacy |
Buffer 全套方法 |
2.3 Layer 3:JIT 编译层
详见第 6 节。
2.4 Layer 4:CUDA Kernel + 后端层
详见第 5、7 节。
3. V1 vs V2 架构对比(Legacy vs Elastic)
DeepEP 存在两套完整实现,代表两种截然不同的设计哲学:
flowchart TB
subgraph V1["V1: Legacy (NVSHMEM-based)"]
V1P["Buffer (Python)"]
V1C["legacy::Buffer (C++)"]
V1K["legacy/internode.cu<br/>legacy/intranode.cu"]
V1B["NVSHMEM Backend<br/>+ IBGDA"]
V1P --> V1C --> V1K --> V1B
end
subgraph V2["V2: Elastic (NCCL Gin-based)"]
V2P["ElasticBuffer (Python)"]
V2C["elastic::ElasticBuffer (C++)"]
V2K["elastic/dispatch.hpp<br/>elastic/combine.hpp"]
V2B["NCCL Gin Backend<br/>+ Symmetric Memory"]
V2P --> V2C --> V2K --> V2B
end
3.1 核心设计差异
| 维度 | V1 (Legacy) | V2 (Elastic) |
|---|---|---|
| 后端 | NVSHMEM + IBGDA | NCCL Gin (Header-only) |
| 内存模型 | 显式 NVLink buffer + RDMA buffer | NCCL Symmetric Memory (统一 VA) |
| Buffer 大小 | 较小 | 较大(README 明示) |
| SM 占用 | 20-24 SMs (normal) | 4-64 SMs (分析式计算) |
| QP 数量 | 手动配置 | 分析式自动计算 (17 / 65 / 129) |
| Scale 能力 | EP ≤ 160 | EP 可达 2048 |
| Config 调优 | 查表法 (config_map) | 无需 auto-tuning |
| 0 SM RDMA EP | 支持 (low_latency mode) | 不再支持 |
| 通信模式 | 显式 internode/intranode 分支 | hybrid/direct 统一 |
3.2 V1 Legacy 架构细节
legacy::Buffer 的核心设计:
- 双缓冲结构:NVLink buffer(节点内)+ RDMA buffer(节点间),物理分离
- NVSHMEM team 管理:通过
NVSHMEM_IBGDA_NUM_RC_PER_PE配置 QP 数,NVSHMEM_QP_DEPTH控制 WQ 深度 - IPC handle 交换:
sync()阶段交换 device_id、IPC handle、NVSHMEM unique_id - Low-latency mode:基于 IBGDA 的 RDMA 直连,双 buffer ping-pong(只能同时持有 2 个结果 tensor)
- Config 查表:
get_dispatch_config(num_ranks)和get_combine_config(num_ranks)是硬编码映射表
// V1 的 Config 查表是硬编码的
config_map = {
2: Config(Buffer.num_sms, 24, 256, 6, 128),
4: Config(Buffer.num_sms, 6, 256, 6, 128),
8: Config(Buffer.num_sms, 6, 256, 6, 128),
16: Config(Buffer.num_sms, 36, 288, 20, 128),
// ... 一直到 160
}
3.3 V2 Elastic 架构细节
elastic::ElasticBuffer 的核心设计(详见第 4 节):
- 统一 Symmetric Memory:通过 NCCL Gin 的
ncclWindow_t实现跨 rank 对称内存 - 逻辑/物理域分离:物理域(RDMA ranks × NVLink ranks)→ 逻辑域(scaleout × scaleup)
- 分析式 SM/QP 计算:基于带宽建模自动推导最优 SM 数
- Hybrid 模式:scaleout 走 RDMA,scaleup 走 NVLink,分层 reduce
4. 核心抽象
4.1 ElasticBuffer(V2 核心)
ElasticBuffer 是 V2 的统一接口,封装了所有 EP 操作:
classDiagram
class ElasticBuffer {
+group: ProcessGroup
+rank_idx, num_ranks: int
+num_bytes, num_cpu_buffer_bytes: int64
+allow_hybrid_mode: bool
+allow_multiple_reduction: bool
+prefer_overlap_with_compute: bool
+num_allocated_qps: int
+runtime: C++ ElasticBuffer
+nccl_comm_handle: NCCLCommHandle
+num_scaleout/num_scaleup_ranks: int
+num_rdma/num_nvl_ranks: int
+__init__(group, num_max_tokens_per_rank, hidden, ...)
+destroy()
+dispatch(x, topk_idx, ...) → (recv_x, handle, event)
+combine(x, handle, ...) → (combined_x, event)
+barrier()
+engram_write(storage) [实验]
+engram_fetch(indices) → hook [实验]
+pp_send/pp_recv [实验]
+all_gather [实验]
+get_theoretical_num_sms() → int
+get_buffer_size_hint() → int
}
class NCCLSymmetricMemoryContext {
+rank_idx, num_ranks: int
+num_scaleout/num_scaleup_ranks: int
+num_rdma/num_nvl_ranks: int
+is_scaleup_nvlink: bool
+comm: ncclComm_t
+dev_comm: ncclDevComm_t
+window: ncclWindow_t
+mapped_window_ptr: void*
+num_allocated_qps: int
}
ElasticBuffer --> NCCLSymmetricMemoryContext : nccl_context
ElasticBuffer --> EPHandle : dispatch returns
ElasticBuffer --> EventOverlap : returns
内存布局
ElasticBuffer 的内存布局是理解其设计的核心:
[Workspace (2MB-aligned)] [GPU Buffer] [CPU Buffer (可选)]
↑ mapped_window_ptr
- Workspace:存放 barrier signals、rank/expert 接收计数、channel metadata、PP 信号、AGRS 信号
- GPU Buffer:dispatch/combine 的 send/recv 区域
- CPU Buffer:Engram 存储区(可选,位于 GPU buffer 之后)
SymmetricMemory 的三种实现:
| 类型 | 类 | 用途 |
|---|---|---|
| 纯 GPU | GPUSymmetricMemory |
ncclMemAlloc,默认模式 |
| GPU + CPU | ElasticSymmetricMemory |
CUDA Driver API cuMemCreate,连续 VA |
| GPU + 多 CPU (Hybrid) | HybridElasticSymmetricMemory |
各 rank CPU segment 通过 POSIX FD import |
4.2 EPHandle(通信句柄)
EPHandle 是 dispatch 返回的路由元数据,combine 消费它来完成反向 reduce:
class EPHandle:
do_expand: bool # 是否使用 expand 布局
num_experts: int
expert_alignment: int # expert token 对齐
num_max_tokens_per_rank: int
num_sms: int
topk_idx: Tensor # [num_tokens, num_topk],克隆的 expert 索引
psum_num_recv_tokens_per_scaleup_rank: Tensor # scaleup rank 前缀和
psum_num_recv_tokens_per_expert: Tensor # expert 前缀和
num_unaligned_recv_tokens_perpert: Tensor # expand 模式下未对齐计数
recv_src_metadata: Tensor # [num_recv_tokens, topk+2],源 token 索引 + slot
dst_buffer_slot_idx: Tensor # 目标 buffer slot
token_metadata_at_forward # hybrid 模式元数据
channel_linked_list # hybrid 模式链表
num_recv_tokens: int
num_expanded_tokens: int
EPHandle 支持 cached mode:当 gating decision 不变时,可复用 handle 跳过 layout 重算与 CPU sync,这对 inference decoding 场景至关重要。
43. EventOverlap(事件重叠)
EventOverlap 封装了 CUDA event,提供通信-计算重叠的 Python 友好接口:
# 典型用法:先做不依赖 recv_x 的计算,再等待通信完成
recv_x, ..., handle, event = buffer.dispatch(...)
with event: # __exit__ 时 current_stream_wait()
do_other_compute() # 与通信重叠执行
# 退出 with 后,recv_x 安全可用
recv_x = forward_experts(recv_x, ...)
内部机制:
- event.current_stream_wait():compute stream 等待 comm stream 上的 event
- record_stream():确保 tensor 在 compute/comm 两个 stream 上的访问安全
- EP_AVOID_RECORD_STREAM 环境变量可切换为 manual tensor tracking(CUDA graph 兼容)
- hook_after_wait:支持 deterministic dispatch 的 post-wait sort
5. Kernel 组织
5.1 目录结构
csrc/kernels/
├── elastic/ # V2 kernels
│ ├── api.hpp # kernel 总入口 (launch 函数)
│ ├── dispatch.hpp # dispatch + dispatch_copy_epilogue
│ ├── combine.hpp # combine + combine_reduce_epilogue
│ ├── engram.hpp # Engram fetch (实验)
│ ├── pp_send_recv.hpp # PP send/recv (实验)
│ ├── barrier.hpp # GPU barrier
│ └── utils.hpp # kernel utilities
├── legacy/ # V1 kernels
│ ├── api.cuh # V1 kernel 入口
│ ├── internode.cu # 节点间 (RDMA + NVLink)
│ ├── intranode.cu # 节点内 (NVLink)
│ ├── buffer.cuh # V1 buffer 管理
│ ├── compiled.cuh # V1 编译配置
│ ├── ibgda_device.cuh # IBGDA low-latency
│ ├── launch.cuh # V1 launch 逻辑
│ ├── layout.cu # V1 layout 计算
│ └── utils.cuh # V1 utilities
└── backend/ # 后端抽象
├── api.cuh # nvshmem / nccl / cuda_driver API
├── symmetric.hpp # SymmetricMemory 分配器
├── cuda_driver.cu # CUDA Driver API 封装
├── nccl.cu # NCCL 初始化/销毁
└── nvshmem.cu # NVSHMEM 初始化/销毁
5.2 Elastic Kernel 结构
每个 Elastic kernel 由两个阶段组成(以 dispatch 为例):
flowchart LR
subgraph Dispatch["dispatch (2-phase)"]
D1["Phase 1: dispatch_impl<br/>notify warps + dispatch warps<br/>写对称内存 + 计数"]
D2["Phase 2: dispatch_copy_epilogue<br/>从对称内存拷贝到用户 tensor<br/>full SMs"]
D1 -->|"named barrier"| D2
end
subgraph Combine["combine (2-phase)"]
C1["Phase 1: combine_impl<br/>从本地 push 数据到远端"]
C2["Phase 2: combine_reduce_epilogue<br/>reduce + bias + 写回"]
C1 -->|"named barrier"| C2
end
Dispatch 的 warp 角色分配
| Warp 角色 | 数量 | 职责 |
|---|---|---|
| Notify warps | 4 (kNumNotifyWarps) | 原子累加 rank/expert 接收计数,通知对端 |
| Dispatch warps | 自动计算 | 通过 Gin QP 执行 RDMA/NVLink 数据传输 |
| Scaleout warps | num_channels_per_sm | hybrid 模式:跨节点 forwarding |
| Forward warps | num_channels_per_sm | hybrid 模式:节点内转发 |
Channel 与 QP 的映射
// comm::get_qp_mode: 决定每个 SM 的每个 channel 使用哪个 QP
if (kNumSMs <= kNumAvailableQPs) {
// SM 少:每个 SM 独占 QP,channel 间 round-robin
// e.g., 3 SMs, 10 QPs: SM0→{0,3,6,9}, SM1→{1,4,7}, SM2→{2,5,8}
} else {
// SM 多:所有 SM 共享所有 QP
qp_idx = kQPStartIdx + (global_channel_idx % kNumAvailableQPs)
}
5.3 Hybrid vs Direct 模式
Elastic 内核根据 num_scaleout_ranks 自动选择:
| 条件 | 模式 | 内核 |
|---|---|---|
num_scaleout_ranks == 1 |
Direct | dispatch_impl / combine_impl |
num_scaleout_ranks > 1 |
Hybrid | hybrid_dispatch_impl / hybrid_combine_impl |
Hybrid 模式的分层通信:
1. Scaleup 层(NVLink):节点内各 GPU 先做局部 reduce
2. Scaleout 层(RDMA):跨节点 forwarding,每个 channel 负责一个 scaleout peer 的数据转发
flowchart TB
subgraph HybridDispatch["Hybrid Dispatch 数据流"]
T0["Token on Rank 0"]
T1["Notify Warps: 计算目标 slot"]
T2["Scaleout Warps: RDMA 写到<br/>远端 scaleout rank 的 buffer"]
T3["Forward Warps: NVLink 转发<br/>到目标 scaleup rank"]
T4["最终到达目标 expert 所在 rank"]
end
T0 --> T1 --> T2 --> T3 --> T4
6. JIT 编译管线
DeepEP 的 JIT 系统是其"安装时无需 CUDA 编译"的关键。整个管线设计精巧:
flowchart TB
subgraph Init["初始化阶段"]
I1["__init__.py: init_jit()"]
I2["jit::init(library_root, cuda_home, nccl_root)"]
I3["Compiler::prepare_init()"]
I4["IncludeParser::prepare_init()"]
I5["DeviceRuntime::prepare_init()"]
I1 --> I2 --> I3 & I4 & I5
end
subgraph Build["内核构建阶段"]
B1["generate_impl(args) → 生成 kernel.cu 源码"]
B2["计算 kernel_signature = name$$sig$$flags$$code"]
B3["计算 cache 路径: ~/.deep_ep/cache/kernel.{name}.{hash}/"]
B4{"KernelRuntimeCache 命中?"}
B5["是: 返回缓存的 KernelRuntime"]
B6["否: NVCCCompiler::compile() → kernel.cubin"]
B7["cuobjdump 反汇编 → kernel.sass (可选)"]
B8["原子 rename tmp_dir → cache_dir"]
B9["KernelRuntime 加载 cubin"]
B1 --> B2 --> B3 --> B4
B4 -->|Yes| B5
B4 -->|No| B6 --> B7 --> B8 --> B9
end
subgraph Launch["内核启动阶段"]
L1["construct_launch_config()"]
L2["launch_kernel(kernel, config, args...)"]
L1 --> L2
end
Init --> Build --> Launch
6.1 源码生成(Code Generation)
每个 kernel 的源码是运行时生成的,通过模板参数实现编译期优化:
// dispatch.hpp 生成的 kernel.cu 示例
#include <deep_ep/impls/dispatch.cuh>
using namespace deep_ep::elastic;
static void __instantiate_kernel() {
auto ptr = reinterpret_cast<void*>(&dispatch_impl<
true, // is_scaleup_nvlink
false, // do_cpu_sync
false, // reuse_slot_indices
12, // num_sms
4, // num_notify_warps
28, // num_dispatch_warps
16, // num_scaleup_ranks
14336, // num_hidden_bytes (7168*2)
0, // num_sf_packs
2048, // num_max_tokens_per_rank
256, // num_experts
8, // num_topk
128, // expert_alignment
65, // num_qps
cycles // num_timeout_cycles
>);
}
关键设计:所有可变参数都变成模板参数,NVCC 可做激进优化(循环展开、分支消除、常量传播)。
6.2 缓存机制
缓存签名由四部分构成:
kernel_signature = "{name}$${signature}$${flags}$${code}"
name:kernel 名称(dispatch, combine, ...)signature:编译器版本(NVCC12.9)flags:编译标志code:生成的源码
缓存目录结构:
~/.deep_ep/
├── cache/
│ └── kernel.dispatch.{hex_digest}/
│ ├── kernel.cu # 生成的源码
│ ├── kernel.cubin # 编译产物
│ └── kernel.sass # 可选的反汇编
└── tmp/ # 临时目录(原子 rename 用)
分布式文件系统安全:
- 先编译到 tmp/{uuid}/,再原子 rename 到 cache 目录
- fsync_dir() 确保 data + directory entry 可见
- 多 rank 并发时,若 rename 失败则复用已有目录
6.3 Include Hash 机制
IncludeParser 解析生成代码中的所有 #include <deep_ep/*>,递归计算头文件内容的哈希值。当任何头文件改变时,哈希变化 → cache miss → 重新编译。这保证了头文件修改后的正确重建。
7. 后端抽象层
7.1 三种后端
flowchart LR
subgraph Backends["DeepEP 后端"]
direction TB
N1["NCCL Gin Backend<br/>(V2 主力)"]
N2["NVSHMEM Backend<br/>(V1 专用)"]
N3["CUDA Driver API<br/>(辅助)"]
end
N1 --- N1D["ncclMemAlloc / ncclCommWindowRegister<br/>ncclGin (QP 抽象)<br/>Header-only, 复用 NCCL communicator"]
N2 --- N2D["nvshmem_init / ibgda<br/>NVSHMEM team / QP<br/>独立初始化"]
N3 --- N3D["cuMemCreate / cuMemMap<br/>cuMemcpyBatchAsync<br/>batched_write_and_wait"]
7.2 NCCL Gin 后端(核心创新)
NCCL Gin 是 NCCL 内部的一个轻量级 GDAKI 接口,DeepEP V2 的全部通信都构建其上:
| 特性 | 说明 |
|---|---|
| Header-only | 只需 nccl.h + nccl_device.h,无额外链接 |
| Communicator 复用 | 直接复用 PyTorch 的 NCCL communicator |
| QP 抽象 | ncclGin handle + ncclGinRequest_t |
| Resource Sharing | CTA 级 / GPU 级 QP 共享模式 |
| Symmetric Memory | ncclWindow_t 提供跨 rank 对称内存视图 |
Gin 的使用模式:
// 1. 获取 Gin handle(per-warp)
const auto [qp_idx, sharing_mode] = comm::get_qp_mode<...>(sm_idx, channel_idx, is_notify);
const auto gin = handle::NCCLGin(nccl_dev_comm, nccl_window, qp_idx, sharing_mode);
// 2. 发起 RDMA 请求
ncclGinRequest_t req;
gin.get(dst_ptr, src_ptr, size, dst_rank, req); // RDMA get
// 3. 等待完成
gin.wait(req);
7.3 Symmetric Memory 分配
symmetric.hpp 定义了三种 SymmetricMemory 实现:
| 场景 | 类 | 分配方式 |
|---|---|---|
| 纯 GPU | GPUSymmetricMemory |
ncclMemAlloc |
| GPU + CPU (单节点) | ElasticSymmetricMemory |
cuMemCreate + cuMemMap |
| GPU + CPU (Hybrid) | HybridElasticSymmetricMemory |
POSIX FD import + cuMemImportFromShareableHandle |
Hybrid 模式的 CPU segment 跨 rank 共享通过 POSIX file descriptor 实现:
1. 每个 rank 创建 NUMA-local CPU allocation → export 为 FD
2. dist.all_gather_object 交换 (pid, fd)
3. 远端 rank 通过 pidfd_open + pidfd_getfd 导入 FD
4. cuMemImportFromShareableHandle 导入 CUDA 并 map 到本地 VA
7.4 CUDA Driver API
cuda_driver 命名空间提供两个关键原语:
- batched_write:批量 CUDA 信号写入
- batched_wait:批量 CUDA 信号等待(用于 AGRS session 同步)
这些使用 CUDA Driver 的 virtual address management 能力,直接操作 symmetric memory 上的信号量。
8. 性能特征
8.1 README 性能数据
基于 V3 配置(8K tokens/batch, 7168 hidden, top-8 experts, FP8 dispatch, BF16 combine):
| Arch | NIC | Topo | Dispatch BW | Combine BW | #SMs |
|---|---|---|---|---|---|
| SM90 | CX7 | EP 8×2 | 90 GB/s (RDMA) | 81 GB/s (RDMA) | 12 |
| SM90 | CX7 | EP 8×4 | 61 GB/s (RDMA) | 61 GB/s (RDMA) | 6 |
| SM100 | CX7 | EP 8×2 | 90 GB/s (RDMA) | 91 GB/s (RDMA) | 12 |
| SM100 | N/A | EP 8 | 726 GB/s (NVLink) | 740 GB/s (NVLink) | 64 (Max) |
| SM100 | N/A | EP 8 | 643 GB/s (NVLink) | 675 GB/s (NVLink) | 24 (Min) |
8.2 V2 vs V1 对比
- 峰值性能提升:最高 1.3x
- SM 节省:最高 4x(V3 训练从 24 SM 降到 4-6 SM)
8.3 性能分析
这些数字揭示了 DeepEP 的设计权衡:
- RDMA 场景 90 GB/s:接近 CX7 100Gb/s 的理论带宽(~100 GB/s),说明 kernel 效率极高
- NVLink 726 GB/s:接近 H100 NVLink 理论带宽(~900 GB/s),效率 ~80%
- SM 数与带宽的 trade-off:
- 64 SM → 726 GB/s(峰值带宽)
- 24 SM → 643 GB/s(~89% 峰值,节省 62.5% SM) - Hybrid 模式 61-90 GB/s:受限于 RDMA 带宽,但 SM 占用仅 6-12
9. SM 占用哲学
DeepEP 的核心设计哲学是"通信不应与计算争抢 SM"。这体现在多个层面:
9.1 分析式 SM 计算
get_theoretical_num_sms() 基于带宽建模:
# 核心公式(简化)
# 假设 balanced gate,计算 expected top-k 数量
num_expected_topk = num_groups * (1 - C(num_experts - num_experts//num_groups, num_topk) / C(num_experts, num_topk))
# 计算各层流量
sm_read += 1 / num_expected_topk
sm_write += ...
rdma_traffic += ...
nvlink_traffic += ...
# SM 数 = ceil(max(sm_read/sm_read_bw, sm_write/sm_write_bw, rdma_traffic/rdma_bw, nvlink_traffic/nvlink_bw))
9.2 SM 分配的运行时决策
// dispatch.hpp: 根据共享内存消耗计算每个 SM 的 channel 数
num_channels_per_sm = min(
(smem_bytes - notify_smem_bytes) / dispatch_token_layout_size,
32 - kNumNotifyWarps
);
num_channels_per_sm = min(
smem_bytes / combine_token_layout_size,
num_channels_per_sm
);
num_channels_per_sm = min(num_channels_per_sm / 2, kNumMaxChannelsPerSM);
9.3 prefer_overlap_with_compute 标志
当该标志开启时,kernel 倾向于使用更少的 SM 和更多的 channel(最多 4 channels/SM),通过时间换空间的方式让出 SM 给计算。
9.4 与 NCCL 的对比
| 维度 | NCCL | DeepEP |
|---|---|---|
| SM 策略 | 尽可能多占用 SM | 分析式最小化 SM |
| 带宽目标 | 饱和硬件带宽 | 饱和硬件带宽(同等) |
| 计算影响 | 可能影响 GEMM | 几乎不影响 GEMM |
| 适用场景 | 通用集合通信 | MoE all-to-all 专用 |
10. 实验性特性
10.1 Engram(0 SM RDMA 远程内存访问)
Engram 提供远程 CPU 内存的 RDMA 读取能力,用于 KV cache 等场景:
flowchart LR
E1["engram_write(storage)<br/>写入 CPU buffer"]
E2["engram_fetch(indices)<br/>发起 RDMA get"]
E3["返回 hook 闭包<br/>延迟等待"]
E4["hook() → 阻塞等待<br/>RDMA 完成"]
E5["返回 (data, sf)"]
E1 --> E2 --> E3 --> E4 --> E5
- 0 SM 哲学:RDMA get 由 NIC 直接完成,不占用 SM
- CPU buffer:通过
HybridElasticSymmetricMemory分配 NUMA-local CPU 内存 - FP8 支持:支持 FP8 storage + globally-replicated scaling factors
10.2 PP Send/Recv(0 SM Pipeline Parallelism)
基于 NVLink 的 pipeline parallelism 原语:
buffer.pp_set_config(num_max_tensor_bytes, num_max_inflight_tensors)
buffer.pp_send(t, dst_rank_idx, num_sms=0) # 0 = 全部 SM
buffer.pp_recv(t, src_rank_idx, num_sms=0)
- 仅支持 ring 上的 prev/next rank
- 使用
num_max_inflight_tensors * 2 * 2的 buffer 空间(send/recv × prev/next)
10.3 AGRS(All-Gather Reduce-Scatter)
基于 NVLink symmetric memory 的 all-gather 实现:
buffer.agrs_set_config(num_max_session_bytes, num_max_all_gathers)
buffer.create_agrs_session()
t = buffer.agrs_get_inplace_tensor(shape, dtype) # zero-copy
gathered, handle = buffer.all_gather(t)
handle() # wait for data arrival
buffer.destroy_agrs_session()
实现使用 cudaMemcpyBatchAsync(CUDA 13+)+ signal-based session 同步。
10.4 CP(Context Parallelism)
README 提到 "0 SM CP (with Copy Engine)",使用 copy engine 而非 SM 完成数据传输。
11. 环境变量体系
DeepEP 提供了丰富的环境变量用于调试与调优:
| 类别 | 关键变量 | 作用 |
|---|---|---|
| 通用 | EP_BUFFER_DEBUG |
buffer 初始化调试 |
| 通用 | EP_SUPPRESS_NCCL_CHECK |
跳过 NCCL 版本校验 |
| 通用 | EP_AVOID_RECORD_STREAM |
避免 CUDA graph 不兼容的 record_stream |
| 网络 | EP_NIC_NAME |
NIC 名称(默认 mlx5_0) |
| 网络 | EP_OVERRIDE_RDMA_SL |
覆盖 RDMA 服务级别 |
| 网络 | EP_DISABLE_GIN |
禁用 NCCL Gin 后端 |
| JIT | EP_JIT_CACHE_DIR |
编译缓存目录(默认 ~/.deep_ep) |
| JIT | EP_JIT_CPP_STANDARD |
C++ 标准(默认 20) |
| JIT | EP_JIT_NVCC_COMPILER |
NVCC 路径 |
| JIT | EP_JIT_DUMP_PTX/SASS |
导出编译产物 |
| JIT | EP_JIT_PTXAS_CHECK |
检查 local memory 使用 |
| 构建 | DISABLE_SM90_FEATURES |
禁用 SM90 特性 |
| 构建 | EP_NCCL_ROOT_DIR |
NCCL 安装路径 |
持久化环境变量(构建时烘焙到包中):EP_JIT_CACHE_DIR、EP_JIT_PRINT_COMPILER_COMMAND、EP_NUM_TOPK_IDX_BITS、EP_NCCL_ROOT_DIR。
12. 数据流全景
以 V2 的 dispatch + combine 为例,完整数据流如下:
flowchart TB
subgraph DispatchFlow["Dispatch 数据流"]
direction TB
D0["输入: x[num_tokens, hidden], topk_idx, topk_weights"]
D1["notify warps: 解析 topk_idx<br/>原子累加 rank/expert 计数<br/>写入 host_workspace (mapped)"]
D2["dispatch warps: 打包 token data + sf<br/>通过 Gin QP 发起 RDMA/NVLink get"]
D3["数据到达远端对称内存 buffer"]
D4["CPU sync 或 cached mode<br/>确定 num_recv_tokens"]
D5["dispatch_copy_epilogue: 从 buffer 拷贝<br/>到 recv_x, recv_topk_idx, recv_topk_weights"]
D6["输出: recv_x, EPHandle, EventOverlap"]
end
subgraph CombineFlow["Combine 数据流"]
direction TB
C0["输入: expert 计算结果 + EPHandle"]
C1["combine warps: 读取 handle 中的路由元数据<br/>通过 Gin push 数据到远端 buffer"]
C2["combine_reduce_epilogue: 从 buffer reduce<br/>按 topk_weights 加权 + bias"]
C3["输出: combined_x, EventOverlap"]
end
D6 -->|"EPHandle 包含路由元数据"| C0
DispatchFlow --> CombineFlow
13. 关键设计洞察
13.1 为什么用 NCCL Gin 而非原生 RDMA
- Communicator 复用:直接复用 PyTorch 已创建的 NCCL communicator,无需独立初始化
- Header-only:无需链接额外库,部署简单
- QP 抽象:Gin 提供比 raw GDAKI 更高层的接口,自动处理 QP 共享与资源管理
- 与 NCCL 协同:复用 NCCL 的 transport 层、congestion control、adaptive routing
13.2 为什么 V2 buffer 更大
V2 的 buffer 需要同时容纳:
- Scaleup recv buffer(节点内)
- Scaleout send buffer(跨节点发送)
- Scaleout recv buffer(跨节点接收,含 channel 元数据 padding)
Hybrid 模式下的分层通信需要多份中间 buffer,但换来了可扩展性(EP2048)和SM 节省。
13.3 Cached Mode 的意义
Inference decoding 场景中,gating decision 可能在多轮迭代间不变。Cached mode 允许:
- 跳过 layout 重算(topk_idx 解析、prefix sum)
- 跳过 CPU sync(直接复用 handle 中的计数)
- 减少 CPU-GPU 交互延迟
这对低延迟 decoding 至关重要。
13.4 Deterministic Mode
DeepEP 支持 deterministic dispatch:
- 对 recv_src_metadata[:, 0](源 token 全局索引)排序
- 根据排序结果重排 recv_x, recv_sf, recv_topk_weights, recv_topk_idx
- Expand 模式:按 expert 分组排序,更新 slot 指针
14. 总结
DeepEP 是一个设计精良的 MoE 通信运行时,其核心优势在于:
- 极致的 SM 效率:4-64 SM 即可饱和硬件带宽
- 统一的架构:V2 的 ElasticBuffer 统一了 high-throughput / low-latency / hybrid / direct 多种模式
- 运行时 JIT:安装时无需编译,自动缓存编译产物
- 分析式调参:SM 数、QP 数均自动计算,无需 auto-tuning
- NCCL Gin 创新:利用 NCCL 内部接口实现轻量级高性能通信
其代价是 buffer 占用增大、0 SM RDMA low-latency 模式不再支持,但这些 trade-off 在大规模 MoE 训练场景下是合理的。
附录:关键文件索引
| 文件 | 职责 |
|---|---|
deep_ep/__init__.py |
包入口,NCCL 检查,JIT 初始化 |
csrc/python_api.cpp |
pybind11 模块注册 |
deep_ep/buffers/elastic.py |
V2 ElasticBuffer Python 接口 |
deep_ep/buffers/legacy.py |
V1 Buffer Python 接口 |
csrc/elastic/buffer.hpp |
V2 C++ 核心实现 |
csrc/legacy/buffer.hpp |
V1 C++ 核心实现 |
csrc/jit/compiler.hpp |
JIT 编译器(NVCC) |
csrc/jit/cache.hpp |
KernelRuntimeCache |
csrc/jit/kernel_runtime.hpp |
CUBIN 加载与符号解析 |
csrc/jit/launch_runtime.hpp |
Launch 框架(CRTP) |
csrc/jit/include_parser.hpp |
头文件依赖哈希 |
csrc/jit/device_runtime.hpp |
GPU 属性查询 |
csrc/kernels/elastic/dispatch.hpp |
V2 dispatch 内核 |
csrc/kernels/elastic/combine.hpp |
V2 combine 内核 |
csrc/kernels/elastic/engram.hpp |
Engram fetch 内核 |
csrc/kernels/elastic/pp_send_recv.hpp |
PP send/recv 内核 |
csrc/kernels/backend/symmetric.hpp |
SymmetricMemory 分配器 |
csrc/kernels/backend/api.cuh |
后端 API 抽象 |
deep_ep/include/deep_ep/common/layout.cuh |
Workspace/Buffer 布局 |
deep_ep/include/deep_ep/common/comm.cuh |
Gin 通信原语 |
deep_ep/include/deep_ep/impls/dispatch.cuh |
dispatch_impl 定义 |
deep_ep/utils/event.py |
EventOverlap 封装 |
deep_ep/utils/envs.py |
环境检测与分布式初始化 |
tests/elastic/test_ep.py |
EP 正确性与性能测试 |