blogvLLM Blog·21 Oct 2026·Ant Group and the vLLM Community

Announcing Humming: Fast, Flexible Kernels for Quantized Inference

Ant Group and vLLM open-sourced Humming, a quantized inference kernel library that covers NVIDIA GPUs from Turing on and other CUDA-compatible chips.

Read original on vllm.ai

llm-inference quantization gpu-kernels

Depth 4/5Score 4.51

Extract

Ant Group and the vLLM community are open-sourcing Humming, a general-purpose, high-performance kernel library for quantized inference. Humming offers the following features: Broad device support: Supports NVIDIA GPU architectures from Turing onward, along with CUDA-compatible devices such as T-Head's Zhenwu 810E. Extensive quantization support: A unified architecture supports diverse combinations of weight and activation formats for dense and MoE workloads, covering a wide range of integer and floating-point types, bit widths, and scaling schemes. Strong execution performance: Significantly outperforms existing implementations from vLLM Marlin, FlashInfer, and DeepGEMM across many representative workloads. Repository: github.com/vllm-project/humming Background As models and inference workloads grow, memory capacity, memory bandwidth, and compute costs increasingly constrain deployment. This is especially apparent in MoE models: each token activates only a subset of experts, but the full set of weights still needs to 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 reduce memory usage and data movement, and can also improve compute efficiency. Quantization requirements vary 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. Kernels must adapt to these 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. Optimizations must then be carried across multiple implementations. Humming addresses this problem by separating quantization descriptions from execution choices and composing kernels from shared components. General-Purpose …

Explained

What happened

Ant Group and the vLLM community open-sourced Humming (pip install humming-kernels, ~500 KB wheel), one kernel library for quantized inference across NVIDIA GPUs from Turing onward, plus CUDA-compatible chips like T-Head's Zhenwu 810E.

Headline results vs existing kernels (vLLM Marlin, FlashInfer, DeepGEMM):

W4A16 (4-bit weights, 16-bit activations): 3 to 4x cuBLAS at small batch; about 2x Marlin on large-batch H200; up to 5x for MoE on B300.
W4A8 with MXFP4 weights on Hopper: about 4.2x (H20) and 4.5x (H200) DeepGEMM W8A8, even though Hopper has no native FP4 tensor cores.
End to end, DeepSeek-V4.1-Flash on 8x H20: switching MoE from Marlin W4A16 to Humming W4A8 gives +14% to +43% throughput and 25% to 32% lower TTFT. At concurrency 64, output goes from 442 to 634 tokens/s.
JIT compile of one W4A16 kernel: 1.8 s vs 23.4 s for CUTLASS with NVCC.

Why it works

Quantization shrinks weights (and sometimes activations) to fewer bits. That only helps if the kernel can unpack them fast. The problem is a combinatorial explosion: weight format x activation format x scaling scheme x group size x GPU generation x dense vs MoE. Writing one hand-tuned kernel per combination doesn't scale.

Humming splits the problem:

One standard quantization description. Weights, scales and zero points from any export tool are converted into a common format.
Three configs pick the kernel: Layer Config (shapes, types, GPU), Compute Config (GEMM type, options) and Tuning Config (tile sizes, pipeline stages, TMA vs cp.async, MMA instructions).
Shared CUDA components for decoding, data movement and tensor-core math, compiled just in time per workload and cached.

The W4A8 trick: weights stay 4-bit in memory (less bandwidth), get converted to INT8/FP8 inside the kernel as they load, and FP8 tensor cores do the math. Decode is usually memory-bound, so moving fewer weight bytes is the win.

The bigger pattern

This is the compiler approach to a combinatorial problem: describe the variation as data, build from composable parts, specialize at runtime. Same idea as query engines that JIT-compile plans instead of shipping one operator per type combination.

What it means in practice

When choosing a quantized model to self-host, the kernel matters as much as the bit width. Same weights, 14% to 43% different throughput.
For MoE serving, low-bit weights are the big lever because all experts must sit in memory even though few are active per token.
If you run anything on older or cheaper GPUs (H20, consumer SM120), libraries like this make 4-bit models practical without new hardware.