Chapter 05

Down to the metal

The GPU never multiplies by a 4-bit number. There is no 4-bit multiply instruction on an H100 tensor core. Here's what "4-bit inference" actually executes.

4-bit packed weights sit in HBM, two per byte.
A kernel loads them into shared memory as raw bytes.
Each thread unpacks nibbles into registers with shifts and masks.
The nibbles are dequantized to FP16/BF16 in registers.
The tensor core does an ordinary FP16 multiply with FP32 accumulation.

Steps 3 and 4 cost real arithmetic that FP16 inference doesn't pay. And you still win — because step 1 moved a quarter as many bytes, and step 1 was the entire bottleneck.

This is called W4A16. Weights in 4 bits, activations in 16. The compute is FP16; only the storage and memory traffic are quantized. It's the dominant regime for LLM inference precisely because decoding is bandwidth-bound and the dequantization cost hides in the memory shadow.

Dataflow — one output tile

animated · not to timing scale

The actual unpacking

Here is the inner loop, in the spirit of what ggml's CUDA and CPU kernels do for a Q4_0 block. Two 4-bit weights live in one byte:

// one byte holds two weights: hi nibble | lo nibble
const uint8_t b = block->qs[i];      // 0b1011_0010

const int q_lo = (b & 0x0F) - 8;    // 0b0010 = 2  →  -6
const int q_hi = (b >>  4)   - 8;    // 0b1011 = 11 →  +3

const float d = __half2float(block->d);
y[2*i+0] = d * q_lo;
y[2*i+1] = d * q_hi;

The −8 is the symmetric zero-point: nibbles are stored unsigned in [0,15] and shifted to [−8,7] on read. On a GPU this runs vectorised on 32-bit words (8 nibbles at a time) and is often fused with the __hfma2 that follows, so the dequantize costs a handful of ALU slots that would otherwise be idle.

Nibble unpack — step through it

tap a bit to change the byte

The memory hierarchy is why block size is 32

Level	Bandwidth	Latency	Relevance
Registers	~100 TB/s	~0 cyc	Where dequantization happens
Shared mem / L1	~20 TB/s	~30 cyc	Weight tiles staged here
L2 cache	~5 TB/s	~200 cyc	Rarely helps — weights stream once
HBM3	~3.35 TB/s	~500 cyc	The bottleneck. This is what quantization shrinks
NVLink	~900 GB/s	µs	Tensor-parallel cost you avoid by fitting on one GPU
PCIe 5 ×16	~64 GB/s	µs	Offload cliff — avoid at all costs
DDR5 (CPU)	~80 GB/s	ns	Why CPU inference is ~40× slower than an H100

Block sizes in real formats are not arbitrary — they line up with hardware:

32 weights is one CUDA warp — each of the 32 threads handles one weight and the block's scale broadcasts across the warp for free.
32 × 4 bits = 128 bits = exactly one ARM NEON register, or half an AVX-256 register. The whole block payload is one vector load.
256-weight superblocks in k-quants match a natural tile height and let the 8 sub-block scales load in one 128-bit read.
Group size 128 in GPTQ/AWQ aligns with the K dimension of the tensor-core MMA tile, so the dequantize hoists out of the innermost loop.

This is also why Unsloth's Dynamic 2.0 adds Q4_0/Q4_1/Q5_0/Q5_1/Q4_NL variants specifically for Apple Silicon and ARM: those chips lack the wide integer-shuffle instructions that make k-quant unpacking cheap on x86, and flat 32-weight blocks map directly onto NEON.

Where lookup tables come in

NF4 and the IQ codebooks can't be reconstructed with arithmetic — they need a table lookup. On CPU that's a vpshufb (x86) or tbl (ARM) instruction: a 16-entry table lives in one SIMD register and 16 lookups happen in a single instruction. That's why NF4 and IQ4_NL cap their codebooks at 16 entries. The constraint on codebook size is literally the width of a shuffle instruction.

// x86 AVX2: 32 parallel 4-bit lookups in one instruction
__m256i lut  = _mm256_broadcastsi128_si256(nf4_table);  // 16 entries
__m256i idx  = unpack_nibbles(packed);                  // 32 indices
__m256i vals = _mm256_shuffle_epi8(lut, idx);           // 1 cycle

The other thing quantization does to your GPU

Shrinking weights frees VRAM, and VRAM is where the KV cache lives. The KV cache grows linearly with context length and batch size, and at long context can exceed the weights entirely:

KV bytes = 2 × layers × kv_heads × head_dim × ctx × batch × bytes

// Llama-3-70B, 8k context, batch 1, FP16 KV:
// 2 × 80 × 8 × 128 × 8192 × 1 × 2  ≈  2.7 GB
// at 128k context: ≈ 43 GB — larger than a Q4 copy
// of the model itself.

So the real deployment win compounds: 4-bit weights free ~100 GB on a 70B model, which buys either much longer context or a much larger batch — and batching is what moves you rightward on the roofline into compute-bound territory where the GPU is actually efficient. Quantizing the KV cache itself to 8 or 4 bits is a separate and increasingly standard technique.