vLLM 宣布 Humming 后端可运行 Novita 开源的 Chord 内核，称 H200 TP8 提速 1.33x、B300 EP8 decode 提速 2.15x

9 月 16 日，vLLM 的 X 账号 @vllm_project 发帖宣布，vLLM 的 Humming 后端可以运行 Chord，即 Novita AI 开源、面向 Kimi K2.x 的 W4A16（INT4 权重、BF16 激活）MoE（混合专家）内核（帖子）。帖文给出两组数字：内核增益在 H200 TP8 上达到 1.33x（对比调优后的公共 Humming），在 B300 EP8 decode 上达到 2.15x（对比未调优的默认配置）。

帖子同时说明了集成状态：indexed 路径可在兼容的 vLLM revision 上运行，grouped 集成仍在进行中（WIP）。更完整的细节来自帖文链接的技术博客（vllm.ai），该文 9 月 15 日由 Novita AI 与 vLLM 团队联合署名发布。文中的数字由 Novita AI 与 vLLM 团队自行测量并公布；博客也写明，这是内核级测量，不承诺每个工作负载都能获得同样的端到端增益。

六个场景的逐层测量与限定条件

博客标题给两组数字加的限定词并不一样：H200 那组标的是「最高」，B300 那组标的是「未调优」。测量方法是逐层比较 gate/up 与 down 两个 GEMM（通用矩阵乘法）调用的耗时之和，提速值为 Humming 的 gate_up + down 耗时除以 Chord 的同一项，不含路由、激活与通信；每个场景与同一版公共 Humming（inclusionAI/humming，博客记为 commit 4351af3）在相同 GPU 上对比，indexed 场景使用相同形状与相同路由抽样，grouped 场景匹配每专家的行数，计时用 triton.testing.do_bench。

博客总结表里六个场景的数据：

其中 grouped 两行对比的是 Humming 自己的 grouped_contiguous / grouped_masked 路径，其余为 indexed 路径。

B300 那一行需要单独说明：公共 Humming 没有 SM100/SM103 的调优表，所以 2.15x 对比的是一个未调优的默认配置；H200 的所有面板都是调优对调优。vLLM 发帖约一分钟后，Novita AI 也发布了 X 长文，标题是「我们把 INT4 MoE 提速 2 倍，然后把它开源了」（长文），文里主动点出了这一点：「差距是真实的，但它既是内核设计的差距，也是调优覆盖的差距。与其让你们以后自己发现，不如我们先讲出来。」按博客给出的区间，B300 EP8 decode 的增益为 1.81～2.15x；H200 EP8 prefill 为 1.11～1.20x，H200 TP8 为 1.17～1.33x，H200 EP8 decode 为 1.16～1.24x（down 段最高 1.31x）；grouped 的 prefill 与 decode 分别是 1.00～1.31x 和 1.16～1.35x。

调度由每个专家收到的 token 数决定

博客概括的设计动机是：一个 W4A16 MoE 内核不可能对每个请求都最优。prefill（预填充）阶段，每个专家可能收到几百行 token；decode（解码）阶段只剩个位数，博客的 decode 用例是 9～15 行/专家。两种负载下，决定哪种调度更快的量是「路由到每个专家的 token 数」，而不是总 token 数；Chord 的做法是从实际拿到的形状里选调度。

具体到代码，Chord 仓库（novitalabs/chord）当前 main 分支有两套独立的内核家族。indexed（索引式）路径沿用 Humming 的路由接口，直接消费 vLLM 的 sorted_ids / expert_ids / num_tokens_padded 张量，覆盖 H200 EP8 prefill、H200 TP8 单实例服务、H200 EP8 decode 以及 B200/B300 EP8 decode；grouped_contiguous（prefill）与 grouped_masked（decode）是另一套 SM90 家族，派生自 DeepGEMM，使用不同的打包权重布局与路由约定，只用于 EP 部署，覆盖 EP8/EP16/EP32。

博客列举的优化里，几处能直接看到「按形状选策略」：H200 prefill 路径把 WGMMA 主循环的依赖管理批量化，提交后让一个 WGMMA 组保持 in-flight，同时进行下一块的加载与反量化，输出保持逐位一致（bit-identical）；H200 EP8 prefill 的解析器按 tok_e（每个专家的路由 token 数）选 block-M，低于约 80 时沿用基线搜索，高于该值则按填充后的专家行数与寄存器上限定尺寸；decode 路径把 MMA 操作数对调（反量化后的权重占据 MMA-M），使用 m16n8k16 并支持 4 CTAs/SM，其半静态 token 切片调度在 9～15 tokens/expert 下测得 186 µs，动态调度为 216 µs。

profile 与权重布局在模型加载时固定，运行时不切换；token 数只调整每个 profile 内部的调度。

接入方式：用 Chord 替代公共 Humming 包

Chord 以 Apache-2.0 许可开源；博客说明，文中的代码与内核表格对应 commit 7ca91d8（2026 年 9 月 14 日）。安装与启动是两条命令：

pip install git+https://github.com/novitalabs/chord.gitvllm serve <kimi-k2.x-int4-model> --quantization humming

机制在 import 名上：Chord 的分发包同时提供 chord 与 humming 两个模块根，vLLM 的懒加载门面会解析 humming.{dtypes, config, layer, schema, utils.weight} 这些模块，indexed 路径因此可以直接复用既有的 Humming 集成，在兼容的 vLLM revision 上不需要 Chord 专属的框架补丁。「兼容」有具体所指：所在分支需要已包含通用 WNA16 group-scale 支持（PR #48918，2026 年 7 月 17 日创建、8 月 19 日合入）；博客提示，更旧的分支可能需要把 group-32 的键补进 _supports_quant_scheme。

接入还有几个前提：不要和上游的 inclusionAI/humming 同时安装（Chord 有意占用这个 import 名）；要显式选择 Humming（moe_backend="humming" 或 --quantization humming），否则 vLLM 的 WNA16 自动优先级可能先选中别的后端；VLLM_HUMMING_USE_F16_ACCUM 与 VLLM_BATCH_INVARIANT 要保持关闭，因为两个后端都没有实现这些计算选项。已发布的 schema 支持 uint4、group-32、BF16 scales，以及 Kimi K2.x 使用的 compressed-tensors pack-quantized INT4 group-32 检查点格式；不支持的量化方案会在加载时报错，而不是静默选到错误的内核。

vLLM 帖文里写道：「很高兴看到这些内核与基准测试开源。」

端到端报告：TTFT 降 8.6%，吞吐升 9.6%

博客引用的一份早期服务报告，给出了 indexed TP8 路径的端到端数据。这份报告发布在 vLLM 仓库 PR #51815 的描述里（PR 于 2026 年 8 月 11 日创建，目前已关闭、未合并）。测试环境为 8×H200、Kimi-K2.6、TP8 + DCP8、FP8 KV cache，请求来自 ShareGPT，Humming 与 Chord 用的是同一条 --quantization humming 命令：

• 平均 TTFT（首 token 延迟）：2,022 ms 降至 1,849 ms（-8.6%）；
• prefill 输入+输出吞吐：20,716 tok/s 提升至 22,712 tok/s（+9.6%）；
• decode 输出吞吐：batch 8 从 483 提升至 503 tok/s（+4.1%），batch 64 从 1,650 提升至 1,740 tok/s（+5.5%），batch 128 从 2,514 提升至 2,715 tok/s（+8.0%）。

精度方面，报告给出的对照（Humming 对 Chord）是：OCRBench FINAL（1,000 题、官方配置）941 对 947；GSM8K 5-shot（lm-eval，n=500）flex 0.312 对 0.332、strict 0.238 对 0.254，结论是相对 Humming 没有精度回退。报告同时注明，这组数据只覆盖 TP8 端到端；EP8 分离式部署（PD-disaggregated）是该算子重点瞄准的形态，端到端结果未出现在这份报告中。

还没做完：grouped 集成与 B200/B300 prefill 内核

博客「下一步」清单的第一项，是完成 grouped 算子与 vLLM Humming 后端的集成，让 grouped_contiguous 与 grouped_masked 也能通过既有框架集成使用；这两套分组算子目前已经随仓库发布、有独立的算子 API，但还没有接进 vLLM 的 Humming 后端。

第二项是发布 B200/B300 的 EP8 prefill 内核。博客称，双方已有一个工作实现、内部测试性能「有前景」，计划在后续发布中把内核与基准测试一并放出；Novita 在长文里的表述是「内部已经在跑，基准测试即将公布」。

参考链接

• vLLM 帖文（@vllm_project，2026 年 9 月 16 日）
• vLLM x Novita AI 技术博客：Chord, Faster INT4 MoE for Kimi K2.x（2026 年 9 月 15 日）
• 公共 Humming 内核库（inclusionAI/humming）
• Novita AI 的 X 长文：We made INT4 MoE 2x faster. Then we open-sourced it.（2026 年 9 月 16 日）
• Chord 仓库（novitalabs/chord）
• Chord commit 7ca91d8（2026 年 9 月 14 日）
• vLLM PR #48918：[CT] Support Humming for WNA16 MoE
• vLLM PR #51815：[MoE][Humming] Enable Chord W4A16 MoE as a humming backend drop-in