蚂蚁集团与 vLLM 社区开源量化推理内核库 Humming

蚂蚁集团与 vLLM 社区正在开源 Humming，这是一个面向量化推理的通用高性能内核库。Humming 具备以下特性：

• 广泛的设备支持：支持从 Turing 起的 NVIDIA GPU 架构，以及平头哥镇岳 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 的执行路径分为三部分：表示量化计算、为硬件和工作负载选择配置，以及从 CUDA 组件编译内核。

统一的量化表示

不同量化工具导出的模型在数值格式、参数组织和权重布局上各不相同。Humming 将支持的导出格式转换为 Humming 标准格式：一种通用的量化描述，权重、缩放、零点和其他张量均按一致方式组织。

该表示在保留每种方案量化语义的同时，将权重格式、缩放粒度、缩放数据类型、零点和二级缩放描述为独立参数。激活有其自身的数值格式和量化设置。这些参数可在格式和硬件约束下进行组合。

例如，分组 INT4 权重可以与 BF16 激活值配对使用。采用不同缩放方案的低精度浮点权重使用同一套参数体系。这为推理框架提供了面向不同来源模型的统一接口。

针对每种工作负载的调优与 JIT 编译

标准格式为各种导出格式提供了统一的表示。随后，预处理会校验参数、填充维度，并将这些张量重新打包为目标内核所适配的布局。Humming 将量化描述与权重形状、GEMM 类型、目标 GPU 以及计算选项结合起来，以配置执行过程。

三种配置描述内核：

• 层配置描述张量形状、数据类型、量化参数以及目标架构。
• 计算配置指定 GEMM 类型和计算选项，例如批次不变性和输入缩放布局。
• 调优配置选择瓦片形状、流水线阶段数、warp 特化、内存访问以及 MMA 指令。它还控制并发与调度；TMA 可以按张量进行配置。

默认情况下，Humming 使用针对具体设备的启发式规则来选择调优配置。它可以为不同范围的输入尺寸准备不同的配置，并在运行时分派到合适的内核。

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 编译会根据这三种配置组合这些组件，并消除未使用的实现分支。因此，可以通过扩展相关组件来支持新的格式或架构。对共享组件的改进会使兼容的配置受益，从而减少对各自专用内核的重复开发和维护。

这一 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 已全面支持平头哥镇岳 810E，将其共享内核与调优能力带到更多硬件平台。

下图展示了 Humming 在各设备上支持的高性能激活类型：这些类型能够充分利用设备的峰值计算能力。图中省略的部分类型仍可通过其他指令路径运行。例如，B300（SM103）上的 UMMA 指令不再支持 INT8 计算，但更早的 MMA 指令仍可执行 INT8 GEMM，只是性能大幅降低。

量化支持

Humming 对量化类型的支持主要取决于激活类型与权重类型之间的兼容性。其统一的量化表示与内核架构使多种整数和浮点权重格式能够与不同精度的激活配合使用，为模型精度和部署性能提供灵活选择。

下图展示了 Humming 支持的权重与激活类型组合。行表示权重类型，列表示激活类型，勾选表示支持的组合。结合上文的设备支持矩阵，可以勾勒出各硬件架构上可用的高性能量化计算路径。

Humming 还支持一种指数位为零的特殊浮点格式，其可表示的值与相同位宽的对称整数一致。例如，排除 -128 和 float8e0m7 的 INT8 与二者都表示从 -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 位计算能力。这使其适用于许多 GPU 世代和模型部署场景。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 上的 W8A8 结果。在小批量和中等批量尺寸下，Humming 分别达到 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 结果。权重流量的减少转化为更低的 GEMM 延迟。FlashInfer CUTLASS W4A8 也是基于 Humming 方法的实现，而 Humming 当前的实现在此处展示的基准测试中取得更好的性能。

Blackwell 为 MXFP8 和 MXFP4 等微缩放格式增加了原生 Tensor Core 支持。从数据中心 SM100/SM103 设备到面向工作站和个人计算的 SM120/SM121 设备，Humming 利用各架构的低精度计算能力，在单一内核和调优框架内支持 W8A8、W4A8 和 W4A4。

在 B300（SM103）上，Humming 在全部三种 MX 格式组合下均展现出显著的性能提升。下方每个面板都以对应精度的 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（逐 token FP8 激活），端到端吞吐提升约 14%–43%，平均首 token 时延（TTFT）降低约 25%–32%。吞吐优势随并发增加而扩大：在并发 64 时，输出吞吐从 442.37 tokens/s 升至 633.94 tokens/s，而平均每输出 token 耗时（TPOT）下降约 31%。

用法

安装

使用 pip 安装 Humming：

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 起所有架构的 NVIDIA GPU（SM75 及以上）。

--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。例如，通过此兼容路径加载 GPTQ INT3 检查点：

vllm serve hfl/Qwen2.5-VL-3B-Instruct-GPTQ-Int3 --quantization humming

选择激活类型

默认情况下，Humming 使用 checkpoint 中指定的激活类型，并根据设备调整执行方式。如果 GPU 不支持该计算类型，Humming 会回退到受支持的更高精度类型。例如，不支持原生 MXFP4 的设备可以自动使用 FP8、BF16 或 FP16 计算路径，从而让同一个量化 checkpoint 跨 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 编程的设备，将现有量化格式、kernel 组件和调优能力带到更多硬件平台。

添加 Autotune

通过针对性调优，Humming 在大多数场景下已经能够达到最先进的性能。然而，模型形状、量化配置和设备资源的范围很广，使得启发式方法难以在每种情况下都选出最优配置。

因此，Humming 计划引入 autotune，在目标设备上对候选配置进行基准测试，为每个工作负载找到更好的执行参数。Autotune 将补充现有的启发式方法，减少手动配置搜索，帮助更多工作负载接近 kernel 可达到的最佳性能。

与推理框架更深度集成

Humming 将继续深化与 vLLM 等推理框架的集成，扩大在量化模型加载、后端选择、并行推理、运行时调度和设备适配等方面的支持。我们的目标是让更多量化方法能够配合框架的完整能力使用，包括面向专家并行（EP）的量化 All-to-All（A2A）通信等优化。这些工作将使高性能量化 kernel 在多样化部署场景中更易于使用，并帮助将 kernel 性能提升转化为端到端推理收益。

致谢

我们感谢蚂蚁集团的 Jinzhen Lin、Zekai Gu 和 Zhilin Huang，以及 Michael Goin、Roberto L. Castro、Chauncey Jiang 和 vLLM 社区的其他成员，感谢他们在 Humming 开发过程中的贡献和支持。

我们还感谢 vLLM 社区提供开发、测试和 CI 资源，使我们能够在广泛的硬件和工作负载上验证和改进 Humming。

Humming 的设计和优化从 Marlin、vLLM Marlin、LMDeploy、DeepGEMM、FlashInfer、CUTLASS 及其他开源项目中获得灵感。我们感谢这些项目的作者和贡献者分享他们的想法和实现，帮助推动整个社区的高性能量化推理发展。