618ZXW

DeepSeek の 2 番目のオープンソースリリース: MoE と EP 向けにカスタマイズされた通信ライブラリ。現在は NVIDIA グラフィック カードに関連付けられています。

約束通り、朗報です!DeepSeek Open Source Week の第 2 回目が始まりました!

DeepEP は、MoE モデルのトレーニングと推論のための最初のオープンソース EP 通信ライブラリ (エキスパート並列処理) です。

これは、MoE ディスパッチおよびコンビンとも呼ばれる、高スループットと低レイテンシの全対全 GPU カーネルを提供します。

ライブラリは、FP8 を含む低精度演算もサポートしています。

慣例に従い、最も許容度の高いオープンソース ライセンスは MIT です。

本日、DeepSeek はまず GitHub でローンチし、その後公式 Twitter アカウントで通知を投稿することを選択しました

予想通り、観客からは拍手が起こりました。

DeepSeek オープンソースの列車は決して止まりません。

DeepEP のパフォーマンスはどの程度ですか?

DeepSeek の公式 Twitter アカウントは、DeepEP の主要な要素を次のようにまとめています。

  • 高効率で最適化された全対全通信
  • NVLinkとRDMAのノード内およびノー​​ド間サポート
  • トレーニングと推論の事前充填のための高スループットカーネル
  • 推論とデコードのための低レイテンシカーネル
  • ネイティブFP8スケジューリングサポート
  • コンピューティング通信のオーバーラップに対する柔軟な GPU リソース制御。

まず、パフォーマンスの 2 つの重要な側面を見てみましょう。

(注: DeepEP での実装は、DeepSeek-V3 の論文とは若干異なる可能性があります。)

NVLinkとRDMA転送を備えた標準カーネル

DeepSeek-V3 論文で提案されたグループ制限ゲーティング アルゴリズムとの一貫性を維持するために、 DeepEP は、NVLink ドメインから RDMA ドメインへのデータ転送など、非対称ドメイン帯域幅転送に最適化されたカーネル セットを提供します

これらのカーネルは高いスループットを提供するため、トレーニングや推論の事前入力タスクに適しています。

さらに、SM(ストリーミングマルチプロセッサ)数制御もサポートしています。

DeepEP チームは、H800 (最大帯域幅約 160 GB/秒の NVLink) 上で通常のカーネルをテストし、各カーネルを CX7 InfiniBand 400 Gb/秒の RDMA ネットワーク カード (最大帯域幅約 50 GB/秒) に接続しました。

また、DeepSeek-V3/R1 の事前トレーニング設定 (バッチあたり 4096 トークン、非表示トークン 7168 個、最初の 4 つのグループ、最初の 8 人のエキスパート、FP8 スケジューリング、BF16 の組み合わせ) にも従います。

純粋なRDMAによる低レイテンシカーネル

レイテンシの影響を受けやすい推論およびデコードのシナリオ向けに、DeepEP にはレイテンシを最小限に抑えるための純粋な RDMA を備えた低レイテンシ カーネルのセットが含まれています。

このライブラリでは、SM リソースを消費しないオーバーラップ通信計算用のフックベースのメソッドも導入されています。

DeepEP チームは、各カーネルを CX7 InfiniBand 400 Gb/s RDMA ネットワーク カード (最大帯域幅 ~50 GB/s) に接続して、H800 上で低レイテンシ カーネルをテストしました。

これは、一般的な DeepSeek-V3/R1 の製品セットアップ (バッチあたり 128 トークン、7168 個の隠しトークン、最初の 8 人のエキスパート、FP8 スケジューリング、BF16 の組み合わせ) に従います。

コンシューマーグレードのグラフィックカードは現在サポートされていません。最適な自動最適化設定を使用することをお勧めします。

GitHub では、DeepSeek チームがさまざまな適応環境と構成要件を網羅した DeepEP の使用方法を明確に説明しています。

まず、DeepEP に必要なソフトウェアとハ​​ードウェア環境のバージョンは次のとおりです。

  • Hopper GPU (将来的にはさらに多くのアーキテクチャやデバイスをサポートする可能性があります)
  • Python 3.8以降
  • CUDA 12.3以降
  • PyTorch 2.1以降
  • ノード内通信用のNVLink
  • ノード内通信用のRDMAネットワーク

次に、DeepEP を使用するには、チームが修正した NVSHMEM 依存関係をダウンロードしてインストールする必要があります (手順については、DeepSeek チームの NVSHMEM インストール ガイドを参照してください)。

次に、 deep_ep を Python プロジェクトにインポートして、「楽しんで」みましょう。

ネットワーク構成に関しては、DeepEP は InfiniBand ネットワークの包括的なテストに合格しています。

理論的には、統合イーサネットに基づく RDMA (RoCE) とも互換性があります。

InfiniBand は、仮想レーン (VL) を通じてトラフィックの分離をサポートします。

異なるタイプのトラフィック間の干渉を防ぐために、DeepEP チームは、次に示すように、ワークロードを異なる仮想チャネルに分離することを推奨しています。

  • 通常のカーネルを使用するワークロード
  • 低レイテンシカーネルを使用するワークロード
  • その他の作業負荷

DeepEP の場合、開発者は NVSHMEM、_IB、および _SL 環境変数を設定することで仮想チャネルの割り当てを制御できます。

注目すべきは、アダプティブ ルーティングは、トラフィックを複数のパスに均等に分散できる InfiniBand スイッチによって提供される高度なルーティング機能であるということです。

現在、低レイテンシ カーネルは Adaptive Routing をサポートしていますが、通常のカーネルはサポートしていません (サポートはすぐに追加される可能性があります)。

通常のノード間カーネルに対して適応ルーティングを有効にすると、デッドロックやデータ破損の問題が発生する可能性があります

低レイテンシカーネルの場合、アダプティブルーティングを有効にすると、ルーティングの競合によって発生するネットワーク輻輳を完全に排除できますが、追加のレイテンシも発生します。

DeepEP チームは、最適なパフォーマンスを得るために次の構成を推奨しています。

  • ネットワーク負荷の高い環境で適応ルーティングを有効にします。
  • ネットワーク負荷の低い環境では静的ルーティングを使用します。

ちなみに、DeepEP は実稼働環境で重大な混雑が観察されなかったため、混雑制御を無効にしました

DeepEPチームからの最後のアドバイスは

最高のパフォーマンスを実現するために、チームはドキュメント外の PTX 命令 ld.global.nc.L1::no_allocate.L2::256B を発見して使用しました。

この命令により、非コヒーレントな読み取り専用 PTX 修飾子 .nc を使用して揮発性 GPU メモリにアクセスするという、未定義の動作が発生します。

ただし、正確性はテスト済みで保証されています。L1::no_allocate は、Hopper アーキテクチャでより優れたパフォーマンスを発揮します。

カーネルが他のプラットフォームで動作しないことがわかった場合は、`setup.py` に `DISABLE_AGGRESSIVE_PTX_INSTRS=1` を追加してこの機能を無効にするか、問題を送信してください。

クラスターでより優れたパフォーマンスを実現するために、DeepSeek ではすべてのテストを実行し、最適な自動最適化構成を使用することを推奨しています。

デフォルト構成が DeepSeek の内部クラスターで最適化されているためです。

もう一つ

DeepSeek は、このオープンソース ウィークのために GitHub に新しいリポジトリを作成しました。

https://github.com/deepseek-a...

過去 2 日間のリリースに基づくと、この Open Source Week 中にリリースされたすべてのコンテンツは AI インフラストラクチャに関連していると推測されます

しかし、あまり良くないニュースもあります。DeepSeek のオープンソースの毎週の更新スケジュールはやや不安定なようです。

昨日は午前9時34分、今日は午前10時24分、明日は...

したがって、DeepSeek の絶え間ない進歩に遅れずについていくために、QuantumBit は、できるだけ早く参加していただくよう皆様に心からお招きします (doge)。

QR コードをスキャンし、メッセージに「DeepSeek - 職業/名前」を追加してグループ チャットに参加し、DeepSeek Open Source Week を一緒にフォローしましょう。

DeepEP GitHub:

https://github.com/deepseek-a...