Kernel math / Q4_K, MXFP4, and NVFP4

Start with a number in memory, not a model name. A computer stores that number as bits. Eight bits make a byte; memory systems transfer bytes in larger chunks; the CPU or GPU must bring the chunk close enough to a compute unit before a multiplication can happen. For large AI models, moving the weights can cost more time and energy than the multiply itself. That is the practical reason to care about four-bit precision.

Why call four-bit quantization king? It can put a much larger model within a machine's memory budget while preserving useful output quality. But the crown is conditional: Q4_K, MXFP4, and NVFP4 encode and execute weights differently, and the best choice depends on the model, hardware, workload, and quality target. We will compare those contracts by opening the bytes and doing the math.

Before the format: count the bytes

Suppose a dense model has 7 billion weights. BF16 uses 16 bits = two bytes per weight, so the raw weight payload is 14 billion bytes, or about 14 GB in decimal units. An ideal four-bit code uses half a byte per weight: 3.5 GB. Four codes do not fit in one byte; two do. At the level of a conventional 64-byte CPU cache line, 32 BF16 values fit, versus 128 raw four-bit codes. That fourfold density sounds dramatic, but a real quantized line or block also needs scales and perhaps offsets, and some tensors stay at higher precision.

Take the actual Q4_K block we will open below: 256 weights occupy 144 bytes including their metadata, which is 4.5 bits per weight. Seven billion weights all stored in Q4_K would occupy about 3.94 GB, not 3.5 GB. That is still much smaller than 14 GB BF16, but the model may contain embeddings, norms, and other tensors in different formats. The runtime also needs activations, scratch buffers, and a KV cache. Do not equate the checkpoint's “4-bit” label with total process memory.

Imagine an idealized device able to stream 100 GB/s of weights from main memory. One uncached sweep of 14 GB takes at least 140 ms; one sweep of 3.94 GB takes at least 39 ms. These are bandwidth-only lower bounds, not measured token latencies: cache hits, MoE sparsity, scale reads, dequantization, arithmetic, synchronization, and attention all change the outcome. The useful question is whether a kernel can save more time moving fewer bytes than it spends decoding them.

The path matters: DRAM (or GPU HBM) holds the packed weights; caches hold recently fetched lines; vector registers receive bytes and scales; the kernel extracts four-bit fields into lanes, reconstructs arithmetic values or maps them to a supported dot instruction, and accumulates in a wider type. A “four-bit model” does not imply a CPU has a native four-bit multiply. On supported Blackwell GPUs, a particular FP4 Tensor Core path exists; on many CPUs, four-bit is chiefly the storage precision while arithmetic proceeds through INT8, BF16, or FP32 lanes.

Now we can ask why a model can speak coherently after such compression. It is not because four bits preserve every weight. They do not. A code plus a nearby scale can preserve enough useful structure, and selective higher precision can protect sensitive tensors. The quality is conditional on the model, the layer, the quantizer, and the task. We will derive one dot product, inspect each block format, and connect the bytes to the kernel. If wx+b and its gradient is already familiar, this is the next step: changing how the w values are represented before the multiply.

1. A four-number example you can calculate by hand

Take the illustrative signed-INT4 weight block w = [-1.000, 0.250, 0.875, 1.750]. Signed four-bit integers run from -8 to 7. Choose scale s = 0.25 and calculate q_i = clip(round(w_i / s), -8, 7). We obtain [-4, 1, 4, 7]. Two's-complement four-bit encodings are C, 1, 4, and 7 in hexadecimal. If the format places the earlier code in the low nibble, the two payload bytes are 0x1C and 0x74. The scale must also be stored or derivable; the payload alone cannot recover a weight.

Dequantization gives w_hat_i = s q_i = [-1.000, 0.250, 1.000, 1.750]. The 0.875 weight becomes 1.000, an error of +0.125. With x = [2, 1, -1, 0.5], the original dot product is -2 + 0.25 - 0.875 + 0.875 = -1.750; the reconstructed dot product is -2 + 0.25 - 1 + 0.875 = -1.875. Output error is -0.125. One dot product makes the loss visible; a network can sometimes tolerate such errors, but not automatically.

Notice what the runtime need not do: build a complete BF16 copy of the matrix before every multiply. A useful kernel loads packed bytes and scales, extracts nibbles, turns them into arithmetic values, multiplies by activations, and accumulates in a wider type. That is the difference between compressed storage and fast quantized execution.

2. Why scale each block instead of one whole matrix?

Suppose one region of a matrix has weights around 0.02 and another has weights around 2.0. A single scale chosen for the 2.0 outlier may map the 0.02 weights to zero. A smaller group gets its own scale and therefore finer local resolution. The cost is metadata: one scale per 16 or 32 values takes more bytes than one scale per tensor. Good formats search for the right tradeoff among local error, metadata bandwidth, and hardware-friendly block geometry.

A toy row makes this concrete. For x = [0.02, 0.04, 0.06, 0.08], a coarse global step of 0.25 maps all four to zero. A local step of 0.02 maps them to [1, 2, 3, 4] and reconstructs them exactly. Real model distributions and fitted scales are less tidy, but this is why a shared scale is part of the number system, not a footnote.

3. Four things commonly called “4-bit”

NVFP4 step by step. E2M1 is a tiny floating-point code with one sign, two exponent, and one mantissa bit; its positive magnitudes include 0, 0.5, 1, 1.5, 2, 3, 4, and 6. In the documented Transformer Engine recipe, x_hat = q_E2M1 x s_block x s_global. A block of 16 numbers shares one E4M3 scale; the tensor has an FP32 global scale. For a simplified illustration, if q = 3, s_block = 0.5, and s_global = 2, the reconstructed value is 3. The real recipe computes and rounds those scales subject to FP8 range and may use additional training techniques. NVIDIA documents the exact recipe and supported devices.

MXFP4 step by step. The element is also E2M1, but the OCP MX specification gives each group of 32 elements an E8M0 scale, a power of two. If q = 3 and the group scale is 2^-1, x_hat = 1.5. A power-of-two scale is cheap to represent and use, but may fit an awkward group maximum less closely than a finer FP8 scale. A different group size also changes how outliers contaminate neighbours. The OCP specification, not the four-bit label, is the contract.

Q4_K step by step. CKE's dequantizer states its GGML-compatible formula explicitly: w_hat = d x s_j x q - dmin x m_j, where d and dmin are superblock factors and s_j and m_j are subblock codes. A 256-weight superblock contains eight 32-weight subblocks and packed four-bit values. This is an affine integer reconstruction, not E2M1 floating point. The suffix _M is a model-level mixture choice: inspect the actual GGUF tensor inventory rather than assuming every tensor has the same type. See llama.cpp's encoding table and CKE's dequantization code.

Unsloth Dynamic step by step. First estimate which tensors are sensitive to precision loss. Quantize robust tensors aggressively; keep sensitive ones at more bits or a different quantizer. That can spend a small amount of extra memory to avoid a disproportionately large quality loss. The original Unsloth Dynamic 4-bit explanation describes opting out for selected parameters; Dynamic v2 extends the selection and evaluation process. “UD-Q4_K_M” therefore names a selection strategy plus a GGUF recipe, not a new four-bit primitive. Do not equate it with Unsloth's fused GPU kernels or its NF4 training path.

4. Open the Q4_K superblock and count every byte

The toy signed-INT4 example taught the mechanism, but Q4_K is the production-format lesson. Its 256 weights are divided into eight groups of 32. In the CKE/GGML block layout, 128 bytes hold the four-bit codes; two FP16 values d and dmin take four bytes; 12 bytes hold the eight 6-bit scale/minimum pairs. Total: 128 + 4 + 12 = 144 bytes. Divide by 256 and multiply by eight: 4.5 bits per weight, not a literal four bits per complete stored weight. See the CKE byte-level format reference for the neighbouring Q4_0, Q5_K, Q6_K, and Q8_K layouts.

Here is the dequantization rule, with subblock j and code q in 0..15. The minus before the minimum is essential:

\[\hat w_{j,i}=q_{j,i}(d\,s_j)-d_{\min}m_j.\]

Try one subblock with d = 0.10, s_j = 3, dmin = 0.05, and m_j = 2. Its effective step is 0.30 and its offset is -0.10. A stored code q = 7 reconstructs 7(0.30) - 0.10 = 2.00; q = 2 reconstructs 0.60 - 0.10 = 0.50; q = 0 reconstructs -0.10. Those are illustrative decoded header values, not a claim about how this exact block was optimized from source weights. The point is that all 32 codes share the same effective step and offset, while the next 32 may have different s_j and m_j.

Now place this in a dot product. For activation values x = [0.5, -1, 2] and quantized codes q = [7, 2, 0] from this same subblock, the reconstructed contribution is 2.00(0.5) + 0.50(-1) + (-0.10)(2) = 0.30. A fused kernel can unpack the nibbles, apply these two header-derived factors, and accumulate that 0.30 without writing 32 expanded floats to memory. CKE's reference dequantizer is the executable version of this formula.

5. Decode the E2M1 nibble before comparing FP4 formats

Both NVFP4 and MXFP4 use the same E2M1 element concept, but it is not a signed integer from -8 to 7. One sign bit, two exponent bits, and one fraction bit produce a small, non-uniform magnitude set: 0, 0.5, 1, 1.5, 2, 3, 4, and 6, plus their negatives. The spacing widens at larger magnitudes. This is why copying the signed-INT4 rounding equation into an FP4 decoder is wrong; the nearest code must be chosen from these floating-point values after applying the block scale.

For the ambiguous 1.25 row, 2 x 0.5 = 1.0 and 3 x 0.5 = 1.5 are equally distant; exact tie-breaking is a format/converter rule. The table is meant to expose the candidate values, not prescribe that tie. CKE's E2M1 lookup table makes the 16 possible nibble codes explicit: codes 0..7 map to 0, 0.5, 1, 1.5, 2, 3, 4, 6; codes 8..15 map to signed zero and the corresponding negatives. The C path stores twice these values as small integers and folds the compensating 0.5 into its E4M3 scale lookup. That is an implementation trick, not a change to the mathematical format.

NVIDIA's documented NVFP4 recipe reconstructs q_E2M1 x s_E4M3_block x s_FP32_global. For an example with q = 1.5, block scale 0.5, and global scale 2, the output is 1.5. The global scale widens usable tensor range; the per-16 scale adapts locally. MXFP4 instead uses an E8M0 power-of-two scale per 32 elements, so the same q = 1.5 at scale 2^-1 reconstructs 0.75. The group size and scale number system differ even when the four element bits are identical. See the OCP MX specification and NVIDIA's NVFP4 recipe.

CKE's current NVFP4 storage block is especially concrete: four unsigned E4M3 subscales and 32 bytes of packed E2M1 codes hold 64 expert weights, for 36 bytes / 64 = 4.5 bits per weight; the FP32 global/expert scale is a separate runtime operand. Within each 16-weight subblock, the low nibble of eight bytes maps lanes 0-7 and the high nibble maps lanes 8-15: split-half order, not the adjacent-pair order of our toy INT4 example. The CKE deep dive and NVFP4 implementation PR document this path. Today, CKE's NVFP4 maps are candidate MoE-expert paths, not blanket dense-model NVFP4 support.

This is also where a converter can silently ruin an otherwise correct kernel: if it interleaves the low and high nibbles rather than using split-half order, the arithmetic may be mathematically correct on the wrong lanes. CKE's packing contract test checks the bytes and scale transform before model-level parity is attempted.

6. The execution path is the rest of the algorithm

Read that diagram left to right: a checkpoint supplies bytes and tensor dtypes; a converter preserves the layout contract; the generated runtime chooses a provider for each tensor's dtype and operation; the kernel consumes a tile; accumulation proceeds in a wider type. A recipe such as Q4_K_M or an Unsloth UD GGUF can change which tensor uses which existing dtype without requiring a new four-bit arithmetic kernel. CKE's format census gives a concrete mixed-format example and reports a Gemma3 Unsloth UD artifact with a still-open longer-context boundary. That is better evidence than declaring all UD quants bit-exact.

The hardware question now becomes precise: how many bytes does this kernel load, how many instructions does it spend unpacking, which arithmetic unit performs the dot product, and what intermediate precision is used? A weight-only path may use vector FP32 FMA; a W4A8 path may unpack to INT8 and use VNNI or AMX-INT8; a Blackwell FP4 path may use native Tensor Cores. These are different end-to-end costs even if their checkpoint files are similar sizes.

7. From packed weights to a real matrix multiply

For C = A x B, a weight-only W4 path often keeps B packed while A remains BF16/FP16 or is dynamically quantized. A kernel processes tiles: load a block of A, load packed B and its scales, decode two four-bit values from each byte, apply the right scale/offset, accumulate partial dot products, and write a higher-precision output. If the implementation materializes the full dequantized B first, it gives back much of the bandwidth advantage. If it fuses decode with GEMM/GEMV, it pays some extra instructions but reads fewer weight bytes from memory.

W4A8 means four-bit weights and eight-bit activations; it is not another file suffix. One implementation may unpack INT4 weights to INT8 and use an INT8 dot-product instruction; another may dequantize to BF16 and use BF16 matrix hardware. The output/accumulator precision is a separate decision. Activation quantization needs its own scales and outlier policy. W4A16 weight-only inference is different from W4A8 even when both store the same weight file.

We can extend the Q4_K subblock example into a W4A8 calculation. Keep q_w = [7, 2, 0], effective weight step d s_j = 0.30, and minimum offset dmin m_j = 0.10. Quantize x = [0.5, -1, 2] with activation scale s_x = 0.5, producing eight-bit codes q_x = [1, -2, 4]. Then:

\[\sum_i\hat w_i\hat x_i = s_x(d s_j)\sum_i q_{w,i}q_{x,i} - s_x(d_{\min}m_j)\sum_i q_{x,i}.\]

The integer dot is 7(1) + 2(-2) + 0(4) = 3. The activation-code sum is 1 - 2 + 4 = 3. Rescale and correct the offset: 0.5(0.30)(3) - 0.5(0.10)(3) = 0.45 - 0.15 = 0.30, exactly the dequantize-then-multiply result above for this toy row. The extra sum is why an affine weight format cannot always be treated as a plain INT4 x INT8 dot with one final scale. An optimized kernel may precompute, fuse, or algebraically reorganize this correction, but it must preserve the same numerical contract.

On x86, AVX-512 can unpack nibbles and vectorize integer or floating-point arithmetic; AVX-512 VNNI accelerates suitable INT8 dot products. Intel AMX-INT8 and AMX-BF16 accelerate tile products, but a four-bit packed weight does not become an AMX native INT4 operand just because the CPU has AMX. The software must unpack/reformat and schedule tiles. A carefully fused CPU kernel may win on memory traffic; a poor conversion path may not. CKE's Q4_K x Q8_K VNNI and NVFP4 GEMM paths show two distinct engineering approaches, but file existence alone is not a speed or parity result. Intel's AMX/VNNI tuning guide describes the hardware instructions actually available.

On supported NVIDIA Blackwell GPUs, Tensor Cores can consume native FP4 block-scaled layouts. CUDA cores may do auxiliary conversion and address work; they are not interchangeable with Tensor Core FP4 matrix instructions. Layout and scale tiles must match the kernel's contract. On an older GPU, the same logical FP4 weights may need unpack-and-convert, so “4-bit” can save memory without providing the expected FP4 compute speed. Do not transfer a GB200 throughput chart to a CPU, a pre-Blackwell GPU, or a different model.

7a. Four rows of weights: what is actually stored?

Think of a GEMM weight matrix as many rows of weights. We will follow one row across four representations; a row is a 1D slice of the 2D matrix, not four different 1D matrix formats. If an input row is x = [x0, x1, ...], one output is y = sum(wi xi). Every format below must ultimately contribute to that same sum. The bracketed arrays are teaching slices; real blocks must have the full lengths shown in the second column.

7b. Decode one lane, then repeat across a tile

Step 1: read the header. For Q4_K, decode d and dmin from FP16 and unpack subgroup j's six-bit sj and mj. Our earlier illustrative subgroup had d sj = 0.30 and dmin mj = 0.10. For MXFP4, read the group's E8M0 exponent and obtain its power-of-two scale. For NVFP4, read the E4M3 subscale and the separately supplied global scale. Those are three different scale contracts, not three spellings of INT4.

Step 2: read and unpack codes. In a Q4_K pair, byte & 0x0f exposes the low nibble and byte >> 4 the high nibble. CKE's Q4_K block maps the first 32 low nibbles and next 32 high nibbles of a 32-byte span to two adjacent subgroups. NVFP4 also needs its specified split-half lane order, but its nibble is an E2M1 value, not an integer 0..15. MXFP4 uses E2M1 too, with a different group scale. A valid dequantizer must know both the bit meaning and the lane order.

Step 3: reconstruct or keep the dot in integers. For Q4_K codes [7, 2, 0], the example subgroup formula produces [2.00, 0.50, -0.10]. With x = [0.5, -1, 2], the dot is 1.00 - 0.50 - 0.20 = 0.30. For an MXFP4 example, E2M1 codes representing [1, 1.5, -2] and an E8M0 scale of 2^-1 give [0.5, 0.75, -1]. For an NVFP4 example, the same codes with E4M3 scale 0.5 and global scale 2 give [1, 1.5, -2]. These are independent teaching examples, not one shared physical checkpoint.

Step 4: accumulate the matrix product. A scalar reference may reconstruct FP32 values and multiply them by FP32 activations. A vector kernel need not. In the Q4_K x Q8_K route, it can keep unpacked weights as small unsigned integers and activations as signed INT8, accumulate their dot into INT32, then apply the subgroup scales and minimum correction in FP32. The result becomes one contribution to a GEMV output or one tile of a GEMM. The packed file stays small in memory even though registers temporarily hold wider lanes and accumulators.

float sum = 0.0f;
float step = fp16_to_f32(block.d) * subgroup_scale;
float offset = fp16_to_f32(block.dmin) * subgroup_min;
for (int i = 0; i < 32; ++i) {
    unsigned q = unpack_q4_k_lane(block.qs, subgroup, i);
    float w = step * (float)q - offset;
    sum += w * x[i];
}
// This illustrates the math; real Q4_K lane order matters.

Compare that reference with CKE's scalar dequant_q4_k_block. It reads d and dmin, unpacks the eight scale/minimum pairs, and handles the low/high halves in their prescribed order. Its output array is useful as a correctness oracle; a fast inference path need not create that array.

7c. What AVX-512 VNNI does, instruction by instruction

The source path is CKE's Q4_K x Q8_K VNNI kernel. In its prepared-weight route, one 512-bit load brings 64 packed bytes into a vector register. _mm512_and_si512(packed, 0x0f) extracts the low four bits of each byte. _mm512_srli_epi16(packed, 4) followed by the same mask exposes the high nibbles. The result is byte-sized integer lanes in registers, not a 64-element FP32 weight array written back to memory.

Next, _mm512_dpbusd_epi32 multiplies groups of four unsigned weight bytes by signed activation bytes and adds each four-product sum to a signed INT32 lane. The Q4 codes are unsigned, so this pairing is deliberate. The kernel runs low and high halves separately because they correspond to different Q4_K subgroups. It then multiplies those integer sums by the subgroup scales, subtracts the minimum term using precomputed Q8_K activation sums, applies the activation scale, and accumulates a FP32 output. It is not a single instruction that understands a Q4_K file header.

__m512i packed = _mm512_loadu_si512(weight_bytes);
__m512i qlo = _mm512_and_si512(packed, nibble_mask);
__m512i qhi = _mm512_and_si512(
    _mm512_srli_epi16(packed, 4), nibble_mask);
sum_lo = _mm512_dpbusd_epi32(sum_lo, qlo, q8_lanes);
sum_hi = _mm512_dpbusd_epi32(sum_hi, qhi, q8_lanes_next);
// Then scale, subtract Q4_K minimum correction, and accumulate FP32.

This is a reading sketch, not a drop-in replacement for CKE's tiled kernel: the real code handles prepared layouts, Q8 lane broadcasts, tails, row scheduling, and subgroup indices. It explains where SIMD helps. AVX-512 handles many nibbles at once; VNNI reduces several byte multiplies into INT32 sums; FP32 arithmetic restores the Q4_K scale/offset contract. In a separate NVFP4 CPU route, CKE maps E2M1 codes through a tiny signed-integer table and folds a factor into its E4M3 scale lookup. Neither route implies that CKE supports native MXFP4 execution today.

7d. When the smaller archive becomes a faster computation

For one token's matrix-vector decode, the kernel often must stream a large fraction of the weight matrix for relatively little arithmetic per weight. If it reads 144 Q4_K bytes instead of 512 BF16 bytes for each 256-weight block, it may save enough DRAM/HBM traffic to pay for unpacking and scaling. That is why the same device can generate faster after weight quantization. But with a large prefill batch, each loaded weight may be reused across many tokens: arithmetic intensity rises and decode overhead matters more. Native FP4 matrix hardware can change that tradeoff again. Do not infer a fourfold tokens/s speedup from a roughly fourfold payload reduction.

A useful mental model is a roofline: time is bounded below by both bytes moved / sustainable bandwidth and operations / sustainable compute, plus conversion, launch, scheduling, and other model work. Quantization attacks the first term; fused kernels try not to inflate the others. Measure memory footprint, memory bandwidth, prefill and decode separately, and accuracy against the same unquantized baseline.

7e. What CKE has actually implemented

8. What improves, and what can break?

Memory: weight payload drops toward one quarter of BF16, but scales and mixed-precision tensors reduce the headline savings. The KV cache grows with sequence length and is not shrunk by weight-only Q4. At long context, KV memory and attention work can dominate; changing the weight format alone may stop helping much. KV-cache quantization is a different operation with different error modes.

Speed: on bandwidth-bound decode, fewer weight bytes can be valuable. On prefill or a compute-bound GEMM, a native FP4 Tensor Core path may win, but a dequantization-heavy path may lose. Benchmark end-to-end tokens/s and first-token latency, not just the packed GEMM kernel. Batch size, sequence length, memory bandwidth, and model shape change the answer.

Quality: do not infer “imperceptible until very long context” from four bits alone. A sensitive projection, output head, small model, code task, or tool-calling format can fail at short context. A long document may expose an error that a short prompt did not, but long-context failures may also arise from positional generalization or KV-cache quantization. Compare against the same BF16 model with identical prompts, tokenizer, sampling, and KV precision. Check perplexity or KL divergence on held-out tokens, then exact task metrics for math, code, retrieval at several lengths, and structured tool calls. Repeat across seeds where sampling matters.

A practical test grid has two axes: weight format (BF16, Q4_K_M, UD-Q4_K_M, NVFP4 or MXFP4 where supported) and context length (say 2k, 8k, 32k, and a long case the model truly supports). Hold KV precision fixed for the first sweep. Report model size, peak RAM/VRAM, prefill throughput, decode throughput, first-token latency, answer correctness, and task-specific failures. Then run a second sweep for KV-cache precision. Without that separation, it is easy to blame weight quantization for a cache or context-window issue.

9. The kernel-engine lesson

A quantized checkpoint is a contract across converter, layout, dequantizer, GEMM/GEMV dispatch, and model outputs. CKE's quantization overview and format reference are useful companions. Its NVFP4 contract test checks packing and scale handling, while Q4_K code implements a different reconstruction. That is why “supports 4-bit” is too vague a server claim. Ask which format, which tensors, which ISA, which parity tests, and which measured workload.

The punch line is simple: four-bit models are powerful because they move less data and because smart scales and mixed precision preserve useful information. The bit width is only the first line of the design. The scaling math, packing contract, hardware path, and evaluation tell us whether this particular four-bit model is actually the right one to run.

Keep reading

If the dot product felt fast, revisit Backpropagation From Wx Plus B for the original training loop and Muon Optimizer: From wx+b to Matrix Updates for a matrix-level follow-on. For nearby precision engineering, read What Is Quantization and Why Backprop Gets Hard, the earlier CKE quantization format deep dive, and What Numerical Parity Actually Requires. To understand why decode speed differs from prefill, see Prefill vs Decode in CKE. The CKE quantization tutorial and visual format guide contain the zoomable byte diagrams behind this article.