好消息如约而至，DeepSeek 开源周第二弹来了！

DeepEP， 第一个用于 MoE 模型训练和推理的开源 EP 通信库（expert parallelism，专家并行）。

它提供高吞吐量和低延迟的 all-to-all GPU 内核，也称为 MoE dispatch 和 combine。

该库还支持低精度运算，包括 FP8。

同时按惯例，开源协议用的是最为宽松的 MIT。

今天的 DeepSeek选择了先在 GitHub 上线，然后再在官推发上新通知。

不出所料，底下一片叫好：

DeepSeek 开源列车永不停止。

DeepEP 性能如何？

DeepSeek 官推对 DeepEP 进行了要素提炼：

高效和优化的 all-to-all 通信

NVLink 和 RDMA 的节点内和节点间支持

用于训练和推理预填充的高吞吐量内核

用于推理解码的低延迟内核

原生 FP8 调度支持

灵活的 GPU 资源控制，用于计算通信重叠

我们先来看看性能方面的两个重点。

（注：DeepEP 中的实现可能与 DeepSeek-V3 论文有一些细微的差异）

具有 NVLink 和 RDMA 转发的普通内核

为了与 DeepSeek-V3 论文中提出的组限制门控算法保持一致，DeepEP 提供了一组针对非对称域带宽转发进行了优化的内核，例如将数据从 NVLink 域转发到 RDMA 域。

这些内核提供高吞吐量，使其适用于训练和推理预填充任务。

此外，它们还支持 SM（Streaming Multiprocessors）号码控制。

DeepEP 团队在在 H800（~160 GB/s NVLink 最大带宽）上测试普通内核，每个内核都连接到 CX7 InfiniBand 400 Gb/s RDMA 网卡（~50 GB/s 最大带宽）。

且遵循 DeepSeek-V3/R1 预训练设置（每批 4096 个 tokens，隐藏 7168 个，前 4 组，前 8 个专家，FP8 调度和 BF16 组合）。

具有纯 RDMA 的低延迟内核

针对延迟敏感型推理解码场景，DeepEP 包括一组具有纯 RDMA 的低延迟内核，以最大限度地减少延迟。

该库还引入了一种基于 hook 的通信计算重叠方法，不占用任何 SM 资源。

DeepEP 团队在 H800 上测试低延迟内核，每个内核都连接到 CX7 InfiniBand 400 Gb/s RDMA 网卡（~50 GB/s 最大带宽）。

且遵循典型的 DeepSeek-V3/R1 生产设置（每批 128 个 tokens，7168 个隐藏，前 8 个专家，FP8 调度和 BF16 组合）。

暂不支持消费级显卡，建议使用最佳自动优化配置

在 GitHub 上，DeepSeek 团队明确写出了关于 DeepEP 的使用方式，涵盖各种适配环境、配置要求等。

首先是 DeepEP 需要的软硬件环境版本：

Hopper GPUs（以后可能支持更多架构或设备）

Python 3.8 及更高版本

CUDA 12.3 及更高版本

PyTorch 2.1 及更高版本

用于节点内通信的 NVLink

用于节点内通信的 RDMA 网络

然后，将 deep_ep 导入到 Python 项目中，就开始"尽情享受吧"！

至于网络配置方面，DeepEP 已通过 InfiniBand 网络的全面测试。

但理论上，它也与基于融合以太网的 RDMA（RoCE）兼容。

其中，InfiniBand 通过虚拟通道（Virtual Lanes， VL）支持流量隔离。

为了防止不同类型流量之间的干扰，DeepEP 图男队建议将工作负载隔离到不同的虚拟通道中，如下所示：

使用普通内核的工作负载

使用低延迟内核的工作负载

其它工作负载

对于 DeepEP，开发者可以通过设置 NVSHMEM_IB_SL 环境变量来控制虚拟通道分配。

值得注意的是，自适应路由是 InfiniBand 交换机提供的一项高级路由功能，可以在多个路径之间均匀分配流量。

目前，低延迟内核支持 Adaptive Routing，而普通内核不支持（可能很快就会添加支持）。

为普通的节点间内核启用自适应路由，可能会导致死锁或数据损坏问题。

对于低延迟内核，启用 Adaptive routing 可以完全消除路由冲突导致的网络拥塞，但也会带来额外的延迟。

DeepEP 团队建议使用以下配置以获得最佳性能：

在网络负载较重的环境中启用自适应路由

在网络负载较轻的环境中使用静态路由

BTW，DeepEP已禁用拥塞控制（Congestion control），因为团队在生产环境中没有观察到明显的拥塞。

最后一点来自 DeepEP 团队的叮嘱——

为了获得极致性能，团队发现并使用了一条 out-of-doc PTX 指令 ld.global.nc.L1::no_allocate.L2::256B 。

此指令将导致未定义的行为：使用非相干只读 PTX 修饰符 .nc 访问易失性 GPU 内存。

但是，正确性已经过测试，以保证 。L1：：no_allocate 在 Hopper 架构上，性能会好得多。

如果您发现内核在某些其他平台上无法运行，您可以添加到 DISABLE_AGGRESSIVE_PTX_INSTRS=1 setup.py 并禁用此功能，或提交问题。

为了在集群上获得更好的性能，DeepSeek 建议运行所有测试并使用最佳的自动优化配置。

因为默认配置在 DeepSeek 的内部集群上进行了优化～

One More Thing

DeepSeek 为了本次开源周专门在 GitHub 上新开了一个库：

https://github.com/deepseek-ai/open-infra-index

根据这两天的发布，猜测本次开源周发布内容 maybe 均与 AI Infra 有关。

不过一个不那么好的消息，DeepSeek 的开源周更新时间，好像不太稳定。

昨天是上午 9:34，今天是 10:24，明天……

于是乎，为了追踪 DeepSeek 碌碌向前的车轮，量子位诚邀大家一起第一时间上车（doge）。

扫码备注「DeepSeek- 职业 / 姓名」加入群聊，一起追踪 DeepSeek 开源周！

DeepEP GitHub：

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