Do Quantization and Sparsity make LLM Inference Faster ?
Before reading: This article assumes that you know the basic shape of a Transformer. If terms such as decode, Q/K/V projection, MLP projection, and LM head are unfamiliar, start with The Illustrated Transformer for the matrix operations, then read NVIDIA’s introduction to LLM inference for the prefill and decode phases. The Hugging Face encoder-decoder guide also shows how the LM head turns the final model vector into scores for all possible next tokens.
We can compress the weights of Large Language Models by storing each value with fewer bits, which is called quantization, or by removing values in a regular pattern, which is called structured sparsity. Both approaches reduce the amount of data stored in GPU memory, allowing us to run larger models in lower VRAM GPUs.
However, this does not always translate into more speed. The reason is that the GPU runs a kernel (a small program that tells the GPU how to perform an operation, such as matrix multiplication) to perform various operations (like Attention, Matrix Multiplication) in a Large Language Model, and different kernels can have very different speeds. The kernel decides how values are fetched from memory, reused from cache, divided across GPU cores, converted back to a usable form, and combined into an answer. Different quantization and sparsity methods therefore normally have different kernels written for them.
A representation with fewer bits gives the GPU fewer bytes to move, but a slow kernel can spend more time decoding those bytes or organizing the work than it saves on memory traffic. The question was: when do compressed representations turn fewer bytes into less execution time?
Why this matters
Quantization is an excellent way to reduce the size of a model’s weight representation. Smaller weights use less VRAM, so a larger model, or a larger KV cache can fit on the same GPU. Moving fewer bytes from memory can also make an operation faster.
But lower memory use and higher speed are different benefits. Even when the weights load faster, the kernel may need extra instructions to unpack them, apply scales, read sparsity metadata, or create a temporary full-precision matrix. That extra work can erase the time saved by loading fewer bytes. To predict performance, we therefore need to consider the representation and the kernel that executes it.
Main finding
Compression did not guarantee faster inference. NF4 was 1.75× faster than dense FP16 at batch 1, while INT8 was 1.55× faster at batch 32 and 1.77× faster at batch 128. The standard library’s 2:4 sparse path was slower across the full decode step at every tested batch. A custom CUDA 2:4 GEMV kernel changed that batch-1 result from a 2.88× slowdown into a small 1.06× speedup.
The custom kernel is important because it holds the sparse representation roughly constant while changing how the GPU executes it. GEMV simply means matrix-vector multiplication, the form used by these projections when the batch contains one token. The result shows that compression creates an opportunity; the kernel determines how much of that opportunity is realized.
What we did
I extracted the eight matrix-multiplication shapes used in one real Qwen2.5-1.5B decode step. Together, they represent 197 matrix multiplications per generated token. I tested them at batches 1, 8, 32, and 128 using:
- dense FP16 through cuBLAS as the baseline
- INT8 through PyTorch’s integer matrix-multiplication path
- NF4 through bitsandbytes
- 2:4 sparse FP16 through cuSPARSELt
- a custom CUDA 2:4 sparse GEMV kernel for batch 1
The tests ran on an SM86 RTX 3060 12 GB. In this setup, the card delivered a 341.4 GB/s streaming-read ceiling.
The chart below replaces the full heatmap with four representative batch-1 matrix multiplications: a small K projection, a medium Q projection, a large down projection, and the much larger language-model head.
Values above the dashed 1× line are faster than dense FP16. The same compressed format can lose on one matrix and win on another because the useful kernel path depends on the shape.
Dense FP16: the baseline
FP16 stores every weight using 16 bits. It does not compress or skip any weights, so it gives us a stable reference for judging the other approaches. The GPU used the highly optimized cuBLAS matrix-multiplication kernels.
Across the 197 matrix multiplications, the FP16 path took 10.86 ms at batch 1, 11.33 ms at batch 8, 12.42 ms at batch 32, and 19.86 ms at batch 128. Every speedup below compares against the FP16 time at the same batch size.
| batch | FP16 time | use in this article |
|---|---|---|
| 1 | 10.86 ms | latency baseline |
| 8 | 11.33 ms | small-batch baseline |
| 32 | 12.42 ms | larger-batch baseline |
| 128 | 19.86 ms | highest tested batch |
INT8: faster once the batch was large enough
INT8 stores each weight using 8 bits rather than 16. The tested path multiplies the integer values and then runs a scaling step to convert the result back into a useful form.
At batch 32, INT8 reduced the weighted decode-step time from 12.42 ms to 8.01 ms, a 1.55× speedup. At batch 128, it reduced the time from 19.86 ms to 11.24 ms, a 1.77× speedup. These are the two larger batches in this experiment; batch 64 was not part of the recorded sweep.
The specific PyTorch INT8 kernel required more than 16 input rows, so it did not produce measurements for batches 1 and 8. Once the input was supported and large enough, moving fewer weight bytes translated cleanly into speed.
| batch | FP16 time | INT8 time | result |
|---|---|---|---|
| 32 | 12.42 ms | 8.01 ms | 1.55× faster |
| 128 | 19.86 ms | 11.24 ms | 1.77× faster |
NF4: fast at batch 1, slower at larger batches
NF4 stores weights in 4 bits, together with small scaling values needed to reconstruct them. At batch 1, the input is effectively one vector, so bitsandbytes selected a fused GEMV (matrix-vector multiplication) kernel. This kernel read the compressed weights and performed the multiplication directly, without first rebuilding the entire FP16 matrix. The full set of projections therefore fell from 10.86 ms to 6.20 ms, a 1.75× speedup.
At larger batches, the operation became GEMM (matrix-matrix multiplication), and bitsandbytes 0.49.2 switched to a different execution path containing two main kernels:
- a dequantization kernel rebuilt the complete weights in FP16;
- a dense GEMM kernel multiplied the batch by those rebuilt weights.
This was the problem. The weights occupied less VRAM while stored, but the GPU still had to reconstruct and process the much larger FP16 matrix during every operation. The extra work erased the advantage of loading the 4-bit representation. NF4 used the same 4-bit representation at every batch size, yet its performance changed from faster to slower when the kernel path changed. This directly reinforces the experiment’s main point: compression alone does not determine speed; the kernel used to execute it does.
| batch | FP16 time | NF4 time | result |
|---|---|---|---|
| 1 | 10.86 ms | 6.20 ms | 1.75× faster |
| 8 | 11.33 ms | 23.71 ms | 2.09× slower |
| 32 | 12.42 ms | 24.78 ms | 2.00× slower |
| 128 | 19.86 ms | 32.33 ms | 1.63× slower |
Library 2:4 sparsity: fewer values, but too much overhead
In 2:4 structured sparsity, only two values are kept from every group of four. Metadata records where those two values belong. Our stored sparse weights and metadata occupied 56.25% as many bytes as dense FP16 weights.
Despite moving fewer weight bytes, the cuSPARSELt path took 31.28 ms at batch 1, compared with 10.86 ms for FP16. It remained slower across the complete decode step at batches 8, 32, and 128. The library kernel had a fixed cost of roughly 0.10–0.20 ms even for small projections, so those small operations could not do enough useful work to recover that cost.
| batch | FP16 time | library 2:4 time | result |
|---|---|---|---|
| 1 | 10.86 ms | 31.28 ms | 2.88× slower |
| 8 | 11.33 ms | 24.47 ms | 2.16× slower |
| 32 | 12.42 ms | 24.56 ms | 1.98× slower |
| 128 | 19.86 ms | 28.43 ms | 1.43× slower |
Matrix size still mattered. On the very large LM-head matrix, the same library path was 1.41× faster than FP16 at batch 1. The library could use the sparse hardware effectively there; it simply was not a good general kernel for the many smaller projections in this decode step.
Custom CUDA 2:4 GEMV: changing the kernel changed the result
To separate the sparse representation from the library implementation, I wrote a custom CUDA kernel for batch-1 2:4 matrix-vector multiplication. It stores two FP16 values from every group of four and packs their positions into small metadata codes. Each GPU block then produces one output value directly from that compressed data.
On the down projection, the custom kernel reduced the time from 0.101 ms to 0.070 ms, a 1.45× speedup over dense FP16. Across the complete decode step, it reduced the weighted total from 10.86 ms to 10.26 ms, a smaller 1.06× speedup.
| comparison | FP16 | custom CUDA 2:4 | result |
|---|---|---|---|
| down projection | 0.101 ms | 0.070 ms | 1.45× faster |
| complete decode step | 10.86 ms | 10.26 ms | 1.06× faster |
The custom implementation removed most of the library kernel’s fixed overhead, but it did not win everywhere. It matched FP16 on the smaller examples and was slower on the enormous LM head. A better kernel improved the same sparse format, but the best execution strategy still depended on the matrix shape.
Optional: what does the custom kernel actually do?
The GPU kernel itself is only 44 lines in the source file. The complete CUDA file is 83 lines because it also contains a small addition helper and the PyTorch code that launches the kernel. Below is the complete kernel with plain-language comments added. Reading it is optional; the benchmark result above does not require CUDA knowledge.
__global__ void sparse_2to4_gemv_kernel(
const half* __restrict__ input,
const half* __restrict__ values,
const unsigned char* __restrict__ metadata,
half* __restrict__ output,
int k) {
// One thread block calculates one output number (one matrix row).
const int row = blockIdx.x;
const int groups = k / 4;
const int values_per_row = k / 2;
const int metadata_per_row = k / 8;
// Jump to this row's saved weights and position codes.
const half* row_values = values +
static_cast<long long>(row) * values_per_row;
const unsigned char* row_metadata = metadata +
static_cast<long long>(row) * metadata_per_row;
float sum = 0.0f;
// The threads share all the groups of four weights.
for (int group = threadIdx.x; group < groups; group += blockDim.x) {
// Unpack the two positions that survived pruning.
const unsigned char packed = row_metadata[group >> 1];
const unsigned char code =
(group & 1) ? (packed >> 4) : (packed & 0x0f);
const int index0 = code & 0x03;
const int index1 = (code >> 2) & 0x03;
// Load only the two weights we kept, rather than all four weights.
const half2 pair = reinterpret_cast<const half2*>(row_values)[group];
const float2 weight = __half22float2(pair);
// Multiply them by the matching inputs and add the two answers.
const int input_base = group * 4;
sum += weight.x * __half2float(input[input_base + index0]);
sum += weight.y * __half2float(input[input_base + index1]);
}
// First combine each group of 32 threads (called a warp).
sum = warp_sum(sum);
__shared__ float warp_sums[8];
const int lane = threadIdx.x & 31;
const int warp = threadIdx.x >> 5;
if (lane == 0) warp_sums[warp] = sum;
__syncthreads();
// Then the first warp combines those smaller answers into the final one.
if (warp == 0) {
float block_sum =
lane < (blockDim.x / 32) ? warp_sums[lane] : 0.0f;
block_sum = warp_sum(block_sum);
// One thread writes the completed result back to GPU memory.
if (lane == 0) output[row] = __float2half(block_sum);
}
}The key idea is that the kernel never recreates the two removed weights. It
loads the two surviving weights, checks their saved positions, multiplies them
by the corresponding input values, and combines the partial answers.
warp_sum is the six-line helper outside this function; it lets a group of 32
threads add their partial answers together. For a gentle introduction to
kernels, threads, blocks, and the <<<...>>> launch syntax, see
NVIDIA’s easy introduction to CUDA C and C++.
Conclusion
Quantization and sparsity can make LLM inference faster, but bit count alone cannot predict the result. On this GPU and software stack, NF4 was the best tested choice for batch-1 latency, while INT8 was the strongest choice at the larger tested batches of 32 and 128. The library 2:4 path helped only on a very large matrix; the custom sparse kernel showed that a kernel designed for the batch-1 shape could recover a modest full-step win.
The practical rule is simple: use compression to save VRAM, but benchmark the actual kernel at the matrix shapes, batch sizes, hardware, and library versions you plan to serve. Fewer bytes create the possibility of speed, not a guarantee. This experiment measured kernel execution on one model and GPU; it did not measure model-quality changes caused by quantization or pruning.
GitHub repository
The complete report, raw measurements, benchmark code, custom CUDA kernel, and reproduction instructions are available in the LLM Inference Optimization repository.