•
【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網絡
其次，使用DeepEP需要下載並安裝團隊修改後的NVSHMEM相依性（有關說明，請參閱DeepSeek團隊的NVSHMEM安裝指南）。
然後，將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，明天…（量子位元）
DeepEP GitHub：https://github.com/deepseek-ai/DeepEP