蚂蚁开源 Humming：量化推理内核库实测，vLLM 端到端吞吐最高提升 43%
开场：量化格式这么多，为什么内核还是各写各的
量化推理这几年最大的变化，不是某个格式赢了，而是”格式太多、谁也没法一家通吃”：W4A16、W8A8、W4A8、W4A4，MXFP4、NVFP4，GPTQ、AWQ……每家量化工具导出的权重格式、参数组织、权重布局都不一样，内核库只能一个格式配一套实现，编译还越来越重。
10 月 21 日，蚂蚁集团和 vLLM 社区开源了 Humming——一个面向量化推理的通用高性能内核库，思路不是”再多写一个内核”，而是把量化描述和执行选择拆开：量化格式统一成一套标准表示，执行侧用可复用的 CUDA 组件按需拼内核。它同时是 vLLM 官方博客发布的重点项目，仓库在 vllm-project/humming。
Humming 是什么：统一表示 + 配置选择 + JIT 编译
Humming 的执行路径分三步：把量化计算表示出来、为当前硬件和负载选配置、用 CUDA 组件把内核编译出来。
不同工具导出的模型，先统一转换成 Humming 标准格式——一套通用的量化描述，权重、scale、零点等张量组织方式保持一致；预处理再校验参数、补齐维度、按目标内核重排布局。量化描述与矩阵形状、GEMM 类型、目标 GPU、计算选项组合在一起，决定最终执行方式。
▲ Humming 把各家导出格式统一为标准量化表示，参数、scale、零点按统一结构组织
默认情况下，Humming 用设备专属启发式挑选调优配置，可以按输入规模区间准备不同配置、运行时分发到对应内核。小 batch 时，基于异步拷贝（cp.async）+ MMA 的多级流水线更快；大 batch 时，TMA + UMMA 流水线吞吐更好——两种执行路径都放进同一套配置与组件框架里，按负载自动选。
▲ 表示、选配置、编译三步：同一框架内容纳 cp.async/MMA 与 TMA/UMMA 两条流水线
在 CUDA 层面，Humming 把任务调度、数据搬运、反量化与缩放、Tensor Core 计算、输出处理拆成可复用组件，提供标准、warp 特化、专用 UMMA 三种流水线。
▲ 内存与计算后端按 GPU 架构选择实现，支持异步拷贝、TMA、组播
编译速度：从 23 秒到 2 秒以内
这是 Humming 最”工程向”的卖点之一。在 H200 服务器、CUDA 13.4、关闭编译缓存的条件下，单个 W4A16 内核编译耗时：CUTLASS + NVCC 需要 23.35 秒，Humming + NVCC 只要 1.82 秒，Humming + NVRTC 为 1.88 秒。
它的 CUDA 实现刻意做得轻量，不依赖重型内核框架，所以 JIT 编译天然快。模型 warmup 阶段还能并行编译：实测把单层 GEMM 全部 M 尺寸的内核编译并加载完，只需要约 2–4 秒——为新负载准备多个专用内核的成本被压到可用级，迭代调优也变得现实。
硬件支持：Turing 起全架构，含国产芯片
NVIDIA 侧，Humming 完整支持从 Turing（SM75）开始的所有架构，等于给多代硬件统一了一套量化计算框架。
更关键的是扩展性：统一 JIT 架构让”适配一块 CUDA 兼容芯片”变成工程问题而非重写问题——Humming 已经完整支持平头哥真武 810E，路线图里还包括真武 M890 以及更多 CUDA 兼容设备。
高性能激活类型支持矩阵（原文）显示各设备能用满峰值算力的类型各不相同：例如 B300（SM103）的 UMMA 指令已不再支持 INT8，但早期 MMA 指令仍可跑 INT8 GEMM，只是性能明显更低。
量化类型上，Humming 支持 WnA16（含 INT1/INT2/INT3 到 FP4–FP8 权重）、W8A8、W4A8、W4A4 等组合，还支持一种”零指数位”特殊浮点格式（如 float8e0m7），可表示范围与同位宽对称整数一致，符合条件的对称整数权重会被自动转换以加速反量化。社区还挖出 Blackwell Tensor Core 的两个隐藏能力——PTX 类型系统未暴露的 float4e0m3 和 float8e3m4，Humming 在全部 Blackwell 架构（SM100/103/120/121）上支持它们；Rubin 不支持这两个格式。
性能实测：内核层与端到端
内核微基准由 vLLM 社区跑：M 覆盖 1 到 8192 的 2 的幂，内核充分预热、L2 冷缓存独立计时。
W4A16 稠密（基线 cuBLAS）：小 batch 多台设备约 3–4 倍加速；大 batch 下在 H200 上约为 Marlin 的 2 倍，部分 B300 场景超过 3 倍。
WnA16 全家族：小中 batch 普遍有明显收益，SM120 上最低位宽配置超过 BF16 cuBLAS 的 5 倍；H20 大 batch 下多个位宽接近 BF16 cuBLAS 效率。
W8A8 / W4A8（Hopper）：W8A8 块级缩放是 Hopper 上的主流路径，小中 batch 下 Humming 约为 DeepGEMM Contiguous 的 2.6×（H20）和 2.5×（H200）。纯 MXFP4 权重路径更猛：小中 batch 约 4.2×（H20）、4.5×（H200）——靠的是把 MXFP4 快速转成 INT8/FP8 计算格式，让没有原生 MXFP4 Tensor Core 的老架构也能吃到 4-bit 权重的存储红利。
Blackwell MX 路径：B300（SM103）上 W8A8/W4A8/W4A4 分别约 1.25×/1.6×/1.4× 基线；SM121 小 batch 下 W4A8 约 +19%、W4A4 约 +9%（对 FlashInfer CUTLASS）。
▲ MoE 路径：H20/H200 上 MXFP4 W4A16 多个中大型 batch 约为 Marlin 的 2 倍；B300 上 Grouped 路径最高约 5×
端到端是最有说服力的一段：8× H20、TP=8 跑 DeepSeek-V4.1-Flash，模型 MXFP4 权重不动，只把 MoE 执行从 vLLM Marlin W4A16 切到 Humming W4A8（tokenwise FP8 激活），吞吐与延迟全表改善：
| 并发 | 后端 | 输出吞吐 (tokens/s) ↑ | TTFT (ms) ↓ | TPOT (ms) ↓ |
|---|---|---|---|---|
| 1 | Marlin W4A16 | 88.30 | 543.13 | 9.24 |
| 1 | **Humming W4A8** | **100.77 (+14.1%)** | **383.80 (-29.3%)** | 8.46 (-8.5%) |
| 8 | Marlin W4A16 | 287.04 | 2049.81 | 19.90 |
| 8 | **Humming W4A8** | **383.04 (+33.4%)** | **1386.91 (-32.3%)** | 15.49 (-22.2%) |
| 32 | Marlin W4A16 | 398.83 | 3546.74 | 66.44 |
| 32 | **Humming W4A8** | **560.49 (+40.5%)** | **2465.47 (-30.5%)** | 47.48 (-28.5%) |
| 64 | Marlin W4A16 | 442.37 | 5467.35 | 123.25 |
| 64 | **Humming W4A8** | **633.94 (+43.3%)** | **3882.94 (-29.0%)** | **85.70 (-30.5%)** |
并发越高收益越大：端到端吞吐提升 14%–43%，平均 TTFT 降低 25%–32%，并发 64 时 TPOT 降约 31%。
怎么用：pip 装包，vLLM 两个参数
pip install humming-kernels # 或 humming-kernels[cu12] / [cu13]wheel 包只有约 500 KB，按需编译不预打包内核；原则上能跑 PyTorch 2.7 的环境就能跑它。vLLM 0.29.0 起 CUDA 安装已把 Humming 作为依赖，两个参数启用：
vllm serve deepseek-ai/DeepSeek-V4.1-Flash-1017 \
--tensor-parallel-size 8 \
--linear-backend humming --moe-backend humming--quantization humming 仅作为兼容通道（如 GPTQ INT3 这类 vLLM 不再原生支持的格式）。默认按 checkpoint 声明的激活类型执行，设备不支持时自动回退到更高精度路径；也可通过 --quantization-config.moe.activation 手动指定。
