Why architecture matters here
Inference performance on GPUs is governed by the roofline: a kernel is either limited by how fast it can move bytes or by how fast it can do arithmetic, and the crossover depends on arithmetic intensity — flops per byte. For a weight-only-quantized LLM at batch size 1, each weight is used exactly once per token, so arithmetic intensity is low and the matmul is firmly memory-bound; the entire justification for INT4 is that you move a quarter of the bytes. But the moment you batch requests or process a prompt's many tokens at once, the same weights are reused across the batch, arithmetic intensity climbs, and the kernel becomes compute-bound — at which point tensor-core throughput, not memory, is the ceiling. A kernel that wins in one regime by specializing loses in the other.
This is why architecture matters here rather than being a mere implementation detail. A quantization scheme that produces beautiful accuracy but ships with a kernel that only speeds up batch-1 decode is nearly useless for a serving system, which lives at moderate batch sizes to amortize cost. Marlin's design target is explicitly the sustained region: it aims to hold near-4x speedup from batch 1 up through batch sizes of a few dozen, the range where real serving operates, so that the memory saving of INT4 becomes a throughput saving instead of evaporating at the batch size you actually run.
The second reason the architecture is load-bearing is that dequantization is not free. Every INT4 weight must be turned back into an FP16 value before it can enter a tensor-core multiply, and if that unpacking-and-scaling costs more than the memory it saves, the whole exercise is pointless. A slow dequant path can make INT4 slower than just reading FP16 weights directly. Marlin's layout and register-level bit manipulation exist precisely so that dequant is nearly invisible — folded into the load pipeline, done with a handful of instructions per group of weights, and overlapped with the matmul so the tensor cores never wait for it. Get that overlap wrong and the kernel is memory-bound and dequant-bound at the same time, the worst of both worlds.
There is a third, quieter reason the kernel architecture is decisive: the tensor cores are the most expensive silicon on the die, and leaving them idle is the cardinal sin of inference engineering. A kernel that stalls those units waiting on memory or dequantization is burning the chip's most valuable resource. The entire pipeline design — async loads, double buffering, register-level unpacking, warp tiling — exists to keep the tensor cores in a state of continuous useful work, so that the metric that actually matters, achieved tensor-core utilization, stays high across the whole batch range. When you profile a good INT4 kernel you are really asking one question: are the tensor cores busy, or are they waiting?
The architecture: every piece explained
Top row: the input layout and the load pipeline. Weights arrive as packed INT4 — eight 4-bit values crammed into each 32-bit word, in a permuted order chosen so that unpacking maps cleanly onto the register lanes each thread needs for the tensor-core instruction. Alongside them sit group scales: one FP16 scale per group of weights (commonly 128), the quantization metadata that turns integer codes back into real magnitudes. The async global load uses the GPU's cp.async mechanism to copy the next weight tile from global memory into shared memory without occupying registers or stalling the compute — the loads for tile N+1 are issued while tile N is being multiplied, and shared memory is double-buffered so producer and consumer never collide.
Middle row: the compute core. As a tile is consumed, each thread performs dequantization in registers: it unpacks the INT4 codes to FP16 using bitwise operations and a small lookup trick, then multiplies by the group scale. Those FP16 weights feed the tensor-core MMA (matrix-multiply-accumulate) instructions, which multiply the dequantized weight tile against the FP16 activation tile and accumulate in FP16/FP32 registers. Warp tiling organizes the work so that each warp computes a sub-tile of the output and there are enough independent MMAs in flight to hide instruction and memory latency — the essential trick for keeping tensor cores busy. When the reduction dimension is large, split-K divides it across multiple thread blocks that each compute a partial sum and then perform a global reduction to combine partials, which improves occupancy at small output sizes.
Bottom rows: the boundaries with the rest of the model. Activations remain FP16 — Marlin is weight-only quantization, so the activations flowing in carry full precision and the only lossy component is the stored weights, which is why accuracy holds up. The output is FP16, handed straight back to the next layer with no dequant boundary visible to the surrounding model. The ops strip names what you validate: a batch-size sweep to confirm the speedup holds across the serving range, a roofline check to confirm the kernel is near its bandwidth or compute ceiling rather than stalling, and a numerics comparison against FP16 to confirm the quantization plus kernel path stays within accuracy tolerance.
The permuted weight layout deserves emphasis because it is where much of Marlin's cleverness hides. Tensor-core MMA instructions require their operand fragments in a specific, non-obvious arrangement across the threads of a warp. Rather than shuffle data into that arrangement at runtime — which would cost instructions on the hot path — Marlin pre-permutes the INT4 weights offline, during quantization, so that when a thread unpacks its 32-bit word the eight resulting values land in exactly the register lanes the MMA expects. The layout is thus a contract computed once and exploited on every forward pass, trading a one-time repacking cost for a permanently cheaper inner loop — and it is exactly why a checkpoint packed for one kernel cannot be fed to another without repacking.
End-to-end flow
Trace one output tile through the kernel during prefill of a batched request. The serving engine has loaded the model with INT4 weights in Marlin's packed, permuted layout and per-128 group scales, computed offline by the quantization step (GPTQ-style or similar). A batch of eight sequences arrives; the attention and MLP layers issue linear projections, each of which is a Marlin matmul: FP16 activations of shape (tokens x in-features) times INT4 weights of shape (in-features x out-features).
The kernel launches a grid of thread blocks, each owning an output tile. A block begins by issuing cp.async loads for the first weight tile and the first activation tile into shared memory, then — without waiting — issues the loads for the second tile into the other half of the double buffer. The moment the first tile lands, the warps start work: each thread reads its slice of packed INT4 weights from shared memory into registers, unpacks eight codes with a few bit operations, multiplies by the relevant group scale to recover FP16 magnitudes, and issues tensor-core MMA instructions against the FP16 activation fragments. While those MMAs execute, the async loads for the next-next tile are already in flight — the pipeline is full, and the tensor cores are the bottleneck, which is exactly where you want to be.
Because the reduction dimension (in-features) is large, the kernel runs split-K: several blocks each accumulate a partial sum over a slice of the contraction dimension. When they finish, a lightweight global reduction adds the partials into the final FP16 output tile. That output flows immediately into the next operation — a bias add, an activation, the next projection — with no visible dequant step, because from the model's perspective this was just a matmul that happened to read cheaper weights.
Now the stress case that reveals the design. Suppose you sweep batch size from 1 to 64 and plot throughput. At batch 1 the kernel is memory-bound and delivers close to 4x over FP16 because it moves a quarter of the bytes; as batch grows the arithmetic intensity rises and the kernel smoothly transitions toward compute-bound, and — the whole point — the speedup does not collapse in the middle, because the async-load-plus-double-buffer pipeline keeps memory movement hidden and the warp tiling keeps the tensor cores fed throughout. A poorly designed INT4 kernel shows a deep valley at moderate batch sizes; Marlin's curve stays flat, and that flatness is the difference between a research result and a production serving kernel. The whole path is logged and profiled in the roofline check, so when a new GPU or a new model shape enters, you re-run the sweep and confirm the kernel still lives near its ceiling rather than silently falling off it.
It is worth naming what the profiler shows when this works, because that is your diagnostic vocabulary. In the compute-bound regime, the tensor-core pipes report near-full utilization and the memory subsystem sits comfortably below its ceiling — the sign that loads are fully hidden behind compute. In the memory-bound regime at batch 1, the memory-bandwidth counter approaches the card's peak while the tensor cores idle a fraction of the time — expected, and still roughly 4x faster than FP16 because a quarter of the bytes move. The failure signature you are hunting for is the opposite of both: memory below peak and tensor cores idle, which means the pipeline stalled — usually a dequant serialization or a shared-memory occupancy problem — and is precisely the state Marlin's design exists to prevent.