From 463b6ee753c6aa26845b20976a85e7ed46b5ca3d Mon Sep 17 00:00:00 2001 From: Alexander Fedora Date: Wed, 21 Oct 2026 00:04:06 -0700 Subject: [PATCH 1/2] perf: Q3_K's zero point in the operands of the Apple7/8 register decode tile The F32 register decode tile ran Q3_K through operand_f32, each 16-input chain seeded from the table, as Q6_K was. Its 3-bit codes now enter the MMA as the exact half (code - 4) / 8: the two low bits at bits 7-8 and the hmask bit at 9, 1 + code / 8, less 1.5. The chains start from zero, read no seed and close with fma(dot, 8 s). CenteredChunk now holds each centered format's operand and scale. The Apple9 instances are unchanged. M2 Max, ms at one to four lanes: q3k 17408x5120 gate/up 0.586 1.005 1.300 1.767 -> 0.575 0.975 1.218 1.595 q3k 5120x17408 down 0.308 0.529 0.677 0.920 -> 0.300 0.515 0.630 0.842 q6k (unchanged) within 0.2% Q3_K is 5.6% of UD-IQ4_XS's bytes and 0.7-1.6% of the UD-Q4 files'. --- .../kernels/decode/linear_gguf_sgmatrix.metal | 28 +++++++++++++++---- 1 file changed, 23 insertions(+), 5 deletions(-) diff --git a/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal b/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal index b79679cff..2fca91316 100644 --- a/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal +++ b/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal @@ -149,6 +149,14 @@ inline half2 nibble6_operand(uint2 chunk, uint f) { const uint high = f == 3 ? chunk.y >> 4 : chunk.y << (8 - 4 * f); return as_type((at & 0x00F000F0u) | (high & 0x03000300u) | 0x3C003C00u) - half2(1.5h); } +// The exact half operand (code - 4) / 8 of pair f of a Q3_K chunk (uint2: the low bits as quant_spread2 places them, +// pair p's at bits 4p and 16 + 4p, then the hmask bits as quant_spread1 does, at 2p and 16 + 2p): the low bits at bits +// 7-8 and the hmask bit at 9 of each halfword, 1 + code / 8, less 1.5, so a chain of Q3_K's zero point 4 needs no +// seed. On an M2 Max this takes 2% off the 27B's Q3_K projections at one and two lanes and 6-10% at three and four. +inline half2 q3_operand(uint2 chunk, uint f) { + const uint low = f == 0 ? chunk.x << 7 : f == 1 ? chunk.x << 3 : chunk.x >> (4 * f - 7); + return as_type((low & 0x01800180u) | ((chunk.y << (9 - 2 * f)) & 0x02000200u) | 0x3C003C00u) - half2(1.5h); +} // Apple7/8's (F32) view of Q4_0 and Q5_0: the zero point taken as a min, m = -zero d, so their chains close once per // 32 inputs against the table's sums, as Q4_1's and Q5_1's do, instead of once per 16 against the seeds. On an M2 @@ -186,9 +194,19 @@ template <> struct NibbleChunk> { static constexpr constant bool template <> struct NibbleChunk { static constexpr constant bool value = true; }; template <> struct NibbleChunk { static constexpr constant bool value = true; }; template <> struct NibbleChunk> { static constexpr constant bool value = true; }; -// Q6_K, whose F32 chains per 16 inputs take nibble6_operand, unseeded. +// Q6_K and Q3_K, whose F32 chains per 16 inputs take nibble6_operand and q3_operand, (code - zero) / Scale, +// unseeded. template struct CenteredChunk { static constexpr constant bool value = false; }; -template <> struct CenteredChunk { static constexpr constant bool value = true; }; +template <> struct CenteredChunk { + static constexpr constant bool value = true; + static constexpr constant float Scale = 64; + static half2 operand(uint2 chunk, uint f) { return nibble6_operand(chunk, f); } +}; +template <> struct CenteredChunk { + static constexpr constant bool value = true; + static constexpr constant float Scale = 8; + static half2 operand(uint2 chunk, uint f) { return q3_operand(chunk, f); } +}; // The exact bf16 operand of pair f of a chunk. template @@ -404,7 +422,7 @@ inline void decode(device const TableElem *table, device const float *sums, #pragma unroll for (uint f = 0; f < 4; ++f) { if constexpr (nibbles) a[nf][f] = nibble_operand(cur[nf][q], f); - else if constexpr (centered) a[nf][f] = nibble6_operand(cur[nf][q], f); + else if constexpr (centered) a[nf][f] = CenteredChunk::operand(cur[nf][q], f); else if constexpr (F32) a[nf][f] = operand_f32(cur[nf][q], f, lut); else a[nf][f] = operand(cur[nf][q], f, lut); } @@ -439,8 +457,8 @@ inline void decode(device const TableElem *table, device const float *sums, acc[r][nf] = fma(dot, cs[h][nf].x, acc[r][nf]); acc[r][nf] = fma(S::Seeded ? seed[h] : sum, cs[h][nf].y, acc[r][nf]); } else if constexpr (centered) { - // code - 32 = 64 operand. - acc[r][nf] = fma(dot, 64 * float(cs[h][nf]), acc[r][nf]); + // code - zero = Scale operand. + acc[r][nf] = fma(dot, CenteredChunk::Scale * float(cs[h][nf]), acc[r][nf]); } else { acc[r][nf] = fma(dot, float(cs[h][nf]), acc[r][nf]); } From 86a2c7f545412ed9e233e8f3ee15d605a49ccf69 Mon Sep 17 00:00:00 2001 From: Alexander Fedora Date: Wed, 21 Oct 2026 00:58:39 -0700 Subject: [PATCH 2/2] perf: nibble and centered chains' epilogue factors once per coefficient on Apple7/8 The F32 per-32 nibble chains (Q4_K, Q4_1, Q4_0 and Q5_0 as ZeroAsMin, Q5_K, Q5_1, MLX affine 4-bit in groups of 32) closed each chain with fma(dot, 64 s) and fma(sum, 64 s + (m - 128 s)), two operations per chain and column fragment that every chain of a coefficient repeated. The coefficients of those formats now hold 64 s and m - 64 s, computed once when they are decoded, with the same rounding, and the centered chains' (Q6_K, Q3_K) hold Scale s rather than multiplying each chain's s. The span chains (MLX affine in groups of 64 or 128) and the Apple9 instances are unchanged. M2 Max, ms at one to four lanes, two ABBA passes: q4k 17408x5120 gate/up 0.434 0.771 1.089 1.404 -> 0.429 0.758 1.067 1.408 q4k 5120x17408 down 0.228 0.411 0.577 0.729 -> 0.225 0.401 0.564 0.732 q5k 17408x5120 gate/up 0.520 0.907 1.193 1.606 -> 0.517 0.897 1.180 1.581 q5k 5120x17408 down 0.275 0.478 0.622 0.835 -> 0.273 0.471 0.615 0.831 q40 17408x5120 gate/up 0.423 0.767 1.103 1.416 -> 0.420 0.756 1.079 1.421 q6k 17408x5120 gate/up 0.533 0.863 1.164 1.531 -> 0.531 0.862 1.151 1.530 q6k 5120x17408 down 0.280 0.457 0.608 0.808 -> 0.278 0.457 0.605 0.803 q3k 17408x5120 gate/up 0.574 0.975 1.220 1.593 -> 0.571 0.963 1.223 1.588 af4g64 (unchanged) within 0.1% --- .../kernels/decode/linear_gguf_sgmatrix.metal | 21 +++++++++++++------ 1 file changed, 15 insertions(+), 6 deletions(-) diff --git a/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal b/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal index 2fca91316..4853051ed 100644 --- a/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal +++ b/runtime/metal/kernels/decode/linear_gguf_sgmatrix.metal @@ -350,7 +350,16 @@ inline void decode(device const TableElem *table, device const float *sums, #pragma unroll for (uint i = 0; i < I; ++i) { const uint e = lane + 32 * i; - if (e < NC * S::J / JS) cu[JS * (e / NC) * NC + e % NC] = coefficient(src[i], u, JS * (e / NC)); + if (e < NC * S::J / JS) { + C k = coefficient(src[i], u, JS * (e / NC)); + // The F32 per-32 nibble chains' epilogue factors, 64 s and m - 64 s (code = 64 (operand - 1)), once per + // coefficient rather than per chain: on an M2 Max, 1-2% off Q4_K's, Q4_0's and Q5_K's 27B projections at one to + // three lanes (Q4_K and Q4_0 0.3% more at four). + if constexpr (F32 && NibbleChunk::value) k = float2(64 * k.x, fma(64.0f, k.x, k.y)); + // The centered chains' Scale s. + if constexpr (F32 && CenteredChunk::value) k = CenteredChunk::Scale * k; + cu[JS * (e / NC) * NC + e % NC] = k; + } } simdgroup_barrier(mem_flags::mem_threadgroup); #pragma unroll @@ -450,15 +459,15 @@ inline void decode(device const TableElem *table, device const float *sums, for (uint f = h * FG; f < (h + 1) * FG; ++f) sgmatrix::mma_mixed(dot, a[nf][f], b[f]); if constexpr (nibbles) { - // code = 64 (operand - 1): s code + m = 64 s operand + m - 64 s, cs.y being m - 128 s. - acc[r][nf] = fma(dot, 64 * cs[h][nf].x, acc[r][nf]); - acc[r][nf] = fma(sum, fma(64.0f, cs[h][nf].x, cs[h][nf].y), acc[r][nf]); + // code = 64 (operand - 1): s code + m = 64 s operand + m - 64 s. + acc[r][nf] = fma(dot, cs[h][nf].x, acc[r][nf]); + acc[r][nf] = fma(sum, cs[h][nf].y, acc[r][nf]); } else if constexpr (S::HasMin) { acc[r][nf] = fma(dot, cs[h][nf].x, acc[r][nf]); acc[r][nf] = fma(S::Seeded ? seed[h] : sum, cs[h][nf].y, acc[r][nf]); } else if constexpr (centered) { - // code - zero = Scale operand. - acc[r][nf] = fma(dot, CenteredChunk::Scale * float(cs[h][nf]), acc[r][nf]); + // code - zero = Scale operand, the coefficient being Scale s. + acc[r][nf] = fma(dot, float(cs[h][nf]), acc[r][nf]); } else { acc[r][nf] = fma(dot, float(cs[h][nf]), acc[r][nf]); }