Rust in Everything
基于 candle-core 0.11.0

QTensor 与 QMatMul

主线位置：推理主线里每个投影层的那次矩阵乘。ModelWeights::from_gguf 把每个权重包成 QMatMul，forward 里 q、k、v、o 四个投影和 MLP 的三个投影、最后的输出层都是 QMatMul::forward（见量化 llama 的一次前向）。本页从 QMatMul 往下走到 CPU 的点积循环和 Metal 的内核。块格式见量化格式。

flowchart TB
  qmm["QMatMul forward"] --> v{"哪个变体"}
  v -->|QTensor| op["apply_op1_no_bwd"]
  v -->|Tensor| mm32["普通 f32 matmul"]
  v -->|TensorF16| mm16["转 f16 再 matmul"]
  op --> dev{"QStorage"}
  dev -->|Cpu| cpu["量化激活 逐行 vec_dot"]
  dev -->|Metal| rows{"行数是 1"}
  rows -->|是| mv["kernel_mul_mv"]
  rows -->|否| mmk["kernel_mul_mm"]

实验 crate 是 exp-candle/quant-exp（candle-core / candle-transformers 0.11.0，开 metal、accelerate），全部 release 构建；平台 M2 Max、macOS 27、rustc 1.99.0。Metal 的计时都是热缓存（.metal 库已在 macOS 的缓存里，冷热的区别见 Metal 后端）。

QTensor 和 QStorage

QTensor 不是 Tensor，是另一套类型：

pub struct QTensor {
    storage: QStorage,
    shape: Shape,
    /// Lazily initialized storage for repacked quantized data. Currently raw bits, could be `QStorage` in the future.
    /// Not always used.
    #[allow(dead_code)]
    repacked_qs: OnceLock<Option<Vec<u8>>>,
}

和 Tensor 比，它没有 Arc、没有 RwLock、没有布局（只有形状，数据一定是连续的块数组），也没有 op 和 is_variable：量化张量不参与求导。存储 QStorage 和 Storage 一样按设备分：

pub enum QStorage {
    Cpu(Box<dyn QuantizedType>),
    Metal(metal::QMetalStorage),
    Cuda(cuda::QCudaStorage),
}

CPU 上是 Box<dyn QuantizedType>，里面其实是 Vec<BlockQ8_0> 这样的块数组，dtype 在运行时通过 trait 对象分派（为什么是 trait 对象见量化格式）。Metal 上的 QMetalStorage 是 dtype、设备和一个 Arc<Buffer>，块的字节原样放在 MTLBuffer 里，由内核自己解释。

QTensor::new 只检查一件事：最后一维能被块大小整除（quantized tensor must have their last dim divisible by block size）。块沿着最后一维切，一行的数据不跨块。

几个转换方法：

repacked_qs 只给 CPU 上的 Q4K 用，下面讲。

QMatMul：三种形态

QMatMul 是一个枚举：

#[derive(Clone, Debug)]
pub enum QMatMul {
    QTensor(std::sync::Arc<QTensor>),
    Tensor(Tensor),
    TensorF16(Tensor),
}

QTensor 形态保留量化权重，每次乘法用量化内核；Tensor 形态在构造时就反量化成 f32，之后是普通 matmul；TensorF16 反量化成 f16，乘之前把输入转成 f16、乘完转回来。选哪个在 QMatMul::from_arc 里决定，看 dtype 和两个环境变量：

    pub fn from_arc(qtensor: std::sync::Arc<QTensor>) -> Result<Self> {
        let dequantize = match qtensor.dtype() {
            GgmlDType::F32 | GgmlDType::F16 | GgmlDType::BF16 => true,
            _ => DEQUANTIZE_ALL.with(|b| *b),
        };
        let t = if dequantize {
            let tensor = qtensor.dequantize(&qtensor.device())?;
            Self::Tensor(tensor)
        } else if DEQUANTIZE_ALL_F16.with(|b| *b) {
            let tensor = qtensor.dequantize_f16(&qtensor.device())?;
            Self::TensorF16(tensor)
        } else {
            Self::QTensor(qtensor)
        };
        Ok(t)
    }
  • 本来就不是量化类型的（F32、F16、BF16）总是反量化成 f32 Tensor，走普通 matmul。
  • 设了 CANDLE_DEQUANTIZE_ALL（非空、不是 "0"）就全部反量化成 f32；否则设了 CANDLE_DEQUANTIZE_ALL_F16 就反量化成 f16；都没设才保留量化。
  • 两个变量放在 thread_local! 里，每个线程第一次用到时读一次环境变量，之后不再变。from_arc 在建模型时调用，所以要在加载模型之前设好。

没有按设备或形状自动选择：在 Metal 上、在 prefill 时，只要没设环境变量，就都是 QTensor 形态。模型代码也可以绕过 from_arc 直接构造某个变体，下面的实验就是这样把同一个权重做成三种形态的。

forward 里，Tensor 和 TensorF16 两支把权重 broadcast_left 到输入的批维再转置，交给 Tensor::matmul（走一次 matmul 那条路）；QTensor 一支是：

            Self::QTensor(t) => xs.apply_op1_no_bwd(t.as_ref()),

QTensor 实现了 CustomOp1（名字 "qmatmul"），量化矩阵乘就是一个自定义算子：不记反向，按输入所在的设备调 cpu_fwd 或 metal_fwd。权重在算子对象里，输入是那个唯一的张量参数。权重按 [n, k] 存（和 PyTorch 的 Linear 一样，每个输出一行），输入的最后一维必须是 k，输出把最后一维换成 n。

CPU：先量化激活，再逐行点积

cpu_fwd 要求输入连续，把除最后一维之外的都拍平成 m 行，按输入 dtype 分两支：f32 走 matmul_t，f16 走 matmul_t_f16。f32 这一支调到 k_quants::matmul，核心是两步：

        // f32, f16, and bf16 support direct copy
        if T::DIRECT_COPY {
            T::VecDotType::direct_copy(lhs, lhs_b);
        } else {
            for row_idx in 0..m {
                let lhs_b_mut = &mut lhs_b[row_idx * k_in_blocks..(row_idx + 1) * k_in_blocks];
                let lhs = &lhs[row_idx * k..(row_idx + 1) * k];
                T::VecDotType::from_float(lhs, lhs_b_mut)
            }
        }
        // ...
        for row_idx in 0..m {
            // ...
            pool.execute(|tid| {
                // ... 本线程负责的列，每 4 列一组
                for quad_idx in start..end {
                    let col = quad_idx * 4;
                    let (d0, d1, d2, d3) = T::vec_dot_4(
                        k,
                        &rhs_t[col * k_in_blocks..(col + 1) * k_in_blocks],
                        // ... 另外三列
                        lhs_row,
                    );
                    // ... 写进 dst
                }
            });
            // ... 剩下不足 4 列的用 vec_dot_2 / vec_dot
        }

第一步把输入的每一行量化成权重类型的 VecDotType（Q8_0 权重配 Q8_0 激活，K-quants 配 Q8K），放进一个线程局部、跨调用复用的 scratch 缓冲区（#3627 加的，避免每次乘法都分配）。第二步对每一行输入，把 n 个输出按 4 列一组分给 barrier_pool() 的线程（线程池和 CANDLE_NUM_THREADS 见 CPU 后端），每组用 vec_dot_4 让一行激活和四行权重各做一次点积。

这个结构决定了两件事：

  • 一行输入就是一次 pool.execute。 decode 时 m = 1，一次并行段把整个权重读一遍，每个权重字节只读一次，正是量化省带宽的地方。prefill 时 m = 13，就是 13 次并行段，权重被完整地读 13 遍，没有分块复用。
  • 激活也被量化了。 结果不等于「反量化后的权重乘原始 f32 输入」，多了激活量化的误差。下面的实测里这是 CPU 和 Metal 结果不一样的根源。

f16 输入的 matmul_f16 是另一份更简单的实现：每次新分配 lhs_b，两层循环逐个 vec_dot，不用线程池，也不用 vec_dot_4。

Q4K 的 8 列重排

在 aarch64 且开着 dotprod 时（Apple 芯片的默认目标就是），cpu_fwd 对 Q4K 权重、n 是 8 的倍数、输入 f32 的情况先走另一条路：

                // Try the 8-column BlockQ4Kx8 repacked path.
                #[cfg(all(target_arch = "aarch64", target_feature = "dotprod"))]
                if self_storage.dtype() == GgmlDType::Q4K && n.is_multiple_of(8) {
                    // ...
                    let repacked = self.repacked_qs.get_or_init(|| {
                        // ...
                        let packed = k_quants::pack_to_q4kx8(blocks, n);
                        Some(packed.as_bytes().to_vec())
                    });

pack_to_q4kx8 把相邻 8 行权重同一位置的 Q4K 块交错存成一个 1152 字节的 BlockQ4Kx8，matmul_q4k_x8 用 vec_dot_8_q4k_q8k 一次算 8 个输出（这是 llama.cpp repack.cpp 里同一种做法的移植，#3643 加进来）。重排在第一次 forward 时做，结果存进 repacked_qs 这个 OnceLock，原来的块不释放：Q4K 张量在 CPU 上用过一次之后占两份内存。实测一个 8192 × 4096 的 Q4K 权重，第一次单行乘法 6.21 ms，之后 0.13 ms。

fast_mmq / fast_mmvq 只属于 CUDA

quantized/ 下的 fast_mmq.rs（400 行）和 fast_mmvq.rs（323 行）不是 CPU 代码。mod.rs 里：

#[cfg(feature = "cuda")]
pub mod fast_mmq;
#[cfg(feature = "cuda")]
pub mod fast_mmvq;

文件头分别写着「CUDA fast path for GGUF tiled matmul (prompt/prefill phase)」和「CUDA fast path for GGUF matmul with BF16/F32 activations」：fast_mmvq 管 1 到 8 行输入，fast_mmq 管更多行，都由 cuda.rs 里的 QCudaStorage::fwd 先试（#3463、#3465，2026 年 4 月）。本卷没有 CUDA 环境，这部分只读了源码。CPU 上近期的量化性能工作是 #3627 的 BarrierPool 加 scratch 缓冲区、#3643 的 vec_dot_4 和 Q4K 重排，都在上面那段 k_quants::matmul 里。

Metal：一行用 mv，多行用 mm

QMetalStorage::fwd 看输入倒数第二维：

        if src_shape.dim(D::Minus2)? == 1 {
            return self.fwd_mv(self_shape, storage, layout);
        }

一行（decode）走 fwd_mv，对每个批元素编码一次 call_quantized_matmul_mv_t，内核名是 kernel_mul_mv_{dtype}_f32；多行（prefill）编码一次 call_quantized_matmul_mm_t，内核是 kernel_mul_mm_{dtype}_f32。两组内核都在 quantized.metal（7741 行，从 llama.cpp 的 ggml-metal.metal 移植）里。

两种内核对数据的处理不同。mv 内核每个线程组负责几行权重，直接读量化块，在寄存器里解出整数乘 f32 的输入，按 simdgroup 求和，kernel_mul_mv_q8_0_f32 的签名里输入就是 device const float * src1。mm 内核是分块矩阵乘：

template<typename T, typename T4x4, typename simdgroup_T8x8, typename block_q, short nl, void (*dequantize_func)(device const block_q *, short, thread T4x4 &)>
kernel void kernel_mul_mm(device const  uchar * src0,
                          device const  uchar * src1,
                          device        float * dst,
                          // ...
    threadgroup T     * sa = (threadgroup T     *)(shared_memory);
    threadgroup float * sb = (threadgroup float *)(shared_memory + 4096);
    // ...
        T4x4 temp_a;
        dequantize_func(x, il, temp_a);

每个线程组把一块权重 dequantize_func 成 half（T 实例化为 half）放进线程组内存，输入保持 float，用 simdgroup_multiply_accumulate 做 8 × 8 的矩阵乘加，结果 f32。两种内核都不量化激活。

Metal 这边有几处限制，实验程序 qedge 跑出来的：

[metal] QTensor Q8K m=1: Err("Metal error Error while loading function: kernel_mul_mv_q8_K_f32")
[metal] QTensor Q8K m=3: Err("Metal error Unsupported dtype Q8K for operation qmatmul")
[metal] QTensor Q8_1 m=1: Err("Metal error Error while loading function: kernel_mul_mv_q8_1_f32")
[metal] QTensor Q8_1 m=3: Err("Metal error Unsupported dtype Q8_1 for operation qmatmul")
[metal] QTensor BF16 m=1: Err("Metal error Error while loading function: kernel_mul_mv_bf16_f32")
[metal] QTensor BF16 m=3: Ok(9.980648)

Rust 一侧给 Q8K、Q8_1、BF16 都起了 mv 内核的名字，.metal 里却没有这三个函数。Q8K 和 Q8_1 不会出现在 GGUF 权重里，BF16 权重经 from_arc 会被反量化成 f32 Tensor，正常路径碰不到；直接把 QTensor 当算子用才会撞上。输入必须是 f32：mm 那一支有一句 assert_eq!(storage.dtype(), DType::F32)，mv 那一支没有检查：

[metal] Q8_0 with f16 input m=1: Ok, first [-0.87291944, 0.30382147, -0.40540972] vs f32 input [0.05893158, -0.08843393, -0.23894244]

thread 'main' (3887975) panicked at ~/.cargo/registry/src/index.crates.io-1949cf8c6b5b557f/candle-core-0.11.0/src/quantized/metal.rs:390:9:
assertion `left == right` failed
  left: F16
 right: F32

f16 输入一行时，内核把 f16 的字节当 f32 读，返回一组错误的数而不报错；三行时 panic。CPU 上同样的 f16 输入走 matmul_f16，结果正确（[cpu] Q8_0 with f16 input m=1: Ok, first [0.49047852, ...] vs f32 input [0.49050874, ...]）。

📝NOTE
冷知识：Metal 上的 prefill 曾经是一行一个内核

fwd 里「多行走 mm」这个分支来自 #2615「Major T/s improvement Use the Metal qmatmul MM kernels」（2025-07-18 合并，首次出现在 0.9.2-alpha.1，第一个正式版是 0.10.0）。之前 Metal 上只有现在的 fwd_mv，名字就叫 fwd：把输入拍成 m 行，for batch_id in 0..m 每行编码一次 mv 内核，13 个 token 的 prefill 就是每个投影 13 次内核调用，每次把整个权重读一遍。这个 PR 把旧函数改名为 fwd_mv，新写一个 fwd 用移植来的 kernel_mul_mm 一次算完所有行，只在 dim(D::Minus2) == 1 时退回 mv。提交说明里有一条「Try out always using mm」，紧接着是两条「Mirror llama.cpp metric」：最终没有总用 mm，单行仍走 mv。下面的实测能看出原因。

实测：同一个乘法的三种形态

程序 qmv 用一个 8192 × 4096 的随机权重（约是 7B 模型一个投影的大小）量化成 Q8_0、Q4K、Q4_0，各建 QTensor、Tensor（f32）、TensorF16 三种 QMatMul，输入 m 行。每种先跑 3 次预热，再测 7 批（一批 200 次或 40 次 forward，批尾 synchronize 一次），取最快的一批算每次的平均。

Metal（热缓存）：

Q8_0 m=1    quantized 0.091 ms  f32 0.390 ms  f16 0.196 ms
Q8_0 m=2    quantized 0.344 ms  f32 0.618 ms  f16 0.355 ms
Q8_0 m=4    quantized 0.346 ms  f32 0.619 ms  f16 0.357 ms
Q8_0 m=8    quantized 0.347 ms  f32 0.629 ms  f16 0.356 ms
Q8_0 m=16   quantized 0.350 ms  f32 0.489 ms  f16 0.612 ms
Q8_0 m=64   quantized 0.660 ms  f32 0.763 ms  f16 0.588 ms
Q8_0 m=256  quantized 2.213 ms  f32 2.349 ms  f16 2.118 ms
Q4K m=1    quantized 0.065 ms  f32 0.395 ms  f16 0.199 ms
Q4K m=2    quantized 0.388 ms  f32 0.619 ms  f16 0.356 ms
...
Q4K m=256  quantized 2.550 ms  f32 2.340 ms  f16 2.123 ms
Q4_0 m=1    quantized 0.049 ms  f32 0.394 ms  f16 0.196 ms
Q4_0 m=2    quantized 0.333 ms  f32 0.615 ms  f16 0.355 ms
...
Q4_0 m=256  quantized 2.194 ms  f32 2.384 ms  f16 2.124 ms
  • 一行时时间和权重字节数成正比。 f32 权重 134 MB 用 0.39 ms，f16 67 MB 用 0.20 ms，Q8_0 35.7 MB 用 0.091 ms，Q4_0 18.9 MB 用 0.049 ms，折算都在 340 到 390 GB/s，接近这台机器的内存带宽。decode 是读权重的速度决定的，量化在这里直接换成了速度：Q4_0 比 f32 快 8 倍。
  • 两行起换成 mm 内核，时间跳到 0.33–0.39 ms，比两次 mv 还慢，和 f16 的普通 matmul 差不多。mm 内核按 32 行输入一个分块，2 行和 16 行花的时间几乎一样。
  • 64 行以上三种形态差不多（256 行都在 2.1–2.6 ms），这时是算力决定的，f16 的 MLX gemm 略快。

CPU（qmv cpu，第二次运行；f16 那一列在开了 accelerate 的构建里不能用）：

Q8_0 m=1    quantized 0.156 ms  f32 2.374 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q8_0 m=2    quantized 0.325 ms  f32 8.740 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q8_0 m=4    quantized 0.628 ms  f32 4.494 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q8_0 m=8    quantized 1.511 ms  f32 4.732 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q8_0 m=16   quantized 3.236 ms  f32 4.661 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q8_0 m=64   quantized 11.969 ms  f32 5.467 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q8_0 m=256  quantized 54.466 ms  f32 10.657 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q4K m=1    quantized 0.101 ms  f32 2.496 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
...
Q4K m=16   quantized 1.792 ms  f32 4.736 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q4K m=64   quantized 7.498 ms  f32 5.514 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q4K m=256  quantized 32.418 ms  f32 10.634 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
Q4_0 m=1    quantized 0.196 ms  f32 2.336 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
...
Q4_0 m=256  quantized 91.568 ms  f32 15.453 ms  f16 Err(the accelerate backend does not support f16 matmul) ms
  • 一行时量化快 15 到 24 倍（0.10–0.20 ms 对 2.3–2.5 ms）：权重只读一遍，i8 点积。
  • 量化的时间和行数成正比，每行 0.1–0.2 ms，因为每行都要把权重读一遍；Accelerate 的 f32 sgemm 会分块复用，64 行和 1 行只差一倍多。交叉点在 16 到 64 行之间，256 行时量化慢 3 到 6 倍。所以 CPU 上的长 prompt prefill，量化形态是吃亏的。
  • f16 形态在 CPU 上不能用：TensorF16 的乘法是普通 f16 matmul，开了 accelerate 时 CpuStorage::matmul 直接报 the accelerate backend does not support f16 matmul（f16 matmul 的限制见 CPU 后端）。设 CANDLE_DEQUANTIZE_ALL_F16=1 跑 CPU 生成，第一次 forward 就是这个 Error。

反量化本身的代价在建模型时付：8192 × 4096 的 Q8_0 权重，CPU 上反量化成 f32 约 13 ms，Metal 上因为要绕 CPU 一圈约 26–32 ms。SmolLM2 的小形状（576 × 1536 之类）在 Metal 上每次乘法只有 0.06–0.2 ms，三种形态的差别淹没在提交开销里，qmm 程序的结果见实验目录的 qmm.out。

实测：Q8_0 和 Q4_K_M 生成

程序 qgen 用和生成一个 token相同的 prompt（13 个 token）贪心生成 20 个 token，读文件用 BufReader，每个设置一个新进程、连跑三轮：

[q8-cpu round 2] load 65.2 ms, prefill (13 tokens) 32.8 ms, 20 decode steps 135.3 ms (mean 6.76 ms)
[q8-cpu round 3] load 62.3 ms, prefill (13 tokens) 29.3 ms, 20 decode steps 140.6 ms (mean 7.03 ms)
[q8-metal round 2] load 106.1 ms, prefill (13 tokens) 13.1 ms, 20 decode steps 55.4 ms (mean 2.77 ms)
[q8-metal round 3] load 104.3 ms, prefill (13 tokens) 11.2 ms, 20 decode steps 51.8 ms (mean 2.59 ms)
[q4-cpu round 2] load 60.2 ms, prefill (13 tokens) 84.0 ms, 20 decode steps 238.1 ms (mean 11.91 ms)
[q4-cpu round 3] load 56.0 ms, prefill (13 tokens) 91.7 ms, 20 decode steps 260.7 ms (mean 13.04 ms)
[q4-metal round 2] load 99.9 ms, prefill (13 tokens) 11.5 ms, 20 decode steps 58.3 ms (mean 2.91 ms)
[q4-metal round 3] load 92.5 ms, prefill (13 tokens) 11.4 ms, 20 decode steps 59.7 ms (mean 2.99 ms)

生成的文本：

[q8-cpu] text: "Rust is a modern, high-performance, and object-oriented programming language developed by Brian Kern"
[q8-metal] text: "Rust is a modern, high-performance, and object-oriented programming language developed by Brian Kern"
[q4-cpu] text: "Rust is a modern, high-performance, and efficient language developed by the Rust community. It was"
[q4-metal] text: "Rust is a modern, high-performance, and efficient programming language developed by Rust Core. It was"
  • Metal 上两个文件一样快（decode 每步 2.6–3.0 ms），这个模型太小，时间主要在内核调用次数上，不在读权重上。
  • CPU 上 Q4_K_M 比 Q8_0 慢：prefill 84–92 ms 对 29–33 ms，decode 12–13 ms 对 7 ms。文件小了 27%，却慢了将近一倍。
  • Q8_0 在 CPU 和 Metal 上生成一字不差；Q4_K_M 在两个设备上从第 13 个 token 起就不一样。
📝NOTE
没有蠢问题

问： Q4_K_M 每个权重的位数比 Q8_0 少，CPU 上要读的字节更少，为什么反而慢一倍？

答： 因为这个文件里 166 个张量是 Q5_0（行长 576 用不了 K-quants，见 GGUF 文件），而 Q5_0 的 vec_dot 在任何平台上都只有标量实现，也没有 vec_dot_4。qmm 程序里同一层的 MLP 投影，Q8_0 一行 0.006 ms，Q5_0 一行 0.025 ms；13 行时 0.068 ms 对 0.673 ms。注意力和 MLP 的 7 个投影里有 5 个是 Q5_0，速度就由它决定了。量化格式的好坏要看实现里有没有对应的 SIMD 路径，位数只是一半。Metal 上每种格式都有内核，就看不出这个差别。

两个设备为什么不一致

qgen 把第一步（prefill 之后）的 49152 个 logits 存下来。再用 CANDLE_DEQUANTIZE_ALL=1 跑一遍（-deq32，权重反量化成 f32、走普通 matmul），对比各组之间 logits 的最大差：

max|q4-metal-deq32 - q4-cpu-deq32| = 0.0000
max|q4-metal-deq32 - q4-cpu| = 0.8369
max|q4-metal-deq32 - q4-metal| = 0.0260
max|q4-metal-deq32 - q8-metal-deq32| = 1.9615
max|q8-metal-deq32 - q8-cpu| = 0.5388
max|q8-metal-deq32 - q8-metal| = 0.0154
  • 把权重反量化之后，CPU 和 Metal 的 logits 一致（差小于 0.00005）。
  • Metal 的量化形态离它只有 0.015（Q8_0）和 0.026（Q4_K_M）：mv / mm 内核不量化激活，差别只是 f16 中间值和求和顺序。
  • CPU 的量化形态差 0.54（Q8_0）和 0.84（Q4_K_M）：激活被量化成 Q8_0 或 Q8K，多了一层误差。
  • 对照：Q4_K_M 和 Q8_0 两种权重本身的差是 1.96。

所以「CPU 和 Metal 生成的文本不同」是 CPU 量化矩阵乘的特性，不是 Metal 算错了。Q8_0 时的 0.54 还没改变贪心选出的 token，Q4_K_M 时 0.84 加上权重误差，在第 13 个 token 处翻了过来（「efficient language」对「efficient programming language」）。

❗IMPORTANT
动脑筋：在 CPU 上设 CANDLE_DEQUANTIZE_ALL=1 跑 Q4_K_M，生成的文本会和上面哪一行一样？decode 每步会变快还是变慢？先猜，再看文末。

动脑筋答案

[q4-cpu-deq32 round 1] load 150.9 ms, prefill (13 tokens) 39.4 ms, 20 decode steps 313.0 ms (mean 15.65 ms)
[q4-cpu-deq32] text: "Rust is a modern, high-performance, and efficient programming language developed by Rust Core. It was"
[q4-cpu-deq32 round 2] load 137.4 ms, prefill (13 tokens) 38.3 ms, 20 decode steps 327.2 ms (mean 16.36 ms)
[q4-cpu-deq32 round 3] load 136.7 ms, prefill (13 tokens) 39.0 ms, 20 decode steps 373.5 ms (mean 18.68 ms)
  • 文本和 Metal 的一样（「developed by Rust Core」）：去掉激活量化之后，CPU 和 Metal 算的是同一个数学，logits 差小于 0.00005。
  • decode 变慢：每步 15.7–18.7 ms，比量化形态的 12–13 ms 还慢。一行输入时 f32 权重要读 4 倍多的字节（上面 8192 × 4096 的实验里是 15 到 24 倍的差距），Q5_0 的标量点积再慢也比不过。prefill 反而快了一倍多（38–39 ms 对 84–92 ms）：13 行时 Accelerate 的 sgemm 能复用权重，量化路径要读 13 遍。
  • 建模型从约 60 ms 变成约 137 ms，多出来的是 272 个张量的反量化。Q8_0 文件同样设置时 decode 是 15 ms、文本不变。

小结

  • QTensor 是没有布局、不求导的另一套张量；QStorage 在 CPU 上是 Box<dyn QuantizedType>，Metal 上是一个装着原始块的缓冲区。Metal 上的量化、反量化都绕 CPU。
  • QMatMul 三种形态：保留量化、反量化成 f32、反量化成 f16。选择只看 dtype 和 CANDLE_DEQUANTIZE_ALL / CANDLE_DEQUANTIZE_ALL_F16，在建模型时定死；QTensor 形态经 apply_op1_no_bwd 当自定义算子跑。
  • CPU：先把激活逐行量化成 Q8_0 / Q8K，再每行一次 BarrierPool 并行段、4 列一组 vec_dot_4；Q4K 在 aarch64 上第一次用时重排成 8 列交错（多占一份内存）。一行时比 f32 快 15 倍以上，行数多时按行线性变慢，64 行以上不如 Accelerate。fast_mmq / fast_mmvq 是 CUDA 专用。
  • Metal：一行用 mv 内核，读带宽决定速度，Q4_0 比 f32 快 8 倍；两行起用 mm 内核，权重反量化成 half 进线程组内存，激活保持 f32。f16 输入一行时静默算错，多行时 panic。
  • SmolLM2 的 Q4_K_M 在 CPU 上比 Q8_0 慢一倍（Q5_0 只有标量点积），在 Metal 上一样快；CPU 的量化路径因为量化激活，和 Metal 的结果差 0.5–0.8 个 logit，Q4_K_M 的生成文本因此在两个设备上不同。

💬 关联讨论

发起讨论
加载中...