vLLM 和 Novita AI 发布 Chord：适用于 Kimi K2.x 的更快 INT4 MoE 内核

TL;DR

Novita AI 已开源 Chord，这是一个高性能的 W4A16（INT4 权重，BF16 激活）MoE CUDA 操作符，在 NVIDIA B300 GPU 上相比公开的 Humming 后端最高可实现 2.15× 的加速，在 H200 GPU 上最高可实现 1.33× 的加速。该操作符通过现有的 humming 导入根路径与 vLLM 集成，尽管完整的分组内核集成仍在开发中。

核心贡献

Chord 提供了两个独立的内核族——索引 和 分组——分别针对 Kimi K2.x 模型的特定服务形状进行了优化。索引内核针对 H200 EP8 预填充、H200 TP8 单实例服务和 B300 EP8 解码工作负载，而分组内核（连续预填充和掩码解码）针对 SM90 GPU，并采用基于 DeepGEMM 的实现。

"这些数字背后的理念是，一个 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 解码（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 个令牌） 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 传入的投影形状推断得出；索引内核无需额外的分片参数。
  • 索引内核可在 vLLM 的过度分配 moe_align_block_size 缓冲区上运行，并在禁用可信路由验证时保持 CUDA 图可捕获性。
  • 请保持环境变量 VLLM_HUMMING_USE_F16_ACCUM 和 VLLM_BATCH_INVARIANT 关闭，因为两个后端均未实现这些计算选项。
  • 较旧的 vLLM 分支可能需要将 group-32 量化键添加到 _supports_quant_scheme 以识别 Chord 格式。

内核级优化

索引内核技术

  • 批量 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 许可证发布。感谢 Novita AI 团队开发 Chord，也感谢 vLLM 维护者和社区对集成的支持。

Sources