上周五，DeepSeek 发推说本周将是开源周（OpenSourceWeek），并将连续开源五个软件库。
昨天，他们开源了第一个代码库 ——FlashMLA。这是一款用于 Hopper GPU 的高效型 MLA 解码核，仅用了 24 小时就达到了接近 8k 的 star 量（详情请参见《刚刚，DeepSeek 开源 FlashMLA，推理加速核心技术，Star 量飞涨中》）。
今天 DeepSeek 继续开源底层架构的创新，今天开源的项目是首个用于 MoE 模型训练和推理的 EP 通信库 DeepEP。
在分布式系统中（如多 GPU 训练环境），所有处理单元之间需要高效地传递数据。在 MoE 中，这点尤为重要，因为不同「专家」需要频繁交换信息。并且 MoE 模型容易在「专家并行」中出现负载不均衡，导致每个「专家」分到的算力不均，不重要的「专家」难以发挥应有的性能。
此次开源的 DeepEP 做到了：
1. 高效优化的 All-to-All 通信
2. 支持 NVLink 和 RDMA 的节点内 / 跨节点通信
3. 训练及推理预填充阶段的高吞吐量计算核心
4. 推理解码阶段的低延迟计算核心
5. 原生支持 FP8 数据分发
6. 灵活控制 GPU 资源，实现计算与通信的高效重叠
高效通信减少了数据传输的瓶颈，计算核心的优化提升了处理速度，灵活的资源调度让计算和通信不互相等待。
MLA 和 MoE 架构改进可以说是 DeepSeek 的两大重要创新点。昨天是对 MLA 解码内核的优化，今天就公开了另一张王牌 MoE 如何高效通信和并行处理，DeepSeek 可真是太 Open 了！
项目链接：https://github.com/deepseek-ai/DeepEP
至于火到了什么程度？
机器之心文章还没写完，DeepEP 的 Star 量已超 1000 了：
该项目开源后，有人评价说：DeepSeek 为 MoE 模型所达到的优化水平令人印象深刻，这类模型因其规模和复杂性而充满挑战性。DeepEP 能够利用 NVLink 和 RDMA 等尖端硬件技术，并支持 fp8 精度，以如此精确的方式处理这些挑战，简直是突破性的成就。
还有人说，「NVLink 和 RDMA 支持对大规模 MoE 模型来说是革命性的突破。看来 DeepSeek 再次在 AI 基础设施的可能性方面推动了技术边界。」
之前，有人曾质疑 DeepSeek-R1 只是通过模型蒸馏来实现其性能，而非真正的技术创新。还有人怀疑 DeepSeek 低报了训练所需的 GPU 数量。开源周发布的这些内容可以从某些角度证明，DeepSeek 确实通过技术创新实现了真正的训练效率提升和成本降低。
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 内部集群优化的。
更多信息请参见 GitHub 代码库。
结尾必须再强调一句：Real OPENAI has born!
