Crate ferrox_quant 

ferrox-quant: dequantization kernels for the block-quantized tensor formats used by GGUF files (Q4_0, Q8_0, Q4_K, Q5_K, Q6_K).

These block layouts are a public, widely documented convention (originated in ggml). The functions here are independent implementations written against that public layout description, not copied from any other project’s source. Q4_K and Q6_K in particular are the dominant real-world GGUF quantization formats (most published checkpoints ship as Q4_K_M or similar K-quant mixes, not the legacy Q4_0/Q8_0 formats). These are checked against independent Python cross-validation, following the same discipline as ferrox-models’ GGUF-roundtrip tests.

Re-exports

pub use repack::gemm_q4_0x4_group;

pub use repack::gemm_q4_0x4_group_x4;

pub use repack::gemm_q4_kx8_group;

pub use repack::gemm_q4_kx8_group_x4;

pub use repack::gemm_q5_kx8_group;

pub use repack::gemm_q5_kx8_group_x4;

pub use repack::gemm_q6_kx8_group;

pub use repack::gemm_q6_kx8_group_x4;

pub use repack::gemm_q8_0x4_group;

pub use repack::gemm_q8_0x4_group_x4;

pub use repack::gemv_q4_0x4_group;

pub use repack::gemv_q4_kx8_group;

pub use repack::gemv_q4_kx8_q8_k;

pub use repack::gemv_q5_kx8_group;

pub use repack::gemv_q5_kx8_q8_k;

pub use repack::gemv_q6_kx8_group;

pub use repack::gemv_q6_kx8_q8_k;

pub use repack::gemv_q8_0x4_group;

pub use repack::gemv_q8_0x4_q8_0;

pub use repack::make_block_q4_0x4;

pub use repack::make_block_q4_kx8;

pub use repack::make_block_q5_kx8;

pub use repack::make_block_q6_kx8;

pub use repack::make_block_q8_0x4;

pub use repack::pack_q4_0_matrix_x4;

pub use repack::pack_q4_k_matrix_x8;

pub use repack::pack_q5_k_matrix_x8;

pub use repack::pack_q6_k_matrix_x8;

pub use repack::pack_q8_0_matrix_x4;

pub use repack::prepare_q8_acts_x4;

pub use repack::prepare_q8_k_acts_x4;

pub use repack::q4_0x4_gemm_uses_acts_x4;

pub use repack::q4_0x4_interleave;

pub use repack::q4_kx8_gemm_uses_acts_x4;

pub use repack::q4_kx8_interleave;

pub use repack::q5_kx8_gemm_uses_acts_x4;

pub use repack::q5_kx8_interleave;

pub use repack::q6_kx8_gemm_uses_acts_x4;

pub use repack::q6_kx8_interleave;

pub use repack::q8_0x4_gemm_uses_acts_x4;

pub use repack::q8_0x4_interleave;

pub use repack::Q8ActsX4;

pub use repack::Q8KActsX4;

pub use repack::Q4_0X4_BLOCK_BYTES;

pub use repack::Q4_0X4_GEMM_NC;

pub use repack::Q4_0X4_INTERLEAVE;

pub use repack::Q4_0X4_NROWS;

pub use repack::Q4_KX8_BLOCK_BYTES;

pub use repack::Q4_KX8_GEMM_NC;

pub use repack::Q4_KX8_NROWS;

pub use repack::Q5_KX8_BLOCK_BYTES;

pub use repack::Q5_KX8_GEMM_NC;

pub use repack::Q5_KX8_NROWS;

pub use repack::Q6_KX8_BLOCK_BYTES;

pub use repack::Q6_KX8_GEMM_NC;

pub use repack::Q6_KX8_NROWS;

pub use repack::Q8K_ACTS_X4_NC;

pub use repack::Q8_0X4_BLOCK_BYTES;

pub use repack::Q8_0X4_GEMM_NC;

pub use repack::Q8_0X4_INTERLEAVE;

pub use repack::Q8_0X4_NROWS;

Modules

iq_tables

Codebook/grid tables for the IQ1_S / IQ1_M / IQ2_XXS / IQ2_XS / IQ2_S / IQ3_XXS / IQ3_S quantization formats.

repack

Interleaved Q4_K / Q5_K × Q8_K and Q8_0 × Q8 GEMV (llama.cpp repack layouts).

Structs

Q8Activations

An activation vector quantized to signed 8-bit in 32-element blocks, each with its own f32 scale (d), so it can feed the integer vec_dot paths against Q8_0 weights. This mirrors llama.cpp’s quantize_row_q8_1 (minus the block sum, which is only needed for asymmetric weight formats): quantizing the shared activation once per matvec turns every weight-row dot into an int8×int8 → int32 reduction (vdotq_s32 / _mm256_maddubs-class ops) plus a single scale, which is what lets llama.cpp’s CPU matmul stay in integer SIMD.

Q8KActivations

ggml block_q8_K activations for K-quant int-dot (Q4_K/Q5_K/Q6_K). Super-blocks of 256 elements with 16-wide bsums for the min term.

Enums

QuantError

Constants

FP8_KV_BLOCK_BYTES

FP8_KV_GROUP

Metal FERROX_CTK=fp8 KV block: 32 elems → f16 scale + 32 E4M3-ish bytes. Codes are absmax-scaled int8 in [-127,127] (portable stand-in for E4M3).

IQ1_M_BLOCK_BYTES

IQ1_M: 32 low-index bytes + 16 qh bytes (3 high index bits + a sign-of-delta bit per 8-element group) + 8 scale bytes. 1.75 bpw. The only IQ format with no f16 scale field – see for_each_iq1_m.

IQ1_M_BLOCK_ELEMS

IQ1_S_BLOCK_BYTES

IQ1_S: d(f16) + 32 low-index bytes + 8 u16 (3 high index bits + 3 scale bits + sign-of-delta per 32-element group). 1.5625 bpw.

IQ1_S_BLOCK_ELEMS

IQ2_S_BLOCK_BYTES

IQ2_S: d(f16) + 32 low-index bytes + 32 literal sign bytes + 8 qh bytes (2 high index bits per group of 8) + 8 scale bytes. 2.5625 bpw.

IQ2_S_BLOCK_ELEMS

IQ2_XS_BLOCK_BYTES

IQ2_XS: d(f16) + 32 u16 codes (9-bit grid index + 7-bit sign index)

IQ2_XS_BLOCK_ELEMS

IQ2_XXS_BLOCK_BYTES

IQ2_XXS: d(f16) + 32 u16 codes (grid indices + packed scale/signs). 2.0625 bpw.

IQ2_XXS_BLOCK_ELEMS

IQ3_S_BLOCK_BYTES

IQ3_S: d(f16) + 64 low-index bytes + 8 qh bytes (one 9th index bit per grid code) + 32 literal sign bytes + 4 scale bytes (two 4-bit scales per pair of 32-element groups). 3.4375 bpw.

IQ3_S_BLOCK_ELEMS

IQ3_XXS_BLOCK_BYTES

IQ3_XXS: d(f16) + 64 grid-index bytes + 8 u32 scale/sign words. 3.0625 bpw.

IQ3_XXS_BLOCK_ELEMS

IQ4_NL_BLOCK_BYTES

IQ4_NL_BLOCK_ELEMS

IQ4_XS_BLOCK_BYTES

IQ4_XS_BLOCK_ELEMS

MXFP4_GGUF_BLOCK_BYTES

GGUF block-MXFP4 (ggml type tag 39): one 17-byte block = 1 E8M0 scale byte + 16 nibble bytes covering 32 elements, low nibble -> element j, high nibble -> element j+16. Same E2M1 codebook and E8M0 scale math as the Kimi safetensors two-buffer MXFP4 path above (dot_mxfp4_row_f32) – ggml expresses it as doubled-integer kvalues times a half scale (2^(e-128)), this module as true E2M1 values times the full 2^(e-127) scale; the products are identical across the whole E8M0 range including the e < 2 denormal patterns. Only the byte layout differs: interleaved 17-byte blocks in one stream here, two separate packed/scale tensors there.

MXFP4_GGUF_BLOCK_ELEMS

MXFP4_GROUP_SIZE

Elements per MXFP4 scale group (real, confirmed both from ggml’s QK_MXFP4 and directly from a real Kimi K3 shard’s own tensor shapes: *.weight_scale is in_dim/32 bytes, *.weight_packed is in_dim/2 bytes).

Q2_K_BLOCK_BYTES

Q2_K: a 256-element super-block, 16 sub-blocks of 16, each with its own 4-bit scale and 4-bit min packed one byte per sub-block (not Q4_K’s cross-byte 6-bit packing – a real, verified difference, not assumed), plus one shared f16 super-block scale d and f16 super-block min-scale dmin. Layout: scales(16) + qs(64) + d(2) + dmin(2) = 84 bytes – note d/dmin come after scales/qs, the opposite field order from every other K-quant format here, verified directly against real ggml-common.h/ggml-quants.c source (block_q2_K, dequantize_row_q2_K).

Q2_K_BLOCK_ELEMS

Q3_K_BLOCK_BYTES

Q3_K: a 256-element super-block, 16 sub-blocks of 16, each with its own signed 6-bit scale (packed via a byte-wise interleaving scheme across 12 bytes, verified against dequantize_row_q3_K’s real aux[] unpacking – see q3_k_unpack_scales’s doc comment), a 3-bit value per element (2 low bits from qs, 1 high bit from hmask, centered by -4 when the high bit is clear), scaled by one shared f16 d. Layout: hmask(32) + qs(64) + scales(12) + d(2) = 110 bytes.

Q3_K_BLOCK_ELEMS

Q4_0_BLOCK_BYTES

Q4_0: 32 packed 4-bit values (16 bytes) sharing one f16 scale. 18 bytes per block.

Q4_0_BLOCK_ELEMS

Q4_1_BLOCK_BYTES

Q4_1: like Q4_0 but asymmetric – an f16 scale d and an f16 min m (value = q*d + m, no -8 bias), 32 packed 4-bit values. Layout: d(2) + m(2) + qs(16) = 20 bytes. Verified against real ggml-common.h/ggml-quants.c source, not guessed.

Q4_1_BLOCK_ELEMS

Q4_K_BLOCK_BYTES

Q4_K: a 256-element super-block, split into 8 32-element sub-blocks, each with its own 6-bit scale and 6-bit min (packed into 12 bytes), plus one shared f16 scale-of-scales d and scale-of-mins dmin. Layout: d(2) + dmin(2) + scales(12) + qs(128) = 144 bytes.

Q4_K_BLOCK_ELEMS

Q5_0_BLOCK_BYTES

Q5_0: like Q4_0 (single f16 scale d, symmetric -16 bias) but each element gets a 5th bit from a 4-byte qh bitplane. Layout: d(2) + qh(4) + qs(16) = 22 bytes.

Q5_0_BLOCK_ELEMS

Q5_1_BLOCK_BYTES

Q5_1: Q5_0’s 5th-bit scheme combined with Q4_1’s asymmetric d+m (no bias subtraction). Layout: d(2) + m(2) + qh(4) + qs(16) = 24 bytes.

Q5_1_BLOCK_ELEMS

Q5_K_BLOCK_BYTES

Q5_K: the same 8-sub-blocks-of-32 / 6-bit-scale-and-min layout as Q4_K (same 12-byte packed scales, same unpacking), but each element gets a 5th bit from a separate 32-byte qh bitplane (one bit per element, 256 bits total) instead of Q4_K’s plain 4-bit nibble. Layout: d(2) + dmin(2) + scales(12) + qh(32) + qs(128) = 176 bytes.

Q5_K_BLOCK_ELEMS

Q5_K_GEMM_NC

How many activations one gemm_q5_k_q8_row / gemm_q6_k_q8_row keeps in flight. Amortizes weight-block scale/qh/qs loads over the batch (Phi-4 Q5_K qkv / Q6_K ffn_down) without full Kx8 repack.

Q6_K_BLOCK_BYTES

Q6_K: a 256-element super-block, split into 16 16-element sub-blocks each with its own signed 8-bit scale, plus one shared f16 super-block scale d. Layout: ql(128) + qh(64) + scales(16) + d(2) = 210 bytes.

Q6_K_BLOCK_ELEMS

Q6_K_GEMM_NC

Q8_0_BLOCK_BYTES

Q8_0: 32 int8 values sharing one f16 scale. 34 bytes per block.

Q8_0_BLOCK_ELEMS

Q8_1_BLOCK_BYTES

Q8_1: like Q8_0 (32 signed 8-bit values, one f16 scale d) plus an extra f16 field s that upstream ggml uses only as a precomputed per-block sum for its own fused SIMD dot-product kernels – not needed for correct dequantization, since y = qs*d is unaffected by it. Layout: d(2) + s(2) + qs(32) = 36 bytes.

Q8_1_BLOCK_ELEMS

TURBO4_KV_BLOCK_BYTES

TURBO4_KV_GROUP

Metal FERROX_CTK=turbo4 KV block: 32 elems → f16 scale + 16 nibble bytes.

Functions

dequant_bf16

BF16 isn’t a block-quantized format at all – it’s IEEE-754 binary32 truncated to its sign bit + 8 exponent bits + 7 mantissa bits (the upper 16 bits of an f32), so widening it back to f32 is an exact, lossless bit shift: f32::from_bits((bits as u32) << 16), zero- padding the low 16 mantissa bits rather than any real dequantization math. Included here anyway (rather than as a one-off in ferrox-models::loader) so every real element type ferrox recognizes has one obvious home.

dequant_f16

F16 (IEEE-754 binary16) widened to f32. Like dequant_bf16 this is a plain element type, not a block format: every f16 value is exactly representable in f32, so the widening is lossless. GgmlType::F16 is what llama-quantize --pure-free conversions and every *-f16.gguf carry, and it is also the dtype ggml uses for token_embd in some mixed checkpoints.

dequant_iq1_m

dequant_iq1_s

dequant_iq2_s

dequant_iq2_xs

dequant_iq2_xxs

dequant_iq3_s

dequant_iq3_xxs

dequant_iq4_nl

dequant_iq4_xs

dequant_mxfp4_gguf

dequant_mxfp4_row

Dequantizes one row of Kimi K3’s MXFP4-packed expert weights. Unlike every other kernel in this module, MXFP4 here is NOT a single interleaved byte stream – Kimi K3’s real safetensors checkpoint stores the packed 4-bit codes and the per-group E8M0 scales as two separate tensors (*.weight_packed, *.weight_scale; confirmed directly against a real shard header’s tensor shapes, not ggml’s own combined-block GGUF convention), so this takes both buffers directly rather than one combined block stream. packed is in_dim/2 bytes (2 nibble-packed E2M1 codes per byte, low-nibble-first-half / high-nibble-second-half within each 32-element group – same convention as this module’s other nibble-packed formats); scales is in_dim/MXFP4_GROUP_SIZE bytes (one E8M0 scale byte per group).

dequant_q2_k

Dequantize a Q2_K buffer into f32. Formula verified against real ggml-quants.c::dequantize_row_q2_K: 16 sub-blocks of 16 elements, each sub-block’s (scale, min) packed one byte per sub-block (sc & 0xF = 4-bit scale, sc >> 4 = 4-bit min – much simpler than Q4_K’s cross-byte 6-bit packing), value = d*scale*raw2bit - dmin*min, raw2bit in 0..=3 (2 bits per element from qs, 4 elements packed per byte).

dequant_q3_k

Dequantize a Q3_K buffer into f32. Formula verified against real ggml-quants.c::dequantize_row_q3_K: 16 sub-blocks of 16 elements, value = d_all*(scale-32)*(raw3bit-bias), raw3bit = 2 bits from qs plus 1 high bit from hmask (bit m, m sweeping all 8 bit positions across the whole block – hmask is indexed the same way regardless of which half of qs is active, only the bit tested changes), bias = 4 when the high bit is clear, 0 when set.

dequant_q4_0

Dequantize a Q4_0 buffer into f32. Each byte packs two 4-bit nibbles (low nibble = element i, high nibble = element i+16), each nibble biased by -8 before scaling, matching the public Q4_0 convention.

dequant_q4_1

Dequantize a Q4_1 buffer into f32. Formula verified against real ggml-quants.c::dequantize_row_q4_1: y = q*d + m, no bias subtraction (unlike Q4_0’s symmetric q-8).

dequant_q4_k

Dequantize a Q4_K buffer into f32. See the module doc comment and Q4_K_BLOCK_BYTES for the block layout.

dequant_q5_0

Dequantize a Q5_0 buffer into f32. Formula verified against real ggml-quants.c::dequantize_row_q5_0: symmetric, y = (q-16)*d where q is the 4-bit nibble with the 5th bit from qh ORed in.

dequant_q5_1

Dequantize a Q5_1 buffer into f32. Formula verified against real ggml-quants.c::dequantize_row_q5_1: Q5_0’s 5th-bit scheme, but asymmetric like Q4_1 (y = q*d + m, no -16 bias).

dequant_q5_k

Dequantize a Q5_K buffer into f32. See the module doc comment and Q5_K_BLOCK_BYTES for the block layout. Shares Q4_K’s scale/min packing (q4_k_scale_min) and 4-outer-iteration structure; the only difference is each nibble gets a 5th bit from qh, whose 32 bytes are reused across all 4 outer iterations at different bit positions (u1/u2, doubling by 4 each iteration) rather than being consumed sequentially the way qs is.

dequant_q6_k

Dequantize a Q6_K buffer into f32. See the module doc comment and Q6_K_BLOCK_BYTES for the block layout.

dequant_q8_0

Dequantize a Q8_0 buffer into f32.

dequant_q8_1

Dequantize a Q8_1 buffer into f32. Formula verified against real ggml-quants.c::dequantize_row_q8_1: identical to Q8_0 (y = q*d) – the extra s field (upstream: a precomputed per-block sum used only by ggml’s own fused SIMD dot kernels) doesn’t change the dequantized value and is intentionally unread here.

dot_iq1_m_f32

dot_iq1_m_f32_scalar

dot_iq1_s_f32

dot_iq1_s_f32_scalar

dot_iq2_s_f32

dot_iq2_s_f32_scalar

dot_iq2_xs_f32

IQ2_XS / IQ2_S / IQ3_S / IQ1_M dispatch: scalar only. These landed for coverage – before them, tags 17/21/22/29 fell to GgmlType::Other and the tensor could not be decoded at all, which silently ruled out 5 of the 16 published Unsloth UD-* variants. They deliberately match the state of their older siblings’ NEON/GPU story (none), rather than growing a vectorized path that no golden vector would then be able to distinguish from the scalar one.

dot_iq2_xs_f32_scalar

dot_iq2_xxs_f32

dot_iq2_xxs_f32_scalar

dot_iq3_s_f32

dot_iq3_s_f32_scalar

dot_iq3_xxs_f32

dot_iq3_xxs_f32_scalar

dot_iq4_nl_f32

Fused IQ4_NL dequant+dot, same math as dequant_iq4_nl. Dispatches to AVX2+FMA or NEON when available, same mechanism as dot_q4_0_f32.

dot_iq4_nl_f32_scalar

dot_iq4_xs_f32

Fused IQ4_XS dequant+dot, same math as dequant_iq4_xs. Dispatches to AVX2+FMA or NEON when available, same mechanism as dot_q4_0_f32.

dot_iq4_xs_f32_scalar

dot_mxfp4_gguf_f32

GGUF block-MXFP4 dispatch: scalar only so far (the two-buffer safetensors MXFP4 form has AVX2/NEON kernels above; this block form hasn’t needed one yet).

dot_mxfp4_gguf_f32_scalar

dot_mxfp4_row_f32

Fused MXFP4 dequant+dot, same math as dequant_mxfp4_row. Dispatches to AVX2+FMA or NEON when available (see simd_x86::dot_mxfp4_row_f32_avx2/ simd_aarch64::dot_mxfp4_row_f32_neon), same mechanism as dot_q4_0_f32 – this is the hot path for every routed expert’s FFN in a real Kimi K3 forward pass, so unlike Q4_0/Q8_0’s optional legacy-format status, keeping this scalar-only directly costs real inference speed.

dot_mxfp4_row_f32_scalar

dot_q2_k_f32

Fused Q2_K dequant+dot, same math as dequant_q2_k. Dispatches to AVX2+FMA or NEON when available, same mechanism as dot_q4_k_f32.

dot_q2_k_f32_scalar

dot_q3_k_f32

Fused Q3_K dequant+dot, same math as dequant_q3_k. Dispatches to AVX2+FMA or NEON when available, same mechanism as dot_q4_k_f32.

dot_q3_k_f32_scalar

dot_q4_0_f32

Same idea for Q4_0: fused dequant + dot, no intermediate f32 buffer. Dispatches to AVX2+FMA when available, same mechanism as dot_q8_0_f32.

dot_q4_0_f32_scalar

dot_q4_0_q8

Integer vec_dot of a Q4_0 weight row against pre-quantized Q8 activations (llama.cpp ggml_vec_dot_q4_0_q8_0). Opt-in via FERROX_CPU_INT_DOT for Q4_0 matvecs.

dot_q4_0_q8_2row

Two contiguous Q4_0 rows × one Q8 act (shared act loads). Faster than two dot_q4_0_q8 calls on Apple DotProd.

dot_q4_0_q8_scalar

dot_q4_1_f32

Fused Q4_1 dequant+dot, same math as dequant_q4_1. Dispatches to AVX2+FMA or NEON when available, same mechanism as dot_q4_0_f32.

dot_q4_1_f32_scalar

dot_q4_k_f32

Fused Q4_K dequant+dot: identical math to dequant_q4_k, but accumulated directly against x instead of materializing a dequantized row. Dispatches to SIMD when the host CPU supports it, same mechanism as dot_q8_0_f32.

dot_q4_k_f32_scalar

dot_q4_k_q8

Integer vec_dot of a Q4_K weight row against Q8KActivations (llama.cpp ggml_vec_dot_q4_K_q8_K). Opt-in via FERROX_CPU_INT_DOT.

dot_q4_k_q8_scalar

dot_q5_0_f32

Fused Q5_0 dequant+dot, same math as dequant_q5_0. Dispatches to AVX2+FMA or NEON when available, same mechanism as dot_q4_0_f32.

dot_q5_0_f32_scalar

dot_q5_1_f32

Fused Q5_1 dequant+dot, same math as dequant_q5_1. Dispatches to AVX2+FMA or NEON when available, same mechanism as dot_q4_0_f32.

dot_q5_1_f32_scalar

dot_q5_k_f32

Fused Q5_K dequant+dot: identical math to dequant_q5_k, but accumulated directly against x instead of materializing a dequantized row. Dispatches to SIMD when available, same mechanism as dot_q8_0_f32.

dot_q5_k_f32_scalar

dot_q5_k_q8

Integer vec_dot of a Q5_K weight row against Q8KActivations (llama.cpp ggml_vec_dot_q5_K_q8_K). Opt-in via FERROX_CPU_INT_DOT.

dot_q5_k_q8_scalar

dot_q6_k_f32

Fused Q6_K dequant+dot: identical math to dequant_q6_k, but accumulated directly against x instead of materializing a dequantized row. Dispatches to SIMD when available, same mechanism as dot_q8_0_f32.

dot_q6_k_f32_scalar

dot_q6_k_q8

Integer vec_dot of a Q6_K weight row against Q8KActivations (llama.cpp ggml_vec_dot_q6_K_q8_K). Opt-in via FERROX_CPU_INT_DOT.

dot_q6_k_q8_scalar

dot_q8_0_f32

Fused dot product between one Q8_0-quantized row (stored as raw block bytes) and an f32 activation vector, without ever materializing a dequantized f32 copy of the row. This is the memory-bandwidth-saving trick llama.cpp’s quantized matmul kernels rely on: for large weight matrices, bandwidth (not FLOPs) dominates inference cost, and Q8_0 moves 4x fewer bytes than a dequant-then- matmul approach that expands every weight to f32 up front.

dot_q8_0_f32_scalar

dot_q8_0_q8

Integer vec_dot of a Q8_0 weight row against pre-quantized Q8 activations: Σ_blocks d_w * d_a * Σ_i (q_w · q_a). Dispatches to a NEON dotprod / AVX2 kernel when available, else the scalar loop. Numerically ≈ dot_q8_0_f32 up to activation-quant error.

dot_q8_0_q8_scalar

dot_q8_1_f32

Fused Q8_1 dequant+dot, same math as dequant_q8_1. Dispatches to AVX2+FMA or NEON when available – mathematically identical to Q8_0 (y = q*d), so the SIMD kernels are Q8_0’s kernels with the quantized bytes read from offset 4 instead of offset 2 (Q8_1’s block has an extra 2-byte field between d and the int8 values).

dot_q8_1_f32_scalar

gemm_q5_k_q8_row

One Q5_K weight row × acts.len() Q8_K activations → out[j].

gemm_q5_k_q8_row_scalar

gemm_q6_k_q8_row

One Q6_K weight row × acts.len() Q8_K activations → out[j].

gemm_q6_k_q8_row_scalar

pack_fp8_kv_blocks

Pack f32 into Metal fp8-style KV blocks (scaled int8, Q8_0-compatible layout).

pack_turbo4_kv_blocks

Pack f32 into Metal turbo4 KV blocks (no WHT).

quantize_activations_q8

Quantize an activation row to Q8Activations (32-element blocks, ggml quantize_row_q8_0 rounding: d = amax/127, q = round(x/d)). x.len() must be a multiple of 32.

quantize_activations_q8_k

Quantize activations to ggml Q8_K (256-elem super-blocks). Positive scale convention (d = amax/127) matching our Q8_0 path; bsums enable the Q4_K min correction without re-scanning q.

quantize_q8_0

Quantize an f32 slice into Q8_0 blocks (used by test fixtures and by the CPU reference “quantize activations for a symmetric int8 matmul” path). Not performance tuned; correctness-first reference only.

unpack_fp8_kv_blocks

Unpack pack_fp8_kv_blocks.

unpack_turbo4_kv_blocks

Unpack pack_turbo4_kv_blocks.