DeepSeek 开源活动第二天，新开源库如约而至。
今天发布的是 DeepEP，第一个用于混合专家（MoE，Mixture of Experts）模型训练和推理的开源 EP 通信库。
GitHub 库地址：https://github.com/deepseek-ai/DeepEP
其优势包括：
✅高效、优化的 all-to-all 通信
✅节点内和节点间（intranode and internode）均支持 NVLink 和 RDMA
✅用于训练和推理预填充的高吞吐量内核
✅用于推理解码的低延迟内核
✅原生 FP8 调度支持
✅灵活的 GPU 资源控制，实现计算-通信重叠
（来源：DeepSeek）
EP 的全名是 Expert Parallelism，又称专家并行。它是在分布式学习中专门针对 MoE 场景的并行策略，其核心思想就是将不同专家放在不同计算节点上进行并行计算。
据 DeepSeek 介绍，DeepEP 是一个为 MoE 和 EP 量身定制的通信库。它能提供高吞吐量和低延迟的 all-to-all GPU 内核，也称为 MoE 调度和组合。该库还支持低精度操作，包括 FP8。性能基准测试显示吞吐量显著提升，尤其是在训练和推理预填充任务中。
为了与 DeepSeek-V3 论文中提出的组限制门控算法（group-limited gating algorithm）保持一致，DeepEP 提供了一组针对非对称域带宽转发（例如将数据从 NVLink 域转发到 RDMA 域）进行优化的内核。
这些内核提供了高吞吐量，使其适合训练和推理预填充任务。此外，它们还支持 SM（Streaming Multiprocessors，流式多处理器）数量控制。
（来源：DeepSeek）
对于延迟敏感的推理解码，DeepEP 包含一组具有纯 RDMA 的低延迟内核，以最大限度地减少延迟。该库还引入了一种基于 hook 的通信-计算重叠方法，该方法不占用任何流式多处理器资源。
值得注意的是，DeepSeek 特意指出，该库中的实现可能与 DeepSeek-V3 论文略有不同。
在性能方面，针对具有 NVLink 和 RDMA 转发的普通内核，DeepSeek 在 H800（~160GB/s NVLink 最大带宽）上测试了它们，每个内核都连接到 CX7 InfiniBand 400Gb/s RDMA 网卡（~50GB/s 最大带宽）。他们遵循的是 DeepSeek-V3/R1 预训练设置。
（来源：DeepSeek）
针对具有纯 RDMA 的低延迟内核，DeepSeek 同样在 H800 上测试了它们，每个内核都连接到 CX7 InfiniBand 400Gb/s RDMA 网卡（最大带宽约 50GB/s）。他们遵循的是 DeepSeek-V3/R1 生产设置。
（来源：DeepSeek）
在网络配置方面，DeepEP 已在 InfiniBand 网络上进行了全面测试。不过，理论上它也兼容融合以太网上的 RDMA (RoCE)。
InfiniBand 通过虚拟通道支持流量隔离。为了防止不同类型的流量之间发生干扰，DeepSeek 建议在不同的虚拟通道之间划分工作负载，包括使用普通内核的工作负载、使用低延迟内核的工作负载，以及其他工作量。
此外，自适应路由是 InfiniBand 交换机提供的一项高级路由功能，可以将流量均匀地分布在多条路径上。目前，低延迟内核支持自适应路由，而普通内核则不支持（可能很快会添加支持）。为普通节点间内核启用自适应路由可能会导致死锁或数据损坏问题。
对于低延迟内核，启用自适应路由可以完全消除路由冲突导致的网络拥塞，但也会引入额外的延迟。DeepSeek 建议采用以下配置以获得最佳性能：在网络负载较重的环境中启用自适应路由，在网络负载较轻的环境中使用静态路由。
最后，DeepSeek 还表示，自己发现并使用了未在英伟达官方文档中披露的 PTX 指令 ld.global.nc.L1::no_allocate.L2::256B，由此可以获得“极致性能”。
（来源：DeepSeek）
从指令来推测，它似乎是一个特殊的 PTX 指令，用于从全局内存加载数据，绕过 L1 缓存，以 256 字节的事务大小通过 L2 缓存处理。
这种指令的使用反映了 DeepSeek 对 GPU 内存层次结构的深度理解，旨在通过创新工程和精细控制缓存行为来实现性能极致化。
“为了让集群获得更好的性能，我们建议运行所有测试并使用最佳的自动调整配置。默认配置在 DeepSeek 的内部集群上进行了优化。”DeepSeek 表示。
参考资料：
https://x.com/deepseek_ai/status/1894211757604049133
https://github.com/deepseek-ai/DeepEP
运营/排版：何晨龙
