vLLM 與 Novita AI 發布 Chord：針對 Kimi K2.x 的更快 INT4 MoE 核心

TL;DR

Novita AI 已將 Chord 開源，這是一個高性能的 W4A16（INT4 权重，BF16 激活）MoE CUDA 操作器，與公開的 Humming 後端相比，在 NVIDIA B300 GPU 上可實現最高 2.15× 的加速，在 H200 GPU 上可實現最高 1.33× 的加速。該操作器透過現有的 humming 導入根目錄與 vLLM 整合，但完整的分組核心整合仍在進行中。

核心貢獻

Chord 提供兩個獨立的核心家族——索引式 和 分組式，各自針對 Kimi K2.x 模型的特定服務形狀進行優化。索引式核心針對 H200 EP8 預填入、H200 TP8 單實例服務以及 B300 EP8 解碼工作負載，而分組式核心（連續預填入與遮罩解碼）則針對 SM90 GPU，並採用 DeepGEMM 衍生的實作方式。

"這些數字背後的想法是，一個 W4A16 MoE 核心無法適用於所有請求。在預填入與解碼之間，每專家的路由 token 數量差異可達數個數量級，而正是這個數量，而非總 token 數，決定了哪種排程勝出。"

性能亮點

索引式核心在 H200 與 B300 上超越 Humming

場景 GPU 加速比（Chord 對 Humming）
H200 EP8 預填入 H200 1.11–1.20×
H200 TP8 單實例服務 H200 1.17–1.33×
H200 EP8 解碼（gate/up） H200 1.16–1.24×（down 最高達 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 個 token） 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 的惰性外觀會解析 humming.{dtypes,config,layer,…}，無需在支援 W4A16 分組尺度格式的分支上進行 Chord 專用修補。
  • 不支援的量化方案會在載入時拋出錯誤，而非靜默降回錯誤的核心。
  • 分組整合目前仍在進行中；僅索引式路徑可直接使用。

執行時考量

  • 優化選擇（EP8 對 TP8）由 vLLM 傳遞的投影形狀推斷；索引式核心無需額外的 shard 參數。
  • 索引式核心可在 vLLM 的過度配置 moe_align_block_size 缓衝區上運作，且在停用信任路由驗證時仍可保持 CUDA-Graph 可捕捉性。
  • 請將環境變數 VLLM_HUMMING_USE_F16_ACCUM 與 VLLM_BATCH_INVARIANT 關閉，因為兩者後端皆未實作這些運算選項。
  • 較舊的 vLLM 分支可能需要將 group-32 量化鍵加入 _supports_quant_scheme 以識別 Chord 格式。

核心層級優化

索引式核心技術

  • 批次 wait<1> WGMMA 流水線 重疊權重載入、反量化與前一組 WGMMA，使 up-投影提升 3–6 %，down-投影提升 1–5 %。
  • 每專家 token 數（tok_e）啟發式 根據每專家路由 token 數選擇 block-M 大小，使用 80 token/專家的簡單閾值切換基線 block 數搜尋與填充行感知佈局。
  • 受限的 2-CTA/SM 窗口 提升中等大小圖塊的駐留波長，隱藏 cp.async 收集延遲。
  • 形狀感知的 stream-K 開關 當 M×N 網格已飽和時停用 stream-K，避免不必要的分割/還原開銷。

分組 SM90 核心技術

  • 連續預填入 將每個專家補齊至 128 行邊界，並使用 BLOCK_K=64 與轉置的 BF16 比例。
  • 遮罩解碼 為每專家保留固定的行預算，以 BLOCK_K=128 壓縮權重，並根據預期 token 數選擇 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 %
解碼吞吐量（batch 8） 483 tok/s 503 tok/s +4.1 %
解碼吞吐量（batch 64） 1650 tok/s 1740 tok/s +5.5 %
解碼吞吐量（batch 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 許可證釋出。感謝 Novita AI 團隊開發 Chord，以及 vLLM 維護者與社群對整合的支援。

Sources