ggml Q4_0 vec_dot Walkthrough

dst[12288] = W_gate[12288 × 4096] · x[4096]      W_gate: Q4_0, x: F32

This is the FFN gate projection of Qwen3-8B for one token: decode handles one token per step, so every weight matrix is multiplied by a vector. Each row of W holds K = 4096 weights, and there are N = 12288 rows, one per output.

Setup:

• Apple M4 Pro: 10 performance and 4 efficiency cores; NEON with dotprod and i8mm.
• ggml 83835ff (Aug 2025), Release build, CPU backend only.
• The thread figures use the same ggml plus a trace in mul_mat that logs which rows each thread took, and when.

01Quantization Types in GGUF#

A GGUF file stores each tensor with a ggml_type. Quantized types are block formats: a row is cut into fixed-size blocks, and each block stores one or more scales and a low-bit integer code per weight.

1.1Q8_0 and Q4_0#

From src/ggml-common.h:

#define QK8_0 32
typedef struct {
    ggml_half d;           // delta
    int8_t  qs[QK8_0];     // quants
} block_q8_0;              // 2 + 32 = 34 bytes per 32 values -> 8.5 bit/weight

#define QK4_0 32
typedef struct {
    ggml_half d;           // delta
    uint8_t qs[QK4_0 / 2]; // nibbles / quants
} block_q4_0;              // 2 + 16 = 18 bytes per 32 values -> 4.5 bit/weight

• Block: 32 consecutive values of a row.
• d: one fp16 scale per block, called delta in the comment because it is the step between two neighboring codes.
• qs: one integer code per value, which decodes as .

Neither format stores an offset, so code always decodes to . That is what _0 means, while Q4_1 and Q5_1 add a per-block offset and decode as .

Q8_0 is quantized by quantize_row_q8_0_ref, which sets , where is the largest absolute value in the block, and stores , which falls in .

Q4_0 stores each code as a 4-bit nibble, two per byte, and subtracts 8 when decoding: , so . quantize_row_q4_0_ref takes the value with the largest magnitude, keeps its sign as , and sets , so that value decodes exactly, as code .

1.2Q2_K: two levels of scales#

The _K types (“k-quants”) split a row into super-blocks of 256 weights, and each super-block into 16 groups of 16. Q2_K gives each weight a 2-bit code, so a group has only 4 levels.

A weight decodes in two steps:

1. Per group. Group has a scale and an offset, and a weight is , with . The offset moves the 4 levels to where the group’s values are; without it they would always be 0, 1, 2 and 3 times the scale.
2. Per super-block. and are 4-bit integers and , times two fp16 values shared by the whole super-block:
   ggml calls the offsets mins, hence the names dmin and .

Together:

Figure 2 draws this with the numbers of the example below.

In memory (src/ggml-common.h):

#define QK_K 256
typedef struct {
    uint8_t scales[QK_K/16]; // scales and mins, quantized with 4 bits
    uint8_t qs[QK_K/4];      // quants
    ggml_half d;             // super-block scale for quantized scales
    ggml_half dmin;          // super-block scale for quantized mins
} block_q2_K;                // 16 + 64 + 4 = 84 bytes per 256 weights -> 2.625 bit/weight

The offsets have no array of their own: scales[j] holds in its low 4 bits and in its high 4 bits, and the decoder splits them with sc & 0xF and sc >> 4. Per weight, that is 2 bits of code, 0.5 bits for and (8 bits per 16 weights), and 0.125 bits for d and dmin (32 bits per 256): 2.625 bits.

1.3Types vs presets: Q4_K_M, Q2_K, Q2_K_S#

Q4_K is a tensor type, a block format like the ones above, while Q4_K_M is a preset for llama-quantize (a llama_ftype in llama.cpp). The preset stores most tensors in its base type, q4_K, and a few sensitive ones, such as attn_v and ffn_down, in a type with more bits.

The suffix sets how many tensors get the extra bits: fewest for _S, more for _M, most for _L; more bits mean a larger file and a more accurate model. llama-quantize --help lists the numbers for Llama-3-8B, where perplexity measures how well a model predicts text (lower is better) and the last column is how much it rises over the fp16 model:

Not every base type has all three: Q3_K has _S, _M and _L; Q4_K and Q5_K have _S and _M; Q2_K has Q2_K and Q2_K_S; Q6_K has none. A preset with no suffix, such as Q4_K, means its _M.

Qwen3-8B-Q4_K_M.gguf, read with ggml’s gguf_init_from_file:

• Most tensors are q4_K.
• attn_v and ffn_down are q6_K on 18 of the 36 layers: the first 4, the last 5, and every third one in between (use_more_bits() in llama-quant.cpp).
• output is also q6_K; the norm vectors stay f32.

There is no Q2_K_M, only Q2_K and Q2_K_S. For Qwen3-8B (32 query heads and 8 KV heads, so n_gqa = 4), their rules give:

On the CPU, each type has its own dot-product kernel and its own format for the activation: Q4_0 and Q8_0 pair with Q8_0, and the _K types with Q8_K (8-bit codes in blocks of 256). The rest of the post uses Q4_0, whose kernel is the shortest, but how mul_mat splits the work across threads doesn’t depend on the type.

02One Row: vec_dot#

The CPU backend’s entry for Q4_0 (src/ggml-cpu/ggml-cpu.c):

[GGML_TYPE_Q4_0] = {
    .from_float   = quantize_row_q4_0,
    .vec_dot      = ggml_vec_dot_q4_0_q8_0,
    .vec_dot_type = GGML_TYPE_Q8_0,
    .nrows        = 2,   // 1 without __ARM_FEATURE_MATMUL_INT8
},

• vec_dot_type: the format x is converted to first, here Q8_0, so the inner loop multiplies int8 by int8. The conversion runs once per mul_mat and covers 4096 values, against 50 million weights.
• from_float: the function that converts x.
• vec_dot: computes one output, the dot product of one row of W and x.
• nrows: rows per vec_dot call; see the note at the end of this section.

Both operands are now in blocks of 32 with one scale each, and the blocks line up along K. For one output, row r of W against x:

• wq[b][j] is the code of value j in block b of the row, as an int8 (for Q4_0, the nibble minus 8), and wd[b] is that block’s scale.
• xq[b][j] and xd[b] are the same for x, after its conversion to Q8_0.

float sum = 0;
for (int i = 0; i < 4096; i++) {
  int b = i / 32, j = i % 32;
  float w = wd[b] * wq[b][j];
  float a = xd[b] * xq[b][j];
  sum += w * a;
}

float sum = 0;
for (int b = 0; b < 128; b++) {
  int isum = 0;
  for (int j = 0; j < 32; j++)
    isum += wq[b][j] * xq[b][j];
  sum += wd[b] * xd[b] * isum;
}

The left version turns each weight and each activation into a float before multiplying: 3 float multiplies per value, 12288 per row. But w * a is the scale product wd[b] * xd[b] times the code product wq[b][j] * xq[b][j], and the scale product is the same for all 32 values of a block. So the right version multiplies the codes as integers, which is exact, and applies the two scales once per block: 128 times per row instead of 4096.

One Q4_0 block and one Q8_0 block, from a small probe (K = 64):

Byte qs[j] holds element j in its low nibble and element j + 16 in its high nibble. One AND gives elements 0–15 in order and one shift gives 16–31, with no shuffle, and each half lines up with one half of the Q8_0 block. The portable C version, where x is the weight row and y the quantized activation:

for (int ib = 0; ib < nb; ++ib) {
    int sumi0 = 0;
    int sumi1 = 0;

    for (int j = 0; j < qk/2; ++j) {
        const int v0 = (x[ib].qs[j] & 0x0F) - 8;
        const int v1 = (x[ib].qs[j] >>   4) - 8;

        sumi0 += (v0 * y[ib].qs[j]);
        sumi1 += (v1 * y[ib].qs[j + qk/2]);
    }

    int sumi = sumi0 + sumi1;
    sumf += sumi*GGML_CPU_FP16_TO_FP32(x[ib].d)*GGML_CPU_FP16_TO_FP32(y[ib].d);
}

2.1The NEON kernel#

A NEON register is 128 bits, split into lanes: 16 int8 values, or 4 int32 or fp32 values. The ARM version (src/ggml-cpu/arch/arm/quants.c) reads the codes of a whole Q4_0 block with one 16-byte load:

const uint8x16_t v0_0 = vld1q_u8(x0->qs);

const int8x16_t v0_0ls = vsubq_s8(vreinterpretq_s8_u8(vandq_u8  (v0_0, m4b)), s8b);  // elements 0-15
const int8x16_t v0_0hs = vsubq_s8(vreinterpretq_s8_u8(vshrq_n_u8(v0_0, 4)),   s8b);  // elements 16-31

const int8x16_t v1_0l = vld1q_s8(y0->qs);
const int8x16_t v1_0h = vld1q_s8(y0->qs + 16);

const int32x4_t p_0 = ggml_vdotq_s32(ggml_vdotq_s32(vdupq_n_s32(0), v0_0ls, v1_0l), v0_0hs, v1_0h);

sumv0 = vmlaq_n_f32(sumv0, vcvtq_f32_s32(p_0), GGML_CPU_FP16_TO_FP32(x0->d)*GGML_CPU_FP16_TO_FP32(y0->d));

ggml_vdotq_s32 is SDOT from the dotprod extension, which multiplies 16 pairs of int8 values and adds each run of 4 consecutive products into one of 4 int32 lanes. Two SDOTs cover the block’s 32 products. The 4 lanes are converted to fp32, multiplied by the two scales and added lane by lane into the accumulator sumv0, without being summed together yet.

Each iteration takes two blocks, each into its own accumulator, so one block’s multiply-add doesn’t wait for the previous one. A row of 4096 values is 128 blocks, so 64 iterations:

At the end of the row, the two accumulators are added and vaddvq_f32 sums their 4 lanes.

03Who Computes What#

vec_dot runs on one thread and produces one output, so the parallelism comes from a thread pool: ggml starts nth threads once, and for each op every thread calls the same function, here ggml_compute_forward_mul_mat, with its own index ith, and the function decides which part of the work is thread ith’s. For the gate projection:

• Chunk: 64 rows of W, the unit a thread takes; 12288 / 64 = 192 chunks.
• Tile: 16 rows; a chunk is 4 tiles. vec_dot writes a tile’s 16 results into a small buffer, which is then copied into dst. In prefill a tile is 16 rows × 16 columns of x, so each column stays in cache while 16 rows use it. In decode there is one column, and a tile is just 16 vec_dot calls.
• vec_dot: one row, 64 NEON iterations.

Only chunks are shared between threads: once a thread takes a chunk, it computes all of it.

3.1The two phases#

// phase 1: every thread quantizes its own slice of x (along K) into wdata
int64_t ne10_block_start = (ith * ne10/bs) / nth;
int64_t ne10_block_end   = ((ith + 1) * ne10/bs) / nth;
from_float(x + ne10_block_start*bs, wdata + ne10_block_start*nbw0,
           (ne10_block_end - ne10_block_start) * bs);

if (ith == 0) {
    atomic_store_explicit(&params->threadpool->current_chunk, nth, memory_order_relaxed);
}
ggml_barrier(params->threadpool);

// phase 2: chunks of dst rows
int chunk_size = 16;
if (nr0 == 1 || nr1 == 1) {
    chunk_size = 64;                       // decode: nr1 = 1
}
int64_t nchunk0 = (nr0 + chunk_size - 1) / chunk_size;    // 12288 / 64 = 192
...
int current_chunk = ith;                   // first chunk: my own index
while (current_chunk < nchunk0 * nchunk1) {
    ggml_compute_forward_mul_mat_one_chunk(params, dst, ..., ir0_start, ir0_end, ...);
    current_chunk = atomic_fetch_add_explicit(&params->threadpool->current_chunk, 1, memory_order_relaxed);
}

The loops over the batch dimensions are left out.

Phase 1 splits x along K, the same way every time. x has 4096 / 32 = 128 blocks, and thread ith quantizes blocks [ith · 128 / nth, (ith + 1) · 128 / nth) into wdata, a scratch buffer shared by all threads. Every thread then waits at ggml_barrier until all nth have arrived, since in phase 2 any thread may read any block of x.

Phase 2 splits the output into chunks of up to chunk_size rows by chunk_size columns of dst. Decode has one column (nr1 = 1), where 16 × 16 would mean only 16 outputs per chunk, so ggml raises chunk_size to 64: 192 chunks of 64 rows. Threads take chunks as they go:

• Thread i starts on chunk i, so chunks 0 to nth − 1 are already taken, and thread 0 sets the shared counter current_chunk to nth before the barrier.
• When a thread finishes a chunk, it calls atomic_fetch_add on the counter, which returns the old value and adds 1 in one atomic step. The old value is the thread’s next chunk, so no two threads get the same one.
• Once the counter reaches 192, threads leave the loop.

With fewer than 4 · nth chunks, or on a NUMA system, ggml skips this and gives each thread one equal range. That doesn’t happen here: 192 chunks is well above 4 · 14 = 56.

One traced mul_mat with 1, 2 and 4 threads, where each square is a chunk, 16 per line in row order, coloured by the thread that computed it: