vLLM and Novita AI Release Chord: Faster INT4 MoE Kernels for Kimi K2.x

TL;DR

Novita AI has open‑sourced Chord, a high‑performance W4A16 (INT4 weights, BF16 activations) MoE CUDA operator that delivers up to 2.15× speed‑up on NVIDIA B300 GPUs and up to 1.33× on H200 GPUs compared with the public Humming backend. The operator integrates with vLLM via the existing humming import root, though full grouped‑kernel integration is still work‑in‑progress.

Core Contributions

Chord provides two independent kernel families—indexed and grouped—that are each tuned for specific serving shapes of Kimi K2.x models. The indexed kernels target H200 EP8 prefill, H200 TP8 single‑instance serving, and B300 EP8 decode workloads, while the grouped kernels (contiguous prefill and masked decode) target SM90 GPUs and use a DeepGEMM‑derived implementation.

"The idea behind these numbers is that one W4A16 MoE kernel cannot be right for every request. Routed tokens per expert varies by orders of magnitude between prefill and decode, and it is that quantity, not total token count, that decides which schedule wins."

Performance Highlights

Indexed kernels beat Humming on H200 and B300

Scenario GPU Speed‑up (Chord vs. Humming)
H200 EP8 prefill H200 1.11–1.20×
H200 TP8 single‑instance serving H200 1.17–1.33×
H200 EP8 decode (gate/up) H200 1.16–1.24× (down up to 1.31×)
B300 EP8 decode (untuned Humming) B300 1.81–2.15×

These figures are per‑layer kernel timings, not end‑to‑end latency guarantees. Full tables, shape definitions, and benchmarking methodology are documented in the repository’s docs/performance.md and docs/benchmarking.md.

Grouped kernels provide comparable gains on SM90

Scenario GPU Speed‑up
H200 EP8 grouped prefill (128 rows/expert) H200 1.31×
H200 EP8 grouped decode (32 tokens/expert) H200 1.35×

Grouped kernels use a different weight layout (contiguous or masked) and a persistent, warp‑specialized main loop that pipelines TMA loads, dequantization, and WGMMA execution.

Integration with vLLM

Installation and activation

pip install git+https://github.com/novitalabs/chord.git
vllm serve <kimi-k2.x-int4-model> --quantization humming
# or set moe_backend="humming" in the vLLM config
  • The package installs both chord and humming module roots. vLLM’s lazy facade resolves humming.{dtypes,config,layer,…} without requiring a Chord‑specific patch on branches that support the W4A16 group‑scale format.
  • Unsupported quantization schemes cause a load‑time error rather than silently falling back to an incorrect kernel.
  • Grouped integration is currently a work‑in‑progress; only the indexed path works out‑of‑the‑box.

Runtime considerations

  • Profile selection (EP8 vs. TP8) is inferred from the projection shapes passed by vLLM; no extra shard argument is needed for indexed kernels.
  • Indexed kernels can operate on vLLM’s over‑allocated moe_align_block_size buffers and remain CUDA‑Graph capturable when trusted routing validation is disabled.
  • Keep the environment variables VLLM_HUMMING_USE_F16_ACCUM and VLLM_BATCH_INVARIANT off, as neither backend implements those compute options.
  • Older vLLM branches may need to add the group‑32 quantization keys to _supports_quant_scheme to recognize the Chord format.

Kernel‑Level Optimizations

Indexed kernel techniques

  • Batched wait<1> WGMMA pipelining overlaps weight loading, dequantization, and the previous WGMMA group, yielding a 3–6 % gain on the up‑projection and 1–5 % on the down‑projection.
  • Tokens‑per‑expert (tok_e) heuristic selects block‑M sizes based on routed tokens per expert, using a simple threshold of 80 tokens/expert to switch between a baseline block‑count search and a padded‑row‑aware layout.
  • Bounded 2‑CTA/SM window raises resident warps for mid‑size tiles, hiding cp.async gather latency.
  • Shape‑aware stream‑K gating disables stream‑K when the M×N grid is saturated, avoiding unnecessary split/reduction overhead.

Grouped SM90 kernel techniques

  • Contiguous prefill pads each expert to a 128‑row boundary and uses BLOCK_K=64 with transposed BF16 scales.
  • Masked decode reserves a fixed per‑expert row budget, packs weights with BLOCK_K=128, and selects BLOCK_M from the expected token count.
  • Wave‑aware BN selection chooses BN256 only when enough N‑tiles exist; otherwise BN128 is used to keep occupancy high.
  • Latency‑tuned buffered‑K depth targets ~512 buffered K elements for typical decode, scaling up to ~768 for large masked tiles without exhausting shared memory.

End‑to‑End Serving Impact

A separate serving report on an 8×H200 configuration (Kimi‑K2.6, FP8 KV cache, ShareGPT workload) showed modest but consistent improvements:

Metric Humming Chord Δ
Mean TTFT 2022 ms 1849 ms –8.6 %
Prefill throughput 20,716 tok/s 22,712 tok/s +9.6 %
Decode throughput (batch 8) 483 tok/s 503 tok/s +4.1 %
Decode throughput (batch 64) 1650 tok/s 1740 tok/s +5.5 %
Decode throughput (batch 128) 2514 tok/s 2715 tok/s +8.0 %
The report observed no accuracy regression on OCRBench and GSM8K.

Roadmap

  1. Complete grouped integration with vLLM’s Humming backend, exposing contiguous and masked operators through the existing framework.
  2. Release EP8 prefill kernels for B200/B300 GPUs; internal tests show promising performance and will be shared in a follow‑up release.

Getting Started

Chord is available at https://github.com/novitalabs/chord. The repository includes:

  • Getting started guide – installation and basic usage.
  • Optimizations documentation – detailed description of the kernel heuristics.
  • Performance tables – full per‑scenario latency numbers.
  • Tuning and benchmarking methodology – reproducible test scripts. Contributions, issue reports, and benchmark results from the community are welcomed.

Acknowledgements

Chord’s indexed path builds on inclusionAI/Humming (commit 4351af3). The grouped SM90 backend adapts DeepGEMM’s Hopper GEMM infrastructure. The project is released under the Apache‑2.0 license. Thanks to the Novita AI team for developing Chord and to the vLLM maintainers and community for supporting the integration.

Sources