Humming：vLLM与蚂蚁集团开源通用量化推理内核库
蚂蚁集团与vLLM社区开源了Humming，这是一个面向量化推理的通用高性能内核库。Humming的重要性在于，它用统一的量化表示、JIT编译路径和共享CUDA组件取代了“一种量化格式对应一个内核”的方法——因此，一个库即可覆盖从Turing起的所有NVIDIA GPU以及兼容CUDA的设备（如T-Head的Zhenwu 810E）上的稠密和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类型和计算选项，例如batch不变性和输入缩放布局。
- Tuning Config选择tile形状、流水线阶段数、warp specialization、内存访问和MMA指令。它还控制并发和调度；TMA可按张量配置。
默认情况下，Humming使用设备特定的启发式方法选择调优配置，并可为输入大小范围准备不同配置，在运行时调度到适当的内核。公告指出，在小batch size下，使用异步拷贝（cp.async）和MMA的多级流水线可能优于基于TMA和UMMA的流水线，而在较大batch size下，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秒。
Humning还在模型预热期间并行编译内核。在报告的测试中，编译并加载覆盖单层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 specialization、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在小batch size的几个设备上提供约3–4倍的加速。随着batch size的增长，在H200大batch工作负载上达到约Marlin性能的2倍，在部分B300情况下超过其性能的3倍。
这些增益扩展到MoE。Humming为token分布在专家之间且每个专家的工作量随负载变化的工作负载提供高效的Indexed和Grouped路径。H20和H200上的MXFP4 W4A16在多个中等至大batch情况下达到约Marlin性能的2倍。在B300上，MXINT4和NVFP4路径在大batch size下取得更大增益，Grouped路径达到约5倍。SM120和SM121上的结果也显示Humming在多种工作负载上优于Marlin和相应的FlashInfer实现。
Humming的优化不仅限于W4A16，还扩展到其他WnA16组合。其共享的反量化和计算路径在不同位宽的整数和浮点权重上提供强大性能，包括INT1、INT2、INT3以及FP4到FP8。SM120上的最低位配置超过BF16 cuBLAS性能的5倍，在H20的大batch size下，多种位宽保持接近BF16 cuBLAS的计算效率。
W8A8 / W4A8 / W4A4
随着Hopper的出现，具有blockwise缩放的FP8 W8A8成为高性能量化推理的主流方法。Humming通过共同优化数据移动、缩放和专家调度来改进这一成熟路径。在H20和H200的中小batch size下，Humming的W8A8分别达到约DeepGEMM Contiguous性能的2.6倍和2.5倍，Indexed和Grouped布局均具有竞争力。
随着更多模型采用原生MXFP4权重量化，将较小的权重占用转化为实际推理增益变得越来越有价值。Humming的W4A8路径在加载时快速反量化MXFP4权重，并使用FP8 Tensor Core进行计算，从而在Hopper上实现高效的W4A8-MXFP4执行。以DeepGEMM W8A8为基线，Humming的W4A8在中小batch size下达到约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在中小batch size下分别达到约基线性能的1.25倍、1.6倍和1.4倍，并在许多工作负载上优于相应的FlashInfer实现。Indexed和Grouped路径都受益于原生低精度计算。
这些增益也扩展到SM120和SM121。Humming的MXFP4 × MXFP8 W4A8和NVFP4 × NVFP4 W4A4在大多数测试点上优于FlashInfer CUTLASS。在SM121的小batch size下，性能增益分别达到约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和其他开源项目的灵感。
总结
蚂蚁集团与vLLM社区开源了Humming，这是一个面向量化推理的JIT编译内核库，支持从Turing起的NVIDIA GPU以及兼容CUDA的设备（如T-Head的Zhenwu 810E），并在稠密和MoE工作负载上报告了相较于vLLM Marlin、FlashInfer和DeepGEMM的加速。
标题
Humming：vLLM与蚂蚁集团开源通用量化推理内核库