发布 Humming：面向量化推理的高性能灵活 Kernel 库（19 分钟阅读）

# 宣布推出 Humming：面向量化推理的快速、灵活 Kernel 库 来源：https://vllm.ai/blog/2026-10-21-humming 蚂蚁集团与 vLLM 社区开源了 Humming——一个通用、高性能的量化推理 kernel 库。Humming 提供以下特性： - **广泛的设备支持**：支持从 Turing 架构起的 NVIDIA GPU，以及兼容 CUDA 的设备，如平头哥的镇岳 810E。 - **丰富的量化支持**：统一架构支持密集（Dense）与 MoE 负载下多样化的权重量化与激活量化格式组合，覆盖广泛的整数与浮点类型、位宽及缩放方案。 - **强大的执行性能**：在众多代表性负载上显著超越 vLLM Marlin、FlashInfer 和 DeepGEMM 的现有实现。 仓库：github.com/vllm-project/humming (https://github.com/vllm-project/humming) ## 背景 随着模型和推理负载不断增大，内存容量、内存带宽和计算成本对部署的制约日益凸显。这在 MoE 模型中尤为明显：每个 token 只激活部分专家，但全部权重仍需存储。例如，DeepSeek-V3 拥有 6710 亿参数，每个 token 激活 370 亿参数。量化通过降低权重和激活的精度来减少内存占用和数据搬运，同时也能提升计算效率。 不同模型和 GPU 世代的量化需求各不相同。常见配置包括 WNA16、W8A8、W4A8 和 W4A4，其中 W 和 A 分别指定权重与激活的位宽，N 表示可变的权重位宽。即使精度相同，INT4 和 NVFP4 等方案在数值格式和缩放方式上也存在差异。缩放粒度、缩放数据表示、分组大小以及对称/非对称量化进一步扩大了配置空间。 这些存储和计算上的节省只有通过高效的 kernel 才能转化为更低的延迟和更高的吞吐。Prefill、Decode 以及 MoE 专家负载的变化对计算和内存提出了不同要求。kernel 必须适配这些负载，同时保持转换、反量化和调度开销尽可能低。 为每种组合开发和集成专用 kernel 会带来不断增长的维护负担。每新增一种组合都可能需要独立的权重预处理、kernel 选择、框架适配和测试。优化也必须在多个实现之间同步。 Humming 通过将量化描述与执行选择分离、并从共享组件中组合 kernel 来解决这一问题。 ## 通用架构 Humming 的执行路径分为三个部分：表示量化计算、为硬件和负载选择配置、以及从 CUDA 组件编译 kernel。 ### 统一的量化表示 不同量化工具导出的模型在数值格式、参数组织和权重布局上各不相同。Humming 将支持的导出格式统一转换为 Humming 标准格式：一种通用的量化描述，其中权重、缩放（scales）、零点（zero points）等张量的组织方式保持一致。 该表示在保留每种方案量化语义的同时，将权重格式、缩放粒度、缩放数据类型、零点和二次缩放（secondary scales）作为独立参数进行描述。激活拥有自己的数值格式和量化设置。这些参数在格式和硬件约束下可以自由组合。 例如，分组 INT4 权重可以与 BF16 激活配对。采用不同缩放方案的低精度浮点权重使用同一套参数体系。这为推理框架提供了来自不同来源模型的统一接口。 ### 针对每种负载的调优与 JIT 编译 标准格式为不同导出格式提供了统一表示。预处理阶段随后验证参数、填充维度，并将这些张量重新打包为适合目标 kernel 的布局。Humming 将量化描述与权重形状、GEMM 类型、目标 GPU 和计算选项结合，来配置执行方式。 三种配置共同描述 kernel： - **Layer Config**：描述张量形状、数据类型、量化参数和目标架构。 - **Compute Config**：指定 GEMM 类型和计算选项，如 batch 不变性（batch invariance）和输入缩放布局。 - **Tuning Config**：选择 tile 形状、流水线阶段数、warp 特化（warp specialization）、内存访问方式和 MMA 指令。它还控制并发和调度；TMA 可以按张量进行配置。 默认情况下，Humming 使用设备特定的启发式规则选择调优配置。它可以为不同输入尺寸范围准备不同的配置，并在运行时分发到合适的 kernel。 Humming 的统一架构使得针对不同负载调优执行参数和选择流水线变得简单。在小 batch 尺寸下，使用异步拷贝（`cp.async`）和 MMA 的多阶段流水线可能优于围绕 TMA 和 UMMA 构建的流水线。在较大 batch 尺寸下，TMA 和 UMMA 可以提供更好的吞吐。Humming 在同一配置和组件框架内容纳了这两种执行路径，从而能够为每种负载选择更有效的策略。 这三种配置共同特化通过 NVRTC 或 NVCC 编译的 kernel。编译后的 kernel 会被缓存以供复用。推理时将选定的 kernel 与预处理后的权重和运行时激活结合使用。 ### 跨格式与架构的共享 CUDA 组件 在 CUDA 层面，Humming 将任务调度、数据搬运、反量化与缩放、Tensor Core 计算和输出处理组织为可复用的组件。它提供标准流水线、warp 特化流水线和专用 UMMA 流水线。这些流水线共享兼容的组件，同时使用各自执行策略所需的数据流和同步方式。 每组组件处理一类不同的变化来源： - **数值处理组件处理量化格式。**解码和算术组件考虑位宽、整数或浮点表示、缩放和零点。反量化和缩放被集成到 kernel 的计算中。 - **内存与计算后端处理 GPU 架构。**Humming 为目标架构选择数据搬运和 Tensor Core 实现，在支持的情况下使用异步拷贝、TMA 和多播（multicast）。 - **任务映射与调度处理密集和 MoE 负载。**每种 GEMM 类型使用相应的调度逻辑。缩放和偏置加法也可以融合到输出阶段。 JIT 编译根据三种配置组合这些组件，并消除未使用的实现分支。因此，新增一种格式或架构只需扩展相关组件即可支持。共享组件的改进会惠及所有兼容配置，减少了为单独专用 kernel 重复开发和维护的成本。 该 CUDA 实现轻量且避免了对重量级 kernel 框架的依赖，因此 JIT 编译也十分快速。在 H200 服务器上使用 CUDA 13.4 进行单个 W4A16 kernel 编译测试，在禁用编译缓存的情况下，编译时间从 **CUTLASS 配合 NVCC 的 23.35 秒** 降至 **Humming 配合 NVCC 的 1.82 秒** 或 **配合 NVRTC 的 1.88 秒**。这降低了为新负载准备 kernel 的成本，使迭代调优更加可行。Humming 还能在模型预热期间并行编译 kernel。在我们的测试中，编译并加载覆盖单层 GEMM 所有支持 M 尺寸的 kernel 仅需 **约 2–4 秒**，这使得在推理前准备多个专用 kernel 成为可能。 ### 基于设备参数的性能调优 Humming 利用 GEMM 配置和设备参数对 kernel 进行深度优化。量化格式、矩阵形状、GPU 指令支持、SM 数量、共享内存和寄存器资源都会影响执行选择，将低位存储的优势转化为实际的推理性能。 **在现有 GPU 上高效运行 W4A8-MXFP4 推理。**Blackwell 之前的架构缺少原生 MXFP4 Tensor Core 支持，传统 W4A8-MXFP4 kernel 在较大 M 下可能受限于反量化和缩放开销。Humming 通过一条快速的 MXFP4 到 INT8/FP8 反量化路径来解决这一问题，该路径在 kernel 内将 MXFP4 权重转换为受支持的计算格式。在此场景下，该路径在小 M 下大幅超越传统 W4A8 实现，在大 M 下接近 W8A8 的性能。因此，H200 等 GPU 可以运行高性能 W4A8-MXFP4 kernel，同时保留 4 位权重的存储优势。 **与每种设备匹配的反量化路径。**相同的量化格式在不同 GPU 上可以从不同实现中受益。Humming 为设备选择解码、数值转换和缩放路径，并将其与合适的内存访问和计算流水线配对，以减少指令和数据搬运开销。共享组件使这些优化可以在兼容的密集和 MoE kernel 之间复用。 **丰富的调优控制与设备特定启发式规则。**Humming 提供 tile 形状、流水线阶段数、warp 特化、TMA、MMX 指令、CTA 并发和 Stream-K 调度等控制。设备特定的启发式规则将 GEMM 形状与资源约束结合，在并行度、数据复用和流水线开销之间取得平衡，随着输入尺寸变化选择不同的执行策略。开发者可以直接使用默认配置，也可以针对特定模型和部署场景进行调优。 ## 结果 ### 广泛的设备支持 Humming 的统一架构实现了广泛的 GPU 支持。对于 NVIDIA GPU，Humming 完全支持从 Turing 起的每一种架构，为多代硬件带来了统一的量化计算框架。 统一的 JIT 架构也使得将 Humming 扩展到其他兼容 CUDA 编程的 GPU 变得简单。Humming 已经完全支持平头哥的镇岳 810E，将共享 kernel 和调优能力带到更多硬件平台。 下图展示了 Humming 在每种设备上支持的高性能激活类型：即可以充分利用设备峰值计算能力的类型。图中省略的一些类型仍可通过其他指令路径运行。例如，B300（SM103）上的 UMMA 指令不再支持 INT8 计算，但较早的 MMA 指令仍可执行 INT8 GEMM，但性能会大幅降低。 ### 量化支持 Humming 对量化类型的支持主要取决于激活类型与权重类型之间的兼容性。其统一的量化表示和 kernel 架构使得广泛的整数和浮点权重格式可以与不同精度的激活配合使用，为模型精度和部署性能提供了灵活的选择。 下图展示了 Humming 支持的权重与激活类型组合。行代表权重类型，列代表激活类型，勾选表示支持的组合。结合上方的设备支持矩阵，可以概览每种硬件架构上可用的高性能量化计算路径。 Humming 还支持一种零指数位的特殊浮点格式，其可表示的值与同位宽对称整数相同。例如，排除 `-128` 的 INT8 和 `float8e0m7` 都可以表示从 `-127` 到 `127` 的整数。该格式的权重与同位宽的 INT 权重具有相同的激活类型兼容性。Humming 会自动将符合条件的对称整数权重转换为该格式以加速反量化。 社区还发现了 Blackwell Tensor Core 的两项[隐藏能力](https://zhuanlan.zhihu.com/p/2059376150565089368)：原生支持 `float4e0m3` 和 `float8e3m4`——这两种特殊浮点格式在 PTX 类型系统中并未暴露。Humming 在所有 Blackwell 架构上支持这些特性，包括数据中心的 SM100/SM103 以及工作站和个人计算的 SM120/SM121 设备。两种格式已集成到统一的量化和计算框架中，为低精度推理提供更多原生计算选项。Rubin 不支持这两个隐藏 Tensor Core 格式，因此这些原生计算路径的支持不会延伸到 Rubin。 ### 性能 以下 kernel 微基准测试由 vLLM 社区运行，测量 GEMM 性能，M 覆盖从 1 到 8192 的所有 2 的幂次。Kernel 完全预热后，在 L2 缓存冷启动状态下独立计时，样本之间有短暂间隔。MoE 加速比使用 Gateup 和 Down GEMM 延迟之和。每个图表均注明了其基线、量化格式、矩阵维度和软件版本。 #### WnA16 WnA16 是最通用、应用最广泛的量化推理方法之一。它将低位权重与 FP16/BF16 激活结合，在减少权重存储和内存流量的同时利用广泛可用的 16 位计算能力。这使其适用于多种 GPU 世代和模型部署场景。W4A16 系列尤其受到社区关注，也是 Humming 优化工作的重点。 对于密集 W4A16，Humming 在大多数测试配置下达到或超越 vLLM Marlin，在若干 Hopper 和 Blackwell 负载上取得显著增益。下图以 cuBLAS 作为统一起始基线。在小 batch 尺寸下，Humming 在多种设备上实现约 3–4 倍的加速。随着 batch 尺寸增大，它继续高效利用硬件，在大批量 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 架构的出现，采用块状缩放（blockwise scaling）的 FP8 W8A8 成为主流的高性能量化推理方案。Humming 通过同时优化数据搬运、缩放和专家调度，进一步提升了这一成熟的计算路径。下图左列展示了 H20 和 H200 上的 W8A8 结果。在中小批量下，Humming 分别达到 DeepGEMM Contiguous 约 2.6 倍和 2.5 倍的性能，同时基于 bf...