4.8 Quantization II: INT4/INT8/FP8, GGUF, bitsandbytes & QAT¶

Post-training quantization (PTQ) compresses a trained model by changing the numeric format of its weights and, optionally, its activations. The previous chapter, Quantization I: Post-Training Quantization (GPTQ, AWQ, SmoothQuant), covered the calibration-based algorithms that decide how to pick the right scale factors. This chapter covers the formats those algorithms produce — INT4, NF4, INT8, FP8, and the GGUF k-quant family — and the inference runtimes and training techniques built around them. We also cover quantization-aware training (QAT) and QLoRA, which push accuracy further at the cost of more compute during training.

The core tension throughout is this: modern LLMs are memory-bandwidth-bound during autoregressive decoding (see The Roofline Model & Performance Engineering), so using fewer bits per weight almost always helps throughput, but every bit removed increases the chance that the rounding error degrades generation quality beyond an acceptable threshold.

The Quantization Landscape: A Taxonomy¶

Before diving into formats, it helps to fix terminology.

Granularity refers to how many weights share a single scale (and zero-point). Options range from per-tensor (one scalar), to per-row/per-column, to per-group (e.g., a block of 128 consecutive weights). Finer granularity costs more metadata but reduces quantization error dramatically — INT4 per-group-128 loses far less information than INT4 per-tensor.

Scope refers to what gets quantized:

• Weight-only quantization (W-only): Weights are stored in low precision; activations remain in BF16/FP16 at runtime. The kernel dequantizes weights on-the-fly and performs the GEMM in FP16. Memory footprint shrinks; arithmetic intensity of the GEMM itself is unchanged, but you win because you now move fewer bytes from HBM.
• Weight + activation quantization (W+A): Both weights and activations are quantized, usually to INT8 or INT8+INT4. The entire matrix multiply happens in low-precision integer arithmetic on hardware integer units, which can deliver higher TFLOP/s than FP16 on some GPU generations. The tradeoff is that activation distributions are much more dynamic and harder to quantize accurately.

Symmetric vs. asymmetric: Symmetric quantization maps \([-\alpha, +\alpha]\) linearly to \([-2^{b-1}, 2^{b-1}-1]\) — the zero-point is always 0, which simplifies dequantization math. Asymmetric allows a nonzero zero-point \(z\) to shift the representable range, accommodating one-sided activation distributions (e.g., post-ReLU activations that are all positive).

The linear quantize-dequantize pair for a weight \(w\) with scale \(s\) and zero-point \(z\) is:

The quantization error per element is bounded by \(\frac{s}{2}\), and if we model the rounding residual as uniform on \([-s/2, s/2]\) its variance is \(s^2/12\) — the standard round-off noise model, and a useful sanity check when you measure reconstruction MSE in code. It also makes the scale/clipping tradeoff precise. Setting \(s\) from the group’s absolute maximum (“absmax”, what every snippet in this chapter does) guarantees no clipping, but a single outlier weight then inflates \(s\) — and hence the error on all \(g-1\) well-behaved weights — for the whole group. The MSE-optimal scale is generally smaller than absmax: you accept clipping error on one or two extreme weights to shrink rounding error everywhere else. This is why real PTQ implementations do not use absmax directly but grid-search a clipping ratio \(\alpha \in (0, 1]\) with \(s = \alpha \max(|\mathbf{w}|)/q_\text{max}\), minimizing either reconstruction MSE or, better, the output error \(\lVert \mathbf{X}\mathbf{W}^\top - \mathbf{X}\hat{\mathbf{W}}^\top \rVert^2\) on calibration activations. Minimizing output error rather than weight error is exactly what separates GPTQ, AWQ, and SmoothQuant from naive round-to-nearest — see Quantization I: Post-Training Quantization (GPTQ, AWQ, SmoothQuant).

INT8: The Safe Harbor¶

INT8 is the most widely deployed quantization format because the accuracy penalty is usually negligible and both Tensor Core (via mma.sync) and integer ALU paths are mature.

INT8 Weight-Only (LLM.int8)¶

Tim Dettmers et al. introduced LLM.int8() as part of bitsandbytes. The key insight was that large language models have a small fraction of outlier activation channels — typically 0.1–1 % of channels depending on model size — that take values far outside the typical range. Quantizing these outliers with per-tensor INT8 causes catastrophic error.

The solution is mixed-precision decomposition: identify the handful of outlier columns at runtime, keep those multiplications in FP16, and quantize the rest as INT8 per-column. At 6.7 B parameters and above, this approach nearly eliminates the accuracy gap with FP16 while halving the memory footprint.

# bitsandbytes INT8 weight-only quantization (load_in_8bit)
# Requires: pip install bitsandbytes transformers accelerate

import torch
from transformers import AutoModelForCausalLM, AutoTokenizer

model_id = "meta-llama/Meta-Llama-3-8B-Instruct"

# `load_in_8bit=True` triggers LLM.int8() decomposition via bitsandbytes
model = AutoModelForCausalLM.from_pretrained(
    model_id,
    load_in_8bit=True,       # weight-only INT8; activations remain FP16
    device_map="auto",       # spread layers across available GPUs
    torch_dtype=torch.float16,
)
tokenizer = AutoTokenizer.from_pretrained(model_id)

# The model is now ~8 GB instead of ~16 GB for BF16
# Linear layers become bitsandbytes.nn.Linear8bitLt modules
for name, module in model.named_modules():
    if "Linear8bitLt" in type(module).__name__:
        print(f"{name}: weight dtype={module.weight.dtype}")
        break

INT8 Weight + Activation (SmoothQuant / TensorRT-LLM)¶

For weight+activation INT8, the serving stack typically uses per-token dynamic quantization of activations: at each layer, collect the max absolute value of the current token’s activation vector, compute a scale, quantize to INT8, run the INT8 GEMM on Tensor Cores, and dequantize the output. This is done via CUTLASS or TensorRT-LLM’s int8_sq plugin.

SmoothQuant (Xiao et al., 2022) migrates some of the per-channel outlier variance from activations into weights by applying a per-channel scaling factor \(s_j\) before quantization:

After rescaling, both sides are smoother and quantize more cleanly. This is detailed further in Quantization I: Post-Training Quantization (GPTQ, AWQ, SmoothQuant).

INT4 & NF4: Pushing to 4 Bits¶

4-bit quantization halves memory relative to INT8, enabling a 70 B model to fit on a single 80 GB A100 at around 35 GB. The challenge is that rounding errors are much larger with only 16 representable values.

Per-Group INT4¶

The standard INT4 scheme uses group quantization: every \(g\) consecutive weights (often \(g = 128\)) share one scale (and optionally one zero-point). The metadata overhead is \(\frac{16}{g}\) bits per weight (one FP16 scale per group), negligible at \(g = 128\).

For a weight matrix \(\mathbf{W} \in \mathbb{R}^{d_\text{out} \times d_\text{in}}\), each row is split into \(\frac{d_\text{in}}{g}\) groups. Each group is quantized to INT4 values in \([-8, 7]\) (signed) or \([0, 15]\) (unsigned with zero-point). GPTQ and AWQ both output this format; the difference is only in how the scales are found.

Dequantization at inference time:

where \(s\) and \(z\) are the group scale and zero-point. The kernel unpacks two 4-bit values from each byte, dequantizes to FP16, and then calls the FP16 GEMM.

NF4: Normal Float 4¶

NF4 (Dettmers et al., QLoRA, 2023) is a 4-bit data type whose code points are placed at the quantiles of a standard normal \(\mathcal{N}(0,1)\) rather than at uniform spacing. The idea: each of the 16 codes should be used about equally often, so no code is wasted on a region of the distribution where weights never land. The QLoRA paper calls this “information-theoretically optimal” for normally distributed data — read that as equal-probability bins, not as MSE-optimal. A Lloyd–Max quantizer (which minimizes expected squared error for a given density) would place its levels slightly differently; equal-probability quantiles are the entropy-maximizing choice, and the two coincide only in the high-resolution limit. In practice the difference at 4 bits is small, and NF4’s real advantage over plain INT4 comes from spending resolution near zero where most weights actually are.

The construction has two details worth knowing, because they explain the odd-looking constants in the code table below:

1. The grid is asymmetric, with an exact zero. A symmetric 16-level quantile grid has no code at \(0\), but a code for exactly zero is valuable (most weights are near zero, and it makes an exactly-zero weight exactly representable). NF4 therefore builds the two halves separately — 8 levels on the positive side, 7 on the negative — and inserts \(0\) as the 16^th code. That is why the table below has 7 negative entries, a zero, and 8 positive entries.
2. The tails are clipped at a finite quantile. \(Q_\mathcal{N}(1) = \infty\), so the extreme levels are taken at an offset \(p_{\max} = 0.9677083\) rather than at \(1\) — the average of the two “half-bin-in-from-the-end” quantiles for a 15-level and a 16-level grid, \(\tfrac{1}{2}\big[(1 - \tfrac{1}{30}) + (1 - \tfrac{1}{32})\big]\). The whole set is then divided by its largest magnitude so the codes span exactly \([-1, +1]\).

At runtime the group scale is simply the group’s absolute maximum, \(s = \max(|\mathbf{w}|)\), so the largest-magnitude weight in each block maps onto the extreme code \(\pm 1\); the nearest code point index is then stored as 4 bits.

NF4 is the weight format used by QLoRA for the frozen base model. It achieves slightly lower perplexity than INT4 per-group on the same model at the same 4-bit budget, because its code points are better matched to the actual weight distribution.

# bitsandbytes NF4 quantization (load_in_4bit with bnb_4bit_quant_type="nf4")
import torch
from transformers import AutoModelForCausalLM, BitsAndBytesConfig

bnb_config = BitsAndBytesConfig(
    load_in_4bit=True,
    bnb_4bit_quant_type="nf4",       # vs "fp4" (less accurate but faster unpacking)
    bnb_4bit_compute_dtype=torch.bfloat16,  # dtype for the dequantized compute
    bnb_4bit_use_double_quant=True,  # double quantization (see below)
)

model = AutoModelForCausalLM.from_pretrained(
    "meta-llama/Meta-Llama-3-8B-Instruct",
    quantization_config=bnb_config,
    device_map="auto",
)
# Model is now ~4 GB instead of ~16 GB

Double Quantization¶

Double quantization (also from QLoRA) applies a second quantization step to the scale factors themselves. bitsandbytes’ NF4 uses a default block size of \(g = 64\), and each block’s absmax scale is an FP32 value — that is \(32/64 = 0.5\) bits per weight of pure metadata, which is a lot when the payload is only 4 bits. Double quantization quantizes those FP32 absmaxes to 8 bits (bitsandbytes first subtracts their mean so the values are centered, then applies an 8-bit dynamic-exponent map), grouping \(g_2 = 256\) scales under one FP32 meta-scale. The overhead becomes

a saving of about \(0.373\) bits/weight — which the QLoRA paper reports as roughly 3 GB on a 65 B model. Note the metadata is quantized to 8-bit integers, not FP8; FP8 hardware types are not involved here at all.

FP8 Inference¶

FP8 introduces a floating-point format at 8 bits. Two variants exist:

• E4M3: 4 exponent bits, 3 mantissa bits, exponent bias 7. Representable magnitudes run from the smallest subnormal \(2^{-9} \approx 1.95 \times 10^{-3}\) up to a maximum of \(448\) (the OCP FP8 spec reclaims the all-ones exponent pattern for NaN instead of infinities, which is how you get 448 rather than 240). Better for weights and activations, which need mantissa bits more than range.
• E5M2: 5 exponent bits, 2 mantissa bits, exponent bias 15. Range up to \(57344\) with a smallest subnormal of \(2^{-16} \approx 1.5 \times 10^{-5}\) — much wider dynamic range, one fewer mantissa bit. Better for gradients during training, whose magnitudes span many more orders of magnitude.

NVIDIA Hopper GPUs (H100, H200) introduced native FP8 Tensor Core support, and the Blackwell generation (B200, GB200) carries FP8 forward while adding native 4-bit floating-point Tensor Cores as well. The hardware performs FP8 × FP8 multiplications and accumulates in FP32, then scales and stores results in FP8 or FP16.

FP8 differs from INT8 in one crucial way: because the exponent field already gives each element its own scale, the explicit scale factor can be much coarser — per-tensor or per-row rather than per-group-128. On Hopper the scaling happens around the GEMM, not inside it: the Tensor Core multiplies FP8 operands and accumulates in FP32, and the scale is folded into the epilogue (this is exactly what torch._scaled_mm exposes). Blackwell changes that by adding block-scaled MMA instructions that consume a per-block scale factor directly from the instruction operands, which is what makes MXFP8/MXFP4/NVFP4 cheap in hardware. The practical consequence of coarse scaling is that FP8 shines on GEMM-heavy forward passes (attention and FFN projections) but can degrade more than per-group INT4/INT8 on outlier-heavy activations unless paired with SmoothQuant-style smoothing or with fine-grained FP8 — the recipe DeepSeek-V3 popularized, quantizing activations in \(1\times128\) tiles and weights in \(128\times128\) blocks so that no single outlier channel sets the scale for a whole tensor.

On Blackwell, NVFP4 pushes this idea one step further into 4-bit floating point (E2M1: 1 sign, 2 exponent, 1 mantissa bit). It uses small 16-value micro-blocks, each with its own FP8 (E4M3) scale factor plus a single FP32 scale for the whole tensor — finer-grained than the 32-value blocks used by the open MXFP4 standard. NVIDIA reports roughly 1.8× lower memory footprint than FP8 with well under 1% accuracy degradation on language-modeling benchmarks; vLLM and TensorRT-LLM both support NVFP4 checkpoints (e.g., vLLM’s quantization="modelopt_fp4"), and it is now being used for pretraining as well as inference (see Pretraining Large Language Models with NVFP4 below), not just as a serving-time format.

FP8 training is discussed in Mixed Precision, bf16 & FP8 Training and FlashAttention 2 & 3: Work Partitioning, Warp Specialization & FP8.

# FP8 inference via TensorRT-LLM (conceptual API sketch)
# In practice, TRT-LLM's quantization workflow handles the conversion
# The snippet below shows the key concepts using transformer_engine directly.

import torch
import transformer_engine.pytorch as te
from transformer_engine.common.recipe import Format, DelayedScaling

# Create FP8 recipe: E4M3 for forward, E5M2 for backward (if training)
fp8_recipe = DelayedScaling(
    fp8_format=Format.E4M3,
    amax_history_len=16,      # track amax over last 16 iters to set scale
    amax_compute_algo="max",
)

# A TransformerEngine linear layer that uses FP8 GEMM on Hopper
linear = te.Linear(4096, 4096, bias=False)

with te.fp8_autocast(enabled=True, fp8_recipe=fp8_recipe):
    x = torch.randn(1, 512, 4096, device="cuda", dtype=torch.bfloat16)
    y = linear(x)  # internally uses FP8 GEMM
    print(f"Output dtype: {y.dtype}")  # still BF16 after cast-back

For practical FP8 inference with vLLM, the quickest path is on-the-fly quantization at load time — no calibration, per-tensor dynamic scales:

# Quantize weights to FP8 as the model loads (needs Ada/Hopper/Blackwell).
# Zero setup, but you pay the conversion cost on every server start and you
# get no activation calibration.
vllm serve meta-llama/Meta-Llama-3-70B-Instruct \
    --quantization fp8 \
    --dtype bfloat16 \
    --gpu-memory-utilization 0.92

The Checkpoint Format Layer: compressed-tensors and llm-compressor¶

On-the-fly quantization is convenient but wasteful. In production you quantize once, offline, and ship a checkpoint the server can mmap straight into the right kernel. The format that has consolidated this in the vLLM/SGLang ecosystem is compressed-tensors — a safetensors-compatible checkpoint format that stores packed low-bit weights alongside a quantization_config describing scheme, group size, symmetry, and which layers were left in BF16. It is what replaced the proliferation of one-format-per-algorithm checkpoints (gptq, awq, marlin, …) with a single serialization that a runtime can inspect and route to the best available kernel.

The tool that writes those checkpoints is llm-compressor (the vLLM Project’s toolkit, successor to Neural Magic’s SparseML). The algorithms it runs — GPTQ, AWQ, SmoothQuant — are covered in Quantization I; what matters here is that one oneshot() call selects the format you deploy:

# pip install llmcompressor
# Produce an offline FP8 W8A8 checkpoint that vLLM loads with no conversion cost.
from llmcompressor import oneshot
from llmcompressor.modifiers.quantization import QuantizationModifier

MODEL_ID = "meta-llama/Meta-Llama-3-8B-Instruct"

recipe = QuantizationModifier(
    targets="Linear",              # quantize every nn.Linear ...
    scheme="FP8_DYNAMIC",          # ... to FP8 weights + dynamic per-token FP8 acts
    ignore=["lm_head"],            # ... except the output head, which is sensitive
)

# FP8_DYNAMIC needs no calibration data: weight scales come from the weights,
# activation scales are computed per token at runtime.
oneshot(model=MODEL_ID, recipe=recipe, output_dir="Llama-3-8B-Instruct-FP8")

# Swap `scheme` for other formats this chapter covers:
#   "FP8"        -> static per-tensor FP8 activations (needs calibration data)
#   "W8A8"       -> INT8 weights + INT8 activations (add a SmoothQuant modifier)
#   "W4A16"      -> INT4 group-128 weight-only (pair with GPTQModifier + calib set)
#   "NVFP4"      -> 4-bit float, Blackwell block-scaled kernels
# Then simply: vllm serve Llama-3-8B-Instruct-FP8

NVIDIA’s parallel tool for its own stack is TensorRT Model Optimizer (nvidia-modelopt), which produces the NVFP4/FP8 checkpoints TensorRT-LLM consumes and that vLLM loads via --quantization modelopt / modelopt_fp4.

Which Kernel Actually Runs¶

A W4A16 checkpoint is useless without a fast kernel, and this is where most of the throughput spread in the table below comes from. vLLM’s INT4/FP8 weight-only path runs on Marlin (Frantar et al.), a mixed-precision GEMM that hides dequantization behind the memory pipeline and holds near-4× speedup out to batch sizes of 32–64, where naive dequantize-then-FP16-GEMM has long since collapsed to BF16 speed. Machete is its Hopper successor, built on CUTLASS 3.x and wgmma with weights pre-shuffled at load time. The lesson generalizes: at batch size 1 any correct INT4 kernel wins because you are bandwidth-bound; at moderate batch the kernel, not the format, decides whether weight-only quantization is still a win.

GGUF & llama.cpp K-Quants¶

GGUF (GPT-Generated Unified Format) is the binary container format used by llama.cpp, the most widely deployed CPU/edge inference engine. It replaces the older GGML format and stores model weights alongside all necessary metadata (tokenizer, architecture hyperparameters, rope parameters, etc.) in a single self-describing file.

The K-Quant Family¶

llama.cpp’s k-quants are a family of mixed-precision block quantization schemes contributed by Iwan Kawrakow (hence the “k”). The key innovation is two-level block scaling: each super-block of 256 weights carries fp16 super-block scales, and those in turn scale a set of cheap low-bit sub-block scales, one per 16 or 32 weights. You get near-per-32 scale granularity for a fraction of the metadata cost of storing an fp16 scale every 32 weights.

The block layouts and their exact bits/weight (derived from the block_* structs in ggml-quants.h):

The _S / _M / _L suffixes you see on filenames (Q4_K_M, Q3_K_L) are not different block layouts — they are tensor-level mixes. llama-quantize applies the named type to most tensors but promotes the ones empirically most sensitive to quantization error (typically attn_v and ffn_down, plus the output/embedding matrices) to a higher k-quant. Q4_K_S is Q4_K nearly everywhere; Q4_K_M stores some of those sensitive tensors at Q6_K. That is why a real Q4_K_M file measures a few tenths of a bit per weight above the 4.5-bit block figure, and why _M is the recommended default: the extra bits go exactly where they buy the most quality.

Importance-matrix (imatrix) quantization and the IQ family. Plain k-quants minimize weight reconstruction error, treating all weights as equally important. llama.cpp’s llama-imatrix tool fixes that by running calibration text through the fp16 model and recording, per weight column, the mean squared activation that multiplies it — an importance matrix, the same activation-awareness idea behind AWQ (Quantization I). Passing --imatrix to llama-quantize weights the rounding search by that importance and is close to mandatory below 4 bits. Built on top of it is the IQ family (IQ1_S … IQ4_XS), which replaces scalar rounding with lookup into a fixed codebook of lattice points — vector quantization in small groups — reaching genuinely usable quality at 2–3 bits/weight where Q2_K struggles.

GGUF has also picked up native 4-bit floating-point support alongside the integer k-quants: OpenAI’s open-weight gpt-oss models ship natively in MXFP4, and llama.cpp added first-class MXFP4 loading (in collaboration with NVIDIA) so those weights run directly without a lossy conversion to a k-quant format.

Q4_K internal structure: Each super-block covers 256 weights. It stores two fp16 super-block scales — d (applied to the sub-block scales) and dmin (applied to the sub-block mins) — plus 8 sub-block scales and 8 sub-block mins covering 8 sub-blocks of 32 weights each. Those 16 values (8 scales + 8 mins) are each quantized to 6 bits and packed into a 12-byte array. Individual weights are stored as 4-bit unsigned integers. This matches ggml’s block_q4_K struct, which is the source of truth: ggml_half d, dmin; (the two fp16 super-block scales), uint8_t scales[12]; (the sixteen 6-bit sub-block scales and mins), and uint8_t qs[128]; (the 256 4-bit weights). The dequantization formula per weight is:

where \(d\) and \(d_\text{min}\) are the two fp16 super-block scales (for the sub-block scales and mins respectively), \(s_j\) and \(m_j\) are the 6-bit scale and min of the sub-block \(j\) that weight \(\hat{w}\) belongs to, and \(q \in [0, 15]\) is the stored 4-bit weight.

# Convert a Hugging Face model to GGUF Q4_K_M using llama.cpp's converter.
# First clone llama.cpp and build it (the project is CMake-only; the old
# top-level Makefile was removed, and binaries land in ./build/bin/):
#   git clone https://github.com/ggml-org/llama.cpp && cd llama.cpp
#   pip install -r requirements.txt
#   cmake -B build -DGGML_CUDA=ON && cmake --build build --config Release -j

# Step 1: Convert HF model to GGUF F16 (lossless intermediate)
# (Run from the llama.cpp directory)

# Step 1: Lossless F16 conversion
python convert_hf_to_gguf.py \
    /path/to/meta-llama/Meta-Llama-3-8B-Instruct \
    --outfile llama3-8b-f16.gguf \
    --outtype f16

# Step 2 (optional but recommended below 5 bits): build an importance matrix
# from a few MB of calibration text, so the quantizer knows which weight
# columns actually matter.
./build/bin/llama-imatrix \
    -m llama3-8b-f16.gguf \
    -f calibration.txt \
    -o llama3-8b.imatrix

# Step 3: Quantize to Q4_K_M (the recommended default for CPU/edge)
./build/bin/llama-quantize \
    --imatrix llama3-8b.imatrix \
    llama3-8b-f16.gguf llama3-8b-Q4_K_M.gguf Q4_K_M

# Step 4: Run inference
./build/bin/llama-cli \
    -m llama3-8b-Q4_K_M.gguf \
    -n 256 \
    -p "Explain quantization to a 5 year old:" \
    --n-gpu-layers 33   # offload 33 layers to GPU; rest runs on CPU RAM

The --n-gpu-layers flag enables GPU+CPU split inference: the first \(n\) layers run on the GPU (fast), remaining layers on the CPU (using system RAM). This allows running a 70 B model with 24 GB VRAM + 32 GB system RAM — unthinkable with any other stack.

Why GGUF for Edge Deployment?¶

GGUF’s portability is unmatched: the same binary runs on macOS (Metal), Linux (CUDA or CPU), Windows (DirectML or CUDA), and even Android/iOS via llama.cpp bindings. For edge deployment, the k-quant Q4_K_M format on a 7 B model typically results in a ~4.1 GB file that runs at 20–40 tokens/second on a modern CPU — no GPU required.

Exporting Your Own Model to GGUF¶

convert_hf_to_gguf.py does not read arbitrary PyTorch checkpoints — it dispatches on the architectures field of config.json to a registered Model subclass that knows how to rename and reshape that architecture’s tensors. So for a model you trained yourself there are two honest paths:

1. Borrow a supported architecture. If your model is structurally a Llama (RMSNorm, RoPE, SwiGLU, GQA — the Stack-100M recipe), the cheapest route is to write your state dict out under HF’s LlamaForCausalLM naming with a matching config.json, and let the existing converter handle it. The one trap is RoPE layout: HF’s Llama implementation stores Q/K in the “split-half” permutation and the converter un-permutes it, so a model trained with interleaved RoPE must have its Q and K projection rows permuted at export time or generation will be silently garbage.
2. Register a new architecture. Subclass Model in convert_hf_to_gguf.py, declare your tensor-name mapping, and add the matching graph-building code on the C++ side in llama.cpp. This is the real cost of a custom architecture — worth knowing before you invent a novel block.

A related warning for small models: quantization damage does not scale down with parameter count, it scales up. A 100 M-parameter model has far less redundancy to absorb rounding error than a 7 B one, so the 0.1–0.5 perplexity penalty quoted below for Q4_K_M at 7 B can become several points at 100 M. At that scale prefer Q8_0 or Q6_K — the file is still tiny (a ~100 M model is ~110 MB at Q8_0) and you keep essentially all of your quality. The capstone walks the whole export-and-measure loop for a model you built yourself in Evaluation & Serving: Honest Benchmarks, int4 Quantization, and Running on a Laptop.

bitsandbytes: The PyTorch-Native Quantization Library¶

bitsandbytes (bnb) provides drop-in quantized linear layers for PyTorch. It is the primary quantization backend for Hugging Face Transformers and the PEFT library (used by QLoRA).

Architecture of bnb Linear Layers¶

Both Linear8bitLt and Linear4bit are weight-only: the dequantized GEMM still runs in FP16 hardware. The bandwidth saving is in loading weights from HBM; once on-chip (in L2 or registers), the weights are converted to FP16 before multiply-accumulate.

Implementing a Minimal NF4 Layer From Scratch¶

This reconstruction shows the exact mechanism — not production-ready, but pedagogically complete:

import torch
import torch.nn as nn
import numpy as np

# The 16 NF4 code points (from QLoRA paper, normalized to [-1, 1])
NF4_CODES = torch.tensor([
    -1.0,       -0.6961928,  -0.5250730,  -0.3954816,
    -0.2849375, -0.1832600,  -0.0911578,  0.0,
     0.0795761,  0.1609030,   0.2461331,   0.3379990,
     0.4407979,  0.5626170,   0.7229568,   1.0,
], dtype=torch.float32)

def quantize_nf4(weight: torch.Tensor, group_size: int = 64):
    """
    Quantize a 1-D weight tensor to NF4 per-group.
    Returns: (packed_indices, scales) where packed_indices is uint8
    with two 4-bit indices per byte.
    """
    weight = weight.float()
    n = weight.numel()
    assert n % group_size == 0
    n_groups = n // group_size
    w_groups = weight.view(n_groups, group_size)

    # Scale each group so its max absolute value maps to 1.0
    scales = w_groups.abs().max(dim=1).values  # (n_groups,)
    scales = scales.clamp(min=1e-8)
    w_norm = w_groups / scales.unsqueeze(1)    # (n_groups, group_size) in [-1, 1]

    # Find nearest NF4 code point for each weight
    # Broadcast: (n_groups, group_size, 1) vs (16,)
    codes = NF4_CODES.to(weight.device)
    dists = (w_norm.unsqueeze(-1) - codes).abs()  # (n_groups, group_size, 16)
    indices = dists.argmin(dim=-1).byte()           # (n_groups, group_size), dtype=uint8

    # Pack two 4-bit indices into one byte
    indices_flat = indices.view(-1)  # (n,)
    packed = (indices_flat[0::2] << 4) | indices_flat[1::2]  # (n//2,) uint8

    return packed, scales

def dequantize_nf4(packed: torch.Tensor, scales: torch.Tensor, group_size: int = 64):
    """Unpack NF4 indices and reconstruct FP32 weights."""
    # Unpack nibbles
    hi = (packed >> 4).byte()
    lo = (packed & 0xF).byte()
    indices_flat = torch.stack([hi, lo], dim=1).view(-1)  # interleaved back

    n = indices_flat.numel()
    n_groups = n // group_size
    codes = NF4_CODES.to(packed.device)
    w_norm = codes[indices_flat.long()].view(n_groups, group_size)

    # Re-apply group scales
    w_reconstructed = w_norm * scales.unsqueeze(1)
    return w_reconstructed.view(-1)

# --- Demo ---
torch.manual_seed(42)
w = torch.randn(256)           # simulate a weight vector (one row of a linear layer)
packed, scales = quantize_nf4(w, group_size=64)

print(f"Original size:    {w.numel() * 4} bytes (FP32)")
print(f"Quantized size:   {packed.numel()} bytes (NF4 packed)")
print(f"Scales overhead:  {scales.numel() * 4} bytes (FP32 scales)")

w_hat = dequantize_nf4(packed, scales, group_size=64)
mse = ((w - w_hat) ** 2).mean().item()
snr = (w.var() / mse).item()
print(f"MSE:  {mse:.6f}")
print(f"SNR:  {snr:.1f}  (higher is better; >100 is practically lossless)")

Quantization-Aware Training (QAT)¶

Post-training quantization (PTQ) is cheap — no training required — but QAT can close the accuracy gap for aggressive bit widths (INT4 and below) by teaching the model to be robust to quantization noise during training.

The Straight-Through Estimator¶

The core challenge is that the rounding operation \(\operatorname{round}(\cdot)\) has zero gradient almost everywhere. QAT works by using a straight-through estimator (STE) in the backward pass: the forward pass rounds normally, but the backward pass pretends the rounding did not happen and passes gradients through unchanged.

For a quantized weight \(q = \operatorname{round}(w/s)\), the forward pass uses \(q\), and the backward pass computes:

This is a biased estimator, but empirically it works well and allows the model to adjust its weights so that rounding hurts less.

import torch
import torch.nn as nn
import torch.nn.functional as F

class STEQuantize(torch.autograd.Function):
    """
    Quantize to b bits with straight-through estimator in backward.
    """
    @staticmethod
    def forward(ctx, x: torch.Tensor, scale: float, bits: int):
        # Quantize: clamp to representable range, then round
        qmin = -(2 ** (bits - 1))
        qmax =  (2 ** (bits - 1)) - 1
        x_scaled = x / scale
        x_clipped = x_scaled.clamp(qmin, qmax)
        x_quant = x_clipped.round()
        # Store nothing for backward — STE passes gradient directly
        return x_quant * scale  # dequantized immediately (fake-quant)

    @staticmethod
    def backward(ctx, grad_output):
        # STE: pass gradient through unchanged
        return grad_output, None, None

class FakeQuantLinear(nn.Linear):
    """
    A drop-in replacement for nn.Linear that applies fake-quantization
    to weights during the forward pass (simulates INT4 weight quantization).
    """
    def __init__(self, *args, bits=4, group_size=128, **kwargs):
        super().__init__(*args, **kwargs)
        self.bits = bits
        self.group_size = group_size

    def get_scale(self, w: torch.Tensor) -> torch.Tensor:
        """Per-group symmetric scale: s = max(|w|) / (2^(b-1) - 1)"""
        w_groups = w.view(-1, self.group_size)
        s = w_groups.abs().max(dim=1).values / (2 ** (self.bits - 1) - 1)
        return s.unsqueeze(1).expand_as(w_groups).reshape_as(w)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        # Apply fake-quantization to weights
        scale = self.get_scale(self.weight)
        w_fq = STEQuantize.apply(self.weight, 1.0, self.bits)  # simplified
        return F.linear(x, w_fq, self.bias)

# Minimal QAT training loop sketch
def qat_finetune_step(model, batch, optimizer):
    """Replace all Linear layers with FakeQuantLinear, then fine-tune."""
    optimizer.zero_grad()
    outputs = model(**batch)
    loss = outputs.loss
    loss.backward()   # gradients flow through STE
    optimizer.step()
    return loss.item()

Doing QAT For Real: torchao¶

You would not hand-roll the module swap above for a real model. The PyTorch-native library for this is torchao (pytorch/ao) — the same package that provides the float8 training path used by torchtitan (Mixed Precision, bf16 & FP8 Training). Its quantization API is a single in-place quantize_(model, config) that walks the module tree and swaps weights for tensor subclasses carrying the packed data plus scales, so the model stays a normal nn.Module and composes with torch.compile, FSDP2, and torch.export.

QAT in torchao is a two-phase flow, and the phase boundary is the important idea to carry away regardless of library:

1. Prepare. Insert fake-quantization (quantize → dequantize in the same dtype, with an STE backward) into the forward pass. Weights are still BF16 and still trainable; the model is merely simulating the target numerics. Fine-tune for a small number of steps — QAT after pretraining is a fine-tuning-scale cost, typically well under 1 % of the pretraining budget.
2. Convert. Replace the simulation with the real low-bit representation, producing the deployable checkpoint.

# pip install torchao
import torch
from torchao.quantization import quantize_, Int4WeightOnlyConfig
from torchao.quantization.qat import QATConfig

base_config = Int4WeightOnlyConfig(group_size=32)   # the target deployment numerics

# --- Phase 1: prepare (insert fake-quant + STE), then fine-tune normally ---
quantize_(model, QATConfig(base_config, step="prepare"))
for batch in qat_dataloader:                 # a short fine-tune, not a full retrain
    loss = model(**batch).loss
    loss.backward()
    optimizer.step(); optimizer.zero_grad()

# --- Phase 2: convert to the real 4-bit representation ---
quantize_(model, QATConfig(base_config, step="convert"))

# Straight PTQ (no fine-tuning) is the same call without the QAT wrapper:
#   quantize_(model, Int4WeightOnlyConfig(group_size=32))
# torchao's API is versioned; older releases spell the configs as factory
# functions (`int4_weight_only(group_size=32)`) — check the installed version.

The payoff is real but bounded: QAT mainly earns its keep at 4 bits and below, or for small models where PTQ damage is largest. At INT8, or at W4A16 on a 7 B+ model with a good PTQ algorithm, the remaining gap is usually too small to justify the training run. The decision rule: reach for QAT only after GPTQ/AWQ with a proper calibration set has failed to hit your quality bar.

QLoRA: Quantization + Low-Rank Adaptation¶

QLoRA (Dettmers et al., 2023) is arguably the most impactful combination of quantization and fine-tuning. The recipe:

1. Freeze the base model weights in NF4 (4-bit, per-group-64, double quantization).
2. Add LoRA adapters (small rank-\(r\) matrices \(A, B\) in BF16) alongside the frozen quantized layers.
3. Fine-tune only the LoRA adapters. Gradients flow through the NF4-dequantized base weights using STE, then into the BF16 LoRA params.
4. 4-bit NF4 paged optimizer states: instead of keeping FP32 Adam states for the base model, only LoRA params have optimizer states — since they are tiny (\(r \ll d\)), this is cheap.

The key trick: paged optimizers (bnb’s PagedAdamW32bit) keep optimizer states in CPU RAM and page them to GPU only when needed, preventing OOM on long sequences.

# QLoRA fine-tuning with bitsandbytes + PEFT
from transformers import AutoModelForCausalLM, BitsAndBytesConfig
from peft import LoraConfig, get_peft_model, TaskType
import bitsandbytes as bnb
import torch

# 1. Load base model in 4-bit NF4
bnb_config = BitsAndBytesConfig(
    load_in_4bit=True,
    bnb_4bit_quant_type="nf4",
    bnb_4bit_compute_dtype=torch.bfloat16,
    bnb_4bit_use_double_quant=True,
)
base_model = AutoModelForCausalLM.from_pretrained(
    "meta-llama/Meta-Llama-3-8B",
    quantization_config=bnb_config,
    device_map="auto",
)

# 2. Add LoRA adapters (only these parameters will be trained)
lora_config = LoraConfig(
    r=16,                        # rank — try 8–64 depending on task
    lora_alpha=32,               # scaling: effective_lr_scaling = alpha/r
    target_modules=["q_proj", "k_proj", "v_proj", "o_proj",
                    "gate_proj", "up_proj", "down_proj"],
    lora_dropout=0.05,
    bias="none",
    task_type=TaskType.CAUSAL_LM,
)
model = get_peft_model(base_model, lora_config)
model.print_trainable_parameters()
# Trainable params: ~20M / 8B total (0.25%) — massive memory saving

# 3. Use paged optimizer to handle memory spikes
optimizer = bnb.optim.PagedAdamW32bit(
    model.parameters(),
    lr=2e-4,
    weight_decay=0.01,
)

LoRA and the PEFT framework are covered in depth in PEFT I: LoRA, QLoRA, DoRA & The Adapter Family. The memory-efficient training angle is covered in Memory-Efficient Training: Checkpointing, Offloading & LoRA Math.

KV-Cache Quantization¶

In long-context inference, the KV cache can easily rival or exceed model weight memory. For a model with 32 layers, 32 heads, head dimension 128, generating a 32 K-token sequence, each of K and V is:

(That figure assumes multi-head attention; every modern model uses grouped-query attention, which divides it by the query-to-KV head ratio — 4× for Llama 3 8B’s 32 query heads over 8 KV heads. GQA is the first KV-memory lever; quantization stacks on top of it. See Multi-Head Attention, MQA, GQA & MLA.)

Quantizing the KV cache to INT8 halves this to 4.2 GB; INT4 reduces it to 2.1 GB. KV quantization is more delicate than weight quantization because keys and values are computed dynamically (they change every sequence), have heavier-tailed distributions than model weights, and — unlike weights — cannot be calibrated offline against the tensor you will actually quantize.

Per-Token Dynamic Quantization of KV¶

Two axes exist, and which one you quantize along matters more than the bit width:

1. Per-token (per-position) symmetric quantization: one scale per (layer, head, position), computed from the max absolute value across that position’s head_dim elements. This is what the code below does and what most runtimes ship.
2. Per-channel quantization: one scale per (layer, head, channel), shared across positions.

The KIVI result (Liu et al., ICML 2024) is that keys and values want different axes. Key tensors have persistent outlier channels — a few dimensions of \(K\) that are large at every position — so per-token scaling lets one bad channel set the scale for the whole vector; keys should be quantized per-channel. Value tensors show no such channel structure, and because \(V\) is contracted along the position axis by the attention weights, per-token quantization keeps each value’s error independent; values should be quantized per-token. Getting this asymmetry right is what lets KIVI reach 2-bit KV cache without fine-tuning. The practical wrinkle is that per-channel K quantization needs a full channel’s worth of positions before the scale is known, so implementations keep the most recent group of tokens in full precision as a sliding “residual” window and quantize only the settled prefix.

In production, the widely deployed setting is FP8 KV cache rather than INT8: vLLM exposes --kv-cache-dtype fp8 (E4M3 by default, with E5M2 available), which halves KV memory and, on Hopper/Blackwell, feeds the FP8 attention kernel directly with no dequantization step. E4M3’s exponent field absorbs the heavy tails that hurt INT8 here, which is why FP8 is the better default for KV even where INT8 wins for weights.

# Simplified KV cache quantization kernel (conceptual)
import torch

def quantize_kv_int8(kv: torch.Tensor):
    """
    Quantize a KV tensor of shape (batch, heads, seq_len, head_dim) to INT8.
    Per-token (per-position) symmetric quantization.
    Returns int8 tensor + FP16 per-token scale tensor.
    """
    # kv: (B, H, T, D)
    # Compute per-position max-abs across head_dim
    scale = kv.abs().max(dim=-1, keepdim=True).values / 127.0  # (B, H, T, 1)
    scale = scale.clamp(min=1e-8)
    kv_int8 = (kv / scale).round().clamp(-128, 127).to(torch.int8)
    return kv_int8, scale.to(torch.float16)

def dequantize_kv(kv_int8: torch.Tensor, scale: torch.Tensor):
    """Recover approximate FP16 KV from INT8 + scale."""
    return kv_int8.to(torch.float16) * scale

# Memory comparison for a 32-layer, 32-head, 128-dim model at 8K context
B, H, T, D = 1, 32 * 32, 8192, 128  # flattened heads
kv = torch.randn(B, H, T, D)
kv_int8, scale = quantize_kv_int8(kv.view(B, 32, 32, T, D).view(B, H, T, D))

bf16_size = kv.numel() * 2  # bytes
int8_size  = kv_int8.numel() * 1 + scale.numel() * 2
print(f"BF16 KV size: {bf16_size / 1e9:.2f} GB")
print(f"INT8 KV size: {int8_size  / 1e9:.2f} GB  ({100*int8_size/bf16_size:.0f}% of BF16)")

INT4 KV cache (used in FlexGen for offloading) quantizes per-group-20 along the token dimension. Accuracy impact on generation quality is measurable on long-context tasks; for short contexts INT4 KV is essentially lossless. PagedAttention (discussed in PagedAttention & KV-Cache Memory Management) already manages KV memory in blocks; each block can independently carry a scale, making per-block quantization natural.

Accuracy–Performance Tradeoffs Across the Zoo¶

Choosing a quantization format involves balancing three axes: model quality (perplexity or downstream task score), inference memory, and inference throughput.

The Perplexity Cost Hierarchy¶

Roughly, the accuracy ordering from best to worst for a well-implemented scheme on a 7 B model is:

The gap between Q4_K_M and BF16 is typically 0.1–0.5 perplexity points on Wikitext-2 for a 7 B model — usually imperceptible in downstream quality. The gap grows for smaller models (3 B and below) and for tasks requiring precise factual recall.

Throughput and Latency¶

For a single-user, memory-bandwidth-bound decode scenario on an A100:

The large spread in INT4 throughput reflects kernel quality, not format: hand-tuned mixed-precision GEMMs (Marlin/Machete in vLLM, ExLlama’s kernels for single-user local inference) significantly outperform naive dequantize-then-FP16-GEMM. Note also that these numbers are for batch size 1. As batch size grows, the GEMM becomes compute-bound rather than bandwidth-bound and weight-only quantization’s advantage decays toward 1× — which is precisely when W8A8 (FP8 or INT8), whose win comes from arithmetic throughput, takes over.

Practical Decision Tree¶

Summary: Format Comparison Reference¶

Further Reading¶

• Dettmers et al., “LLM.int8(): 8-bit Matrix Multiplication for Transformers at Scale”, NeurIPS 2022 — the mixed-precision decomposition that made INT8 practical for very large models.
• Dettmers et al., “QLoRA: Efficient Finetuning of Quantized LLMs”, NeurIPS 2023 — introduces NF4, double quantization, paged optimizers, and the QLoRA recipe.
• Xiao et al., “SmoothQuant: Accurate and Efficient Post-Training Quantization for Large Language Models”, ICML 2023 — the per-channel migration trick that enables W8A8.
• Frantar et al., “GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers”, ICLR 2023 — the second-order OBC-based algorithm that produces the INT4 weights used by many GGUF conversions.
• Frantar et al., “MARLIN: Mixed-Precision Auto-Regressive Parallel Inference on Large Language Models” — the mixed-precision GEMM kernel that keeps W4A16 fast past batch size 1; the default INT4 path in vLLM, with Machete as its Hopper/CUTLASS-3 successor.
• bitsandbytes library (Tim Dettmers / Hugging Face) — github.com/bitsandbytes-foundation/bitsandbytes — the production Python/CUDA implementation of LLM.int8() and NF4.
• llama.cpp (Georgi Gerganov and contributors) — github.com/ggml-org/llama.cpp — the canonical k-quant and GGUF implementation; ggml-quants.c and the block_* structs in ggml-quants.h are the reference for Q4_K/Q5_K internals.
• Sheng et al., “FlexGen: High-Throughput Generative Inference of Large Language Models with a Single GPU”, ICML 2023 — demonstrates INT4 KV cache quantization and CPU offloading.
• NVIDIA Transformer Engine (github.com/NVIDIA/TransformerEngine) — reference implementation of FP8 training and inference on Hopper GPUs, including delayed scaling and amax history.

Exercises¶

1. Weight-only INT4 quantization does not change the arithmetic intensity of the GEMM — the kernel dequantizes each weight back to FP16 and runs the same FP16 multiply-accumulate as the BF16 baseline. Yet the throughput table lists INT4 W-only at roughly \(2.5\)–\(3.5\times\) the decode throughput of BF16. Using the roofline picture from the chapter, explain why decode gets faster even though the FLOP count and the compute precision are unchanged.

2. A group of weights is quantized with symmetric INT4 (signed, so the positive code limit is \(2^{b-1}-1 = 7\)) using the chapter’s per-group scale rule \(s = \max(|\mathbf{w}|) / (2^{b-1}-1)\). Suppose the group’s largest-magnitude weight is \(\max(|\mathbf{w}|) = 0.84\).

(a) Compute the scale \(s\). (b) Give the worst-case absolute quantization error for any weight in this group. © A single weight in the group has value \(w = 0.30\). What integer code \(q\) does it map to, and what is its dequantized value \(\hat{w}\) and error?

3. The chapter states that the Q4_K block layout costs exactly \(4.5\) bits/weight (a real Q4_K_M file averages a few tenths more, because the _M mix promotes attn_v and ffn_down to Q6_K), and gives the exact block_q4_K layout: two fp16 super-block scales d, dmin; a 12-byte array scales[12] holding the sixteen 6-bit sub-block scales and mins; and qs[128], the 256 packed 4-bit weights. Each super-block covers 256 weights. Derive the \(4.5\) bits/weight figure from this struct, and identify how many bits of that total are “pure payload” (the 4-bit weights) versus metadata overhead.

4. Consider the chapter’s KV-cache setting: a model with \(L = 32\) layers, \(H = 32\) heads, head dimension \(D = 128\), decoding a single sequence (\(B=1\)) out to \(T = 8192\) tokens. The cache stores both \(K\) and \(V\).

(a) Compute the BF16 KV-cache size in GB. (b) Now quantize to INT8 per-token, exactly as quantize_kv_int8 does: one INT8 byte per element, plus one fp16 scale per (layer, head, token) for each of \(K\) and \(V\). Compute the INT8 size and express it as a percentage of the BF16 size. (Use \(1\text{ GB} = 10^9\) bytes.)

5. The chapter’s FakeQuantLinear.forward contains a deliberately “simplified” bug: it computes a proper per-group scale with get_scale, then throws it away by calling STEQuantize.apply(self.weight, 1.0, self.bits) with a hard-coded scale of 1.0. Explain why passing 1.0 produces near-useless quantization for typical weights, then fix forward so it actually fake-quantizes on the per-group scale grid (still using the straight-through estimator). Include a short test that shows the fixed version has much lower weight-reconstruction MSE.

def forward(self, x: torch.Tensor) -> torch.Tensor:
    # Per-group symmetric scale, one value per group of `group_size` weights
    scale = self.get_scale(self.weight)            # same shape as weight
    # Fake-quantize on the correct grid; STE lets gradients pass through
    w_fq = STEQuantize.apply(self.weight, scale, self.bits)
    return F.linear(x, w_fq, self.bias)

import torch

torch.manual_seed(0)
layer = FakeQuantLinear(512, 512, bits=4, group_size=128)
# Give it realistic small-magnitude weights (std ~ 0.05)
with torch.no_grad():
    layer.weight.normal_(0.0, 0.05)

w = layer.weight.detach()

# Broken version: scale hard-coded to 1.0
w_broken = STEQuantize.apply(w, 1.0, 4)
mse_broken = ((w - w_broken) ** 2).mean().item()

# Fixed version: use the real per-group scale
scale = layer.get_scale(w)
w_fixed = STEQuantize.apply(w, scale, 4)
mse_fixed = ((w - w_fixed) ** 2).mean().item()

print(f"MSE with scale=1.0 (broken): {mse_broken:.6e}")
print(f"MSE with per-group scale:    {mse_fixed:.6e}")
print(f"Improvement factor:          {mse_broken / mse_fixed:.1f}x")

6. The Interview Corner argues that QLoRA does not need a straight-through estimator, even though its base weights are stored in 4-bit NF4 — yet the QAT section insists the STE is essential precisely because \(\operatorname{round}(\cdot)\) has zero gradient. Reconcile these two claims: under what condition is an STE required, and why does QLoRA escape it while plain INT4 QAT does not?