Rust in Everything
基于 candle-core 0.11.0

写一个迷你 Q8_0 量化

主线位置：把推理主线的前半段自己走一遍。读 GGUF 头部、按偏移量取出一个张量的原始字节（对应 Content::read 和 Content::tensor）→ 把字节解释成 Q8_0 的块、反量化成 f32（对应 QTensor::dequantize）→ 量化一行输入、和权重的每一行做量化点积（对应 QMatMul 在 CPU 上的内层 vec_dot）。每一步都拿 candle 的结果逐字节、逐位核对，candle 多出来的东西放在最后对照。

flowchart TB
  gguf["GGUF 文件"] --> hdr["Gguf read 头部"]
  hdr --> info["TensorInfo"]
  info --> raw["tensor_bytes"]
  raw --> blk["BlockQ8_0 块"]
  xs["f32 向量"] --> qz["quantize"]
  qz --> blk
  blk --> deq["dequantize"]
  blk --> dot["vec_dot"]
  deq --> cmp["和 candle 核对"]
  dot --> cmp
  qz --> cmp

mini 是一个只依赖 half（f16）的库 mini-quant，外加一个列文件的 main.rs。candle 只出现在 [dev-dependencies] 里，只有 tests/ 和 examples/ 用到：

[dependencies]
half = "2"

# 只有 tests/ 和 examples/ 用 candle 对照；库和 src/main.rs 不依赖它
[dev-dependencies]
candle-core = { version = "0.11", features = ["metal", "accelerate"] }

平台 M2 Max、macOS 27、rustc 1.99.0。除了「头部为什么要 35 ms」一节是 release，其余都是 debug 构建。库里没有 unwrap / expect，出错一律返回 Result，测试里随便用。模型文件是 bartowski 的 SmolLM2-135M-Instruct-Q8_0.gguf（144811360 字节），只读。

范围

文件做什么代码行（不含空行和注释）
src/q8_0.rsBlockQ8_0、量化、反量化、字节布局、vec_dot76
src/gguf.rsGGUF v2 / v3 头部：元数据（所有值类型，大数组只留前 3 个）、张量目录、对齐、数据区起点；按偏移读张量230
src/lib.rsError、导出19
src/main.rs列出一个 GGUF 文件的元数据和张量，反量化一个张量46
tests/candle_compare.rs五个对照测试179
tests/scale_choice.rs动脑筋用：scale 的倒数取 f32 还是 f1642
examples/crafted.rs手工拼的坏头部，分别交给 mini 和 candle58
examples/header_time.rs读头部计时，数 candle 的 seek 次数50
examples/metal_mv.rs同一次矩阵乘向量：mini、candle CPU、candle Metal48

库一共 325 行代码，加上空行和注释 393 行。

Q8_0 的块

Q8_0 把张量按行优先展平，每 32 个 f32 一块。每块存一个 f16 的 scale d 和 32 个 i8，还原时 x ≈ q * d。一块 34 字节，每个权重 8.5 bit。candle 的定义：

#[derive(Debug, Clone, PartialEq)]
#[repr(C)]
pub struct BlockQ8_0 {
    pub(crate) d: f16,
    pub(crate) qs: [i8; QK8_0],
}
const _: () = assert!(std::mem::size_of::<BlockQ8_0>() == 34);

（BlockQ8_0）#[repr(C)] 加上编译期断言，保证内存布局就是文件布局：f16 对齐到 2，后面的 i8 数组不需要填充。candle 读 GGUF 时把一段字节直接当成 [BlockQ8_0]，不逐块解析。mini 不用 unsafe，用两个函数在块和字节之间换：

#[derive(Debug, Clone, Copy, PartialEq)]
pub struct BlockQ8_0 {
  /// scale：块内绝对值最大的元素除以 127。
  pub d: f16,
  pub qs: [i8; QK8_0],
}

impl BlockQ8_0 {
  /// 文件里的布局（ggml 的 `block_q8_0`）：f16 小端，然后 32 个 i8。
  pub fn to_bytes(&self) -> [u8; BLOCK_BYTES] {
    let mut out = [0u8; BLOCK_BYTES];
    out[..2].copy_from_slice(&self.d.to_le_bytes());
    for (b, &q) in out[2..].iter_mut().zip(&self.qs) {
      *b = q as u8;
    }
    out
  }
  // ... from_bytes 反过来
}

整段字节切块用 as_chunks::<BLOCK_BYTES>()：它返回 &[[u8; 34]] 和剩下的零头，长度不是 34 的倍数就报错。拿到的是定长数组，from_bytes(&[u8; 34]) 里按下标取 bytes[0]、bytes[1] 不会越界，类型已经保证了长度。

量化和反量化

  pub fn quantize(xs: &[f32; QK8_0]) -> Self {
    let amax = xs.iter().fold(0f32, |m, x| m.max(x.abs()));
    let d = amax / 127.0;
    // 倒数用 f32 的 d 算，不用存进去的 f16：和 llama.cpp、candle 一致
    let id = if d != 0.0 { 1.0 / d } else { 0.0 };
    let mut qs = [0i8; QK8_0];
    for (q, &x) in qs.iter_mut().zip(xs) {
      // round 远离 0 取整；as 在越界时饱和，|x * id| 不会超过 127.5
      *q = (x * id).round() as i8;
    }
    Self { d: f16::from_f32(d), qs }
  }

  pub fn dequantize(&self, out: &mut [f32; QK8_0]) {
    let d = self.d.to_f32();
    for (y, &q) in out.iter_mut().zip(&self.qs) {
      *y = q as f32 * d;
    }
  }

这是 candle 的 from_float 和 to_float 的逐行翻译（candle 写 ((1 << 7) - 1) as f32，就是 127；注释写明照的是 ggml 的 quantize_row_q8_0）。三处细节决定了输出能不能逐字节一致：

  • scale 是对称的。 只看绝对值最大的元素，块里最大的那个数量化后正好是 ±127；Q8_0 没有零点（zero point），-128 永远用不到。
  • 量化用 f32 的 d 求倒数，存的是 f16 的 d。 量化时 x * id 里的 id 来自没舍入的 f32，反量化时乘的却是舍入到 f16 之后的 d。两者差一点，后面的动脑筋就问这一点。
  • round 远离 0 取整。 2.5 变成 3，不是银行家舍入的 2。

整个向量用 as_chunks::<32>() 切块，长度不是 32 的倍数就返回 Err。candle 的 from_float 在这里只有 debug_assert!，长度检查在更外层的 QTensor::quantize（最后一维必须是块大小的倍数）。

测试数据是 4096 块固定种子的随机数，再接 6 块特殊的。quantize_bytes_match_candle 把 mini 的字节和 QTensor::quantize(.., GgmlDType::Q8_0)?.data()? 比，再把两边的反量化结果按 to_bits 比：

test quantize_bytes_match_candle ... 131264 values -> 4102 blocks, 139468 bytes (candle 139468 bytes)
bytes that differ: 0
dequantized bit-identical: true
zeros d = 0          qs[..5] = [0, 0, 0, 0, 0]        back[..5] = [0.0, 0.0, 0.0, 0.0, 0.0]
tiny  d = 0          qs[..5] = [0, 4, 8, 12, 16]      back[..5] = [0.0, 0.0, 0.0, 0.0, 0.0]
huge  d = inf        qs[..5] = [0, 4, 8, 12, 16]      back[..5] = [NaN, inf, inf, inf, inf]
inf   d = inf        qs[..5] = [0, 0, 0, 0, 0]        back[..5] = [NaN, NaN, NaN, NaN, NaN]
nan   d = 0.24414063 qs[..5] = [0, 4, 8, 0, 16]       back[..5] = [0.0, 0.9765625, 1.953125, 0.0, 3.90625]
half  d = 1          qs[..5] = [127, 2, 3, 4, 5]      back[..5] = [127.0, 2.0, 3.0, 4.0, 5.0]
ok

字节一个不差，特殊块在两边的行为也一样（反量化逐位相同包括 NaN 和 inf）。特殊块各自说明一件事：

块内容结果
tiny1e-9 * if32 的 d 约 2.4e-10，不是 0，qs 照常算出来；f16 最小的次正规数约 6e-8，存进去的 d 变成 0，整块还原成 0
huge1e7 * id 约 2.4e6，超过 f16 的最大值 65504，存成 inf；0 * inf 得 NaN
inf第 3 个是 infd = inf，id = 0，其余元素全量化成 0，整块还原成 NaN
nan第 3 个是 NaNf32::max 忽略 NaN，scale 照其他元素算；NaN as i8 是 0，NaN 悄悄变成 0
half127 和 i + 0.5d = 1，1.5 → 2、2.5 → 3：远离 0 取整

也就是说，Q8_0 对块内的数值范围有要求：最大值不能超过 127 × 65504（约 832 万），整块都很小时会整块归零，一个 inf 毁掉整块。正常的权重不会碰到这些；量化前没有检查，candle 也没有。

❗IMPORTANT
动脑筋：把 quantize 里的倒数换成存进去的 f16 scale 的倒数，即 1.0 / f16::from_f32(d).to_f32()，重跑 quantize_bytes_match_candle。139468 个字节里会有多少个和 candle 不同？反量化的均方误差变大还是变小？答案在页末，平台同上，debug 构建。

量化点积

/// 两个 Q8_0 向量的点积：块内 i8 x i8 累加成 i32，每块乘一次两个 scale。
pub fn vec_dot(xs: &[BlockQ8_0], ys: &[BlockQ8_0]) -> Result<f32> {
  if xs.len() != ys.len() {
    return err(format!("vec_dot: {} blocks vs {} blocks", xs.len(), ys.len()));
  }
  let mut sum = 0f32;
  for (x, y) in xs.iter().zip(ys) {
    let sum_i: i32 = x.qs.iter().zip(&y.qs).map(|(&a, &b)| a as i32 * b as i32).sum();
    sum += sum_i as f32 * x.d.to_f32() * y.d.to_f32();
  }
  Ok(sum)
}

量化点积省的是乘法的位宽：一块 32 次乘加全在整数上做，每块只有一次浮点乘法。32 个 i8 * i8 最大 32 × 128 × 128 = 524288，i32 装得下，块内求和没有舍入误差。

candle 的 trait GgmlType 给每种块两个点积：vec_dot_unopt 是标量的通用实现，和上面这段一模一样；vec_dot 按 target_feature 选 avx2、neon 或 simd128，都没有才退回 vec_dot_unopt。aarch64 默认开着 neon，所以 M2 上走的是 neon::vec_dot_q8_0_q8_0：

            let p0 = vdotq_s32(x0_0, y0_0);
            let p1 = vdotq_s32(x0_1, y0_1);

            sumv0 = vmlaq_n_f32(
                sumv0,
                vcvtq_f32_s32(vaddq_s32(p0, p1)),
                x0.d.to_f32() * y0.d.to_f32(),
            );
        }
        vaddvq_f32(sumv0)

整数部分（vdotq_s32，有 dotprod 时是一条 sdot 指令）结果和标量完全相同；不同的是浮点累加：neon 版把每块的整数和留在 4 个通道里分别乘 dx * dy、分别累加，最后 vaddvq_f32 才把 4 个通道加起来；标量版每块先把 32 个乘积加成一个数，再按 (sum_i * dx) * dy 的顺序乘。vec_dot_matches_candle 用两个 4096 维的随机向量：

test vec_dot_matches_candle ... f64 dot of the f32 inputs  -15.479279577534399
mini vec_dot               -15.470183
candle vec_dot_unopt       -15.470183
candle vec_dot (neon)      -15.470185
mini - neon = 1.9073486e-6
ok

mini 和 vec_dot_unopt 逐位相同（测试断言 to_bits 相等），和 neon 版差最后两位。三者和「f32 输入、f64 累加」的精确点积都差 0.009，这是两个向量各量化一次带来的误差，比累加顺序的差别大三个数量级。

拼成一次矩阵乘向量

QMatMul 在 CPU 上就是一圈 vec_dot：k_quants::matmul 先把输入的每一行用 from_float 量化成权重的 VecDotType（Q8_0 权重的 VecDotType 就是 Q8_0），再让 BarrierPool 的线程按列分工，每次用 vec_dot_4 算四个输出（线程池见 CPU 后端）。mini 照这个流程算 blk.0.attn_q.weight（[576, 576]）乘一行随机输入：

  // y = x Wᵀ：先把 x 量化成 Q8_0，再和 W 的每一行做 vec_dot
  let x = lcg(9, k);
  let qx = q8_0::quantize(&x).unwrap();
  let ours: Vec<f32> = w.chunks(k / QK8_0).map(|row| q8_0::vec_dot(row, &qx).unwrap()).collect();

test one_row_matmul_matches_qmatmul ... blk.0.attn_q.weight [576, 576] x [576]: 135/576 outputs bit-identical, max |mini - QMatMul| = 3.8146973e-6
first 4  mini [-17.927952, 5.6350007, -12.815998, -6.592821]
first 4  QMatMul [-17.927954, 5.6350007, -12.815999, -6.592821]
ok

576 个输出里 135 个逐位相同，其余差在最后一两位，原因同上：vec_dot_4_q8_0_q8_0 和 vec_dot 一样按通道累加。

Metal 上的同一个 QMatMul 不走这条路。只有一行输入时它调 kernel_mul_mv_q8_0_f32，输入保持 f32，不量化：

        for (int row = 0; row < nr; row++) {
            device const int8_t * qs = x[ib+row*nb].qs + NB_Q8_0*il;
            float sumq = 0.f;
            for (int iq = 0; iq < NB_Q8_0; ++iq) {
                sumq += qs[iq] * yl[iq];
            }
            sumf[row] += sumq*x[ib+row*nb].d;
        }

（kernel_mul_mv_q8_0_f32_impl，yl 是 f32 的输入。）examples/metal_mv.rs 把三种算法和「反量化后的 W 乘 f32 的 x、f64 累加」比：

max |y - reference| over 576 outputs
  mini (x as Q8_0)     8.475e-2
  candle CPU QMatMul   8.475e-2
  candle Metal QMatMul 1.791e-6
y[0]: reference -17.933317  mini -17.927952  cpu -17.927954  metal -17.933317

CPU 的误差主要来自把输入量化成 Q8_0 这一步，Metal 没有这一步，误差小四个数量级。同一个 GGUF、同一行输入，CPU 和 Metal 的 QMatMul 输出在第三位有效数字上就不同，生成的 token 也就可能不同。哪种输入走哪个内核，见 QTensor 与 QMatMul。

📝NOTE
冷知识：vec_dot 以前返回 Result<f32>

PR #3109（「Quantization use debug_assert in hot paths」，首次进入 0.9.2-alpha.1，第一个正式版是 0.10.0）之前，GgmlType 的 vec_dot、vec_dot_unopt、from_float、to_float 都返回 Result，开头先检查 n 是不是 32 的倍数，不是就 bail!。k_quants::matmul 里每个输出元素都是 T::vec_dot(k, rhs_col, lhs_row).map(|value| *dst = value)，经 rayon 收集成一个 Result<Vec<_>> 再 ?。提交说明里的一条是「Use debug_assert instead of if-else + bail! in hot paths」，另一条是「Add compile time verification of block sizes being equal to vec dot type block sizes」：检查改成 debug_assert!，release 构建里不再执行，返回值变成裸的 f32，matmul 里直接 *dst = T::vec_dot(..)。mini 的 vec_dot 还返回 Result，但只在入口比一次两个切片的块数，每块的循环里没有检查。

读 GGUF 头部

GGUF 是 llama.cpp 用的单文件格式，所有整数小端：

部分内容
魔数4 字节 GGUF
版本u32，SmolLM2 这个文件是 3
两个计数张量个数、元数据个数，v2 起都是 u64（v1 是 u32）
元数据每条：u64 长度的键 + u32 类型码 + 值；类型码 0–12，9 是数组：u32 元素类型 + u64 个数 + 元素
张量目录每个张量：名字 + u32 维数 + 每维 u64 + u32 ggml 类型码 + u64 偏移
填充补到 general.alignment 的倍数，没有这个键时是 32
数据区张量的原始字节，偏移从数据区起点算

mini 的 Reader 包一个 Read，自己记读了多少字节，所以读头部不需要 Seek：

/// 记着读到哪儿了，这样头部只需要 `Read`，不需要 `Seek`。
struct Reader<R> {
  inner: R,
  pos: u64,
}

impl<R: Read> Reader<R> {
  fn bytes<const N: usize>(&mut self) -> Result<[u8; N]> {
    let mut buf = [0u8; N];
    self.inner.read_exact(&mut buf)?;
    self.pos += N as u64;
    Ok(buf)
  }

  /// u64 长度加 UTF-8 字节。不按声明的长度预先分配：长度是假的，读到文件尾就报错。
  fn string(&mut self) -> Result<String> {
    let len = self.u64()?;
    let mut buf = Vec::new();
    let got = (&mut self.inner).take(len).read_to_end(&mut buf)? as u64;
    self.pos += got;
    if got != len {
      return err(format!("string of {len} bytes at offset {}, file ends after {got}", self.pos - got));
    }
    String::from_utf8(buf).map_err(|e| Error(format!("string is not utf-8: {e}")))
  }
  // ... u32、u64、value
}

值的 13 种类型在一个 match 里。数组逐个读完，只留前 3 个：

      9 => {
        let elem_type = self.u32()?;
        let len = self.u64()?;
        let mut head = Vec::new();
        for i in 0..len {
          // 嵌套数组靠递归读；深度没有上限
          let v = self.value(elem_type)?;
          if (i as usize) < ARRAY_HEAD {
            head.push(v);
          }
        }
        Value::Array { elem_type, len, head }
      }

SmolLM2 的元数据里有三个大数组：49152 个 token 字符串、49152 个 token 类型、48900 条 merge 规则，这是整份分词器。candle 的 Value::read 把它们全部读成 Value::Array(Vec<Value>) 留在 Content 里；mini 读过就扔，只记个数。字符串的处理也不同：candle 的 read_string 去掉末尾的 \0、用 from_utf8_lossy 容忍非法 UTF-8，mini 遇到非法 UTF-8 直接报错。

张量目录之后是对齐：

    let alignment = match metadata.iter().find(|(k, _)| k == "general.alignment") {
      None => DEFAULT_ALIGNMENT,
      Some((_, Value::U32(a))) if *a != 0 => *a as u64,
      Some((_, v)) => return err(format!("general.alignment must be a non-zero u32, got {v:?}")),
    };
    let data_offset = r.pos.div_ceil(alignment) * alignment;

candle 在 Content::read 里（第 551 行）接受六种整数类型的 general.alignment，其余类型悄悄退回 32，但没有排除 0。后面「坏文件」一节会用到这一点。

每维的长度读出来以后要倒过来：

      // 文件里先写变化最快的那一维（ggml 的 ne[0]），倒过来才是行优先的形状
      dims.reverse();
📝NOTE
没有蠢问题

问： token_embd.weight 在文件里写的是 [576, 49152]，倒过来变成 [49152, 576]。数据本身没动，只倒形状，不会把矩阵弄成转置的吗？

答： 不会，两种写法描述的是同一块内存。ggml 的 ne[0] 是变化最快（相邻元素间隔最小）的那一维，ne[1] 次之，所以文件里的 [576, 49152] 意思是「每 576 个元素一行，共 49152 行」。candle 和 mini 的形状按行优先写，变化最快的放在最后，同一个意思写成 [49152, 576]。数据从不重排，倒的只是记法。Q8_0 的块沿着变化最快的那一维切，所以 candle 的 check_shape 要求的是最后一维（576）是 32 的倍数。

列出 SmolLM2 的张量

cargo run --offline（debug），节选：

version 3  kv 37  tensors 272
alignment 32  data_offset 1785952  header read in 38.427708ms
--- metadata
general.architecture                     "llama"
...
llama.block_count                        30u32
llama.context_length                     8192u32
llama.embedding_length                   576u32
llama.feed_forward_length                1536u32
llama.attention.head_count               9u32
llama.attention.head_count_kv            3u32
llama.rope.freq_base                     100000f32
llama.attention.layer_norm_rms_epsilon   0.00001f32
general.file_type                        7u32
...
tokenizer.ggml.tokens                    ["<|endoftext|>", "<|im_start|>", "<|im_end|>", ...] (len 49152)
tokenizer.ggml.token_type                [3i32, 3i32, 3i32, ...] (len 49152)
tokenizer.ggml.merges                    ["Ġ t", "Ġ a", "i n", ...] (len 48900)
...
tokenizer.chat_template                  "{% for message in messages %}{% if loop.first and messages[0"... (368 bytes)
general.quantization_version             2u32
...
--- tensors
token_embd.weight            Q8_0  [49152, 576]   offset          0   30081024 B
blk.0.attn_norm.weight       F32   [576]          offset   30081024       2304 B
blk.0.ffn_down.weight        Q8_0  [576, 1536]    offset   30083328     940032 B
blk.0.ffn_gate.weight        Q8_0  [1536, 576]    offset   31023360     940032 B
blk.0.ffn_up.weight          Q8_0  [1536, 576]    offset   31963392     940032 B
blk.0.ffn_norm.weight        F32   [576]          offset   32903424       2304 B
blk.0.attn_k.weight          Q8_0  [192, 576]     offset   32905728     117504 B
blk.0.attn_output.weight     Q8_0  [576, 576]     offset   33023232     352512 B
blk.0.attn_q.weight          Q8_0  [576, 576]     offset   33375744     352512 B
blk.0.attn_v.weight          Q8_0  [192, 576]     offset   33728256     117504 B
...
output_norm.weight           F32   [576]          offset  143023104       2304 B
F32     61 tensors      140544 B
Q8_0   211 tensors   142884864 B
--- blk.0.attn_q.weight [576, 576]: 331776 values, max |x| 5.2516174
first 8: [-0.073791504, 0.16603088, -0.10146332, -0.7286911, 1.1714401, -0.119911194, 0.08301544, 0.6825714]

几个数可以互相核对：

  • 头部 1785952 字节，加上最后一个张量的偏移和大小（143023104 + 2304），正好是文件大小 144811360。头部里绝大部分是分词器。
  • token_embd.weight 有 49152 × 576 = 28311552 个元素，除以 32 再乘 34，得 30081024 字节。
  • 61 个 F32 是 30 层各两个 RmsNorm 的权重加 output_norm.weight；211 个 Q8_0 是 30 层各 7 个投影加 embedding。没有 output.weight：输出层和 embedding 共用权重，quantized_llama 的处理见生成一个 token。
  • 张量按文件里的顺序列出。candle 的 Content 把元数据和张量目录都放进 HashMap，按名字查很方便，原来的顺序就没了。

和 candle 对照

tests/candle_compare.rs 一共五个测试，上面已经出现了三个，剩下两个核对头部和整个文件的反量化。cargo test --offline --test candle_compare -- --nocapture --test-threads=1（debug，节选）：

running 5 tests
test every_q8_0_tensor_dequantizes_like_candle ... 211 Q8_0 tensors, 134479872 values: bit-identical to QTensor::dequantize
ok
test gguf_header_matches_candle ... 37 metadata entries, 272 tensors, data_offset 1785952: same as candle
ok
...
test result: ok. 5 passed; 0 failed; 0 ignored; 0 measured; 0 filtered out; finished in 7.08s
  • gguf_header_matches_candle：37 条元数据逐条比（标量比值，浮点比 to_bits；数组比长度和前 3 个元素），272 个张量比形状、偏移和字节数，数据区起点相同。
  • every_q8_0_tensor_dequantizes_like_candle：211 个 Q8_0 张量、1.34 亿个值，mini 的反量化和 content.tensor(..)?.dequantize(&Device::Cpu)? 逐位相同。反量化只有一次乘法 q as f32 * d，没有累加，没有可以不同的地方。

反过来验证一下测试抓得住错：把 quantize 的倒数改成取自 f16 的 d，quantize_bytes_match_candle 立刻红（见动脑筋答案）；同一次运行里 every_q8_0_tensor_dequantizes_like_candle 仍然通过，因为反量化不经过 quantize。

坏文件

GGUF 常常是从网上下载的，头部里的每个数都可能是假的。examples/crafted.rs 手工拼了四个坏头部，分别交给 mini 和 candle（debug）：

--- alignment 0 (57 bytes)
mini   Err(Error("general.alignment must be a non-zero u32, got U32(0)"))

thread 'main' (3812439) panicked at ~/.cargo/registry/src/index.crates.io-1949cf8c6b5b557f/candle-core-0.11.0/src/quantized/gguf_file.rs:560:43:
attempt to divide by zero
note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace
candle panicked
--- string len 2^40 (64 bytes)
mini   Err(Error("string of 1099511627776 bytes at offset 32, file ends after 32"))
candle Err("gguf: string length 1099511627776 exceeds max 1073741824")
--- 2^40 tensors (24 bytes)
mini   Err(Error("io: failed to fill whole buffer"))
candle Err("gguf: tensor_count 1099511627776 exceeds max 1073741824")
--- nested arrays (1249 bytes)
mini   Ok(1)
candle Err("gguf: value nesting depth exceeds max 64")
  • general.alignment = 0：candle 在第 560 行 position.div_ceil(alignment) 除以 0，panic，release 构建同样 panic。Content::read 返回 Result，这个输入却让它 panic；本地仓库里 0.11.0 之后的 main（12bb5d1f）这一行没变。
  • 假的字符串长度：candle 先和上限 2^30 比，再和文件剩下的字节数比，都通过了才 vec![0u8; len]。mini 没有上限，靠的是不预先分配：take(len).read_to_end 按实际读到的字节增长，文件只剩 32 字节，就只分配 32 字节，然后报错。
  • 假的张量个数：同理，candle 在读之前就拒绝；mini 不 with_capacity(count)，读第一个张量名就碰到文件尾。错误信息差一些。
  • 一层套一层的数组：candle 限制深度 64，mini 没有限制，100 层照样读完。套 100 万层（12 MB 的头部），cargo run --example crafted -- deep：
mini, 1_000_000 nested arrays (12000049 bytes)

thread 'main' (3812461) has overflowed its stack
fatal runtime error: stack overflow, aborting

这是递归解析器的老问题，进程直接中止，Result 和 catch_unwind 都接不住。candle 的这些上限是 PR #3556 和 PR #3585 加的，都首次进入 0.11.0，详见 GGUF 文件。

头部为什么要 35 ms

生成一个 token 测过：candle 直接对 File 读头部要 130 ms，套一层 BufReader 降到 31 ms。mini 的 debug 构建在上面的输出里是 38 ms（几次运行在 35–38 ms 之间）。examples/header_time.rs 在 release 下把两边各读 5 次，再把整个文件读进内存、用 Cursor 各读 3 次，最后数一数 candle 往 File 上调了多少次 seek：

run 0: mini    10.16ms (37 kv)   candle    35.15ms (37 kv)
run 1: mini     9.74ms (37 kv)   candle    35.08ms (37 kv)
run 2: mini     9.59ms (37 kv)   candle    34.90ms (37 kv)
run 3: mini     9.57ms (37 kv)   candle    34.60ms (37 kv)
run 4: mini     9.43ms (37 kv)   candle    34.54ms (37 kv)
cursor 0: mini     9.37ms   candle     6.48ms
cursor 1: mini     9.42ms   candle     6.31ms
cursor 2: mini     9.47ms   candle     6.27ms
candle Content::read: 98384 seeks on the File

单看解析，candle 更快（Cursor 上 6.3 ms 对 9.4 ms；mini 慢在哪里本页没有细查）。慢的是 BufReader<File> 这一层：candle 在每个字符串和数组之前调一次 remaining_bytes，它调 reader.stream_position()。BufReader 的 stream_position 要问底下的 File 当前位置，也就是一次 lseek 系统调用，缓冲区再满也省不掉。98384 次一次不差地对得上：每个字符串之前一次（两个大数组里 49152 + 48900 个，37 个键，272 个张量名，13 个字符串值，general.languages 里 1 个），每个数组之前一次（4 个），再加 Content::read 开头和结尾的 5 次。就是这些系统调用把 6 ms 的解析拉长到 35 ms。remaining_bytes 是 PR #3556 为了「长度不能超过文件剩下的字节」加的。这个 PR 自己已经发现过一次代价：它的第二个提交「Avoid re-seeking on every GGUF length check」把每次检查从「跳到文件末尾再跳回来」改成只在开头取一次文件大小，说明里写原来的做法「roughly doubled load time」；剩下的这一次 stream_position 在 BufReader<File> 上仍然是系统调用。mini 自己数 pos，同一个检查不需要问操作系统。

candle 有而这里没有的

candle 有mini 的做法讲它的页
SIMD 点积：vec_dot 按 target_feature 选 neon / avx2 / simd128，vec_dot_4 一次算四列，dotprod 时用 sdot只有标量版，结果和 vec_dot_unopt 逐位相同量化格式
GgmlDType 的 15 种（含 F32、F16、BF16）：Q4_0 到 Q8_1，还有 256 个元素一块的 K-quants（Q2K 到 Q8K），各有量化、反量化、点积；部分类型还有 imatrix 加权量化（from_float_imatrix）只有 Q8_0；ggml_type_info 只用来算其他类型张量的字节数量化格式
QMatMul：from_arc 按 dtype 和 CANDLE_DEQUANTIZE_ALL / CANDLE_DEQUANTIZE_ALL_F16 选「保持量化」「反量化成 f32」「反量化成 f16」三种形态；CPU 上量化输入后用 BarrierPool 多线程算一个 for 循环一行一行 vec_dot，单线程QTensor 与 QMatMul、CPU 后端
Metal 内核：quantized.metal 7741 行，矩阵乘向量（mul_mv，输入保持 f32）、矩阵乘矩阵（mul_mm）、按行取（get_rows）；Metal 上的量化和反量化则拷回 CPU 做只有 CPUQTensor 与 QMatMul、Metal 后端
QTensor：带形状、放在某个设备上的量化张量，能 embedding、能当自定义算子接进 TensorVec<BlockQ8_0> 加一个 dimsQTensor 与 QMatMul
分配上限：字符串和数组 2^30、张量维数 4、嵌套深度 64，读之前和文件剩余字节比不按声明的长度预先分配；没有深度上限，深层嵌套会栈溢出GGUF 文件
宽容：GGUF v1、反过来的魔数、字符串末尾的 \0、非法 UTF-8、六种整数类型的 general.alignment只认 v2 / v3、GGUF、合法 UTF-8、u32 对齐，其余报错GGUF 文件
写 GGUF：gguf_file::write只读GGUF 文件

mini 的反量化和 candle 逐位相同、量化逐字节相同，做不到的是速度、种类和设备。burn 走的是另一条路：后端的量化算子默认先反量化成浮点、调浮点算子，再决定要不要量化回去（Backend trait），没有 GGUF 这一套按块的整数点积。

动脑筋答案

把 let id = ... 一行换成 1.0 / f16::from_f32(d).to_f32()（0 时仍取 0），cargo test --offline --test candle_compare quantize -- --nocapture（节选，断言打印的整个字节数组已截掉）：

131264 values -> 4102 blocks, 139468 bytes (candle 139468 bytes)
bytes that differ: 905

thread 'quantize_bytes_match_candle' (3829290) panicked at tests/candle_compare.rs:61:3:
assertion `left == right` failed
...
test quantize_bytes_match_candle ... FAILED

tests/scale_choice.rs 把两种做法放在一起，只看 4096 块随机数据：

4096 blocks; blocks that differ 745, qs that differ 842
mean squared error  f32 id 1.8196e-4   f16 id 1.8178e-4
largest |qs difference| Some(1)
  • 139468 个字节里 905 个不同：随机部分 842 个 qs（分布在 745 块里），tiny 块 31 个（f16 的 d 是 0，倒数取 0，整块量化成 0），huge 块 31 个（f16 的 d 是 inf，倒数是 0），nan 块 1 个。scale 字节本身一个没变，变的全是 qs，每处只差 1。
  • 均方误差反而变小了：1.8196e-4 → 1.8178e-4。反量化乘的是 f16 的 d，用它的倒数去量化，round 选出的就是离 x / d_f16 最近的整数；用 f32 的 d 量化，在 x / d 落在 .5 附近时会选到另一边。
  • 所以 ggml 的写法在误差上不是最优的，candle 照抄它，换来的是和现有的 GGUF 文件、和 llama.cpp 量化出来的字节一致。对一个格式来说，「所有实现产出同样的字节」比万分之一的误差更重要。

小结

  • Q8_0 每 32 个值一块：一个 f16 scale（块内最大绝对值除以 127）加 32 个 i8，34 字节，8.5 bit 一个权重。mini 的量化和 candle 逐字节相同，SmolLM2 的 211 个 Q8_0 张量反量化逐位相同。
  • 逐字节一致靠三个细节：scale 对称取 amax / 127，量化用 f32 的 scale 求倒数而存 f16，round 远离 0。改成用 f16 求倒数误差更小，但 905 个字节和 candle 不同。
  • 量化点积块内做整数乘加、每块一次浮点乘法。mini 的标量版和 vec_dot_unopt 逐位相同，和 neon 版差在累加顺序；CPU 的 QMatMul 还要先把输入量化，Metal 的内核不量化输入，同一行输入两者的误差差四个数量级。
  • GGUF 头部是魔数、版本、元数据、张量目录、对齐填充；形状要倒过来读。mini 自己记位置，只需要 Read；candle 每个字符串都问一次 stream_position，BufReader<File> 上是 98384 次 lseek，占了读头部 35 ms 里的大头。
  • general.alignment = 0 让 candle 0.11.0 的 Content::read 除以 0 panic；mini 没有分配上限，靠「不按声明的长度预先分配」挡住假长度，但深层嵌套的数组会让它栈溢出。

💬 关联讨论

发起讨论
加载中...