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

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

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

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

此次開(kāi)源的 DeepEP 做到了：

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

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

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

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

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

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

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

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

項(xiàng)目鏈接：https://github.com/deepseek-ai/DeepEP

至于火到了什么程度？

機(jī)器之心文章還沒(méi)寫(xiě)完，DeepEP 的 Star 量已超 1000 了：

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

還有人說(shuō)，「NVLink 和 RDMA 支持對(duì)大規(guī)模 MoE 模型來(lái)說(shuō)是革命性的突破。看來(lái) DeepSeek 再次在 AI 基礎(chǔ)設(shè)施的可能性方面推動(dòng)了技術(shù)邊界。」

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

DeepEP 是什么？

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

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

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

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

DeepEP 性能如何？

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

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

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

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

注意事項(xiàng)

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

更多信息請(qǐng)參見(jiàn) GitHub 代碼庫(kù)。

結(jié)尾必須再?gòu)?qiáng)調(diào)一句：Real OPENAI has born!