TA的每日心情 | 开心 2020-4-8 10:45 |
|---|
签到天数: 227 天 [LV.7]分神
|
本帖最后由 xiejin77 于 2025-2-26 09:45 编辑 ' w" f) l) W, y/ Z- r" @ R
2 J" q$ r6 U/ y M. c$ _
DeepSeek DeepEP:MoE 训练/推理加速,开源通信库背后的工程哲学
3 ?) [9 w" s, |$ s7 v) k, W在发布了 DeepSeek-V3 相关的研究成果之后,DeepSeek 团队开源了其关键组件之一:DeepEP。DeepEP,全称 Deep Expert Parallelism,是一个专为混合专家(MoE)模型和专家并行(EP)训练及推理而设计的通信库。它的发布,不仅展示了 DeepSeek 团队在分布式系统和高性能计算方面的深厚积累,更体现了一种以实际问题为导向、开放协作的工程哲学。
4 N9 \7 U9 m& B3 a: G9 ^4 }
/ x$ L2 i7 \ ~- K' Z一、DeepSeek 开源第二天,DeepEP 亮相
6 Z' N# v" p( b( {3 n7 |8 f( ~DeepSeek-V3 相关研究成果,在经济性和性能之间取得了卓越的平衡。其背后离不开两个关键因素:一是精妙的 Mixture-of-Experts (MoE) 架构设计;二是高效的底层系统支持。DeepEP 正是后者中的关键一环,它为 MoE 模型的大规模分布式训练和低延迟推理提供了必要的通信基础设施。
" i! l4 P9 J% y6 c5 ]1 f a$ n! C
8 P! f I* X$ }' i) _; iMoE 模型的核心思想是将一个大型模型分解为多个“专家”网络,每个专家负责处理输入的不同方面。这种结构使得模型可以在保持计算效率的同时,拥有庞大的参数容量。然而,这也带来了新的挑战:专家网络通常分布在不同的 GPU 上,模型参数和中间激活值需要在 GPU 之间频繁交换。通信效率直接影响着训练速度和推理性能,成为 MoE 模型能否成功的关键。DeepEP 的出现,正是为了解决这一核心问题。: }( W, ^ e" ~) ?" P) m! M
8 v% g% Z+ v' E6 |1 m; K; q- f二、DeepEP 技术亮点:超越速度的深度优化
/ X0 K% Q! u# s$ o" Y, s' r7 GDeepEP 的设计目标不仅仅是提供高速通信,更是在多个层面上进行了深度优化,使其成为 MoE 模型训练和推理的理想选择。这个其中尤其是第三点,对于非延迟内核的RDMA支持,这个其实涉及到一个非常隐蔽的设定。从本质上来说,是规避了对于大模型训练的英伟达的IB体系,甚至是Mellanox的依赖。这一点对于禁运都是有非常大意义的(利好菊厂、中兴之类的甚至还有国内生产RDMA网卡的创业企业)。
' W& }( y7 X/ p3 D& m" C
/ s c. p" d$ O0 ?5 T- z. @全到全 GPU 内核:奠定通信基石( J. q( G' n* M3 y$ {6 Y
9 p; `, c3 m; W! Q5 q7 L& M) }
MoE 模型和专家并行训练都依赖于高效的全到全 GPU 通信。DeepEP 提供了针对此场景专门优化的通信内核。这些内核充分利用 GPU 之间的互连带宽,实现了高吞吐量的数据交换,这是所有后续优化的基础。
+ C) j1 w9 `- i7 i, z
1 k& R/ \& B5 X& x9 I除了支持常规的 FP32/BF16 数据类型,DeepEP 还特别支持 FP8 等低精度操作。在许多情况下,低精度计算可以在不显著影响模型精度的情况下,显著减少通信量,从而进一步提高训练和推理效率,特别是在带宽受限的环境中。
% {- u& m& u1 q9 [% R/ M& q6 U9 Z
) E9 s5 ]8 O7 `# v h. ~0 O0 U"DeepEP 是一个为混合专家 (MoE) 和专家并行 (EP) 量身定制的通信库。它提供高吞吐量和低延迟的全到全 GPU 内核,也称为 MoE 调度和组合。该库还支持低精度操作,包括 FP8。"
& S u, h3 T% V
7 `5 J h6 z8 H! e% ]3 E4 m8 M6 g2 l非对称带宽优化内核:精细化资源管理) u1 v" t4 }) p, W4 T' u/ c
& t6 u1 b5 e5 {7 I5 fDeepSeek-V3 论文中提出的组限制门控算法(Group-Restricted Gating)是 相关的核心创新之一。为了充分发挥这一算法的优势,DeepEP 提供了一组专门针对非对称带宽环境优化的内核。2 h! j$ b" r0 G. g0 z3 O! C
0 e7 L4 D( V" b+ Z% i9 j
在实际的硬件部署中,NVLink 和 RDMA 之间的带宽往往存在差异。NVLink 通常用于 GPU 之间的直接高速互连,提供极高的带宽;而 RDMA 则用于跨节点通信,带宽相对较低。DeepEP 的这些内核能够智能地感知并利用这种非对称性,高效地将数据在 NVLink 域和 RDMA 域之间进行转发。0 i$ @% ~6 E2 ?. [' X" X
( {2 `2 P/ d3 h6 [这种精细化的资源管理不仅提高了吞吐量,还支持 SM(流式多处理器)数量控制。这意味着用户可以根据实际任务的需求,精细地调整计算资源的分配,进一步优化性能。
: l' S: _* a0 e" u7 K4 V. g" D3 r8 S7 I
"为了与DeepSeek-V3论文中提出的组限制门控算法保持一致,DeepEP 提供了一组针对非对称域带宽转发(例如将数据从 NVLink 域转发到 RDMA 域)进行优化的内核。这些内核提供高吞吐量,使其适合训练和推理预填充任务。此外,它们还支持 SM(流式多处理器)数量控制。"
8 b: v3 W4 S" u$ w0 Q
* R( W9 ?1 C% L& a/ y p; z低延迟内核:推理性能的保障" s. F8 [" H. B- s
6 q. U" D$ q) _ a& V
对于在线推理服务而言,低延迟是至关重要的性能指标。DeepEP 专门为延迟敏感的推理解码任务设计了一组低延迟内核。这些内核通过直接使用 RDMA 通信,绕过 NVLink,从而尽可能地减少通信延迟。
5 S' T: O# ~/ J
, }7 O% ^1 ?) n: Q8 m3 w& V在推理阶段,模型逐个生成 token,每次生成都需要进行通信。DeepEP 的低延迟内核能够将每次通信的时间压缩到微秒级别,从而显著提升整体的推理速度,降低用户感知的响应时间。& c7 Y$ |0 F! g& t0 X
& v8 J7 u2 D" i
"对于延迟敏感的推理解码,DeepEP 包含一组具有纯 RDMA 的低延迟内核,以最大限度地减少延迟。"
: l: [( ? e3 \: u1 l6 I% p4 R$ K U/ e u8 l4 e; K: L
通信计算重叠:系统级优化
# \: F8 K9 w8 m% F& p9 m# |; f8 J: O% A" j& ^) M# i1 R$ t5 W' y
DeepEP 引入了一种基于钩子的通信计算重叠机制,这是一种系统级的优化策略。传统的通信库通常需要占用一定的 SM 资源来进行通信调度和管理,这在一定程度上会影响计算效率。DeepEP 的创新之处在于,它的通信机制完全不占用任何 SM 资源,将宝贵的计算资源全部用于模型本身的计算。( \! k% c. [! z9 f
; B" I- m% h! ~1 s* G: N+ _这种机制通过精心设计的钩子函数,将通信操作与计算操作异步地交织在一起。当 GPU 执行计算任务时,通信操作在后台并行进行,从而实现了计算和通信的高度并行,最大化了硬件资源的利用率。
2 q+ M( k0 d$ }( y
( W0 s0 q' i$ m, Y8 o5 i$ |( x. @6 O"该库还引入了一种基于钩子的通信计算重叠方法,该方法不占用任何 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控制,允许用户精细的控制通信和计算的时间。
0 d$ U% ` w" m, s$ h0 w* f
( @1 ]) f5 i% q1 S1 d+ B9 E9 O三、DeepEP 性能数据解读:以实测为依据
8 ?8 W% \/ ~7 l2 _2 V: A. x! ^4 [% S/ W. ]DeepSeek 团队在 H800 GPU 上对 DeepEP 进行了详尽的性能测试,并公开了测试数据。这些数据有力地证明了 DeepEP 在不同场景下的卓越性能。 z3 Y- m4 N* ~3 T5 v; i9 [3 @
6 o5 u+ F) @0 ~) @2 [
普通内核性能:逼近理论极限! G o: m, F; o; k: h0 E
- C {4 P' G- `3 U. y6 ?' [, ]! s
在 H800 硬件环境(NVLink 最大带宽约 160 GB/s,连接到 CX7 InfiniBand 400 Gb/s RDMA 网卡,最大带宽约 50 GB/s)下,DeepEP 的普通内核展现出了接近理论带宽极限的吞吐量。( L! x( s. F: N' o# W# U8 I
: j6 s* i' D2 K6 n8 q% b
9 ], o5 V& N8 L. Q( J这些数据清晰地表明:5 g! G$ \2 S$ g, T& ^
, Q# ~1 I6 \- e; ]( D" z( q; j7 ^
* DeepEP 能够充分利用 NVLink 和 RDMA 的带宽资源。- {/ R, J. C7 R+ E
* 在节点内部,DeepEP 的吞吐量几乎达到了 NVLink 的理论上限。: {& i( K" l- z- O. L$ R
* 在节点之间,DeepEP 的吞吐量也接近了 RDMA 的理论上限。+ s: a' |' N3 Q8 I) u, x
* 无论是在节点内还是节点间通信,DeepEP 都展现出了极高的效率。
n+ j' w( o" m5 X, S低延迟内核性能:微秒级延迟5 L2 z0 x0 N+ u4 q0 u
" I* Q; q& h0 E g3 @! O7 I: H低延迟内核的测试数据同样证明了 DeepEP 的出色设计:. _1 }1 I0 c) h5 _" T! l
& y; \8 ~% T H( O3 N( l& }. z3 e# p7 r3 n
这些数据表明,DeepEP的低延迟内核在保持高带宽的同时,将通信延迟控制在微秒级别,为实时推理应用提供了强有力的支持。
. n; j K0 d6 M( p# U, Y
, ?+ L |+ C* i7 H X四、DeepEP 的工程哲学:实用主义与开放协作
# c6 t/ n* H1 Y# Z5 WDeepEP 的设计和实现,体现了 DeepSeek 团队一贯的工程哲学:
% Z$ y' Q: j) k# }9 k0 c1 ^9 b% K3 K4 Q: j- } D. ]
以问题为导向,实用至上7 r+ M+ }; f1 |0 ?( t% L, z
' v$ a, W- A& Y) b
DeepSeek 团队并不追求理论上的完美或“炫技”,而是始终以解决实际问题为出发点。他们深入理解 MoE 模型训练和推理的痛点,针对性地设计和优化 DeepEP。3 }& v$ c, g; o6 o: b
& g9 ]; \+ B/ G6 s8 P8 P4 z7 f: Y
一个典型的例子是,DeepSeek 团队发现并使用了行为超出文档范围的 PTX 指令 ld.global.nc.L1::no_allocate.L2::256B。虽然这个指令在官方文档中没有明确定义,但 DeepSeek 团队通过充分的实验验证了其在 Hopper 架构上的正确性和性能优势,并将其应用到 DeepEP 中。* a3 ?2 S6 Y7 ^% i- t
4 j3 y, t; C+ P8 M6 ~5 W
"为了实现极致性能,我们发现并使用了行为超出文档范围的 PTX 指令:ld.global.nc.L1::no_allocate.L2::256B。此指令将导致未定义的行为:使用非连贯只读 PTX 修饰符访问易失性 GPU 内存 .nc。但正确性已 .L1::no_allocate 在 Hopper 架构上测试以保证,并且性能会好得多。如果您发现内核在其他一些平台上无法运行,您可以添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 并 setup.py 禁用此功能,或者提交问题。"
- Y& K$ }' j0 L/ V5 h. E+ ~- Y/ Q a/ U3 c) g K6 A
PTX是底层的汇编,与硬件结构相关;发现其未公开的隐藏指令,事实上只能是尝试出来的。在实践中基于频繁的使用迭代,发现隐藏指令,这种不拘泥于常规、以实用为导向的精神,正是 DeepEP 能够实现卓越性能的关键。它体现了一种“黑客”精神:在现有条件下,充分利用一切可利用的资源,解决实际问题。- B/ y4 \& E$ W! a5 B0 N
M% X" m* Z1 ]! q开放协作,共同进步
9 {' _" U0 Z j8 q- z$ c. P: k: `( y1 J0 {/ x# u
DeepSeek 团队不仅在技术上精益求精,更有着开放、协作的精神。他们将经过实践验证的 DeepEP 开源,与整个 AI 社区分享他们的成果。. X1 L& U# _& t8 p) n7 P5 H, @+ k
0 i- d/ K2 b l3 y8 T( T6 r
DeepEP 采用 MIT 许可协议,这意味着任何人都可以自由地使用、修改和分发 DeepEP 的代码,无需担心版权问题。这种开放性将极大地促进 MoE 模型的研究和应用,降低 MoE 模型的开发门槛。. X+ \$ m; N! s/ o; W! _3 T' Y
% r0 C9 ^4 D ?! G& l$ |DeepSeek 的做法与一些闭门造车的做法形成了鲜明对比。他们不仅公布了结果,更重要的是公开了实现这一结果的关键技术(DeepEP),让整个社区都能从中受益。# P) g9 v2 _# ]8 T$ Y) ]( q' K
! S+ s' M7 @7 s+ t) q- u
软硬件协同,深入底层' [* V/ f. I2 e J2 T
# z3 i, e- W/ ~$ A- b( V# kDeepEP 的成功,也体现了 DeepSeek 团队对软硬件协同优化的深刻理解。他们不仅仅停留在算法层面,而是深入到底层硬件和系统,充分挖掘硬件的潜力。1 n0 u+ \+ x0 ~' u) |
8 `( c& ^& h$ Z# P: R# E; Z+ nDeepEP 对 NVLink 和 RDMA 的精细化利用,对 SM 数量的控制,以及对底层 PTX 指令的使用,都表明了 DeepSeek 团队对硬件特性的深刻理解。这种软硬件协同的优化思路,是实现极致性能的必由之路。: D% T3 j0 v+ L& i' \2 D
2 s( B( V3 [% i8 z6 P: }五、DeepEP 的网络配置与优化
% A- w }+ R5 ]3 v+ `( QDeepEP 在网络层面也进行了多项优化,以适应复杂的集群环境,并充分利用网络带宽。这里也有一个隐蔽设定,还是接续第二部分的,在剥离所谓的IB功能依赖。DS的工程师团队在用英伟达体系训练的时候,在HPC的论文中可以说是边用边骂,但一边骂一边还得用……开源这个项目,事实上是剥离了IB的不少复杂功能依赖,尤其是流量隔离和拥塞控制之类的。这也是工程师思维,花里胡哨的功能依赖是需要简化的。. }0 ]: _ S0 _" H2 [# ^2 ~
; M/ B7 n) A% H( p% m0 w
流量隔离( D: v! s( ?. Z2 E
, m, l7 c W* J" {: b5 ^) J IDeepEP 利用 InfiniBand 的虚拟通道 (VL) 特性,实现不同类型流量的隔离。通过将不同工作负载(如使用普通内核的工作负载、使用低延迟内核的工作负载,以及其他工作负载)分配到不同的虚拟通道,可以有效避免相互干扰,提高整体网络性能。DeepEP 通过环境变量 NVSHMEM_IB_SL 来控制虚拟通道的分配。 这种精细化的流量管理,可以确保关键任务(如 MoE 训练)获得足够的带宽资源,不受其他任务的影响。; S1 ^% [+ | ^+ \. A6 m' r
; C9 B' i( H( Y, r3 O
"为了防止不同类型的流量之间发生干扰,我们建议在不同的虚拟通道之间划分工作负载,如下所示: * 使用普通内核的工作负载 * 使用低延迟内核的工作负载 * 其他工作量 对于 DeepEP,您可以通过设置环境变量来控制虚拟通道分配NVSHMEM_IB_SL。"* b0 f |4 r; I, D! `7 C
7 L H% c0 q& a1 H: q0 _$ X
自适应路由
6 Y; P" m g8 t7 ^7 P
' E! u4 O4 D0 v B/ c8 i自适应路由是 InfiniBand 交换机提供的一项高级功能,可以将流量动态地分布在多条路径上,从而提高网络的鲁棒性和吞吐量。DeepEP 的低延迟内核支持自适应路由,可以有效消除因路由冲突导致的网络拥塞,降低延迟。* \ k7 i2 y- r) y
V' J. ?0 N d8 \1 g* ~
DeepSeek 团队根据实践经验,建议在网络负载较重的环境中启用自适应路由,以获得更好的稳定性和吞吐量;而在网络负载较轻的环境中,则可以使用静态路由,以减少路由计算的开销,进一步降低延迟。
" p# ^6 a! ?) N! K, P
: Y. @, r! {3 O"对于低延迟内核,启用自适应路由可以完全消除路由冲突导致的网络拥塞,但也会引入额外的延迟。我们建议采用以下配置以获得最佳性能: * 在网络负载较重的环境中启用自适应路由 * 在网络负载较轻的环境中,使用静态路由"
8 U X# F! s$ O4 s6 R1 U- h4 F5 ?; |/ N, M
拥塞控制$ d' P( k: ]6 A- k( z* n
! F; u8 }. t% ]1 {" _* C
DeepSeek 团队在生产环境中没有观察到明显的网络拥塞,因此 DeepEP 目前默认禁用了拥塞控制。这反映了 DeepSeek 团队务实的态度,避免引入不必要的复杂性。当然,如果未来需要,DeepEP 也保留了启用拥塞控制的灵活性。8 H& B, _/ G# ~1 I# S$ q" e
* O& `. X+ A7 r
总结:DeepEP 的深远意义
0 x- }6 I' j$ Y$ l- m V
# J3 A2 T" }2 Z: Z' o" i/ ], D9 CDeepEP 不仅仅是一个高性能的通信库,它更代表了一种先进的工程理念:7 b( p8 x% E3 ]8 t1 P
! A8 `* w* F4 K; O; k* g
以实际问题为导向: DeepEP 的所有设计决策都紧密围绕 MoE 模型训练和推理的实际需求,不追求“炫技”,而是注重解决实际问题。
9 x6 K V* k6 M; `# e软硬件协同优化: DeepEP 深入底层硬件,充分利用硬件特性,实现极致的性能。( w# L! Y' t2 Q6 t; a- k0 i# j
开放协作: DeepEP 的开源,将降低 MoE 模型的开发门槛,促进整个 AI 社区的发展。
( b6 V, j: K2 U+ R% z+ @2 WDeepEP 的发布,是 DeepSeek 团队的重要贡献。它将加速 MoE 模型的研究和应用,推动 AI 技术向更高效、更经济的方向发展。DeepEP 的成功,也再次证明了开放协作的力量,以及实用主义工程哲学的重要性。我们有理由相信,DeepSeek 团队未来将继续秉持这种精神,为 AI 社区带来更多的惊喜。* g2 N3 p' [/ G
7 u. j0 S: A, J4 X+ U' E0 o* T# s2 O
原文链接 |
评分
-
查看全部评分
|