DeepSeek 发推说本周将是开源周(OpenSourceWeek),并将连续开源五个软件库。DeepSeek 开源底层架构的创新,项目是首个用于 MoE 模型训练和推理的 EP 通信库 DeepEP。在分布式系统中(如多 GPU 训练环境),所有处理单元之间需要高效地传递数据。在 MoE 中,这点尤为重要,因为不同「专家」需要频繁交换信息。并且 MoE 模型容易在「专家并行」中出现负载不均衡,导致每个「专家」分到的算力不均,不重要的「专家」难以发挥应有的性能。
此次开源的 DeepEP 做到了:
1. 高效优化的 All-to-All 通信
2. 支持 NVLink 和 RDMA 的节点内 / 跨节点通信
3. 训练及推理预填充阶段的高吞吐量计算核心
4. 推理解码阶段的低延迟计算核心
5. 原生支持 FP8 数据分发
6. 灵活控制 GPU 资源,实现计算与通信的高效重叠
高效通信减少了数据传输的瓶颈,计算核心的优化提升了处理速度,灵活的资源调度让计算和通信不互相等待。
DeepEP定义
DeepEP 是一个专为混合专家系统(MoE)和专家并行(EP)定制的通信库。它提供高吞吐量和低延迟的 all-to-all GPU 内核, 这些内核也被称为 MoE 分发和合并。该库还支持低精度操作,包括 FP8。
为了与 DeepSeek-V3 论文中提出的 group-limited gating 算法保持一致,DeepEP 提供了一套针对非对称域带宽 forwarding 进行优化的内核,例如从 NVLink 域到 RDMA 域的数据 forwarding。这些内核提供高吞吐量,适用于训练和推理预填充(prefilling)任务。此外,它们还支持 SM(流式多处理器,Streaming Multiprocessors)数量控制。
DeepEP性能测试
具有 NVLink 和 RDMA forwarding 的常规内核:DeepSeek 在 H800 上测试常规内核(NVLink 最大带宽约 160 GB/s),每个 H800 连接到一个 CX7 InfiniBand 400 Gb/s RDMA 网卡(最大带宽约 50 GB/s)。他们遵循 DeepSeek-V3/R1 预训练设置(每批次 4096 个 token,7168 隐藏维度,top-4 组,top-8 专家,FP8 分发和 BF16 合并)。

注意事项
为了极致性能,DeepSeek 发现并使用了一个未记录在文档中的 PTX 指令:ld.global.nc.L1::no_allocate.L2::256B。这个指令会导致一个未定义的行为:使用非一致性只读 PTX 修饰符「.nc」访问易变的 GPU 内存。但在 Hopper 架构上,通过「.L1::no_allocate」已测试确保了正确性,且性能会大幅提升。如果你发现内核在某些其他平台上不 work,你可以在 setup.py 中添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 来禁用此功能,或提交 issue。为了在集群上获得更好的性能,DeepSeek 建议运行所有测试并使用最佳的自动调优配置。默认配置是针对 DeepSeek 内部集群优化的。