DeepSeek開源第二天：拉爆MoE訓練和推理

2025-05-13 22:26 夕小瑤科技說

今天，在追求效能極致的路上，DeepSeek 又祭出新神器——
專為 MoE 模型打造的 DeepEP 通訊庫。
昨天主打幹崩推理效能，今天雙管齊下，訓練和推理一起拿下，依舊是所到之處寸草不生的野蠻收割。
（PS：這還是我印象裡的以天下為公的謙謙君子版 deepseek 麼，這兩天開源直接變身不擇手段的效能狂魔）
這次開源的是——DeepEP，是首個用於 MoE 模型訓練和推理的開源 EP 通訊庫，用於訓練和推理的高吞吐量和低延遲
依舊是先說結論：
  1. 實現高效的 all-to-all 通訊
  2. 提供高吞吐（NVLink + RDMA）與低延遲（純 RDMA）兩套通訊核心，兼顧大批次訓練與即時推理場景。
  3. 支援 NVLink 和 RDMA 的節點內 / 跨節點通訊。
  4. 提供 SM 數量控制介面，可在計算與通訊之間靈活分配 GPU 資源。
  5. 整合可以重疊通訊和計算的 hook 機制，允許在解碼時後臺並行接收資料，不佔用任何 SM 資源。
  6. 原生 FP8 排程支援
這是倉庫地址：
https://github.com/deepseek-ai/DeepEP
別急，MoE、EP、NVLink + RDMA、GPU、SM 這些烏七八糟的名詞我爭取都能講地明白點。

MoE 與 EP 是什麼

專家並行(Expert Parallelism，EP)是一種平行計算的方式，主要用在混合專家模型(Mixture of Experts, MoE)中。
  • 將不同的"專家"(模型的子網路)分配到不同的計算裝置(如 GPU)上
  • 根據輸入資料的特點，選擇合適的專家來處理
  • 多個專家可以同時並行工作，提高整體計算效率
MoE 模型在每次推理時，並不是所有專家都參與計算，而是根據輸入資料的需求，只啟用一部分。
比如 DeepSeek-R1，它的實際引數大小為 671B（實際還有 14B 的 MTP 投機解碼模組），在每次推理的時候，只啟用 37B 的引數量，256 個專家裡啟用 8 個。
在這個過程中，每個裝置會根據路由規則，將自身的資料傳送到相應專家所在的裝置上，然後等待專家完成計算後再將結果返回到原裝置。
這個過程中，每兩個專家之間需要進行通訊來同步各自的計算結果，這個通訊過程被稱為All-to-All 通訊。
打個比方——
想象一個橫跨多個城市的巨型物流中心，在每一個城市就是一個節點 Node，每一個城市裡有多個倉庫（GPU），而每個倉庫裡都有物流專家負責接收包裹。
物流中心每天要處理巨量的包裹（token），每個包裹需要同時派發給指定的 8 個物流專家（top-8 experts）。
專家通常分佈在不同的倉庫（ GPU）上，我們需要依賴高速通訊來完成包裹在“物流中心”以及“專家”之間的流轉。
此時，
  • NVLink 相當於“城內高速路”專門負責同一個城市內不同倉庫的聯絡，也就是同一節點內 GPU-GPU 之間的高速通訊，頻寬可達百 GB/s 級，極大提升單機多卡之間的資料交換效率。
  • RDMA（Remote Direct Memory Access）相當於“跨城高速公路”用於跨城市通訊，也就是跨節點通訊，直接讓一臺伺服器的 GPU 與另一臺伺服器的 GPU 之間進行遠端讀寫，跳過 CPU，有效避免多層網路的額外延遲，頻寬通常在幾十 GB/s 級。
在 MoE（混合專家）模型中，專家並行（Expert Parallelism）需要不停地對 token 進行分發（dispatch）和聚合（combine）。
即每一次輸入的 token 都需要傳送給不同的專家，專家處理結束後再將各自的結果組合在一起。
這就像大型物流中心“分揀包裹”，同時要在站內（NVLink）和跨節點（RDMA）兩大“運輸通道”上來回調配。
當 Token 數量暴增時，如果通訊方案不夠高效，就會像人工分揀一樣堵成“長龍”，讓價值百萬的 H800 叢集陷入“等傳輸”的尷尬局面。
所以，DeepEP 就是讓 H800 高效通訊的解決方案。
那麼 DeepEP 具體做了哪些改動呢？
它設計了兩種核心，分別應對高吞吐和低延遲兩種場景，前者用於訓練和快速處理使用者輸入的文字，後者用於加速大模型一個字一個字生成的速度。

高吞吐核心最佳化

針對高吞吐場景，DeepEP 提供了一組同時適用於訓練和推理預填充任務的通用核心。
它可以直接將資料從 NVLink 域轉發到 RDMA 域，也就是可以直接“跨城高速公路”轉“城內高速路”，提供非常高的吞吐量。
已知 NVLink 的最高頻寬是 160GB/s，DeepEP 在實現專家並行時，實測頻寬最高可以達到 153GB/s！無限接近極限值。

低延遲核心

而在對延遲敏感的大模型解碼環節，DeepEP 提供了另外一組低延遲核心，它只使用 RDMA（遠端直接視訊記憶體訪問）以最小化延遲。
可以看到，完全依照 RDMA 後的延遲在各個專家並行的程度下均達到了微秒級，最高達到了 46GB/s 的頻寬（理論極限 50GB/s）。
這倆恐怖的數字，估計連英偉達都沒想到自己的卡還能這麼用。
由於在實際分發階段，往往無法提前知道當前這個階段究竟會接收多少 tokens。
因此，這裡會涉及一個隱式的 CPU 等待 GPU 接收計數訊號的過程。
圖中可以看到 CPU 端先發起相應的操作，然後在等待狀態下監聽 GPU 的進度。
當 GPU 收到分發或合併指令並完成相應的工作後，會將結果或狀態告知 CPU。CPU 收到訊號後，再啟動下一階段的排程和計算。
透過這種顯式和隱式的等待-通知機制，能夠保證更高效的並行與資源利用。
針對該核心還專門引入了一種基於hook的通訊-計算重疊方法，不會佔用任何流式多處理器（SM）資源。
在 GPU 內部，還有一個重要概念：流式多處理器（Streaming Multiprocessor，SM）。
它是 GPU 的基本計算單元，每個 SM 都能並行執行成百上千個執行緒，為模型運算提供驚人的吞吐力。然而，如果通訊速度無法跟上 Token 和專家並行度的需求，SM 等不到資料就會“閒置”，從而導致 GPU 資源浪費。
正因如此，如何在保證大規模分發 / 聚合的同時，讓 SM 最大化利用，成為 MoE 並行實踐中的一大挑戰。
DeepSeek 在用 NVIDIA 的 H800 GPU 訓練 V3 時，對 GPU 的核心計算單元（SM，流多處理器）進行了定製化調整,他們把 132 個 SM 中的 20 個專門用來處理伺服器之間的通訊任務，而不是計算任務。
它做到了以下幾點：
  1. 降低通訊開銷：RDMA 在後臺傳輸資料，減少了因資料分發和收集帶來的阻塞或等待。
  2. 釋放計算資源：通訊不佔用 GPU SM，可讓更多計算核心專注於模型的前向或後向計算過程。
  3. 雙批次（或多批次）重疊：利用接收鉤子進一步將不同批次的計算與通訊交錯執行，在不干擾計算的情況下完成資料傳輸，從而提升整體吞吐量。
  4. 靈活適配：可根據實際負載大小或場景需求，調整通訊與計算的重疊比例，從而獲得更好的效能表現。

PTX 指令挖掘

最後，DeepEp 發現並使用了一個未記錄在英偉達文件之中的PTX 指令：
ld.global.nc.L1::no_allocate.L2::256B
雖然此指令會導致未定義行為：使用非一致性只讀 PTX 修飾符.nc 訪問不穩定 GPU 記憶體。
但是在 Hopper 架構上，DeepEp 使用_._
L1::no_allocate
DeepEP 強調這條指令在 Hopper 架構（NVIDIA 最新的 GPU 架構之一）上經過了測試，保證了正確性，並且效能更好。
順便科普一下——
  • PTX (Parallel Thread Execution)：是 NVIDIA 為其 GPU 設計的一種低階虛擬指令集架構（ISA）。它類似於 CPU 的組合語言，但更接近硬體。開發者可以直接編寫 PTX 程式碼，或者透過 CUDA 編譯器（如 nvcc）將 CUDA C/C++ 程式碼編譯成 PTX 程式碼。
這意味著什麼，給 NVIDIA 的護城河開了個口子，並非完全“黑盒”，有可能透過逆向工程或其他手段搞點事情出來。
到這裡，我只想說：
文能開源惠天下，武能壓榨每一滴算力！
🐮

相關文章