vLLM 推理内核深度解析

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

作者杨艺韬 · 8,609 字

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

“Perfection is not attainable, but if we chase perfection we can catch excellence.” — Vince Lombardi

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

本章要点

从量化的数学原理出发：为什么把 FP16 的 65536 种取值压缩到 INT4 的 16 种还能保持模型能力
理解为什么量化在 memory-bound 的 decode 阶段可能同时省显存和加速 —— 不是两个独立的事，是同一个字节数下降现象的两面
读懂 per-tensor / per-channel / per-group 三种粒度的精度-速度权衡曲线
走一遍 Marlin / Machete / CUTLASS FP8 三类量化 kernel，理解为什么权重布局、反量化位置和 GPU 代际会改变结果
掌握 FP8 / GPTQ / AWQ / W8A8 / bitsandbytes 等格式的工程取舍：精度、速度、硬件要求、校准成本
理解 vLLM 的可插拔量化注册系统：当前源码中 27 种内置量化名字如何映射到具体 Config
知道 KV Cache 量化（FP8-KV）的收益、代价、scale 来源与适用场景
避开 5 类量化陷阱：高并发收益缩水、group_size 不匹配、重复指定量化、KV scale 缺失、精度 drift
拿到决策树式选型指南 + A100/H100 对比时应该怎么测

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:#ef4444,color:#fff,stroke:none
    style F8 fill:#f59e0b,color:#fff,stroke:none
    style I4 fill:#10b981,color:#fff,stroke:none
    style I4_bad fill:#94a3b8,color:#fff,stroke:none

注意最后一行——“INT4 朴素反量化”的路径。如果你的做法是：先读 INT4 → 展开成 FP16 → 再做 FP16 GEMM，那读取确实只有 35 GB，但中间 FP16 展开占了 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：

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 个值

实际用得很少，主要用于激活量化（下文）。

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:#fff,stroke:none
    style Note fill:#3b82f6,color:#fff,stroke:none

13.3.5 一张对比表

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

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 浮点矩阵乘。

13.4.1 Marlin：Ampere 上的 INT4 极致优化

Marlin (arXiv:2406.09792) 是 Frantar 等人为 Ampere GPU (A100) 设计的 INT4 × FP16 GEMM kernel。关键技巧：

权重重排：把 INT4 权重 pack 成 GPU 友好的形状，load 一次 128-bit 拿 32 个 INT4 值
寄存器内反量化：INT4 → FP16 在寄存器完成，不走 shared memory
异步 mma.sync：计算与下一块权重的加载 overlap
Persistent threads：一个 CUDA thread block 跑到最后不退出，减少 launch overhead

在 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（kernels/mixed_precision/machete.py:22-24），也就是 H100 这一代开始的能力级别。它的关键不只是”更快”，还包括一组明确约束：

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

这类源码约束比一句”某 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 支持 BF16/FP16 activation，最低 capability 返回 80（fp8.py:89-99），而 CUTLASS FP8 是否可用还要通过 cutlass_scaled_mm_supports_fp8(capability) 运行时判断（utils/w8a8_utils.py:35-42）。

所以 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:#3b82f6,color:#fff,stroke:none
    style MC fill:#10b981,color:#fff,stroke:none
    style FP8K fill:#8b5cf6,color:#fff,stroke:none

13.4.5 kernel 演进的启示

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

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

Volta（2017）：FP16 TensorCore
Ampere（2020）：BF16 TensorCore + INT8 TensorCore + 稀疏
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/8 bit。步骤：

收集一小批校准数据（通常 128-512 条文本）
对每个线性层按列顺序量化
每量化一列后，用 Hessian 信息把误差”补偿”到后续还没量化的列
输出 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 / s) · (s · W_quant)
     = x · (s · W_quant) / s

$s$ 是 per-channel 的缩放因子，基于激活的 absmax 估计。关键是 $s$ 的放大让 W 的重要通道获得更大 scale，量化精度提升。

优点：

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

缺点：

同样需要校准数据
量化算法比 GPTQ 稍复杂

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 工具路径仍然存在，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 的可插拔量化注册

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

# vllm/model_executor/layers/quantization/base_config.py
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, Ampere=80, Hopper=90）"""

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

注册一个新量化方法：

# vllm/model_executor/layers/quantization/my_quant.py
from typing import Any, Dict, List, Optional

import torch

from vllm.model_executor.layers.linear import LinearBase
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):
        self.bits = bits
        self.group_size = group_size

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

    @classmethod
    def get_supported_act_dtypes(cls) -> List[torch.dtype]:
        return [torch.float16, torch.bfloat16]

    @classmethod
    def get_min_capability(cls) -> int:
        return 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(QuantizeMethodBase):
    def create_weights(self, layer, ...):
        # 创建量化权重张量
        ...

    def apply(self, layer, x):
        # 量化矩阵乘
        ...

然后用户跑：

vllm serve my-model \
    --quantization my_fancy_quant

vLLM 自动找到注册表里的 MyFancyQuantConfig，配合配置文件里的参数初始化，然后为支持的层创建量化权重。

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

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

上面演示了用户视角怎么写新量化方法，但 vllm/model_executor/layers/quantization/__init__.py 本身只有 153 行、每一行都对应一个工程决策。完整看一遍 register_quantization_config：

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

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

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

4. get_quantization_config 里的”懒导入”（line 82-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
    # ... 25 个 import
    method_to_config: dict[str, Type[QuantizationConfig]] = {
        "aqlm": AQLMConfig, "awq": AWQConfig, ...
    }
    method_to_config.update(_CUSTOMIZED_METHOD_TO_QUANT_CONFIG)
    return method_to_config[quantization]

为什么懒？源码注释写得很直接：避免过早触发 torch.compile（__init__.py:83）。只有当用户指定某个量化方法、get_quantization_config 被调用时才 import 具体 Config。这里不要脑补固定的冷启动秒数；真正可靠的结论是，量化模块的 CUDA/torch 依赖不应在顶层无条件加载。

这 153 行注册表是”可插拔架构 + 名字安全 + 启动边界”三件一体的范本。新增内置格式通常需要实现自己的 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 元素字节数从 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 参数（kv_cache.py:27-43）
如果 kv_cache_dtype 不是 auto 且没有动态计算 scale，vLLM 优先加载 checkpoint 里的 k_scale / v_scale；缺失时退回 1.0，并对 fp8_e4m3 的 1.0 scale 给出准确性警告（kv_cache.py:54-101）
当前实现要求 FP8 KV cache 的 k/v scale 是 per-tensor float，而不是任意细粒度 scale（kv_cache.py:79-82）
calculate_kv_scales 的配置默认是 False；开启后会动态计算 k/v scale，否则从 checkpoint 读取，缺失时默认为 1.0（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 支持矩阵。

13.8.3 双重量化

常见错误是 checkpoint 已经带量化配置，却又手动指定了不匹配的 --quantization，让 vLLM 的自动识别和用户覆盖互相打架。

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

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

如果 checkpoint 配置完整，通常先让 vLLM 自动识别。只有配置缺失、需要显式选择运行时量化，或你明确知道要覆盖默认识别时，才手动指定 --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 配置
双重量化	加载报错或输出异常	不随意覆盖 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 -->|是| End[评估 FP8 KV<br/>验证 scale 与质量]
    KV -->|否| End[部署候选中最稳的一项]
    NoQ --> End
    INT8 --> End
    GPTQ --> End
    AWQ --> End
    BnB --> End

    style Q1Y fill:#10b981,color:#fff,stroke:none
    style NoQ fill:#3b82f6,color:#fff,stroke:none
    style GPTQ fill:#8b5cf6,color:#fff,stroke:none
    style AWQ fill:#f59e0b,color:#fff,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	可能受 kernel 支持限制	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 对本地源码统计，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
`gptq_marlin_24.py`	295
`gptq.py / awq.py / awq_triton.py / aqlm / bitsandbytes / fbgemm_fp8 / hqq_marlin / ipex_quant / marlin / neuron_quant / ptpc_fp8 / qqq / experts_int8 / deepspeedfp / gguf` 等	(各 100~400)
顶层合计	8621

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

18160 行不是一个小 feature。27 个内置量化名字背后，是 checkpoint 格式、kernel、硬件能力和模型结构的组合复杂度。
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/ 合计 18160 行；__init__.py 里有 27 个内置量化名字；fp8.py 832 行，compressed_tensors_moe.py 1100 行。量化在 vLLM 里是一套完整子系统，不是一个 CLI 参数。

源码导航

量化注册基类：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

论文

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: a Mixed-Precision LLM Inference Kernel Optimized for Modern GPUs”, 2024 (arXiv:2406.09792)

NVIDIA, “FP8 Formats for Deep Learning”, 2022 (arXiv:2209.05433)

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

13.1 量化的数学本质

13.1.1 “过剩的精度”

13.1.2 为什么 FP16 是”过剩”的

13.1.3 量化公式

13.1.4 量化误差的量化

13.1.5 离群值（Outlier）的双刃剑

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

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

13.2.2 一个简化的 roofline 视角

13.3 量化的三种粒度

13.3.1 Per-Tensor

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

13.3.3 Per-Group

13.3.4 精度-速度曲线对比

13.3.5 一张对比表

13.4 Kernel 演进：Marlin → Machete → Hopper FP8

13.4.1 Marlin：Ampere 上的 INT4 极致优化

13.4.2 Machete：Hopper 上的下一代

13.4.3 CUTLASS FP8：原生 FP8 TensorCore

13.4.4 三代 kernel 对比

13.4.5 kernel 演进的启示

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

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

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

13.5.3 AWQ：激活感知的量化

13.5.4 INT8 / W8A8：传统分水岭

13.5.5 格式对比总表

13.6 vLLM 的可插拔量化注册

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

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

13.7.1 --kv-cache-dtype fp8

13.7.2 工程实现要点

13.7.3 何时开启

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

13.8 量化的五个陷阱

13.8.1 高并发场景失效

13.8.2 group_size 选错

13.8.3 双重量化

13.8.4 KV 精度错配

13.8.5 精度 drift 的监测

13.8.6 五大陷阱速查表

13.9 选型决策树

13.10 A100 vs H100：应该怎么实测

13.10.1 源码统计：quantization/ 18160 行

13.11 本章小结

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

13.7.1 `--kv-cache-dtype fp8`