上周五，DeepSeek 發推說本周將是開源周（OpenSourceWeek），并將連續開源五個軟件庫。

昨天，他們開源了第一個代碼庫 ——FlashMLA。這是一款用于 Hopper GPU 的高效型 MLA 解碼核，僅用了 24 小時就達到了接近 8k 的 star 量（詳情請參見《剛剛，DeepSeek 開源 FlashMLA，推理加速核心技術，Star 量飛漲中》）。

今天 DeepSeek 繼續開源底層架構的創新，今天開源的項目是首個用于 MoE 模型訓練和推理的 EP 通信庫 DeepEP。

在分布式系統中（如多 GPU 訓練環境），所有處理單元之間需要高效地傳遞數據。在 MoE 中，這點尤為重要，因為不同「專家」需要頻繁交換信息。并且 MoE 模型容易在「專家并行」中出現負載不均衡，導致每個「專家」分到的算力不均，不重要的「專家」難以發揮應有的性能。

此次開源的 DeepEP 做到了：

1. 高效優化的 All-to-All 通信

2. 支持 NVLink 和 RDMA 的節點內 / 跨節點通信

3. 訓練及推理預填充階段的高吞吐量計算核心

4. 推理解碼階段的低延遲計算核心

5. 原生支持 FP8 數據分發

6. 靈活控制 GPU 資源，實現計算與通信的高效重疊

高效通信減少了數據傳輸的瓶頸，計算核心的優化提升了處理速度，靈活的資源調度讓計算和通信不互相等待。

MLA 和 MoE 架構改進可以說是 DeepSeek 的兩大重要創新點。昨天是對 MLA 解碼內核的優化，今天就公開了另一張王牌 MoE 如何高效通信和并行處理，DeepSeek 可真是太 Open 了！

項目鏈接：https://github.com/deepseek-ai/DeepEP

至于火到了什么程度？

機器之心文章還沒寫完，DeepEP 的 Star 量已超 1000 了：

該項目開源后，有人評價說：DeepSeek 為 MoE 模型所達到的優化水平令人印象深刻，這類模型因其規模和復雜性而充滿挑戰性。DeepEP 能夠利用 NVLink 和 RDMA 等尖端硬件技術，并支持 fp8 精度，以如此精確的方式處理這些挑戰，簡直是突破性的成就。

還有人說，「NVLink 和 RDMA 支持對大規模 MoE 模型來說是革命性的突破。看來 DeepSeek 再次在 AI 基礎設施的可能性方面推動了技術邊界。」

之前，有人曾質疑 DeepSeek-R1 只是通過模型蒸餾來實現其性能，而非真正的技術創新。還有人懷疑 DeepSeek 低報了訓練所需的 GPU 數量。開源周發布的這些內容可以從某些角度證明，DeepSeek 確實通過技術創新實現了真正的訓練效率提升和成本降低。 

DeepEP 是什么？

DeepEP 是一個專為混合專家系統（MoE）和專家并行（EP）定制的通信庫。它提供高吞吐量和低延遲的 all-to-all GPU 內核， 這些內核也被稱為 MoE 分發和合并。該庫還支持低精度操作，包括 FP8。

為了與 DeepSeek-V3 論文中提出的 group-limited gating 算法保持一致，DeepEP 提供了一套針對非對稱域帶寬 forwarding 進行優化的內核，例如從 NVLink 域到 RDMA 域的數據 forwarding。這些內核提供高吞吐量，適用于訓練和推理預填充（prefilling）任務。此外，它們還支持 SM（流式多處理器，Streaming Multiprocessors）數量控制。

對于對延遲敏感的推理解碼，DeepEP 包含一套使用純 RDMA 的低延遲內核，以最小化延遲。該庫還引入了一種 hook-based 的通信 - 計算重疊方法，不占用任何 SM 資源。

注意：本庫中的實現可能與 DeepSeek-V3 論文有一些細微差異。

DeepEP 性能如何？

具有 NVLink 和 RDMA forwarding 的常規內核

DeepSeek 在 H800 上測試常規內核（NVLink 最大帶寬約 160 GB/s），每個 H800 連接到一個 CX7 InfiniBand 400 Gb/s RDMA 網卡（最大帶寬約 50 GB/s）。他們遵循 DeepSeek-V3/R1 預訓練設置（每批次 4096 個 token，7168 隱藏維度，top-4 組，top-8 專家，FP8 分發和 BF16 合并）。

具有純 RDMA 的低延遲內核

DeepSeek 在 H800 上測試低延遲內核，每個 H800 連接到一個 CX7 InfiniBand 400 Gb/s RDMA 網卡（最大帶寬約 50 GB/s）。他們遵循典型的 DeepSeek-V3/R1 生產設置（每批次 128 個 token，7168 隱藏維度，top-8 專家，FP8 分發和 BF16 合并）。

注意事項

• 為了極致性能，DeepSeek 發現并使用了一個未記錄在文檔中的 PTX 指令：ld.global.nc.L1::no_allocate.L2::256B。這個指令會導致一個未定義的行為：使用非一致性只讀 PTX 修飾符「.nc」訪問易變的 GPU 內存。但在 Hopper 架構上，通過「.L1::no_allocate」已測試確保了正確性，且性能會大幅提升。如果你發現內核在某些其他平臺上不 work，你可以在 setup.py 中添加 DISABLE_AGGRESSIVE_PTX_INSTRS=1 來禁用此功能，或提交 issue。
• 為了在你的集群上獲得更好的性能，DeepSeek 建議運行所有測試并使用最佳的自動調優配置。默認配置是針對 DeepSeek 內部集群優化的。

更多信息請參見 GitHub 代碼庫。

結尾必須再強調一句：Real OPENAI has born!