11 — Marlin, Machete, and the Quantized-GEMM Kernel Reality

A quantized model with a bad kernel is a slow fp16 model that also lost quality. The kernel is the deliverable. This is where Phase 2 (GPU architecture) and Phase 5 fuse into one skill.

The core trick, said once, correctly

Every fast quantized GEMM in <phone_number_or_numberic_id_or_random_id_152> follows the same three-line recipe:

1. Store weights compressed in HBM. INT4 or FP4 packed. 4× less bandwidth than fp16.

2. Dequantize in registers or shared memory, just before the tensor-core MMA. Never write dequantized weights back to HBM — that defeats the entire purpose.

3. Feed the tensor cores with fp16/bf16/int8 (whatever the MMA supports natively). Accumulate in fp32.

The HBM bandwidth win is real (weights are ~half the traffic in decode). The tensor-core precision win is real (fp16/bf16 accumulate cleanly). The only question is how you overlap dequant with MMA so the dequant doesn’t become the new bottleneck.

Read that recipe three times. If you can articulate it under study pressure, you’re 80% of the way to sounding like a kernel person.

Marlin — the paper and the paradigm

Marlin, arxiv <phone_number_or_numberic_id_or_random_id_153> (Frantar, Castro, Chen, Hoefler, Alistarh, Aug 2024). Repo: github.com/IST-DASLab/marlin. Read the README twice — it is a masterclass in Ampere optimization prose.

The problem Marlin solved: existing W4A16 GEMMs (pre-2024) were fast at batch=1 (bandwidth-bound anyway, easy) but slower than fp16 at batch=8+ because the dequantization got in the way. Speculative decoding, multi-user batching, medium-batch serving — all the settings where you’d want a quantized win — didn’t get one.

Marlin’s contribution: a W4A16 GEMM that stays memory-bound up to batch 16–32, delivering ~4× speedup vs fp16 in that regime, and 2.8× end-to-end vLLM improvement on quantized models.

The key optimizations (each of these deserves a line in your study cheatsheet):

1. Async global-to-shared memory copies (cp.async, Ampere feature). Overlaps weight-load latency with MMA compute. This is the Ampere trick.

2. Double-buffered pipeline stages. Two shared-memory buffers alternating: while stage A feeds tensor cores, stage B loads the next tile.

3. INT4-in-INT32 packed layout. 8 four-bit weights per int32 element. Load is coalesced, unpack is a bit-shift.

4. Register-resident dequantization. Weight fp16s live only in registers during MMA; never spilled to shared memory.

5. Warp-specialized epilogue. Some warps do dequant, some do MMA, producer/consumer style. Not as extreme as Hopper warp specialization, but the same idea.

6. Interleaved weight layout. Weights stored in an unusual order (“striped” layout) that makes the shared-memory writes bank-conflict-free after unpack.

The README shows Nsight Compute screenshots proving the kernel achieves ~95% of theoretical memory bandwidth at batch=1, holds ~80% at batch=16, and only starts to drop past batch=64. That is the achievement.

2:4 sparsity extension: sparse-marlin adds NVIDIA structured sparsity (mma.sp instruction) on top of INT4 quant. Effectively 8-bit compression (INT4 × 0.5 density). Gets used less than it should because the accuracy cost of enforcing 2:4 during PTQ is real, but the kernel is there.

Machete — the Hopper successor

Machete is Neural Magic’s next-generation kernel for Hopper (H100/H200). Not a standalone arxiv paper — documented in vLLM PRs and the compressed-tensors library.

Key differences vs Marlin:

• Uses wgmma (warp-group MMA) instructions. Hopper’s async MMA that lets a warp group issue a large matmul and continue doing other work while it runs. Marlin doesn’t have this on Ampere; Machete lives on it.

• Tensor Memory Accelerator (TMA) for bulk async copies. TMA is Hopper’s replacement for cp.async, moving whole tiles in one instruction with better latency hiding.

• Thread-block clusters — groups of thread blocks that can share distributed shared memory across SMs. Machete uses this for larger effective tiles.

• Producer-consumer warp specialization — explicit; goes further than Marlin’s on Ampere.

• Faster than Marlin on H100 (as expected) and, critically, stays memory-bound at higher batch (up to 32–64) because the wgmma pipeline overlaps better.

In vLLM: Machete is the default W4A16 kernel on Hopper starting late 2024. You get it automatically when you serve a GPTQ or AWQ model on H100. On Ampere you still get Marlin. This is transparent unless you go look.

The kernel path table (memorize)

Format	Ampere (A100, 3090, 4090)	Hopper (H100, H200)	Blackwell (B200, 5090)
W4A16 (GPTQ/AWQ)	Marlin	Machete	Machete (Blackwell-tuned soon)
W8A8 INT8 (SmoothQuant)	cuBLAS int8 GEMM	cuBLAS int8 GEMM	cuBLAS int8 GEMM
W8A8 FP8	~software emul (skip)	cuBLAS fp8 GEMM	Native, faster
W4A8 (QServe)	Marlin-style + int8 act	Machete-adapted	Blackwell native path
NVFP4	Emulation (skip)	Emulation (skip)	Native tcgen05.mma
MXFP4	Emulation (skip)	Emulation (skip)	Native (with MR-GPTQ)
INT4 KV attention	FlashInfer paged path	FlashInfer paged path	FlashInfer paged path

The study-friendly summary: INT4 weight on Ampere = Marlin; INT4 weight on Hopper = Machete; FP4 anywhere = Blackwell-only in practice.

LiquidGEMM — the current W4A8 SOTA

LiquidGEMM, arxiv <phone_number_or_numberic_id_or_random_id_154> (Sep 2025). Repo: LiquidAI/LiquidGEMM (external). Key contributions:

• LiquidQuant: a 2-instruction dequantization kernel per 4 elements (down from Marlin’s ~4-5). This changes the dequant/MMA balance and lets W4A8 stay memory-bound at higher batch.

• Implicit fine-grained pipelining: overlap dequant with the next tile’s MMA at the register level, not just at the block level.

• Results: 2.9× over previous SOTA W4A8 kernels, 1.12–1.63× over TensorRT-LLM’s own W4A8 path.

Why this matters: W4A8 is the QServe (omniserve) target format, and it’s arguably the most balanced quant scheme (4-bit weights for capacity, 8-bit activations for compute-bound prefill). If the kernel is fast enough to justify the extra activation cost, W4A8 becomes competitive with FP8 for medium-batch serving. LiquidGEMM is the kernel that closes that gap.

Triton W4A16 SplitK

Hoque et al., arxiv <phone_number_or_numberic_id_or_random_id_155> (Jan 2024). Not a production kernel — a research Triton implementation. Contribution: for skinny-M shapes (small batch, common in decode), SplitK parallelization along the K dimension gives 65% average speedup on A100 and 124% on H100 over existing Triton W4A16 baselines.

Why you care: Marlin and Machete are C++/CUTLASS heroic engineering, and reimplementing them in Triton for learning is hard. This paper (and its accompanying repo) is your reference “aspirational Triton W4A16 kernel” for the naive-vs-Marlin gap analysis project. You can write a Triton W4A16 in a weekend; hitting Marlin’s numbers is a multi-month project. This paper shows what’s achievable in pure Triton and where the gaps are.

llama.cpp CPU quantized kernels

Already introduced in 10_gguf_k_quants.md. Same three-line recipe, adapted to CPU:

1. Store Q4_K_M super-blocks in RAM. 4.5 bits/weight, sequential DDR reads.

2. Dequantize into f32 SIMD registers (AVX-512 _mm512_cvtepi8_ps, or AMX tile ops on Sapphire Rapids+).

3. VNNI int8 dot products (or AMX bf16 matmul) with the quantized activations (Q8_K working type).

Canonical file: ggml/src/ggml-cpu/ggml-cpu-quants.c. Pick ggml_vec_dot_q4_K_q8_K and annotate every SIMD intrinsic. This is the CPU equivalent of reading Marlin.

AMX vs AVX-512 on Xeon Sapphire Rapids: AMX gives ~2–4× speedup for large-enough matmuls, but has warm-up cost. llama.cpp uses AMX above a threshold M, AVX-512 below it. This kind of runtime dispatch is standard — read the dispatch table in ggml-cpu.c.

Marlin’s Nsight Compute lessons (worth internalizing)

The Marlin README shows profiling data that maps directly to your Phase 2 skills:

• Achieved memory bandwidth: ~95% of HBM peak at batch=1. This is what “memory-bound done right” looks like.

• Tensor-core utilization: low at small batch (that’s correct — you’re memory-bound), climbs with batch, peaks around 70% at batch=32.

• Warp stalls: long-scoreboard stalls (waiting on HBM) dominate at batch=1; short-scoreboard (waiting on shared-mem/register hazards) starts appearing at batch=8.

• Occupancy: deliberately kept moderate. Marlin does not max occupancy — it maxes ILP per warp, which is the Volkov “Better Performance at Lower Occupancy” principle you’ll internalize in Phase 2.

When you run your own W4A16 Triton kernel in 14_projects.md project 2, these are the exact same numbers you’ll compare against.

Where kernels are NOT the bottleneck (equally important)

• Small batch, small model, single query. You’re bandwidth-bound on weights; any reasonable kernel is within 20% of optimal. Marlin vs a decent Triton kernel matters little here.

• Very long prefill. Compute-bound, so tensor-core utilization matters more than dequant efficiency. FP8/int8 activation quant is the win, not fancier W4A16.

• MoE with wide expert parallelism. All-to-all communication dominates; kernel choice fades.

Knowing when not to reach for Marlin is part of the skill.

Anti-patterns

• Don’t write a naive dequant-then-GEMM. Materializing dequantized weights in HBM is the anti-pattern that killed early W4A16 speedups. If you’re not fused, you’re wrong.

• Don’t compare Marlin to fp16 cuBLAS at batch=1. Both are bandwidth-bound on the same weights (fp16 loads 2× more bytes, so Marlin is 2–4× faster — that’s the honest ratio). Compare at batch=8/16/32 where the difference is engineering, not format.

• Don’t try to write a production kernel in Triton. Triton is great for learning, prototyping, fused custom ops. Production kernels for W4A16 belong in CUTLASS/CuTe (or hand-CUDA) because you need the last 20% on register allocation and instruction scheduling.

• Don’t ignore Ampere. “Everyone’s on Hopper” is false. Zoho customer racks with A100 or L40S will be doing quantized serving for years. Marlin is your friend.

The two-sentence study answer

Fast quantized GEMMs all share one recipe: weights compressed in HBM → dequantized in registers → tensor-core MMA accumulated in fp32, fused in one kernel — the trick is how you hide the dequant behind the MMA. Marlin (Ampere, cp.async + double-buffered pipeline, arxiv <phone_number_or_numberic_id_or_random_id_156>) stays memory-bound to batch ~32 with ~4× speedup; Machete is its Hopper successor using wgmma + TMA + thread-block clusters, and it’s the default in vLLM on H100.

Homework

1. Read the Marlin README (github.com/IST-DASLab/marlin) end-to-end. Note every Nsight Compute metric they cite. Predict what the same metrics would look like on your GPU before running.

2. Write a Triton W4A16 dequant+GEMM kernel following the Triton tutorial + the arxiv <phone_number_or_numberic_id_or_random_id_157> paper’s SplitK idea. Measure at batch=1, 8, 32 vs Marlin (via vLLM). Write the gap analysis. This is project 2 in 14_projects.md.

3. Pick one CUDA file in the Marlin repo (marlin_cuda_kernel.cu) and annotate the double-buffered pipeline section line by line. This is the closest thing to a graduate thesis you can do in a weekend.

4. Read the Machete PR history in vLLM (search vllm-project/vllm PRs for “machete”). Note which optimizations Marlin lacks and which are Hopper-only.