DeepSeek 开源狂欢周（二）DeepEP深度技术解析 | 解锁 MoE 模型并行加速

在大模型时代，Mixture-of-Experts (MoE) 模型凭借其强大的容量和高效的计算能力，成为研究和应用的热点。然而，MoE 模型的训练和推理面临着巨大的专家并行通信挑战。近日，DeepSeek 开源了 DeepEP 项目，为解决这一难题提供了强大的武器。OpenCSG社区将深入剖析 DeepEP 的技术原理、架构设计和性能表现，帮助开发者更好地理解和使用 DeepEP，加速 MoE 模型的落地应用。

DeepSeek DeepEP

DeepEP 是一个专为专家混合（MoE）和专家并行（EP）量身定制的通信库。它提供高吞吐量和低延迟的全到全 GPU 内核，也称为 MoE 调度和组合。该库还支持低精度操作，包括 FP8。

为了与 DeepSeek-V3 论文中提出的组限制门控算法对齐，DeepEP 提供了一组针对非对称域带宽转发优化的内核，例如将数据从 NVLink 域转发到 RDMA 域。这些内核提供高吞吐量，使其适用于训练和推理预填充任务。此外，它们支持 SM（流式多处理器）数量控制。

• 提速！ 🚀 提供高吞吐量、低延迟的 All-to-All GPU 内核，让 MoE 调度和组合更快更强！

• 优化！ ⚙️ 针对非对称域带宽转发进行优化（NVLink 到 RDMA），充分利用硬件性能！

• 精简！ 🔢 原生支持 FP8 等低精度操作，降低计算和通信开销！

• 灵活！ 🧰 支持 SM 数量控制，让资源分配更智能！

技术亮点抢先看：

• 与 DeepSeek-V3 论文对齐： 采用组限制门控算法，优化非对称域带宽转发。

• 训练推理全覆盖： 高吞吐内核适用于训练和推理预填充任务。

• 极低延迟： 针对延迟敏感的推理解码，采用纯 RDMA 技术。

• 资源友好： 基于钩子的通信-计算重叠方法，不占用任何 SM 资源。

重要提示：

DeepEP 库中的实现可能与 DeepSeek-V3 论文存在一些细微差异，使用前请仔细阅读官方文档哦！

MoE 与 EP：分布式训练的挑战 🤯

MoE 模型，通过引入专家模型和路由机制，能够扩展模型规模，提升性能。 然而，MoE 模型的训练和推理对通信提出了更高的要求。 在分布式训练环境中，尤其是 EP 并行模式下，数据需要在不同 GPU 之间进行频繁交换。

• 核心问题：

  • All-to-All 通信： 不同“专家”之间的数据交换，需要高效的 All-to-All 通信。

  • 负载不均衡： MoE 模型容易出现负载不均衡，导致部分“专家”的算力无法充分发挥。

DeepEP 的目标，正是解决这些问题，提高 EP 通信的效率和可靠性。

DeepEP 的核心特性：高性能 EP 通信 🚀

• DeepEP 提供了以下关键特性，助力 MoE 模型实现高性能 EP 通信：

• 高效优化的 All-to-All 通信： DeepEP 提供了优化的 All-to-All 通信内核，针对 MoE 模型的特点进行了专门设计。

• 节点内 / 跨节点通信支持： 灵活支持 NVLink 和 RDMA 的节点内和跨节点通信，充分利用硬件互连优势。

• 训练 / 推理预填充阶段高吞吐量计算核心： 针对训练和推理的预填充阶段，DeepEP 提供了高吞吐量计算核心，加速数据处理。

• 推理解码阶段低延迟计算核心： 对于延迟敏感的推理解码任务，DeepEP 提供了低延迟计算核心，降低推理时间。

• 原生 FP8 数据分发： 支持 FP8 数据类型，降低计算和通信开销，提高效率。

• 灵活的 GPU 资源控制： 可以灵活控制 GPU 资源，实现计算与通信的高效重叠，避免资源浪费。

技术内幕：DeepEP 的关键技术点 🧐

• 组限制门控算法与非对称域带宽优化：

  • 为了与 DeepSeek-V3 论文中提出的组限制门控算法对齐，DeepEP 针对非对称域带宽进行了优化。

  • DeepEP 针对 NVLink 域到 RDMA 域的数据转发进行了优化，最大化利用 NVLink 和 RDMA 的优势。

    2. 性能优化的内核：

  • 提供了高吞吐量内核，适用于训练和推理预填充任务，并支持 SM 数量控制。

  • 提供了低延迟内核，采用纯 RDMA，降低了延迟。

    3. 通信-计算重叠：

  • DeepEP 引入基于钩子的通信-计算重叠方法，不占用 SM 资源。

DeepEP 的性能表现：数据说话 📊

DeepSeek 在 H800 上进行了性能测试：

常规内核 (NVLink)： DeepEP 在 H800 上使用常规内核，NVLink 最大带宽约 160 GB/s，每个 H800 连接一个 CX7 InfiniBand 400 Gb/s RDMA 网卡（最大带宽约 50 GB/s）。

测试结果如下：

Type	Dispatch #EP	Bottleneck bandwidth	Combine #EP	Bottleneck bandwidth
Intranode	8	153 GB/s (NVLink)	8	158 GB/s (NVLink)
Internode	16	43 GB/s (RDMA)	16	43 GB/s (RDMA)
Internode	32	44 GB/s (RDMA)	32	47 GB/s (RDMA)
Internode	64	46 GB/s (RDMA)	64	45 GB/s (RDMA)

纯 RDMA 的低延迟内核

在 H800 上测试低延迟内核，每个内核连接到一个 CX7 InfiniBand 400 Gb/s RDMA 网络卡（最大带宽约为 50 GB/s）。遵循典型的 DeepSeek-V3/R1 生产设置（每批 128 个标记，7168 个隐藏层，前 8 个专家，FP8 调度和 BF16 组合）。

测试结果如下：

Dispatch #EP	Latency	RDMA bandwidth	Combine #EP	Latency	RDMA bandwidth
8	163 us	46 GB/s	8	318 us	46 GB/s
16	173 us	43 GB/s	16	329 us	44 GB/s
32	182 us	41 GB/s	32	350 us	41 GB/s
64	186 us	40 GB/s	64	353 us	41 GB/s
128	192 us	39 GB/s	128	369 us	39 GB/s
256	194 us	39 GB/s	256	360 us	40 GB/s

注意事项与最佳实践 ⚠️

• PTX 指令： DeepEP 中使用了未记录在文档中的 PTX 指令: ld.global.nc.L1::no_allocate.L2::256B。 请仔细阅读官方文档，并在 setup.py 中添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 来禁用此功能， 如果遇到问题，请积极提交 issue。

  自动调优： 为了在您的集群上获得最佳性能，DeepSeek 建议运行所有测试并使用最佳的自动调优配置。 默认配置是针对 DeepSeek 内部集群优化的。

• 快速上手 DeepEP

  使用 DeepEP 需要 Hopper GPU、Python 3.8 以上、CUDA 12.3 以上、PyTorch 2.1 以上，以及用于内节点通信的 NVLink 和跨节点通信的 RDMA 网络。库依赖于一个修改版的 NVSHMEM，安装前需先配置此依赖。DeepEP 提供了清晰的 Python API，使开发者能轻松地将其集成到现有 MoE 模型中。对于推理场景，DeepEP 的低延迟模式特别有价值，提供不占用 SM 资源的通信-计算重叠功能。

  Development

• # Build and make symbolic links for SO filesNVSHMEM_DIR=/path/to/installed/nvshmem python setup.py build# You may modify the specific SO names according to your own platformln -s build/lib.linux-x86_64-cpython-38/deep_ep_cpp.cpython-38-x86_64-linux-gnu.so# Run test cases# NOTES: you may modify the `init_dist` function in `tests/utils.py`# according to your own cluster settings, and launch into multiple nodes python tests/test_intranode.pypython tests/test_internode.pypython tests/test_low_latency.py

  Install

• NVSHMEM_DIR=/path/to/installed/nvshmem python setup.py install

OpenCSG 社区开源加速计划

作为OpenCSG社区的一部分，我们一直致力于为开发者提供优质的开源资源。此次DeepSeek的DeepEP项目已同步到OpenCSG社区，欢迎大家访问并使用该项目。

DeepEP项目原始GitHub地址：

https://github.com/deepseek-ai/DeepEP

OpenCSG社区同步的DeepEP项目地址：

https://opencsg.com/codes/deepseek-ai/DeepEP

如果您遇到网络问题无法快速访问GitHub，可以通过我们的服务轻松同步该项目，确保不受网络限制影响。

OpenCSG为您提供了DeepSeek R1和V3系列模型的万兆网络高速下载服务，帮助您快速获取所需模型，避免因文件过大造成下载困难。

DeepSeek R1下载：

https://opencsg.com/models/DeepseekAI/DeepSeek-R1 

DeepSeek V3下载：

https://opencsg.com/models/deepseek-ai/DeepSeek-V3

同时，我们还提供了各种蒸馏版、量化版，您可以访问我们的awesome DeepSeek合集来找到最适合的模型版本。

awesome-deepseek-r1-collection：

https://opencsg.com/collections/85/ 

awesome-deepseek-v3-collection：

https://opencsg.com/collections/86/ 

awesome-deepseek-Janus-collection：

https://opencsg.com/collections/87/

开源狂欢 继续期待

DeepEP 不仅仅是一个技术工具，更是 DeepSeek 对开源社区的诚意回馈。作为全球首个开源的 EP 通信库，DeepEP 已经在 DeepSeek 内部经历了生产级超大规模并发的严苛考验，其性能和质量毋庸置疑！

长期以来，广大企业和社区用户都苦于缺乏 EP 通信库的开源实现，而传统 DP/TP 在大规模推理 MoE 场景下又存在性能瓶颈。DeepEP 的横空出世，有望造福无力自行研发 EP 技术的厂商，加速 MoE 模型在各行各业的落地。

更重要的是，随着 DeepEP 的普及，可以预见未来将涌现出更多低价甚至免费的 DeepSeek R1 系列模型，最终惠及广大用户，推动 AI 技术的 democratization！

OpenCSG 社区与您同行 🤝

OpenCSG 社区将继续与您一起，第一时间为您带来 DeepSeek 的开源震撼，让我们共同期待更多激动人心的创新成果！