蚂蚁集团与 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 类型和计算选项，例如批次不变性和输入缩放布局。
• 调优配置选择 tile 形状、流水线阶段数、warp 特化、内存访问和 MMA 指令。它还控制并发和调度；TMA 可以按张量进行配置。

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

Humming 的统一架构使得调优执行参数、为不同工作负载选择流水线变得简单直接。在小 batch size 下，使用异步拷贝（cp.async）和 MMA 的多级流水线可以胜过围绕 TMA 和 UMMA 构建的流水线。在较大的 batch size 下，TMA 和 UMMA 可以提供更好的吞吐量。Humming 在同一配置和组件框架内同时容纳这两种执行路径，使其能够为每种工作负载选择更有效的策略。

这三种配置共同特化了用 NVRTC 或 NVCC 编译的 kernel。编译后的 kernel 会被缓存以便复用。推理将选定的 kernel 与预处理后的权重和运行时激活值结合起来。

跨格式和架构共享的 CUDA 组件

在 CUDA 层面，Humming 将任务调度、数据移动、反量化和缩放、Tensor Core 计算以及输出处理组织为可复用的组件。它提供标准流水线、warp 特化流水线和专用 UMMA 流水线。这些流水线共享兼容的组件，同时采用各自执行策略所需的数据流和同步方式。

每组组件处理不同来源的差异：

• 数值处理负责量化格式。解码和算术组件负责位宽、整数或浮点表示、缩放因子和零点。反量化和缩放被集成到 kernel 的计算中。
• 内存和计算后端负责 GPU 架构。Humming 为目标架构选择数据移动和 Tensor Core 实现，在支持的地方使用异步拷贝、TMA 和多播。
• 任务映射和调度负责稠密和 MoE 工作负载。每种 GEMM 类型都使用相应的调度逻辑。缩放和偏置加法也可以融合到输出阶段。

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

这个 CUDA 实现轻量，避免了重量级内核框架依赖，同时也保持了 JIT 编译的快速。在我们于搭载 CUDA 13.4 的 H200 服务器上进行的单 W4A16 内核编译测试中，编译时间从使用 NVCC 的 CUTLASS 的 23.35 秒降至使用 NVCC 的 Humming 的 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 在若干中大型 batch 场景下达到了约 Marlin 两倍的性能。在 B300 上，MXINT4 和 NVFP4 路径在大 batch 尺寸下取得了更大的收益，其中 Grouped 路径达到约 5 倍。SM120 和 SM121 上的结果也显示，Humming 在一系列工作负载中优于 Marlin 及对应的 FlashInfer 实现。

Humming 的优化还从 W4A16 扩展到了其他 WnA16 组合。其共享的反量化与计算路径为不同位宽的整数和浮点权重带来了强劲性能，包括 INT1、INT2、INT3 以及 FP4 到 FP8。下图展示了这些格式在中小 batch 尺寸下的广泛收益，其中 SM120 上的最低位宽配置超过 BF16 cuBLAS 性能的 5 倍。在 H20 上的大 batch 尺寸下，多种位宽仍保持接近 BF16 cuBLAS 的计算效率。这让模型可以在精度、存储和吞吐之间做出更细粒度的选择。

W8A8 / W4A8 / W4A4

在 Hopper 上，带块级缩放的 FP8 W8A8 成为高性能量化推理的主流方案。Humming 通过同时优化数据搬运、缩放和专家调度，进一步改进了这条成熟的计算路径。下图左列展示了 H20 和 H200 上的 W8A8 结果。在中小 batch 尺寸下，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 提升到 633.94 tokens/s，而平均每输出 token 耗时（TPOT）下降约31%。

用法

安装

使用 pip 安装 Humming：

pipinstallhumming-kernels

Humming 使用 CUDA 编译器按需生成内核。如果系统上没有可用的 CUDA 编译器，请选择以下安装选项之一，以包含你所在环境的 CUDA 主版本对应的编译依赖：

pipinstall"humming-kernels[cu12]"# CUDA 12
pipinstall"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 及以上）。

The--linear-backend和--moe-backend选项分别选择 Humming 作为稠密 GEMM 和 MoE GEMM 的计算后端。例如，在张量并行大小为 8 的情况下为两者启用 Humming：

vllmservedeepseek-ai/DeepSeek-V4.1-Flash-1017\
--tensor-parallel-size8\
--linear-backendhumming--moe-backendhumming

该选项仅作为兼容路径，用于 Humming 支持但 vLLM 不再支持或尚未原生支持的格式，例如 GPTQ INT3。例如，通过此兼容路径加载 GPTQ INT3 检查点：--quantization hummingvllm

servehfl/Qwen2.5-VL-3B-Instruct-GPTQ-Int3--quantizationhumming选择激活类型

默认情况下，Humming 使用检查点中指定的激活类型，并根据设备调整执行方式。如果 GPU 不支持该计算类型，Humming 会回退到受支持的更高精度类型。例如，不支持原生 MXFP4 的设备可以自动使用 FP8、BF16 或 FP16 计算路径，从而让同一个量化检查点跨 GPU 代际运行。

用户可以通过 vLLM 的量化配置接口覆盖默认设置，为部署选择激活类型。以下示例将 MoE 层的激活类型设置为 MXFP8：

vllm

serveopenai/gpt-oss-20b--moe-backendhumminghumming--quantization-config.moe.activationmxfp8

未来计划

更广泛、更深入的设备支持

Humming 将继续扩大硬件覆盖范围，同时适配新设备的计算与内存能力，帮助更多平台释放量化推理的性能潜力。优先事项包括：

• 支持更多 Rubin 架构特性，包括 Locality Domains，以进一步优化任务映射和数据访问。
• 将支持扩展到更多 T-Head PPU 设备，包括 Zhenwu M890。
• 支持更多兼容 CUDA 编程的设备，将现有的量化格式、内核组件和调优能力带到更多硬件平台。

新增 Autotune

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

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

与推理框架的更深层集成

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

致谢

我们感谢蚂蚁集团的 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 及其他开源项目。我们感谢这些项目的作者和贡献者分享他们的想法和实现，帮助推动整个社区的高性能量化推理发展。