Humming: vLLM and Ant Group Open-Source a General-Purpose Kernel Library for Quantized Inference
Ant Group and the vLLM community have open-sourced Humming, a general-purpose, high-performance kernel library for quantized inference. Humming matters because it replaces the one-kernel-per-quantization-format approach with a unified quantization representation, a JIT compilation path, and shared CUDA components — so a single library covers dense and MoE workloads across NVIDIA GPUs from Turing onward and CUDA-compatible devices such as T-Head's Zhenwu 810E, while reporting significant speedups over vLLM Marlin, FlashInfer, and DeepGEMM on many representative workloads.
The code is available at github.com/vllm-project/humming.
Why quantized inference needs a general-purpose kernel library
Quantization lowers the precision of weights and activations to reduce memory usage and data movement, and can also improve compute efficiency. As models and inference workloads grow, memory capacity, memory bandwidth, and compute costs increasingly constrain deployment.
The pressure is most visible in MoE models: each token activates only a subset of experts, but the full set of weights still has to be stored. DeepSeek-V3, cited in the announcement, has 671 billion parameters with 37 billion activated per token.
Quantization requirements vary widely across models and GPU generations. Common configurations include WNA16, W8A8, W4A8, and W4A4, where W and A specify weight and activation bit widths and N denotes a variable weight bit width. Even at the same precision, schemes such as INT4 and NVFP4 differ in numeric format and scaling. Scaling granularity, scale representation, group size, and symmetric or asymmetric quantization further expand the configuration space.
These storage and compute savings only translate into lower latency and higher throughput with efficient kernels. Prefill, decode, and varying MoE expert loads place different demands on compute and memory, and kernels must adapt to those workloads while keeping conversion, dequantization, and scheduling overhead low.
Developing and integrating a dedicated kernel for every combination creates a growing maintenance burden: each addition can require its own weight preprocessing, kernel selection, framework adapters, and tests, and optimizations must then be carried across multiple implementations. Humming addresses this by separating quantization descriptions from execution choices and composing kernels from shared components.
Humming's architecture: represent, configure, compile
Humming's execution path has three parts: represent the quantized computation, select configurations for the hardware and workload, and compile a kernel from CUDA components.
A unified quantization representation
Humming converts models exported by different quantization tools into the Humming standard format: a common quantization description with consistently organized weights, scales, zero points, and other tensors. This representation preserves each scheme's quantization semantics while describing weight format, scaling granularity, scale data type, zero points, and secondary scales as separate parameters. Activations have their own numeric format and quantization settings, and these parameters can be combined subject to format and hardware constraints.
For example, groupwise INT4 weights can be paired with BF16 activations, and low-precision floating-point weights with different scaling schemes use the same parameter system. This gives inference frameworks a common interface for models from different sources.
Tuning and JIT compilation per workload
Preprocessing validates parameters, pads dimensions, and repacks tensors into layouts suited to the target kernel. Humming then combines the quantization description with the weight shape, GEMM type, target GPU, and compute options to configure execution. Three configurations describe the kernel:
- Layer Config describes tensor shapes, data types, quantization parameters, and the target architecture.
- Compute Config specifies the GEMM type and computation options, such as batch invariance and input scale layout.
- Tuning Config selects tile shapes, pipeline stage counts, warp specialization, memory access, and MMA instructions. It also controls concurrency and scheduling; TMA can be configured per tensor.
By default, Humming selects tuning configurations using device-specific heuristics, and it can prepare different configurations for ranges of input sizes and dispatch to the appropriate kernel at runtime. The announcement notes that at small batch sizes a multistage pipeline using asynchronous copies (cp.async) and MMA can outperform a pipeline built around TMA and UMMA, while at larger batch sizes TMA and UMMA can offer better throughput. Humming accommodates both execution paths within the same configuration and component framework.
The three configurations specialize kernels compiled with NVRTC or NVCC, and compiled kernels are cached for reuse. Inference combines the selected kernel with the preprocessed weights and runtime activations.
Shared CUDA components across formats and architectures
At the CUDA level, Humming organizes task scheduling, data movement, dequantization and scaling, Tensor Core computation, and output processing into reusable components. It provides standard, warp-specialized, and dedicated UMMA pipelines that share compatible components while using the dataflow and synchronization each execution strategy requires.
Each group of components handles a different source of variation:
- Numeric processing handles quantization formats. Decoding and arithmetic components account for bit width, integer or floating-point representation, scales, and zero points. Dequantization and scaling are integrated into the kernel's computation.
- Memory and compute backends handle GPU architectures. Humming selects data movement and Tensor Core implementations for the target architecture, using asynchronous copies, TMA, and multicast where supported.
- Task mapping and scheduling handle dense and MoE workloads. Each GEMM type uses the appropriate scheduling logic, and scaling and bias addition can be fused into the output stage.
JIT compilation combines these components according to the three configurations and eliminates unused implementation branches. A new format or architecture can therefore be supported by extending the relevant components, and improvements to a shared component benefit compatible configurations — reducing the duplicated development and maintenance of separate specialized kernels.
JIT compilation is fast enough for iterative tuning
Humming's CUDA implementation is lightweight and avoids heavyweight kernel framework dependencies, which keeps JIT compilation fast. In a single W4A16 kernel compilation test on an H200 server with CUDA 13.4, and with compilation caches disabled, compilation time dropped from 23.35 seconds for CUTLASS with NVCC to 1.82 seconds for Humming with NVCC or 1.88 seconds with NVRTC.
Humming also compiles kernels in parallel during model warmup. In the reported tests, compiling and loading the kernels covering all supported M sizes for a single layer's GEMM takes only about 2–4 seconds, making it practical to prepare multiple specialized kernels before inference.
Device-specific performance tuning
Humming optimizes kernels using both the GEMM configuration and device parameters. Quantization formats, matrix shapes, GPU instruction support, SM count, shared memory, and register resources all inform execution choices.
Efficient W4A8-MXFP4 inference on existing GPUs. Architectures before Blackwell lack native MXFP4 Tensor Core support, and conventional W4A8-MXFP4 kernels can be limited by dequantization and scaling overhead at larger M. Humming addresses this with a fast MXFP4-to-INT8/FP8 dequantization path that converts MXFP4 weights into a supported compute format within the kernel. In this setting, the path substantially outperforms conventional W4A8 implementations at small M and approaches W8A8 performance at large M, so GPUs such as H200 can run high-performance W4A8-MXFP4 kernels while retaining the storage benefits of 4-bit weights.
Dequantization paths matched to each device. The same quantization format can benefit from different implementations across GPUs. Humming selects decoding, numeric conversion, and scaling paths for the device, pairing them with suitable memory access and compute pipelines to reduce instruction and data movement overhead. Shared components make these optimizations reusable across compatible dense and MoE kernels.
Extensive tuning controls with device-specific heuristics. Humming exposes tile shapes, pipeline stage counts, warp specialization, TMA, MMA instructions, CTA concurrency, and Stream-K scheduling. Device-specific heuristics combine GEMM shapes with resource constraints to balance parallelism, data reuse, and pipeline overhead, selecting different execution strategies as input sizes change. Developers can use the default configurations directly or refine them for particular models and deployment scenarios.
Results: device support
Humming's unified architecture enables broad GPU support. For NVIDIA GPUs, Humming fully supports every architecture from Turing onward, bringing a common quantized computing framework to multiple hardware generations. The unified JIT architecture also makes it straightforward to extend Humming to other GPUs compatible with CUDA programming, and Humming already fully supports T-Head's Zhenwu 810E.
Device support is not uniform at the instruction level. UMMA instructions on the B300 (SM103) no longer support INT8 computation, but earlier MMA instructions can still execute INT8 GEMM, with substantially lower performance.
Results: quantization support
Humming's support for quantization types depends primarily on compatibility between activation and weight types. Its unified quantization representation and kernel architecture allow a wide range of integer and floating-point weight formats to work with activations at different precisions.
Humming also supports a special floating-point format with zero exponent bits, whose representable values match those of symmetric integers at the same bit width. For example, INT8 excluding -128 and float8e0m7 both represent integers from -127 to 127. Weights in this format have the same activation type compatibility as INT weights of the same bit width, and Humming automatically converts eligible symmetric integer weights to this format for faster dequantization.
The community has also uncovered two hidden capabilities in Blackwell Tensor Cores: native support for float4e0m3 and float8e3m4, two special floating-point formats that are not exposed in the PTX type system. Humming supports these features across all Blackwell architectures, including datacenter SM100/SM103 and SM120/SM121 devices for workstations and personal computing. Rubin does not support these two hidden Tensor Core formats, so support for these native compute paths does not extend to Rubin.
Results: kernel microbenchmarks
The kernel microbenchmarks were run by the vLLM community and measure GEMM performance with M covering every power of two from 1 to 8192. Kernels are fully warmed up and timed independently with a cold L2 cache, with a brief pause between samples. MoE speedups use the sum of Gateup and Down GEMM latencies.
WnA16
WnA16 combines low-bit weights with FP16/BF16 activations, reducing weight storage and memory traffic while using widely available 16-bit compute capabilities. The W4A16 family is a major focus of Humming's optimization work.
For dense W4A16, Humming matches or outperforms vLLM Marlin in most tested configurations, with substantial gains on several Hopper and Blackwell workloads. Using cuBLAS as a common baseline, Humming delivers roughly 3–4× speedups on several devices at small batch sizes. As batch size grows, it reaches approximately twice Marlin's performance on large-batch H200 workloads and more than three times its performance in some B300 cases.
These gains extend to MoE. Humming provides efficient Indexed and Grouped paths for workloads where tokens are distributed across experts and the amount of work per expert changes with the load. MXFP4 W4A16 on H20 and H200 reaches roughly twice Marlin's performance in several medium-to-large batch cases. On B300, the MXINT4 and NVFP4 paths achieve even larger gains at large batch sizes, with the Grouped path reaching approximately 5×. Results on SM120 and SM121 also show Humming outperforming Marlin and the corresponding FlashInfer implementations across a range of workloads.
Humming's optimizations extend beyond W4A16 to other WnA16 combinations. Its shared dequantization and compute paths deliver strong performance for integer and floating-point weights at different bit widths, including INT1, INT2, INT3, and FP4 through FP8. The lowest-bit configurations on SM120 exceed 5× the performance of BF16 cuBLAS, and at large batch sizes on H20, multiple bit widths maintain compute efficiency close to BF16 cuBLAS.
W8A8 / W4A8 / W4A4
With Hopper, FP8 W8A8 with blockwise scaling became a mainstream approach to high-performance quantized inference. Humming improves this established path by optimizing data movement, scaling, and expert scheduling together. On H20 and H200 at small and medium batch sizes, Humming's W8A8 reaches approximately 2.6× and 2.5× the performance of DeepGEMM Contiguous, respectively, with competitive performance from both Indexed and Grouped layouts.
As more models adopt native MXFP4 weight quantization, converting the smaller weight footprint into practical inference gains becomes increasingly valuable. Humming's W4A8 path rapidly dequantizes MXFP4 weights as they are loaded and uses FP8 Tensor Cores for computation, enabling efficient W4A8-MXFP4 execution on Hopper. Keeping DeepGEMM W8A8 as the baseline, Humming's W4A8 reaches approximately 4.2× baseline performance on H20 and 4.5× on H200 at small and medium batch sizes, also substantially outperforming Humming's own W8A8 results in the same figure. FlashInfer CUTLASS W4A8 is also an implementation based on Humming's approach, while Humming's current implementation achieves better performance in the benchmarks shown.
Blackwell adds native Tensor Core support for microscaling formats such as MXFP8 and MXFP4. From datacenter SM100/SM103 devices to SM120/SM121 devices for workstations and personal computing, Humming uses each architecture's low-precision compute capabilities to support W8A8, W4A8, and W4A4 within one kernel and tuning framework. On B300 (SM103), using DeepGEMM Contiguous at the corresponding precision as the baseline, W8A8, W4A8, and W4A4 reach approximately 1.25×, 1.6×, and 1.4× baseline performance at small and medium batch sizes, and outperform the corresponding FlashInfer implementations across many workloads. Both Indexed and Grouped paths benefit from native low-precision computation.
These gains also extend to SM120 and SM121. Humming's MXFP4 × MXFP8 W4A8 and NVFP4 × NVFP4 W4A4 outperform FlashInfer CUTLASS at most tested points. On SM121 at small batch sizes, performance gains reach approximately 19% for W4A8 and 9% for W4A4; on SM120, Humming also leads across most tested sizes.
End-to-end inference performance
Humming's W4A8 gains extend to model serving. In a DeepSeek-V4.1-Flash test on eight H20 GPUs with TP=8, keeping the model's MXFP4 weights unchanged and switching MoE execution from Marlin W4A16 (BF16 activations) to Humming W4A8 (tokenwise FP8 activations) improved end-to-end throughput by approximately 14%–43% and reduced mean time to first token (TTFT) by approximately 25%–32%. The throughput advantage grows with concurrency: at concurrency 64, output throughput rises from 442.37 to 633.94 tokens/s, while mean time per output token (TPOT) falls by approximately 31%.
| Concurrency | MoE backend | Input throughput (tokens/s) ↑ | Output throughput (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%) |
Installation
Humming installs from PyPI:
pip install humming-kernels
Humming uses a CUDA compiler to generate kernels on demand. If no CUDA compiler is available on the system, choose one of the following installation options to include the compilation dependencies for your environment's CUDA major version:
pip install "humming-kernels[cu12]" # CUDA 12
pip install "humming-kernels[cu13]" # CUDA 13
Humming is extremely lightweight, with simple dependencies and a wheel package of only about 500 KB. On-demand compilation avoids bundling a large collection of precompiled kernels. In principle, any environment capable of running PyTorch 2.7 can also run Humming.
Using Humming in vLLM
Humming is deeply integrated into vLLM. Starting with vLLM 0.29.0, CUDA installations include Humming as a dependency, so no separate Humming installation is needed. Humming supports NVIDIA GPUs across all architectures from Turing onward (SM75 and above).
The --linear-backend and --moe-backend options select Humming as the compute backend for dense GEMM and MoE GEMM, respectively. For example, enable Humming for both with tensor parallel size 8:
vllm serve deepseek-ai/DeepSeek-V4.1-Flash-1017 \
--tensor-parallel-size 8 \
--linear-backend humming --moe-backend humming
The --quantization humming option is intended only as a compatibility path for formats that Humming supports but vLLM no longer supports or does not yet support natively, such as GPTQ INT3:
vllm serve hfl/Qwen2.5-VL-3B-Instruct-GPTQ-Int3 --quantization humming
Selecting activation types
By default, Humming uses the activation type specified in the checkpoint and adapts execution to the device. If the GPU does not support that compute type, Humming falls back to a supported higher-precision type. For example, devices without native MXFP4 support can automatically use FP8, BF16, or FP16 compute paths, allowing the same quantized checkpoint to run across GPU generations.
Users can override the default through vLLM's quantization configuration interface to select an activation type for their deployment. The following example sets the activation type for MoE layers to MXFP8:
vllm serve openai/gpt-oss-20b --moe-backend humming --quantization-config.moe.activation mxfp8
Future plans
Broader and deeper device support
Humming will continue expanding hardware coverage while adapting to the compute and memory capabilities of new devices. Stated priorities include:
- Supporting more Rubin architecture features, including Locality Domains, to further optimize task mapping and data access.
- Extending support to more T-Head PPU devices, including the Zhenwu M890.
- Supporting more devices compatible with CUDA programming, bringing existing quantization formats, kernel components, and tuning capabilities to additional hardware platforms.
Adding autotune
With targeted tuning, Humming can already achieve state-of-the-art performance in most scenarios. However, the wide range of model shapes, quantization configurations, and device resources makes it difficult for heuristics to select the optimal configuration in every case. Humming therefore plans to introduce autotune, benchmarking candidate configurations on the target device to find better execution parameters for each workload. Autotune will complement existing heuristics, reducing manual configuration search.
Deeper integration with inference frameworks
Humming will continue deepening its integration with inference frameworks such as vLLM, expanding support across quantized model loading, backend selection, parallel inference, runtime scheduling, and device adaptation. The stated aim is to make more quantization methods work with the full range of framework capabilities, including optimizations such as quantized All-to-All (A2A) communication for expert parallelism (EP).
Acknowledgments
The announcement thanks Jinzhen Lin, Zekai Gu, and Zhilin Huang from Ant Group, as well as Michael Goin, Roberto L. Castro, Chauncey Jiang, and other members of the vLLM community, for their contributions and support throughout Humming's development. It also thanks the vLLM community for providing development, testing, and CI resources used to validate and improve Humming across a wide range of hardware and workloads.
Humming's design and optimization draw inspiration from Marlin, vLLM Marlin, LMDeploy, DeepGEMM, FlashInfer, CUTLASS, and other open-source projects.