爱吱声

标题: DeepSeek DeepEP:MoE 训练/推理加速,开源通信库背后的工程哲学 [打印本页]

作者: xiejin77    时间: 2025-2-26 09:43
标题: DeepSeek DeepEP:MoE 训练/推理加速,开源通信库背后的工程哲学
本帖最后由 xiejin77 于 2025-2-26 09:45 编辑 ; [( n# d0 i" W: R, [4 [- X

$ O( o! E  T" s2 {$ U# ODeepSeek DeepEP:MoE 训练/推理加速,开源通信库背后的工程哲学
* u9 g, f+ L( X- I, Y在发布了 DeepSeek-V3 相关的研究成果之后,DeepSeek 团队开源了其关键组件之一:DeepEP。DeepEP,全称 Deep Expert Parallelism,是一个专为混合专家(MoE)模型和专家并行(EP)训练及推理而设计的通信库。它的发布,不仅展示了 DeepSeek 团队在分布式系统和高性能计算方面的深厚积累,更体现了一种以实际问题为导向、开放协作的工程哲学。
+ ?9 L% ~& k, ]' M
3 K. X  b# m+ ]: l$ E6 W一、DeepSeek 开源第二天,DeepEP 亮相3 e- g! T6 @( B: k1 s* F- {7 o$ E
DeepSeek-V3 相关研究成果,在经济性和性能之间取得了卓越的平衡。其背后离不开两个关键因素:一是精妙的 Mixture-of-Experts (MoE) 架构设计;二是高效的底层系统支持。DeepEP 正是后者中的关键一环,它为 MoE 模型的大规模分布式训练和低延迟推理提供了必要的通信基础设施。5 e- _" U- B: M7 n; C: Y

; l! I0 x: @3 _: t& O. P8 B% lMoE 模型的核心思想是将一个大型模型分解为多个“专家”网络,每个专家负责处理输入的不同方面。这种结构使得模型可以在保持计算效率的同时,拥有庞大的参数容量。然而,这也带来了新的挑战:专家网络通常分布在不同的 GPU 上,模型参数和中间激活值需要在 GPU 之间频繁交换。通信效率直接影响着训练速度和推理性能,成为 MoE 模型能否成功的关键。DeepEP 的出现,正是为了解决这一核心问题。
: a2 {+ }5 i7 q% F+ `4 g& ^8 R+ \7 J) q2 `+ P. U( d+ Z2 q' G. q% C
二、DeepEP 技术亮点:超越速度的深度优化; A& o! }, B+ C/ o& V; A' O1 t
DeepEP 的设计目标不仅仅是提供高速通信,更是在多个层面上进行了深度优化,使其成为 MoE 模型训练和推理的理想选择。这个其中尤其是第三点,对于非延迟内核的RDMA支持,这个其实涉及到一个非常隐蔽的设定。从本质上来说,是规避了对于大模型训练的英伟达的IB体系,甚至是Mellanox的依赖。这一点对于禁运都是有非常大意义的(利好菊厂、中兴之类的甚至还有国内生产RDMA网卡的创业企业)。0 I) g/ m) j, D1 T6 `1 t
/ Y* e3 [; O* g/ o
全到全 GPU 内核:奠定通信基石
# w. w" ?. ?% o: f) y
# g1 V+ E1 S) l1 E. A# CMoE 模型和专家并行训练都依赖于高效的全到全 GPU 通信。DeepEP 提供了针对此场景专门优化的通信内核。这些内核充分利用 GPU 之间的互连带宽,实现了高吞吐量的数据交换,这是所有后续优化的基础。
- n. X! s; n1 w8 @2 v1 a/ b9 n  p( B) p$ l( j# q8 l
除了支持常规的 FP32/BF16 数据类型,DeepEP 还特别支持 FP8 等低精度操作。在许多情况下,低精度计算可以在不显著影响模型精度的情况下,显著减少通信量,从而进一步提高训练和推理效率,特别是在带宽受限的环境中。
3 ]9 Q2 K$ D3 F& G  w: X: y2 `$ K2 ?: x) n+ c& m3 c
"DeepEP 是一个为混合专家 (MoE) 和专家并行 (EP) 量身定制的通信库。它提供高吞吐量和低延迟的全到全 GPU 内核,也称为 MoE 调度和组合。该库还支持低精度操作,包括 FP8。"
3 H) o3 ]; C' @0 a* [4 _; c9 v! W& s# i% D# H7 h5 N" h/ ~, Q" [" a/ c
非对称带宽优化内核:精细化资源管理4 L1 G/ h0 c8 p' F7 y5 f# w. b
4 i4 t- o$ U9 m1 j1 x7 X1 D" V7 F
DeepSeek-V3 论文中提出的组限制门控算法(Group-Restricted Gating)是 相关的核心创新之一。为了充分发挥这一算法的优势,DeepEP 提供了一组专门针对非对称带宽环境优化的内核。1 W! v, m3 ]; F) ?  B

. @$ B) \$ t' l# ?在实际的硬件部署中,NVLink 和 RDMA 之间的带宽往往存在差异。NVLink 通常用于 GPU 之间的直接高速互连,提供极高的带宽;而 RDMA 则用于跨节点通信,带宽相对较低。DeepEP 的这些内核能够智能地感知并利用这种非对称性,高效地将数据在 NVLink 域和 RDMA 域之间进行转发。) m% q; h2 G  B; z

' r, n. T  W7 }7 Q* I这种精细化的资源管理不仅提高了吞吐量,还支持 SM(流式多处理器)数量控制。这意味着用户可以根据实际任务的需求,精细地调整计算资源的分配,进一步优化性能。
% X' O7 T1 f3 x4 P0 g! n5 \, C2 x8 K  o% K1 `) S, l5 Y6 i
"为了与DeepSeek-V3论文中提出的组限制门控算法保持一致,DeepEP 提供了一组针对非对称域带宽转发(例如将数据从 NVLink 域转发到 RDMA 域)进行优化的内核。这些内核提供高吞吐量,使其适合训练和推理预填充任务。此外,它们还支持 SM(流式多处理器)数量控制。"
" }, z9 s, r: F% P0 p
/ K- O- t/ n' h4 i3 P  o8 B' `低延迟内核:推理性能的保障+ P; D3 |0 L, U
0 B3 q4 D' G+ B, Y% I
对于在线推理服务而言,低延迟是至关重要的性能指标。DeepEP 专门为延迟敏感的推理解码任务设计了一组低延迟内核。这些内核通过直接使用 RDMA 通信,绕过 NVLink,从而尽可能地减少通信延迟。
1 o5 O2 C& Y! w- M! q& Z# v3 d9 B/ W0 j& g# z1 N
在推理阶段,模型逐个生成 token,每次生成都需要进行通信。DeepEP 的低延迟内核能够将每次通信的时间压缩到微秒级别,从而显著提升整体的推理速度,降低用户感知的响应时间。7 M. T8 C4 C% E7 w& {
8 N& p1 Q9 U  x* Q5 R5 A
"对于延迟敏感的推理解码,DeepEP 包含一组具有纯 RDMA 的低延迟内核,以最大限度地减少延迟。"' b7 E/ z- j8 \4 I5 a8 }
( S& P. q4 P' r$ @' s
通信计算重叠:系统级优化
8 j) W0 h7 F% @8 L+ A
6 F  B+ C( z% r+ @. s: EDeepEP 引入了一种基于钩子的通信计算重叠机制,这是一种系统级的优化策略。传统的通信库通常需要占用一定的 SM 资源来进行通信调度和管理,这在一定程度上会影响计算效率。DeepEP 的创新之处在于,它的通信机制完全不占用任何 SM 资源,将宝贵的计算资源全部用于模型本身的计算。- Y: F  Z" A- [( Y2 c' z

; i8 |  x& q# X) p这种机制通过精心设计的钩子函数,将通信操作与计算操作异步地交织在一起。当 GPU 执行计算任务时,通信操作在后台并行进行,从而实现了计算和通信的高度并行,最大化了硬件资源的利用率。5 O  x+ T; u  M# _5 k4 \8 S' G, Q
/ D6 @( G( s$ Z/ {0 P4 c
"该库还引入了一种基于钩子的通信计算重叠方法,该方法不占用任何 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控制,允许用户精细的控制通信和计算的时间。
7 Y( H4 U, A6 R) P) T) [' e4 @6 J4 B9 N. i3 p
三、DeepEP 性能数据解读:以实测为依据
( F: U/ u# q+ x( p: y0 dDeepSeek 团队在 H800 GPU 上对 DeepEP 进行了详尽的性能测试,并公开了测试数据。这些数据有力地证明了 DeepEP 在不同场景下的卓越性能。$ C7 A" Z% \# y! {
0 g# Z  f% L8 j( e
普通内核性能:逼近理论极限
6 G( h% }% Z3 x( n0 |
! D" w1 B- r% Z+ L9 [' X在 H800 硬件环境(NVLink 最大带宽约 160 GB/s,连接到 CX7 InfiniBand 400 Gb/s RDMA 网卡,最大带宽约 50 GB/s)下,DeepEP 的普通内核展现出了接近理论带宽极限的吞吐量。
7 a- \+ E1 q  f) y, s0 Z( H! {& g! O- z

9 L! G5 o5 i( g" ?) {这些数据清晰地表明:
8 B' _9 q6 h6 U) q9 J& i0 s
+ P1 W  a# T, e4 A*   DeepEP 能够充分利用 NVLink 和 RDMA 的带宽资源。2 {! U% F4 \4 X9 s1 U1 K- ]  l
*   在节点内部,DeepEP 的吞吐量几乎达到了 NVLink 的理论上限。3 T9 v# U/ e8 p+ j$ [1 }
*   在节点之间,DeepEP 的吞吐量也接近了 RDMA 的理论上限。
1 j$ W* U( K3 T: `2 W- J3 V! s* M+ g*   无论是在节点内还是节点间通信,DeepEP 都展现出了极高的效率。; c7 @: S1 D3 _/ g7 T8 Z' L
低延迟内核性能:微秒级延迟( t$ x! }3 B; X/ ]  \0 ~

+ ?7 `# X$ ?; @# c! F2 Q2 C+ @2 M低延迟内核的测试数据同样证明了 DeepEP 的出色设计:
5 w# g4 W& U, p# S. G8 V4 S( t& ]
$ J3 }: _2 g  G/ l( J+ R
这些数据表明,DeepEP的低延迟内核在保持高带宽的同时,将通信延迟控制在微秒级别,为实时推理应用提供了强有力的支持。
: X; t) P2 |1 B/ z/ N
2 |9 f: e0 L  b, v- V1 b4 t2 T四、DeepEP 的工程哲学:实用主义与开放协作" U; \: ]# e  z6 g8 C% k; V
DeepEP 的设计和实现,体现了 DeepSeek 团队一贯的工程哲学:
1 {0 T4 b  o3 t  i1 q$ H
6 @$ M8 t: B2 G+ [& ?$ H" G1 q以问题为导向,实用至上' L, K4 ^7 b" _1 {$ S
5 h# M6 O$ t' H' {% X' q. ?* u
DeepSeek 团队并不追求理论上的完美或“炫技”,而是始终以解决实际问题为出发点。他们深入理解 MoE 模型训练和推理的痛点,针对性地设计和优化 DeepEP。
9 X5 ^, J* [) x2 q' ^7 Y# @: w' F" ]! s- E) Y  _+ l" Q, r
一个典型的例子是,DeepSeek 团队发现并使用了行为超出文档范围的 PTX 指令 ld.global.nc.L1::no_allocate.L2::256B。虽然这个指令在官方文档中没有明确定义,但 DeepSeek 团队通过充分的实验验证了其在 Hopper 架构上的正确性和性能优势,并将其应用到 DeepEP 中。
2 r) z0 h, v- O. N: U2 ?( H4 k3 ?8 e3 q$ ^
"为了实现极致性能,我们发现并使用了行为超出文档范围的 PTX 指令:ld.global.nc.L1::no_allocate.L2::256B。此指令将导致未定义的行为:使用非连贯只读 PTX 修饰符访问易失性 GPU 内存 .nc。但正确性已 .L1::no_allocate 在 Hopper 架构上测试以保证,并且性能会好得多。如果您发现内核在其他一些平台上无法运行,您可以添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 并 setup.py 禁用此功能,或者提交问题。"
: |& A4 n: y7 }. Y& ?0 l* E  j. c6 ~( k* ^$ I
PTX是底层的汇编,与硬件结构相关;发现其未公开的隐藏指令,事实上只能是尝试出来的。在实践中基于频繁的使用迭代,发现隐藏指令,这种不拘泥于常规、以实用为导向的精神,正是 DeepEP 能够实现卓越性能的关键。它体现了一种“黑客”精神:在现有条件下,充分利用一切可利用的资源,解决实际问题。; v4 T6 h$ S( l5 ~  F5 o  k

$ m2 B% S7 F6 w6 p开放协作,共同进步  z: s( j  d6 k8 l* h. u

9 }; @1 L5 L. v; A) U% ~2 @5 o* `* @, y) pDeepSeek 团队不仅在技术上精益求精,更有着开放、协作的精神。他们将经过实践验证的 DeepEP 开源,与整个 AI 社区分享他们的成果。
' A& w6 }5 t% k
1 N4 U& d( s6 w6 U" eDeepEP 采用 MIT 许可协议,这意味着任何人都可以自由地使用、修改和分发 DeepEP 的代码,无需担心版权问题。这种开放性将极大地促进 MoE 模型的研究和应用,降低 MoE 模型的开发门槛。
0 \3 r! K2 x8 {4 \5 i* s0 A4 u7 e# }$ J
DeepSeek 的做法与一些闭门造车的做法形成了鲜明对比。他们不仅公布了结果,更重要的是公开了实现这一结果的关键技术(DeepEP),让整个社区都能从中受益。
1 |& |" T' \9 D! A6 |" B; D* d8 O- ]# X3 O8 c
软硬件协同,深入底层4 V' j: V+ v9 Q
# W! G# \6 N2 n3 ]
DeepEP 的成功,也体现了 DeepSeek 团队对软硬件协同优化的深刻理解。他们不仅仅停留在算法层面,而是深入到底层硬件和系统,充分挖掘硬件的潜力。: h8 \# M- R2 @" y
/ s! X, r' z& o
DeepEP 对 NVLink 和 RDMA 的精细化利用,对 SM 数量的控制,以及对底层 PTX 指令的使用,都表明了 DeepSeek 团队对硬件特性的深刻理解。这种软硬件协同的优化思路,是实现极致性能的必由之路。1 u. q5 x4 R4 [4 ]9 _$ `) h

4 y/ d" s- x+ [' q3 [; U6 Y# V五、DeepEP 的网络配置与优化
! @5 {6 v* H+ V" g% |DeepEP 在网络层面也进行了多项优化,以适应复杂的集群环境,并充分利用网络带宽。这里也有一个隐蔽设定,还是接续第二部分的,在剥离所谓的IB功能依赖。DS的工程师团队在用英伟达体系训练的时候,在HPC的论文中可以说是边用边骂,但一边骂一边还得用……开源这个项目,事实上是剥离了IB的不少复杂功能依赖,尤其是流量隔离和拥塞控制之类的。这也是工程师思维,花里胡哨的功能依赖是需要简化的。
: _: X% g: Q0 h% D# `
  h' [- P2 J- ?, A; o+ d' _流量隔离: f" n% a; L+ c" J: N0 D7 J7 H

+ K# W( c- ]0 b4 ~DeepEP 利用 InfiniBand 的虚拟通道 (VL) 特性,实现不同类型流量的隔离。通过将不同工作负载(如使用普通内核的工作负载、使用低延迟内核的工作负载,以及其他工作负载)分配到不同的虚拟通道,可以有效避免相互干扰,提高整体网络性能。DeepEP 通过环境变量 NVSHMEM_IB_SL 来控制虚拟通道的分配。 这种精细化的流量管理,可以确保关键任务(如 MoE 训练)获得足够的带宽资源,不受其他任务的影响。7 V- `0 }! D) a% @- J
* X+ O$ Z4 H8 ^; F. X
"为了防止不同类型的流量之间发生干扰,我们建议在不同的虚拟通道之间划分工作负载,如下所示: * 使用普通内核的工作负载 * 使用低延迟内核的工作负载 * 其他工作量 对于 DeepEP,您可以通过设置环境变量来控制虚拟通道分配NVSHMEM_IB_SL。"
$ C+ q0 J# T5 B7 k; }4 p! D7 T: M* }& b7 C+ A0 P  |( E2 Q
自适应路由
/ [( {! i( `% J: f7 l$ j9 ~, a6 U% z. y2 r
自适应路由是 InfiniBand 交换机提供的一项高级功能,可以将流量动态地分布在多条路径上,从而提高网络的鲁棒性和吞吐量。DeepEP 的低延迟内核支持自适应路由,可以有效消除因路由冲突导致的网络拥塞,降低延迟。. S. q) F7 d% P: V! b7 P0 ~. ?/ N

# ]3 m# h: m" H1 _5 N; IDeepSeek 团队根据实践经验,建议在网络负载较重的环境中启用自适应路由,以获得更好的稳定性和吞吐量;而在网络负载较轻的环境中,则可以使用静态路由,以减少路由计算的开销,进一步降低延迟。/ u* e9 L6 ?# m* K4 ^+ y

* F+ W1 K5 T. c7 T4 p! |7 y"对于低延迟内核,启用自适应路由可以完全消除路由冲突导致的网络拥塞,但也会引入额外的延迟。我们建议采用以下配置以获得最佳性能: * 在网络负载较重的环境中启用自适应路由 * 在网络负载较轻的环境中,使用静态路由"
: W& P( O& m# Y
& q6 u; I) U/ z1 \拥塞控制
6 v+ U9 `* j' s1 F1 n' ~. O
% K0 Z- B2 z! |; y/ |+ q" CDeepSeek 团队在生产环境中没有观察到明显的网络拥塞,因此 DeepEP 目前默认禁用了拥塞控制。这反映了 DeepSeek 团队务实的态度,避免引入不必要的复杂性。当然,如果未来需要,DeepEP 也保留了启用拥塞控制的灵活性。7 u2 W9 u, l, R" g3 I& X. @
& I& C" I+ d$ q  S
总结:DeepEP 的深远意义& p+ C* K7 y6 O6 F8 \" {

7 F  _+ k1 y6 d) _5 R  V3 iDeepEP 不仅仅是一个高性能的通信库,它更代表了一种先进的工程理念:0 w" o  D$ Z6 m$ k3 S6 G

2 V0 B8 I/ ~9 y* ~3 Q以实际问题为导向: DeepEP 的所有设计决策都紧密围绕 MoE 模型训练和推理的实际需求,不追求“炫技”,而是注重解决实际问题。& H5 X" C8 F% V  |
软硬件协同优化: DeepEP 深入底层硬件,充分利用硬件特性,实现极致的性能。- y* y. J$ D; t# a; ]8 ^
开放协作: DeepEP 的开源,将降低 MoE 模型的开发门槛,促进整个 AI 社区的发展。
/ F( v/ o; S+ m+ `! jDeepEP 的发布,是 DeepSeek 团队的重要贡献。它将加速 MoE 模型的研究和应用,推动 AI 技术向更高效、更经济的方向发展。DeepEP 的成功,也再次证明了开放协作的力量,以及实用主义工程哲学的重要性。我们有理由相信,DeepSeek 团队未来将继续秉持这种精神,为 AI 社区带来更多的惊喜。
) s: O; L  s# f9 o& Q5 _2 x* U( d0 `) Y9 E1 U$ T/ a
原文链接
作者: 马鹿    时间: 2025-2-26 22:53
分布控制是必走之路,没见到这个DeepEP之前, 我就知道不管是哪家, 这个分布控制早晚的事。 我不是it的, 但是现代工业控制用的DCS就是分散/分布控制, 兼顾了速度和规模。
作者: xiejin77    时间: 2025-2-27 13:36
马鹿 发表于 2025-2-26 22:53( p5 n8 E( g( Q+ ]$ @3 s, g! E
分布控制是必走之路,没见到这个DeepEP之前, 我就知道不管是哪家, 这个分布控制早晚的事。 我不是it的,  ...
' h3 ]+ Q9 C( k4 C' M1 y& x
马鹿老师说的大势是非常准确的。
, Q4 d3 ?; Q9 T. {9 ^$ A
  R+ }6 {- c, e+ J9 h/ {) ]5 G只是这个分布式是需要平衡代价和效用的,这几年的创新大都有分布式的影子,物联网,分布式数据库,分布式存储,分布式机器学习甚至是区块链、加密数字货币都是如此。
( v% `) R- I# S% L1 g" v
) `5 r7 n! k9 `' C( S2 r但是对于大模型训推这个领域,计算密集、数据密集、通讯密集;分布式到底能做成啥,还需要进一步演化。
作者: 马鹿    时间: 2025-2-27 21:20
xiejin77 发表于 2025-2-27 00:36# @: Z, P4 \- \  K; `
马鹿老师说的大势是非常准确的。" o% q. f) k+ q: T) u* }

. U# H/ ?4 H, C  l  J7 U% U, k, q只是这个分布式是需要平衡代价和效用的,这几年的创新大都有分布式的影 ...
5 A9 E8 e' n3 c
正常, 不同DCS厂家的底层也不一定是一样的(没用过所有的,但是到目前我没见过一样的。。。), 而且一直在发展, 至少前几年我看到某家的DCS又进步了。。。而且随着芯片能力的增强,以前写代码要特别注意的事(比如执行时间), 现在都不用那么抠门地考虑了。0 w4 g7 Q: L, c

1 q" h; w4 c2 V* J* Z- `




欢迎光临 爱吱声 (http://129.226.69.186/bbs/) Powered by Discuz! X3.2