从 Qwen 3.8 谈谈模型量化与硬件支持
1. 前言
最近 Qwen 3.8 27B 火了起来，很容易在网上看到大家纷纷用 24G 显存的显卡或者统一内存的笔记本跑起来。我实测了一下 Coding 和 Debug 能力也确实有接近 Opus 4.6 的水准，而且配合 Speculative Decoding 通常能在显存带宽的 Roofline 上获得 3-4x 的加速，让我们看到了端侧模型真正完成大部分任务的可能性。比如 200GB/s 的有效内存带宽配上每 Token 20GB 的权重 + KV 访存，理论上每秒能 decode 10 Token，而 Speculative Decoding 可以达到 30-40 Token/s，足够日常使用了。
但在本地部署上，llama.cpp (ggml) 这个推理框架可谓是真神，支持了超级多的量化格式，以 Unsloth 的 Qwen3.8-27B GGUF 为例，看看它们的体积：
| 量化 | |
|---|---|
| 1-bit | UD-IQ1_S 6.19 GB UD-IQ1_M 6.73 GB |
| 2-bit | UD-IQ2_XXS 7.27 GB UD-IQ2_S 8.37 GB UD-Q2_K_XL 9.83 GB |
| 3-bit | UD-IQ3_XXS 10.9 GB UD-IQ3_S 12 GB UD-Q3_K_XL 13.1 GB |
| 4-bit | UD-IQ4_XS 14.3 GB UD-Q4_K_S 15.4 GB Q4_0 16.1 GB UD-Q4_K_M 16.5 GB Q4_1 17.5 GB UD-Q4_K_XL 17.6 GB |
| 5-bit | UD-Q5_K_S 18.7 GB UD-Q5_K_M 19.8 GB UD-Q5_K_XL 20.9 GB |
| 6-bit | UD-Q6_K 22 GB UD-Q6_K_M 23.1 GB UD-Q6_K_L 24.2 GB UD-Q6_K_XL 25.3 GB |
| 8-bit | Q8_0 29 GB UD-Q8_K_XL 31.5 GB |
| 16-bit | BF16 54.7 GB |
然后问题来了：我们自己部署的时候，应该用哪个呢？特别是同样的量化 bit 数往往有好几个版本。
但是实际测试我发现，尤其是 IQ 系列明明少搬了一半字节，节省了那么多显存带宽，但为什么引入了巨大的计算开销，甚至在一些显卡上跑得比 Q 系列还慢？这背后的原理是什么？
因此，我跟着 AI 开始探究：浮点格式到底有多少种、4 比特以下为什么必须分 block、Q 系列和 IQ 系列各自在算什么、为什么现代硬件绕不开查找表，以及最后，上面那个文件列表究竟该怎么选，最后跟着 AI 学习了很多很多，也写了这篇文章。我也仔细思考了一下，未来端侧硬件会不会直接用 LUT 来实现高效的量化推理。
2. 先看一张图，量化到底掉了多少精度？
我们先来看看 unsloth 官方给出的实测曲线。横轴是 GGUF 文件大小（GB，已扣掉 MTP），纵轴是 Top-1% Accuracy：
图片来源：unsloth Dynamic 3.0 GGUFs
绿色那条就是 unsloth 的 Dynamic v3.0，几个点的大致位置：
| 量化 | 权重体积 | Top-1% | 量化 | 权重体积 | Top-1% |
|---|---|---|---|---|---|
| UD-IQ1_S | 6.19 GB | ~73% | UD-IQ4_XS | 14.3 GB | ~94% |
| UD-IQ2_XXS | 7.27 GB | ~79% | UD-Q4_K_M | 16.5 GB | ~96% |
| UD-Q2_K_XL | 9.83 GB | ~87% | UD-Q5_K_M | 19.8 GB | ~97% |
| UD-IQ3_S | 12 GB | ~91% | UD-Q6_K | 22 GB | ~98% |
| UD-Q3_K_XL | 13.1 GB | ~93% | Q8_0 | 29 GB | ~99% |
从这张图我们可以看出：
-
拐点确实在 Q4 附近：从 6.19 GB 到 16.5 GB，精度从 73% 涨到 96%；而从 16.5 GB 再涨到 29 GB，只多了 3 个点。由此看出， Q4_K_M 是 sweet spot，Q5_K_M 也还行，Q6_K 以上就没必要了。
-
10 GB 以下精度损失很可观：Q2 及以下每省 1 GB 会让精度掉好几个百分点，IQ1_S 直接掉到 73%。所以至于那些 16GB 内存的 Mac mini 本地运行的情况（macOS 默认只让 GPU 用大约 2/3 的统一内存），基本上只能用 IQ1/IQ2 这一档，性能损失不可避免。
-
同样的文件大小，不同的量化权重发布者能差好几个点： Unsloth 的图对比了 4 家量化权重发布者的同体积模型和精度，差距有 8 个百分点左右，甚至有的还是完全相同的量化格式。究竟差在哪呢？差在逐 tensor 的比特分配方法，和跑 imatrix 用的校准语料。
3. 为什么需要量化？
对于 dense 模型来说，LLM 的自回归 decode，每生成一个 token 就要把全部权重读一遍。比方说我们的 dense 模型权重 10GB，而我们的内存带宽是 256GB/s，那么每秒最多能读 25.6 个 token。对于 MoE 这个问题大幅缓解，虽然权重很大，但推理中，per token per layer 只有约 1/16 的 Expert 被激活，读的量就大幅减少了。
这个比例有多难看呢？
| 显卡 | 显存 | FP16 算力 | 显存带宽 | 硬件想要的 flops/byte |
|---|---|---|---|---|
| RTX 3090 | 24 GB | 71 TFLOP/s | 936 GB/s | ~76 |
| RTX 4090 | 24 GB | 165 TFLOP/s | 1008 GB/s | ~164 |
| RTX 5090 | 32 GB | 210 TFLOP/s | 1792 GB/s | ~117 |
| RX 7900 XTX | 24 GB | 123 TFLOP/s | 960 GB/s | ~128 |
| Arc Pro B70 | 32 GB | 未公布 | 608 GB/s | 见下 |
| Strix Halo (Radeon 8060S) | 128 GB 统一内存 | 59 TFLOP/s | 256 GB/s | ~232 |
| batch=1 decode 实际能提供的 | 2 flops / 2 bytes = 1 |
注1：考虑 Speculative Decoding 的话，batch=1 decode 的吞吐量可以达到 3-4 flops/byte，但仍然远低于硬件的潜力。
注2：算力一列取 FP16 Tensor Core 的 dense 值。GeForce 卡在 FP32 累加时是半速，表里给的就是 FP32 累加的数，llama.cpp 的 mma 路径走的正是 f32 累加。
Arc Pro B70 很复杂，官方数据是 22.94 TFLOPS FP32 和 367 TOPS INT8，没有给 XMX 的 FP16 数。不过按 INT8 算，367 TOPS / 608 GB/s ≈ 600 ops/byte，结论不受影响，而且对我们的 scope 来说 INT8 反而更合适，因为量化推理的内积本来就走整数。
几件有意思的事：
- 几张卡的比例都在 76 到 232 之间，而 decode 只能给 1。 差了两个数量级。
- 5090 的比例反而比 4090 好看。 因为 GDDR7 让带宽涨了 78%，而算力只涨了 27%。对 batch=1 decode 来说，5090 的提升幅度比跑分上看起来的更大。
Strix Halo 是这里面最特别的一个：128 GB 是 CPU 和 GPU 共用的 LPDDR5X 统一内存。带宽只有 256 GB/s，比 3090 还低不少，但容量是其他几张的四倍以上。换句话说，它能装下别人装不下的模型，只是跑得慢。
顺便拿前面那张文件列表对一下，就知道这些卡各自能吃到哪一档：
- 24 GB（3090 / 4090 / 7900 XTX）：UD-Q4_K_XL 17.6 GB 很舒服，还剩 6 GB 给 KV；UD-Q5_K_XL 20.9 GB 开始紧张；UD-Q6_K 22 GB 基本只够短上下文。
- 32 GB（5090 / B70）：可以一路吃到 UD-Q6_K_XL 25.3 GB，Q8_0 29 GB 也塞得下但没多少余量了。
- 128 GB 统一内存（Strix Halo）：连 BF16 的 54.7 GB 都塞得下。顺带一提，那张图里 unsloth 标的硬件兼容性第一个就是 Ryzen AI Max 128 GB，正是这类机器。
所以结论很直接：decode 速度是一个带宽指标，不是算力指标。 每个权重的字节数减半，吞吐大致翻倍。这就是推理侧量化存在的全部理由，也是为什么 Q4_K_M 这个档位的性价比最高。
KV Cache 是同一个故事的另一种形态：它随上下文长度线性增长，而且每生成一个 token 都要被完整重读一遍。
KV cache 字节数 = 2 * n_layer * n_kv_head * head_dim * n_ctx * sizeof(elem)
^ ^
K 和 V 各一份 每元素字节数Qwen3.8-27B 算一下。它是个混合架构：64 层里只有每 4 层一层是全注意力（共 16 层），其余 48 层是 Gated DeltaNet 这种线性注意力，状态大小固定、不随上下文增长。全注意力那 16 层是 GQA 24/4，head_dim 256。
只有 16 层有会长大的 KV：
每 token = 2 * 16 * 4 * 256 * 2 byte = 65536 byte = 64 KB
128k 上下文，fp16：64 KB * 131072 = 8 GB
256k 上下文，fp16：64 KB * 262144 = 16 GB <-- 这是它标称的最大上下文8 GB 是什么概念？前面文件列表里 UD-Q4_K_XL 是 17.6 GB，一张 24 GB 的卡装完权重刚好剩 6 GB 出头，也就是说 fp16 的 KV 撑不到 128k 就先爆了。这时候量化 KV 对于我们能不能开满 256k 上下文就至关重要了。
而且它恰好满足后面要讲的两条硬约束：head_dim = 256，既能被 block 大小 32 整除（所以 Q4_0/Q8_0 都能用），也能被 64 整除（所以 Hadamard 旋转那条路径也走得通）。
这里还有一件事值得提前说：K 和 V 对量化的敏感度并不一样。
K 参与的是 Q·K 内积，结果要过 softmax。softmax 是指数函数，输入端一点点误差会被指数放大，进而改变整个注意力分布。错的不是数值，是“模型看向哪里”。而 V 的误差只是被注意力权重加权平均掉，几十上百个 token 的误差平均下来，反而相对温和。
所以 K 比 V 敏感，KV 量化天然应该是非对称的：
| 配置 | 128k 的 KV | 相对 fp16 |
|---|---|---|
| fp16 / fp16 | 8.00 GB | 100% |
-ctk q8_0 -ctv q8_0 |
4.25 GB | 53% |
-ctk q8_0 -ctv q4_0 |
3.25 GB | 41% |
-ctk q4_0 -ctv q4_0 |
2.25 GB | 28% |
中间那一档往往是最划算的：K 保住 8 bit 尽可能不影响 softmax，V 压到 4 bit 省下的空间最多。第 10 节会讲清楚这背后的机制，以及 llama.cpp 为此做的一件挺妙的事：量化之前先把 K、V 转一个身。
而且这其实是两个问题，性质完全不同：
| 权重量化 | KV Cache 量化 | |
|---|---|---|
| 什么时候做 | 离线，一次 | 在线，每个 token |
| 编码器能做多复杂 | 可以多趟迭代搜索 | 只能单趟闭式计算 |
| 有没有校准数据 | 有，重要性矩阵 | 没有，数据还没产生 |
| 数据分布 | 静态，研究得很透 | 动态，离群值多，和位置相关 |
| 坏掉的表现 | 模型整体变笨 | 模型在长上下文后段变笨 |
后面第 10 节会专门讲 KV，先从浮点格式说起。
4. 浮点格式到底有多少种？
4.1 先拆开看一个浮点数
所有 IEEE-754 风格的浮点数都是同样的三个字段：
+---+-----------+--------------------+
| S | exponent | mantissa |
+---+-----------+--------------------+
1 E bits M bits
符号位 (S) -> 正负
指数位 (E) -> 表示这个数是 2 的多少次幂
尾数位 (M) -> 精度：有多少位有效数字
bias = 2^(E-1) - 1
数字 = (-1)^S * 2^(E-bias) * (1 + M/2^M)按比例画出来是这样：
S EEEEEEEEEEE MMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMM
FP64 |1| 11 | 52 | 64 bit
S EEEEEEEE MMMMMMMMMMMMMMMMMMMMMMM
FP32 |1| 8 | 23 | 32 bit
S EEEEEEEE MMMMMMM
BF16 |1| 8 | 7 | 16 bit
S EEEEE MMMMMMMMMM
FP16 |1| 5 | 10 | 16 bit
S EEEEE MM
E5M2 |1| 5 |2| 8 bit
S EEEE MMM
E4M3 |1| 4 | 3 | 8 bit
EEEEEEEE
E8M0 | 8 | <- 没有符号位，没有尾数。就是一个 2 的幂。 8 bit
S EE M
E2M1 |1|2|1| 4 bit注意 BF16 就是 FP32 把低 16 位尾数砍掉，所以它俩之间的转换只是一次移位。
硬指标：
| 格式 | S/E/M | bias | 最大有限值 | 最小规格数 | 十进制有效位 | inf/NaN |
|---|---|---|---|---|---|---|
| FP64 | 1/11/52 | 1023 | 1.8e308 | 2.2e-308 | ~15.9 | 有 |
| FP32 | 1/8/23 | 127 | 3.4e38 | 1.18e-38 | ~7.2 | 有 |
| TF32 | 1/8/10 | 127 | 3.4e38 | 1.18e-38 | ~3.3 | 有 |
| BF16 | 1/8/7 | 127 | 3.39e38 | 1.18e-38 | ~2.4 | 有 |
| FP16 | 1/5/10 | 15 | 65504 | 6.1e-5 | ~3.3 | 有 |
| FP8 E5M2 | 1/5/2 | 15 | 57344 | 6.1e-5 | ~0.9 | 有 |
| FP8 E4M3 | 1/4/3 | 7 | 448 | 2^-6 | ~1.2 | 只有 NaN |
| FP4 E2M1 | 1/2/1 | 1 | 6 | 1.0 | ~0.6 | 都没有 |
| E8M0 | 0/8/0 | 127 | 2^127 | 2^-127 | 只能表示 2 的整数幂 | 只有 NaN |
4.2 FP64 和 FP32
FP64 计算太贵了，一般用于科学计算的高精度场景，在 LLM 中几乎不使用。H100 上 FP64 Tensor Core 吞吐大约是 FP16 的 1/15，消费卡上更是 FP32 的 1/64，没人拿它训练或推理 LLM。
FP32 是 baseline。到今天它仍然是几乎所有地方的累加所用的类型，也是 ggml 里中间激活和量化器内部计算用的类型，因为累加等场景我们需要足够的尾数位来保证精度。
4.3 从 FP16 到 BF16：我们要更多的指数位，还是更多的尾数位？
FP16 是第一个真正成功的窄格式，也是近十年来痛苦的源头。它有 10 位尾数，好！但是，5 位指数，坏！最大值 65504，最小规格数 6.1e-5，Transformer 训练时的梯度经常直接下溢到 0。混合精度训练里那一整套 loss scaling 机制，存在的唯一目的就是绕开 FP16 指数范围不够这件事。
BF16 是 Google 的解法，现在回头看是显然正确的：保留 FP32 的 8 位指数，砍尾数。用一半的存储换来和 FP32 完全一样的动态范围。转换几乎不要钱：
// ggml/src/ggml-impl.h:611
static inline ggml_bf16_t ggml_compute_fp32_to_bf16(float s) {
ggml_bf16_t h;
union { float f; uint32_t i; } u;
u.f = s;
if ((u.i & 0x7fffffff) > 0x7f800000) { /* nan */
h.bits = (u.i >> 16) | 64; /* force to quiet */
return h;
}
h.bits = (u.i + (0x7fff + ((u.i >> 16) & 1))) >> 16; // 就近偶数舍入，然后右移
return h;
}就这些！加一个舍入常数，右移 16 位，反向就是左移 16 位。BF16 赢下了训练。
但是，BF16 是个糟糕的推理存储格式。2.4 位十进制有效数字比 FP16 还差，而推理根本不需要那么大的指数范围。权重紧密聚集在 0 附近，你要的是尾数位。
4.4 FP8 是啥？
两个变种，OCP 标准化，Hopper/Ada/Blackwell Tensor Core 原生支持：
- E4M3：4 位指数，3 位尾数，最大 448，没有无穷（那段编码被回收来换范围），NaN 只有
S.1111.111一个。通常用于权重和激活。 - E5M2：5 位指数，2 位尾数，最大 57344，IEEE 风格，有 inf 有 NaN。通常用于梯度，因为梯度需要范围。
E4M3 在这份代码里出现的方式很有意思：它不是权重类型，而是缩放因子的类型，等下讲 NVFP4 会看到。
顺带一提，DeepSeek-V4 Flash 发布的开放权重中，非 expert 部分就是 FP8 E4M3 权重，llama.cpp 的转换脚本有个 --fp8-as-q8 开关，在转 GGUF 的路上把它变成 Q8_0（convert_hf_to_gguf.py:156）。
4.5 E8M0：一个只能表示 2 的幂的“格式”
八位指数，没有符号，没有尾数，值严格等于 2^(x - 127)。除了 2 的整数幂，它什么都表示不了。
它存在的唯一目的就是当共享缩放因子：占 1 字节，而且乘它等价于在指数字段上做整数加法。llama.cpp 的实现压根不碰 FPU：
// ggml/src/ggml-impl.h:477
static inline float ggml_e8m0_to_fp32_half(uint8_t x) {
uint32_t bits;
if (x < 2) {
bits = 0x00200000 << x; // 2^-128 / 2^-127 的次规格数位模式
} else {
bits = (uint32_t)(x - 1) << 23; // 0.5 * 2^(x-127)，即指数写成 (x-1)
}
float result;
memcpy(&result, &bits, sizeof(float));
return result;
}直接往指数字段里写位来构造一个 float，不转换，不做乘法。
4.6 FP4 又是啥？
E2M1。它一共能表示 16 个值，全在这儿：
编码： 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
数值： +0 +0.5 +1.0 +1.5 +2.0 +3.0 +4.0 +6.0 -0 -0.5 -1.0 -1.5 -2.0 -3.0 -4.0 -6.0
\___________ 非均匀，接近对数间距 ___________/这张表在 CUDA 编码器里就是个字面量：
// ggml/src/ggml-cuda/common.cuh:881
__device__ __forceinline__ uint8_t ggml_cuda_float_to_fp4_e2m1(float x, float e) {
const uint8_t sign_bit = (x < 0.0f) << 3;
float ax = fabsf(x) * e;
// Positive LUT
static constexpr float pos_lut[8] = { 0.0f, 0.5f, 1.0f, 1.5f, 2.0f, 3.0f, 4.0f, 6.0f };
...
}一个 4 比特浮点数，在任何有意义的算术层面上都 不是 数值格式。它是一张 16 项的 codebook，只不过比特排布刚好让浮点单元能用移位解码出来。FP4 这个名字是比特怎么摆的历史产物。从功能上说，E2M1 和 llama.cpp 的 IQ4_NL 是同一类东西：一张 16 项查找表。
4.7 把它们画在一根对数轴上
2^-308 2^-38 2^-14 2^-6 2^0 2^6 2^16 2^38 2^308
|-------------|----------|------|-----|-----|------|-----|-----------|
FP64 <=========================================================================>
FP32 <==============================================>
BF16 <==============================================> (范围完全一样!)
FP16 <======================>
E5M2 <======================>
E4M3 <=========>
E2M1 <=>
^
可以看到 FP4 能表达的范围超级小，这就是为什么 FP4
离开"每个 block 一个缩放因子"根本没法工作。E2M1 覆盖 0.5 到 6.0，而一个 Transformer 权重矩阵的跨度远不止于此。
于是就有了真正的那一步创新。
5. 4 比特及以下，还能用传统的 IEEE 754 浮点数吗？
如果元素格式本身没有动态范围，那就把动态范围单独拎出来，让一小组元素共享它。这就是 MX（Microscaling）、NVFP4 的全部思想。而且下面会看到，也是 llama.cpp 从 2023 年起每一个量化类型的思想。
x[i] ~= scale(block) * codebook[ q[i] ]
\___________/ \____________/
宽，共享 窄，每元素一个
(16~32 个共一个) (4 bit，从 16 项表里取)block 大小就是全部的权衡：block 小，能更好地跟住局部离群值，但缩放因子的字节开销大；block 大，便宜，但一个离群值就能毁掉旁边 31 个数。
5.1 MXFP4 是怎么做的？
OCP Microscaling 规范：32 个元素共享一个 E8M0 缩放因子。
block_mxfp4：17 字节存 32 个权重 = 4.25 bit/权重
byte: 0 1 2 3 ... 16
+----+----+----+----+---------------+----+
| e | q1 | q3 | q5 | ... | q31|
|E8M0| q0 | q2 | q4 | | q30|
+----+----+----+----+---------------+----+
^ ^
| +-- 每字节两个 4 bit 的 E2M1 编码
+-------- 共享的 2 的幂指数// ggml/src/ggml-common.h:214
#define QK_MXFP4 32
typedef struct {
uint8_t e; // E8M0
uint8_t qs[QK_MXFP4/2];
} block_mxfp4;编码器挑缩放因子的原则是：让 block 内绝对值最大的元素落到 codebook 的最高档。
// ggml/src/ggml-quants.c:350
const uint8_t e = amax > 0.0f ? (uint8_t) (floorf(log2f(amax)) - 2 + 127) : 0;
const float d = GGML_E8M0_TO_FP32_HALF(e);
for (int j = 0; j < qk/2; ++j) {
const uint8_t x0 = best_index_mxfp4(x[i*qk + 0 + j], d);
const uint8_t x1 = best_index_mxfp4(x[i*qk + qk/2 + j], d);
y[i].qs[j] = x0 | (x1 << 4);
}这一行里那个 - 2 + 127 值得拆开看，它是两件不同的事挤在一起：
e = floorf(log2f(amax)) - 2 + 127
\__________________/ \_/ \_/
L：amax 的指数 | E8M0 的指数 bias
|
E2M1 最大值 6 = 1.5 x 2^2，
所以缩放因子该取 2^(L-2)- 2 是因为 E2M1 的最大 magnitude 是 6 = 1.5 * 2^2：想让 block 内最大的那个数落到 codebook 顶端，缩放因子就得是 2^(L-2)。
+ 127 则纯粹是存储格式的要求。E8M0 只有 8 位指数，没有符号位，字段本身是个无符号整数；想表示负指数就只能整体平移。它表示的值定义为 2^(e - 127)，所以要存 2^L 就得写 e = L + 127。这跟 FP32 用 bias 127 是同一件事、同一个数字。回头看图 1 就会发现，E8M0 本来就是从 FP32 里把指数那一段单独抠出来用。
对一下解码端，确实闭环：
- 存进去
e = L - 2 + 127 - 解码是
GGML_E8M0_TO_FP32_HALF(e)=2^(e-128)=2^(L-3) kvalues_fp4是翻倍的 E2M1，最大 12- 于是最大可表示 =
12 * 2^(L-3)=1.5 * 2^L
而 amax 落在 [2^L, 2^(L+1)) 里，1.5 * 2^L 正好覆盖到它的下半段。这就是 OCP MX 规范那条定义：X = 2^(floor(log2(amax)) - emax_elem)，E2M1 的 emax_elem 就是 2。
（解码用 _half 也就是 2^(e-128) 而不是 2^(e-127)，是因为 kvalues 存的是翻倍值，那个除以 2 在这里抵掉了。至于三目运算符里 amax == 0 时写 e = 0 的那一支，是给全零 block 用的，反正所有 q 也都是 0，乘出来还是 0。）
用 bias 存储还有个顺手的好处：E8M0 字节可以直接当无符号整数比大小，缩放因子的大小顺序就是字节的大小顺序，不用解码。这也是 IEEE 当年选 bias 而不是补码的原因之一。
另外注意 best_index_mxfp4 干的事情：把 16 个码字全线性扫一遍，取绝对误差最小的那个。这是 codebook 搜索，不是舍入运算！
5.2 亲自尝试：32 个数怎么变成 17 字节
光看代码不够直观，直接调 ggml 的 ggml_quantize_chunk 和 to_float 跑一遍，把中间状态全打出来。输入这 32 个数：
0.0312 -0.1875 0.4531 0.8300 -0.0090 0.2500 -0.6250 0.1100
-0.3300 0.0700 0.5900 -0.4400 0.1500 -0.0400 0.7100 0.2100
0.0200 -0.2800 0.3600 -0.5100 0.6600 0.0900 -0.1300 0.4000
-0.7400 0.1800 0.0500 -0.3900 0.2300 0.5500 -0.0600 0.3100第一步，定标。
amax = 0.8300
L = floor(log2(0.8300)) = -1
e = L - 2 + 127 = 124 = 0x7C <- 这一个字节就是整个 block 的缩放因子
缩放 = 2^(L-2) = 0.125第二步，每个数找最近的码字。 拿 0.4531 举例：0.4531 / 0.125 = 3.62，E2M1 的码字表是 0, 0.5, 1, 1.5, 2, 3, 4, 6，离 3.62 最近的是 4，编码 6。
第三步，两个 nibble 塞进一个字节。 注意 llama.cpp 的排布不是相邻两个元素凑一起，而是 qs[j] 的低 4 位放元素 j、高 4 位放元素 j+16：
qs[3] = 0xE7
│└─ 低 4 位 = 7 -> 码字 +6.0 -> 元素 3 (0.8300)
└── 高 4 位 = E -> 码字 -4.0 -> 元素 19 (-0.5100)第四步，解码。 kvalues_fp4[nibble] * 2^(e-128)，一次查表加一次乘法。前 8 个元素的完整往返：
| i | 原值 | nibble | 码字 | 解码值 | 误差 |
|---|---|---|---|---|---|
| 0 | 0.0312 | 0 | 0.0 | 0.0000 | -0.0312 |
| 1 | -0.1875 | 11 | -1.5 | -0.1875 | 0.0000 |
| 2 | 0.4531 | 6 | 4.0 | 0.5000 | +0.0469 |
| 3 | 0.8300 | 7 | 6.0 | 0.7500 | -0.0800 |
| 4 | -0.0090 | 0 | 0.0 | 0.0000 | +0.0090 |
| 5 | 0.2500 | 4 | 2.0 | 0.2500 | 0.0000 |
| 6 | -0.6250 | 14 | -4.0 | -0.5000 | +0.1250 |
| 7 | 0.1100 | 2 | 1.0 | 0.1250 | +0.0150 |
整个 block RMSE = 0.0436，相对误差 11.13%。存储：1 + 16 = 17 字节存 32 个权重。
有三件事这张表说得比什么都清楚：
第一，i=3 那一行暴露了 MXFP4 的一个固有缺陷。 原值 0.8300 是 block 内最大的数，却被削到了 0.7500。为什么？因为缩放因子只能是 2 的幂：amax 落在 [2^-1, 2^0) 里，而最大可表示是 6 * 2^(L-2) = 0.75。只要 amax 超过 1.5 * 2^L，最大的那个元素就一定被削顶，最坏情况能削掉 25%。
这正是 NVFP4 要把缩放因子换成 E4M3 的原因：E4M3 有 3 位尾数，缩放因子不必是 2 的幂，就能贴着 amax 定标，这一刀就不用挨了。
第二，i=1 和 i=5 的误差是 0。 落在码字上的值完全无损。E2M1 的非均匀间距在零点附近密，所以小值反而准。
第三，i=6 的误差最大（0.125）。 -0.6250 卡在码字 -4.0（-0.5）和 -6.0（-0.75）正中间，怎么选都差 0.125。这就是 4 比特的分辨率极限。而这个位置恰好是权重分布里样本不少的区间，也正是 imatrix 和 IQ 系列想改善的地方。
这里有一个实现细节必须知道，不然读 kernel 会读懵：llama.cpp 把 E2M1 的值翻倍存成整数，再用减半的缩放因子补回来。
// ggml/src/ggml-common.h:1126
// e2m1 values (doubled), shared by MXFP4 and NVFP4
GGML_TABLE_BEGIN(int8_t, kvalues_fp4, 16)
0, 1, 2, 3, 4, 6, 8, 12, 0, -1, -2, -3, -4, -6, -8, -12,
GGML_TABLE_END()0, 0.5, 1, 1.5, 2, 3, 4, 6 翻倍就是 0, 1, 2, 3, 4, 6, 8, 12，全是小整数。
这就让 MXFP4 的点积可以完全在 int8 上算完，整个 block 里只在应用缩放因子时碰一次浮点。
5.3 NVFP4 又改了什么？
NVIDIA 的变体，Blackwell Tensor Core 原生支持：16 个元素共享一个 UE4M3 缩放因子。
block_nvfp4：36 字节存 64 个权重 = 4.5 bit/权重
d[0..3] 4 字节，每字节 1 个 UE4M3 缩放因子
qs[0..31] 32 字节，每字节装 2 个 4 bit 的 E2M1 编码，合计 64 个编码
d[0] d[1] d[2] d[3]
| | | |
qs[0..7] qs[8..15] qs[16..23] qs[24..31] <- 各 8 字节
元素 0-15 元素 16-31 元素 32-47 元素 48-63 <- 各 16 个权重（qs[0..31] 是数组下标 0 到 31，也就是 32 个字节；因为一个字节装两个 4 bit 编码，所以对应 64 个权重。）
这里有个容易卡住的地方：UE4M3 是 4 位指数 + 3 位尾数，明明只要 7 bit，为什么占满一个字节？
确实只用了 7 bit。编码器 ggml_fp32_to_ue4m3 的输出最大是 0x7E，从不置最高位；解码器 ggml_ue4m3_to_fp32 取指数时写的是 (x >> 3) & 0xF，也把最高位屏蔽掉了。第 8 位就是填充。
顺带对比一下：MXFP4 的 E8M0 是 8 位指数、无符号无尾数，那一个字节是满的，一位不浪费。NVFP4 用 UE4M3 换来的是缩放因子不必是 2 的幂（下面会看到这有多值钱），代价就是每 16 个权重白扔 1 个 bit。
sub-block 内部的 nibble 排布和 MXFP4 同构，只是窗口从 32 缩到了 16。第 s 个 sub-block 占 qs[s*8 .. s*8+7] 这 8 个字节：
qs[s*8 + j] 低 4 位 -> 元素 s*16 + j
qs[s*8 + j] 高 4 位 -> 元素 s*16 + j + 8 (j = 0..7)
解码：kvalues_fp4[nibble] * ue4m3_to_fp32(d[s])// ggml/src/ggml-common.h:220
#define QK_NVFP4 64
#define QK_NVFP4_SUB 16 // sub-block size for per-group scales
typedef struct {
uint8_t d[QK_NVFP4/QK_NVFP4_SUB]; // UE4M3 scales, one per 16-element sub-block
uint8_t qs[QK_NVFP4/2]; // packed 4-bit E2M1 values
} block_nvfp4;和 MXFP4 有两处差别，都是有意的：
- block 大小 16，不是 32。 离群值的影响范围被缩小了一半。
- 缩放因子是 E4M3，不是 E8M0。 E8M0 只能表示 2 的幂，所以缩放因子本身被量化到了一个“倍数为 2”的粗网格上，codebook 顶端最多可能浪费掉一半量程；E4M3 有 3 位尾数，缩放因子本身能落在理想值 6% 以内。
代价是每 16 个元素 8 位缩放因子 = 0.5 bpw，而 MXFP4 是每 32 个 8 位 = 0.25 bpw。所以 NVFP4 是 4.5 bpw，MXFP4 是 4.25 bpw，前者精度明显更好。
// ggml/src/ggml-quants.c:405
// UE4M3 scale: amax / 6.0 maps the max E2M1 value (6.0) to amax
const uint8_t ue = ggml_fp32_to_ue4m3(amax / 6.0f);5.4 同一组数，NVFP4 跑一遍
把 MXFP4 那节用的 32 个数原样拿过来（再补 32 个小一个量级的凑满 64），跑 NVFP4：
=== 4 个 sub-block 各自的 UE4M3 缩放因子 ===
d[0]=0x21 sub-block 元素 0-15 amax=0.8300
d[1]=0x20 sub-block 元素 16-31 amax=0.7400
d[2]=0x0B sub-block 元素 32-47 amax=0.1245
d[3]=0x09 sub-block 元素 48-63 amax=0.1110第一件事就看出来了：四个 sub-block 拿到了四个完全不同的缩放因子。后两个 sub-block 小了一个量级，d 也跟着降到 0x0B、0x09。要是按 MXFP4 那样 32 个共用一个，后半段的精度就被前半段的大数拖垮了。这就是 block 大小从 32 减到 16 的意义。
再看字节排布，和上面推的完全对得上：
sub-block 0 -> qs[ 0.. 7]: C0 1B 65 D7 20 94 7E 32
qs[ 0]=0xC0 低4位= 0-> 0.0 (元素 0: +0.0312 -> +0.0000) 高4位=12-> -2.0 (元素 8: -0.3300 -> -0.2812)
qs[ 1]=0x1B 低4位=11-> -1.5 (元素 1: -0.1875 -> -0.2109) 高4位= 1-> 0.5 (元素 9: +0.0700 -> +0.0703)
qs[ 2]=0x65 低4位= 5-> 3.0 (元素 2: +0.4531 -> +0.4219) 高4位= 6-> 4.0 (元素10: +0.5900 -> +0.5625)最关键的是元素 3，就是那个被 MXFP4 削顶的 0.8300：
| 缩放因子 | 最大可表示 | 0.8300 解码成 | 误差 | |
|---|---|---|---|---|
| MXFP4 | 2^-3 = 0.125（只能是 2 的幂） |
0.7500 | 0.7500 | -0.0800 |
| NVFP4 | 0.140625（UE4M3，带尾数） | 0.8438 | 0.8438 | +0.0138 |
MXFP4 因为缩放因子必须是 2 的幂，够不着 0.83，只能削；NVFP4 的 UE4M3 有 3 位尾数，amax/6 = 0.1383 能round到 0.140625，于是留出了余量，一刀都不用挨。
整组下来，同样这 32 个数：MXFP4 相对误差 11.13%，NVFP4 10.01%。差距不算悬殊，但别忘了这只是一个小样本，而且真实收益主要来自前面那条：block 小了，离群值的影响范围也小了。代价是 4.5 bpw vs 4.25 bpw。
有个坑要提醒一下：如果你拿 NVIDIA 的规范来对照，会发现参考版 NVFP4 比这里多一级。它的完整层级是这样的：
tensor：整个权重矩阵，比如 4096 x 14336，约 5900 万个数
|
+-- 1 个 FP32 全局缩放 <- llama.cpp 没有这一级
|
+-- sub-block：16 个数
|
+-- 1 个 E4M3 缩放
|
+-- 元素：1 个数 -> 1 个 E2M1 码注意 tensor 和 block 差着六个数量级，别把两者混为一谈。
为什么要多这一级？因为 E4M3 自己的量程有限。 它的最小规格数是 2^-6 = 0.0156，低于这个值就掉进次规格区，精度断崖式下跌。拿 llama.cpp 的 ggml_fp32_to_ue4m3 实测一下就很清楚：
| block 内 amax | 理想 d = amax/6 | 实际编码出的 d | 相对误差 |
|---|---|---|---|
| 2.0 | 0.3333 | 0.3438 | +3.1% |
| 0.83 | 0.1383 | 0.1406 | +1.7% |
| 0.094 | 0.01567 | 0.01563 | -0.3% |
| 0.05 | 0.00833 | 0.00781 | -6.3%（次规格） |
| 0.02 | 0.00333 | 0.00391 | +17.2%（次规格） |
| 0.004 | 0.00067 | 0 | -100%，整个 block 归零 |
amax 掉到 0.094 以下就开始失真，掉到 0.004 以下整个 sub-block 直接变成零。而 LLM 权重的 block 内 amax 落在 0.01 到 0.1 这个区间是相当常见的。
全局 FP32 缩放干的就是这件事：先把整个 tensor 搬到 E4M3 好用的区间，保证每个 per-16 的缩放因子都落在规格数范围内。NVIDIA 的 solution 大致是：
s_global = amax_tensor / (448 * 6) // 448 是 E4M3 上限，6 是 E2M1 上限
d_block = amax_block / (6 * s_global) // 编成 E4M3，此时必然落在好区间
重建：x ~= s_global * d_block * e2m1_code代价呢？每个 tensor 32 bit，对 5900 万权重来说是 32 / 58720256 ≈ 0.0000005 bpw，完全可以忽略。
llama.cpp 的 block_nvfp4 只实现了 per-16 这一级。bpw 一样是 4.5，但这不是等价简化。对整体幅度偏小的 tensor，per-16 的缩放因子会退化到次规格区。
6. 凭什么相邻的数能共用一个缩放因子？
写到这里应该有人要问了：MXFP4 让 32 个相邻的权重共用一个缩放因子，NVFP4 是 16 个。凭什么？相邻的权重就一定尺度相近吗？这件事有理论或者实验支撑吗？
这个问题其实是两个问题，而且它们的答案来自完全不同的地方：
- 凭什么能共享：这是统计问题。
- 凭什么是“相邻”：这是计算问题。
6.1 为什么 block 必须沿 K 方向？
矩阵乘是 y_j = Σ_k w_jk · x_k。如果缩放因子在 k 方向的一段连续区间内是常数，它就能被提到求和外面：
Σ_{k∈block} w_jk · x_k = d_block · Σ_{k∈block} q_jk · x_k
└─ 整数 MAC，一个 block 只碰一次浮点这正是第 9 节那段 NEON 代码在干的事：4 条 TBL 加 SDOT 全走整数，浮点只在 block 边界出现一次。如果缩放因子不是沿归约维度分段常数，这个提取就不成立，你得每个元素乘一次浮点，4 bit 也就白量化了。
所以 block 沿 K 连续，是算术约束，跟“相邻权重更像”没有关系。
6.2 可是转置了怎么办？
这里马上会冒出一个疑问：矩阵有时当左乘、有时当右乘，转置 K 就换了个轴，方向定死不就出问题了吗？
关键前提是：推理永远不转置权重。 训练必须用 W^T 做反向传播，所以任何沿 K 分 block 的格式在训练里都会立刻撞上这个问题；但推理只做 y = W x，W 的用法在量化那一刻就定死了、终生不变。
这就是推理专用格式敢把方向固定的底气。而 ggml 更进一步，它要求两个操作数都以 K 为连续维：
// ggml/src/ggml.c:3270
static inline bool ggml_can_mul_mat(const struct ggml_tensor * t0, const struct ggml_tensor * t1) {
return (t0->ne[0] == t1->ne[0]) && ... // 两边的 ne[0] 都是 K
}
...
GGML_ASSERT(!ggml_is_transposed(a)); // 权重不许转置所以 ggml 的 mul_mat 本质上是 BLAS 里的 A^T B 形态：权重离线沿 K 分 block，激活在线量化成 Q8_0/Q8_1 也沿同一个 K，两边天然对齐，各自的缩放因子都能提到求和外面。压根不存在“有时当左矩阵有时当右矩阵”的情况。
但确实有一个地方会被反着用：attention 的 V。 scores @ V 需要沿另一个方向读它，所以 llama.cpp 的非 FA 路径干脆把 V 转置着存：
// src/llama-kv-cache.h:238
bool v_trans = true; // the value tensor is transposed而分 block 量化的 tensor 是没法转置的，转置会把 block 打散。于是就有了那条硬性规定：量化 V 必须开 flash attention，因为 FA 按 V 的自然布局消费它。llama.cpp 在这里的选择是限制用法，而不是改格式。
6.3 那 tile 呢？格式是考虑了的
既然真实的矩阵乘都是切成 tile 做的，格式有没有管这件事？管了，而且是明着写进指令的。
llama.cpp 这边，MMQ 的 tile 在 K 方向的跨度是 block 大小的整数倍：
// ggml/src/ggml-cuda/mmq.cuh:113
// In other words, the size of the quantized data in the K dimension is a multiple of MMQ_TILE_NE_K.
#define MMQ_TILE_NE_K 32但最有说服力的证据，直接写在 Blackwell 那两条 block scaling MMA 指令的编码里：
mma.sync.aligned.kind::mxf4 .block_scale.scale_vec::2X.m16n8k64...ue8m0
mma.sync.aligned.kind::mxf4nvf4.block_scale.scale_vec::4X.m16n8k64...ue4m3m16n8k64 表示每条 MMA 沿 K 吃 64 个元素。于是：
- MXFP4 block 大小 32 ->
64 / 32 = 2个缩放因子 ->scale_vec::2X - NVFP4 block 大小 16 ->
64 / 16 = 4个缩放因子 ->scale_vec::4X
“一条 MMA 要用几个缩放因子”是指令编码里的一个字段。 block 大小选 32 和 16 不是拍脑袋，是照着 tensor 核心的 K 跨度设计的。llama.cpp 的注释也印证了打包方式：
// ggml/src/ggml-cuda/mmq.cuh:49
// mxfp4 has block size 32, each int32 of d4 contains 2 e8m0 scales in the lower 16 bits
// nvfp4 has block size 16, each int32 of d4 contains 4 ue4m3 scales一个 int32 寄存器，正好装下一条 MMA 需要的全部缩放因子。
所以完整的分工是这样的：
| 层次 | 谁说了算 | 能不能改 |
|---|---|---|
| block 沿 K、block 大小 16/32/256 | 文件格式，量化时定死 | 不能，改了要重新量化 |
| tile 形状、交织顺序、SRAM 布局 | 后端，加载时 repack | 能，纯排列，无损 |
这两层能解耦，前提是 tile 的 K 跨度必须是 block 大小的整数倍，否则一个 block 会被劈到两个 tile 里，那才是真麻烦。硬件厂商选 k64 配 block 32/16，软件选 MMQ_TILE_NE_K 32，都是在维护这个整除关系。
6.4 block 大小到底在防什么？
回到第一个问题：凭什么相邻的数能共用一个缩放因子？
直觉上的答案是“相邻权重尺度相近”，但这个说法其实站不住。LLM 权重在一行之内相当接近高斯分布，并没有“相邻的更像”这回事。随便挑 32 个还是随便挑 1024 个，分布是一样的。
真正的答案是：分 block 缩放防的是离群值。
设想一个 block 里混进一个比邻居大几十倍的数。absmax 定标会让缩放因子被这一个数拉满，剩下 31 个全部挤到 codebook 最底下的一两档，等于集体报废。block 越大，被一个离群值拖下水的邻居就越多；block 越小，损失就被关在一个更小的笼子里。
所以 block 大小买到的不是“更贴合的尺度”，而是离群值的影响半径。这也顺带解释了第 10 节那个 Hadamard 旋转为什么有效：它把分布强行揉回高斯，等于让离群值消失，于是 block 大小也就不那么重要了。
6.5 为什么 E2M1 对 block 大小不敏感，INT4 却很敏感？
这里还有一层，解释了一个常让人困惑的现象：为什么浮点 codebook 敢用 block 32，而 GPTQ、AWQ 那类 INT 方案通常要 group size 128 甚至更小？
关键在codebook 的间距是不是均匀的。
E2M1 的码字是 0, 0.5, 1, 1.5, 2, 3, 4, 6，接近对数间距，相对精度在整个量程上大致恒定。这种 codebook 本身就近似尺度不变：block 内数值整体大一点小一点，相对误差不怎么变，所以对定标是否精确没那么挑剔。
均匀间距的 INT4 没有这个性质。它的绝对精度是固定的，一旦缩放因子偏离最优，小数值那一头的相对误差立刻恶化。所以它必须靠更小的 group 把尺度框准。
不是 INT 方案更讲究，是它们的 codebook 更需要。
6.6 所以 MXFP4 到 NVFP4，必须同时改两件事
理解了上面两点，就能看出 NVIDIA 那两处改动为什么是绑在一起的。
只把 block 从 32 缩到 16、缩放因子还是 E8M0，其实是负收益。因为 2 的幂缩放每个 block 都要付一次“够不着 amax”的浪费，就是前面演示里 0.83 被削到 0.75 那一刀，最坏能削掉 25%。block 越小，付这个代价的次数就越多。
必须同时把缩放因子换成带尾数的 E4M3，让它能贴着 amax 定标，那一刀才不用挨（同一个 0.83，NVFP4 给出的是 0.8438）。这时候 block 变小才真正兑现成精度。
所以 32 到 16 和 E8M0 到 E4M3 是一对，拆开任何一个都不成立。
6.7 那为什么不切成 tile？
既然是二维矩阵，为什么不按 tile 切，让一个 4×8 的小 block 共享一个缩放因子？同样是 32 个元素，为什么非要 1×32 地沿 K 排？
因为权重矩阵里，不同的行整体大小本来就不一样。
权重矩阵的一行就是一个输出神经元的全部输入权重。有的神经元这一行普遍在 0.5 上下，有的普遍在 0.01 上下，差几十倍很常见。
沿 K 切，一个 block 装的是同一行里连续的一段，也就是同一个神经元的一段输入权重，这组数天然是一个量级；一旦跨行，就把几个量级不同的神经元混进了同一个 scale。假如相邻两行差 5 倍，那么跨这两行的 tile 里，小的那一行就等价于遇到了一个 5 倍的离群值。
换个说法：跨行本质上是在人为制造离群值。绕了一圈，又回到上一节那个结论。
反过来说，如果所有行的量级恰好相同，tile 和 1D block 就没有区别了，代价完全来自行与行之间的大小差异，不来自 tile 这个形状本身。
6.8 llama.cpp 真的试过把 tile 放进文件格式
这段历史在类型表里还留着删除的痕迹：
// ggml/src/ggml.c:893
[31] = { .type_name = "TYPE_Q4_0_4_4 REMOVED, use Q4_0 with runtime repacking" },
[33] = { .type_name = "TYPE_Q4_0_8_8 REMOVED, use Q4_0 with runtime repacking" },
// IQ4_NL_4_4 / IQ4_NL_4_8 / IQ4_NL_8_8 同样被删Q4_0_4_4、Q4_0_8_8 就是交织好的 tile 布局，为不同 SIMD 宽度各做一份。问题很直接：一个 GGUF 只能喂一种 CPU。 所以后来全删了，改成加载时 repack。
而这里有个关键性质，它让布局和精度彻底解耦：
// ggml/src/ggml-cpu/repack.cpp:2742
static block_q4_0x4 make_block_q4_0x4(block_q4_0 * in, int blck_size_interleave) {
for (int i = 0; i < 4; i++) out.d[i] = in[i].d; // 缩放因子原样搬
...
memcpy(&elems, &in[src_id].qs[src_offset], sizeof(uint64_t));
elems ^= xor_mask; // 0x8888... 只翻 nibble 的第 3 位
memcpy(&out.qs[dst_offset], &elems, sizeof(uint64_t));输入是已经量化好的 block_q4_0：d 逐个拷贝，量化值 memcpy 搬运，那个 ^= 0x8888... 是个可逆位变换（把无符号 nibble 变成便于 SIMD 隐式减 8 的形式），不改数值。
repack 是纯排列，不重新量化。 所以精度在量化那一刻就定死了，之后所有布局变换都是无损的。CUDA 的 MMQ SRAM layout、CPU 的 x4/x8 交织，都只是在搬字节。这也是为什么一个 GGUF 能跑遍 CPU、CUDA、Metal、SYCL。
顺带，这条规则还解释了后面两件事：
- KV cache 那条
head_dim % blck_size == 0，就是“block 不能跨 head”，因为不同 head 的尺度差异极大，正是上表里“跨行”的情形。 - 量化 V 必须开 flash attention，根源也在这里：
V^T要沿另一个方向读，而 block 是沿 K 排的，转置读等于跨 block，不解码就取不出来。
6.9 这方面的文献
上面这些说法都不是我拍脑袋想的，MX 格式的设计本身就有公开的 ablation 数据：
- Microscaling Data Formats for Deep Learning（Rouhani 等）：MX 标准的原始论文，block 大小和缩放格式的选择就是在这里做的 ablation，并在二十多个 benchmark 上验证了 MX 可以直接替换 FP32。
- Is Finer Better? The Limits of Microscaling Formats in Large Language Models（Fasoli 等）：直接追问“block 更细是不是更好”，结论是并不一致地更好，收益递减很明显。
- AMXFP4: Taming Activation Outliers with Asymmetric Microscaling Floating-Point：专门对付离群值，佐证了“分 block 是为了关住 outlier”这个定位。
- MX+: Pushing the Limits of Microscaling Formats for LLM Serving
离群值那条线的经典工作：LLM.int8()（Dettmers 等，发现激活的 emergent outlier features）、SmoothQuant、AWQ，以及 QuIP# 和 QuaRot。后两者用随机正交/Hadamard 旋转做 incoherence processing，正是 llama.cpp KV cache 那套做法的理论来源，第 10 节会给出完整出处。
7. llama.cpp 里到底有多少种量化？
7.1 通用结构长什么样？
ggml 里每一个量化类型都是一个block（block）：固定数量的权重加上它们的元数据，放在一个连续的 struct 里。ggml_type_traits（ggml/src/ggml.c:668）给每个类型定义了 blck_size 和 type_size，每权重比特数就是 8 * type_size / blck_size。
+------------------+--------------------------------------------------+
| scale(s) | packed quantized codes |
| fp16 / e8m0 / | 4 / 3 / 2 / 1.x bit each |
| 6-bit / 8-bit | |
+------------------+--------------------------------------------------+
^ 被整个 block 摊薄的元数据 ^ 真正的载荷
bpw = 8 * sizeof(block) / block_size7.2 Q4_0：最朴素的做法
先用一句话说清楚 Q4_0 在干什么：
把连续 32 个权重分成一组，在组内找出绝对值最大的那个数，用它定出一个公共的缩放因子；然后把每个数除以这个因子，四舍五入到 16 个等距档位里最近的一个，只存档位编号。
所以每组存下来的东西就两样：
- 1 个 fp16 缩放因子，2 字节
- 32 个 4 比特的档位编号，16 字节
合计 18 字节装 32 个权重，也就是 4.5 bpw。
// ggml/src/ggml-common.h:194
#define QK4_0 32
typedef struct {
ggml_half d; // delta
uint8_t qs[QK4_0 / 2]; // nibbles / quants
} block_q4_0; // 18 字节 / 32 权重 = 4.5 bpw编码器全文就这么长：
// ggml/src/ggml-quants.c:113
float amax = 0.0f, max = 0.0f;
for (int j = 0; j < qk; j++) {
const float v = x[i*qk + j];
if (amax < fabsf(v)) { amax = fabsf(v); max = v; }
}
const float d = max / -8;
const float id = d ? 1.0f/d : 0.0f;
...
const uint8_t xi0 = MIN(15, (int8_t)(x0 + 8.5f));反量化（dequantize_row_q4_0，ggml-quants.c:459）就更简单了，整个公式就一行：
x = (q - 8) * d减 8 是把无符号的档位编号搬回 0 两侧：编号 0 到 15 减去 8 之后变成 -8 到 +7，编号 8 正好对应 0。然后乘上缩放因子还原量级。每个权重只要一次减法加一次乘法。
字节排布和 MXFP4 一样，也不是相邻两个元素凑一起，而是 qs[j] 的低 4 位放元素 j、高 4 位放元素 j+16。
这里有个细节值得留意，第 9 节会用到：Q4_0 的反量化不需要查表。因为 16 个档位是等距的，(q-8)*d 现算就行，codebook 是被公式隐含表达的。等到 codebook 变得不等距（比如 IQ4_NL 或者 E2M1），这一步就只能老老实实查表了。
拿前面 MXFP4 演示用的同一组 32 个数跑一遍，amax = 0.8300，d = max/-8 = -0.10376：
| i | 原值 | 码字 | 反量化 | 误差 |
|---|---|---|---|---|
| 0 | 0.0312 | 8 | 0.0000 | -0.0312 |
| 1 | -0.1875 | 10 | -0.2075 | -0.0200 |
| 2 | 0.4531 | 4 | 0.4150 | -0.0381 |
| 3 | 0.8300 | 0 | 0.8301 | +0.0001 |
| 4 | -0.0090 | 8 | 0.0000 | +0.0090 |
| 5 | 0.2500 | 6 | 0.2075 | -0.0425 |
| 6 | -0.6250 | 14 | -0.6226 | +0.0024 |
| 7 | 0.1100 | 7 | 0.1038 | -0.0062 |
三件事值得注意：
- i=3 那个 amax 几乎无损（0.8300 到 0.8301），因为定标时它就被钉在端点上。
- 小数被压没了：0.0312 和 -0.0090 都归零，因为最小档距就是 0.1038，比它们还大。
- 整组相对误差 7.25%。顺带一提，同一组数 MXFP4 是 11.13%。Q4_0 的缩放因子是自由的 fp16，不像 E8M0 只能取 2 的幂，也就不用挨削顶那一刀。当然这只是一个 32 数的样本，别当成普遍结论。
而 Q4_0 有两个硬伤，后面所有的改进都是冲着它们去的：
硬伤一，它假设权重关于 0 对称。 整组只有一个缩放因子、没有偏移量，所以 16 个档位必须对称地铺在 0 两侧。如果这组权重整体偏向正数，负的那半边档位就白白浪费了。
硬伤二，元数据不便宜。 每 32 个权重要花 16 比特存那个 fp16 缩放因子，折合 0.5 bpw，在 4.5 bpw 的总预算里占了 11%。
7.3 Q4_1 修了第一个硬伤，但代价太大
Q4_1 的做法很直接：再存一个每组的最小值，重建公式从 x = d*q 变成 x = d*q + m。这样档位就能整体平移到数据实际所在的区间，不再假设对称。
代价是每 32 个权重的元数据从 16 比特翻倍到 32 比特，bpw 从 4.5 涨到 5.0。为了修一个硬伤多付 11% 体积，不太划算。
（顺带交代一下这一家子的其他成员：Q5_0/Q5_1 把第 5 位溢出到单独的 qh 字段里；Q8_0 是 8 比特，同时也是激活侧的主力格式。当你要把 Q4_K 权重乘上激活时，llama.cpp 会先把激活量化成 Q8_0/Q8_1，好让内层循环走整数。）
7.4 K-quants：把两件事组合起来
K-quants 的想法很妙：既然元数据这么贵，那就把元数据本身也量化掉。
具体做法是加一层间接。scale 和 min 不再各用一个 fp16 直接存，而是各压成 6 比特的整数；这些 6 比特的值再由每 256 个权重共享的两个 fp16 super-block 缩放因子还原回真实数值。
block_q4_K：一个 256 权重的 super-block
super-block (256 权重, 144 字节, 4.5 bpw)
+--------+--------+--------------------+-----------------------------+
| d | dmin | scales[12] | qs[128] |
| fp16 | fp16 | 8 x (6b scale + | 256 x 4-bit codes |
| | | 6b min) | |
+--------+--------+--------------------+-----------------------------+
| | |
| | +-- 8 个 sub-block，每 sub-block 32 个权重
| +--- 所有 min 的super-block 缩放因子
+------------- 所有 scale 的super-block 缩放因子
第 j 个 sub-block 的重建： x = (d * scale_j) * q - (dmin * min_j)
\_________/ \__________/
有效的 d 有效的 m算一下这笔账。以每 32 个权重为单位：
| 每 32 权重的元数据 | 有偏移量吗 | bpw | |
|---|---|---|---|
| Q4_0 | fp16 scale = 16 bit | 没有 | 4.5 |
| Q4_1 | fp16 scale + fp16 min = 32 bit | 有 | 5.0 |
| Q4_K | 6b scale + 6b min = 12 bit，加上super-block 缩放因子摊下来的 4 bit = 16 bit | 有 | 4.5 |
这就是 K-quants 的全部意义：同样 16 比特的元数据预算，Q4_0 只买到一个缩放因子，Q4_K 买到了缩放因子加偏移量。 等于白拿了 Q4_1 的能力，体积却和 Q4_0 一样。
代价有两个：一是多了一层间接，反量化时要先还原 scale 和 min，稍微复杂一点；二是 block 从 32 变成 256 的 super-block，最小处理单元变大了。
而这还只是 K-quants 的一半。另一半藏在编码器里。
7.5 大家忽略的地方：K-quants 做的是搜索，不是舍入
所谓“高级量化算法”其实就藏在这里。这里根本没用就近舍入！
make_qkx2_quants（以及它带 imatrix 的兄弟 make_qkx3_quants，Q4_K/Q5_K 实际调用的是后者）用加权最小二乘去拟合仿射模型 x ~= scale*L + min，然后扫一遍候选缩放因子找更优解：
// ggml/src/ggml-quants.c:799 (节选)
for (int is = 0; is <= nstep; ++is) { // Q4_K 传入 nstep = 36
iscale = (rmin + rdelta*is + nmax)/(max - min); // 候选缩放因子
float sum_l = 0, sum_l2 = 0, sum_xl = 0;
for (int i = 0; i < n; ++i) {
int l = nearest_int(iscale*(x[i] - min));
l = MAX(0, MIN(nmax, l));
Laux[i] = l;
float w = weights[i]; // <-- 重要性权重
sum_l += w*l; sum_l2 += w*l*l; sum_xl += w*l*x[i];
}
float D = sum_w * sum_l2 - sum_l * sum_l;
if (D > 0) {
// 给定这组码字，(scale, min) 的加权最小二乘闭式解
float this_scale = (sum_w * sum_xl - sum_x * sum_l)/D;
float this_min = (sum_l2 * sum_x - sum_l * sum_xl)/D;
...
if (cur_error < best_error) { /* 留下它 */ }
}
}这是标准的交替优化：猜一个缩放因子 -> 分配码字 -> 在这组码字下反解最优的 (scale, min) -> 量加权误差 -> 留最好的。
Q4_K 传 nstep = 36，也就是每 32 个权重要试 37 个候选缩放因子（ggml-quants.c:1583）。这就是为什么量化一个 70B 模型要几分钟而不是几秒，而且这几分钟花得非常值。
对称版本 make_qx_quants（ggml-quants.c:628）做的是同一件事，is 从 -9 扫到 9，最大化 sumlx^2 / suml2，也就是 x 在重建向量上的投影。对于无偏移的纯缩放拟合，这才是正确的目标函数。
整条离线流水线长这样：
flowchart TD
A["fp16/bf16 safetensors"] --> B["convert_hf_to_gguf.py<br/>输出 fp16/bf16 GGUF"]
B --> C{"有 imatrix 吗？"}
C -->|"没有"| D["启发式权重<br/>w_i = 平均绝对值 + 该元素绝对值"]
C -->|"有"| E["llama-imatrix 跑一遍语料<br/>acc_j += 激活_j 的平方"]
E --> F["w_i = imatrix_i * sqrt(sigma2 + x_i^2)"]
D --> G["llama_tensor_get_type_impl<br/>逐 tensor 分配比特 budget"]
F --> G
G --> H["make_qkx3_quants / make_qx_quants<br/>扫缩放因子 + 加权最小二乘"]
H --> I["打包进 block struct"]
I --> J["量化后的 GGUF"]
style E fill:#2d5a3d,stroke:#4a8,color:#fff
style H fill:#3d3d6b,stroke:#88a,color:#fff7.6 imatrix 是啥？
上面代码里的 weights[i] 不是常数。有 imatrix 的时候它是这样算的：
// ggml/src/ggml-quants.c:1574 (quantize_row_q4_K_impl)
float sigma2 = 2*sum_x2/QK_K; // super-block 内 256 个权重的均方，再乘 2
float av_x = sqrtf(sigma2);
...
if (quant_weights) {
const float * qw = quant_weights + QK_K*i + 32*j;
for (int l = 0; l < 32; ++l)
weights[l] = qw[l] * sqrtf(sigma2 + x[32*j + l]*x[32*j + l]);
} else {
for (int l = 0; l < 32; ++l)
weights[l] = av_x + fabsf(x[32*j + l]);
}这个式子是两个因子相乘，它们衡量的是完全不同的两件事。拆开看。
因子一：qw[l]，这个输入通道有多重要
qw 就是 imatrix，它是真的把模型跑一遍收集来的：
// tools/imatrix/imatrix.cpp
for (int64_t j = 0; j < ne0; ++j) {
acc[j] += x[j] * x[j]; // x 是进入这个矩阵乘法的激活
}累加之后除以样本数，得到的就是每个输入通道的激活平方均值，也就是 E[x_j^2]。
为什么是平方，而不是绝对值？这可以直接推出来。矩阵乘是 y_i = Σ_j w_ij · x_j，如果把权重 w_ij 扰动 dw_ij，输出的变化就是：
dy_i = dw_ij · x_j在很多 token 上取平方期望：
E[dy_i^2] = dw_ij^2 · E[x_j^2]
\_______/
这就是 imatrix所以要最小化输出误差的平方，加权最小二乘里的权重就该取 E[x_j^2]。这不是拍脑袋的启发式，是从目标函数直接落下来的。
结论也就很直白：乘在大激活上的权重重要，输入通道常年接近 0 的权重不重要。把量化误差花在激活小的地方。
这里有个容易误解的点：imatrix 不是每个权重一个值，而是每个输入通道一个值。 它是一个长度等于输入维度的向量，整个 tensor 的所有输出行共用同一份。源码里能看出来：逐行量化的循环里，src 每次前进一行，但 quant_weights 原地不动：
// ggml/src/ggml-quants.c:1631
for (int64_t row = 0; row < nrow; ++row) {
quantize_row_q4_K_impl(src, (block_q4_K*)qrow, n_per_row, quant_weights);
src += n_per_row; // 权重往下走一行
qrow += row_size;
// quant_weights 没有推进，所有行复用同一份
}因子二：sqrt(sigma2 + x_l^2)，这个权重本身有多大
第二个因子跟激活无关，只看权重自己。sigma2 是这个 256 权重 super-block 的均方（乘了个 2），所以：
- 当
|x_l|远小于sqrt(sigma2)时，整项 ≈sqrt(sigma2)，是个常数，所有小权重一视同仁 - 当
|x_l|远大于sqrt(sigma2)时，整项 ≈|x_l|：大权重按 magnitude 加权
也就是一个 floor 加上 magnitude 的软过渡，sqrt(sigma2) 就是那个拐点。
为什么要这一项？看没有 imatrix 的那个分支就明白了：weights[l] = av_x + fabsf(x[l])，也就是 sqrt(sigma2) + |x_l|：同样是“floor 加 magnitude”的形状，只是用加法而不是平方和开根。
所以这一项是 llama.cpp 自己的默认启发式，imatrix 是乘上去的，不是取代它。
两半各自的道理：
- magnitude 项：同样大小的绝对误差，发生在大权重上更亏，因为它对输出的贡献本来就大。
- floor 项：防止接近 0 的权重被完全忽略。如果权重是纯粹的
|x_l|，那么x_l ≈ 0的位置权重也≈0，拟合会完全不管它们，可能把它们舍到很离谱的档位上去。
合起来
最终权重 = 这个输入通道有多重要 x 这个权重本身有多大
\________________/ \_______________/
来自校准语料的激活 来自权重自身的分布两个维度的信息相乘，缺一不可。
这和 GPTQ、AWQ 属于同一族思路，只不过这里实现成最小二乘拟合里的一个权重，而不是基于 Hessian 的误差反馈。代价是在几十万 token 上做一次前向，精度收益基本是白捡的。
而 IQ1/IQ2/IQ3 这几个类型，干脆拒绝在没有 imatrix 的情况下量化（tensor_requires_imatrix，src/llama-quant.cpp:785），而且 imatrix 必须跑一遍语料收集激活平方均值，这也是不同发布者的相同量化方式的 GGUF 质量差异的根源。
7.7 为什么 tensor 之间不平等？
llama_tensor_get_type_impl（src/llama-quant.cpp:424）是 300 行攒出来的经验之谈，讲的是哪些 tensor 值得多给几位。摘几段：
// attn_v 很小，但影响很大，尤其在 GQA 下
if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K) {
new_type = qs.model.hparams.n_gqa() >= 4 ? GGML_TYPE_Q4_K : GGML_TYPE_Q3_K;
}
// 输出层永远不低于 Q5_K/Q6_K
else if (new_type != GGML_TYPE_Q8_0) {
new_type = GGML_TYPE_Q6_K;
}
// 前 1/8 的 ffn_down 层升级
if (i_layer < n_layer/8) new_type = GGML_TYPE_Q4_K;旁边还留了一句注释，直接告诉你这个结论是怎么来的：
// In the 70B model we have 8 heads sharing the same attn_v weights. As a result, the
// attn_v.weight tensor is 8x smaller compared to attn_q.weight. Hence, we can get a
// nice boost in quantization accuracy at a very small cost.这就是文件名里 _S / _M / _L 后缀的含义。Q4_K_M 不是“Q4_K”，它是一套 solution：主体 Q4_K，attn_v 和部分 ffn_down 升到 Q6_K，输出 tensor Q6_K。
回到开头那个文件列表，unsloth 的 UD- 前缀（Unsloth Dynamic）也是同一回事：底层的 ggml block 格式就是本文这些，不一样的是逐 tensor 的分配策略和他们自己的校准语料。所以同名的 IQ2_M 和 UD-IQ2_M，格式相同，质量可以差不少。
7.8 表：代码树里的全部类型
由 sizeof(block) 按 QK_K = 256 算出，并和 ggml-common.h 里的 static_assert 核对过。
| 类型 | block 大小 | 字节/block | bpw | 结构 |
|---|---|---|---|---|
| F32 | 1 | 4 | 32.00 | |
| F16 / BF16 | 1 | 2 | 16.00 | |
| Q8_0 | 32 | 34 | 8.50 | fp16 缩放 + int8 |
| Q6_K | 256 | 210 | 6.56 | 8 比特 sub-block 缩放 |
| Q5_1 | 32 | 24 | 6.00 | scale + min |
| Q5_K | 256 | 176 | 5.50 | 两级，6b scale+min |
| Q5_0 | 32 | 22 | 5.50 | |
| Q4_1 | 32 | 20 | 5.00 | scale + min |
| Q4_K | 256 | 144 | 4.50 | 两级，6b scale+min |
| Q4_0 | 32 | 18 | 4.50 | absmax |
| IQ4_NL | 32 | 18 | 4.50 | 16 项非均匀 codebook |
| NVFP4 | 64 | 36 | 4.50 | E2M1 + 每 16 个一个 UE4M3 |
| IQ4_XS | 256 | 136 | 4.25 | codebook + 6b sub-block 缩放 |
| MXFP4 | 32 | 17 | 4.25 | E2M1 + 每 32 个一个 E8M0 |
| Q3_K | 256 | 110 | 3.44 | 6 比特缩放 |
| IQ3_S | 256 | 110 | 3.44 | 512 项 4 维 codebook |
| IQ3_XXS | 256 | 98 | 3.06 | 256 项 4 维 codebook |
| Q2_K | 256 | 84 | 2.62 | 4b scale+min |
| IQ2_S | 256 | 82 | 2.56 | 1024 项 8 维 codebook |
| IQ2_XS | 256 | 74 | 2.31 | 512 项 8 维 codebook |
| IQ2_XXS | 256 | 66 | 2.06 | 256 项 8 维 codebook |
| TQ2_0 | 256 | 66 | 2.06 | 三值，2b/权重 |
| IQ1_M | 256 | 56 | 1.75 | 2048 项 8 维 codebook |
| TQ1_0 | 256 | 54 | 1.69 | 三值，5 个塞一字节 (3^5=243) |
| IQ1_S | 256 | 50 | 1.56 | 2048 项 8 维 codebook |
注意加粗那些行：全部是 codebook。
这不是巧合。
8. IQ 系列：2 比特是怎么变成可能的？
8.1 标量量化输在哪？
2 bpw 意味着每个权重只有 4 种取值。把 Transformer 权重就近舍入到 4 种取值，模型直接废掉。
但关键在于：谁规定必须一个一个权重地量化呢？
rate-distortion 理论说，在相同码率下矢量量化优于标量量化，而且维度越高差距越大。具体到这里：8 个权重、每个 2 比特，一共 16 比特 = 65536 种可能的 8 维向量。可真实的权重 8 维向量并不是均匀撒在这个空间里的，它们聚集在原点附近，大致是独立同分布的高斯。
所以，与其允许全部 65536 个点，不如挑出最有用的 256 个，存一个 8 比特索引，把省下来的比特花到符号位和缩放因子上。
用 2 维画个直觉：
标量量化：每个轴独立取 4 档 -> 16 个格点
w2
^ . . . .
| . . . .
| . . . .
| . . . .
+-----------> w1
角落上的码字几乎永远用不到
矢量量化：直接在数据密集的地方摆 16 个码字
w2
^ . .
| . . . .
| . . . .
| . .
+-----------> w1
角落不分配码字，码字密度跟随数据分布真实情况是 8 维。8 维立方体的内切球体积只占立方体的 pi^4/(24 * 256)，约 1.6%。也就是说，标量网格 98.4% 的码字都浪费了。矢量量化可以把码字集中在高密度区域，节省下来的比特就能用来存符号和缩放因子。
8.2 IQ2_XXS 的 2.06 bit 是怎么凑出来的？
// ggml/src/ggml-common.h:381
typedef struct {
ggml_half d; // 2 字节：fp16 super-block 缩放因子
uint16_t qs[QK_K/8]; // QK_K = 256（super-block 大小），即 32 个 uint16 = 64 字节
} block_iq2_xxs; // 66 字节 / 256 权重 = 2.0625 bpw精彩的地方在于这 64 字节是怎么花的。uint16 只是存储声明（block 开头是 2 字节的 fp16，声明成 uint32 会引入 2 字节对齐 padding，66 字节就变 68 字节了），真正的字段边界跟它对不上。解码 kernel（dequantize_row_iq2_xxs，ggml-quants.c:2488）每处理 32 个权重（4 组，每组 8 个），就把 qs 里连续 4 个 uint16 也就是 8 字节 memcpy 进两个叫 aux32 的 uint32，再用移位和掩码取字段。按 aux32 这个视图画出来：
aux32[0]: +--------+--------+--------+--------+
| idx3 | idx2 | idx1 | idx0 | 4 个 8 比特 codebook 索引
+--------+--------+--------+--------+ (每 8 个权重一个)
aux32[1]: +----+-------+-------+-------+-------+
|scl | sgn3 | sgn2 | sgn1 | sgn0 |
| 4b | 7b | 7b | 7b | 7b |
+----+-------+-------+-------+-------+
^ ^
| +-- 每组 8 个权重只存 7 个符号位，
| 第 8 个由偶校验约束反推出来
+--------- 4 比特缩放因子：这 32 个权重（4 个小组）共享一个，
解码时 db = d * (0.5 + scl) * 0.25
取字段就是纯移位：scale = aux32[1] >> 28，第 l 组符号 =
(aux32[1] >> 7*l) & 127，第 l 组索引 = aux32[0] 的第 l 个字节
合计：32 个权重 64 比特 = 2.0 bpw，再加 fp16/256 = 2.0625 bpw
拆到每 8 个权重：
8 bit magnitude 模式（256 项 8 维 codebook 的索引）
7 bit 符号（第 8 个由奇偶性还原）
1 bit 分摊那 4 比特缩放因子（4 bit / 32 权重）
-------
16 bit 存 8 个权重这里反复出现的“codebook 索引”指向的是一张 256 项的全局表 iq2xxs_grid，它里面装的是什么、这 256 项是怎么挑出来的，是 8.3 节的主角；这一节先只看比特怎么摆。
那个符号位的技巧值得停下来看一眼。符号是不可压缩的，它就是 8 个实打实的比特。所以 codebook 里只存 magnitude，编码器强行让符号模式满足偶校验，不满足就翻转其中最不重要的那个元素：
// ggml/src/ggml-quants.c:3341
if (nflip%2) {
int imin = 0; float min = weight[8*k+imin]*xb[8*k+imin]*xb[8*k+imin];
for (int i = 1; i < 8; ++i) {
float ax = weight[8*k+i]*xb[8*k+i]*xb[8*k+i];
if (ax < min) { min = ax; imin = i; } // 找出最不重要的那个元素
}
xval[8*k+imin] = -xval[8*k+imin]; // 翻它，把奇偶性凑对
s ^= (1 << imin);
}主动在最不重要的权重上制造一个错误，来节约 1 个比特。 而“最不重要”是用 imatrix 权重衡量的。
约束闭环的另外两半在这里：存的时候只留低 7 位（block_signs[k] = s & 127，ggml-quants.c:3360）；解码端的“反推”则预先烘焙进了一张 128 项的表（ksigns_iq2xs，ggml-common.h:513），第 i 项就是 i 拼上它自己的奇偶校验位，于是还原完整 8 位符号也只是一次查表，运行时连 popcount 都不用算。
8.3 codebook 里藏着什么？
先明确一点：codebook 本身不在模型文件里。GGUF 里逐 block 存的只有上一节那些 8 比特索引，表本体是编译进 llama.cpp 的全局常量，256 项 x 8 字节 = 2 KB，所有 tensor 的所有 block 共享同一张。它在 ggml-common.h 里就是一个 uint64_t[256]，每一项拆成 8 个无符号字节，就是一个 8 维 magnitude 向量，一个字节一维（符号不在这里，是解码时另外乘上去的）：
// ggml/src/ggml-common.h:560
GGML_TABLE_BEGIN(uint64_t, iq2xxs_grid, 256)
0x0808080808080808, 0x080808080808082b, 0x0808080808081919, ...我把 256 项里每一个字节都 dump 出来数过，只有三个不同的值：8、25、43。解码是 db * grid[j]，其中 db = d * (0.5 + scale4) * 0.25，所以换算成缩放单位后，重建结果是 1.0, 3.125, 5.375。
现在对比一下编码器用的格点，它是运行时从另一张表构造的：
// ggml/src/ggml-quants.c:3119
for (int k = 0; k < grid_size; ++k) {
int8_t * pos = (int8_t *)(the_grid + k);
for (int i = 0; i < 8; ++i) {
int l = (kgrid[k] >> 2*i) & 0x3;
pos[i] = 2*l + 1; // {1, 3, 5, 7}
}
}可以看出，编码器在干净的奇整数格点 1,3,5 上搜索，解码器却在 1, 3.125, 5.375 上重建。
这 256 项的对应关系是严格一一对应且保序的。也就是说，解码 codebook 是刻意做成非均匀的：上面几档被往外拉伸了，因为高斯分布的尾部需要把外侧重建点推出去才划算。
IQ3_XXS 也玩同一手：它的 codebook 字节是 4,12,20,28,36,44,52,62，一个标准的 8k+4 等差阶梯，唯独最后一档是 62 而不是 60！
而 IQ4_NL（NL = non-linear）干脆连伪装都不要了：
// ggml/src/ggml-common.h:1120
GGML_TABLE_BEGIN(int8_t, kvalues_iq4nl, 16)
-127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113,
GGML_TABLE_END()16 种取值，间距依次是 23、21、18、16、14、13、12、11、12、12、13、15、16、20、24。零点附近密，尾部稀，是一条按神经网络权重分布拟合出来的 companding 曲线。
和 Q4_0 位宽相同、block 大小相同、同样 18 字节，误差却明显更低，纯粹因为那 16 个重建点摆在了更好的位置。
8.4 为什么编码那么慢，解码那么快？
编码 IQ2_XXS 意味着：对每一组 8 个权重、每一个候选缩放因子，都要在 8 维空间里对一个 256 点 codebook 做最近邻搜索。llama.cpp 在启动时预先建好 inverted index：
// ggml/src/ggml-quants.c:3108 - kmap_q2xs
const int kmap_size = 43692; // 8 个取值在 {0,1,2} 的 2 比特分量的最大索引
for (int i = 0; i < kmap_size; ++i) kmap_q2xs[i] = -1;
for (int i = 0; i < grid_size; ++i) {
...
kmap_q2xs[index] = i; // >= 0：这个格点本身就在 codebook 里
}
// 负值编码的是预计算好的邻居列表的偏移于是流程变成：就近舍入到格点，查表。如果它正好是 256 个码字之一，直接完事，O(1)；如果不是（而这才是常态，6561 个格点里只留了 256 个），就顺着那个负值找到预计算好的邻居列表，做一次小范围的加权搜索：
// ggml/src/ggml-quants.c:3387
int grid_index = kmap_q2xs[u];
if (grid_index < 0) {
const uint16_t * neighbours = kneighbors_q2xs - kmap_q2xs[u] - 1;
grid_index = iq2_find_best_neighbour(neighbours, kgrid_q2xs, xval + 8*k,
waux + 8*k, this_scale, Laux + 8*k);
}两边的成本差距是这样的：
flowchart LR
subgraph ENC["编码：离线一次，几分钟"]
E1["权重"] --> E2["扫 13 个候选缩放因子"]
E2 --> E3["舍入到格点"]
E3 --> E4{"在 codebook 里吗？"}
E4 -->|"不在，约 96%"| E5["邻居搜索<br/>8 维加权 L2"]
E4 -->|"在"| E6["得到索引"]
E5 --> E6
E6 --> E7["重新拟合缩放因子，迭代"]
end
subgraph DEC["解码：每个 token，片上，纳秒级"]
D1["8 比特索引"] --> D2["grid[idx]"]
D2 --> D3["拿到 8 个字节，结束"]
end
ENC -.->|"刻意做成不对称"| DEC
style E5 fill:#5a2d2d,stroke:#a44,color:#fff
style D2 fill:#2d5a3d,stroke:#4a8,color:#fff这个不对称就是全部要点：所有的“聪明”都在离线时一次性付掉，最终交给推理 kernel 的东西是一个指向表的索引。
于是我们终于走到了硬件。
9. 冰山之下：硬件为什么需要查找表？
9.1 先看一眼它们的解码操作
| 格式 | 解码操作 |
|---|---|
| MXFP4 / NVFP4 | kvalues_fp4[nibble]：16 项表 |
| IQ4_NL / IQ4_XS | kvalues_iq4nl[nibble]：16 项表 |
| IQ2_XXS / IQ2_XS | iq2xxs_grid[byte]：256 项表，每项 8 字节 |
| IQ3_XXS / IQ3_S | iq3xxs_grid[byte]：256/512 项表，每项 4 字节 |
| IQ1_S / IQ1_M | iq1s_grid[11 bit]：2048 项表 |
| Q4_0 / Q4_K | q - 8：算术，因为它的 codebook 恰好是均匀的 |
均匀量化器不过是一张“可以现算、不必存储”的 codebook。
而一旦放弃均匀性（4 比特以下你必须放弃，因为均匀间距会把码字浪费在高斯 tail 上），解码就变成了货真价实的查表，没有任何算术捷径。
所以硬件层面的问题就变成了：16 到 256 项的查表，一秒钟几十亿次，并行地做，能做多快？
9.2 CPU 在 2006 年就解决了
x86 SSSE3 有 PSHUFB，ARM NEON 有 TBL。它俩干的是同一件事：给一个寄存器里的 16 字节表和另一个寄存器里的 16 个索引，一条指令、约 1 个周期，吐出 16 个查表结果。
当年加它们是为了视频编解码和字符串处理，结果在今天可以被用作 LLM 反量化的 primitive。
// ggml/src/ggml-cpu/arch/arm/quants.c:765，MXFP4 x Q8_0 点积
const int8x16_t values = vld1q_s8(kvalues_mxfp4); // 16 项 codebook，只加载一次
...
q4b.val[0] = ggml_vqtbl1q_s8(values, vandq_u8 (q4bits.val[0], m4b)); // TBL: 16 次查表
q4b.val[1] = ggml_vqtbl1q_s8(values, vshrq_n_u8(q4bits.val[0], 4)); // TBL: 16 次查表
q4b.val[2] = ggml_vqtbl1q_s8(values, vandq_u8 (q4bits.val[1], m4b));
q4b.val[3] = ggml_vqtbl1q_s8(values, vshrq_n_u8(q4bits.val[1], 4));
prod_1 = ggml_vdotq_s32(ggml_vdotq_s32(vdupq_n_s32(0), q4b.val[0], q8b.val[0]),
q4b.val[1], q8b.val[1]);内层循环里完全没有浮点转换，4 条 TBL 把 64 个打包的 FP4 编码变成 64 个 int8（还记得吗，kvalues_fp4 存的是翻倍后的 E2M1，全是小整数），然后 SDOT 直接和 int8 激活累加。浮点只在每个 block 出现一次，用来应用 E8M0 缩放因子。
x86 路径的形状完全一样：
// ggml/src/ggml-cpu/arch/x86/quants.c:936
const __m128i values128 = _mm_loadu_si128((const __m128i*)kvalues_fp4);
...
const __m256i q4b_1 = MM256_SET_M128I(
_mm_shuffle_epi8(values128, _mm_and_si128(_mm_srli_epi16(q4bits_1, 4), m4b)),
_mm_shuffle_epi8(values128, _mm_and_si128(q4bits_1, m4b)));在 AArch64 上，ggml_vqtbl1q_s8 字面上就是 #define ggml_vqtbl1q_s8 vqtbl1q_s8（ggml/src/ggml-cpu/ggml-cpu-impl.h:302）；ARMv7 没有这条指令，所以那边手写了一段模拟。
9.3 那 GPU 呢？
// ggml/src/ggml-cuda/vecdotq.cuh:56
// CUDA does not have an instruction for selecting bytes with 4 bit indices.
// However, __byte_perm is an instruction that selects bytes with 3 bit indices that can be used instead.__byte_perm 对应 PTX 的 prmt：从8 字节源（两个 32 位寄存器）里用 3 比特索引挑出 4 个输出字节。而 16 项表需要 4 比特索引。
于是 llama.cpp 只能用 8 项查表拼出 16 项查表：
// ggml/src/ggml-cuda/vecdotq.cuh:33，一张 16 项表，8 次查表
static __device__ __forceinline__ int2 get_int_from_table_16(const int & q4, const int8_t * table) {
const uint32_t * table32 = (const uint32_t *) table;
uint32_t tmp[2];
const uint32_t low_high_selection_indices = (0x32103210 | ((q4 & 0x88888888) >> 1));
#pragma unroll
for (uint32_t i = 0; i < 2; ++i) {
const uint32_t shift = 16 * i;
const uint32_t low = __byte_perm(table32[0], table32[1], q4 >> shift); // 第 0-7 项
const uint32_t high = __byte_perm(table32[2], table32[3], q4 >> shift); // 第 8-15 项
tmp[i] = __byte_perm(low, high, low_high_selection_indices >> shift); // 按第 3 位选
}
return make_int2(__byte_perm(tmp[0], tmp[1], 0x6420), __byte_perm(tmp[0], tmp[1], 0x7531));
}8 条 prmt 指令，完成 8 次 4 比特查表。
AMD 那条路径（__builtin_amdgcn_perm，即 V_PERM_B32）要 6 条外加掩码运算。对比一下 2006 年的 CPU：一条 PSHUFB 干 16 次！
总结来说：
| 指令集 | primitive | 每条指令的查表次数 | 备注 |
|---|---|---|---|
| SSSE3 (2006) | PSHUFB |
16 | 16 项表，1 周期 |
| AVX2 | VPSHUFB |
32 | 两个独立的 128 位 lane |
| AVX-512 VBMI | VPERMB |
64 | 完整的跨 lane 64 项表 |
| NEON (2011) | TBL |
16 | 最多 4 个寄存器拼 64 项表；llama.cpp 只用到单寄存器 16 项 |
| SVE2 | TBL |
随向量长度 | |
| CUDA (Volta~Hopper) | PRMT |
1 | 只支持 3 比特索引，要 8 条才凑出一张 16 项表 |
| CDNA/RDNA | V_PERM_B32 |
~1.3 | 6 条指令做 8 次查表 |
| Rubin sm_107 | mma 的 LUT B 模式 |
在 Tensor Core 内部 | 见下文 |
GPU 那一列就是问题所在：它有 CPU 百倍的裸算力，却没有一张够宽的字节置换网络来当真正的查表引擎。Rubin 之前的每一次 FP4/IQ4 反量化，都在把发射槽烧在地址运算上，产出零 FLOP 的工作。
更大的 codebook 情况更糟。IQ2_XXS 那张 256 项 x 8 字节的表塞不进任何置换网络，只能待在内存里：
// ggml/src/ggml-cuda/vecdotq.cuh:1043
const uint2 grid_pos = ((const uint2*)iq2xxs_grid)[aux8[k0/2]];这是一个 2KB 的 __device__ const 数组，用数据相关的下标去读，也就是一次 gather。最好情况下 warp 内地址一致，最坏情况下整个 warp 串行。
那这个 gather 到底吃掉多少性能呢？
9.4 自己测测看
llama.cpp 的 test-backend-ops 可以直接测单个算子，不用下模型。矩阵形状我取的是 DeepSeek-V4-Flash 的 lm_head：4096 -> 129280，5.3 亿个权重。
为什么挑它？因为在 MoE 模型里，单个最大的矩阵往往不是某个 expert，而是 lm_head。DeepSeek-V4-Flash 有 256 个 expert，每个 expert 的矩阵只有 4096×2048，而 lm_head 是 4096×129280，比任何一个 expert 大 60 倍，而且每个 token 都必须完整过一遍。它也足够大，能甩开各家的末级缓存，这一点等下在 Strix Halo 上很关键。
./test-backend-ops perf -o MUL_MAT -p "type_a=q4_K,type_b=f32,m=129280,n=1,k=4096"n=1 就是 batch=1 decode。把耗时换算成实际达到的显存带宽（权重字节数除以耗时），先看 RTX 3090，理论峰值 936 GB/s。类型一列加粗的是解码要查表的 codebook 类型，没加粗的是均匀量化，解码走纯运算：
| 类型 | bpw | 权重 | 耗时 | 相对速度 | 实测带宽 | 带宽利用率 |
|---|---|---|---|---|---|---|
| f16 | 16.00 | 1059 MB | 1191 us | 100% | 889 GB/s | 94.9% |
| q8_0 | 8.50 | 563 MB | 637 us | 187% | 883 GB/s | 94.3% |
| q6_K | 6.56 | 434 MB | 563 us | 212% | 772 GB/s | 82.4% |
| q5_K | 5.50 | 364 MB | 425 us | 280% | 857 GB/s | 91.5% |
| q4_0 | 4.50 | 298 MB | 342 us | 349% | 871 GB/s | 93.1% |
| q4_K | 4.50 | 298 MB | 355 us | 335% | 839 GB/s | 89.6% |
| iq4_nl | 4.50 | 298 MB | 348 us | 342% | 856 GB/s | 91.5% |
| iq4_xs | 4.25 | 281 MB | 329 us | 362% | 854 GB/s | 91.2% |
| mxfp4 | 4.25 | 281 MB | 372 us | 321% | 757 GB/s | 80.9% |
| iq3_s | 3.44 | 228 MB | 398 us | 299% | 572 GB/s | 61.1% |
| iq3_xxs | 3.06 | 203 MB | 362 us | 329% | 560 GB/s | 59.8% |
| iq2_s | 2.56 | 170 MB | 348 us | 343% | 488 GB/s | 52.1% |
| iq2_xs | 2.31 | 153 MB | 335 us | 355% | 457 GB/s | 48.8% |
| iq2_xxs | 2.06 | 137 MB | 316 us | 377% | 432 GB/s | 46.1% |
| iq1_s | 1.56 | 103 MB | 287 us | 415% | 360 GB/s | 38.5% |
我们发现，IQ2_XXS 比 Q4_K 少搬了 54% 的字节，耗时却只少了 11%。如果它能像 Q4_K 那样跑到 89.6% 的带宽利用率，这一次矩阵乘只要 163 us，而不是 316 us。将近一半的性能被 gather 吃掉了。
再说印证的部分。
- 4 bit 那一档全部挤在 81% 到 93%。 q4_0（纯算术，不查表）、q4_K（纯算术）、iq4_nl（16 项 codebook，走 8 条 PRMT）、iq4_xs、mxfp4，利用率几乎一样。也就是说，在 4 bit 这一档，llama.cpp 那套 PRMT 模拟查表的开销小到测不太出来，iq4_nl 只比 q4_0 低 1.6 个百分点。这个 trick 是成功的。
- 一旦 codebook 大到装不进置换网络，利用率就崩了。 iq3 掉到 60%，iq2 掉到 46% 到 52%，iq1_s 只剩 38%。它们的共同点正是那句
iq2xxs_grid[byte]：一次数据相关的显存 gather。 - f16 的 94.9% 是这台卡的实际上限。 拿它当参照，4 bit 那一档只损失了几个点，而 IQ2 损失了一半。
换句话说：codebook 查表在 4 bit 时几乎免费，因为 16 项表塞得进寄存器；一旦 codebook 变成 256 项 x 8 字节，它就只能待在显存里，于是查表本身成了新的瓶颈。
而硬件对这个问题的正面回应，就是 9.6 节要讲的 Rubin。
9.5 换两台机器，结论还成立吗？
同一组测试我又在另外两台机器上跑了一遍：Strix Halo（Radeon 8060S，ROCm 后端，256 GB/s）和 Intel Arc Pro B70（SYCL 后端，608 GB/s）。
每格两个数：相对速度 / 带宽利用率。相对速度以各机器自己的 f16 为 100%，回答的是“换成这个格式，这一次矩阵乘快了几倍”；带宽利用率则是“有没有把内存喂饱”。数据格子里加粗的是比 f16 还慢的；类型一列的加粗和上一张表一样，标的是 codebook 类型。
| 类型 | bpw | RTX 3090 936 GB/s |
Strix Halo 256 GB/s |
Arc Pro B70 608 GB/s |
|---|---|---|---|---|
| f16 | 16.00 | 100% / 95% | 100% / 85% | 100% / 98% |
| q8_0 | 8.50 | 187% / 94% | 202% / 91% | 172% / 90% |
| q6_K | 6.56 | 212% / 82% | 260% / 91% | 243% / 98% |
| q5_K | 5.50 | 280% / 92% | 312% / 92% | 278% / 94% |
| q4_0 | 4.50 | 349% / 93% | 383% / 92% | 350% / 97% |
| q4_K | 4.50 | 335% / 90% | 376% / 90% | 327% / 90% |
| mxfp4 | 4.25 | 321% / 81% | 406% / 92% | 57% / 15% |
| iq4_nl | 4.50 | 342% / 91% | 383% / 92% | 61% / 17% |
| iq4_xs | 4.25 | 362% / 91% | 405% / 92% | 207% / 54% |
| iq3_s | 3.44 | 299% / 61% | 441% / 81% | 114% / 24% |
| iq3_xxs | 3.06 | 329% / 60% | 566% / 92% | 88% / 16% |
| iq2_s | 2.56 | 343% / 52% | 451% / 62% | 173% / 27% |
| iq2_xs | 2.31 | 355% / 49% | 679% / 84% | 99% / 14% |
| iq2_xxs | 2.06 | 377% / 46% | 466% / 51% | 135% / 17% |
| iq1_s | 1.56 | 415% / 38% | 862% / 72% | 207% / 20% |
三个发现：
第一，Q 家族在哪台机器上都稳。 q4_0、q4_K、q5_K、q8_0 的利用率在三台机器上全都在 89% 到 98% 之间，相对速度也整齐地跟着 bpw 走，4.5 bpw 的 q4_0 稳定在 3.5 倍左右。
它们的 codebook 是均匀的，解码是纯算术，没有任何查表动作，所以走到哪都能吃满带宽。这就是 Q4_K_M 作为默认推荐的底气：它不挑后端。
第二，带宽越高，gather 越显眼。 对比 3090 和 Halo：iq3_xxs 在 3090 上利用率只有 60%，在 Halo 上却有 92%；iq2_xs 从 49% 变成 84%。相对速度也跟着走：iq1_s 在 3090 上只有 415%，在 Halo 上能到 862%。
原因是 Halo 每搬一个字节要花的时间是 3090 的 3.7 倍，gather 有足够的空档躲在慢吞吞的内存后面；3090 的显存快得多，gather 就藏不住了。同一个格式在不同机器上的相对代价可以差好几倍，选量化类型不能只看 bpw。
（不过 iq2_s 和 iq2_xxs 在两台机器上都掉得厉害，说明 codebook 大到一定程度，谁也救不了。）
第三，B70 那一列是个问题。 有四个格子的相对速度低于 100%，也就是说这些 4 bit / 3 bit / 2 bit 格式在这台机器上比不量化的 f16 还慢：
mxfp457%（3096 us，而 f16 只要 1777 us）iq4_nl61%iq3_xxs88%iq2_xs99%
搬的字节少了 3.7 倍，跑出来反而慢 1.7 倍。
这跟 codebook 大小无关。证据是 iq4_xs（207% / 54%）和 iq4_nl（61% / 17%）用的是同一张 kvalues_iq4nl 表，只是 block 结构不同，性能却差 3 倍多。所以这是 SYCL 后端对某些类型写了优化 kernel、对另一些走了通用回退。
Halo 上也有同样的问题，只是轻一些：iq3_xxs 92% 而 iq3_s 81%，这个模式同样和 codebook 大小对不上。
结论：选量化类型的时候，先确认你的后端为这个类型写过 kernel。 在 B70 上它让好几个低比特格式跑得比 fp16 还慢。注意，这个 Benchmark 并非测试硬件的潜力，随着算子的优化，还有很大的提升空间。
另外，自己测试时，请让矩阵选得足够大，比如 Strix Halo 有 32 MB Infinity Cache，如果矩阵小于 32 MB，gather 可以躲在 cache 里，就难以反映出真实的带宽利用率。
9.6 未来硬件（如 NVIDIA Rubin）还能支持什么呢？
先把两条路径并排画出来，差别一目了然：
flowchart TB
subgraph SW["软件查表：Rubin 之前的所有 GPU"]
A1["寄存器里的 4 比特编码"] --> A2["掩码 + 移位"]
A2 --> A3["8 条 PRMT 模拟<br/>一张 16 项表"]
A3 --> A4["int8 数值"]
A4 --> A5["转换 / 应用缩放"]
A5 --> A6["Tensor Core MMA<br/>fp16 或 int8"]
end
subgraph HW["硬件查表：Rubin"]
B1["寄存器里的 3 bit 索引<br/>+ 软件装载的 8 项表"] --> B2["Tensor Core MMA<br/>LUT B 模式"]
B2 --> B3["fp32 累加器"]
end
A6 --> B1
linkStyle 7 opacity:0
style A3 fill:#5a2d2d,stroke:#a44,color:#fff
style B2 fill:#2d5a3d,stroke:#4a8,color:#fff上面那条就是 9.3 节的 8 条 PRMT 加显存 gather。下面那条，是随 CUDA 13.4 开发者预览一起公开的 PTX ISA 9.4 里，Rubin（sm_107）的 Tensor Core 新增的 LUT 模式。它做的事一句话就能说完：让 MMA 自己去查一张软件装载的表。
具体拆开是这样：
B 操作数（权重一侧）：
每个权重位置 只存一个 3 bit 索引
每 8x64 block 共享一张 8 项的表，每项 1 字节 E4M3，共 64 bit
体积 3 + 64/512 = 3.125 bpw
重建 发生在 MMA 内部：读索引、查表、乘加一条龙，
kernel 从头到尾不接触解压后的权重矩阵表项选 E4M3 顺带定出了整个格式的精度天花板：重建值只能落在 E4M3 的网格上，所以它最好情况就是逐权重 MXFP8 的精度，NVIDIA 官方博客的说法也是 LUT 表示“最多可保住 MXFP8 的精度”。不过在 3 bit 这个码率上，主导误差的是 512 个权重只有 8 档可挤，单个码字 3% 上下的 E4M3 摆放误差反而是小头；E4M3 换来的是约 2^18 的动态范围，一张表跨好几个量级也兜得住。
关键就在“表由软件装载”这几个字。E2M1、IQ4_NL 这些格式的重建值都是定死的，而这张表里的 8 个 E4M3 值想放哪就放哪：非均匀、非对称、按校准数据逐 tensor 拟合，全都可以。这是 NVIDIA 第一个在 MMA 内部重建非均匀、可编程 codebook 的输入格式，第 8 节 IQ 系列“按数据拟合码字”的思路，第一次被请进了 Tensor Core。
和 CPU 那边的通用查表指令摆在一起看：
| CPU LUT 指令（PSHUFB / TBL） | Rubin 的 LUT B 模式 | |
|---|---|---|
| 表的内容 | 任意内容，运行时装载：PSHUFB 一张 16 字节，TBL 最多 4 个寄存器拼 64 字节 | 任意 8 项 E4M3，运行时装载 |
| 查表发生在 | SIMD 通路，查完还得自己做乘加 | MMA 内部，查表和乘加一体 |
| 谁能用 | 任何 codebook 格式 | 3 bit 标量 codebook 的权重 |
以前没人这么做的原因也不难猜：固定解码只是乘法器输入级的一小撮门电路，可编程解码则意味着在每一次 MAC 的关键路径上插入一次真正的表读取。Rubin 还是把它做了。
第 6 节讲过的设计约束在它身上全都对得上：只支持 B（权重）一侧，而且不许转置，正是“推理永远不转置权重”那条前提的硬件化；8×64 的表共享粒度，也是照着 MMA 的 tile 形状定的。
表长在 Tensor Core 内部，只有走 MMA 的大 batch 路径才用得上：llama.cpp 在 batch <= 8 时对所有量化类型走的都是 MMVQ 向量 kernel（ggml_cuda_should_use_mmvq，mmvq.cu:289），那里没有 Tensor Core，batch=1 decode 的查表还是得靠 PRMT 软件模拟。好在 decode 是带宽瓶颈，小表的开销藏得住（9.4 节实测过），真正吃到红利的是算力瓶颈的 prefill 阶段。至于 llama.cpp 对 LUT 模式本身，目前没有任何适配，现有 FP4 路径的 gating 甚至明确写着到 Rubin 一代为止（common.cuh:286）。
但方向已经明确：表进了 Tensor Core，表的内容留在了软件手里。
9.7 codebook 住在哪里，这些年变了什么？
2019 | fp16 权重 | 没有表。从头到尾都是算术。
|
2023 | Q4_0, Q4_K | 均匀 codebook = 算术 (q - 8)。免费。
|
2024 | IQ2_XXS ... IQ4_NL | 表在内存里。
| | CPU: PSHUFB/TBL，1 条指令 16 次查表。
| | GPU: 8 条 PRMT，或者从 __device__ const 里 gather。
| | ^ 反量化第一次成为 kernel 里不可忽略的开销
|
2024 | MXFP4 / OCP MX | 整个行业统一到同一张 16 项 codebook，
| | 就是为了让它能被做进硬件
|
2025 | Blackwell sm_120 | 表在硅片里。
| mma...block_scale | 4 比特编码成为 Tensor Core 的原生操作数。
| | 反量化开销 -> 0
|
2026 | Rubin sm_107 | 表变成可编程的。
| mma ... LUT B | 权重只存 3 bit 索引，MMA 内部查一张
| | 软件装载的 8 项 E4M3 表，每 8x64 block 一张。9.8 还剩什么问题没解决？
对照第 8 节的清单，Rubin 的 LUT B 其实只接住了一小半。先回顾拟合 codebook 到底强在哪：
- IQ4_NL vs Q4_0：同样 4.5 bpw、同样 block 大小，误差更低，差别纯粹来自码字位置。
- IQ2_XXS：2.06 bpw，任何标量格式在这个码率上都到不了可用质量，靠的就是 8 个权重共享一个 256 项矢量 codebook。
- IQ3_XXS 那个被拉伸的顶档（62 而不是 60），存在的理由就是有人实测发现它更好。
把它们和硬件现状放进同一张表里：
| codebook | 查表发生在 | 同码率精度 | 吞吐（9.4/9.5 节实测） | |
|---|---|---|---|---|
| Q4_0 / Q4_K | 均匀，公式现算 | 不用查 | Baseline | 吃满带宽 |
| MXFP4 / NVFP4 | 固定 16 项 | 寄存器，PRMT 模拟 | 固定表的天花板 | 4 bit 档几乎免费 |
| IQ4_NL / IQ4_XS | 拟合 16 项 | 寄存器，PRMT 模拟 | 高于同位宽的固定表 | 几乎免费 |
| IQ3 / IQ2 / IQ1，AQLM，QuIP# | 拟合矢量表，256~2048 项 | 显存 gather | 同码率最高 | 带宽利用率掉到 38%~61% |
| Rubin LUT B | 拟合 8 项，软件装载 | MMA 内部 | 未公布 | 零开销 |
最后一行就是 Rubin：拟合出来的表第一次拿到了零开销的查表。但对照上面几行，它还差三件事：
- 表太小。 8 项对应 3 bit，是 IQ3 档；IQ4_NL 的 16 项就已经装不下。
- 只有标量，没有矢量。 IQ2_XXS 那种 256 项 x 8 维的矢量 codebook，也就是 2 bpw 档的全部依仗，依然没有任何硬件对应；8.1 节讲的 rate-distortion 收益，恰恰是标量表拿不到的那部分。
- 只有权重，没有激活。 LUT 模式只作用于 B 操作数，A 侧还是老老实实的数值格式。
与此同时，极低比特那一端正在发生一个有趣的反转：在 1~2 比特下，一些设计干脆不再查权重，而是查部分点积：把一组激活的全部 2^k 种可能和预先算好，于是“k 比特权重组的乘法”退化成一次查表。
这把 LUT 从解压装置直接变成了运算单元本身，而且天然贴合本来就由 LUT 堆出来的 FPGA 和 NPU 架构。llama.cpp 的 TQ1_0/TQ2_0 三值类型是这个世界的格式侧，运算侧目前主要还在论文里，代表作是 Microsoft 的 LUT Tensor Core（ISCA 2025）：给 Tensor Core 配一张可编程查找表，预计算激活与码字的部分积，把可编程 LUT 这个方向推得比 Rubin 还要远。
一句话总结这一部分：8 比特以下，“反量化”不是算术，是解码；而解码就是查表。CPU 从 2006 年起就有查表指令，这正是它在量化推理上显得格外能打的原因；GPU 长年没有，只能用大约 8 倍的代价去模拟；直到 Rubin，MMA 才第一次学会查一张软件装载的表。
10. KV Cache 又是另一回事
10.1 为什么权重那套搬不过来？
权重量化有重要性矩阵、有缩放因子扫描、时间随便花。KV 量化一样都没有：tensor 在 token 到来之前根本不存在，而且必须在把它写进显存的那点时间里编码完。
所以 llama.cpp 只允许 KV 用那些便宜的、闭式的类型：
// common/arg.cpp:305
const std::vector<ggml_type> kv_cache_types = {
GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16,
GGML_TYPE_Q8_0,
GGML_TYPE_Q4_0, GGML_TYPE_Q4_1, GGML_TYPE_IQ4_NL,
GGML_TYPE_Q5_0, GGML_TYPE_Q5_1,
};注意名单里没有 K-quants（它们的编码器要跑 37 轮搜索），也没有 IQ2/IQ3（它们需要一个“针对未来激活”的 imatrix，而那玩意儿不可能存在）。IQ4_NL 能进这个名单，是因为它的编码只是 16 路扫描，解码只是一条 PSHUFB。
还有两条硬约束，都在创建 context 时检查：
// src/llama-context.cpp:3597
if (ggml_is_quantized(params.type_v) && params.flash_attn_type != LLAMA_FLASH_ATTN_TYPE_ENABLED) {
// enabling flash_attn since it is required for quantized V cache
}
// src/llama-context.cpp:3609
if (model->hparams.n_embd_head_k(il) % blck_size != 0) {
// "K cache type %s with block size %u does not divide n_embd_head_k=%u"
}- 量化 V 必须开 flash attention。 不开 FA 的话，
V^T @ softmax需要转置读取 V，而分 block 量化的 tensor 不解码是没法转置的；FA 则按 V 的自然布局消费它。 - block 大小必须整除
head_dim。 一个量化 block 绝不能跨两个 head，因为不同 head 的数值尺度差异极大。head_dim = 128配 block 大小 32 没问题；head_dim = 80的话 Q4_0 直接被拒。
10.2 量化之前，先转个身
Attention 的 key 有非常严重的离群通道，少数几个维度的 magnitude 是其余维度的 10~100 倍，而且在所有 token 上稳定出现。
分 block 量化对这种分布的处理极其糟糕：一个离群值决定了整个 block 的缩放因子，其余的数全被压到剩下那一两档上。
这份代码树实现的是 incoherence processing 那一路的修法，理论脉络是三篇论文：QuIP（Chee 等，2023，用随机正交旋转把权重变得“不相干”，2 bit 量化第一次有了理论保证）、QuIP#（Tseng 等，2024，把旋转换成可以 O(n log n) 计算的 Hadamard 变换）、QuaRot（Ashkboos 等，2024，把旋转从权重推广到激活和 KV cache，llama.cpp 这套做法和它最接近）。具体动作是：量化之前，先把 K 和 V 乘上一个正交的 Walsh-Hadamard 矩阵。
// src/llama-kv-cache.cpp:20
// orthonormal Walsh-Hadamard rotation matrix
// note: res^2 == I
static void ggml_gen_hadamard(ggml_tensor * tensor) {
...
data[0*n + 0] = 1.0 / sqrtf(n);
for (int s = 1; s < n; s *= 2) { // Sylvester 递推
for (int i = 0; i < s; i++) {
for (int j = 0; j < s; j++) {
const float val = data[i*n + j];
data[(i + s)*n + (j )] = val;
data[(i )*n + (j + s)] = val;
data[(i + s)*n + (j + s)] = -val;它为什么有效呢？Hadamard 矩阵是一个“最大程度打散”的正交变换：每个输出坐标都是全部输入的 +/-1/sqrt(n) 倍求和，于是单个特别大的输入通道会被均匀抹开到全部 n 个输出上，结果向量接近高斯。block 内动态范围随之塌缩，4 比特分 block 量化器一下子就能用了。
而且推理时它是免费的，因为 H 正交：(HQ) . (HK) = Q^T H^T H K = Q . K。把 query 用同一个矩阵旋转，attention 分数在数学上完全不变！
计算图里干的正是这件事：
// src/llama-graph.cpp:2778
q_cur = llama_mul_mat_hadamard(ctx0, q_cur, inp->self_k_rot);
k_cur = llama_mul_mat_hadamard(ctx0, k_cur, inp->self_k_rot);
...
v_cur = llama_mul_mat_hadamard(ctx0, v_cur, inp->self_v_rot);V 这边不会自动抵消，因为 V 是一个值、不是内积的一半，所以旋转改在 attention 输出上还原（llama-graph.cpp:2814）。能这么做的理由是一样的：H 是线性作用的，而且 H^2 = I。
整条路径：
flowchart TD
Q["q_cur"] -->|"乘 H"| QR["旋转后的 q"]
K["k_cur"] -->|"乘 H"| KR["旋转后的 k"]
V["v_cur"] -->|"乘 H_v"| VR["旋转后的 v"]
KR --> KQ["量化 -> Q4_0 / IQ4_NL"]
VR --> VQ["量化 -> Q4_0 / IQ4_NL"]
KQ --> KC[("K cache")]
VQ --> VC[("V cache")]
QR --> FA["flash attention"]
KC --> FA
VC --> FA
FA --> OUT["attn 输出"]
OUT -->|"乘 H_v 还原"| FIN["最终输出"]
KC -.->|"上下文平移"| SH["反量化到 f32<br/>乘 H 解旋转<br/>RoPE<br/>乘 H 再旋转<br/>重新量化"]
SH -.-> KC
style KQ fill:#3d3d6b,stroke:#88a,color:#fff
style VQ fill:#3d3d6b,stroke:#88a,color:#fff
style SH fill:#5a2d2d,stroke:#a44,color:#fff那条虚线就是代价。上下文平移（context shift）需要对已经躺在 cache 里的 key 施加 RoPE，而 RoPE 和 Hadamard 旋转不可交换，于是：
// src/llama-kv-cache.cpp:1864
if (ggml_is_quantized(cur->type)) {
// dequantize to f32 -> RoPE -> quantize back
tmp = ggml_cast(ctx, cur, GGML_TYPE_F32);
tmp = llama_mul_mat_hadamard(ctx, tmp, rot); // 转回去
tmp = ggml_rope_ext(ctx, tmp, shift, factors, n_rot, ...);
tmp = llama_mul_mat_hadamard(ctx, tmp, rot); // 再转过来
tmp = ggml_cpy(ctx, tmp, cur);
}完整地过一遍 fp32、两次矩阵乘、再重新量化，而每一次往返都会叠加量化误差。这是在线编解码器的根本代价：数据必须既可解码又可修改。
10.3 KV 实际该怎么选？
-ctk q8_0 -ctv q8_0：cache 减半，质量损失基本可以忽略。长上下文场景下这应该是默认选项。-ctk q4_0 -ctv q4_0：减到四分之一。没有 Hadamard 旋转时长上下文召回会有明显退化，有旋转则好得多。-ctk q8_0 -ctv q4_0：开头算过的那个非对称平衡点，K 保 8 bit 不动摇 softmax，V 压到 4 bit 省空间。Qwen3.8-27B 128k 下这一档是 3.25 GB，比全 fp16 的 8 GB 省了六成。注意这种 K、V 不同类型的配法只适用于 Qwen 这类 GQA 模型：DeepSeek 这类 MLA 架构的 K、V 共用一份 cache，llama.cpp 要求两者类型一致（llama-context.cpp:3590），会直接拒绝启动。-ctk iq4_nl -ctv iq4_nl：和 q4_0 同样大小但 codebook 更好。GPU 上更慢（原因见上一节），CPU 上通常划算。
11. 回到最初的问题：那个文件列表该怎么选？
先学会读名字：
Q4_K_M
| | |
| | +-- solution：S = small（统一处理），M = medium（attn_v 和部分
| | ffn_down 升级），L = large（升级更多）
| +----- K = K-quant 家族：256 权重 super-block，两级 6 比特缩放
+------- 名义 4 bit/权重（Q4_K 实际 4.5）
IQ3_XXS
| | |
| | +-- XXS < XS < S < M：codebook 索引和缩放因子给多少位
| +----- 名义 3 bit（实际 3.06）
+------- I = imatrix 系：对固定 codebook 做矢量量化，
量化时**必须**有 imatrix
UD-* unsloth 自己的逐 tensor solution + 自己的校准语料，
底层 block 格式还是上面这些然后按硬件和 budget 挑：
| 情况 | 选 | 理由 |
|---|---|---|
| 显存充足 | Q8_0 或 Q6_K | 和 fp16 基本无法区分 |
| 最常见的情况 | Q4_K_M | 精度/体积的拐点，很难被打败 |
| GPU，要速度 | Q4_K_M 或 MXFP4 | 均匀 codebook，kernel 里不用 gather |
| Blackwell 显卡 | NVFP4 / MXFP4 | 走原生 Tensor Core，查表免费 |
| 纯 CPU | IQ4_XS 或 IQ4_NL | PSHUFB/TBL 让 codebook 变成零成本 |
| 勉强塞得下 | IQ3_M，再不行 IQ2_M | 需要好的 imatrix，自己测 perplexity |
| 实在没辙 | IQ1_M | 大模型跑 IQ1_M 常常仍然强于小模型跑 Q8_0 |
| KV cache | q8_0，超长上下文用 q4_0 | 见上一节 |
所以回到网上的两派言论：
- “显存够就 Q8，不够就 Q4_K_M”：对，因为 Q4_K 的 codebook 是均匀的，GPU kernel 里连查表都不需要，它的速度是所有 4 bit 方案里最稳的。
- “IQ 同样大小精度更高”：也对，但要知道你为此付了什么。实测下来 IQ2/IQ3 在 GPU 上仍然比 Q4_K 快，只是远达不到字节数该给的加速：少搬一半字节，只快两成。纯 CPU 推理请放心用 IQ4_XS（PSHUFB 让 codebook 免费）；GPU 上用 IQ2/IQ3 是拿吞吐换容量，值不值取决于你是不是真的塞不下。
最后两个提醒。
第一，永远要拿“大模型低 bpw”去和“小模型高 bpw”对比，通常前者赢，这本来就是量化的意义所在。
第二，量化的伤害在各项能力上分布极不均匀！perplexity（PPL，测试语料上逐 token 负对数似然的指数平均，越低越好）几乎不动的时候，多步推理、代码生成和长上下文检索已经退化得很明显了。所以请测你真正在乎的那个指标，别只看 PPL。
12. 想自己核对的话，从哪读起？
本文引用的所有源码路径和行号，都对应 llama.cpp 的 commit 9ee9fc04c（2026-08-19 的 master）。行号会随版本漂移，对不上的时候请以这个 commit 为准，或者直接搜函数名和表名。
| 文件 | 里面有什么 |
|---|---|
ggml/src/ggml-common.h |
全部 block 结构、全部 codebook 表。从第 185 行开始看 |
ggml/src/ggml-quants.c |
参考编解码器。第 799 行的 make_qkx2_quants 是 K-quants 的心脏 |
ggml/src/ggml-impl.h:477-560 |
E8M0 和 UE4M3 的位操作 |
ggml/src/ggml-cpu/arch/arm/quants.c |
NEON kernel，第 765 行是 MXFP4 的 TBL 循环 |
ggml/src/ggml-cpu/arch/x86/quants.c |
AVX2 对应实现，第 936 行 |
ggml/src/ggml-cuda/vecdotq.cuh:33 |
PRMT 模拟查表，注释一定要读 |
ggml/src/ggml-cuda/mmvq.cu:289 |
batch 大小如何决定走不走 Tensor Core 路径 |
src/llama-quant.cpp:424 |
哪个 tensor 分到哪个类型，以及为什么 |
tools/imatrix/imatrix.cpp |
重要性矩阵的采集 |
src/llama-kv-cache.cpp:20 |
Hadamard 矩阵的生成 |
src/llama-graph.cpp:2778 |
旋转在计算图里的施加点 |
本文用到的外部资料：
- unsloth Dynamic 3.0 GGUFs 文档：开头那张精度曲线的出处，页面上还有 KL Divergence、MMLU 等完整表格
- OCP Microscaling Formats (MX) v1.0 规范：MXFP4 / E8M0 的标准定义
- Inside NVIDIA Rubin GPU Architecture：Rubin LUT 权重格式的官方介绍
- PTX ISA 9.4（CUDA 13.4 开发者预览）：Rubin LUT 模式的指令级定义
第 9 节那组吞吐数据的复现命令。注意 test-backend-ops 的矩阵形状是硬编码在 tests/test-backend-ops.cpp 里的，文中用的 DeepSeek-V4-Flash lm_head 形状需要自己加一行：
// tests/test-backend-ops.cpp，加在原有的 mul_mat 循环之后
for (ggml_type type_a : all_types) {
test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 129280, 1, 4096, {1, 1}, {1, 1}));
}然后 cmake --build build --target test-backend-ops 重编（只改这一个文件，十几秒就好），再跑：
./test-backend-ops perf -o MUL_MAT -p "type_a=q4_K,type_b=f32,m=129280,n=1,k=4096"常用命令：
# 采集重要性矩阵
llama-imatrix -m model-f16.gguf -f calibration.txt -o model.imatrix
# 带 imatrix 量化
llama-quantize --imatrix model.imatrix model-f16.gguf model-IQ3_M.gguf IQ3_M
# 量化了多少伤害
llama-perplexity -m model-IQ3_M.gguf -f wiki.test.raw
# 量化 KV cache
llama-cli -m model.gguf -c 131072 -ctk q8_0 -ctv q4_0 -fa on13. 写在最后
整条线索三句话就能讲完：
浮点格式本质上是在赌“比特该怎么在范围和精度之间分”，而到了 8 比特以下，剩的比特已经不够精度了，于是范围被挪出去变成每个 block 共享的缩放因子，元素本身退化成一个小小的索引。
一旦元素变成索引，重建值就成了自由参数，而“按数据拟合”必然打得过任何固定的算术规则。这就是为什么 IQ 系列在相同码率下比 Q 系列更准。
而当每一种格式在本质上都是一次 codebook 查表之后，瓶颈就不再是计算，而是查表。这解释了为什么带 PSHUFB 的 CPU 多年来在这件事上悄悄地表现优异，为什么 GPU 要花 8 条 PRMT 去假装自己有查表指令，以及为什么 NVIDIA 走到 Rubin，终于让 MMA 学会了查一张软件装载的表。
而未来的端侧硬件，有没有可能上一个像 FPGA 那样的 LUT 运算单元，直接把 ~3072 bit 的输入映射到 64 个 8 bit 的输出？它有没有可能直接让端侧超低开销地查一张 codebook？这条路上还有很多问题没解决，也许硬件会和算法一起，给我们带来更多惊喜。