刚刚，DeepSeek开源MoE训练、推理EP通信库DeepEP，真太Open了！

上周五，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 了！

• 为了极致性能，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 内部集群优化的。