Ant Group and vLLM open-source Humming, a kernel library for quantized inference

Ant Group and the vLLM community are open-sourcing Humming, a general-purpose, high-performance kernel library for quantized inference. According to the project, it supports NVIDIA GPUs from Turing onward and many weight and activation formats for dense and MoE workloads.
- Ant Group and the vLLM community open-sourced Humming, a kernel library for quantized inference.
- It supports NVIDIA GPUs from Turing onward and CUDA-compatible devices such as T-Head's Zhenwu 810E.
- According to the project, it significantly outperforms vLLM Marlin, FlashInfer and DeepGEMM on many representative workloads.
- In a W4A16 test on an H200, compile time fell from 23.35 seconds (CUTLASS with NVCC) to 1.82 seconds (Humming with NVCC).
- A fast MXFP4-to-INT8/FP8 path enables W4A8-MXFP4 inference on pre-Blackwell GPUs such as H200.
What Humming offers
The vLLM blog lists three main features. The first is broad device support: NVIDIA GPU architectures from Turing onward, plus CUDA-compatible devices such as T-Head's Zhenwu 810E. The second is extensive quantization support, with a unified architecture covering combinations of weight and activation formats for dense and MoE workloads, across integer and floating-point types, bit widths and scaling schemes.
The third is performance. According to the announcement, Humming significantly outperforms existing implementations from vLLM Marlin, FlashInfer and DeepGEMM across many representative workloads. The code is available at github.com/vllm-project/humming.
Why quantized kernels matter
The authors point out that memory capacity, memory bandwidth and compute costs increasingly constrain deployment as models grow. This is especially visible in MoE models, where each token activates only some experts but all weights must still be stored. DeepSeek-V3, for example, has 671 billion parameters, with 37 billion activated per token.
Quantization lowers the precision of weights and activations to cut memory use and data movement. Configurations vary widely: common ones include WNA16, W8A8, W4A8 and W4A4, and even at the same precision, schemes such as INT4 and NVFP4 differ in numeric format and scaling. Group size, scale representation and symmetric or asymmetric quantization add more variation. The savings only become lower latency and higher throughput with efficient kernels, and writing a dedicated kernel for each combination creates a growing maintenance burden, from weight preprocessing to framework adapters and tests.
How the architecture works
Humming separates quantization descriptions from execution choices and builds kernels from shared components. It converts supported export formats into a common Humming standard format with consistently organized weights, scales, zero points and other tensors. Preprocessing then validates parameters, pads dimensions and repacks tensors for the target kernel.
Three configurations describe each kernel: a Layer Config for shapes, data types, quantization parameters and target architecture; a Compute Config for GEMM type and options such as batch invariance; and a Tuning Config for tile shapes, pipeline stages, warp specialization, memory access and MMA instructions. By default, device-specific heuristics choose the tuning configuration, and Humming can dispatch different kernels for different input size ranges at runtime.
The team says that at small batch sizes a multistage pipeline using cp.async and MMA can outperform one built around TMA and UMMA, while at larger batch sizes TMA and UMMA can offer better throughput. Humming supports both paths in one framework. Kernels are compiled with NVRTC or NVCC and cached.
Compile times and device-specific tuning
The CUDA layer is organized into reusable components for scheduling, data movement, dequantization and scaling, Tensor Core computation and output processing. A new format or architecture can be supported by extending the relevant components, and improvements to a shared component benefit compatible configurations.
The implementation is lightweight and avoids heavyweight kernel framework dependencies. In a single W4A16 kernel compilation test on an H200 server with CUDA 13.4, with compilation caches disabled, compile time fell from 23.35 seconds for CUTLASS with NVCC to 1.82 seconds for Humming with NVCC or 1.88 seconds with NVRTC. Compiling and loading kernels for all supported M sizes of one layer's GEMM takes about 2 to 4 seconds, according to the team's tests.
The post also describes W4A8-MXFP4 inference on GPUs older than Blackwell, which lack native MXFP4 Tensor Core support. Humming uses a fast MXFP4-to-INT8/FP8 dequantization path inside the kernel. According to the authors, this substantially outperforms conventional W4A8 implementations at small M and approaches W8A8 performance at large M, so GPUs such as H200 can run W4A8-MXFP4 while keeping the storage benefits of 4-bit weights.
What it means for inference teams
For teams serving quantized models, Humming offers a single interface across export formats and GPU generations, instead of separate kernels for each scheme. The faster compilation is meant to make iterative tuning and preparing specialized kernels before inference more practical. The performance claims come from the project itself.


