DeepSeek DeepEP：MoE 训练/推理加速，开源通信库背后的工程哲学

xiejin77 · 发表于 2025-2-26 09:43:30

本帖最后由 xiejin77 于 2025-2-26 09:45 编辑

DeepSeek DeepEP：MoE 训练/推理加速，开源通信库背后的工程哲学
在发布了 DeepSeek-V3 相关的研究成果之后，DeepSeek 团队开源了其关键组件之一：DeepEP。DeepEP，全称 Deep Expert Parallelism，是一个专为混合专家（MoE）模型和专家并行（EP）训练及推理而设计的通信库。它的发布，不仅展示了 DeepSeek 团队在分布式系统和高性能计算方面的深厚积累，更体现了一种以实际问题为导向、开放协作的工程哲学。

一、DeepSeek 开源第二天，DeepEP 亮相
DeepSeek-V3 相关研究成果，在经济性和性能之间取得了卓越的平衡。其背后离不开两个关键因素：一是精妙的 Mixture-of-Experts (MoE) 架构设计；二是高效的底层系统支持。DeepEP 正是后者中的关键一环，它为 MoE 模型的大规模分布式训练和低延迟推理提供了必要的通信基础设施。

MoE 模型的核心思想是将一个大型模型分解为多个“专家”网络，每个专家负责处理输入的不同方面。这种结构使得模型可以在保持计算效率的同时，拥有庞大的参数容量。然而，这也带来了新的挑战：专家网络通常分布在不同的 GPU 上，模型参数和中间激活值需要在 GPU 之间频繁交换。通信效率直接影响着训练速度和推理性能，成为 MoE 模型能否成功的关键。DeepEP 的出现，正是为了解决这一核心问题。

二、DeepEP 技术亮点：超越速度的深度优化
DeepEP 的设计目标不仅仅是提供高速通信，更是在多个层面上进行了深度优化，使其成为 MoE 模型训练和推理的理想选择。这个其中尤其是第三点，对于非延迟内核的RDMA支持，这个其实涉及到一个非常隐蔽的设定。从本质上来说，是规避了对于大模型训练的英伟达的IB体系，甚至是Mellanox的依赖。这一点对于禁运都是有非常大意义的（利好菊厂、中兴之类的甚至还有国内生产RDMA网卡的创业企业）。

全到全 GPU 内核：奠定通信基石

MoE 模型和专家并行训练都依赖于高效的全到全 GPU 通信。DeepEP 提供了针对此场景专门优化的通信内核。这些内核充分利用 GPU 之间的互连带宽，实现了高吞吐量的数据交换，这是所有后续优化的基础。

除了支持常规的 FP32/BF16 数据类型，DeepEP 还特别支持 FP8 等低精度操作。在许多情况下，低精度计算可以在不显著影响模型精度的情况下，显著减少通信量，从而进一步提高训练和推理效率，特别是在带宽受限的环境中。

"DeepEP 是一个为混合专家 (MoE) 和专家并行 (EP) 量身定制的通信库。它提供高吞吐量和低延迟的全到全 GPU 内核，也称为 MoE 调度和组合。该库还支持低精度操作，包括 FP8。"

非对称带宽优化内核：精细化资源管理

DeepSeek-V3 论文中提出的组限制门控算法（Group-Restricted Gating）是相关的核心创新之一。为了充分发挥这一算法的优势，DeepEP 提供了一组专门针对非对称带宽环境优化的内核。

在实际的硬件部署中，NVLink 和 RDMA 之间的带宽往往存在差异。NVLink 通常用于 GPU 之间的直接高速互连，提供极高的带宽；而 RDMA 则用于跨节点通信，带宽相对较低。DeepEP 的这些内核能够智能地感知并利用这种非对称性，高效地将数据在 NVLink 域和 RDMA 域之间进行转发。

这种精细化的资源管理不仅提高了吞吐量，还支持 SM（流式多处理器）数量控制。这意味着用户可以根据实际任务的需求，精细地调整计算资源的分配，进一步优化性能。

"为了与DeepSeek-V3论文中提出的组限制门控算法保持一致，DeepEP 提供了一组针对非对称域带宽转发（例如将数据从 NVLink 域转发到 RDMA 域）进行优化的内核。这些内核提供高吞吐量，使其适合训练和推理预填充任务。此外，它们还支持 SM（流式多处理器）数量控制。"

低延迟内核：推理性能的保障

对于在线推理服务而言，低延迟是至关重要的性能指标。DeepEP 专门为延迟敏感的推理解码任务设计了一组低延迟内核。这些内核通过直接使用 RDMA 通信，绕过 NVLink，从而尽可能地减少通信延迟。

在推理阶段，模型逐个生成 token，每次生成都需要进行通信。DeepEP 的低延迟内核能够将每次通信的时间压缩到微秒级别，从而显著提升整体的推理速度，降低用户感知的响应时间。

"对于延迟敏感的推理解码，DeepEP 包含一组具有纯 RDMA 的低延迟内核，以最大限度地减少延迟。"

通信计算重叠：系统级优化

DeepEP 引入了一种基于钩子的通信计算重叠机制，这是一种系统级的优化策略。传统的通信库通常需要占用一定的 SM 资源来进行通信调度和管理，这在一定程度上会影响计算效率。DeepEP 的创新之处在于，它的通信机制完全不占用任何 SM 资源，将宝贵的计算资源全部用于模型本身的计算。

这种机制通过精心设计的钩子函数，将通信操作与计算操作异步地交织在一起。当 GPU 执行计算任务时，通信操作在后台并行进行，从而实现了计算和通信的高度并行，最大化了硬件资源的利用率。

"该库还引入了一种基于钩子的通信计算重叠方法，该方法不占用任何 SM 资源。" "NOTES: the actual tensor will not be received only if you call hook(), it is useful for double-batch overlapping, but without any SM occupation" DeepEP还提供了双批次（Double-Batch）重叠的方法，通过return_recv_hook控制，允许用户精细的控制通信和计算的时间。

三、DeepEP 性能数据解读：以实测为依据
DeepSeek 团队在 H800 GPU 上对 DeepEP 进行了详尽的性能测试，并公开了测试数据。这些数据有力地证明了 DeepEP 在不同场景下的卓越性能。

普通内核性能：逼近理论极限

在 H800 硬件环境（NVLink 最大带宽约 160 GB/s，连接到 CX7 InfiniBand 400 Gb/s RDMA 网卡，最大带宽约 50 GB/s）下，DeepEP 的普通内核展现出了接近理论带宽极限的吞吐量。

这些数据清晰地表明：

* DeepEP 能够充分利用 NVLink 和 RDMA 的带宽资源。
* 在节点内部，DeepEP 的吞吐量几乎达到了 NVLink 的理论上限。
* 在节点之间，DeepEP 的吞吐量也接近了 RDMA 的理论上限。
* 无论是在节点内还是节点间通信，DeepEP 都展现出了极高的效率。
低延迟内核性能：微秒级延迟

低延迟内核的测试数据同样证明了 DeepEP 的出色设计：

这些数据表明，DeepEP的低延迟内核在保持高带宽的同时，将通信延迟控制在微秒级别，为实时推理应用提供了强有力的支持。

四、DeepEP 的工程哲学：实用主义与开放协作
DeepEP 的设计和实现，体现了 DeepSeek 团队一贯的工程哲学：

以问题为导向，实用至上

DeepSeek 团队并不追求理论上的完美或“炫技”，而是始终以解决实际问题为出发点。他们深入理解 MoE 模型训练和推理的痛点，针对性地设计和优化 DeepEP。

一个典型的例子是，DeepSeek 团队发现并使用了行为超出文档范围的 PTX 指令 ld.global.nc.L1::no_allocate.L2::256B。虽然这个指令在官方文档中没有明确定义，但 DeepSeek 团队通过充分的实验验证了其在 Hopper 架构上的正确性和性能优势，并将其应用到 DeepEP 中。

"为了实现极致性能，我们发现并使用了行为超出文档范围的 PTX 指令：ld.global.nc.L1::no_allocate.L2::256B。此指令将导致未定义的行为：使用非连贯只读 PTX 修饰符访问易失性 GPU 内存 .nc。但正确性已 .L1::no_allocate 在 Hopper 架构上测试以保证，并且性能会好得多。如果您发现内核在其他一些平台上无法运行，您可以添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 并 setup.py 禁用此功能，或者提交问题。"

PTX是底层的汇编，与硬件结构相关；发现其未公开的隐藏指令，事实上只能是尝试出来的。在实践中基于频繁的使用迭代，发现隐藏指令，这种不拘泥于常规、以实用为导向的精神，正是 DeepEP 能够实现卓越性能的关键。它体现了一种“黑客”精神：在现有条件下，充分利用一切可利用的资源，解决实际问题。

开放协作，共同进步

DeepSeek 团队不仅在技术上精益求精，更有着开放、协作的精神。他们将经过实践验证的 DeepEP 开源，与整个 AI 社区分享他们的成果。

DeepEP 采用 MIT 许可协议，这意味着任何人都可以自由地使用、修改和分发 DeepEP 的代码，无需担心版权问题。这种开放性将极大地促进 MoE 模型的研究和应用，降低 MoE 模型的开发门槛。

DeepSeek 的做法与一些闭门造车的做法形成了鲜明对比。他们不仅公布了结果，更重要的是公开了实现这一结果的关键技术（DeepEP），让整个社区都能从中受益。

软硬件协同，深入底层

DeepEP 的成功，也体现了 DeepSeek 团队对软硬件协同优化的深刻理解。他们不仅仅停留在算法层面，而是深入到底层硬件和系统，充分挖掘硬件的潜力。

DeepEP 对 NVLink 和 RDMA 的精细化利用，对 SM 数量的控制，以及对底层 PTX 指令的使用，都表明了 DeepSeek 团队对硬件特性的深刻理解。这种软硬件协同的优化思路，是实现极致性能的必由之路。

五、DeepEP 的网络配置与优化
DeepEP 在网络层面也进行了多项优化，以适应复杂的集群环境，并充分利用网络带宽。这里也有一个隐蔽设定，还是接续第二部分的，在剥离所谓的IB功能依赖。DS的工程师团队在用英伟达体系训练的时候，在HPC的论文中可以说是边用边骂，但一边骂一边还得用……开源这个项目，事实上是剥离了IB的不少复杂功能依赖，尤其是流量隔离和拥塞控制之类的。这也是工程师思维，花里胡哨的功能依赖是需要简化的。

流量隔离

DeepEP 利用 InfiniBand 的虚拟通道 (VL) 特性，实现不同类型流量的隔离。通过将不同工作负载（如使用普通内核的工作负载、使用低延迟内核的工作负载，以及其他工作负载）分配到不同的虚拟通道，可以有效避免相互干扰，提高整体网络性能。DeepEP 通过环境变量 NVSHMEM_IB_SL 来控制虚拟通道的分配。这种精细化的流量管理，可以确保关键任务（如 MoE 训练）获得足够的带宽资源，不受其他任务的影响。

"为了防止不同类型的流量之间发生干扰，我们建议在不同的虚拟通道之间划分工作负载，如下所示： * 使用普通内核的工作负载 * 使用低延迟内核的工作负载 * 其他工作量对于 DeepEP，您可以通过设置环境变量来控制虚拟通道分配NVSHMEM_IB_SL。"

自适应路由

自适应路由是 InfiniBand 交换机提供的一项高级功能，可以将流量动态地分布在多条路径上，从而提高网络的鲁棒性和吞吐量。DeepEP 的低延迟内核支持自适应路由，可以有效消除因路由冲突导致的网络拥塞，降低延迟。

DeepSeek 团队根据实践经验，建议在网络负载较重的环境中启用自适应路由，以获得更好的稳定性和吞吐量；而在网络负载较轻的环境中，则可以使用静态路由，以减少路由计算的开销，进一步降低延迟。

"对于低延迟内核，启用自适应路由可以完全消除路由冲突导致的网络拥塞，但也会引入额外的延迟。我们建议采用以下配置以获得最佳性能： * 在网络负载较重的环境中启用自适应路由 * 在网络负载较轻的环境中，使用静态路由"

拥塞控制

DeepSeek 团队在生产环境中没有观察到明显的网络拥塞，因此 DeepEP 目前默认禁用了拥塞控制。这反映了 DeepSeek 团队务实的态度，避免引入不必要的复杂性。当然，如果未来需要，DeepEP 也保留了启用拥塞控制的灵活性。

总结：DeepEP 的深远意义

DeepEP 不仅仅是一个高性能的通信库，它更代表了一种先进的工程理念：

以实际问题为导向： DeepEP 的所有设计决策都紧密围绕 MoE 模型训练和推理的实际需求，不追求“炫技”，而是注重解决实际问题。
软硬件协同优化： DeepEP 深入底层硬件，充分利用硬件特性，实现极致的性能。
开放协作： DeepEP 的开源，将降低 MoE 模型的开发门槛，促进整个 AI 社区的发展。
DeepEP 的发布，是 DeepSeek 团队的重要贡献。它将加速 MoE 模型的研究和应用，推动 AI 技术向更高效、更经济的方向发展。DeepEP 的成功，也再次证明了开放协作的力量，以及实用主义工程哲学的重要性。我们有理由相信，DeepSeek 团队未来将继续秉持这种精神，为 AI 社区带来更多的惊喜。

原文链接

马鹿 · 发表于 2025-2-26 22:53:09

分布控制是必走之路，没见到这个DeepEP之前，我就知道不管是哪家，这个分布控制早晚的事。我不是it的，但是现代工业控制用的DCS就是分散/分布控制，兼顾了速度和规模。

xiejin77 · 发表于 2025-2-27 13:36:47

马鹿发表于 2025-2-26 22:53
* W/ d) p$ B% E5 F分布控制是必走之路，没见到这个DeepEP之前，我就知道不管是哪家，这个分布控制早晚的事。我不是it的， ...

马鹿老师说的大势是非常准确的。

只是这个分布式是需要平衡代价和效用的，这几年的创新大都有分布式的影子，物联网，分布式数据库，分布式存储，分布式机器学习甚至是区块链、加密数字货币都是如此。

但是对于大模型训推这个领域，计算密集、数据密集、通讯密集；分布式到底能做成啥，还需要进一步演化。

马鹿 · 发表于 2025-2-27 21:20:54

xiejin77 发表于 2025-2-27 00:36" }5 g7 o h2 B9 F
马鹿老师说的大势是非常准确的。: B$ h% g- F d+ U2 l5 \1 L0 u

* q& j; c% `5 m只是这个分布式是需要平衡代价和效用的，这几年的创新大都有分布式的影 ...

正常，不同DCS厂家的底层也不一定是一样的（没用过所有的，但是到目前我没见过一样的。。。），而且一直在发展，至少前几年我看到某家的DCS又进步了。。。而且随着芯片能力的增强，以前写代码要特别注意的事（比如执行时间），现在都不用那么抠门地考虑了。

		自动登录	找回密码
密码			注册

[科技前沿] DeepSeek DeepEP：MoE 训练/推理加速，开源通信库背后的工程哲学

评分

评分

浏览过的版块