KV-cache kernels

01 · Beyond smaller numbers

Quantization moves day 11's ceiling

Storing weights in INT8 or FP8 instead of FP32 quarters the bytes — and through day 11's lens, that's not a storage footnote, it's a different machine: memory-bound kernels get a ~4× higher ceiling, because the thing that was limiting them — bytes across DRAM — shrank 4×. Most quantization wins in practice are bandwidth wins, which is why they show up most dramatically in memory-bound paths: weight streaming, KV-cache reads (day 26), elementwise chains.

But a quantized kernel is not the FP32 kernel with a different dtype. Three things change structurally: values must be dequantized on the fly, accumulation switches to wider integer or float types, and somebody has to decide the scales — the per-group conversion factors that make low-bit storage represent real weight distributions.

02 · Scales and dequantization

Store small, compute big, convert in registers

The universal pattern: quantized data travels with scale factors — per tensor, per output channel, or per small group of channels — and the kernel widens values into compute precision at load time:

# per-group scale, dequantized in registers after the load
w_q = tl.load(w_ptr + offs)              # int8 values
s   = tl.load(scale_ptr + offs // GROUP) # one scale per group
w   = w_q.to(tl.float16) * s             # dequant: cheap, in-register

# int8 GEMMs: products go to int32, requantize at the end
acc = tl.dot(a_q, b_q)                   # accumulates int32, not int8
out = (acc.to(tl.float16) * a_scale * b_scale)

Two accuracy details separate working kernels from broken ones. First, accumulation never happens in the narrow format — int8 products accumulate in int32 (products of 8-bit values need up to 16 bits immediately, and sums grow from there); FP16 math typically accumulates in FP32. Second, the two dominant FP8 encodings trade differently: e4m3 spends bits on mantissa (precision) while e5m2 spends them on exponent (range) — gradients prefer the range, weights prefer the precision. "FP8" is a family, not a format.

03 · Where scales come from

Calibration, group sizes, and outlier channels

Scales are chosen before the kernel ever runs: run representative data through the model, record each tensor's range (min–max, or a percentile to resist outliers), and set the scale to map that range onto the format's representable span. The granularity is the accuracy lever: per-group scales (every 64–128 channels) track local structure and preserve quality, at the cost of extra scale traffic — the kernel loads one small scale vector per group, which day 10's L2 happily caches. Coarser scales are cheaper and lossier. The failure mode has a name in every quantization writeup: outlier channels — a handful of activations with huge magnitudes that a global scale would flatten everyone else to accommodate for. Per-group and per-channel scaling exist largely to quarantine them.

04 · What changes in the kernel's shape

The skeleton survives; the plumbing changes

Strip a quantized kernel to its structure and you'll recognize it instantly: same tile mapping (days 14–15), same reduction discipline (day 16), same fusion instincts (day 23). What's new is plumbing — interleaved scale loads, in-register dequant, widened accumulators, requantization at the epilogue — plus one new roofline fact: the ceiling moved, so yesterday's "compute-bound" GEMM may be memory-bound again in INT8, and re-tiling after quantization is normal, not a regression. One honest scope note: the fastest quantized GEMMs run on tensor cores with format-specific instructions, which — like multi-GPU and PTX — this journey deliberately leaves for Part 2. What you can do now is read every quantized kernel you meet and name its moving parts.

Next: day 26 — KV-cache kernels: the memory bottleneck behind every LLM server, and how paged attention and FlashDecoding reshape kernels around it.