DeepSeek、DeepEPをオープンソース公開（5本中2本目）

• Mixture-of-Experts（MoE）および Expert Parallelism（EP）のための高性能通信ライブラリ
• GPUベースのAll-to-Allカーネルを提供し、MoEのディスパッチおよび結合演算を高速に処理
• FP8のような低精度演算をサポート
• DeepSeek-V3論文で提案されたグループ制限ゲーティング（group-limited gating）アルゴリズムを適用し、非対称ドメイン帯域幅フォワーディングを最適化
  • 例: NVLink → RDMAデータ転送の最適化
  • 学習および**推論プリフィル（prefilling）**処理に適した高スループットを提供
• レイテンシに敏感な推論デコーディング向けにRDMA専用の低レイテンシカーネルを含む
• 通信-演算オーバーラップ手法を提供（SMリソースを占有しない）

性能

一般カーネル（NVLinkおよびRDMA転送）

• DeepEPはH800 GPUとCX7 InfiniBand 400Gb/s RDMAネットワーク環境で性能をテスト
• DeepSeek-V3/R1構成をベースに、バッチあたり4096トークン、7168個の隠れノード、top-4グループ、top-8エキスパート構成を適用し、FP8ディスパッチおよびBF16結合を使用
• 性能テストの結果、ノード内（NVLinkベース）通信は約150GB/s以上の帯域幅を示し、ノード間（RDMAベース）通信はエキスパート数に応じて40〜47GB/s水準の帯域幅を記録
• エキスパート数が増えるほどRDMA帯域幅がわずかに増加する傾向を示した（例: エキスパート16人のとき43GB/s、エキスパート64人のとき46GB/s）

低レイテンシカーネル（純粋なRDMA）

• 低レイテンシカーネルの性能を測定した結果、一般カーネルよりレイテンシ（latency）が大幅に短縮
• バッチあたり128トークンを処理する環境では、エキスパート数に応じてレイテンシは増加したが、RDMA帯域幅は比較的一定に維持
• たとえば、エキスパート8人のとき163マイクロ秒（us）、エキスパート256人のとき**194マイクロ秒（us）**程度まで増加
• 結合（combine）演算ではディスパッチよりも高いレイテンシが発生し、エキスパート数が増えるほどRDMA帯域幅が40GB/s以下へ徐々に低下する傾向を示す
• つまり、低レイテンシカーネルは小規模なエキスパートグループでは非常に高速に動作するが、エキスパート数が多くなるとレイテンシが増えるため、適切なバランスが必要

ネットワーク設定

トラフィック分離（Traffic Isolation）

• InfiniBandの**Virtual Lanes（VL）**を活用してトラフィックを分離可能
• 推奨される分離方式:
  • 一般カーネルを使用するジョブ
  • 低レイテンシカーネルを使用するジョブ
  • その他のジョブ
• NVSHMEM_IB_SL環境変数を通じてVLを設定可能

アダプティブルーティング（Adaptive Routing）

• InfiniBandスイッチのアダプティブルーティングをサポート
• 低レイテンシカーネルでは有効化可能、一般カーネルでは無効化が必要（有効化するとデータ破損のリスク）
• 設定の推奨事項:
  • ネットワーク負荷が高い場合: アダプティブルーティングを有効化
  • ネットワーク負荷が低い場合: 静的ルーティングを維持

輻輳制御（Congestion Control）

• DeepEPは輻輳制御機能を無効化した状態で運用
• 実環境でネットワーク輻輳が深刻ではないことを確認

主な技術的考慮事項

• 非公式PTX命令の使用: ld.global.nc.L1::no_allocate.L2::256Bを活用して性能を最適化
  • Hopperアーキテクチャでは正常に動作するが、他のプラットフォームではDISABLE_AGGRESSIVE_PTX_INSTRS=1を設定して無効化可能
• 自動チューニング推奨: 最適な性能のため、クラスターごとに性能テストを行ったうえで設定を適用する必要あり

DeepSeek Open Infraとして公開される5つのオープンソース のうち2本目