Humming：vLLM 与蚂蚁集团开源通用量化推理内核库
蚂蚁集团与 vLLM 社区已开源 Humming，这是一个面向量化推理的通用高性能内核库。Humming 的重要性在于，它用统一的量化表示、JIT 编译路径和共享的 CUDA 组件取代了“每种量化格式一个内核”的方法——因此，一个库即可覆盖从 Turing 起的 NVIDIA GPU 以及 T-Head 的 Zhenwu 810E 等 CUDA 兼容设备上的稠密和 MoE 工作负载，并在许多代表性工作负载上报告了相较于 vLLM Marlin、FlashInfer 和 DeepGEMM 的显著加速。
代码可在 github.com/vllm-project/humming 获取。
为什么量化推理需要通用内核库
量化降低了权重和激活的精度，以减少内存使用和数据移动，同时也能提高计算效率。随着模型和推理工作负载的增长，内存容量、内存带宽和计算成本日益成为部署的制约因素。
这种压力在 MoE 模型中最为明显：每个 token 仅激活部分专家，但仍需存储全部权重。公告中引用的 DeepSeek-V3 拥有 6710 亿参数，每个 token 激活 370 亿参数。
不同模型和 GPU 代际对量化的需求差异很大。常见配置包括 WNA16、W8A8、W4A8 和 W4A4，其中 W 和 A 分别表示权重和激活的位宽，N 表示可变的权重位宽。即使在相同精度下，INT4 和 NVFP4 等方案在数值格式和缩放方式上也存在差异。缩放粒度、缩放表示、组大小以及对称或非对称量化进一步扩大了配置空间。
这些存储和计算上的节省只有通过高效的内核才能转化为更低的延迟和更高的吞吐量。Prefill、decode 以及不同的 MoE 专家负载对计算和内存提出了不同的要求，内核必须适应这些工作负载，同时保持转换、反量化和调度开销较低。
为每种组合开发和集成专用内核会带来日益增长的维护负担：每次新增都可能需要自己的权重预处理、内核选择、框架适配器和测试，而且优化工作必须在多个实现之间传递。Humming 通过将量化描述与执行选择分离，并从共享组件组合内核来解决这一问题。
Humming 的架构：表示、配置、编译
Humming 的执行路径包含三个部分：表示量化计算、为硬件和工作负载选择配置、从 CUDA 组件编译内核。
统一的量化表示
Humming 将不同量化工具导出的模型转换为 Humming 标准格式：一种通用的量化描述，权重、缩放因子、零点和其他张量以一致的方式组织。这种表示保留了每种方案的量化语义，同时将权重格式、缩放粒度、缩放数据类型、零点和次级缩放因子作为独立参数描述。激活具有自己的数值格式和量化设置，这些参数可以在格式和硬件约束下进行组合。
例如，分组 INT4 权重可以与 BF16 激活配对，而具有不同缩放方案的低精度浮点权重则使用相同的参数系统。这为推理框架提供了来自不同来源模型的通用接口。
针对每个工作负载的调优和 JIT 编译
预处理会验证参数、填充维度，并将张量重新打包为适合目标内核的布局。然后，Humming 将量化描述与权重形状、GEMM 类型、目标 GPU 和计算选项相结合，以配置执行。三种配置描述了内核：
- Layer Config 描述张量形状、数据类型、量化参数和目标架构。
- Compute Config 指定 GEMM 类型和计算选项，例如批量不变性和输入缩放布局。
- Tuning Config 选择 tile 形状、流水线阶段数、warp 特化、内存访问和 MMA 指令。它还控制并发和调度；TMA 可以按张量配置。
默认情况下，Humming 使用设备特定的启发式方法选择调优配置，并且可以为不同范围的输入大小准备不同的配置，并在运行时调度到适当的内核。公告指出，在小批量下，使用异步复制（cp.async）和 MMA 的多级流水线可以优于基于 TMA 和 UMMA 的流水线，而在大批量下，TMA 和 UMMA 可以提供更好的吞吐量。Humming 在同一配置和组件框架内同时支持这两种执行路径。
这三种配置专门化使用 NVRTC 或 NVCC 编译的内核，编译后的内核会被缓存以供重用。推理将所选内核与预处理后的权重和运行时激活相结合。
跨格式和架构的共享 CUDA 组件
在 CUDA 层面，Humming 将任务调度、数据移动、反量化和缩放、Tensor Core 计算以及输出处理组织为可重用组件。它提供了标准、warp 特化和专用 UMMA 流水线，这些流水线共享兼容的组件，同时使用每种执行策略所需的数据流和同步方式。
每组组件处理不同的变化来源：
- 数值处理处理量化格式。 解码和算术组件考虑位宽、整数或浮点表示、缩放因子和零点。反量化和缩放被集成到内核的计算中。
- 内存和计算后端处理 GPU 架构。 Humming 为目标架构选择数据移动和 Tensor Core 实现，在支持的地方使用异步复制、TMA 和多播。
- 任务映射和调度处理稠密和 MoE 工作负载。 每种 GEMM 类型使用适当的调度逻辑，缩放和偏置加法可以融合到输出阶段。
JIT 编译根据三种配置组合这些组件，并消除未使用的实现分支。因此，可以通过扩展相关组件来支持新格式或新架构，而对共享组件的改进将使兼容配置受益——减少了单独专用内核的重复开发和维护。
JIT 编译足够快，可用于迭代调优
Humming 的 CUDA 实现轻量级，避免了对重量级内核框架的依赖，这保持了 JIT 编译的高速度。在 H200 服务器上使用 CUDA 13.4 进行的单个 W4A16 内核编译测试中，禁用编译缓存后，编译时间从 CUTLASS 与 NVCC 的 23.35 秒 降至 Humming 与 NVCC 的 1.82 秒 或 与 NVRTC 的 1.88 秒。
Humming 还在模型预热期间并行编译内核。在报告的测试中，为单层 GEMM 编译和加载覆盖所有支持的 M 大小的内核仅需 约 2–4 秒，这使得在推理前准备多个专用内核变得切实可行。
设备特定的性能调优
Humming 使用 GEMM 配置和设备参数来优化内核。量化格式、矩阵形状、GPU 指令支持、SM 数量、共享内存和寄存器资源都会影响执行选择。
在现有 GPU 上进行高效的 W4A8-MXFP4 推理。 Blackwell 之前的架构缺乏原生 MXFP4 Tensor Core 支持，传统的 W4A8-MXFP4 内核在较大 M 时可能受到反量化和缩放开销的限制。Humming 通过一条快速的 MXFP4 到 INT8/FP8 反量化路径解决了这个问题，该路径在内核内将 MXFP4 权重转换为支持的计算格式。在这种情况下，该路径在小 M 时显著优于传统的 W4A8 实现，在大 M 时接近 W8A8 的性能，因此 H200 等 GPU 可以运行高性能的 W4A8-MXFP4 内核，同时保留 4 位权重的存储优势。
针对每个设备匹配的反量化路径。 相同的量化格式在不同 GPU 上可能受益于不同的实现。Humming 为设备选择解码、数值转换和缩放路径，并将它们与合适的内存访问和计算流水线配对，以减少指令和数据移动开销。共享组件使这些优化可在兼容的稠密和 MoE 内核中重用。
广泛的调优控制与设备特定的启发式方法。 Humming 公开了 tile 形状、流水线阶段数、warp 特化、TMA、MMA 指令、CTA 并发和 Stream-K 调度。设备特定的启发式方法将 GEMM 形状与资源约束相结合，以平衡并行性、数据重用和流水线开销，随着输入大小的变化选择不同的执行策略。开发人员可以直接使用默认配置，或针对特定模型和部署场景进行优化。
结果：设备支持
Humming 的统一架构实现了广泛的 GPU 支持。对于 NVIDIA GPU，Humming 完全支持从 Turing 起的每一代架构，为多个硬件代际带来了通用的量化计算框架。统一的 JIT 架构还使得将 Humming 扩展到其他兼容 CUDA 编程的 GPU 变得简单，Humming 已经完全支持 T-Head 的 Zhenwu 810E。
设备支持在指令层面并不统一。B300（SM103）上的 UMMA 指令不再支持 INT8 计算，但早期的 MMA 指令仍然可以执行 INT8 GEMM，尽管性能大幅降低。
结果：量化支持
Humming 对量化类型的支持主要取决于激活类型和权重类型之间的兼容性。其统一的量化表示和内核架构允许广泛的整数和浮点权重格式与不同精度的激活配合使用。
Humming 还支持一种特殊的浮点格式，其指数位为零，可表示的值与相同位宽下的对称整数相同。例如，排除 -128 的 INT8 和 float8e0m7 都表示从 -127 到 127 的整数。这种格式的权重与相同位宽的 INT 权重具有相同的激活类型兼容性，Humming 会自动将符合条件的对称整数权重转换为这种格式，以实现更快的反量化。
社区还在 Blackwell Tensor Core 中发现了两个隐藏功能：原生支持 float4e0m3 和 float8e3m4，这两种特殊浮点格式未在 PTX 类型系统中公开。Humming 在所有 Blackwell 架构上支持这些功能，包括面向工作站和个人计算的数据中心 SM100/SM103 和 SM120/SM121 设备。Rubin 不支持这两种隐藏的 Tensor Core 格式，因此这些原生计算路径的支持不会扩展到 Rubin。
结果：内核微基准测试
内核微基准测试由 vLLM 社区运行，测量 M 覆盖从 1 到 8192 的所有 2 的幂的 GEMM 性能。内核完全预热，并在冷 L2 缓存下独立计时，样本之间短暂暂停。MoE 加速使用 Gateup 和 Down GEMM 延迟之和。
WnA16
WnA16 将低比特权重与 FP16/BF16 激活相结合，减少了权重存储和内存流量，同时使用广泛可用的 16 位计算能力。W4A16 系列是 Humming 优化工作的主要焦点。
对于稠密 W4A16，Humming 在大多数测试配置中匹配或优于 vLLM Marlin，在多个 Hopper 和 Blackwell 工作负载上取得了显著提升。以 cuBLAS 作为共同基线，Humming 在多个设备的小批量下提供了大约 3–4 倍的加速。随着批量增长，在 H200 大批量工作负载上达到 约为 Marlin 性能的两倍，在某些 B300 情况下 超过三倍。
这些提升也扩展到 MoE。Humming 为 token 分布在专家之间且每个专家的工作量随负载变化的工作负载提供了高效的 Indexed 和 Grouped 路径。H20 和 H200 上的 MXFP4 W4A16 在多个中大批量情况下达到 约为 Marlin 性能的两倍。在 B300 上，MXINT4 和 NVFP4 路径在大批量下实现了更大的提升，Grouped 路径达到约 5 倍。SM120 和 SM121 上的结果也显示 Humming 在一系列工作负载上优于 Marlin 和相应的 FlashInfer 实现。
Humming 的优化不仅限于 W4A16，还扩展到其他 WnA16 组合。其共享的反量化和计算路径在不同位宽的整数和浮点权重上提供了强劲性能，包括 INT1、INT2、INT3 以及 FP4 到 FP8。SM120 上的最低位配置超过了 BF16 cuBLAS 性能的 5 倍，在 H20 的大批量下，多种位宽保持了接近 BF16 cuBLAS 的计算效率。
W8A8 / W4A8 / W4A4
随着 Hopper 的出现，具有块状缩放的 FP8 W8A8 成为高性能量化推理的主流方法。Humming 通过共同优化数据移动、缩放和专家调度来改进这一成熟路径。在 H20 和 H200 的小批量和中批量下，Humming 的 W8A8 分别达到 DeepGEMM Contiguous 性能的约 2.6 倍和 2.5 倍，Indexed 和 Grouped 布局均具有竞争力的性能。
随着更多模型采用原生 MXFP4 权重量化，将较小的权重占用转化为实际的推理收益变得越来越有价值。Humming 的 W4A8 路径在加载时快速反量化 MXFP4 权重，并使用 FP8 Tensor Core 进行计算，从而在 Hopper 上实现高效的 W4A8-MXFP4 执行。以 DeepGEMM W8A8 为基线，Humming 的 W4A8 在小批量和中批量下达到 H20 上约 4.2 倍、H200 上约 4.5 倍的基线性能，也显著优于同一图中 Humming 自身的 W8A8 结果。FlashInfer CUTLASS W4A8 也是基于 Humming 方法的实现，而 Humming 当前的实现在所示基准测试中取得了更好的性能。
Blackwell 增加了对 MXFP8 和 MXFP4 等微缩放格式的原生 Tensor Core 支持。从数据中心 SM100/SM103 设备到面向工作站和个人计算的 SM120/SM121 设备，Humming 利用每种架构的低精度计算能力，在一个内核和调优框架内支持 W8A8、W4A8 和 W4A4。在 B300（SM103）上，以相应精度的 DeepGEMM Contiguous 为基线，W8A8、W4A8 和 W4A4 在小批量和中批量下分别达到 约 1.25 倍、1.6 倍和 1.4 倍的基线性能，并在许多工作负载上优于相应的 FlashInfer 实现。Indexed 和 Grouped 路径都受益于原生低精度计算。
这些提升也扩展到 SM120 和 SM121。Humming 的 MXFP4 × MXFP8 W4A8 和 NVFP4 × NVFP4 W4A4 在大多数测试点上优于 FlashInfer CUTLASS。在 SM121 的小批量下，性能提升达到 W4A8 约 19%，W4A4 约 9%；在 SM120 上，Humming 在大多数测试大小上也处于领先地位。
端到端推理性能
Humming 的 W4A8 提升也扩展到模型服务。在八块 H20 GPU 上使用 TP=8 的 DeepSeek-V4.1-Flash 测试中，保持模型的 MXFP4 权重不变，将 MoE 执行从 Marlin W4A16（BF16 激活）切换到 Humming W4A8（tokenwise FP8 激活），端到端吞吐量提高了约 14%–43%，平均首 token 时间（TTFT）降低了约 25%–32%。吞吐量优势随并发度增长：在并发度 64 时，输出吞吐量从 442.37 上升到 633.94 tokens/s，而平均每输出 token 时间（TPOT）下降约 31%。
| 并发度 | MoE 后端 | 输入吞吐量（tokens/s）↑ | 输出吞吐量（tokens/s）↑ | TTFT（ms）↓ | TPOT（ms）↓ |
|---|---|---|---|---|---|
| 1 | vLLM Marlin W4A16 | 2827.15 | 88.30 | 543.13 | 9.24 |
| 1 | Humming W4A8 | 3226.23（+14.1%） | 100.77（+14.1%） | 383.80（-29.3%） | 8.46（-8.5%） |
| 2 | vLLM Marlin W4A16 | 4781.06 | 149.33 | 794.31 | 10.33 |
| 2 | Humming W4A8 | 5646.98（+18.1%） | 176.38（+18.1%） | 555.66（-30.0%） | 9.20（-10.9%） |
| 4 | vLLM Marlin W4A16 | 7427.64 | 231.99 | 1422.32 | 11.72 |
| 4 | Humming W4A8 | 9160.81（+23.3%） | 286.12（+23.3%） | 971.74（-31.7%） | 10.22（-12.9%） |
| 8 | vLLM Marlin W4A16 | 9190.03 | 287.04 | 2049.81 | 19.90 |
| 8 | Humming W4A8 | 12263.89（+33.4%） | 383.04（+33.4%） | 1386.91（-32.3%） | 15.49（-22.2%） |
| 16 | vLLM Marlin W4A16 | 11124.67 | 348.96 | 2463.69 | 36.28 |
| 16 | Humming W4A8 | 15171.00（+36.4%） | 475.89（+36.4%） | 1844.71（-25.1%） | 26.44（-27.1%） |
| 32 | vLLM Marlin W4A16 | 12669.43 | 398.83 | 3546.74 | 66.44 |
| 32 | Humming W4A8 | 17804.79（+40.5%） | 560.49（+40.5%） | 2465.47（-30.5%） | 47.48（-28.5%） |
| 64 | vLLM Marlin W4A16 | 13937.42 | 442.37 | 5467.35 | 123.25 |
| 64 | Humming W4A8 | 19973.00（+43.3%） | 633.94（+43.3%） | 3882.94（-29.0%） | 85.70（-30.5%） |
安装
Humming 可从 PyPI 安装：
pip install humming-kernels
Humming 使用 CUDA 编译器按需生成内核。如果系统上没有可用的 CUDA 编译器，请选择以下安装选项之一，以包含适合您环境 CUDA 主版本的编译依赖项：
pip install "humming-kernels[cu12]" # CUDA 12
pip install "humming-kernels[cu13]" # CUDA 13
Humming 非常轻量，依赖简单，wheel 包仅约 500 KB。按需编译避免捆绑大量预编译内核。原则上，任何能够运行 PyTorch 2.7 的环境都可以运行 Humming。
在 vLLM 中使用 Humming
Humming 已深度集成到 vLLM 中。从 vLLM 0.29.0 开始，CUDA 安装将 Humming 作为依赖项包含在内，因此无需单独安装 Humming。Humming 支持从 Turing 起（SM75 及以上）所有架构的 NVIDIA GPU。
--linear-backend 和 --moe-backend 选项分别选择 Humming 作为稠密 GEMM 和 MoE GEMM 的计算后端。例如，使用张量并行大小 8 为两者启用 Humming：
vllm serve deepseek-ai/DeepSeek-V4.1-Flash-1017 \
--tensor-parallel-size 8 \
--linear-backend humming --moe-backend humming
--quantization humming 选项仅作为 Humming 支持但 vLLM 不再支持或尚未原生支持的格式（例如 GPTQ INT3）的兼容路径：
vllm serve hfl/Qwen2.5-VL-3B-Instruct-GPTQ-Int3 --quantization humming
选择激活类型
默认情况下，Humming 使用检查点中指定的激活类型，并适应设备执行。如果 GPU 不支持该计算类型，Humming 会回退到支持的高精度类型。例如，没有原生 MXFP4 支持的设备可以自动使用 FP8、BF16 或 FP16 计算路径，从而允许相同的量化检查点跨 GPU 代际运行。
用户可以通过 vLLM 的量化配置接口覆盖默认设置，为其部署选择激活类型。以下示例将 MoE 层的激活类型设置为 MXFP8：
vllm serve openai/gpt-oss-20b --moe-backend humming --quantization-config.moe.activation mxfp8
未来计划
更广泛、更深入的设备支持
Humming 将继续扩展硬件覆盖范围，同时适应新设备的计算和内存能力。已声明的优先事项包括：
- 支持更多 Rubin 架构特性，包括 Locality Domains，以进一步优化任务映射和数据访问。
- 将支持扩展到更多 T-Head PPU 设备，包括 Zhenwu M890。
- 支持更多兼容 CUDA 编程的设备，将现有量化格式、内核组件和调优能力带到更多硬件平台。
添加自动调优
通过有针对性的调优，Humming 在大多数场景中已经可以达到最先进的性能。然而，模型形状、量化配置和设备资源的广泛范围使得启发式方法难以在每种情况下选择最佳配置。因此，Humming 计划引入自动调优，在目标设备上对候选配置进行基准测试，为每个工作负载找到更好的执行参数。自动调优将补充现有的启发式方法，减少手动配置搜索。
与推理框架更深入的集成
Humming 将继续深化与 vLLM 等推理框架的集成，扩展量化模型加载、后端选择、并行推理、运行时调度和设备适配方面的支持。已声明的目标是让更多量化方法适用于框架的全部功能，包括专家并行（EP）的量化 All-to-All（A2A）通信等优化。
致谢
公告感谢蚂蚁集团的林金镇、顾泽凯和黄志林，以及 Michael Goin、Roberto L. Castro、Chauncey Jiang 和 vLLM 社区的其他成员在 Humming 开发过程中的贡献和支持。它还感谢 vLLM 社区提供的开发、测试和 CI 资源，用于在广泛的硬件和工作负载上验证和改进 Humming。
Humming 的设计和优化借鉴了 Marlin、vLLM Marlin、LMDeploy、DeepGEMM、FlashInfer、CUTLASS 和其他开源项目的灵感。