第13章 量化：精度与速度的工程平衡

量化是 LLM 推理工程里的近似艺术：把数值表示从模型训练时的宽余格式，换成推理时更贴近带宽、显存和 kernel 能力的格式。收益不是白送的，精度、硬件和负载形态必须一起看。

13.1 量化的数学本质

13.1.1 "过剩的精度"

在所有量化方案之前，先理解一个朴素事实：LLM 权重的数值分布有大量冗余，低 bit 表示足以保留核心信息。

一个比喻：FP16 相当于给每个权重配一把带千分位刻度的游标卡尺，而实际上模型权重只需要毫米尺就能量准。多余的刻度不但没用，还在占用 HBM 的带宽。量化就是换一把合适的尺子。

13.1.2 为什么 FP16 是"过剩"的

一个典型的 LLM 权重张量长什么样？如果拿某个大模型线性层的 q_proj.weight 跑直方图，常见现象是：

• 主体分布集中在 0 附近，正负值大体对称
• 绝大多数权重落在一个很窄的区间，少数 outlier 决定全局最大值
• 不同层、不同投影矩阵的尾部形态会明显不同，不能用一个固定 scale 覆盖所有层

FP16 有 65536 种离散取值，动态范围很大。但很多权重张量只使用了其中很小的有效区间，大量 bit pattern 不会被触发——信息密度低。

如果把一个窄区间用 INT4 的 16 个离散值去覆盖，步长会粗很多；但只要 scale 按层、按通道或按 group 选得足够细，误差就不再由少数 outlier 主导。量化真正要解决的不是"INT4 一定够不够"，而是"怎样把有限的 16 个格子分配给真正高频的数值区域"。

13.1.3 量化公式

最朴素的对称线性量化：

$$w_q = \text{round}\left(\frac{w}{s}\right), \quad w_q \in [-2^{b-1}, 2^{b-1}-1]$$

其中 $s = \frac{\max(|w|)}{2^{b-1} - 1}$ 是缩放因子，$b$ 是 bit 数。

反量化：$\hat{w} = w_q \cdot s$

引入零点（zero point）后变成非对称量化，更好处理非对称分布：

$$w_q = \text{round}\left(\frac{w}{s}\right) + z$$

$z$ 是把非对称分布平移到量化区间中心的偏置。

13.1.4 量化误差的量化

量化误差 $e = \hat{w} - w$，在 round-to-nearest 的均匀量化里，上界约为 $\frac{s}{2}$。如果一个 group 的有效范围越窄，$s$ 越小，单个权重的最大舍入误差也越小。

关键是这个误差会不会在模型前向传播中放大。Transformer 的每层是矩阵乘、残差和归一化的组合：矩阵乘可能传播误差，残差会让误差不必逐层完全累乘，归一化又会重新约束激活尺度。不同模型、校准集和评测集的结果差异很大，所以工程上必须用目标任务的离线集验证，而不能把某个论文或 checkpoint 的 perplexity 变化当成通用承诺。

量化之所以可行，靠的是模型对这类有界舍入误差有相当的容忍度；容忍到什么程度，只能靠上面说的离线验证来回答。

13.1.5 离群值（Outlier）的双刃剑

有一个特别的挑战：离群值。某些层的权重或激活尾部会明显比主体分布更长。朴素 per-tensor 量化会因为这些少数值拉大 scale，导致主体区域的离散格子变稀，其他权重的有效精度骤降。

策略	如何处理 outlier
朴素 per-tensor	被离群值主导，精度差
Per-channel	把离群值限制在一行，其他行不受影响
Per-group	进一步细化，outlier 只影响 128 个值
SmoothQuant	在量化前把 outlier 的 activation 转移到 weight
AWQ	识别重要通道并放大 scale 补偿

现代量化方法的大量设计，都可以理解成在限制 outlier 对 scale 的破坏：要么把 scale 切细，要么移动 activation/weight 的尺度，要么在 kernel 里支持更复杂的元数据。

13.2 量化为什么能加速：memory-bound 视角

第 12 章讲过 LLM Decode 很容易落在 memory-bound 区域。量化减少的正是"搬的字节数"。下面这张表只算权重读取的理论下限，不代表真实 TPOT；真实延迟还包括 KV 读写、attention、调度、通信、kernel launch 和反量化开销。

格式	bit/param	70B 参数的权重字节数	在 2 TB/s HBM 下的仅读权重下限
FP16/BF16	16	约 140 GB	约 70 ms
FP8	8	约 70 GB	约 35 ms
INT4 (GPTQ/AWQ)	4	约 35 GB	约 17.5 ms

字节数会近似线性下降，但端到端延迟不会自动线性下降。只有当权重读取是主瓶颈、kernel 能把反量化和矩阵乘融合、并且 batch 没有把系统推向 compute-bound 时，量化的加速才会接近字节数下降的比例。

graph TB
    subgraph "Decode step 的量化收益路径"
        F16["FP16/BF16<br/>2 bytes / param"]
        F8["FP8<br/>1 byte / param<br/>依赖 FP8 kernel 支持"]
        I4["INT4 + fused kernel<br/>0.5 byte / param<br/>边读边反量化边算"]
        I4_bad["INT4 朴素路径<br/>先展开成 FP16<br/>收益被中间读写吞掉"]
    end
    F16 --> |权重字节减半| F8
    F16 --> |权重字节减 3/4 + 融合 kernel| I4
    F16 --> |"权重虽小<br/>但中间张量变大"| I4_bad
    style F16 fill:#dc2626,color:#fff,stroke:none
    style F8 fill:#f59e0b,color:#333,stroke:none
    style I4 fill:#10b981,color:#333,stroke:none
    style I4_bad fill:#94a3b8,color:#333,stroke:none

注意最后一行——"INT4 朴素反量化"的路径。如果你的做法是：先读 INT4 → 展开成 FP16 → 再做 FP16 GEMM，那 INT4 权重的读取确实只有 35 GB，但逐层展开出的 FP16 临时权重累计要写出 140 GB、再被 GEMM 读回 140 GB，净收益大打折扣。这正是 Marlin 这类"融合反量化 + GEMM" kernel 存在的价值（下文详述）。

13.2.1 高并发场景：量化的"折扣"

batch size 足够大时，整个系统从 memory-bound 向 compute-bound 过渡。此时权重读取不再是瓶颈，量化的加速收益减小、甚至可能出现负收益（因为反量化本身是 FLOPs 开销）。

正确做法不是背一个固定加速比，而是按自己的服务形态测一组小矩阵：

维度	为什么要测
batch size / 并发	判断系统是否从 memory-bound 转向 compute-bound
prompt 长度 / decode 长度	分离 prefill、decode、KV 压力
TPOT / ITL / 吞吐	避免只看 tok/s 掩盖单请求延迟
显存峰值 / KV block 使用率	判断量化节省的是权重、KV，还是二者都有
输出质量指标	确认校准集和生产分布没有明显漂移

生产建议：低并发 chat / 单机部署 / 边缘推理 → 开量化；高并发 API 网关 → 先 benchmark 再决定。

13.2.2 一个简化的 roofline 视角

量化的收益可以用 Roofline 模型解释：

arithmetic_intensity = FLOPs / Bytes

对 GEMM（batch=1 时）：

• FP16：每个权重参与 2 FLOPs，读取 2 字节 → AI = 1.0
• INT4：每个权重参与 2 FLOPs，读取 0.5 字节 → AI = 4.0

batch size 小 → AI 小 → memory-bound → 量化赚 batch size 大 → AI 大 → compute-bound → 量化效益递减

这个视角让你能预判量化收益的方向，但上线前仍然要 benchmark。Roofline 只能告诉你瓶颈可能在哪里，不能替你验证 kernel 是否走到了预期路径。

13.3 量化的三种粒度

量化的"粒度"决定了一个 scale（和 zero point）覆盖多少个权重。粒度越细精度越好，但元数据越多。

13.3.1 Per-Tensor

整个权重张量共享一个 scale。元数据最少，但精度最差——任何一个离群值都会拉大 scale，让其他权重的量化精度下降。

weight.shape = [8192, 8192]  # 67M 个值
scale.shape  = []            # 1 个值

INT4/INT8 权重量化里很少用它，更多用于激活量化（下文）；FP8 权重则常用 per-tensor——vLLM 对未量化 checkpoint 做在线 FP8 量化时，ops.scaled_fp8_quant(layer.weight, scale=None) 为每个权重张量算出的就是一个 scale（vllm/model_executor/layers/quantization/fp8.py:326-328）。

13.3.2 Per-Channel（Per-Row / Per-Column）

每行（或每列）一个 scale。元数据增加一个维度：

weight.shape  = [8192, 8192]
scale_row.shape = [8192]         # row = output channel

这对线性层尤其自然——每个输出通道有自己的 scale。精度比 per-tensor 好很多，GPTQ / AWQ 都至少用 per-channel。

13.3.3 Per-Group

行内再切成若干 group（典型 group_size=128），每 group 一个 scale：

weight.shape = [8192, 8192]
# row 内每 128 列一个 scale
# 8192 / 128 = 64 groups per row
scale.shape  = [8192, 64]  # 共 524,288 个 scale

这是目前 4-bit 量化里很常见的选择。精度进一步提升，元数据成本也更高。group_size=128 常被作为默认起点；再小（如 32）可能提升精度但增加 scale 读写和 kernel 复杂度，再大（如 256）减少元数据但更容易被局部 outlier 影响。

13.3.4 精度-速度曲线对比

graph TB
    Coarse["Per-Tensor<br/>1 scale / tensor<br/>开销: 几字节<br/>精度: ★★"]
    Channel["Per-Channel<br/>1 scale / row<br/>开销: KB 级<br/>精度: ★★★★"]
    G128["Per-Group 128<br/>1 scale / 128 值<br/>开销: MB 级<br/>精度: ★★★★★"]
    G32["Per-Group 32<br/>1 scale / 32 值<br/>开销: 数 MB<br/>精度: ★★★★★+<br/>但 kernel 开销显著"]
    Coarse --> Channel --> G128 --> G32
    Note["常见起点: Per-Group 128"]
    G128 -.-> Note
    style G128 fill:#10b981,color:#333,stroke:none
    style Note fill:#2563eb,color:#fff,stroke:none

13.3.5 一张对比表

粒度	元数据开销	典型用途	精度排名
Per-Tensor	~几字节	激活量化、FP8 权重	⭐⭐
Per-Channel	KB 级（约 16 KB）	INT8 权重	⭐⭐⭐⭐
Per-Group 128	MB 级（约 1 MB）	INT4 权重（GPTQ/AWQ 常见配置）	⭐⭐⭐⭐⭐
Per-Group 32	数 MB（约 4 MB）	极致精度需求	⭐⭐⭐⭐⭐+

（括号内按上文 8192×8192 权重、FP16 scale 估算，不含 zero point。）

13.4 Kernel 演进：Marlin → Machete → Hopper FP8

量化对 kernel 的要求极为苛刻——不能简单"反量化 + GEMM"，必须尽量"一边读一边解一边算"。vLLM 源码里能看到几条并行路线：Marlin 处理 GPTQ/AWQ 等 mixed-precision 权重，Machete 给 Hopper 级别设备提供另一套 mixed-precision kernel，FP8 路线则通过 CUTLASS 或其他 FP8 kernel 走 8-bit 浮点矩阵乘。

先说清楚代码在哪。 Python 侧（vllm/model_executor/layers/quantization/）只负责"选哪条路、怎么摆权重"，真正的 CUDA/CUTLASS 实现全部在仓库根目录的 csrc/quantization/ 下，一条路线一个子目录：

csrc/quantization/
├── gptq_marlin/    gptq_marlin.cu + gptq_marlin_repack.cu + awq_marlin_repack.cu
├── marlin/         dense/ + sparse/（2:4 结构化稀疏，对应 gptq_marlin_24）+ qqq/
├── machete/        machete_mm_kernel.cuh / machete_prepacked_layout.cuh / generate.py …
├── cutlass_w8a8/   CUTLASS 的 INT8 / FP8 scaled_mm（文件最多的一个目录）
├── fp8/            common.cu（量化/scale）+ fp8_marlin.cu + nvidia/amd 各自的 quant_utils
├── awq/ gptq/ aqlm/ gguf/ fp4/ compressed_tensors/ …

注意 machete/ 里那个 generate.py——Machete 的 kernel 实例是用 Python 脚本生成 C++ 源码再编译的，因为 mixed-precision GEMM 的模板参数组合（激活类型 × 权重类型 × scale/zero point 类型 × 输出类型 × tile schedule）太多，手写不现实。group_size 反倒不在模板参数里：它作为运行时参数传进 kernel（csrc/quantization/machete/machete_mm_kernel.cuh:215-217），支持哪些由 Python 侧的 MACHETE_SUPPORTED_GROUP_SIZES = [-1, 128] 把关（vllm/model_executor/layers/quantization/utils/machete_utils.py:9）。所以"某个 group_size 不被支持"表现为选 kernel 时被拒绝、换走别的实现，而不是编译报错。

13.4.1 Marlin：Ampere 上的 INT4 极致优化

Marlin (arXiv:2408.11743) 是 Frantar 等人为 Ampere GPU 设计的 FP16 × INT4 GEMM kernel（论文主要在 A10 上评测，也测了 A100、RTX 3090 等）。关键技巧：

1. 权重离线重排：让每个线程一次 128-bit（16 字节）读取恰好拿到后续 4 个 Tensor Core 块所需的 32 个 INT4 值
2. 寄存器内反量化：权重经 cp.async 从全局内存异步搬进 shared memory，再读入寄存器，在寄存器里解成 FP16 并直接排成 Tensor Core 需要的寄存器布局
3. 加载与计算重叠：cp.async 的异步加载与 Tensor Core 计算流水重叠，下一个权重操作数的反量化与当前列的 mma.sync 交错执行
4. 条带式划分（striped partitioning）：按 SM 数把 tile 均匀分配，条带可以跨列，思路类似 stream-K，减少 wave 量化和全局归约

在 vLLM 里，GPTQMarlinLinearMethod 会按量化配置创建 packed 权重和 scale，并通过 mixed-precision kernel 包装选择可用实现；核心价值不是"用了 INT4"这四个字，而是权重布局、scale 布局和 kernel 的访存模式配套。只把权重存成 INT4、再在外面展开成 FP16，无法拿到 Marlin 这类 kernel 的主要收益。

13.4.2 Machete：Hopper 上的下一代

Machete 是面向 Hopper 级别能力的 mixed-precision kernel 路线。源码里 MacheteLinearKernel.get_min_capability() 返回 90（vllm/model_executor/layers/quantization/kernels/mixed_precision/machete.py:22-24），也就是 H100 这一代开始的能力级别。它的关键不只是"更快"，还包括一组明确约束：

• activation reordering 在跨设备切分输入特征时暂不支持（vllm/model_executor/layers/quantization/kernels/mixed_precision/machete.py:30-34）
• zero points 暂不支持（vllm/model_executor/layers/quantization/kernels/mixed_precision/machete.py:35-36）
• weight type、group size、shape 都要通过 query_machete_supported_quant_types、MACHETE_SUPPORTED_GROUP_SIZES 和 shape 检查（vllm/model_executor/layers/quantization/kernels/mixed_precision/machete.py:38-50）

不满足时并不报错：gptq_marlin 和 compressed-tensors 的 WNA16 方案经 choose_mp_linear_kernel 按候选列表 _POSSIBLE_KERNELS（vllm/model_executor/layers/quantization/kernels/mixed_precision/__init__.py:21-27）的顺序 Machete → AllSpark → Marlin → BitBLAS → Exllama 逐个尝试，取第一个算力达标、且自身 can_implement 判定能实现的 kernel；awq_marlin 不走这套选择，直接用 Marlin。

这类源码约束比一句"某 kernel 更快"更重要：生产环境里经常不是 GPU 不支持，而是某个 checkpoint 的 group size、zero point 或并行切分方式让它走不到目标 kernel。

13.4.3 CUTLASS FP8：原生 FP8 TensorCore

Hopper / Ada 之后的 GPU 让 FP8 成为推理工程里的重要路径。理想状态下，8-bit 浮点可以直接作为 matrix operand 进入支持 FP8 的矩阵乘 kernel，累加和输出 dtype 再按实现策略处理。

M = A @ B
A: FP8 E4M3, shape [M, K]
B: FP8 E4M3, shape [K, N]
Accumulator: FP32, shape [M, N]
Output: cast 回 FP8 E4M3 或 BF16

FP8 的优势是权重字节数减半，同时保留浮点动态范围；但它不是自动低误差，也不是自动最快。vLLM 的 Fp8Config.get_supported_act_dtypes() 返回 [torch.bfloat16, torch.half]（vllm/model_executor/layers/quantization/fp8.py:89-91），get_min_capability() 返回 80（vllm/model_executor/layers/quantization/fp8.py:93-95）——注意 80 是 Ampere 而不是 Hopper，也就是说 vLLM 的 FP8 配置在 A100 上也会被接受，但那不等于跑到了原生 FP8 TensorCore——算力低于 8.9 时，Fp8LinearMethod 改走 FP8 Marlin 的 weight-only（W8A16）路径（vllm/model_executor/layers/quantization/fp8.py:174-177）。真正的 CUTLASS FP8 是否可用要靠运行时的 cutlass_fp8_supported() 判断，它内部调 ops.cutlass_scaled_mm_supports_fp8(capability)（vllm/model_executor/layers/quantization/utils/w8a8_utils.py:36-43）。配置被接受，不等于跑到了快路径——这是 FP8 最常见的误判来源。

所以 FP8 的工程判断应当是：先确认 checkpoint 格式、activation scheme、GPU capability 和 vLLM 实际选择的 kernel，再用目标负载测延迟与质量。

13.4.4 三代 kernel 对比

graph TB
    subgraph "Ampere (A100)"
        M["Marlin<br/>INT4 weight + FP16/BF16 act<br/>重点: packed layout + fused dequant"]
    end
    subgraph "Hopper (H100)"
        MC["Machete<br/>capability >= 90<br/>受 zero point / group size / shape 约束"]
        FP8K["FP8 / CUTLASS path<br/>运行时检查 capability<br/>质量依赖 scale 与任务"]
    end
    subgraph "共同目标"
        Target["减少 HBM 字节<br/>避免中间 FP16 展开<br/>让 kernel 走到匹配路径"]
    end
    M --> Target
    MC --> Target
    FP8K --> Target
    style M fill:#2563eb,color:#fff,stroke:none
    style MC fill:#10b981,color:#333,stroke:none
    style FP8K fill:#7c3aed,color:#fff,stroke:none

13.4.5 kernel 演进的启示

从 Marlin → Machete → FP8 TensorCore，量化 kernel 的演进路径告诉我们一条工程规律：

硬件每代都会为"主流优化方向"增加原生指令。

• Volta（2017）：FP16 TensorCore
• Ampere（2020）：BF16/TF32 TensorCore + 2:4 结构化稀疏
• Hopper（2022）：FP8 TensorCore + WGMMA + TMA
• Blackwell（2024）：继续强化低精度矩阵乘路径

今天的软件技巧明天可能变成硬件路径。这也意味着量化不是一个静态格式选择题，而是 checkpoint 格式、GPU 代际、kernel 支持矩阵和服务负载的组合题。

13.5 四大主流格式的工程取舍

13.5.1 FP8：硬件友好，但仍要验证 scale 和质量

FP8 把 FP16 的 16 bit 压缩成 8 bit，两个变体：

• E4M3（4-bit 指数 + 3-bit 尾数 + 1 符号）：动态范围 ±448，精度高
• E5M2（5-bit 指数 + 2-bit 尾数 + 1 符号）：动态范围 ±57344，精度低但范围广

很多 FP8 权重 checkpoint 使用 E4M3；KV Cache 则可能因为范围需求选择不同 FP8 变体。vLLM 在 fp8.py 里根据 checkpoint quant config 解析 activation_scheme 和 weight_block_size，而不是写死一套固定默认值。

优点：

• 权重字节数相对 FP16/BF16 减半
• 浮点格式比整数格式更容易保留动态范围
• 在支持 FP8 的 kernel 上可以减少显式反量化和中间格式转换
• 部分模型（DeepSeek-V3、Llama-3.1）官方直接发 FP8 版本

缺点：

• 需要确认当前 GPU、CUDA/vLLM 构建和 CUTLASS/平台判断是否真正支持目标 FP8 kernel
• 压缩比只有 2× 不如 INT4 的 4×
• scale 缺失或 activation 分布漂移时，质量仍然可能下降

什么时候用：GPU 和 vLLM kernel 都支持 FP8、又希望在 2× 权重压缩和较低校准成本之间取得平衡时，优先把 FP8 纳入候选；最终仍以目标任务评测为准。

13.5.2 GPTQ：训练后量化里的经典

GPTQ（Frantar et al., ICLR 2023）用"逐列量化 + Hessian 信息修正"把权重量化到 3/4 bit（论文还用分组探索了约 2 bit 与三值量化；vLLM 的 gptq 路径接受 2/3/4/8 bit）。步骤：

1. 收集一小批校准数据（论文用 128 条从 C4 随机截取的 2048 token 片段）
2. 对每个线性层按列顺序量化
3. 每量化一列后，用 Hessian 信息把误差"补偿"到后续还没量化的列
4. 输出 INT4 权重 + per-group scale（常见配置 group_size=128）

优点：

• 压缩比 4×（INT4）
• 在很多 decoder-only 模型上是成熟路线
• 配合 Marlin/Machete 等 fused kernel 才能体现主要速度收益
• HuggingFace Hub 上大量预量化 GPTQ 模型可用

缺点：

• 需要校准数据（虽然只要几 MB）
• 量化过程需要几分钟到几小时（一次性）
• 对某些新架构（MoE、MLA）支持可能滞后

13.5.3 AWQ：激活感知的量化

AWQ（Lin et al., MLSys 2024）的核心洞察：激活值很大的那些通道更敏感。量化前把这些通道的权重"放大"、激活"缩小"（逆变换），量化后精度更好。

原式: y = x · W
AWQ:  y ≈ (x · diag(s)⁻¹) · Q(diag(s) · W)    # Q(·) 为分组量化；x 侧的 1/s 可融进前一个算子

$s$ 是按输入通道的缩放因子，论文取 $s = s_X^{\alpha}$：$s_X$ 是校准集上每个通道激活的平均幅值，$\alpha$ 在 $[0, 1]$ 上网格搜索。关键是：把重要通道的权重放大 $s$ 倍后，所在 group 的量化步长基本不变，这些通道的相对量化误差约缩小为原来的 $1/s$。

优点：

• 对 activation outlier 更敏感的模型，AWQ 的思路经常更合适
• 同样 4× 压缩比
• 社区生态活跃，大量预量化模型

缺点：

• 同样需要校准数据（论文称只统计每通道平均幅值、不做 GPTQ 那样的回归重建，对校准集依赖更小）
• 缩放指数 $\alpha$ 与权重裁剪阈值都要在校准集上搜索

GPTQ vs AWQ：两者在 vLLM 里都有 Marlin 路线，但 checkpoint 的 bits、group_size、zero point、desc_act、是否 MoE 都会影响最后走到哪个 kernel。实践里先看社区上哪个预量化 checkpoint 可信，再用自己的验证集比较，而不是假设二者速度和质量必然一致。

13.5.4 INT8 / W8A8：传统分水岭

早期还有一种 W8A8（权重 INT8 + 激活 INT8）路线。它没有被所有场景淘汰，但在支持 FP8 的新硬件上，很多部署会优先尝试 FP8：

• FP8 的浮点动态范围对激活分布更友好
• vLLM 的 W8A8 工具路径仍然存在，vllm/model_executor/layers/quantization/utils/w8a8_utils.py 里同时处理 CUTLASS FP8 支持检查和 scaled-mm 调用
• 对老硬件、特定 checkpoint 或已验证的 INT8 流水线，W8A8 仍可能是可维护的方案

因此 W8A8 更适合作为候选路线之一，而不是简单贴上"过时"标签。

13.5.5 格式对比总表

格式	权重压缩	质量风险	校准/准备成本	关键依赖	适合优先尝试的场景
FP8	约 2×	低到中，取决于 scale 与任务	低到中	GPU/kernel FP8 支持	新硬件、希望少改流水线
GPTQ INT4	约 4×	中，取决于校准集和 group size	中	GPTQ checkpoint + fused kernel	显存紧张、已有可信 GPTQ
AWQ INT4	约 4×	中，activation outlier 处理更强	中	AWQ checkpoint + fused kernel	有 AWQ 预量化模型
INT8 W8A8	约 2×	中	中	scaled-mm / INT8 路线	老硬件或已有 INT8 验证链路
bitsandbytes NF4	约 4×	中到高	低	运行时量化路径	快速实验、离线验证

13.6 vLLM 的可插拔量化注册

v0.8.5 源码在 vllm/model_executor/layers/quantization/__init__.py:8-36 定义了 27 个内置量化名字。新方法的加入不需要改 scheduler、worker 或 attention 核心路径，关键是两个抽象基类：QuantizeMethodBase 定义权重创建和 apply（vllm/model_executor/layers/quantization/base_config.py:11-43），QuantizationConfig 定义方法名、activation dtype、最低硬件能力、配置文件名、反序列化和层到量化方法的映射（vllm/model_executor/layers/quantization/base_config.py:60-145）。

# vllm/model_executor/layers/quantization/base_config.py:11 起（节选、简化：
# 省略了 get_supported_act_dtypes / get_config_filenames / from_config 等抽象方法）
class QuantizeMethodBase(ABC):
    """量化线性层的"怎么算"。"""
    @abstractmethod
    def create_weights(self, layer, *weight_args, **extra_weight_attrs): ...

    @abstractmethod
    def apply(self, layer, *args, **kwargs) -> torch.Tensor: ...

class QuantizationConfig(ABC):
    """某种量化方法的"元信息"。"""
    @abstractmethod
    def get_name(self) -> str: ...

    @classmethod
    @abstractmethod
    def get_min_capability(cls) -> int:
        """最低 GPU 算力（Volta=70, Turing=75, Ampere=80）"""

    @abstractmethod
    def get_quant_method(self, layer, prefix) -> Optional[QuantizeMethodBase]: ...

注册一个新量化方法：

# 示意代码：放在你自己的 Python 包里（不必改 vLLM 源码树），接口按 v0.8.5
from typing import Any, Dict, List, Optional

import torch

from vllm.model_executor.layers.linear import LinearBase, LinearMethodBase
from vllm.model_executor.layers.quantization import register_quantization_config
from vllm.model_executor.layers.quantization.base_config import (
    QuantizationConfig, QuantizeMethodBase
)

@register_quantization_config("my_fancy_quant")
class MyFancyQuantConfig(QuantizationConfig):
    def __init__(self, bits: int, group_size: int):
        super().__init__()  # 基类在这里初始化 packed_modules_mapping，模型加载时会用到
        self.bits = bits
        self.group_size = group_size

    def get_name(self) -> str:
        return "my_fancy_quant"

    def get_supported_act_dtypes(self) -> List[torch.dtype]:
        # 注意：基类里这个是普通实例方法，不是 classmethod
        return [torch.float16, torch.bfloat16]

    @classmethod
    def get_min_capability(cls) -> int:
        return 80  # 基类 docstring: 70=Volta, 75=Turing, 80=Ampere

    @staticmethod
    def get_config_filenames() -> List[str]:
        return ["quant_config.json"]

    @classmethod
    def from_config(cls, config: Dict[str, Any]) -> "MyFancyQuantConfig":
        return cls(bits=config["bits"], group_size=config["group_size"])

    def get_quant_method(self, layer, prefix) -> Optional[QuantizeMethodBase]:
        if isinstance(layer, LinearBase):
            return MyFancyLinearMethod(self)
        return None

class MyFancyLinearMethod(LinearMethodBase):  # 线性层方法的基类，签名见 linear.py:147
    def __init__(self, quant_config: MyFancyQuantConfig):
        self.quant_config = quant_config

    def create_weights(self, layer: torch.nn.Module,
                       input_size_per_partition: int,
                       output_partition_sizes: List[int], input_size: int,
                       output_size: int, params_dtype: torch.dtype,
                       **extra_weight_attrs):
        # 创建量化权重张量，并用 layer.register_parameter 挂到 layer 上
        ...

    def apply(self, layer: torch.nn.Module, x: torch.Tensor,
              bias: Optional[torch.Tensor] = None) -> torch.Tensor:
        # 量化矩阵乘
        ...

然后用户跑：

vllm serve my-model \
    --quantization my_fancy_quant

有两个前提。其一，--quantization 的 choices 在构建命令行解析器时就从 QUANTIZATION_METHODS 取定（vllm/engine/arg_utils.py:600-603），所以注册代码必须在此之前执行：vllm serve 会先调 load_general_plugins()（vllm/engine/arg_utils.py:1635-1638），加载 vllm.general_plugins 入口点组里的插件（vllm/plugins/__init__.py:79），注册一个入口点函数、在函数里 import 上面的模块即可。其二，若模型 config.json 里已有 quantization_config，它的 quant_method 必须与注册的名字一致，否则 _verify_quantization 在启动时报不匹配（vllm/config.py:815-820）。满足这两点后，vLLM 从注册表找到 MyFancyQuantConfig，先读 config.json 里的 quantization_config，没有再在模型目录里找 get_config_filenames() 列出的文件，交给 from_config 初始化（vllm/model_executor/model_loader/weight_utils.py:143-203），然后为支持的层创建量化权重。

零改动核心引擎——这是第 18 章讨论"可插拔接口是生态前提"的最好例证。量化方法可以在 model_executor/layers/quantization/ 内扩展，核心调度、worker 和 attention 路径只消费统一后的 QuantizationConfig/QuantizeMethodBase 合同。

13.6.1 注册表真实实现的四个工程决策

上面演示了用户视角怎么写新量化方法，但 vllm/model_executor/layers/quantization/__init__.py 本身只有 152 行，却集中体现了几个工程决策。完整看一遍 register_quantization_config（节选，省略 docstring）：

# __init__.py:8
QuantizationMethods = Literal[
    "aqlm", "awq", "deepspeedfp", "tpu_int8", "fp8", "ptpc_fp8",
    "fbgemm_fp8", "modelopt", "nvfp4", "marlin", "bitblas", "gguf",
    "gptq_marlin_24", "gptq_marlin", "gptq_bitblas", "awq_marlin",
    "gptq", "compressed-tensors", "bitsandbytes", "qqq", "hqq",
    "experts_int8", "neuron_quant", "ipex", "quark", "moe_wna16",
    "torchao",
]
QUANTIZATION_METHODS: list[str] = list(get_args(QuantizationMethods))

_CUSTOMIZED_METHOD_TO_QUANT_CONFIG = {}

def register_quantization_config(quantization: str):
    def _wrapper(quant_config_cls):
        if quantization in QUANTIZATION_METHODS:
            raise ValueError(
                f"The quantization method `{quantization}` is already exists.")
        if not issubclass(quant_config_cls, QuantizationConfig):
            raise ValueError("The quantization config must be a subclass of "
                             "`QuantizationConfig`.")
        _CUSTOMIZED_METHOD_TO_QUANT_CONFIG[quantization] = quant_config_cls
        QUANTIZATION_METHODS.append(quantization)
        return quant_config_cls
    return _wrapper

四个决策点每个都值得看一眼：

1. 内置列表是 Literal 静态类型、不是普通 list（vllm/model_executor/layers/quantization/__init__.py:8-37）。QuantizationMethods = Literal["aqlm", ...] 让代码侧的类型标注能收敛到一组已知字符串；运行时列表 QUANTIZATION_METHODS 则从 get_args 展开而来。当前统计有 27 种内置方法，说明量化已经不是一条路径，而是一组围绕不同 checkpoint 格式、kernel 和硬件后端的适配层。

2. 重名拒绝而非覆盖（vllm/model_executor/layers/quantization/__init__.py:65-73）。if quantization in QUANTIZATION_METHODS: raise ValueError(...) 意味着内置方法不能被用户注册的同名配置替换。这是重要的安全承诺：第三方扩展不能悄悄劫持 "fp8" 这种标准名字。

3. issubclass 运行时校验（vllm/model_executor/layers/quantization/__init__.py:69-71）。装饰器里还检查 quant_config_cls 必须是 QuantizationConfig 的子类。Python 动态类型下，如果你误把一个普通 class 丢进 @register_quantization_config，把错误提前到注册时比等到模型加载中途再失败更容易定位。

4. get_quantization_config 里的"懒导入"（vllm/model_executor/layers/quantization/__init__.py:79-84）。真正把名字映射到具体 Config 类的 method_to_config 字典是在函数内部构造的、不在模块顶层：

def get_quantization_config(quantization: str):
    if quantization not in QUANTIZATION_METHODS:
        raise ValueError(...)
    # lazy import to avoid triggering `torch.compile` too early
    from .aqlm import AQLMConfig
    from .awq import AWQConfig
    # ... 其余 import（整段共 26 行 import、27 个 Config 类，含 quark）
    method_to_config: dict[str, Type[QuantizationConfig]] = {
        "aqlm": AQLMConfig, "awq": AWQConfig,
        # ... 其余 25 项
    }
    method_to_config.update(_CUSTOMIZED_METHOD_TO_QUANT_CONFIG)
    return method_to_config[quantization]

为什么懒？源码注释写得很直接——「lazy import to avoid triggering torch.compile too early」（vllm/model_executor/layers/quantization/__init__.py:83）。只有当 get_quantization_config 被调用时（用户指定了量化方法，或 checkpoint 带量化配置、_verify_quantization 逐个探测各方法时），才 import 这些具体 Config 模块。这里不要脑补固定的冷启动秒数；真正可靠的结论是，量化模块的 CUDA/torch 依赖不应在顶层无条件加载。

这 152 行注册表是"可插拔架构 + 名字安全 + 启动边界"三件一体的范本。新增内置格式通常需要实现自己的 Config/Method，并在这个文件里加入 Literal、import 和 dict 映射；新增外部自定义格式则可以通过 register_quantization_config 进入 _CUSTOMIZED_METHOD_TO_QUANT_CONFIG。可演化性不是免费的，它来自这套明确的合同边界。

13.7 KV Cache 量化：再省一半显存

量化权重只是一半故事。LLM 推理里的另一大显存占用是 KV Cache。长上下文、高并发、GQA/MQA 头数、block 分配策略一起决定 KV 的峰值；在某些长上下文服务里，KV 占用会接近甚至超过压缩后的权重。

先看公式，而不是背一个固定数字：

KV bytes = 2(K,V) * layers * active_tokens * kv_heads * head_dim * bytes_per_elem

如果一个 80 层、8 个 KV head、head_dim=128 的 GQA 模型同时缓存 128×8192 个 token，FP16 KV 的理论数据量约为 320 GiB。实际服务中的峰值会受平均上下文长度、分页块空洞、并发波动和 prefix 复用影响，不应直接拿这个上限当部署容量。

Context 长了 KV 能压垮一切。KV 量化的意义就凸显出来。

13.7.1 --kv-cache-dtype fp8

vLLM 支持把 KV Cache 放到 FP8 路线：

vllm serve model-name \
    --kv-cache-dtype fp8 \
    --calculate-kv-scales  # 或用 checkpoint 里预存的 scale

--kv-cache-dtype 的合法取值由 CacheDType = Literal["auto", "fp8", "fp8_e4m3", "fp8_e5m2"] 限定（vllm/config.py:1291），默认 "auto"。

效果：

• KV 元素字节数从 FP16/BF16 的 2 字节降到 FP8 的 1 字节，理论上 KV 数据本体减半
• 同等显存下可能换来更高并发或更长上下文，但 block 管理、碎片和其他显存项会削弱线性收益
• 质量风险取决于 scale 来源、任务类型和 attention backend，必须单独评测

13.7.2 工程实现要点

KV 量化比权重量化更微妙：

• 权重是静态的，量化一次用一辈子；KV 是动态的，每步追加新 token 都要量化
• vLLM 的 BaseKVCacheMethod 会给 attention layer 创建 q_scale、k_scale、v_scale、prob_scale 参数（vllm/model_executor/layers/quantization/kv_cache.py:27-43）
• 如果 kv_cache_dtype 不是 auto 且没有动态计算 scale，vLLM 优先加载 checkpoint 里的 k_scale / v_scale；缺失时退回 1.0，并对 fp8_e4m3 的 1.0 scale 给出准确性警告（vllm/model_executor/layers/quantization/kv_cache.py:54-101）
• 当前实现要求 FP8 KV cache 的 k/v scale 是 per-tensor float，而不是任意细粒度 scale（vllm/model_executor/layers/quantization/kv_cache.py:79-82）
• calculate_kv_scales 的配置默认是 False；开启后会动态计算 k/v scale，否则从 checkpoint 读取，缺失时默认为 1.0（vllm/config.py:1347-1350）

13.7.3 何时开启

开：

• 长上下文或高并发让 KV 成为显存主项
• 生产 RAG / 长文档分析
• checkpoint 提供可信 scale，或你愿意承担动态 scale 的额外成本并完成质量评测

不开：

• 极度精度敏感的任务（数学推导、代码严格正确性）
• 短上下文、低并发，KV 不是瓶颈
• scale 缺失且评测显示输出质量不稳定

13.7.4 权重量化 + KV 量化的组合收益

两种量化组合使用，显存收益可以叠加，但加速不一定叠加：

权重量化	KV 量化	主要节省	主要风险
不开	不开	无	显存最大、质量最稳
FP8 权重	不开	权重约减半	依赖 FP8 kernel 和 scale
INT4 权重	不开	权重约减到 1/4	依赖校准集和 fused kernel
FP8 权重	FP8 KV	权重与 KV 都下降	KV scale 和 attention 质量要验证
INT4 权重	FP8 KV	显存压力最低	权重误差和 KV 误差可能叠加

因此不要把"INT4 权重 + FP8 KV"当成默认终极配置。它适合显存压力极高、验证集通过、且实际 kernel 路径确认无误的场景；如果服务更重视稳定质量，FP8 权重或不量化 KV 也可能是更好的折中。

13.8 量化的五个陷阱

13.8.1 高并发场景失效

如 13.2.1 节所见，batch 大了量化收益可能下降。如果你的服务是高并发 API 网关，先做 benchmark 确认量化仍在赚钱，不要盲目加 --quantization 就期待加速。

13.8.2 group_size 选错

GPTQ/AWQ 的 group_size：

• 128 经常是预量化 checkpoint 的常见配置
• 32 元数据更多，可能改善质量，也可能让 kernel 选择变差
• 256 元数据更少，但更容易受局部 outlier 影响
• -1 通常表示整行/整通道粒度，需要确认目标 kernel 是否支持

大部分场景先用 checkpoint 自带配置，不要手改 group_size。group_size 不是纯算法参数，它还决定权重布局和 kernel 支持矩阵：v0.8.5 的 Marlin 要求算力 ≥ 8.0、只接受 group_size ∈ {-1, 32, 64, 128}（vllm/model_executor/layers/quantization/utils/marlin_utils.py:21、40-41），Machete 只接受 {-1, 128}；像 256 这样的值会让 GPTQ/AWQ checkpoint 无法自动升级到 Marlin 路线。

13.8.3 双重量化

常见错误是 checkpoint 已经带量化配置，却又手动指定了不匹配的 --quantization。v0.8.5 不会"互相打架"地硬跑：_verify_quantization 发现两者不一致会在启动时直接抛 ValueError（vllm/config.py:815-820）。

# 可能错误：checkpoint 已带量化配置，手动覆盖成另一路线
vllm serve some-awq-checkpoint --quantization gptq

# 通常先让 vLLM 从 checkpoint 配置识别
vllm serve some-awq-checkpoint

如果 checkpoint 配置完整，通常先让 vLLM 自动识别。还有一个更隐蔽的坑：对 GPTQ/AWQ checkpoint，不指定 --quantization（或指定 marlin/gptq_marlin/awq_marlin）时，只要满足 Marlin 条件就会自动改用 gptq_marlin/awq_marlin；显式写 --quantization gptq 或 awq 反而会强制走较慢的原始 kernel（vllm/model_executor/layers/quantization/gptq_marlin.py:133-151、vllm/model_executor/layers/quantization/awq_marlin.py:104-121）。只有配置缺失、需要显式选择运行时量化，或你明确知道要覆盖默认识别时，才手动指定 --quantization。

13.8.4 KV 精度错配

权重量化到 INT4 + KV 不量化，可能出现"权重已经很小、KV 仍然很大"的现象。解决方法不一定是立刻打开 FP8 KV，而是先用上面的 KV 公式估算峰值，再验证 --kv-cache-dtype fp8 的质量风险。

13.8.5 精度 drift 的监测

量化部署后精度缓慢下降的可能原因：

• 输入分布和校准数据差异过大（比如校准用英文、生产以中文为主）
• Prompt 模板与校准时不同
• tokenizer 或聊天模板版本不一致

生产做法：上线时保留一份原始 FP16/BF16 或上一版本 shadow 服务，抽样对比输出，监控任务相关指标和人工抽检结果。发现 drift 后优先回滚量化 checkpoint、替换校准集或关闭 KV 量化。

13.8.6 五大陷阱速查表

陷阱	症状	解法
高并发失效	加了 quantization 没快	按 batch / prompt / decode 矩阵 benchmark
group_size 选错	精度或速度异常	优先沿用 checkpoint 配置
双重量化	启动报 ValueError，或被迫走较慢的 kernel	不随意覆盖 checkpoint 量化类型
KV 错配	权重小但显存仍爆	估算 KV 后再评估 FP8 KV
精度 drift	上线后用户抱怨	shadow 服务 + 监控

13.9 选型决策树

graph TB
    Start[部署量化决策]
    Start --> Q1{GPU / vLLM 是否支持目标 FP8 kernel?}
    Q1 -->|是| Q1Y[把 FP8 放入候选<br/>测质量与 TPOT]
    Q1 -->|否或不确定| Q2{显存是否紧张?}
    Q2 -->|否, 能放下 FP16| NoQ[不量化<br/>FP16 精度最好]
    Q2 -->|是| Q3{压缩比需求}
    Q3 -->|约 2× 即可| INT8[FP8 或 W8A8<br/>看硬件与 checkpoint]
    Q3 -->|约 4× 更好| Q4{已有可信预量化 checkpoint?}
    Q4 -->|有 GPTQ| GPTQ[GPTQ<br/>确认 group size / kernel]
    Q4 -->|有 AWQ| AWQ[AWQ<br/>确认 zero point / kernel]
    Q4 -->|都没, 想快速实验| BnB[bitsandbytes<br/>NF4 实时量化<br/>加载时间换准备时间]
    Q1Y --> KV{KV 是否成为显存瓶颈?}
    KV -->|是| KVE[评估 FP8 KV<br/>验证 scale 与质量]
    KV -->|否| End[部署候选中最稳的一项]
    KVE --> End
    NoQ --> End
    INT8 --> End
    GPTQ --> End
    AWQ --> End
    BnB --> End
    style Q1Y fill:#10b981,color:#333,stroke:none
    style NoQ fill:#2563eb,color:#fff,stroke:none
    style GPTQ fill:#7c3aed,color:#fff,stroke:none
    style AWQ fill:#f59e0b,color:#333,stroke:none

13.10 A100 vs H100：应该怎么实测

A100 和 H100 的量化对比，不能用一张跨环境数字表定论。应该固定模型、checkpoint、vLLM commit、CUDA/driver、tensor parallel、prompt/decode 分布、并发和采样参数，然后测下面这些项：

维度	A100 上重点看	H100 上重点看
权重量化	Marlin/GPTQ/AWQ 是否走到 fused kernel	Machete 或 FP8 路线是否可用
FP8	无原生 FP8 算力，走 FP8 Marlin（W8A16）	CUTLASS/平台 FP8 检查是否为真
KV 量化	是否真正缓解显存瓶颈	scale 来源和 attention backend 质量
延迟	TPOT、ITL、P95/P99	TPOT、ITL、P95/P99
质量	目标任务指标和人工样本	同一套指标，不能只看 perplexity

测完之后再下结论。常见结果是：低并发 decode 更容易从 INT4/FP8 中受益；高并发或长 prefill 场景收益会被 compute、attention、通信和调度稀释；KV FP8 是否值得，取决于 KV 是否真的是显存瓶颈。

13.10.1 源码统计：quantization/ 18160 行

把 §13.6 抽象提到的"20+ 种量化方法可插拔注册"在源码层统计——

用 find ... -name '*.py' | xargs wc -l 对 v0.8.5 源码统计，vllm/model_executor/layers/quantization/ 合计 18160 行。顶层有 29 个 .py 文件和 4 个子目录：

子目录/类别	行	角色
compressed_tensors/	3310	统一压缩格式适配，含 compressed_tensors_moe.py 1100 行 MoE 路线
utils/	3654	marlin_utils 413 + quant_utils 571 + fp8_utils 523 + int8_utils 459 + machete_utils + w8a8_utils 等共享 kernel helper
kernels/	1528	跨 backend kernel 包装
quark/	1047	AMD Quark 量化框架适配
顶层 .py 共 29 个	8621	量化格式实现 + base_config / schema / kv_cache / init

顶层实现文件按行数排序——

量化格式	行
fp8.py	832
gptq_marlin.py	643
awq_marlin.py	518
bitblas.py	459
moe_wna16.py	447
gptq_bitblas.py	438
modelopt.py	410
gguf.py	408
bitsandbytes.py	396
aqlm.py	374
hqq_marlin.py	328
gptq_marlin_24.py	295
gptq.py / awq.py / awq_triton.py / fbgemm_fp8 / ipex_quant / marlin / neuron_quant / ptpc_fp8 / qqq / experts_int8 / deepspeedfp / base_config / schema / kv_cache / __init__ 等	(各约 70~320)
顶层合计（29 个文件）	8621

两条值得记住的物理事实——

1. 18160 行不是一个小 feature。27 个内置量化名字背后，是 checkpoint 格式、kernel、硬件能力和模型结构的组合复杂度。
2. fp8.py 832 行、compressed_tensors_moe.py 1100 行。前者说明 FP8 不是一个 dtype 开关，后者说明 MoE 量化需要专门处理 expert 权重、scale 和 fused MoE 路径。

§13.6 的结论因此更具体：注册表让量化方法不必侵入 scheduler/worker，但复杂度没有消失，只是被收拢到 quantization 子系统内部。这个边界清晰，维护负担也真实。

13.11 本章小结

量化是 vLLM 生态里最成熟、进化最快的领域之一：

• 数学原理：LLM 权重分布有冗余，但 outlier、scale 粒度和校准集决定量化误差能否被控制
• 加速机制：量化直接减少 HBM 搬字节数，decode memory-bound 场景最容易受益
• 三种粒度：per-tensor / per-channel / per-group 的选择，是精度、元数据和 kernel 支持的三方权衡
• 三类 kernel：Marlin、Machete、FP8/CUTLASS 路线各有硬件和 checkpoint 约束
• 多种格式：FP8、GPTQ、AWQ、W8A8、bitsandbytes 都是候选，不存在无条件最优
• 可插拔注册：27 个内置量化名字通过 QuantizationConfig 映射到具体实现，自定义方法可通过注册表加入
• KV Cache 量化：长 context 高并发时可能显著省显存，但 scale 缺失会带来质量风险
• 五个陷阱：高并发收益缩水、group_size 不匹配、重复指定量化、KV scale 缺失、精度 drift

一句话记忆：

量化不是免费加速开关，而是把数值格式、kernel 和硬件放到同一张账本里重新定价。账算对了，它能省显存、降延迟；账算错了，它会把质量和稳定性一起带走。

---

源码导航

• 量化注册基类：vllm/model_executor/layers/quantization/base_config.py
• FP8 实现：vllm/model_executor/layers/quantization/fp8.py
• GPTQ + Marlin：vllm/model_executor/layers/quantization/gptq_marlin.py
• AWQ + Marlin：vllm/model_executor/layers/quantization/awq_marlin.py
• Machete kernel：vllm/model_executor/layers/quantization/kernels/mixed_precision/machete.py
• Machete 工具：vllm/model_executor/layers/quantization/utils/machete_utils.py
• CUTLASS FP8：vllm/model_executor/layers/quantization/utils/w8a8_utils.py
• Marlin 工具：vllm/model_executor/layers/quantization/utils/marlin_utils.py
• KV Cache 量化：vllm/model_executor/layers/quantization/kv_cache.py
• CUDA kernel 实现（不在 Python 侧）：csrc/quantization/——gptq_marlin/、marlin/、machete/（含代码生成脚本 generate.py）、cutlass_w8a8/、fp8/

论文

• Frantar et al., "GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers", ICLR 2023 (arXiv:2210.17323)
• Lin et al., "AWQ: Activation-aware Weight Quantization for LLM Compression and Acceleration", MLSys 2024 (arXiv:2306.00978)
• Frantar et al., "MARLIN: Mixed-Precision Auto-Regressive Parallel Inference on Large Language Models", 2024 (arXiv:2408.11743)
• Micikevicius et al.（NVIDIA / Arm / Intel）, "FP8 Formats for Deep Learning", 2022 (arXiv:2209.05433)