衡宇 发自 凹非寺
量子位 | 公众号 QbitAI
好消息如约而至，DeepSeek开源周第二弹来了！
DeepEP
， 第一个用于MoE模型训练和推理的开源EP通信库
（expert parallelism，专家并行）
。
它提供高吞吐量和低延迟的all-to-all GPU内核，也称为MoE dispatch和combine。
该库还支持低精度运算，包括FP8。
同时按惯例，开源协议用的是最为宽松的MIT。
今天的DeepSeek
选择了先在GitHub上线，然后再在官推发上新通知
。
不出所料，底下一片叫好：
DeepSeek开源列车永不停止。
DeepEP性能如何？
DeepSeek官推对DeepEP进行了要素提炼：
-
高效和优化的all-to-all通信
-
NVLink和RDMA的节点内和节点间支持
-
用于训练和推理预填充的高吞吐量内核
-
用于推理解码的低延迟内核
-
原生FP8调度支持
-
灵活的GPU资源控制，用于计算通信重叠
我们先来看看性能方面的两个重点。
（注：DeepEP中的实现可能与DeepSeek-V3论文有一些细微的差异）
具有NVLink和RDMA转发的普通内核
为了与DeepSeek-V3论文中提出的组限制门控算法保持一致，
DeepEP提供了一组针对非对称域带宽转发进行了优化的内核
，例如将数据从NVLink域转发到RDMA域。
这些内核提供高吞吐量，使其适用于训练和推理预填充任务。
此外，它们还支持SM
（Streaming Multiprocessors）
号码控制。
DeepEP团队在在H800
（~160 GB/s NVLink最大带宽）
上测试普通内核，每个内核都连接到CX7 InfiniBand 400 Gb/s RDMA网卡
（~50 GB/s 最大带宽）
。
且遵循DeepSeek-V3/R1预训练设置
（每批4096个tokens，隐藏7168个，前4组，前8个专家，FP8调度和BF16组合）
。
具有纯RDMA的低延迟内核
针对延迟敏感型推理解码场景，DeepEP包括一组具有纯RDMA的低延迟内核，以最大限度地减少延迟。
该库还引入了一种基于hook的通信计算重叠方法，不占用任何SM资源。
DeepEP团队在H800上测试低延迟内核，每个内核都连接到CX7 InfiniBand 400 Gb/s RDMA 网卡
（~50 GB/s 最大带宽）
。
且遵循典型的DeepSeek-V3/R1生产设置
（每批128个tokens，7168个隐藏，前8个专家，FP8调度和BF16组合）
。
暂不支持消费级显卡，建议使用最佳自动优化配置
在GitHub上，DeepSeek团队明确写出了关于DeepEP的使用方式，涵盖各种适配环境、配置要求等。
首先是DeepEP需要的软硬件环境版本：
其次，使用DeepEP需要下载并安装团队修改后的NVSHMEM依赖项
（有关说明，请参阅DeepSeek团队的NVSHMEM安装指南）
。
然后，将 deep_ep 导入到Python项目中，就开始“尽情享受吧”！
至于网络配置方面，DeepEP已通过InfiniBand网络的全面测试。
但理论上，它也与基于融合以太网的RDMA
（RoCE）
兼容。
其中，InfiniBand通过虚拟通道
（Virtual Lanes， VL）
支持流量隔离。
为了防止不同类型流量之间的干扰，DeepEP图男队建议将工作负载隔离到不同的虚拟通道中，如下所示：
-
使用普通内核的工作负载
-
使用低延迟内核的工作负载
-
其它工作负载
对于DeepEP，开发者可以通过设置 NVSHMEM_IB_SL 环境变量来控制虚拟通道分配。
值得注意的是，自适应路由是InfiniBand交换机提供的一项高级路由功能，可以在多个路径之间均匀分配流量。
目前，低延迟内核支持Adaptive Routing，而普通内核不支持
（可能很快就会添加支持）
。
为普通的节点间内核启用自适应路由，可能会导致死锁或数据损坏问题
。
对于低延迟内核，启用Adaptive routing可以完全消除路由冲突导致的网络拥塞，但也会带来额外的延迟。
DeepEP团队建议使用以下配置以获得最佳性能：
-
在网络负载较重的环境中启用自适应路由
-
在网络负载较轻的环境中使用静态路由
BTW，DeepEP
已禁用拥塞控制
（Congestion control）
，因为团队在生产环境中没有观察到明显的拥塞。
最后一点来自DeepEP团队的叮嘱——
为了获得极致性能，团队发现并使用了一条out-of-doc PTX指令ld.global.nc.L1::no_allocate.L2::256B 。
此指令将导致未定义的行为：使用非相干只读PTX修饰符 .nc 访问易失性GPU内存。
但是，正确性已经过测试，以保证 。L1：：no_allocate 在 Hopper 架构上，性能会好得多。
如果您发现内核在某些其他平台上无法运行，您可以添加到DISABLE_AGGRESSIVE_PTX_INSTRS=1 setup.py并禁用此功能，或提交问题。
