今天小編分享的科學經驗: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