The Marlin Kernel and Why Quantized Weights Still Run Fast on GPU

There is no tensor core that multiplies 4-bit weights by 16-bit activations. Every quantized weight is expanded back to 16 bits before any arithmetic touches it, which sounds like it should make the whole exercise pointless. It does not, and the reason is the only piece of hardware reasoning you need to predict when quantization speeds inference up and when it does not.

No GPU multiplies 4-bit weights

A W4A16 kernel — 4-bit weights, 16-bit activations — reads packed integers from global memory, unpacks them, multiplies by the group’s scale and subtracts its zero point to recover FP16 values, and feeds those to the same tensor-core instruction an unquantized model would have used. The multiply-accumulate is identical. All the quantization did was change what was in memory.

So the kernel does strictly more arithmetic than the FP16 version: everything the FP16 version did, plus an unpack and an affine transform per weight. If arithmetic were the constraint, quantization would be a pure loss.

Why that is still a win: the roofline

During autoregressive decoding with a single sequence, the model processes one token per forward pass. Each weight in the model is loaded from memory and used in exactly one multiply-accumulate. The arithmetic intensity — operations per byte moved — is therefore about as low as it can be, and the kernel is bound entirely by how fast weights can be read.

Modern accelerators are badly lopsided in that ratio. Red Hat’s 2024 write-up of Marlin puts current GPU FLOP-to-byte ratios in the range of 100 to 200, meaning the hardware can do a hundred-odd floating-point operations in the time it takes to fetch one byte. If a weight is used fewer than a few dozen times per load, moving it is the entire cost, and cutting the bytes by four cuts the time by approximately four.

That is the whole justification for weight-only 4-bit at batch size one: a theoretical 4x ceiling that comes from bytes, not flops. The same reasoning also tells you where it ends. Every additional sequence in the batch reuses the same loaded weight for another multiply-accumulate, so arithmetic intensity rises linearly with batch size, and at some batch size the kernel crosses from bandwidth-bound to compute-bound and the 4x disappears.

Why earlier kernels collapsed at batch size 8

The uncomfortable fact about W4A16 kernels before Marlin is that they lost most of their advantage far earlier than the roofline says they should. Red Hat’s analysis describes prior kernels achieving near-ideal speedups — around 3.87x — at batch size 1 and deteriorating rapidly as the token count rose, well before the compute-bound crossover.

The cause is not the roofline; it is that a kernel tuned for batch size one is a kernel that has stopped trying to keep the GPU busy. At batch one there is almost no arithmetic to overlap the memory loads with, so the natural implementation is a simple load-then-compute sequence over few thread blocks. Grow the batch and the arithmetic arrives, but the kernel has no structure for overlapping it: it is still waiting on loads with idle tensor cores, and it is now doing the dequantization work too. The result sits between the two rooflines, touching neither.

What Marlin does

Marlin, developed at IST-DASLab and integrated into vLLM, is a FP16×INT4 kernel written on the premise that the interesting regime is not batch size 1 but the batch sizes real serving actually runs at. The techniques are all about keeping both the memory pipeline and the tensor cores saturated simultaneously.

• Asynchronous global loads. Weights are copied from global into shared memory with cuda::memcpy_async, which does not block the issuing threads. The warps continue computing on data they already have while the next tile is in flight.
• Double buffering. Shared memory holds more than one tile in a circular queue, so tile N+1 is being fetched while tile N is being consumed. Without this the async copy has nothing to hide behind.
• L2 cache discipline. Quantized weights are streamed in a way that avoids evicting the activations, which are small, reused across the whole matrix, and worth keeping resident. During decoding the output is small enough to sit in L2, which accelerates the reduction of partial results.
• Stream-K striped partitioning. The reduction dimension is split non-uniformly across streaming multiprocessors so that every SM has work, instead of the tail of a uniform split leaving most of the GPU idle on the last wave. The layout is chosen so each SM’s share is contiguous, minimising synchronisation between the partial reductions.
• Dequantization scheduled against the pipeline. The unpack-and-scale work is placed so it fills the gaps where threads would otherwise be waiting, rather than forming a serial stage of its own.

Marlin’s authors report speedups of roughly 3.9x relative to FP16 on an inference-optimised NVIDIA A10 for large matrices at batch sizes up to 16–32, where prior kernels had fallen to around 2x. Those are the authors’ published figures for their own hardware and matrix shapes; treat them as evidence about the technique rather than a prediction for your deployment.

The window where it applies

The batch-size window is the part of this that generalises, and it is worth stating plainly because it determines whether a quantized checkpoint helps you at all.

• Below the window — one user, one stream — weight-only 4-bit is close to its theoretical best, and the reason to quantize is both memory and speed.
• Inside the window — the batch sizes a small serving deployment actually reaches — a good kernel keeps most of the advantage and a bad one does not. This is the range Marlin was built for, and it is why kernel choice, not just format choice, shows up in throughput.
• Above the window — large-batch serving, and prefill of a long prompt, which is a large effective batch — the multiply is compute-bound, the dequantization is pure added work, and 4-bit weights save memory without saving time. At that point the reason to quantize is fitting the model and its KV cache on the card, which is a real reason on its own.