Low-Precision Quantization Kernels in Production: Comparing Marlin, ExLlamaV2, FlashInfer, and BitBLAS Architecture, Memory Bandwidth, and Decoding Throughput

Autoregressive large language model (LLM) serving operates under two distinct compute regimes: a compute-bound prefill phase and a memory-bandwidth-bound decode phase. While processing the initial prompt involves matrix-matrix multiplications (GEMM) with high arithmetic intensity, generating tokens one by one requires matrix-vector multiplications (GEMV) or small-batch GEMMs. In this decoding regime, the GPU spends the vast majority of its clock cycles moving billions of weight parameters from High Bandwidth Memory (HBM or GDDR) into on-chip SRAM caches and registers, leaving streaming multiprocessor (SM) compute units largely idle.

Post-training quantization (PTQ) addresses this memory wall by compressing weight representations from 16-bit floating-point (FP16/BF16) down to 8-bit, 4-bit, 2-bit, or ternary formats. Mathematically, 4-bit weight quantization reduces memory traffic by up to 4x, offering a proportional theoretical ceiling for decoding throughput. However, converting these theoretical memory savings into real-world speedups requires specialized low-level GPU execution kernels. Standard NVIDIA Tensor Cores execute matrix math in FP16, BF16, FP8, or INT8. Consequently, low-precision weights must be fetched, unpacked, dequantized, scaled, and accumulated on the fly without introducing latency bottlenecks.

Over the past two years, the open-source systems ecosystem has evolved beyond naive dequantization loops toward highly specialized kernel architectures. Four frameworks define the current state of the art in low-precision execution: Marlin (IST Austria, Neural Magic, ETH Zurich), ExLlamaV2 (turboderp), FlashInfer (UW SAMPL), and BitBLAS (Microsoft Research). Understanding their internal memory layouts, warp scheduling models, and compilation strategies determines how production inference platforms maximize goodput across diverse serving workloads.

1. Marlin: Asynchronous Warp Pipelining for Batched Serving

Early open-source 4-bit kernels (such as initial AutoGPTQ and AutoAWQ implementations) achieved solid acceleration at batch size $B=1$. However, their speedup relative to FP16 baselines collapsed rapidly as batch sizes expanded to $B \ge 8$. At higher concurrency, these naive kernels suffered from register pressure, thread divergence, and uncoalesced memory accesses, falling back into compute-bound stalls before saturating GPU memory bandwidth.

Developed by Elias Frantar et al. and detailed in MARLIN: Mixed-Precision Auto-Regressive Parallel Inference on Large Language Models, the Marlin kernel was engineered specifically to maintain near-ideal (~4x) speedups across batch sizes ranging from $B=1$ to $B=32$, and sustained acceleration up to $B=128$.

Key Architectural Mechanisms

1. Pre-Permuted Weight Layouts: Standard 4-bit weights are packed as contiguous nibbles in row-major or column-major formats. To feed NVIDIA Tensor Cores via the mma.sync or ldmatrix hardware instructions, weights must be rearranged into specific register layouts across 32 threads in a warp. Naive kernels perform this layout transformation dynamically at runtime, burning ALU cycles. Marlin eliminates runtime permuting entirely by reordering and repacking the quantized weight tensor offline during model preparation. When Marlin loads a 128-bit global memory chunk, the bits are already arranged to land directly in the exact thread registers required by the Tensor Core instructions.
2. Asynchronous Global Memory Pipelining: Marlin leverages hardware asynchronous copy instructions (cp.async introduced in NVIDIA Ampere architectures). Thread blocks issue non-blocking global memory loads directly into shared memory (SRAM), bypassing intermediate register files. By implementing double-buffered or multi-stage shared memory pipelines, Marlin overlaps the latency of fetching the next weight tile with the Tensor Core computation of the current tile.
3. Optimized Thread Block Tiling: Marlin partitions output matrix tiles across Streaming Multiprocessors such that each warp specializes in a dense $16 \times 16 \times 16$ or $16 \times 16 \times 32$ matrix multiply-accumulate (MMA) tile. For intermediate batch sizes (e.g., $B=16$), Marlin dynamically adjusts its K-dimension reduction strategy to ensure that all available SMs on large GPUs (such as the 108 SMs on an RTX 4090 or 132 SMs on an H100) remain fully saturated.

In production engines such as vLLM and SGLang, Marlin serves as the high-throughput backend for both GPTQ and AWQ 4-bit and 8-bit checkpoints (gptq_marlin, awq_marlin), achieving over 80% of peak theoretical GPU memory bandwidth utilization during high-concurrency decode loops.

2. ExLlamaV2: Hand-Tuned Assembly and Variable Bitrate Quantization

While Marlin is architected for concurrent serving engines, ExLlamaV2 (authored by turboderp) is engineered for minimal latency in single-stream and low-concurrency ($B=1\dots 4$) inference. ExLlamaV2 pairs custom handwritten CUDA kernels with the proprietary EXL2 quantization format.

The EXL2 Mixed-Precision Format

Traditional PTQ formats enforce uniform bitwidths across all layers (e.g., every linear projection is quantized to exactly 4 bits). However, transformer layers exhibit highly non-uniform sensitivity to quantization noise: attention output projections ($W_o$) and down-projections ($W_{down}$) typically suffer greater degradation than key/value projections ($W_k, W_v$) or gate projections ($W_{gate}$).

EXL2 allows sub-module and per-tensor precision mixing across 2, 3, 4, 5, 6, and 8 bits. By measuring the second-order quantization error on calibration data across multiple trial bitwidths, ExLlamaV2 solves a global optimization problem to hit an arbitrary target average bitrate (such as 3.2, 4.25, or 5.0 bits per weight) while assigning higher precision to error-sensitive matrices.

Kernel Execution Mechanics

1. Register-Resident Dequantization: ExLlamaV2 avoids intermediate memory allocations. In decode mode ($B=1$), the kernel maps the incoming FP16 activation vector into registers and streams the quantized weights directly through fast arithmetic unpacking pipelines. Nibbles are unpacked into temporary registers, scaled by group-level floating-point factors, and immediately multiplied using SIMD multiply-accumulate operations.
2. Minimal Kernel Launch Overhead: ExLlamaV2 bypasses the heavy dispatch layers of general-purpose frameworks. Its forward pass is written as a tight C++/CUDA execution loop that minimizes CUDA API call overheads, maximizing token generation speeds on consumer hardware (e.g., RTX 3090, RTX 4090, RTX 5090).
3. Fused Output Sampling: ExLlamaV2 integrates logit computation and token sampling directly with layer normalization kernels, avoiding redundant memory roundtrips at the end of the transformer pipeline.

In single-user scenarios, ExLlamaV2 consistently outpaces general-purpose frameworks in raw tokens-per-second per stream, making it the dominant execution runtime for local desktop engines and low-concurrency API proxies like TabbyAPI.

3. FlashInfer: Composable Serving Operators and Split-K Scheduling

Developed by the SAMPL group at the University of Washington and integrated into next-generation serving frameworks, FlashInfer approaches LLM acceleration through a modular operator architecture. Rather than treating quantization as an isolated matrix multiplication problem, FlashInfer unifies attention computation, low-precision GEMM, sampling, and inter-GPU communication into a composable library.

Key Architectural Mechanisms

1. Split-K Parallelization for Small Batches: When executing GEMV or small-batch GEMM on GPUs with over 100 Streaming Multiprocessors, the number of output tiles (governed by batch size $M$ and output dimension $N$) is often smaller than the number of available SMs. Standard matrix multiplication kernels leave dozens of SMs idle. FlashInfer addresses this with split-K reduction: the inner dimension $K$ (e.g., hidden dimension 4096 or 8192) is split across multiple independent thread blocks across different SMs. Each block computes a partial accumulation, followed by an efficient deterministic reduction in global memory.
2. Paged KV Cache Direct Integration: FlashInfer's low-precision and attention operators execute directly over non-contiguous paged memory layouts (BatchDecodeWithPagedKVCacheWrapper). This eliminates the memory copy overheads that traditional runtimes incur when transferring paged memory into temporary contiguous tensors before calling low-level kernels.
3. CUDA Graph and JIT Template Instantiation: To eliminate kernel launch overheads in high-frequency serving, FlashInfer kernels are designed to be fully captured within static NVIDIA CUDA Graphs. FlashInfer uses Just-In-Time (JIT) C++ template compilation to generate architecture-tuned binary artifacts for specific GPU generations (Ada Lovelace, Hopper, Blackwell), optimizing vector register widths and shared memory bank alignment.

FlashInfer serves as the primary attention and low-precision kernel backend for SGLang and is increasingly adopted in vLLM as the default backend for Hopper (H100/H200) and Blackwell (B200) infrastructure.

4. BitBLAS: Microsoft's Auto-Tuning Sub-Byte Tensor Compiler

While Marlin, ExLlamaV2, and FlashInfer rely on handcrafted CUDA/C++ kernels tailored to specific bit configurations (primarily 4-bit and 8-bit), Microsoft Research developed BitBLAS to automate the generation of mixed-precision BLAS operators for arbitrary sub-byte precisions.

Built on top of Apache TVM and Microsoft's TileLang infrastructure (stemming from research presented in Ladder: Enabling Efficient Low-Precision Deep Learning Computing through Hardware-Aware Tensor Transformation), BitBLAS operates as a domain-specific tensor compiler.

Sub-Byte Compilation Paradigm

1. Arbitrary Bitwidth Support: BitBLAS natively supports weight precisions spanning 1-bit (binary), 2-bit (ternary, as used in BitNet b1.58), 4-bit (INT4, FP4), and 8-bit (INT8, FP8), alongside arbitrary activation formats (FP16, BF16, INT8, FP8).
2. Hardware-Aware Layout Transformation: Sub-byte arithmetic (such as 1.58-bit or 2-bit weights) lacks native hardware instructions on standard Tensor Cores. BitBLAS formulates the unpacking and bit-manipulation steps as symbolic tensor expressions. The compiler analyzes the target GPU's instruction set architecture (ISA) and synthesizes optimal SIMD bit-shifting sequences, packing multiple sub-byte weights into 32-bit registers before issuing Tensor Core MMA instructions.
3. Polyhedral Search and Auto-Tuning: Generic CUDA kernels rarely achieve uniform efficiency across different matrix shapes (such as intermediate FFN expansions vs. attention projections). BitBLAS employs an automated search engine that explores:

• Thread block tile shapes ($M_{tile} \times N_{tile} \times K_{tile}$)
• Warp arrangement within thread blocks
• Shared memory double-buffering stages
• Memory access vectorization widths (e.g., 128-bit int4 vectorized loads)
• Register allocation limits to prevent register spilling

By auto-tuning kernels specifically for each matrix dimension in a target model checkpoint on the host hardware, BitBLAS frequently matches or outperforms hand-written kernels for esoteric bitwidths (1-bit, 2-bit, and asymmetric FP4).

5. Architectural Comparison and Roofline Trade-Offs

The selection of a quantization execution kernel dictates where an inference pipeline operates on the GPU roofline curve:

| Feature / Dimension | Marlin | ExLlamaV2 | FlashInfer | BitBLAS | | :--- | :--- | :--- | :--- | :--- | | Primary Design Target | High-concurrency batched serving ($B=8\dots 128$) | Single-stream & low-batch decode ($B=1\dots 4$) | Modular serving operators & paged attention | Arbitrary sub-byte compilation ($1\dots 8$ bits) | | Supported Formats | GPTQ, AWQ (INT4, INT8, FP8) | EXL2 (2–8 bpw variable), GPTQ | INT4, INT8, FP8, FP16 | INT1, INT2, INT4, FP4, INT8, BitNet | | Memory Pipelining | Asynchronous cp.async multi-stage | Direct register unpacking | Split-K reduction & paged KV | Auto-tuned TileLang / TVM schedules | | Weight Transformation | Offline offline pre-permutation | Offline calibration packing | Dynamic / pre-formatted | Compiler-generated layout packing | | Compilation Model | Pre-compiled Ahead-of-Time (AOT) CUDA | Pre-compiled AOT C++/CUDA | JIT template instantiation + CUDA Graphs | JIT polyhedral auto-tuning compiler | | Primary Ecosystem | vLLM, SGLang, TensorRT-LLM | TabbyAPI, ExLlamaV2 CLI, text-generation-webui | SGLang, vLLM, TensorRT-LLM | Custom serving, TVM, PyTorch / TileLang |

Decoding Roofline Dynamics

1. Single-Stream ($B=1$): At batch size 1, memory latency dominates. ExLlamaV2 and BitBLAS minimize instruction overhead and kernel launch latency. ExLlamaV2 achieves the lowest single-stream time-per-output-token (TPOT) because its execution path contains virtually no runtime abstraction layers.
2. Medium Batches ($B=4\dots 32$): As batch size increases, the workload transitions from pure GEMV toward GEMM. Naive kernels experience immediate bandwidth degradation due to warp stall cycles. Marlin's pre-permuted layouts and asynchronous memory double-buffering allow it to maintain linear scaling, extracting up to 4x throughput gains over unquantized FP16 baselines.
3. High Batches ($B \ge 64$): At high concurrency, matrix multiplications become compute-bound. Here, weight-only quantization (W4A16) must perform dequantization while computing at FP16 Tensor Core speeds, causing the throughput advantage over FP16/FP8 to narrow. Under these conditions, FP8 weight-activation (W8A8 / FP8) or FP4 execution paths outpace weight-only INT4 kernels.

6. Production Deployment Strategy

When architecting a production inference cluster, engineering teams should match their quantization kernel backend to their traffic profile and hardware infrastructure:

1. Multi-Tenant Serving Platforms (vLLM / SGLang on A100 / H100): Deploy Marlin-quantized checkpoints (AWQ-Marlin or GPTQ-Marlin). For workloads handling concurrent request streams with continuous batching and prefix caching, Marlin delivers the highest aggregate token throughput while preventing latency degradation during traffic spikes.
2. Local AI Assistants and Desktop Tooling (Consumer RTX 3090 / 4090 / 5090): Deploy ExLlamaV2 with EXL2 format. For single-tenant or low-concurrency workloads where per-token generation latency ($>100\text{ tok/s}$) is the primary metric, EXL2 provides the best Pareto frontier of model perplexity to decoding speed.
3. Advanced Research and Non-Standard Architectures (1-Bit, BitNet, FP4): Deploy BitBLAS. When serving models with custom bitwidths (such as 1.58-bit ternary networks or experimental 2-bit LoRA adapters), BitBLAS generates optimal hardware-tuned kernels without requiring months of manual CUDA engineering.
4. Unified Heterogeneous Clusters (Hopper & Blackwell): Leverage FlashInfer within serving runtimes to manage paged KV caches and split-K kernel execution, ensuring deterministic latencies within CUDA Graph captures.

Sources

• MARLIN: Mixed-Precision Auto-Regressive Parallel Inference on Large Language Models (Frantar et al., 2024)
• ExLlamaV2: Fast Inference Library for Local LLMs (turboderp)
• FlashInfer: High-Performance LLM Serving Kernel Library (UW SAMPL)
• BitBLAS: Mixed-Precision Matrix Multiplication Library (Microsoft Research)
• Ladder: Enabling Efficient Low-Precision Deep Learning Computing through Hardware-Aware Tensor Transformation (OSDI 2024)
• The Era of 1-bit LLMs: All Large Language Models are in 1.58 Bits (Wang et al., 2024)
• vLLM: Easy, Fast, and Cheap LLM Serving with PagedAttention (Kwon et al., 2023)
• SGLang: Fast and Expressive LLM Serving (Zheng et al., 2023)