AutoAWQ+Marlin: Efficient INT4 Inference

• The paper demonstrates that AutoAWQ+Marlin couples AWQ-style INT4 quantization with MARLIN’s packed mixed-precision kernels to accelerate transformer linear layers.
• It employs a tile-based layout that leverages INT4 weights with FP16/BF16 activations to achieve up to 3.9× speedup on Ampere GPUs in batched autoregressive serving.
• Deployment via vLLM shows practical improvements in decode times and performance, while performance gains depend on batch size, kernel tuning, and architecture match.

Searching arXiv for MARLIN and AWQ-related papers to ground the response in current literature. AutoAWQ+Marlin denotes the deployment path in which AutoAWQ- or AWQ-quantized LLM weights are executed through MARLIN, a family of mixed-precision CUDA kernels for autoregressive inference that accelerates the linear layers of transformers by combining INT4 weight storage with FP16 or BF16 activations. In practice, this path is exposed through vLLM’s AWQ+MARLIN integration, which converts AWQ-formatted weights into MARLIN’s packed tile layout at load time and then uses MARLIN kernels for the Q/K/V projections, attention output projection, and MLP projections during decode (Frantar et al., 2024). The resulting system is most relevant to the regime where weight traffic dominates latency: batched autoregressive serving at moderate batch sizes on Ampere, and, more narrowly, batch-1 decode on edge- and physical-AI-style deployments where the realized benefit depends as much on kernel implementation and runtime overheads as on nominal bit-width reduction (Chen, 28 May 2026).

1. Definition and execution model

AutoAWQ provides activation-aware weight quantization, while Marlin is the kernel family that executes the quantized linear algebra in packed INT4 form rather than dequantising weights into BF16 for a separate GEMM. In the studied decode path, only the weight GEMMs move to Marlin’s INT4 packed kernels; attention itself remains PyTorch SDPA in the main matrix, so the path reduces weight traffic and swaps GEMM kernels without introducing a fused SDPA path (Chen, 28 May 2026).

MARLIN, expanded as Mixed-precision Auto-Regressive LINear kernels, targets the large matrix multiplications in attention and MLP blocks. In the core design, the activation matrix $A$ is FP16, while the weight matrix $B$ is stored in INT4 with per-column or per-group scale metadata, typically with group size $G = 128$. Dequantization is performed inside the kernel using register-level bit manipulations and synchronized shared-memory staging, followed by Tensor Core mma.sync on FP16 operands. The paper states that INT4 weights with FP16 activations are supported, and that mixed BF16 activations are conceptually supported by the same path (Frantar et al., 2024).

The AutoAWQ+Marlin interface is therefore not a distinct quantization algorithm in its own right. It is a coupling of AWQ-style weight-only INT4 quantization and MARLIN-style packed mixed-precision GEMM execution. This distinction is important because the quantizer determines the weight representation and accuracy characteristics, whereas the kernel determines whether the theoretical reduction in weight traffic is actually converted into lower step time. A plausible implication is that “AutoAWQ+Marlin” should be understood as a system-level composition rather than a single method.

2. Quantization format, packing, and kernel mechanics

MARLIN supports uniform INT4 weight quantization with either per-column scales or per-group scales with group size $G = 128$, and supports both symmetric and asymmetric formats. The standard dequantization formulas are given as

$$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$

for asymmetric per-group quantization, and

$\hat{w} = s_g q$

for symmetric per-group quantization. The data notes that typical AWQ settings are INT4, group size 128, and symmetric per-group scaling, and that zero-points are not commonly used in AWQ, simplifying metadata (Frantar et al., 2024).

The MARLIN kernel expects weights to be rearranged offline into $16 \times 64$ tiles aligned to 128-bit thread loads and packed so that each thread reads the exact INT4 elements needed for four $16 \times 16$ Tensor Core blocks. Scales are similarly packed; for grouped quantization, the scales used by the same thread role across $16 \times 16$ blocks are concatenated into 16-byte vectors, enabling single-instruction shared-memory loads. In the AWQ+MARLIN path in vLLM, the necessary packing and conversion into this tile layout is performed at load time, including arrangement of INT4 lanes and packing of group scales into 16-byte vectors (Frantar et al., 2024).

The kernel design relies heavily on Ampere instructions and memory primitives. The paper identifies cp.async global-to-shared asynchronous copy with evict_first, ldmatrix.sync for bank-conflict-free movement of activation tiles into per-thread register layouts, and mma.sync as the core instructions; Sparse-MARLIN additionally uses mma.sp for 2:4 Sparse Tensor Cores. Asynchronous global-to-shared copies are used to reduce L2 pollution from one-shot weight loads, while activation tiles are kept resident for reuse. Dequantization uses a single LOP3 bit-manipulation to promote INT4 nibble values into FP16 mantissae, after which scales are applied either once per output or per group in the main loop (Frantar et al., 2024).

At the scheduling level, each SM owns an output tile $C_{\text{sm}}$ of size $B$0, warps compute fixed-width 64-column sub-tiles, and pipeline depth $B$1 is used to overlap memory movement, dequantization, and Tensor Core execution. Striping across $B$2 and $B$3 tile columns is used to keep SMs busy and reduce global reductions on real LLM layer shapes. This suggests that AWQ+MARLIN performance depends not only on weight compression but on a highly specific layout-and-scheduling contract between quantized weights and the execution kernel (Frantar et al., 2024).

3. Performance envelope in batched autoregressive serving

The principal result of MARLIN is that INT4 quantization can preserve near-maximum memory-bound speedups even in batched autoregressive serving. For 4-bit weights with FP16 activations, MARLIN achieves approximately the maximum $B$4–$B$5 speedup versus FP16 up to batch sizes 16–32 on NVIDIA Ampere GPUs, as shown on per-layer benchmarks on A10 for a large $B$6 matmul with group size 128. The speedup then gradually decreases as batch size increases, reaching roughly $B$7 at $B$8 (Frantar et al., 2024).

End-to-end integration into vLLM also yields material decode gains. On Llama-2-7B/A10 with 64 input tokens and 64 generated tokens per sequence, MARLIN reduces generation time by approximately $B$9 relative to FP16. Across models and GPUs, the paper reports $G = 128$0–$G = 128$1 speedups at moderate batch sizes, with representative values of $G = 128$2 at $G = 128$3 on Llama-2-7B/A10, $G = 128$4 at $G = 128$5 on Llama-2-13B/A6000, and $G = 128$6 at $G = 128$7 on Yi-34B/A100. In a serving benchmark on RTX A6000, MARLIN reduces TPOT by approximately $G = 128$8, and Sparse-MARLIN reaches approximately $G = 128$9; TTFT also improves, and speedups remain stable across QPS (Frantar et al., 2024).

The underlying roofline account is explicit. For mixed FP16-activation, INT4-weight matmul, the time estimate is

$G = 128$0

For $G = 128$1 and symmetric quantization without zero-points, the average bytes per weight element are approximately

$G = 128$2

yielding a theoretical memory-bound speedup

$G = 128$3

versus FP16 weights. The paper reports a measured per-layer maximum of $G = 128$4, matching that estimate (Frantar et al., 2024).

The transition to compute-boundedness explains the regime dependence. On an A10 GPU, the compute-to-memory ratio indicates that loading 4-bit weights remains the bottleneck until batch size reaches approximately $G = 128$5, so batches up to 16–32 remain strongly memory-bound and stay close to the $G = 128$6 ceiling; beyond approximately 64, kernels trend compute-bound and speedups decline. For practical deployment, the data therefore recommends decode batch sizes at or below 32 to achieve approximately $G = 128$7 speedups on A10, A6000, and RTX 3090, while noting that prefill at very large batch sizes is already compute-bound and sees much smaller gains (Frantar et al., 2024).

4. Batch-1 decode and the physical-AI inference gap

A later study reframes AutoAWQ+Marlin in a different workload: batch-1 autoregressive decode for “physical AI” systems such as robots, autonomous vehicles, embodied agents, and edge copilots. Its core result is that this workload is memory-dominated but not necessarily bandwidth-limited, because launch-side overhead becomes visible on very fast GPUs. The paper formalizes an analytic lower bound on per-step latency as

$G = 128$8

with KV bytes per step

$G = 128$9

It then defines

$$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$0

which is interpreted as the achieved fraction of peak HBM bandwidth (Chen, 28 May 2026).

The reported pattern is that $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$1 falls as peak HBM bandwidth rises. For Qwen-2.5-7B at context length 2048, L4 reaches $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$2 while H100 reaches $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$3, so faster memory does not translate into proportional latency gains. A CUDA Graphs A/B experiment isolates the missing term: on H100 SXM5 at Qwen-2.5-7B, context 2048, batch 1, CUDA Graphs improves step time by $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$4 with a 95% bootstrap confidence interval $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$5, reducing mean eager decode from 14.83 ms to mean graphed decode of 11.78 ms across $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$6 fresh sessions. On L4, the same intervention yields only $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$7 (Chen, 28 May 2026).

Within this framework, AutoAWQ+Marlin is a weight-only INT4 path that swaps BF16 GEMMs for Marlin’s packed INT4 GEMMs while leaving attention on SDPA. On Ada L4 for Qwen-2.5-7B at context 2048 and batch 1, the paper reports the following controlled matrix:

Path	Step time	Achieved fraction of floor
bf16 SDPA baseline	62.32 ms/step	0.821
AutoAWQ + Marlin	45.24 ms/step	0.289
GPTQ + ExLlamaV2 EXL2 4.25 bpw	17.36 ms/step	0.754

These numbers show that AutoAWQ+Marlin improves over BF16 on L4, reducing step time from 62.32 ms to 45.24 ms, corresponding to a $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$8 speedup over BF16, but remains far above the INT4 analytic floor of 13.09 ms. The study attributes this gap to kernel implementation rather than bit-width alone: Marlin was tuned for Ampere SM80 and is not best-in-class on Ada SM89, whereas ExLlamaV2 uses Ada-tuned INT4 kernels that stream packed 4-bit weights more efficiently on that architecture (Chen, 28 May 2026).

A common misconception addressed directly by the data is that 4-bit quantization should mechanically yield approximately $$\hat{w} = s_g (q - z_g), \quad q \in \{0,\dots,15\}$$9 step-time reduction in batch-1 decode. The L4 matrix contradicts that simplification. The study’s interpretation is that compressing weights is necessary but insufficient: the runtime must actually stream compressed weights efficiently, and on fast GPUs it must also remove host-side launch overhead. This suggests that AutoAWQ+Marlin should be evaluated as a workload-specific and architecture-specific runtime path rather than as a guaranteed realization of nominal INT4 memory savings.

5. Accuracy characteristics and deployment workflow

The accuracy behavior of AutoAWQ+Marlin is attributed primarily to the quantization method rather than to the kernel. The MARLIN paper states that accuracy is determined by the quantization algorithm, naming GPTQ and AWQ, and that MARLIN’s kernels do not introduce numerical deviations compared to standard dequant-GEMM beyond FP16 rounding behavior. Dequantization is executed deterministically in registers and scales are applied as prescribed. The paper further notes that AWQ or AutoAWQ typically uses INT4 with group size 128 and symmetric per-group scales, delivering high accuracy with low activation outlier sensitivity, and that the vLLM AWQ+MARLIN adapter allows retention of AWQ’s accuracy while obtaining MARLIN’s speed (Frantar et al., 2024).

The deployment workflow described in the data is correspondingly concrete. First, quantize with AutoAWQ using INT4 weight quantization with group size $\hat{w} = s_g q$0 and symmetric per-group scales, then validate accuracy on the intended tasks. Second, enable the AWQ+MARLIN path in vLLM so that weights and scales are packed into MARLIN’s $\hat{w} = s_g q$1 tile layout and per-thread contiguous 128-bit loads. Third, use vLLM with the AWQ MARLIN backend for linear layers, which will automatically select MARLIN kernels for supported layers and devices. Fourth, choose batch sizes and sequence lengths that remain near memory-bound execution, especially decode batch sizes at or below 32. Finally, validate TPOT, TTFT, and numerical accuracy under realistic QPS (Frantar et al., 2024).

The practical scope of these recommendations is bounded by workload. For decode, the paper advises targeting batch sizes $\hat{w} = s_g q$2 for approximately $\hat{w} = s_g q$3 speedups on A10, A6000, and RTX 3090, while noting that batch sizes up to approximately 64 still yield significant gains. For prefill, large acceleration should not be expected at batch sizes greater than 512, because MARLIN approaches compute-bound FP16 performance there. In the physical-AI batch-1 setting, the later paper gives a more restrictive interpretation: AutoAWQ+Marlin helps on bandwidth-bound GPUs such as L4, but on H100-class silicon at batch 1 and short contexts, launch overhead becomes a primary lever and weight compression alone is not the binding constraint (Chen, 28 May 2026).

6. Limitations, trade-offs, and comparison with adjacent approaches

Several limitations recur across the sources. Speedups diminish as batch size grows beyond approximately 64 because arithmetic intensity rises and matmul becomes compute-bound. Relative gains are lower on high-bandwidth, high-compute GPUs such as A100, where pipeline startup and reduction overheads weigh more heavily. Irregular layer shapes or sharding across many GPUs can increase global reductions. The current design targets Ampere features, and while the approach is described as extensible to Ada and Hopper, optimal tuning may differ. Some attention variants with nonstandard projections or fused operators may require additional integration work (Frantar et al., 2024).

The batch-1 decode study narrows these limitations further. It reports no AutoAWQ+Marlin results beyond L4 at context 2048, does not combine AutoAWQ+Marlin with CUDA Graphs, and notes that at very long contexts the simple memory floor undercounts growing attention compute. It also emphasizes that operator boundaries and dispatcher heuristics matter materially: in its H100 single-decode experiments, default SDPA outperformed explicit FlashAttention paths for that shape, reinforcing that theoretical IO savings do not by themselves determine realized latency (Chen, 28 May 2026).

A broader comparison is provided by the XFP study, which is written partly for users currently relying on AutoAWQ+Marlin. It contrasts Marlin INT4’s fixed uniform representation with XFP’s learned codebooks, explicit sparse FP16 outlier residual, and cosine-thresholded auto-selection of effective bit-width. On workstation Blackwell for Qwen3.5-122B-A10B at TP=1, the paper reports XFP V2 at 110 tok/s and Marlin INT4 at 73.6 tok/s, with GSM8K strict-match in the same 94–95% band; it summarizes this as XFP being 49% faster than Marlin INT4 at TP=1. At the same time, the paper explicitly states that AutoAWQ+Marlin may still meet the needs of deployments optimized for large-batch throughput with mature INT4 kernels, especially where concurrency rather than single-stream latency is the bottleneck (Witt, 14 May 2026).

The resulting picture is not that AutoAWQ+Marlin is superseded uniformly, but that it occupies a specific operating region. It is most compelling for decode-heavy serving where linear-layer weight traffic dominates and where the available kernels are well matched to the target GPU. It is less compelling when the architecture-kernel match is poor, as in the reported Ada L4 case, or when runtime overheads outside GEMM dominate, as in H100 batch-1 decode. A plausible implication is that the future performance of the path depends less on changes to AWQ itself than on architecture-tuned MARLIN variants, broader fusion, and better collapse of launch-side overhead.