DeepSeek、MoEの学習と推論向けオープンソース DeepEP ライブラリを公開

 (github.com/deepseek-ai)
1 ポイント 投稿者 GN⁺ 2025-02-26 | 1件のコメント | WhatsAppで共有

DeepEP

DeepEP は、Mixture-of-Experts (MoE) と expert parallelism (EP) のための通信ライブラリ。高速・低遅延の all-to-all GPU カーネルを提供し、MoE dispatch と combine として知られている処理を担う。また、FP8 を含む低精度演算をサポートする。DeepSeek-V3 論文で提案された group-limited gating アルゴリズムに合わせて、非対称ドメイン帯域幅転送を最適化したカーネルを提供し、NVLink ドメインから RDMA ドメインへデータを転送する。これらのカーネルは高スループットを実現し、学習と推論プリフィル作業に適している。さらに、SM (Streaming Multiprocessors) 数の制御もサポートする。遅延に敏感な推論デコーディング向けに、DeepEP は純粋な RDMA を用いて遅延を最小化する低遅延カーネルを含む。このライブラリは、SM リソースを占有しない hook ベースの通信・計算オーバーラップ方式を導入している。

性能

NVLink および RDMA 転送を使う通常カーネル

  • H800 で NVLink の最大帯域幅約 160 GB/s、CX7 InfiniBand 400 Gb/s RDMA ネットワークカード（最大帯域幅約 50 GB/s）に接続して通常カーネルをテスト。
  • DeepSeek-V3/R1 の事前学習設定（バッチあたり 4096 トークン、7168 hidden、上位 4 グループ、上位 8 エキスパート、FP8 dispatch および BF16 combine）に従う。

純粋な RDMA を使う低遅延カーネル

  • H800 で CX7 InfiniBand 400 Gb/s RDMA ネットワークカード（最大帯域幅約 50 GB/s）に接続して低遅延カーネルをテスト。
  • 一般的な DeepSeek-V3/R1 本番設定（バッチあたり 128 トークン、7168 hidden、上位 8 エキスパート、FP8 dispatch および BF16 combine）に従う。

クイックスタート

要件

  • Hopper GPU（今後さらに多くのアーキテクチャやデバイスをサポートする可能性あり）
  • Python 3.8 以上
  • CUDA 12.3 以上
  • PyTorch 2.1 以上
  • ノード内通信用の NVLink
  • ノード間通信用の RDMA ネットワーク

NVSHMEM 依存関係のダウンロードとインストール

DeepEP は修正済み NVSHMEM に依存している。インストールガイドを参照して導入する必要がある。

ネットワーク構成

DeepEP は InfiniBand ネットワークで完全にテストされており、理論上は RDMA over Converged Ethernet (RoCE) とも互換性がある。

トラフィック分離

InfiniBand は Virtual Lanes (VL) を通じてトラフィック分離をサポートする。異なる種類のトラフィック間の干渉を防ぐため、次のように仮想レーンへジョブを分離することが推奨される:

  • 通常カーネルを使うジョブ
  • 低遅延カーネルを使うジョブ
  • その他のジョブ

DeepEP では、NVSHMEM_IB_SL 環境変数を設定して仮想レーン割り当てを制御できる。

適応型ルーティング

適応型ルーティングは InfiniBand スイッチが提供する高度なルーティング機能で、複数の経路へトラフィックを均等に分散できる。現在、低遅延カーネルは適応型ルーティングをサポートするが、通常カーネルはサポートしていない（今後対応する可能性がある）。通常のノード間カーネルで適応型ルーティングを有効にすると、デッドロックやデータ破損の問題が発生する可能性がある。低遅延カーネルでは、適応型ルーティングを有効にするとルーティング競合によるネットワーク混雑を完全に取り除ける一方で、追加の遅延が発生する。最適な性能のため、次の構成が推奨される:

  • ネットワーク負荷が高い環境では適応型ルーティングを有効化
  • ネットワーク負荷が低い環境では静的ルーティングを使用

輻輳制御

本番環境で有意な輻輳が観測されなかったため、輻輳制御は無効化されている。

インターフェースと例

モデル学習または推論プリフィルでの使用例

通常カーネルは、モデル学習または推論プリフィル段階（逆伝播部分なし）で使用できる。

推論デコーディングでの使用例

低遅延カーネルは、推論デコーディング段階で使用できる。

注意事項

  • 極限の性能を得るため、文書化されていない PTX 命令 ld.global.nc.L1::no_allocate.L2::256B を発見して使用している。この命令は、非一貫性の読み取り専用 PTX 修飾子を使って揮発性 GPU メモリへアクセスする未定義動作を引き起こす。ただし、Hopper アーキテクチャで .L1::no_allocate をテストしたところ、性能が大幅に向上した。他のプラットフォームでカーネルが動作しない場合は、setup.py に DISABLE_AGGRESSIVE_PTX_INSTRS=1 を追加してこれを無効化するか、issue を報告できる。
  • クラスタでより良い性能を得るため、すべてのテストを実行し、最適な自動チューニング構成を使うことが推奨される。デフォルト構成は DeepSeek の内部クラスタで最適化されている。

ライセンス

このコードリポジトリは MIT ライセンスの下で公開されており、NVSHMEM を参照するコード（csrc/kernels/ibgda_device.cuh および third-party/nvshmem.patch を含む）には NVSHMEM SLA が適用される。

関連記事

1件のコメント

 
GN⁺ 2025-02-26
Hacker Newsのコメント
  • 極限の性能のために、文書化されていないPTX命令を見つけて使用している。この命令は、一貫性のない読み取り専用PTX修飾子を使って揮発性GPUメモリにアクセスするため、未定義動作を引き起こす可能性がある。しかし、Hopperアーキテクチャでは .L1::no_allocate でテストされた正確性が保証され、性能もはるかに向上するはずだ
  • Zuckerbergは、MetaがAIをオープンソース化していると主張するのをやめるべきだ。彼らが公開しているのはコードではなく重みだけだ。本当のオープンソースAIはDeepSeekだけだ
  • まるで駄菓子屋にいる子どものような気分だ。こうしたトリックのいくつかは、論文をもとに正しくリバースエンジニアリングするのにかなり時間がかかるだろう。今週の発表が、MoEを標準的な学術モデルとして使うルネサンスの始まりになることを願う
  • この人たちを好きにならずにはいられない。彼らは私たちみんなのために、本当にオープンソースの限界を押し広げている。共有してくれてありがとう
    • 効率的で最適化されたall-to-all通信
    • NVLinkとRDMAによるノード内およびノード間サポート
    • 学習と推論のプリフィル向け高スループットカーネル
    • 推論デコード向け低遅延カーネル
    • ネイティブFP8ディスパッチ対応
    • 計算と通信のオーバーラップのための柔軟なGPUリソース制御
  • DeepSeekの取り組みの背後にある動機は間違っているかもしれない（たとえば、AIにおける米国の先行優位をなくそうとする国家支援の試みなど）。しかし世界全体にとって、その結果はただただ素晴らしい
    • 最悪の場合（間違った理由でこの取り組みをしているとしても）でも、DeepSeekには感謝している。彼らは、OpenAIが何年も世界中に対して嘘をついてきたことを、実際にやっている
    • 本当にすごい
  • みんなが期待していたPTXが今回含まれているのか気になる
  • 技術レポートで言及されていたPTX命令は、ここのコードに対応づけられるべきだ
  • 米国が、DeepSeekがH800しか使っていなかったことを確認するためにシンガポールでGPUの領収書を追跡しているあいだ、世界の残りの国々はフルのH100でこれらの最適化を実行できる
    • 米国の制裁と、その命令が世界全体を覆っていると信じる傲慢さのせいで、H100を入手したりアクセスしたりするのが難しかったふりをしているのだろうか
  • 本物の"Open AI™"企業による2度目のオープンソースリリースであり、MITライセンスのもとで公開されている
    • DeepSeekは、$157B+をうたう企業よりもオープンだ
    • ほとんど誰もMetaのLlamaについて語っておらず、誰もがLlama 4がそれ相応の理由とともにリリースされることを期待すべきだ
    • 目標は、ゼロへの競争の真ん中で板挟みにならないことだ