AREX Feed Article
vLLM 与蚂蚁集团开源量化内核库 Humming：宣称内核最高提速 5 倍、端到端 MoE 吞吐提升 14%～43%
“量化很容易。为每种格式 × 每种 GPU × 每种工作负载都提供快速量化内核，并不容易。”北京时间 2026 年 10 月 21 日 12 时 17 分，开源大模型推理引擎 vLLM 的官方 X 账号，与蚂蚁集团联合开源面向量化推理的统一内核库 Humming，并称其为“one kernel library for all of it”——一个内核库包住全部组合，同时给出两组核心数字：内核最高提速 5 倍，端到端 MoE（混合专家）吞吐提升 14%～43%。
帖文所链当天发布，已经可以访问，PyPI 也已能装到 。吞吐数字的出处是一条端到端测试：在 8 张 H20、张量并行 TP=8 的环境测 DeepSeek-V4.1-Flash，保持模型 MXFP4 权重不变，只把 MoE 执行从 vLLM Marlin W4A16（BF16 激活）换成 Humming W4A8（逐 token 的 FP8 激活），端到端吞吐提升约 14%～43%，平均首 token 延迟（TTFT）下降约 25%～32%。这些成绩均出自 vLLM 与蚂蚁一方自测，配置、版本与完整数据表都附在博客里。
组合爆炸：同一位宽，换个格式就是另一套内核
模型和推理负载变大后，显存容量、显存带宽与算力成本先后顶到上限，MoE 尤其明显：每个 token 只激活一部分专家，全部权重却都得驻留显存。博客举的例子是 DeepSeek-V3——6710 亿参数，每个 token 只激活 370 亿。量化把权重和激活降到低位宽，省显存、省数据搬运，也可能提高计算效率。
问题出在组合空间。博客列出的常见配置就有 WNA16、W8A8、W4A8、W4A4（W、A 分别指权重与激活的位宽，N 表示可变权重位宽）；同为 4bit，INT4 和 NVFP4 的数值格式与缩放方式并不相同；缩放粒度、scale（缩放系数）的数据类型、group size（分组大小）、对称或非对称量化，每往下一层就多出一批组合。prefill（预填充）、decode（解码）和不同 MoE 专家负载对计算与访存的偏好又各不一样，内核还得把格式转换、反量化和调度的开销压住。博客的判断是：为每一种组合单独写一套内核，各自的权重预处理、内核选择、框架适配与测试都要单独维护，优化也无法在多套实现之间复用。
统一表示加 JIT 编译：23.35 秒的编译缩到 1.82 秒
Humming 把“量化怎么描述”和“内核怎么执行”拆开。不同量化工具导出的模型先统一转成 Humming 标准格式：权重、scale、zero point（零点）等张量按一致方式组织，量化语义原样保留；权重格式、缩放粒度、scale 数据类型、zero point、二级 scale 都是独立参数，激活另有自己的数值格式与量化设置——groupwise INT4 权重可以直接配 BF16 激活，不同缩放方案的低比特浮点权重走同一套参数系统。
执行由三份配置描述。Layer Config 管张量形状、数据类型、量化参数和目标架构；Compute Config 管 GEMM（通用矩阵乘）类型以及 batch invariance（批次不变性）、输入 scale 布局等计算选项；Tuning Config 管 tile 形状、流水线级数、warp specialization、访存与 MMA 指令，TMA 可以按张量配置。默认由设备相关启发式选择，也可以按输入规模区间准备多套配置、运行时分发到对应内核——小批量下 cp.async 加 MMA 的多级流水线可能更快，大批量换 TMA 加 UMMA，两条路径在同一框架里共存。
CUDA 层面，任务调度、数据搬运、反量化与缩放、Tensor Core 计算、输出处理被拆成可复用组件：数值处理组件对应量化格式，访存与计算后端对应 GPU 架构，任务映射与调度对应稠密（dense）和 MoE 两类负载。内核由 NVRTC 或 NVCC 即时编译（JIT），按三份配置组合组件、剪掉用不到的实现分支，编译产物进缓存复用。
编译速度是这套架构的直接收益。博客的测试条件是 H200 服务器、CUDA 13.4、关闭编译缓存：编译单个 W4A16 内核，CUTLASS 加 NVCC 要 23.35 秒，Humming 加 NVCC 只要 1.82 秒，走 NVRTC 为 1.88 秒；模型 warmup（预热）阶段并行编译，为单层 GEMM 备齐覆盖所有支持 M 尺寸的内核并加载，约需 2～4 秒。整个库只有一个约 500 KB 的 wheel 安装包（PyPI 上 x86_64 wheel 文件大小为 517.2 kB），依赖只有 PyTorch 和 CUDA 编译器，原则上能跑 PyTorch 2.7 的环境就能运行。
官方基准拆解：“最高 5 倍”和“+43%”分别从哪来
方法先交代：内核微基准由 vLLM 社区测量，M 取 1～8,192 的全部 2 的幂，内核充分预热、样本间以冷 L2 缓存逐点计时；MoE 加速比取 Gateup 与 Down 两个 GEMM 延迟之和，图中配置为 256 个专家、top_k=8、BF16 激活，图注版本为 Humming 0.1.18、vLLM 0.32.0、FlashInfer 0.7.2。
dense W4A16 上，Humming 在多数受测配置追平或超过 vLLM Marlin：小批量在多款设备上比公共基线 cuBLAS 快约 3～4 倍；大批量在 H200 上约为 Marlin 的 2 倍，B300 部分场景超过 3 倍。MoE 收益更大：MXFP4 W4A16 在 H20、H200 的中大批量约为 Marlin 的 2 倍；B300 上 MXINT4 与 NVFP4 在大批量进一步拉开，Grouped 路径约 5 倍——帖文“最高 5 倍内核提速”对应的就是这类配置。SM120 与 SM121 上，Humming 在一批工作负载中同时超过 Marlin 与对应 FlashInfer 实现。WnA16 家族还覆盖 INT1～INT3 与 FP4～FP8，SM120 上最低比特配置超过 BF16 cuBLAS 的 5 倍。
W8A8、W4A8、W4A4 的数字是：H20 与 H200 小中批量的 FP8 W8A8 达到 DeepGEMM Contiguous 的约 2.6 倍与 2.5 倍；W4A8-MXFP4 在 Hopper 上靠内核内快速反量化，小中批量达到 DeepGEMM W8A8 基线的约 4.2 倍（H20）、4.5 倍（H200）；B300 的三种 MX 组合约为对应精度基线的 1.25、1.6、1.4 倍；SM121 小批量比 FlashInfer CUTLASS 高约 19%（W4A8）与 9%（W4A4）。博客还写明一个细节：FlashInfer 的 CUTLASS W4A8 同样是基于 Humming 方法的实现，而在图示基准中，Humming 当前实现更快。
端到端这条线，博客只测了一组设定：8 张 H20（TP=8）上的 DeepSeek-V4.1-Flash，MXFP4 权重不动，MoE 后端从 Marlin W4A16 换成 Humming W4A8。输出吞吐增益随并发走高，从并发 1 的 +14.1% 到并发 64 的 +43.3%：
| 并发 | Marlin W4A16 输出吞吐（tokens/s） | Humming W4A8 输出吞吐（tokens/s） | 增幅 | TTFT（ms） |
|---|---|---|---|---|
| 1 | 88.30 | 100.77 | +14.1% | 543.13 → 383.80（−29.3%） |
| 16 | 348.96 | 475.89 | +36.4% | 2,463.69 → 1,844.71（−25.1%） |
| 64 | 442.37 | 633.94 | +43.3% | 5,467.35 → 3,882.94（−29.0%） |
并发 64 时，平均每输出 token 延迟（TPOT）从 123.25 ms 降到 85.70 ms，约 −30.5%。
覆盖范围：Turing 起的每一代 NVIDIA GPU，外加 Zhenwu 810E
设备支持是“每种 GPU”承诺的具体内容：NVIDIA 侧从 Turing（SM75）起完整支持每一代架构；GitHub 支持矩阵按激活类型列出门槛——FP16 激活 SM75+、BF16 SM80+、FP8 SM89+、FP4 激活 SM120+、INT8 SM75+、INT4 SM80+。统一 JIT 架构也让移植到其他 CUDA 兼容设备不必重写内核：T-Head（平头哥）的 Zhenwu 810E 已获完整支持，共享同一套内核与调优能力。
老卡也有专门路径。Blackwell 之前的架构没有 MXFP4 Tensor Core 原生支持，传统 W4A8-MXFP4 内核在较大 M 时容易被反量化与缩放开销卡住；Humming 在内核内把 MXFP4 权重快速转成 INT8/FP8 再进 Tensor Core，H200 这类 Hopper 卡因此能跑高性能 W4A8-MXFP4，小批量明显快于常规 W4A8 实现，大批量接近 W8A8，同时保住 4bit 权重的显存收益。
博客还提到社区此前在 Blackwell Tensor Core 里找到的两个隐藏浮点格式：float4e0m3 与 float8e3m4，未在 PTX 类型系统中暴露；Humming 已把它们接入统一量化框架，覆盖 SM100/SM103 数据中心卡与 SM120/SM121 工作站及个人计算设备。这条路径不会延伸到下一代 Rubin——博客明确写了 Rubin 不支持这两个隐藏格式。
上手：pip 一条命令，vLLM 0.29.0 起自带依赖
安装是 pip install humming-kernels 一条命令；系统里没有 CUDA 编译器时，按 CUDA 12 或 13 加装 cu12、cu13 可选依赖组（extra）。vLLM 0.29.0 起，CUDA 安装默认把 Humming 作为依赖带上。在 vLLM 里用 --linear-backend humming 与 --moe-backend humming 分别为 dense GEMM 和 MoE GEMM 指定后端，例如以张量并行 8 启动 DeepSeek-V4.1-Flash；--quantization humming 是另一条兼容路径，服务于 vLLM 原生不再支持或尚未支持的格式，博客给的示例是加载 GPTQ INT3 的 Qwen2.5-VL checkpoint（模型检查点）。激活类型默认跟随 checkpoint，设备不支持所选类型时自动回退到更高精度的计算路径（比如无原生 MXFP4 支持的卡落到 FP8、BF16 或 FP16），也可以显式指定——示例命令把 gpt-oss-20b 的 MoE 层激活设为 MXFP8。
自测数字的边界，与博客开出的下一步
PyPI 显示 0.1.18 于 10 月 21 日上传、维护者为 linjinzhen；博客致谢名单里，蚂蚁集团侧是 Jinzhen Lin、Zekai Gu、Zhilin Huang，vLLM 社区侧是 Michael Goin、Roberto L. Castro、Chauncey Jiang 等人。项目自述设计受益于 Marlin、vLLM Marlin、LMDeploy、DeepGEMM、CUTLASS、FlashInfer 等开源项目。
数字的边界需要按博客原文摆清楚：帖文给出的“最高 5 倍”与“+14%～43%”各自对应前文拆解过的特定配置与单一测试设定，不是全场景的普遍加速。全部成绩均为项目方自测口径，核对它们要依靠博客公开的配置、版本与图表，更广硬件与负载上的检验要等社区动手。
博客列出的下一步都还是计划：适配更多 Rubin 特性（含 Locality Domains）、支持更多 T-Head PPU 设备（含 Zhenwu M890）、引入 autotune（自动调优）在目标设备上实测候选配置、面向专家并行的量化 A2A（All-to-All）通信。至于这两组自测数字能否在第三方硬件与负载上复现，是 Humming 开源后悬着的下一步。