38. GGUF: running llama.cpp’s models

Why GGUF matters

Most people who run open models locally never touch a safetensors checkpoint. They download a single .gguf file, often a “Q4_K_M” quantization a quarter the size of the original, and run it with llama.cpp, Ollama or LM Studio. Hugging Face hosts tens of thousands of them, and many new models appear in GGUF within hours of their release. An engine that can’t read GGUF can’t run what most local users have on disk.

Chapter 20 explained quantization from first principles with one simple format of its own. GGUF is what quantization looks like after years of engineering for CPUs and consumer GPUs: a dozen block formats with different bit budgets, and a file format that bundles them with everything else a runtime needs. This chapter reads all of it.

The file

A GGUF file is one header followed by the tensor data:

"GGUF"  version (u32 = 3)  tensor count (u64)  metadata count (u64)
metadata  x count:   key (u64 length + UTF-8)  type (u32)  value
tensor info x count: name  n_dims (u32)  dims (u64 each, innermost first)  ggml type (u32)  offset (u64)
padding to general.alignment (32 by default)
tensor data, each tensor at its offset, aligned

Metadata values are typed: integers of every width, floats, booleans, strings, and arrays of any of them. That’s where a GGUF file keeps what a Hugging Face checkpoint spreads over config.json, tokenizer.json, tokenizer_config.json and generation_config.json:

class GGUFReader:
    """Parses the header and memory-maps the data; tensors are read lazily, one at a time."""

    def __init__(self, path):
        self.path = Path(path)
        self.file = open(self.path, "rb")
        self.data = mmap.mmap(self.file.fileno(), 0, access=mmap.ACCESS_READ)
        self.pos = 0
        if self.read(4) != MAGIC:
            raise ValueError(f"{path} is not a GGUF file")
        self.version = self.scalar("<I")
        if self.version not in (2, 3):
            raise ValueError(f"Unsupported GGUF version {self.version}")
        tensor_count, kv_count = self.scalar("<Q"), self.scalar("<Q")
        self.metadata = {}
        for _ in range(kv_count):
            key = self.string()
            self.metadata[key] = self.value(self.scalar("<I"))
        self.tensors = {}
        for _ in range(tensor_count):
            name = self.string()
            dims = [self.scalar("<Q") for _ in range(self.scalar("<I"))]
            kind, offset = self.scalar("<I"), self.scalar("<Q")
            self.tensors[name] = (list(reversed(dims)), ggml_quants.NAMES.get(kind, kind), offset)
        alignment = self.metadata.get("general.alignment", 32)
        self.data_start = -(-self.pos // alignment) * alignment

    def read(self, n):
        out = self.data[self.pos:self.pos + n]
        self.pos += n
        return out

    def scalar(self, fmt):
        (value,) = struct.unpack(fmt, self.read(struct.calcsize(fmt)))
        return value

    def string(self):
        return self.read(self.scalar("<Q")).decode("utf-8", errors="replace")

    def value(self, kind):
        """One typed metadata value.  (Your engine: Chapter 38)"""
        if kind in SCALARS:
            return self.scalar(SCALARS[kind])
        if kind == STRING:
            return self.string()
        if kind == ARRAY:
            element, count = self.scalar("<I"), self.scalar("<Q")
            if element in SCALARS and element != 7:            # numeric arrays: one bulk read
                dtype = np.dtype(SCALARS[element])
                values = np.frombuffer(self.read(count * dtype.itemsize), dtype=dtype)
                return values.tolist()
            return [self.value(element) for _ in range(count)]
        raise ValueError(f"Unknown GGUF value type {kind}")

    def raw(self, name):
        """The tensor's bytes, without copying (a view of the memory map)."""
        shape, kind, offset = self.tensors[name]
        weights, size = ggml_quants.BLOCK[kind]
        count = int(np.prod(shape))
        nbytes = count // weights * size
        start = self.data_start + offset
        with warnings.catch_warnings():                  # the map is read-only, and so is our use of it
            warnings.simplefilter("ignore")
            return torch.frombuffer(memoryview(self.data)[start:start + nbytes], dtype=torch.uint8)

    def tensor(self, name):
        """Dequantized float32 tensor in PyTorch's [rows, cols] order."""
        shape, kind, _ = self.tensors[name]
        return ggml_quants.dequantize(self.raw(name), kind, shape)

The reader parses the header and memory-maps the file. A tensor’s bytes are a view into the map, never copied until they’re converted, so opening a 40 GB file is instant and the operating system’s page cache serves repeated loads. Dimensions are stored innermost first (GGML’s ne[0] is the row length), so the reader reverses them into PyTorch’s [rows, cols].

Quantization types

Every quantized GGML type stores a tensor as a sequence of fixed-size blocks, each holding a few scales and the codes of 32 or 256 consecutive weights of a row. A row’s length must be a multiple of the block size, so a writer falls back to another type (usually Q8_0) for tensors whose rows don’t fit.

The legacy types

Q8_0  (34 bytes / 32 weights = 8.5 bits):  d:f16 | q:i8 x 32                 w = d * q
Q4_0  (18 bytes / 32 weights = 4.5 bits):  d:f16 | q:u4 x 32                 w = d * (q - 8)
Q4_1  (20 bytes / 32 weights = 5.0 bits):  d:f16 | m:f16 | q:u4 x 32         w = d * q + m
Q5_0, Q5_1: a fifth bit per weight in a 32-bit mask

Two 4-bit codes share each byte, with llama.cpp’s own order: byte $j$’s low nibble is weight $j$ and its high nibble is weight $j + 16$, not weight $j + 1$. Getting that wrong produces a perfectly plausible-looking tensor with its columns shuffled, the same trap as RoPE’s pairing in Chapter 17.

def dequant_q8_0(b):
    return f16(b[:, :2])[:, None] * b[:, 2:].view(torch.int8).float()


def dequant_q4_0(b):
    return f16(b[:, :2])[:, None] * (nibbles(b[:, 2:]).float() - 8)


def dequant_q4_1(b):
    return f16(b[:, :2])[:, None] * nibbles(b[:, 4:]).float() + f16(b[:, 2:4])[:, None]


def _high_bits(qh_bytes):
    """A little-endian u32 of 32 flags -> [n, 32] values 0/16: bit j belongs to weight j."""
    bits = (qh_bytes[:, :, None] >> torch.arange(8, dtype=torch.uint8)) & 1     # [n, 4, 8]
    return bits.reshape(-1, 32).float() * 16


def dequant_q5_0(b):
    return f16(b[:, :2])[:, None] * (nibbles(b[:, 6:]).float() + _high_bits(b[:, 2:6]) - 16)


def dequant_q5_1(b):
    return f16(b[:, :2])[:, None] * (nibbles(b[:, 8:]).float() + _high_bits(b[:, 4:8])) + f16(b[:, 2:4])[:, None]

K-quants

The K-quants (Q2_K to Q6_K, by Kawrakow, 2023) improved quality at a given size with a two-level scheme. A super-block of 256 weights is split into 8 sub-blocks of 32 (or 16 of 16). Each sub-block has its own scale and minimum, which adapt to its local range, and those scales are themselves quantized, to 6 bits in Q4_K, against one f16 d and dmin per super-block. Q4_K’s scale fields cost 128 bits per 256 weights, 0.5 bits per weight, exactly Q4_0’s overhead; for that price every 32-weight sub-block gets both a scale and a minimum, which Q4_1 needed 5 bits per weight to afford.

Q4_K’s 144-byte block:

d:f16  dmin:f16  scales:12 bytes (8 six-bit scales + 8 six-bit mins)  qs:128 bytes (256 four-bit codes)

weight = (d * scale[s]) * q  -  (dmin * min[s]),     s = sub-block 0..7

The twelve scale bytes pack sixteen 6-bit numbers. Sub-blocks 0-3 keep their 6 bits in bytes 0-3 (scales) and 4-7 (mins); sub-blocks 4-7 keep their low 4 bits in bytes 8-11 and their top 2 bits in the otherwise unused top bits of bytes 0-7:

def scale_min_k4(scales):
    """Q4_K / Q5_K: 12 bytes -> 8 six-bit scales and 8 six-bit mins.  (Your engine: Chapter 38)

    Sub-blocks 0-3 keep their low 6 bits in bytes 0-3 (scales) and 4-7 (mins); sub-blocks 4-7
    keep their low 4 bits in bytes 8-11 and their top 2 bits in the spare top bits of bytes 0-7.
    """
    s = scales.to(torch.int32)
    low_sc, low_m = s[:, 0:4] & 63, s[:, 4:8] & 63
    high_sc = (s[:, 8:12] & 0xF) | ((s[:, 0:4] >> 6) << 4)
    high_m = (s[:, 8:12] >> 4) | ((s[:, 4:8] >> 6) << 4)
    return torch.cat((low_sc, high_sc), 1).float(), torch.cat((low_m, high_m), 1).float()


def dequant_q4_k(b):
    """(Your engine: Chapter 38)"""
    d, dmin = f16(b[:, 0:2])[:, None], f16(b[:, 2:4])[:, None]
    sc, m = scale_min_k4(b[:, 4:16])
    q = b[:, 16:].reshape(-1, 4, 32)                       # 4 chunks of 32 bytes = 8 sub-blocks of 32
    codes = torch.stack((q & 0xF, q >> 4), dim=2).reshape(-1, 8, 32).float()   # chunk i: low -> sub 2i, high -> 2i+1
    return ((d * sc)[:, :, None] * codes - (dmin * m)[:, :, None]).reshape(-1, QK_K)


def dequant_q5_k(b):
    d, dmin = f16(b[:, 0:2])[:, None], f16(b[:, 2:4])[:, None]
    sc, m = scale_min_k4(b[:, 4:16])
    qh = b[:, 16:48]                                       # bit 2i / 2i+1 of qh[l]: the 5th bit for subs 2i, 2i+1
    q = b[:, 48:].reshape(-1, 4, 32)
    low = torch.stack((q & 0xF, q >> 4), dim=2).reshape(-1, 8, 32).to(torch.int32)
    bit = (qh[:, None, :].to(torch.int32) >> torch.arange(8)[None, :, None]) & 1     # [n, 8, 32]
    codes = (low + 16 * bit).float()
    return ((d * sc)[:, :, None] * codes - (dmin * m)[:, :, None]).reshape(-1, QK_K)


def q6_k_codes(b):
    """Q6_K's 256 six-bit codes (0..63) in weight order: 4 low bits from ql, 2 high bits from qh."""
    ql, qh = b[:, :128].to(torch.int32), b[:, 128:192].to(torch.int32)
    out = []
    for half in range(2):                                  # two halves of 128 weights
        l_lo, l_hi = ql[:, 64 * half:64 * half + 32], ql[:, 64 * half + 32:64 * half + 64]
        h = qh[:, 32 * half:32 * half + 32]
        out += [(l_lo & 0xF) | (((h >> 0) & 3) << 4), (l_hi & 0xF) | (((h >> 2) & 3) << 4),
                (l_lo >> 4) | (((h >> 4) & 3) << 4), (l_hi >> 4) | (((h >> 6) & 3) << 4)]
    return torch.cat(out, 1)


def dequant_q6_k(b):
    """(Your engine: Chapter 38)"""
    sc, d = b[:, 192:208].view(torch.int8).float(), f16(b[:, 208:210])[:, None]
    return d * sc.repeat_interleave(16, 1) * (q6_k_codes(b).float() - 32)     # one scale per 16 weights


def dequant_q2_k(b):
    sc, q = b[:, :16].to(torch.int32), b[:, 16:80].to(torch.int32)
    d, dmin = f16(b[:, 80:82])[:, None], f16(b[:, 82:84])[:, None]
    out = []
    for half in range(2):
        chunk = q[:, 32 * half:32 * half + 32]
        for shift in range(4):
            codes = ((chunk >> (2 * shift)) & 3).float()                     # 32 codes: two sub-blocks of 16
            for part in range(2):
                s = sc[:, 8 * half + 2 * shift + part]
                out.append(d * (s & 0xF)[:, None].float() * codes[:, 16 * part:16 * part + 16]
                           - dmin * (s >> 4)[:, None].float())
    return torch.cat(out, 1)


def dequant_q3_k(b):
    hmask, q = b[:, :32].to(torch.int32), b[:, 32:96].to(torch.int32)
    raw, d = b[:, 96:108].to(torch.int32), f16(b[:, 108:110])[:, None]
    # 16 six-bit scales: low 4 bits in bytes 0-7 (two per byte), top 2 bits in bytes 8-11.
    low = torch.cat((raw[:, 0:8] & 0xF, raw[:, 0:8] >> 4), 1)
    high = torch.cat([(raw[:, 8:12] >> (2 * i)) & 3 for i in range(4)], 1)
    scales = (low | (high << 4)).float() - 32
    out, bit = [], 0
    for half in range(2):
        chunk = q[:, 32 * half:32 * half + 32]
        for shift in range(4):
            codes = ((chunk >> (2 * shift)) & 3) - 4 * (1 - ((hmask >> bit) & 1))
            for part in range(2):
                out.append(d * scales[:, 8 * half + 2 * shift + part][:, None] * codes[:, 16 * part:16 * part + 16].float())
            bit += 1
    return torch.cat(out, 1)

Q6_K, which the popular “_M” mixes use for the most sensitive tensors (the output head, some attention and FFN-down layers), splits each 6-bit code into 4 low bits and 2 high bits stored in separate arrays, with an 8-bit signed scale per 16 weights. Q2_K and Q3_K go further down, with 2-bit codes (Q3_K adds a third bit from a mask) and 4-bit or 6-bit sub-block scales.

A name like Q4_K_M isn’t a tensor type. It’s a recipe: mostly Q4_K, with Q6_K for the tensors that hurt most when quantized. A GGUF file can mix types freely, tensor by tensor.

Nothing in this code is checked against prose or memory. The test uses llama.cpp’s own Python package, gguf, as the reference: for each legacy type, gguf.quants.quantize produces blocks, and our decoder must reproduce gguf.quants.dequantize bit for bit; for the K-quants, random blocks (with sane f16 scale fields) must decode identically. Our files must open in llama.cpp’s reader, and its files in ours.

Beyond the K-quants

llama.cpp keeps adding formats. IQ types (“importance” quants, IQ1_S to IQ4_XS) use non-uniform codebooks and lattice codes chosen with an importance matrix from calibration data; IQ4_NL’s 16 values (included here) are a non-uniform table that fits the bell shape of weight distributions better than a uniform grid. MXFP4 is the OCP microscaling format that OpenAI’s gpt-oss models ship in: 32 FP4 (E2M1) values sharing a power-of-two scale (Chapter 39). Each is another block decoder of the same shape.

The tokenizer is in the file

tokenizer.ggml.model = "gpt2" means byte-level BPE (Chapter 35), with the vocabulary in tokens, the merge ranks in merges and each token’s kind in token_type (3 = control, which is special; 4 = user-defined, which is added but ordinary text). One thing is missing: the pre-tokenizer’s regex, which GGUF doesn’t store. Instead, tokenizer.ggml.pre names it ("qwen2", "llama-bpe", …), and every runtime keeps a table from names to patterns, which is why a llama.cpp release is needed for each new model family:

def tokenizer_from_gguf(metadata):
    """Rebuild a byte-level BPE tokenizer.json spec from GGUF metadata.  (Your engine: Chapter 38)

    token_type 3 (control) tokens are special; type 4 (user-defined) tokens are added but
    ordinary text, like <think>.
    """
    from ..serve.tokenizer import Tokenizer
    model = metadata.get("tokenizer.ggml.model")
    if model != "gpt2":
        raise ValueError(f"Only byte-level BPE ('gpt2') GGUF tokenizers are implemented, not {model!r}")
    pre = metadata.get("tokenizer.ggml.pre", "default")
    if pre not in PRE_TOKENIZERS:
        raise ValueError(f"Unknown pre-tokenizer {pre!r}; add its regex to PRE_TOKENIZERS")
    tokens = metadata["tokenizer.ggml.tokens"]
    kinds = metadata.get("tokenizer.ggml.token_type", [1] * len(tokens))
    added = [{"id": i, "content": t, "special": k == 3} for i, (t, k) in enumerate(zip(tokens, kinds)) if k in (3, 4)]
    spec = {"model": {"type": "BPE", "vocab": {t: i for i, t in enumerate(tokens)},
                      "merges": metadata.get("tokenizer.ggml.merges", [])},
            "pre_tokenizer": {"type": "Sequence", "pretokenizers": [
                {"type": "Split", "pattern": {"Regex": PRE_TOKENIZERS[pre]}, "behavior": "Isolated"},
                {"type": "ByteLevel", "add_prefix_space": False, "use_regex": False}]},
            "normalizer": {"type": "NFC"} if pre == "qwen2" else None,
            "decoder": {"type": "ByteLevel"}, "added_tokens": added}
    eos = metadata.get("tokenizer.ggml.eos_token_id")
    config = {"chat_template": metadata.get("tokenizer.chat_template"),
              "eos_token": tokens[eos] if eos is not None else None}
    return Tokenizer(spec, config)

The result is the same Tokenizer class as Chapter 35’s, so the server, the detokenizer and constrained decoding all work unchanged on a GGUF model.

Loading the model

GGUF tensor names are short and architecture-neutral: token_embd, blk.{i}.attn_q, blk.{i}.ffn_down, output_norm, output. Loading Qwen3 is a name table plus the hyperparameters from metadata:

@torch.no_grad()
def load_gguf(path, device="cpu", dtype=torch.bfloat16, keep_quantized=True):
    """A Qwen3 or Qwen3-MoE model from a GGUF file.  (Your engine: Chapter 38)

    Quantized 2-D weights become QuantLinear-style modules that keep their codes (repacked once
    into the engine's group-affine layout, formats/runtime.py) unless keep_quantized=False.
    Returns (model, tokenizer).
    """
    from ..qwen3 import Qwen3, Qwen3Config
    from ..moe import Qwen3Moe, Qwen3MoeConfig
    from .runtime import AffineQuantLinear, repack_ggml
    reader = GGUFReader(path)
    raw = config_from_gguf(reader.metadata)
    arch = raw["model_type"]
    if arch not in ("qwen3", "qwen3_moe", "llama", "qwen2"):
        raise ValueError(f"GGUF architecture {arch!r} is not supported")
    names = set(reader.tensors)
    raw["tie_word_embeddings"] = "output.weight" not in names
    if arch in ("llama", "qwen2"):                                 # Chapter 42's registry
        from ..models import build_model
        with torch.device("meta"):
            model = build_model(raw)
        cfg = model.cfg
    else:
        cfg = (Qwen3MoeConfig if arch == "qwen3_moe" else Qwen3Config).from_hf(raw)
        with torch.device("meta"):
            model = (Qwen3Moe if arch == "qwen3_moe" else Qwen3)(cfg)
    model = model.to_empty(device=device).to(dtype)
    if cfg.tie_word_embeddings:
        model.lm_head.weight = model.model.embed_tokens.weight
    params = dict(model.named_parameters())

    def put(target, name, heads=None):
        """heads: undo llama.cpp's Q/K row permutation (Llama-architecture files only)."""
        shape, kind, _ = reader.tensors[name]
        module_name, _, attr = target.rpartition(".")
        module = model.get_submodule(module_name)
        rows = None if heads is None else unpermute_qk(torch.arange(shape[0])[:, None], heads)[:, 0]
        if keep_quantized and kind not in ("F32", "F16", "BF16") and isinstance(module, torch.nn.Linear) and attr == "weight":
            codes, scales, offsets, group, packed = repack_ggml(reader.raw(name), kind, shape)
            if rows is not None:                                   # quantization is per row: permuting rows is exact
                codes, scales, offsets = codes[rows], scales[rows], offsets[rows]
            parent, _, child = module_name.rpartition(".")
            layer = AffineQuantLinear.from_packed(codes, scales, offsets, group, packed, dtype=dtype)
            if module.bias is not None:
                layer.bias = module.bias
            setattr(model.get_submodule(parent), child, layer.to(device))
            return
        value = reader.tensor(name).to(device=device, dtype=dtype).reshape(params[target].shape)
        params[target].copy_(value if rows is None else value[rows.to(device)])

    for name in names:
        if name == "token_embd.weight":
            put("model.embed_tokens.weight", name)
        elif name == "output_norm.weight":
            put("model.norm.weight", name)
        elif name == "output.weight":
            put("lm_head.weight", name)
        elif name == "rope_freqs.weight":                          # Llama 3's frequency divisors
            model.inv_freq = model.inv_freq / reader.tensor(name).float()
        elif name.startswith("blk."):
            _, layer, kind, part = name.split(".", 3)
            prefix = f"model.layers.{layer}."
            if kind.endswith("_exps"):                     # MoE experts, stacked [E, out, in]
                experts = model.model.layers[int(layer)].mlp.experts
                tensor = reader.tensor(name).to(device=device, dtype=dtype)
                if kind == "ffn_down_exps":
                    experts.down_proj.copy_(tensor)
                else:
                    inter = experts.gate_up_proj.shape[1] // 2
                    part = slice(0, inter) if kind == "ffn_gate_exps" else slice(inter, 2 * inter)
                    experts.gate_up_proj[:, part].copy_(tensor)
            else:
                heads = {"attn_q": cfg.num_attention_heads, "attn_k": cfg.num_key_value_heads}.get(kind)
                put(prefix + GGUF_NAMES[kind] + "." + part, name, heads if arch == "llama" and part == "weight" else None)
        else:
            raise ValueError(f"Unexpected GGUF tensor {name}")
    return model.eval(), tokenizer_from_gguf(reader.metadata)

MoE layers store their experts stacked: ffn_gate_exps is one 3-D tensor of shape [E, I, D], the same layout Chapter 27 chose for grouped kernels.

One more trap waits in Llama-family GGUFs. llama.cpp’s converter permutes the rows of Llama’s Q and K projections, turning Hugging Face’s rotate-half RoPE pairs $(i, i + d/2)$ into llama.cpp’s interleaved pairs $(2i, 2i+1)$. A Llama GGUF loaded into a rotate-half model must undo it (unpermute_qk; Chapter 42’s registry uses it). Qwen models use rotate-half RoPE in llama.cpp too, so their weights aren’t permuted.

Serving quantized weights

Dequantizing every weight to BF16 at load time works, but throws away the reason for quantizing: a Q4_K model would occupy 3.5 times its file size in GPU memory, and decode, being memory-bound, would read 3.5 times more bytes per token. The weights must stay quantized, and the matmul must decode them on the fly, like Chapter 20’s W4A16 kernel.

Writing one kernel per GGML type (llama.cpp has dozens) is a large job. Look at what the types decode to, though:

Every one is code × scale + offset per group. So are GPTQ and AWQ (Chapter 39). The engine therefore has one quantized layer, AffineQuantLinear, and each format is a function that repacks its blocks into that layout once, at load time:

def repack_ggml(raw, kind, shape):
    """GGML blocks -> (codes, scales, offsets, group size, packed). Exactly the same weights, in
    the engine's layout.  (Your engine: Chapter 38)"""
    rows, cols = shape
    weights, size = ggml_quants.BLOCK[kind]
    b = raw.reshape(-1, size)
    f16 = ggml_quants.f16
    if kind == "Q8_0":
        codes = b[:, 2:].view(torch.int8).reshape(rows, cols)
        scales = f16(b[:, :2]).reshape(rows, -1)
        return codes.clone(), scales, torch.zeros_like(scales), 32, False
    if kind == "Q4_0":
        d = f16(b[:, :2]).reshape(rows, -1)
        codes = ggml_quants.nibbles(b[:, 2:]).reshape(rows, cols)
        return pack_nibbles(codes), d, -8 * d, 32, True
    if kind == "Q4_1":
        codes = ggml_quants.nibbles(b[:, 4:]).reshape(rows, cols)
        return pack_nibbles(codes), f16(b[:, :2]).reshape(rows, -1), f16(b[:, 2:4]).reshape(rows, -1), 32, True
    if kind in ("Q4_K", "Q5_K"):
        d, dmin = f16(b[:, 0:2])[:, None], f16(b[:, 2:4])[:, None]
        sc, m = ggml_quants.scale_min_k4(b[:, 4:16])
        scales, offsets = (d * sc).reshape(rows, -1), (-dmin * m).reshape(rows, -1)
        if kind == "Q4_K":
            q = b[:, 16:].reshape(-1, 4, 32)
            codes = torch.stack((q & 0xF, q >> 4), dim=2).reshape(rows, cols)
            return pack_nibbles(codes), scales, offsets, 32, True
        qh = b[:, 16:48]
        q = b[:, 48:].reshape(-1, 4, 32)
        low = torch.stack((q & 0xF, q >> 4), dim=2).reshape(-1, 8, 32).to(torch.int32)
        bit = (qh[:, None, :].to(torch.int32) >> torch.arange(8)[None, :, None]) & 1
        return (low + 16 * bit).to(torch.int8).reshape(rows, cols), scales, offsets, 32, False
    if kind == "Q6_K":                                    # w = d * sc * (q - 32), one scale per 16
        scales = (f16(b[:, 208:210])[:, None] * b[:, 192:208].view(torch.int8).float()).reshape(rows, -1)
        codes = ggml_quants.q6_k_codes(b).to(torch.int8).reshape(rows, cols)
        return codes, scales, -32 * scales, 16, False
    raise NotImplementedError(f"No repacking for {kind}; load it dense (keep_quantized=False)")

class AffineQuantLinear(nn.Module):
    """A linear layer whose weight stays as codes + per-group scale and offset."""

    def __init__(self, codes, scales, offsets, group_size, packed, perm=None, bias=None, dtype=torch.bfloat16):
        super().__init__()
        self.register_buffer("codes", codes)
        self.register_buffer("scales", scales.float())
        self.register_buffer("offsets", offsets.float())
        self.register_buffer("perm", perm)              # GPTQ act-order: input columns grouped by g_idx
        self.bias = None if bias is None else nn.Parameter(bias.to(dtype), requires_grad=False)
        self.group_size, self.packed, self.compute_dtype = group_size, packed, dtype
        self.out_features = codes.shape[0]
        self.in_features = codes.shape[1] * (2 if packed else 1)

    @classmethod
    def from_packed(cls, codes, scales, offsets, group_size, packed, dtype=torch.bfloat16):
        return cls(codes, scales, offsets, group_size, packed, dtype=dtype)

    @property
    def weight_codes(self):
        return unpack_nibbles(self.codes) if self.packed else self.codes

    def dequantized_weight(self):
        g = self.group_size
        return (self.weight_codes.float() * self.scales.repeat_interleave(g, 1)[:, :self.in_features]
                + self.offsets.repeat_interleave(g, 1)[:, :self.in_features])

    def forward(self, x):
        shape = x.shape
        x2 = x.reshape(-1, shape[-1])
        if self.perm is not None:
            x2 = x2[:, self.perm]
        if self.use_kernel(x2):
            from ..kernels.triton_formats import affine_matmul
            y = affine_matmul(x2, self.codes, self.scales, self.offsets, self.group_size, self.out_features, self.packed)
        elif self.use_cpu_kernel(x2):
            from ..kernels import cpu
            y = cpu.affine_gemm(x2, self.codes, self.scales, self.offsets, self.group_size, self.packed).to(x.dtype)
        else:
            y = F.linear(x2, self.dequantized_weight().to(x.dtype))
        if self.bias is not None:
            y = y + self.bias
        return y.reshape(*shape[:-1], self.out_features)

    def use_kernel(self, x):
        from ..kernels import interpreting
        return x.is_cuda or (interpreting() and getattr(self, "force_kernel", False))

    def use_cpu_kernel(self, x):
        """Decode-sized batches on the CPU, when the SIMD extension is built (Chapter 40)."""
        if not CPU_KERNELS["enabled"] or x.device.type != "cpu" or x.shape[0] > 16:
            return False
        if self.in_features % 32 or self.group_size % 32:
            return False
        from ..kernels import cpu
        return cpu.available()

    @property
    def storage_bytes(self):
        return sum(t.numel() * t.element_size() for t in (self.codes, self.scales, self.offsets))

and one Triton kernel that dequantizes tiles in registers, the same structure as Chapter 20’s W4A16 kernel with an offset and an arbitrary group size:

@triton.jit
def affine_matmul_kernel(x_ptr, codes_ptr, scales_ptr, offsets_ptr, y_ptr, M, N, K, groups, G,
                         PACKED: tl.constexpr, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr):
    """(Your engine: Chapter 38)"""
    pid_m, pid_n = tl.program_id(0), tl.program_id(1)
    rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
    for k0 in range(0, K, BLOCK_K):
        rk = k0 + tl.arange(0, BLOCK_K)
        k_ok = rk < K
        x = tl.load(x_ptr + rm[:, None] * K + rk[None, :], mask=(rm[:, None] < M) & k_ok[None, :], other=0.0)
        w_ok = (rn[:, None] < N) & k_ok[None, :]
        if PACKED:                                   # byte k // 2 holds code k in its low (even k) or high nibble
            byte = tl.load(codes_ptr + rn[:, None] * (K // 2) + rk[None, :] // 2, mask=w_ok, other=0).to(tl.int32)
            code = (byte >> ((rk[None, :] % 2) * 4)) & 0xF
        else:
            code = tl.load(codes_ptr + rn[:, None] * K + rk[None, :], mask=w_ok, other=0).to(tl.int32)
        group = rn[:, None] * groups + rk[None, :] // G
        scale = tl.load(scales_ptr + group, mask=w_ok, other=0.0)
        offset = tl.load(offsets_ptr + group, mask=w_ok, other=0.0)
        w = code.to(tl.float32) * scale + offset                      # [BLOCK_N, BLOCK_K], in registers
        acc += tl.dot(x.to(tl.float32), tl.trans(w), input_precision="ieee")
    tl.store(y_ptr + rm[:, None] * N + rn[None, :], acc.to(y_ptr.dtype.element_ty),
             mask=(rm[:, None] < M) & (rn[None, :] < N))


def affine_matmul(x, codes, scales, offsets, group_size, out_features, packed, block_m=16, block_n=32, block_k=32):
    check_device(x, codes, scales, offsets)
    m, k = x.shape
    y = torch.empty((m, out_features), device=x.device, dtype=x.dtype)
    grid = (triton.cdiv(m, block_m), triton.cdiv(out_features, block_n))
    affine_matmul_kernel[grid](x.contiguous(), codes, scales, offsets, y, m, out_features, k, scales.shape[1], group_size,
                               PACKED=packed, BLOCK_M=block_m, BLOCK_N=block_n, BLOCK_K=block_k)
    return y

The repacking is exact. The test checks that the repacked layer’s weights equal llama.cpp’s dequantization for Q8_0, Q4_0, Q4_1, Q4_K, Q5_K and Q6_K, and that the kernel’s output matches a dense matmul. Exactness has a price in memory: the per-group scales and offsets are stored as FP32 (Q4_K’s d·sc products don’t round-trip through FP16), so Q4_K takes 6 bits per weight resident instead of 4.5, and Q8_0 takes 10 instead of 8.5. A kernel that decodes Q4_K’s super-blocks directly, reading the 12 scale bytes as llama.cpp’s CUDA kernels do, recovers the difference (stretch exercise 2). It’s the right next step for a production engine, and the layout above is the right first one, because it makes every format run, correctly, with one kernel.

Writing GGUF

The writer is the reader in reverse, and with it the engine can export: quantize a model it has loaded, fine-tuned (Chapter 22) or edited (Chapter 23), and hand the file to llama.cpp, Ollama or LM Studio.

def fallback_type(kind, row_length):
    """Rows must hold whole blocks: like llama.cpp, fall back to Q8_0 (or F16) when they can't."""
    weights = ggml_quants.BLOCK[kind][0]
    if row_length % weights == 0:
        return kind
    return "Q8_0" if row_length % 32 == 0 else "F16"


def write_gguf(path, metadata, tensors, alignment=32):
    """metadata: {key: python value}; tensors: {name: (float tensor [rows, cols], ggml type name)}."""
    def pack_string(s):
        data = s.encode("utf-8")
        return struct.pack("<Q", len(data)) + data

    def pack_value(v):
        if isinstance(v, bool):
            return struct.pack("<I?", 7, v)
        if isinstance(v, int):
            return struct.pack("<Iq", 11, v) if v < 0 else struct.pack("<IQ", 10, v) if v >= 2 ** 32 else struct.pack("<II", 4, v)
        if isinstance(v, float):
            return struct.pack("<If", 6, v)
        if isinstance(v, str):
            return struct.pack("<I", STRING) + pack_string(v)
        if isinstance(v, (list, tuple)):
            if all(isinstance(x, str) for x in v):
                return struct.pack("<IIQ", ARRAY, STRING, len(v)) + b"".join(pack_string(x) for x in v)
            if all(isinstance(x, float) for x in v):
                return struct.pack("<IIQ", ARRAY, 6, len(v)) + np.asarray(v, dtype="<f4").tobytes()
            return struct.pack("<IIQ", ARRAY, 5, len(v)) + np.asarray(v, dtype="<i4").tobytes()
        raise TypeError(f"Can't store {type(v)} in GGUF metadata")

    metadata = {"general.alignment": alignment, **metadata}
    blobs, infos, offset = [], [], 0
    for name, (tensor, kind) in tensors.items():
        kind = fallback_type(kind, tensor.shape[-1])
        data = ggml_quants.quantize(tensor, kind).numpy().tobytes()
        dims = list(reversed(tensor.shape))
        infos.append(pack_string(name) + struct.pack("<I", len(dims)) + struct.pack(f"<{len(dims)}Q", *dims)
                     + struct.pack("<IQ", ggml_quants.TYPES[kind], offset))
        padded = data + b"\0" * (-len(data) % alignment)
        blobs.append(padded)
        offset += len(padded)
    header = MAGIC + struct.pack("<IQQ", 3, len(tensors), len(metadata))
    header += b"".join(pack_string(k) + pack_value(v) for k, v in metadata.items()) + b"".join(infos)
    header += b"\0" * (-len(header) % alignment)
    Path(path).write_bytes(header + b"".join(blobs))

def quantize_q8_0(w):
    blocks = w.float().reshape(-1, 32)
    d = blocks.abs().amax(1) / 127
    q = torch.where(d[:, None] > 0, blocks / d[:, None].clamp_min(1e-30), torch.zeros_like(blocks)).round().clamp(-127, 127)
    return torch.cat((_f16_bytes(d), q.to(torch.int8).view(torch.uint8)), 1)


def quantize_q4_k(w):
    """A straightforward Q4_K encoder: per 32-weight sub-block, min/max -> scale and min; the
    8 scales and 8 mins are then quantized to 6 bits against the super-block's d and dmin.
    (llama.cpp's encoder also searches each sub-block's range to minimize error.)"""
    x = w.float().reshape(-1, 8, 32)
    lo = x.amin(-1).clamp(max=0)                           # Q4_K stores w = scale*q - min, min >= 0
    hi = x.amax(-1)
    sub_scale = (hi - lo) / 15
    sub_min = -lo
    d = sub_scale.amax(1) / 63
    dmin = sub_min.amax(1) / 63
    sc = torch.where(d[:, None] > 0, sub_scale / d[:, None].clamp_min(1e-30), torch.zeros_like(sub_scale)).round().clamp(0, 63)
    m = torch.where(dmin[:, None] > 0, sub_min / dmin[:, None].clamp_min(1e-30), torch.zeros_like(sub_min)).round().clamp(0, 63)
    d_eff = (d.to(torch.float16).float()[:, None] * sc)
    m_eff = (dmin.to(torch.float16).float()[:, None] * m)
    q = torch.where(d_eff[:, :, None] > 0, (x + m_eff[:, :, None]) / d_eff[:, :, None].clamp_min(1e-30),
                    torch.zeros_like(x)).round().clamp(0, 15).to(torch.uint8)
    sc, m = sc.to(torch.uint8), m.to(torch.uint8)
    packed = torch.zeros(x.shape[0], 12, dtype=torch.uint8)
    packed[:, 0:4] = (sc[:, 0:4] & 63) | ((sc[:, 4:8] >> 4) << 6)
    packed[:, 4:8] = (m[:, 0:4] & 63) | ((m[:, 4:8] >> 4) << 6)
    packed[:, 8:12] = (sc[:, 4:8] & 0xF) | ((m[:, 4:8] & 0xF) << 4)
    q = q.reshape(-1, 4, 2, 32)
    qs = (q[:, :, 0] | (q[:, :, 1] << 4)).reshape(-1, 128)
    return torch.cat((_f16_bytes(d), _f16_bytes(dmin), packed, qs), 1)


QUANTIZE = {"Q8_0": quantize_q8_0, "Q4_K": quantize_q4_k}


def quantize(w, type_name):
    """float tensor -> raw GGML bytes (rows must be multiples of the block size)."""
    if type_name == "F32":
        return w.float().contiguous().view(torch.uint8).reshape(-1)
    if type_name in ("F16", "BF16"):
        return w.to(torch.float16 if type_name == "F16" else torch.bfloat16).contiguous().view(torch.uint8).reshape(-1)
    return QUANTIZE[type_name](w).reshape(-1)

The Q4_K encoder here is the straightforward one: each sub-block’s range from its min and max, then the 6-bit quantization of scales and mins. llama.cpp’s encoder also searches each sub-block’s range to minimize squared error (and, with an importance matrix, weighted error), which buys a few percent of quality. The test checks that llama.cpp’s decoder reads our blocks exactly as ours does.

Run it

python run.py gguf                                   # export the test model and reload it
python run.py gguf --model-dir model.Q4_K_M.gguf     # or run a real GGUF file

{"type": "Q8_0", "file_bits_per_param": 8.54, "resident_bits_per_quantized_weight": 10.0, "relative_logit_error": 0.0071, "same_greedy_next_token": true}
{"type": "Q4_K", "file_bits_per_param": 4.67, "resident_bits_per_quantized_weight": 6.0, "relative_logit_error": 0.0697, "same_greedy_next_token": true}
{"served_from_q4_k": [220, 296, 51, 500, 59, 328, 150, 402]}

The 4-layer test model (random weights, rows of 256) exported as Q8_0 is 8.54 bits per parameter on disk and changes its logits by 0.7%; as Q4_K, 4.67 bits and 7% (random weights have no structure for quantization to exploit, so real models fare better). The Q4_K model is then served by the engine core of Chapter 31, its linear layers running from codes through AffineQuantLinear.

Build it

pytest tests/test_ch38_gguf.py
python run.py gguf --impl engine

Stretch exercises

1. ★ Load a real Qwen3 GGUF (for example Qwen3-0.6B Q8_0 and Q4_K_M) with run.py gguf --model-dir, and compare its next-token probabilities on a few prompts with the BF16 safetensors checkpoint’s: KL divergence per token, by quantization type. Where: experiments/ch38.py (create it), using engine.formats.gguf.load_gguf and engine.evaluation.logit_quality.
2. ★★★ Write a Triton kernel that reads Q4_K super-blocks directly: per (row, super-block), load d, dmin and the 12 scale bytes, unpack the 8 scales and mins in registers, and decode codes from qs. Compare resident memory and decode bandwidth with AffineQuantLinear. Where: add a direct Q4_K kernel/launcher in engine/kernels/triton_formats.py; dispatch it from engine/formats/runtime.py.
3. ★★ Add the SentencePiece-style GGUF tokenizer (tokenizer.ggml.model = "llama"): merges come from token scores, highest first, with byte fallback. Where: tokenizer_from_gguf in engine/formats/gguf.py, with tokenization support in engine/serve/tokenizer.py.
4. ★★ Implement llama.cpp’s importance-weighted Q4_K encoder: given per-column activation statistics from calibration data, choose each sub-block’s scale and min to minimize the importance-weighted squared error. Measure KL against the plain encoder on a real model. Where: quantize_q4_k in engine/formats/ggml_quants.py.

Check your understanding

1. What does a GGUF file contain that a model.safetensors file doesn’t?
2. Why does the reader memory-map the file rather than read it?
3. In Q4_0, which weights share a byte? What would a decoder that paired weights $j$ and $j+1$ produce?
4. Why do K-quants quantize their sub-block scales, and what does that buy over Q4_0’s one f16 scale per 32 weights?
5. What is Q4_K_M, as opposed to Q4_K?
6. Why can one kernel serve GGUF, GPTQ and AWQ weights, and what does the shared layout cost for Q4_K?
7. Why must a Llama GGUF’s Q and K rows be un-permuted for a rotate-half model, but a Qwen GGUF’s not?

Going deeper

• The GGUF specification (docs/gguf.md in the ggml repository) and llama.cpp’s gguf-py package, used as this chapter’s reference.
• llama.cpp’s ggml/src/ggml-quants.c (dequantize_row_q4_K and friends) and ggml-common.h (the block structs); its CUDA dequantization kernels in ggml/src/ggml-cuda/.
• Kawrakow’s pull requests introducing the K-quants (llama.cpp #1684, 2023) and the IQ types, with their perplexity tables.
• convert_hf_to_gguf.py, for how each architecture’s tensors and metadata are named.