今天小编分享的科学经验：DeepSeek开源第二弹，为MoE和EP量身定制的通信库！暂和英伟达显卡绑定，欢迎阅读。
好消息如约而至，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
