在大模型时代，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 nodespython tests/test_intranode.pypython tests/test_internode.pypython tests/test_low_latency.pyInstall
-
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 的开源震撼，让我们共同期待更多激动人心的创新成果！
