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

xiejin77

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

一、DeepSeek 开源第二天，DeepEP 亮相
DeepSeek-V3 相关研究成果，在经济性和性能之间取得了卓越的平衡。其背后离不开两个关键因素：一是精妙的 Mixture-of-Experts (MoE) 架构设计；二是高效的底层系统支持。DeepEP 正是后者中的关键一环，它为 MoE 模型的大规模分布式训练和低延迟推理提供了必要的通信基础设施。
% K/ J) Y, c/ A2 P
# T5 C5 v1 i2 C6 Z/ lMoE 模型的核心思想是将一个大型模型分解为多个“专家”网络，每个专家负责处理输入的不同方面。这种结构使得模型可以在保持计算效率的同时，拥有庞大的参数容量。然而，这也带来了新的挑战：专家网络通常分布在不同的 GPU 上，模型参数和中间激活值需要在 GPU 之间频繁交换。通信效率直接影响着训练速度和推理性能，成为 MoE 模型能否成功的关键。DeepEP 的出现，正是为了解决这一核心问题。
- t  N6 a  }. {' A  x/ s
二、DeepEP 技术亮点：超越速度的深度优化
DeepEP 的设计目标不仅仅是提供高速通信，更是在多个层面上进行了深度优化，使其成为 MoE 模型训练和推理的理想选择。这个其中尤其是第三点，对于非延迟内核的RDMA支持，这个其实涉及到一个非常隐蔽的设定。从本质上来说，是规避了对于大模型训练的英伟达的IB体系，甚至是Mellanox的依赖。这一点对于禁运都是有非常大意义的（利好菊厂、中兴之类的甚至还有国内生产RDMA网卡的创业企业）。
  X) U3 c9 {& Q
全到全 GPU 内核：奠定通信基石

$ C5 X* }6 ^, V9 VMoE 模型和专家并行训练都依赖于高效的全到全 GPU 通信。DeepEP 提供了针对此场景专门优化的通信内核。这些内核充分利用 GPU 之间的互连带宽，实现了高吞吐量的数据交换，这是所有后续优化的基础。
3 t7 w8 E0 ^9 l9 R0 T5 Y. W& T
2 ^( W/ H4 M& t  w& _8 k除了支持常规的 FP32/BF16 数据类型，DeepEP 还特别支持 FP8 等低精度操作。在许多情况下，低精度计算可以在不显著影响模型精度的情况下，显著减少通信量，从而进一步提高训练和推理效率，特别是在带宽受限的环境中。
! S1 l5 J4 D6 a, n2 z: c( e! \
/ |: Y7 g9 L' Q# J"DeepEP 是一个为混合专家 (MoE) 和专家并行 (EP) 量身定制的通信库。它提供高吞吐量和低延迟的全到全 GPU 内核，也称为 MoE 调度和组合。该库还支持低精度操作，包括 FP8。"

4 s7 Q# j$ Y3 U( p5 w; f4 c: B8 q非对称带宽优化内核：精细化资源管理

$ |, g2 S2 S, t# j4 IDeepSeek-V3 论文中提出的组限制门控算法（Group-Restricted Gating）是 相关的核心创新之一。为了充分发挥这一算法的优势，DeepEP 提供了一组专门针对非对称带宽环境优化的内核。
+ \9 b" i# x5 o& X4 H
8 x  f, K1 }! I* j在实际的硬件部署中，NVLink 和 RDMA 之间的带宽往往存在差异。NVLink 通常用于 GPU 之间的直接高速互连，提供极高的带宽；而 RDMA 则用于跨节点通信，带宽相对较低。DeepEP 的这些内核能够智能地感知并利用这种非对称性，高效地将数据在 NVLink 域和 RDMA 域之间进行转发。

2 l0 B& G; _$ F+ n/ a/ S; U5 ?这种精细化的资源管理不仅提高了吞吐量，还支持 SM（流式多处理器）数量控制。这意味着用户可以根据实际任务的需求，精细地调整计算资源的分配，进一步优化性能。
1 U! ~" T+ N# H: E" ~' l
/ z4 x# _& ^2 }. d, h"为了与DeepSeek-V3论文中提出的组限制门控算法保持一致，DeepEP 提供了一组针对非对称域带宽转发（例如将数据从 NVLink 域转发到 RDMA 域）进行优化的内核。这些内核提供高吞吐量，使其适合训练和推理预填充任务。此外，它们还支持 SM（流式多处理器）数量控制。"
1 L9 h5 p+ i1 s& ~( e; F+ e7 r# a2 V
低延迟内核：推理性能的保障

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

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

4 }0 s4 H; i( ?0 i& C/ s) `"对于延迟敏感的推理解码，DeepEP 包含一组具有纯 RDMA 的低延迟内核，以最大限度地减少延迟。"

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

DeepEP 引入了一种基于钩子的通信计算重叠机制，这是一种系统级的优化策略。传统的通信库通常需要占用一定的 SM 资源来进行通信调度和管理，这在一定程度上会影响计算效率。DeepEP 的创新之处在于，它的通信机制完全不占用任何 SM 资源，将宝贵的计算资源全部用于模型本身的计算。
# S; W4 w% s( o- S
, P( _' A  q$ B% D这种机制通过精心设计的钩子函数，将通信操作与计算操作异步地交织在一起。当 GPU 执行计算任务时，通信操作在后台并行进行，从而实现了计算和通信的高度并行，最大化了硬件资源的利用率。
0 j& L! j/ p, x7 ?; Q" \
"该库还引入了一种基于钩子的通信计算重叠方法，该方法不占用任何 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控制，允许用户精细的控制通信和计算的时间。
8 o( U+ W0 b* h
0 K3 _! X% T6 O三、DeepEP 性能数据解读：以实测为依据
& q# U! d# ^# y/ @DeepSeek 团队在 H800 GPU 上对 DeepEP 进行了详尽的性能测试，并公开了测试数据。这些数据有力地证明了 DeepEP 在不同场景下的卓越性能。
6 c1 i8 }( S, }, f4 Y* g. ]$ G% r# s! ?
% _' X5 Q1 c9 o. t( @( [4 s1 P普通内核性能：逼近理论极限
+ |+ Z; q# _: W0 {* `: G" z. L9 I
在 H800 硬件环境（NVLink 最大带宽约 160 GB/s，连接到 CX7 InfiniBand 400 Gb/s RDMA 网卡，最大带宽约 50 GB/s）下，DeepEP 的普通内核展现出了接近理论带宽极限的吞吐量。
, z$ E( Y/ R! q0 K

8 D: A2 |# q8 c  g5 r; P$ g8 C# u这些数据清晰地表明：
7 D4 B# J, [9 W, g) n
  ~  e% s- {" d& L9 [. o7 e: }*   DeepEP 能够充分利用 NVLink 和 RDMA 的带宽资源。
*   在节点内部，DeepEP 的吞吐量几乎达到了 NVLink 的理论上限。
8 g+ x9 W  D7 L9 c*   在节点之间，DeepEP 的吞吐量也接近了 RDMA 的理论上限。
*   无论是在节点内还是节点间通信，DeepEP 都展现出了极高的效率。
/ t$ V# D( k+ D0 T/ O低延迟内核性能：微秒级延迟

低延迟内核的测试数据同样证明了 DeepEP 的出色设计：
/ d! A! s" p5 c& T
, i7 |7 d$ b  X& |, A$ v9 o: w
这些数据表明，DeepEP的低延迟内核在保持高带宽的同时，将通信延迟控制在微秒级别，为实时推理应用提供了强有力的支持。

四、DeepEP 的工程哲学：实用主义与开放协作
0 X- ^+ P4 C; v) L2 J& xDeepEP 的设计和实现，体现了 DeepSeek 团队一贯的工程哲学：
- q: ?, ]. D! b- n% X4 A' }" [
以问题为导向，实用至上

$ T7 r3 I# q) m" `( h% ^" mDeepSeek 团队并不追求理论上的完美或“炫技”，而是始终以解决实际问题为出发点。他们深入理解 MoE 模型训练和推理的痛点，针对性地设计和优化 DeepEP。

一个典型的例子是，DeepSeek 团队发现并使用了行为超出文档范围的 PTX 指令 ld.global.nc.L1::no_allocate.L2::256B。虽然这个指令在官方文档中没有明确定义，但 DeepSeek 团队通过充分的实验验证了其在 Hopper 架构上的正确性和性能优势，并将其应用到 DeepEP 中。
% s' }4 d" P, f% G4 ^
"为了实现极致性能，我们发现并使用了行为超出文档范围的 PTX 指令：ld.global.nc.L1::no_allocate.L2::256B。此指令将导致未定义的行为：使用非连贯只读 PTX 修饰符访问易失性 GPU 内存 .nc。但正确性已 .L1::no_allocate 在 Hopper 架构上测试以保证，并且性能会好得多。如果您发现内核在其他一些平台上无法运行，您可以添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 并 setup.py 禁用此功能，或者提交问题。"
+ @0 |6 D1 }+ }) ^
PTX是底层的汇编，与硬件结构相关；发现其未公开的隐藏指令，事实上只能是尝试出来的。在实践中基于频繁的使用迭代，发现隐藏指令，这种不拘泥于常规、以实用为导向的精神，正是 DeepEP 能够实现卓越性能的关键。它体现了一种“黑客”精神：在现有条件下，充分利用一切可利用的资源，解决实际问题。
8 l0 ?! X, T; x( N& |
开放协作，共同进步

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

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

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

& o5 D* h) O" l1 R* n# F, L$ ?7 ]软硬件协同，深入底层
% D1 x8 ?' [/ m% Q8 v3 ^- P6 ]0 C
. w9 ~0 F: V; [" ~- L7 ]DeepEP 的成功，也体现了 DeepSeek 团队对软硬件协同优化的深刻理解。他们不仅仅停留在算法层面，而是深入到底层硬件和系统，充分挖掘硬件的潜力。
- z$ O* o( z3 W) X2 T
" s6 O4 C$ f# M6 b9 oDeepEP 对 NVLink 和 RDMA 的精细化利用，对 SM 数量的控制，以及对底层 PTX 指令的使用，都表明了 DeepSeek 团队对硬件特性的深刻理解。这种软硬件协同的优化思路，是实现极致性能的必由之路。
) `6 V* i+ ?5 r. B) p$ U
五、DeepEP 的网络配置与优化
6 }+ N+ B+ L/ K- W! l- p( |DeepEP 在网络层面也进行了多项优化，以适应复杂的集群环境，并充分利用网络带宽。这里也有一个隐蔽设定，还是接续第二部分的，在剥离所谓的IB功能依赖。DS的工程师团队在用英伟达体系训练的时候，在HPC的论文中可以说是边用边骂，但一边骂一边还得用……开源这个项目，事实上是剥离了IB的不少复杂功能依赖，尤其是流量隔离和拥塞控制之类的。这也是工程师思维，花里胡哨的功能依赖是需要简化的。

流量隔离
0 V7 D, f" W, [- F+ }
4 T7 ~& ?: V2 q* h& [) gDeepEP 利用 InfiniBand 的虚拟通道 (VL) 特性，实现不同类型流量的隔离。通过将不同工作负载（如使用普通内核的工作负载、使用低延迟内核的工作负载，以及其他工作负载）分配到不同的虚拟通道，可以有效避免相互干扰，提高整体网络性能。DeepEP 通过环境变量 NVSHMEM_IB_SL 来控制虚拟通道的分配。 这种精细化的流量管理，可以确保关键任务（如 MoE 训练）获得足够的带宽资源，不受其他任务的影响。
) j7 j+ n% p! P0 A% i
7 v5 Q# S1 E5 H! h"为了防止不同类型的流量之间发生干扰，我们建议在不同的虚拟通道之间划分工作负载，如下所示： * 使用普通内核的工作负载 * 使用低延迟内核的工作负载 * 其他工作量 对于 DeepEP，您可以通过设置环境变量来控制虚拟通道分配NVSHMEM_IB_SL。"
0 c5 K% E5 u) n: ~$ p$ T
自适应路由
% g, @1 x& A  v/ \: x) k% `
自适应路由是 InfiniBand 交换机提供的一项高级功能，可以将流量动态地分布在多条路径上，从而提高网络的鲁棒性和吞吐量。DeepEP 的低延迟内核支持自适应路由，可以有效消除因路由冲突导致的网络拥塞，降低延迟。

DeepSeek 团队根据实践经验，建议在网络负载较重的环境中启用自适应路由，以获得更好的稳定性和吞吐量；而在网络负载较轻的环境中，则可以使用静态路由，以减少路由计算的开销，进一步降低延迟。
0 t7 }# p% b/ R! G& U
"对于低延迟内核，启用自适应路由可以完全消除路由冲突导致的网络拥塞，但也会引入额外的延迟。我们建议采用以下配置以获得最佳性能： * 在网络负载较重的环境中启用自适应路由 * 在网络负载较轻的环境中，使用静态路由"

, r4 k' a8 S7 U+ o  h拥塞控制

4 v8 i( G, U& _) k1 \8 d* w( q4 aDeepSeek 团队在生产环境中没有观察到明显的网络拥塞，因此 DeepEP 目前默认禁用了拥塞控制。这反映了 DeepSeek 团队务实的态度，避免引入不必要的复杂性。当然，如果未来需要，DeepEP 也保留了启用拥塞控制的灵活性。
9 I& a: Y- }, a# \) Y
' I* ~" b5 u8 w* p/ A9 y6 G总结：DeepEP 的深远意义
6 K- B; \7 M, }  p
& P3 q" D9 n. |9 [. HDeepEP 不仅仅是一个高性能的通信库，它更代表了一种先进的工程理念：
% s* M+ L, E* @! s7 T' R- v  M; {
  O& G6 r  @4 k: s+ y1 t以实际问题为导向： DeepEP 的所有设计决策都紧密围绕 MoE 模型训练和推理的实际需求，不追求“炫技”，而是注重解决实际问题。
软硬件协同优化： DeepEP 深入底层硬件，充分利用硬件特性，实现极致的性能。
$ I* ~$ R+ b* {0 ^开放协作： DeepEP 的开源，将降低 MoE 模型的开发门槛，促进整个 AI 社区的发展。
5 n. ~2 D" n' E; [) ?DeepEP 的发布，是 DeepSeek 团队的重要贡献。它将加速 MoE 模型的研究和应用，推动 AI 技术向更高效、更经济的方向发展。DeepEP 的成功，也再次证明了开放协作的力量，以及实用主义工程哲学的重要性。我们有理由相信，DeepSeek 团队未来将继续秉持这种精神，为 AI 社区带来更多的惊喜。
  ~% j' m3 Y" Z4 l; s2 O) w+ T
8 Z6 W7 ^+ p1 \原文链接

马鹿

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

xiejin77

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

马鹿老师说的大势是非常准确的。
- D2 k  a0 ]1 G6 B
' ~( U5 W$ m0 j0 q8 Q7 u只是这个分布式是需要平衡代价和效用的，这几年的创新大都有分布式的影子，物联网，分布式数据库，分布式存储，分布式机器学习甚至是区块链、加密数字货币都是如此。
6 F4 Z# l7 U+ c' r
; q% F. g4 J2 U但是对于大模型训推这个领域，计算密集、数据密集、通讯密集；分布式到底能做成啥，还需要进一步演化。

马鹿

xiejin77 发表于 2025-2-27 00:36
马鹿老师说的大势是非常准确的。
# N/ u. g: f$ t9 N9 |
只是这个分布式是需要平衡代价和效用的，这几年的创新大都有分布式的影 ...

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