DeepSeek开源DeepEP：MoE训练推理EP通信库，拥抱开源！

上周五，DeepSeek 发布了一则推文，宣称本周为开源周（OpenSourceWeek），期间将连续推出五个软件库。

昨日，他们首度开源了第一个代码库——FlashMLA。这是一款专为Hopper GPU 设计的高效MLA解码核心，发布后仅24小时内便获得了接近8000个星标（具体信息详见《刚刚，DeepSeek开源FlashMLA，推理加速核心技术，Star量飞涨中》）。

今天，DeepSeek 继续推出底层架构的创新，新开源项目为首个针对MoE（混合专家）模型训练及推理的EP通信库——DeepEP。

在分布式系统中（如多GPU训练环境），各处理单元之间需高效传递数据。在MoE架构中，这一需求显得尤为突出，因为不同的“专家”需频繁地交换信息。此外，MoE模型在“专家并行”时容易出现负载不均衡，导致每个“专家”所分配的算力不均，从而影响无关紧要“专家”的性能发挥。

此次开源的DeepEP实现了以下功能：

1. 高效优化的全到全（All-to-All）通信

2. 支持NVLink及RDMA的节点内/跨节点通信

3. 在训练及推理预填充阶段提供高吞吐量的计算核心

4. 在推理解码阶段提供低延迟的计算核心

5. 原生支持FP8数据分发

6. 灵活控制GPU资源，实现计算与通信的高效重叠

高效的通信减少了数据传输过程中的瓶颈，而优化后的计算核心则提升了处理速度，灵活的资源调度使得计算与通信得以并行进行。

MLA与MoE架构的改进可以说是DeepSeek的两项核心创新。昨日对MLA解码核心的优化，今日则公开了另一项“王牌”——MoE模型的高效通信及并行处理，深觉DeepSeek确实展现了极大的开放性！

项目链接：https://github.com/deepseek-ai/DeepEP

至于其受到的关注程度，

截至机器之心文章撰写未完时，DeepEP的星标数量已超过1000个：

项目开源后，不少评论表示：DeepSeek为MoE模型所达到的优化水平非常令人印象深刻，这类模型因其规模和复杂性而面临诸多挑战。DeepEP能够利用NVLink和RDMA等前沿硬件技术，并以FP8精度高效处理这些挑战，堪称是一项突破性成就。

还有人指出，“NVLink和RDMA的支持对大规模MoE模型来说是革命性的突破。看来DeepSeek再次在AI基础设施的可能性方面拓展了技术的边界。”

此前，曾有人质疑DeepSeek-R1仅通过模型蒸馏实现其性能，并非真正的技术创新。同时，也有人对DeepSeek所报GPU训练需求数量表示怀疑。此次开源周发布的内容从某种程度上证实，DeepSeek确实通过技术创新实现了训练效率的提升和成本的降低。

DeepEP的概述：

DeepEP是一个专门为混合专家（MoE）系统和专家并行（EP）设计的通信库。其提供高吞吐量和低延迟的全到全GPU内核，这些内核也被称作MoE的分发和合并。该库同时支持低精度操作，包括FP8。

为了与DeepSeek-V3论文中所提出的group-limited gating算法保持一致，DeepEP提供了针对非对称带宽转发优化的一整套内核，例如从NVLink域到RDMA域的数据转发。这些内核为训练和推理预填充任务提供高吞吐量，并支持流式多处理器（Streaming Multiprocessors，SM）数量的控制。

针对延迟敏感的推理解码，DeepEP引入了一套使用纯RDMA的低延迟内核，旨在最小化延迟。此外，该库还提供了一种基于hook的通信与计算重叠方法，不占用任何SM资源。

注意：本库中的实现可能与DeepSeek-V3论文存在细微差异。

DeepEP的性能表现：

具有NVLink和RDMA转发的常规内核

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合并）。

具有纯RDMA的低延迟内核

DeepSeek在H800上测试了低延迟内核，每个H800同样连接到一个CX7 InfiniBand 400 Gb/s RDMA网卡（最大带宽约50 GB/s）。他们遵循典型的DeepSeek-V3/R1生产设置（每批次128个token，7168个隐藏维度，top-8专家，FP8分发和BF16合并）。

注意事项：

• 为达到最佳性能，DeepSeek发现并使用了一个未记录在文档中的PTX指令：ld.global.nc.L1::no_allocate.L2::256B。此指令可能导致未定义行为，使用非一致性只读的PTX修饰符“.nc”访问可变GPU内存。然而，在Hopper架构上，通过“.L1::no_allocate”已进行了测试以确保其正确性并大幅提升性能。如在某些其他平台上发现该内核无法正常工作，可以在setup.py中添加DISABLE_AGGRESSIVE_PTX_INSTRS=1以禁用此功能，或提交问题反馈。

最后，让我们拭目以待第三天的发布内容。答案将于24小时后揭晓。