vLLM と Novita AI が Chord をリリース：Kimi K2.x 用の高速 INT4 MoE カーネル

TL;DR

Novita AI は Chord をオープンソース化しました。これは、NVIDIA B300 GPU で公開されている Humming バックエンドと比較して最大 2.15×、H200 GPU で最大 1.33× の高速化を実現する、高性能な W4A16（INT4 重み、BF16 活性化）MoE CUDA 演算子です。この演算子は既存の humming インポートルートを介して vLLM と統合されていますが、完全なグループ化カーネル統合はまだ開発中です。

コア貢献

Chord は、Kimi K2.x モデルの特定のサービング形状に最適化された、インデックス付き と グループ化 の2つの独立したカーネルファミリを提供します。インデックス付きカーネルは H200 EP8 プレフィル、H200 TP8 シングルインスタンスサービング、B300 EP8 デコードワークロードを対象としており、グループ化カーネル（連続プレフィルおよびマスク付きデコード）は SM90 GPU を対象とし、DeepGEMM 由来の実装を使用しています。

"これらの数値の背景にある考え方は、1つの W4A16 MoE カーネルがすべてのリクエストに適しているわけではないということです。プレフィルとデコードの間で専用エキスパートごとのルーティングトークン数は桁違いに異なり、その量こそがどのスケジュールが勝つかを決定するのです。"

パフォーマンスのハイライト

インデックス付きカーネルは H200 と B300 で Humming を上回る

シナリオ GPU 速度向上（Chord vs. Humming）
H200 EP8 プレフィル H200 1.11–1.20×
H200 TP8 シングルインスタンスサービング H200 1.17–1.33×
H200 EP8 デコード（ゲート/アップ） H200 1.16–1.24×（ダウンは最大 1.31×）
B300 EP8 デコード（チューニング未実施の Humming） B300 1.81–2.15×

これらの数値は レイヤーごとのカーネル実行時間 であり、エンドツーエンドのレイテンシーガラントリーではありません。完全なテーブル、形状定義、ベンチマーク手法については、リポジトリの docs/performance.md および docs/benchmarking.md を参照してください。

グループ化カーネルは SM90 で同等の利点を提供

シナリオ GPU 速度向上
H200 EP8 グループ化プレフィル（128 行/エキスパート） H200 1.31×
H200 EP8 グループ化デコード（32 トークン/エキスパート） H200 1.35×

グループ化カーネルは異なる重みレイアウト（連続またはマスク付き）と、永続的でワープ専用のメインループを使用しており、TMA 読み込み、デキューント化、WGMMA 実行をパイプライン化しています。

vLLM との統合

インストールと有効化

pip install git+https://github.com/novitalabs/chord.git
vllm serve <kimi-k2.x-int4-model> --quantization humming
# または vLLM 設定で moe_backend="humming" を設定
  • パッケージは chord と humming の両方のモジュールルートをインストールします。vLLM の遅延フェイスは、W4A16 グループスケール形式をサポートするブランチでは、Chord 特有のパッチを必要とせずに humming.{dtypes,config,layer,…} を解決できます。
  • サポートされていない量子化スキームは、誤ったカーネルに静かにフォールバックするのではなく、ロード時にエラーを発生させます。
  • グループ化統合は現在開発中 であり、現時点ではインデックスパスのみが即時利用可能です。

ランタイムの考慮事項

  • プロファイル選択（EP8 と TP8）は、vLLM が渡す投影形状から推論されます。インデックス付きカーネルには、追加のシャード引数は必要ありません。
  • インデックス付きカーネルは、vLLM のオーバーアロケートされた moe_align_block_size バッファ上で動作でき、信頼できるルーティング検証が無効化されている場合、CUDA グラフキャプチャ可能になります。
  • 環境変数 VLLM_HUMMING_USE_F16_ACCUM と VLLM_BATCH_INVARIANT は 無効にしてください。どちらのバックエンドもこれらの計算オプションを実装していません。
  • 古い vLLM ブランチでは、_supports_quant_scheme にグループ化 32 の量子化キーを追加する必要がある場合があります。

カーネルレベルの最適化

インデックス付きカーネル技術

  • バッチ化された wait<1> WGMMA パイプライン化：重み読み込み、デキューント化、および前の WGMMA グループを重ね合わせることで、アッププロージェクションで 3–6 %、ダウンプロージェクションで 1–5 % の向上を実現。
  • エキスパートごとのトークン数（tok_e）ヒューリスティック：ルーティングされたエキスパートごとのトークン数に基づいて block-M サイズを選択。80 トークン/エキスパートという単純な閾値を使って、ベースラインのブロック数探索とパッド付き行対応レイアウトの間で切り替え。
  • 2-CTA/SM ウィンドウの上限付き：中程度のタイルでは、残留ワープ数を増やし、cp.async ゲザーのレイテンシを隠蔽。
  • 形状に応じたストリーム-K ゲート：M×N グリッドが飽和している場合、ストリーム-K を無効化し、不要な分割/還元のオーバーヘッドを回避。

グループ化 SM90 カーネル技術

  • 連続プレフィル：各エキスパートを 128 行の境界にパディングし、BLOCK_K=64 と転置された BF16 スケールを使用。
  • マスク付きデコード：エキスパートごとに固定の行予算を確保し、BLOCK_K=128 で重みをパックし、BLOCK_M を予想されるトークン数から選択。
  • 波に応じた BN 選択：十分な N タイルが存在する場合のみ BN256 を選択。それ以外は、占有率を維持するために BN128 を使用。
  • レイテンシ調整済みのバッファリングK深度：通常のデコードでは約 512 のバッファリングK要素をターゲットとし、大きなマスク付きタイルでは共有メモリを枯渇させずに最大約 768 までスケーリング。

エンドツーエンドのサービングへの影響

8×H200 構成（Kimi‑K2.6、FP8 KVキャッシュ、ShareGPTワークロード）での別個のサービングレポートでは、わずかではあるが一貫した改善が観測されました：

メトリクス Humming Chord Δ
平均 TTFT 2022 ms 1849 ms –8.6 %
プレフィルスループット 20,716 tok/s 22,712 tok/s +9.6 %
デコードスループット（バッチ 8） 483 tok/s 503 tok/s +4.1 %
デコードスループット（バッチ 64） 1650 tok/s 1740 tok/s +5.5 %
デコードスループット（バッチ 128） 2514 tok/s 2715 tok/s +8.0 %
レポートでは OCRBench および GSM8K で 精度の低下は観測されませんでした。

ロードマップ

  1. vLLM の Humming バックエンドとの完全なグループ化統合：既存のフレームワークを通じて連続およびマスク付き演算子を公開。
  2. B200/B300 GPU 用の EP8 プレフィルカーネルのリリース：内部テストでは有望なパフォーマンスが得られており、後続リリースで共有予定。

はじめに

Chord は https://github.com/novitalabs/chord で利用可能です。リポジトリには以下が含まれます：

  • はじめにガイド – インストールと基本的な使い方。
  • 最適化ドキュメント – カーネルヒューリスティックの詳細な説明。
  • パフォーマンステーブル – シナリオごとの完全なレイテンシ数値。
  • チューニングとベンチマーク手法 – 再現可能なテストスクリプト。 コミュニティからの貢献、問題報告、ベンチマーク結果を歓迎しています。

謝辞

Chord のインデックスパスは inclusionAI/Humming（コミット 4351af3）に基づいています。グループ化 SM90 バックエンドは DeepGEMM の Hopper GEMM インフラをアダプトしています。プロジェクトは Apache‑2.0 ライセンスの下でリリースされています。Chord の開発に貢献した Novita AI チームに感謝し、vLLM のメンテナーやコミュニティが統合を支援してくれたことに感謝します。

Sources