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

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

今天 DeepSeek 繼續(xù)開源底層架構(gòu)的創(chuàng)新，今天開源的項目是首個用于 MoE 模型訓(xùn)練和推理的 EP 通信庫 DeepEP。

在分布式系統(tǒng)中（如多 GPU 訓(xùn)練環(huán)境），所有處理單元之間需要高效地傳遞數(shù)據(jù)。在 MoE 中，這點尤為重要，因為不同「專家」需要頻繁交換信息。并且 MoE 模型容易在「專家并行」中出現(xiàn)負載不均衡，導(dǎo)致每個「專家」分到的算力不均，不重要的「專家」難以發(fā)揮應(yīng)有的性能。

此次開源的 DeepEP 做到了：

1. 高效優(yōu)化的 All-to-All 通信

2. 支持 NVLink 和 RDMA 的節(jié)點內(nèi) / 跨節(jié)點通信

3. 訓(xùn)練及推理預(yù)填充階段的高吞吐量計算核心

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

5. 原生支持 FP8 數(shù)據(jù)分發(fā)

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

高效通信減少了數(shù)據(jù)傳輸?shù)钠款i，計算核心的優(yōu)化提升了處理速度，靈活的資源調(diào)度讓計算和通信不互相等待。

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

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

至于火到了什么程度？

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

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

還有人說，「NVLink 和 RDMA 支持對大規(guī)模 MoE 模型來說是革命性的突破?？磥?DeepSeek 再次在 AI 基礎(chǔ)設(shè)施的可能性方面推動了技術(shù)邊界?！?/span>

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

DeepEP 是什么？

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

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

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

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

DeepEP 性能如何？

具有 NVLink 和 RDMA forwarding 的常規(guī)內(nèi)核

DeepSeek 在 H800 上測試常規(guī)內(nèi)核（NVLink 最大帶寬約 160 GB/s），每個 H800 連接到一個 CX7 InfiniBand 400 Gb/s RDMA 網(wǎng)卡（最大帶寬約 50 GB/s）。他們遵循 DeepSeek-V3/R1 預(yù)訓(xùn)練設(shè)置（每批次 4096 個 token，7168 隱藏維度，top-4 組，top-8 專家，F(xiàn)P8 分發(fā)和 BF16 合并）。

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

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

注意事項

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

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

結(jié)尾必須再強調(diào)一句：Real OPENAI has born!