DeepSeek OpenSourceWeek第二弹:MoE训练模型+EP通信库DeepEP

edwin99
edwin99
2025-02-25 07:38
28 阅读
0 评论
文章封面
目录
正在加载目录...
项目链接:https://github.com/deepseek-ai/DeepEP
 

DeepSeek 发推说本周将是开源周(OpenSourceWeek),并将连续开源五个软件库。DeepSeek 开源底层架构的创新,项目是首个用于 MoE 模型训练和推理的 EP 通信库 DeepEP。在分布式系统中(如多 GPU 训练环境),所有处理单元之间需要高效地传递数据。在 MoE 中,这点尤为重要,因为不同「专家」需要频繁交换信息。并且 MoE 模型容易在「专家并行」中出现负载不均衡,导致每个「专家」分到的算力不均,不重要的「专家」难以发挥应有的性能。

此次开源的 DeepEP 做到了:

1. 高效优化的 All-to-All 通信

2. 支持 NVLink 和 RDMA 的节点内 / 跨节点通信

3. 训练及推理预填充阶段的高吞吐量计算核心

4. 推理解码阶段的低延迟计算核心

5. 原生支持 FP8 数据分发

6. 灵活控制 GPU 资源,实现计算与通信的高效重叠

高效通信减少了数据传输的瓶颈,计算核心的优化提升了处理速度,灵活的资源调度让计算和通信不互相等待。

 

DeepEP定义

DeepEP 是一个专为混合专家系统(MoE)和专家并行(EP)定制的通信库。它提供高吞吐量和低延迟的 all-to-all GPU 内核, 这些内核也被称为 MoE 分发和合并。该库还支持低精度操作,包括 FP8。

为了与 DeepSeek-V3 论文中提出的 group-limited gating 算法保持一致,DeepEP 提供了一套针对非对称域带宽 forwarding 进行优化的内核,例如从 NVLink 域到 RDMA 域的数据 forwarding。这些内核提供高吞吐量,适用于训练和推理预填充(prefilling)任务。此外,它们还支持 SM(流式多处理器,Streaming Multiprocessors)数量控制。

对于对延迟敏感的推理解码,DeepEP 包含一套使用纯 RDMA 的低延迟内核,以最小化延迟。该库还引入了一种 hook-based 的通信 - 计算重叠方法,不占用任何 SM 资源。注意:本库中的实现可能与 DeepSeek-V3 论文有一些细微差异。
 

DeepEP性能测试

具有 NVLink 和 RDMA forwarding 的常规内核:DeepSeek 在 H800 上测试常规内核(NVLink 最大带宽约 160 GB/s),每个 H800 连接到一个 CX7 InfiniBand 400 Gb/s RDMA 网卡(最大带宽约 50 GB/s)。他们遵循 DeepSeek-V3/R1 预训练设置(每批次 4096 个 token,7168 隐藏维度,top-4 组,top-8 专家,FP8 分发和 BF16 合并)。


具有纯 RDMA 的低延迟内核:DeepSeek 在 H800 上测试低延迟内核,每个 H800 连接到一个 CX7 InfiniBand 400 Gb/s RDMA 网卡(最大带宽约 50 GB/s)。他们遵循典型的 DeepSeek-V3/R1 生产设置(每批次 128 个 token,7168 隐藏维度,top-8 专家,FP8 分发和 BF16 合并)。

 

注意事项

为了极致性能,DeepSeek 发现并使用了一个未记录在文档中的 PTX 指令:ld.global.nc.L1::no_allocate.L2::256B。这个指令会导致一个未定义的行为:使用非一致性只读 PTX 修饰符「.nc」访问易变的 GPU 内存。但在 Hopper 架构上,通过「.L1::no_allocate」已测试确保了正确性,且性能会大幅提升。如果你发现内核在某些其他平台上不 work,你可以在 setup.py 中添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 来禁用此功能,或提交 issue。为了在集群上获得更好的性能,DeepSeek 建议运行所有测试并使用最佳的自动调优配置。默认配置是针对 DeepSeek 内部集群优化的。

 
 

 

评论区 (0)

登录后参与评论

暂无评论,抢沙发吧!