设为首页收藏本站

爱吱声

 找回密码
 注册
搜索
查看: 911|回复: 3
打印 上一主题 下一主题

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

[复制链接]
  • TA的每日心情
    开心
    2020-4-8 10:45
  • 签到天数: 227 天

    [LV.7]分神

    跳转到指定楼层
    楼主
     楼主| 发表于 2025-2-26 09:43:30 | 只看该作者 回帖奖励 |倒序浏览 |阅读模式
    本帖最后由 xiejin77 于 2025-2-26 09:45 编辑
    ; B  [8 @, X) ?- l/ R
    ! a, f7 a  O# x' n% iDeepSeek DeepEP:MoE 训练/推理加速,开源通信库背后的工程哲学% W* M8 L( |, S7 |% r# a
    在发布了 DeepSeek-V3 相关的研究成果之后,DeepSeek 团队开源了其关键组件之一:DeepEP。DeepEP,全称 Deep Expert Parallelism,是一个专为混合专家(MoE)模型和专家并行(EP)训练及推理而设计的通信库。它的发布,不仅展示了 DeepSeek 团队在分布式系统和高性能计算方面的深厚积累,更体现了一种以实际问题为导向、开放协作的工程哲学。
    3 E3 L2 z7 O- K% I* s
    ' T9 \3 T- j1 ~: F1 w" k0 f+ n4 m6 ~& q一、DeepSeek 开源第二天,DeepEP 亮相3 a  L; h+ X  z; k6 q4 j
    DeepSeek-V3 相关研究成果,在经济性和性能之间取得了卓越的平衡。其背后离不开两个关键因素:一是精妙的 Mixture-of-Experts (MoE) 架构设计;二是高效的底层系统支持。DeepEP 正是后者中的关键一环,它为 MoE 模型的大规模分布式训练和低延迟推理提供了必要的通信基础设施。* k3 _) W: v: {# ]
    & F6 O: _1 S2 f7 |
    MoE 模型的核心思想是将一个大型模型分解为多个“专家”网络,每个专家负责处理输入的不同方面。这种结构使得模型可以在保持计算效率的同时,拥有庞大的参数容量。然而,这也带来了新的挑战:专家网络通常分布在不同的 GPU 上,模型参数和中间激活值需要在 GPU 之间频繁交换。通信效率直接影响着训练速度和推理性能,成为 MoE 模型能否成功的关键。DeepEP 的出现,正是为了解决这一核心问题。
    3 d6 r* L( A6 O% B9 p
    & E6 M; M( Z3 E$ |二、DeepEP 技术亮点:超越速度的深度优化. H' m1 E* h3 t  }& C
    DeepEP 的设计目标不仅仅是提供高速通信,更是在多个层面上进行了深度优化,使其成为 MoE 模型训练和推理的理想选择。这个其中尤其是第三点,对于非延迟内核的RDMA支持,这个其实涉及到一个非常隐蔽的设定。从本质上来说,是规避了对于大模型训练的英伟达的IB体系,甚至是Mellanox的依赖。这一点对于禁运都是有非常大意义的(利好菊厂、中兴之类的甚至还有国内生产RDMA网卡的创业企业)。* M* [" D: v( ]" [
    ( l4 K/ h) j7 Z- Y6 x$ e
    全到全 GPU 内核:奠定通信基石* t) G! |+ k, ^
    7 S( a  E1 I# w6 K; ]
    MoE 模型和专家并行训练都依赖于高效的全到全 GPU 通信。DeepEP 提供了针对此场景专门优化的通信内核。这些内核充分利用 GPU 之间的互连带宽,实现了高吞吐量的数据交换,这是所有后续优化的基础。
    7 o6 O* C! a8 R3 ]" J1 y3 c5 `
    8 w: b- ~. [% d: n5 e5 f8 X# A除了支持常规的 FP32/BF16 数据类型,DeepEP 还特别支持 FP8 等低精度操作。在许多情况下,低精度计算可以在不显著影响模型精度的情况下,显著减少通信量,从而进一步提高训练和推理效率,特别是在带宽受限的环境中。- b% }! s6 Y7 Y: c

    + }( I  _* x! m% Q: W  ]7 r"DeepEP 是一个为混合专家 (MoE) 和专家并行 (EP) 量身定制的通信库。它提供高吞吐量和低延迟的全到全 GPU 内核,也称为 MoE 调度和组合。该库还支持低精度操作,包括 FP8。"
    # X% h) w9 j7 {3 V
    " v! i% d4 U% |9 ?  P非对称带宽优化内核:精细化资源管理% G9 b, w7 D1 Y$ w& Y1 I* l3 o
    ; v# V: m/ O$ G3 q2 _
    DeepSeek-V3 论文中提出的组限制门控算法(Group-Restricted Gating)是 相关的核心创新之一。为了充分发挥这一算法的优势,DeepEP 提供了一组专门针对非对称带宽环境优化的内核。5 V# ^( C  I/ U% V

    " K  `. l( u- W& b" R在实际的硬件部署中,NVLink 和 RDMA 之间的带宽往往存在差异。NVLink 通常用于 GPU 之间的直接高速互连,提供极高的带宽;而 RDMA 则用于跨节点通信,带宽相对较低。DeepEP 的这些内核能够智能地感知并利用这种非对称性,高效地将数据在 NVLink 域和 RDMA 域之间进行转发。# J: n1 Y  l" X+ ?2 _% O% D

    ! Q4 U3 E7 A5 I9 ?0 K( ]# ^这种精细化的资源管理不仅提高了吞吐量,还支持 SM(流式多处理器)数量控制。这意味着用户可以根据实际任务的需求,精细地调整计算资源的分配,进一步优化性能。
    , f0 M  N( v( s2 g' l& [+ v+ ]# L# Y0 ?$ ]
    "为了与DeepSeek-V3论文中提出的组限制门控算法保持一致,DeepEP 提供了一组针对非对称域带宽转发(例如将数据从 NVLink 域转发到 RDMA 域)进行优化的内核。这些内核提供高吞吐量,使其适合训练和推理预填充任务。此外,它们还支持 SM(流式多处理器)数量控制。"6 P8 G9 H# _& I( q2 x* k  F! K

    + m& D; V! ^) P7 |低延迟内核:推理性能的保障* q( V2 V# t: t6 X

    7 _% n, C$ h6 F$ ]- }/ O" H; w' w对于在线推理服务而言,低延迟是至关重要的性能指标。DeepEP 专门为延迟敏感的推理解码任务设计了一组低延迟内核。这些内核通过直接使用 RDMA 通信,绕过 NVLink,从而尽可能地减少通信延迟。" t( I1 E: |& s, r1 B/ ]3 m- h2 U4 ]2 G
    7 E* O+ c" d8 b5 j! b
    在推理阶段,模型逐个生成 token,每次生成都需要进行通信。DeepEP 的低延迟内核能够将每次通信的时间压缩到微秒级别,从而显著提升整体的推理速度,降低用户感知的响应时间。
    . M4 `+ [; I0 T8 g8 B# l" s0 P% E8 Y/ Z; Y4 |5 Y6 P, `7 m6 O
    "对于延迟敏感的推理解码,DeepEP 包含一组具有纯 RDMA 的低延迟内核,以最大限度地减少延迟。"+ ~% M+ D9 [' u1 n8 I9 e7 ^6 i

    ; A2 s! q6 Y7 ?( A: g通信计算重叠:系统级优化
    % a; D1 S, P+ a0 l* p5 y# a
    5 r% N2 c# X4 t2 ?8 zDeepEP 引入了一种基于钩子的通信计算重叠机制,这是一种系统级的优化策略。传统的通信库通常需要占用一定的 SM 资源来进行通信调度和管理,这在一定程度上会影响计算效率。DeepEP 的创新之处在于,它的通信机制完全不占用任何 SM 资源,将宝贵的计算资源全部用于模型本身的计算。9 }- e  ]& I! s! n  ~" f1 }; r

    + P! h. T3 C$ d0 Q1 T这种机制通过精心设计的钩子函数,将通信操作与计算操作异步地交织在一起。当 GPU 执行计算任务时,通信操作在后台并行进行,从而实现了计算和通信的高度并行,最大化了硬件资源的利用率。/ U+ p" o, R7 g. P0 p& o/ u# m
    1 z9 P) t; u) `
    "该库还引入了一种基于钩子的通信计算重叠方法,该方法不占用任何 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控制,允许用户精细的控制通信和计算的时间。
      ]/ m& k0 M  c4 ]- Y; B/ Z4 Z0 N& a: P$ N7 E
    三、DeepEP 性能数据解读:以实测为依据
    6 r' `2 d( e& u$ N7 bDeepSeek 团队在 H800 GPU 上对 DeepEP 进行了详尽的性能测试,并公开了测试数据。这些数据有力地证明了 DeepEP 在不同场景下的卓越性能。
    5 Z: l# G; a& c3 C, _0 E5 _
    8 R, b: `4 H* x$ w普通内核性能:逼近理论极限0 b& ?: E/ D6 _9 K0 d. b

    ; A$ V% b6 ~+ O) p5 _5 n$ a2 M' @在 H800 硬件环境(NVLink 最大带宽约 160 GB/s,连接到 CX7 InfiniBand 400 Gb/s RDMA 网卡,最大带宽约 50 GB/s)下,DeepEP 的普通内核展现出了接近理论带宽极限的吞吐量。
    & m: C2 x. M% g1 @% g. Z7 A* n$ Z0 `$ F8 ^* M# r" q
    : \+ z3 G* n- P; P, K" E) F
    这些数据清晰地表明:
    0 f1 u- w8 J9 Z3 Y, O9 |! A& A% \8 a
    *   DeepEP 能够充分利用 NVLink 和 RDMA 的带宽资源。
    " w& ^& q& \5 g$ F2 R, _8 s, w# B*   在节点内部,DeepEP 的吞吐量几乎达到了 NVLink 的理论上限。
    " A' R' }+ X8 n. W8 p6 l- x4 ]*   在节点之间,DeepEP 的吞吐量也接近了 RDMA 的理论上限。. {$ Y) X7 c! l- c
    *   无论是在节点内还是节点间通信,DeepEP 都展现出了极高的效率。
    . n/ N# n6 Z; @# D低延迟内核性能:微秒级延迟
    ) Y$ w  I- Z& n3 U' Z) L5 J. q- W3 G
    * Z* C6 @4 b9 q1 O4 E# Z低延迟内核的测试数据同样证明了 DeepEP 的出色设计:
    & O& t' _5 G* L8 E  K. A6 _6 [+ r0 Y( s3 j! m3 k0 L9 [1 s: v* y
    8 b% G4 Z! Q# V
    这些数据表明,DeepEP的低延迟内核在保持高带宽的同时,将通信延迟控制在微秒级别,为实时推理应用提供了强有力的支持。
    7 t7 A5 q/ E* g/ {. T% F; z* V" J8 P5 Y4 E/ P
    四、DeepEP 的工程哲学:实用主义与开放协作
    * w, q; Z' v$ E5 eDeepEP 的设计和实现,体现了 DeepSeek 团队一贯的工程哲学:! ^# a, p( K: b2 R
    6 o5 y4 a4 y6 s" M" x' _. H
    以问题为导向,实用至上
    1 \8 c- J* {) z3 \. h8 a& n" q' g* L* y" r
    DeepSeek 团队并不追求理论上的完美或“炫技”,而是始终以解决实际问题为出发点。他们深入理解 MoE 模型训练和推理的痛点,针对性地设计和优化 DeepEP。
    1 J$ s1 T: r* l9 R* l" C/ T2 O1 u- z) Q
    一个典型的例子是,DeepSeek 团队发现并使用了行为超出文档范围的 PTX 指令 ld.global.nc.L1::no_allocate.L2::256B。虽然这个指令在官方文档中没有明确定义,但 DeepSeek 团队通过充分的实验验证了其在 Hopper 架构上的正确性和性能优势,并将其应用到 DeepEP 中。
    % N. c1 |4 {' I4 l# k+ I
    # o/ V7 s& f  o. n"为了实现极致性能,我们发现并使用了行为超出文档范围的 PTX 指令:ld.global.nc.L1::no_allocate.L2::256B。此指令将导致未定义的行为:使用非连贯只读 PTX 修饰符访问易失性 GPU 内存 .nc。但正确性已 .L1::no_allocate 在 Hopper 架构上测试以保证,并且性能会好得多。如果您发现内核在其他一些平台上无法运行,您可以添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 并 setup.py 禁用此功能,或者提交问题。"+ m# E) ?1 t% G) K! r0 p: _

    ; l/ K6 \( J! I. u+ d- UPTX是底层的汇编,与硬件结构相关;发现其未公开的隐藏指令,事实上只能是尝试出来的。在实践中基于频繁的使用迭代,发现隐藏指令,这种不拘泥于常规、以实用为导向的精神,正是 DeepEP 能够实现卓越性能的关键。它体现了一种“黑客”精神:在现有条件下,充分利用一切可利用的资源,解决实际问题。
    , k; G$ P6 a6 v3 B0 S  K
    9 _, D5 Z8 O' L& z8 a+ w开放协作,共同进步
    ! {0 V2 R: C$ }2 X9 x0 t. Z- d4 p6 n2 G4 w: a9 y. W9 F  _
    DeepSeek 团队不仅在技术上精益求精,更有着开放、协作的精神。他们将经过实践验证的 DeepEP 开源,与整个 AI 社区分享他们的成果。
    - O" g6 L, }. d: {# ]9 o
    ; t# B( X4 @; o, XDeepEP 采用 MIT 许可协议,这意味着任何人都可以自由地使用、修改和分发 DeepEP 的代码,无需担心版权问题。这种开放性将极大地促进 MoE 模型的研究和应用,降低 MoE 模型的开发门槛。" z" \/ K5 O$ m, Z
    ' q& s" P$ _! v# x+ t+ q! f" X  w
    DeepSeek 的做法与一些闭门造车的做法形成了鲜明对比。他们不仅公布了结果,更重要的是公开了实现这一结果的关键技术(DeepEP),让整个社区都能从中受益。
    ) S6 s. n6 y' m! u
    + }: j1 g. ~) ^软硬件协同,深入底层
    % k0 }' k- T! @8 V  L8 g
    * n: ?4 E+ S' X( A: M; x7 eDeepEP 的成功,也体现了 DeepSeek 团队对软硬件协同优化的深刻理解。他们不仅仅停留在算法层面,而是深入到底层硬件和系统,充分挖掘硬件的潜力。$ y9 J% E# t5 ^0 Z3 s
    ( N+ U2 Z8 N0 X: p6 V/ O$ `/ R0 x
    DeepEP 对 NVLink 和 RDMA 的精细化利用,对 SM 数量的控制,以及对底层 PTX 指令的使用,都表明了 DeepSeek 团队对硬件特性的深刻理解。这种软硬件协同的优化思路,是实现极致性能的必由之路。
    . {; F$ g$ Y8 |4 k- ?
      i; C+ Q7 F9 y五、DeepEP 的网络配置与优化
    # r5 a' ]2 I4 _3 {* qDeepEP 在网络层面也进行了多项优化,以适应复杂的集群环境,并充分利用网络带宽。这里也有一个隐蔽设定,还是接续第二部分的,在剥离所谓的IB功能依赖。DS的工程师团队在用英伟达体系训练的时候,在HPC的论文中可以说是边用边骂,但一边骂一边还得用……开源这个项目,事实上是剥离了IB的不少复杂功能依赖,尤其是流量隔离和拥塞控制之类的。这也是工程师思维,花里胡哨的功能依赖是需要简化的。7 _" i1 r0 B9 |; G' R9 g% w8 k

    8 F7 M/ A6 O& c. S% m3 O- n流量隔离$ D2 b/ p* I5 I9 t
    & g* f. z7 V) t# S) ^& u
    DeepEP 利用 InfiniBand 的虚拟通道 (VL) 特性,实现不同类型流量的隔离。通过将不同工作负载(如使用普通内核的工作负载、使用低延迟内核的工作负载,以及其他工作负载)分配到不同的虚拟通道,可以有效避免相互干扰,提高整体网络性能。DeepEP 通过环境变量 NVSHMEM_IB_SL 来控制虚拟通道的分配。 这种精细化的流量管理,可以确保关键任务(如 MoE 训练)获得足够的带宽资源,不受其他任务的影响。: y; C3 `4 ]9 Q' Q& A. M; D; w" G+ G
    ( ]9 A; L1 R2 X4 l1 a$ n
    "为了防止不同类型的流量之间发生干扰,我们建议在不同的虚拟通道之间划分工作负载,如下所示: * 使用普通内核的工作负载 * 使用低延迟内核的工作负载 * 其他工作量 对于 DeepEP,您可以通过设置环境变量来控制虚拟通道分配NVSHMEM_IB_SL。"
    3 ^' L9 A' z" a( y9 m$ V+ h1 {" u) G' W1 [7 K+ M$ {7 G& ~  t
    自适应路由
    ' s. L, q" {- b8 O* e1 f- B4 P5 ?' m; D# l9 C6 \
    自适应路由是 InfiniBand 交换机提供的一项高级功能,可以将流量动态地分布在多条路径上,从而提高网络的鲁棒性和吞吐量。DeepEP 的低延迟内核支持自适应路由,可以有效消除因路由冲突导致的网络拥塞,降低延迟。
    ) q" N) L7 O) S% h: Q, ~6 n% v4 y- ]; I) D7 A6 x
    DeepSeek 团队根据实践经验,建议在网络负载较重的环境中启用自适应路由,以获得更好的稳定性和吞吐量;而在网络负载较轻的环境中,则可以使用静态路由,以减少路由计算的开销,进一步降低延迟。
    9 a# p2 E3 b  X. V
    7 K3 h3 v2 c2 k/ G- P; Z2 N"对于低延迟内核,启用自适应路由可以完全消除路由冲突导致的网络拥塞,但也会引入额外的延迟。我们建议采用以下配置以获得最佳性能: * 在网络负载较重的环境中启用自适应路由 * 在网络负载较轻的环境中,使用静态路由"! n6 g% A/ k% W

    9 Q+ o8 l6 H+ A! a5 b: r拥塞控制
    ( P+ n: D3 ?, a
    ; J* w; D. ]4 ^6 J; MDeepSeek 团队在生产环境中没有观察到明显的网络拥塞,因此 DeepEP 目前默认禁用了拥塞控制。这反映了 DeepSeek 团队务实的态度,避免引入不必要的复杂性。当然,如果未来需要,DeepEP 也保留了启用拥塞控制的灵活性。2 x! Q; V3 J9 {/ \

    5 j8 J( n' P: F% M; O9 z总结:DeepEP 的深远意义! {0 d0 F( s. U$ M# w3 S. R" [

    8 o5 Q. v' w! a- F- NDeepEP 不仅仅是一个高性能的通信库,它更代表了一种先进的工程理念:
    * Z- T& B/ ]' ^; c* M
    , ^  _( P: W+ D' A! q6 m# u以实际问题为导向: DeepEP 的所有设计决策都紧密围绕 MoE 模型训练和推理的实际需求,不追求“炫技”,而是注重解决实际问题。1 k9 Q, ~9 |0 N7 [- M; C) \
    软硬件协同优化: DeepEP 深入底层硬件,充分利用硬件特性,实现极致的性能。
    : E7 D3 u6 x+ S# H( M7 ]2 P开放协作: DeepEP 的开源,将降低 MoE 模型的开发门槛,促进整个 AI 社区的发展。
    + j& U* g9 D4 b- J5 \$ QDeepEP 的发布,是 DeepSeek 团队的重要贡献。它将加速 MoE 模型的研究和应用,推动 AI 技术向更高效、更经济的方向发展。DeepEP 的成功,也再次证明了开放协作的力量,以及实用主义工程哲学的重要性。我们有理由相信,DeepSeek 团队未来将继续秉持这种精神,为 AI 社区带来更多的惊喜。
    ( M* Z4 h6 g/ L9 V  U' T2 k
    6 ~" z. \2 G% V0 T& w原文链接

    评分

    参与人数 7爱元 +75 学识 +2 收起 理由
    常挨揍 + 15
    唐家山 + 4
    indy + 10 谢谢分享
    黑洞的颜色 + 10 涨姿势
    老票 + 2 涨姿势

    查看全部评分

  • TA的每日心情
    开心
    11 小时前
  • 签到天数: 3519 天

    [LV.Master]无

    沙发
    发表于 2025-2-26 22:53:09 | 只看该作者
    分布控制是必走之路,没见到这个DeepEP之前, 我就知道不管是哪家, 这个分布控制早晚的事。 我不是it的, 但是现代工业控制用的DCS就是分散/分布控制, 兼顾了速度和规模。

    评分

    参与人数 2爱元 +16 收起 理由
    xiejin77 + 6 精彩
    indy + 10

    查看全部评分

    回复 支持 反对

    使用道具 举报

  • TA的每日心情
    开心
    2020-4-8 10:45
  • 签到天数: 227 天

    [LV.7]分神

    板凳
     楼主| 发表于 2025-2-27 13:36:47 | 只看该作者
    马鹿 发表于 2025-2-26 22:539 q  p) H5 u! N' t8 T  t
    分布控制是必走之路,没见到这个DeepEP之前, 我就知道不管是哪家, 这个分布控制早晚的事。 我不是it的,  ...

    / S. W6 ]4 S$ a5 G# C马鹿老师说的大势是非常准确的。& O5 G5 n, X; g0 s- [
    ( _) \1 }7 K5 h/ q
    只是这个分布式是需要平衡代价和效用的,这几年的创新大都有分布式的影子,物联网,分布式数据库,分布式存储,分布式机器学习甚至是区块链、加密数字货币都是如此。
      R4 z9 b  s& a6 k8 v+ G5 z; N3 _: x; n! g) L1 E4 x
    但是对于大模型训推这个领域,计算密集、数据密集、通讯密集;分布式到底能做成啥,还需要进一步演化。
    回复 支持 反对

    使用道具 举报

  • TA的每日心情
    开心
    11 小时前
  • 签到天数: 3519 天

    [LV.Master]无

    地板
    发表于 2025-2-27 21:20:54 | 只看该作者
    xiejin77 发表于 2025-2-27 00:36
    9 V$ \6 C6 I- D3 h+ C2 O; s$ w) b马鹿老师说的大势是非常准确的。
    ' p$ w2 q) k$ Q. Z3 ~+ H4 X6 z  S0 k1 Q. S1 H
    只是这个分布式是需要平衡代价和效用的,这几年的创新大都有分布式的影 ...

    7 M( M. Y0 }/ `正常, 不同DCS厂家的底层也不一定是一样的(没用过所有的,但是到目前我没见过一样的。。。), 而且一直在发展, 至少前几年我看到某家的DCS又进步了。。。而且随着芯片能力的增强,以前写代码要特别注意的事(比如执行时间), 现在都不用那么抠门地考虑了。
    " |' L) ^" }4 b  T- m( ]. t. n1 A, T, o& p  C9 R
    回复 支持 反对

    使用道具 举报

    手机版|小黑屋|Archiver|网站错误报告|爱吱声   

    GMT+8, 2025-7-23 11:11 , Processed in 0.040339 second(s), 22 queries , Gzip On.

    Powered by Discuz! X3.2

    © 2001-2013 Comsenz Inc.

    快速回复 返回顶部 返回列表