DeepSeek開(kāi)源第二彈，為MoE和EP量身定制的通信庫(kù)！暫和英偉達(dá)顯卡綁定

衡宇 2025-02-25 11:31:49 來(lái)源：量子位

衡宇 發(fā)自 凹非寺

量子位 | 公眾號(hào) QbitAI

好消息如約而至，DeepSeek開(kāi)源周第二彈來(lái)了！

DeepEP， 第一個(gè)用于MoE模型訓(xùn)練和推理的開(kāi)源EP通信庫(kù)（expert parallelism，專(zhuān)家并行）。

它提供高吞吐量和低延遲的all-to-all GPU內(nèi)核，也稱(chēng)為MoE dispatch和combine。

該庫(kù)還支持低精度運(yùn)算，包括FP8。

同時(shí)按慣例，開(kāi)源協(xié)議用的是最為寬松的MIT。

今天的DeepSeek選擇了先在GitHub上線(xiàn)，然后再在官推發(fā)上新通知。

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

DeepSeek開(kāi)源列車(chē)永不停止。

DeepEP性能如何？

DeepSeek官推對(duì)DeepEP進(jìn)行了要素提煉：

• 高效和優(yōu)化的all-to-all通信
• NVLink和RDMA的節(jié)點(diǎn)內(nèi)和節(jié)點(diǎn)間支持
• 用于訓(xùn)練和推理預(yù)填充的高吞吐量?jī)?nèi)核
• 用于推理解碼的低延遲內(nèi)核
• 原生FP8調(diào)度支持
• 靈活的GPU資源控制，用于計(jì)算通信重疊

我們先來(lái)看看性能方面的兩個(gè)重點(diǎn)。

（注：DeepEP中的實(shí)現(xiàn)可能與DeepSeek-V3論文有一些細(xì)微的差異）

具有NVLink和RDMA轉(zhuǎn)發(fā)的普通內(nèi)核

為了與DeepSeek-V3論文中提出的組限制門(mén)控算法保持一致，DeepEP提供了一組針對(duì)非對(duì)稱(chēng)域帶寬轉(zhuǎn)發(fā)進(jìn)行了優(yōu)化的內(nèi)核，例如將數(shù)據(jù)從NVLink域轉(zhuǎn)發(fā)到RDMA域。

這些內(nèi)核提供高吞吐量，使其適用于訓(xùn)練和推理預(yù)填充任務(wù)。

此外，它們還支持SM（Streaming Multiprocessors）號(hào)碼控制。

DeepEP團(tuán)隊(duì)在在H800（~160 GB/s NVLink最大帶寬）上測(cè)試普通內(nèi)核，每個(gè)內(nèi)核都連接到CX7 InfiniBand 400 Gb/s RDMA網(wǎng)卡（~50 GB/s 最大帶寬）。

且遵循DeepSeek-V3/R1預(yù)訓(xùn)練設(shè)置（每批4096個(gè)tokens，隱藏7168個(gè)，前4組，前8個(gè)專(zhuān)家，F(xiàn)P8調(diào)度和BF16組合）。

具有純RDMA的低延遲內(nèi)核

針對(duì)延遲敏感型推理解碼場(chǎng)景，DeepEP包括一組具有純RDMA的低延遲內(nèi)核，以最大限度地減少延遲。

該庫(kù)還引入了一種基于hook的通信計(jì)算重疊方法，不占用任何SM資源。

DeepEP團(tuán)隊(duì)在H800上測(cè)試低延遲內(nèi)核，每個(gè)內(nèi)核都連接到CX7 InfiniBand 400 Gb/s RDMA 網(wǎng)卡（~50 GB/s 最大帶寬）。

且遵循典型的DeepSeek-V3/R1生產(chǎn)設(shè)置（每批128個(gè)tokens，7168個(gè)隱藏，前8個(gè)專(zhuān)家，F(xiàn)P8調(diào)度和BF16組合）。

暫不支持消費(fèi)級(jí)顯卡，建議使用最佳自動(dòng)優(yōu)化配置

在GitHub上，DeepSeek團(tuán)隊(duì)明確寫(xiě)出了關(guān)于DeepEP的使用方式，涵蓋各種適配環(huán)境、配置要求等。

首先是DeepEP需要的軟硬件環(huán)境版本：

• Hopper GPUs（以后可能支持更多架構(gòu)或設(shè)備）
• Python 3.8及更高版本
• CUDA 12.3及更高版本
• PyTorch 2.1及更高版本
• 用于節(jié)點(diǎn)內(nèi)通信的NVLink
• 用于節(jié)點(diǎn)內(nèi)通信的RDMA網(wǎng)絡(luò)

其次，使用DeepEP需要下載并安裝團(tuán)隊(duì)修改后的NVSHMEM依賴(lài)項(xiàng)（有關(guān)說(shuō)明，請(qǐng)參閱DeepSeek團(tuán)隊(duì)的NVSHMEM安裝指南）。

然后，將 deep_ep 導(dǎo)入到Python項(xiàng)目中，就開(kāi)始“盡情享受吧”！

至于網(wǎng)絡(luò)配置方面，DeepEP已通過(guò)InfiniBand網(wǎng)絡(luò)的全面測(cè)試。

但理論上，它也與基于融合以太網(wǎng)的RDMA（RoCE）兼容。

其中，InfiniBand通過(guò)虛擬通道（Virtual Lanes， VL）支持流量隔離。

為了防止不同類(lèi)型流量之間的干擾，DeepEP圖男隊(duì)建議將工作負(fù)載隔離到不同的虛擬通道中，如下所示：

• 使用普通內(nèi)核的工作負(fù)載
• 使用低延遲內(nèi)核的工作負(fù)載
• 其它工作負(fù)載

對(duì)于DeepEP，開(kāi)發(fā)者可以通過(guò)設(shè)置 NVSHMEM_IB_SL 環(huán)境變量來(lái)控制虛擬通道分配。

值得注意的是，自適應(yīng)路由是InfiniBand交換機(jī)提供的一項(xiàng)高級(jí)路由功能，可以在多個(gè)路徑之間均勻分配流量。

目前，低延遲內(nèi)核支持Adaptive Routing，而普通內(nèi)核不支持（可能很快就會(huì)添加支持）。

為普通的節(jié)點(diǎn)間內(nèi)核啟用自適應(yīng)路由，可能會(huì)導(dǎo)致死鎖或數(shù)據(jù)損壞問(wèn)題。

對(duì)于低延遲內(nèi)核，啟用Adaptive routing可以完全消除路由沖突導(dǎo)致的網(wǎng)絡(luò)擁塞，但也會(huì)帶來(lái)額外的延遲。

DeepEP團(tuán)隊(duì)建議使用以下配置以獲得最佳性能：

• 在網(wǎng)絡(luò)負(fù)載較重的環(huán)境中啟用自適應(yīng)路由
• 在網(wǎng)絡(luò)負(fù)載較輕的環(huán)境中使用靜態(tài)路由

BTW，DeepEP已禁用擁塞控制（Congestion control），因?yàn)閳F(tuán)隊(duì)在生產(chǎn)環(huán)境中沒(méi)有觀察到明顯的擁塞。

最后一點(diǎn)來(lái)自DeepEP團(tuán)隊(duì)的叮囑——

為了獲得極致性能，團(tuán)隊(duì)發(fā)現(xiàn)并使用了一條out-of-doc PTX指令ld.global.nc.L1::no_allocate.L2::256B 。

此指令將導(dǎo)致未定義的行為：使用非相干只讀PTX修飾符 .nc 訪問(wèn)易失性GPU內(nèi)存。

但是，正確性已經(jīng)過(guò)測(cè)試，以保證 。L1：：no_allocate 在 Hopper 架構(gòu)上，性能會(huì)好得多。

如果您發(fā)現(xiàn)內(nèi)核在某些其他平臺(tái)上無(wú)法運(yùn)行，您可以添加到DISABLE_AGGRESSIVE_PTX_INSTRS=1 setup.py并禁用此功能，或提交問(wèn)題。

為了在集群上獲得更好的性能，DeepSeek建議運(yùn)行所有測(cè)試并使用最佳的自動(dòng)優(yōu)化配置。

因?yàn)槟J(rèn)配置在DeepSeek的內(nèi)部集群上進(jìn)行了優(yōu)化～

One More Thing

DeepSeek為了本次開(kāi)源周專(zhuān)門(mén)在GitHub上新開(kāi)了一個(gè)庫(kù)：

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

根據(jù)這兩天的發(fā)布，猜測(cè)本次開(kāi)源周發(fā)布內(nèi)容maybe均與AI Infra有關(guān)。

不過(guò)一個(gè)不那么好的消息，DeepSeek的開(kāi)源周更新時(shí)間，好像不太穩(wěn)定。

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

于是乎，為了追蹤DeepSeek碌碌向前的車(chē)輪，量子位誠(chéng)邀大家一起第一時(shí)間上車(chē)（doge）。

掃碼備注「DeepSeek-職業(yè)/姓名」加入群聊，一起追蹤DeepSeek開(kāi)源周！

DeepEP GitHub：

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