vLLM 量化方法详解：原理、算法与实现

本文档聚焦于 vLLM 中与 NVIDIA GPU 兼容的通用量化方法，排除了框架特定方案（GGUF、Compressed-Tensors、Quark、Hummingbird）和硬件厂商专有方案（ModelOpt、INC、NVFP4）。涵盖：数学原理、量化/反量化公式、位宽与分组结构、核心算法步骤、Kernel 级实现细节及硬件要求。

目录

1. FP8 W8A8（动态/静态量化）
2. AWQ（激活感知权重量化）
3. GPTQ / Marlin
4. AWQ Marlin
5. BitsAndBytes（NF4 / QLoRA）
6. MXFP4（OCP MX 规范）
7. Online Quantization（在线量化）
8. TurboQuant
9. FP8 / INT8 KV Cache
10. 对比总结

1. FP8 W8A8（动态/静态量化）

论文

• FP8 格式标准: FP8 Formats for Deep Learning (Micikevicius et al., 2022, arXiv:2209.05433)
• vLLM 实现: 主要面向 H100/H200 GPU，利用 torch._scaled_mm 调用 Tensor Core。

概述

FP8 W8A8 将权重和激活均量化为 FP8 格式。支持两种调度模式：

• 静态量化：校准集上预计算 scale，推理时固定。
• 动态量化：推理时实时统计最大值计算 scale。

量化公式

FP8 有两种指数位宽子格式：

量化核心操作：

q = round_clamp(x / scale, max_fp8, min_fp8)
scale = max(|x|) / max_fp8   (每个量化组独立计算)

scale 粒度策略（vLLM 实现）:

• Per-tensor：整个张量共享一个 scale
• Per-token：每行（token）独立 scale
• Per-block：沿 token 和 channel 维度划分块（如 128×128），每块独立 scale

反量化

x̂ = q × scale

Kernel 实现

// 伪代码：FP8 GEMM 内核
// 输入: A_fp8 (M×K), B_fp8 (K×N), scale_a, scale_b
// 调用 H100 Tensor Core:
//   MMA: FP8 → FP32 累加
//   每 16 个时钟周期完成 16×8×16 (A×B) 的矩阵乘法
//   scale 融合: C = A_fp8 × B_fp8 × scale_a × scale_b

路径：vllm/model_executor/layers/quantization/fp8.py

硬件要求

• NVIDIA H100/H200/B200 (SM90+) 原生 FP8 Tensor Core
• 实测约 2× 性能提升 over FP16
• 动态量化支持 SM75+，但性能下降（无原生 FP8 MMA）

2. AWQ（激活感知权重量化）

论文

• AWQ: Activation-aware Weight Quantization for LLM Compression and Acceleration (Lin et al., 2024, arXiv:2306.00978)

概述

AWQ 是一种仅权重的 4-bit 量化方法，利用激活值分布识别权重中的"显著通道"（约 1%），仅对这部分通道施加逐通道缩放因子 ∆ 来吸收量化误差，而非直接保留为高精度。

核心洞察

权重中约 1% 的通道对 LLM 输出影响最大（通过激活幅度 = |x| 衡量）。与其保留为 FP16，不如对权重施加缩放因子 ∆ > 1，减小最大值范围 → 降低量化误差。

量化公式

Scale 搜索：

1. 将权重 W 按通道划分为组，每组 g 个权重
2. 对每个通道 i，计算激活幅度 s_x = mean(|x_i|)
3. 引入缩放因子 s = s_x^α，α 在 [0,1] 中搜索

伪量化：

Δ_ch = max(|W_ch|) / max(q)          // 原始 scale
Δ'_ch = Δ_ch / s_ch                  // 经缩放因子调整
W_q = quantize(W_ch · s_ch, Δ'_ch)   // 先缩放再量化

数学等价性：

量化误差主要由边界截断引起。设：

Q(W · Δ) ≈ Q(W) · Δ   当 Δ 接近 1

而 AWQ 的核心在于：对于显著通道（s_x 大），使用 s_ch > 1 将权重"压缩"进量化网格的密集区域，以降低 MSE。

反量化与 GEMM 融合

W_deq = W_q · Δ'_ch = W_q · (Δ_ch / s_ch)
Y = X · W_deq^T = X · (W_q · Δ'_ch)^T

AWQ 的关键技巧是在 GEMM 中使用组内共享 scale，并在写回时融合 per-channel 反量化。

权重布局

输入: W_fp16, shape [OC, IC]
输出: W_q, shape [OC, IC/g] 存储为 uint32 pack（每个 uint32 装 8 个 4-bit 权重）
      scale, shape [OC, IC/g]

打包规则：一个 uint32 存储 8 个 4-bit 值，按 0-7 顺序从低到高存放。

Kernel 细节

路径：vllm/model_executor/layers/quantization/awq.py

AWQ 在 vLLM 中通过融合 GEMM 内核执行：

// 每个线程束处理:
//   1. 从 global memory 加载 packed uint32 权重
//   2. 解包为 8× 4-bit 值
//   3. 加载 scale 并反量化: w_fp16 = w_q * scale
//   4. 使用 CUDA Core 执行 FP16 GEMM

硬件要求

• 任何 NVIDIA GPU（SM70+），无需 Tensor Core 支持
• 在 A100/H100 上使用 cuBLAS FP16 GEMM 性能最佳

3. GPTQ / Marlin

论文

• GPTQ: GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers (Frantar et al., 2023, arXiv:2210.17323)
• Marlin: Marlin: A Fast 4-bit Inference Kernel for LLMs (Frantar et al., 2024, arXiv:2402.11764)

概述

GPTQ 是一种仅权重的 4-bit 后训练量化算法，基于 Optimal Brain Quantizer (OBQ) 的扩展。Marlin 是一种专为 GPTQ 格式设计的高性能 GPU Kernel，利用 4×2 组结构和特化 tiling 策略实现接近 FP16 GEMM 的吞吐。

GPTQ 算法细节

问题设定

给定预训练权重矩阵 W ∈ ℝ^{k×d}（IC×OC）和一小组校准数据 X ∈ ℝ^{d×n}（从预训练数据采样），希望找到量化权重 Ŵ 使得：

min || WX - ŴX ||²_F

等价于最小化输出 feature map 的 MSE。

量化顺序

GPTQ 逐列（column-by-column）量化，列顺序由 Hessian 对角线元素决定。

Hessian 计算

对当前需要量化的列 j，计算：

H_j = 2 X_j X_j^T + λI

其中 X_j 是 X 的第 j 行（激活值），λ 是阻尼项。

Cholesky 分解

H_j = LL^T

使用 Cholesky 分解高效求解最优量化增量。

核心迭代

对每一列 j（按 Hessian 优先级排序）：

// 1. 量化当前权重
q_j = round_clamp(w_j / scale, q_min, q_max)

// 2. 计算量化误差
err = w_j - q_j * scale

// 3. 计算误差补偿向量 (OBQ 核心)
δ = -(err / H_j[j,j]) · (H_j)^{-1}[-1, :]

// 4. 将 δ 传播到剩余未量化权重
w_remaining += δ

// 5. 更新 H 矩阵（高斯消元）
H = H - (1/H_j[j,j]) · H_j[:,j] · H_j[j,:]

关键创新：相比于逐个权重独立量化的 OBQ，GPTQ 将列合并为块（block，如 128 列），使用批处理 Cholesky 分解，一次处理整块，显著加速。

分组结构

// 组大小 g = 128
// 每组共享一个 scale 和 zero-point
// W 按 [OC, IC/g] 分块
for each group i:
    group_weights = W[:, i*g : (i+1)*g]
    scale_i = max(|group_weights|) / (2^(bits-1) - 1)
    zero_point_i = 0 (对称量化)
    W_q[:, i] = pack_4bit(round(group_weights / scale_i))

Marlin Kernel 细节

核心设计原则

1. 权重布局重置：将权重重排为 Marlin 专有格式（16×2 块结构），提升缓存局部性
2. 异步加载：通过 cp.async 隐藏内存延迟
3. 共享内存双缓冲：当前块计算的同时预取下个块
4. Warp 级组广播：group-wise scale 在 warp 级别共享

块结构 (16×2)

原始权重: [16 tiles] × [groups_per_tile = 2]
处理方式:
  tile 0: 权重 [16×64] + 对应的 group scales
  tile 1: 权重 [16×64] + 对应的 group scales

中间布局

Marlin 将权重从 [OC, IC] 重排为：

[OC / 16, IC / (64 * 2), 16, 64 * 2]

每个 tile 的 [16, 64] 部分连续存储，scale 预加载到寄存器。

Kernel 伪代码

// 输入: W_packed [OC/16, IC/(128), 16, 128] (uint32)
//        scales [OC/16, IC/128, 16, 2] (fp16)
//        input A [M, IC] (fp16)
// 输出: C [M, OC] (fp16)

// 步骤
for each output_channel_block of size 16:
    for each input_channel_block of size 128:
        // 1. 从 global memory 加载 packed weights
        // 2. 解包为 32 个 4-bit 值
        // 3. 加载 scale (per group)
        // 4. 反量化: w_fp16 = w_q * scale
        // 5. 使用 FP16 FMA 计算局部部分和
        // 6. 将部分和写入 shared memory
    // 跨 warp 累加
    // 写回 global memory

硬件要求

• GPTQ 生成：任何 CUDA GPU
• Marlin 推理：NVIDIA Ampere (SM80+) 或更新架构
• 实测达到 A100 FP16 GEMM 理论带宽的 85%+

4. AWQ Marlin

概述

AWQ Marlin 并不是一种新的量化算法，而是将 AWQ 格式的权重转换为 Marlin 内核可消费的格式。

为什么需要？

AWQ 的权重打包格式和分组结构与 GPTQ 不同，无法被 Marlin kernel 直接读取。AWQ Marlin 模块执行格式转换：

AWQ 格式                   AWQ Marlin (Marlin 格式)
───────                    ────────────────────────
W_q [OC, IC/8]            W_q [OC/16, IC/128, 16, 128]
scale [OC, IC/g]           scale [OC/16, IC/128, 16, 2]
                          + g_idx [OC, IC/128] (group index lookup)

路径：vllm/model_executor/layers/quantization/awq_marlin.py

格式转换过程

def awq_to_marlin(w_q_awq, scale_awq, group_size):
    // 1. 将 AWQ 的 [OC, IC/8] 展开为 [OC, IC]
    // 2. 对每个 4-bit 值应用 group index 映射
    // 3. 重新打包为 Marlin 的 16×2 tile 排列
    // 4. scale 矩阵从 [OC, IC/g] → [OC/16, IC/128, 16, 2]
    // 5. 可选: 应用 AWQ 的 per-channel 缩放因子 ∆
    return w_q_marlin, scales_marlin

性能

• 转换仅需一次，在前处理阶段完成
• 推理时直接使用 Marlin kernel，性能同原生 Marlin
• 系统自动检测模型格式，在 awq_marlin.py 中完成路由

5. BitsAndBytes（NF4 / QLoRA）

论文

• QLoRA: Efficient Finetuning of Quantized Language Models (Dettmers et al., 2023, arXiv:2305.14314)

概述

BitsAndBytes 提供 NF4 （NormalFloat4）和 INT8 两种量化格式，核心创新在于NF4 数据类型和双重量化（Double Quantization）。

NF4 数据类型

设计原则

NF4 是一种信息论最优的 4-bit 数据类型：假设权重服从正态分布 N(0, σ²)，而信息论告诉我们，在给定 bit-width 下，等概率量化（即每个量化值对应的概率质量相等）使量化误差最小。

构造步骤

1. 正态分位数计算： 给定 4-bit → 16 个量化级别。将 N(0,1) 的 CDF 均分为 16 个等概率区间：

   quantiles[i] = Φ^(-1)((i + 0.5) / 16)    for i = 0, 1, ..., 15
   

   其中 Φ 是标准正态 CDF。

2. 归一化： 将 quantiles 归一化到 [-1, 1] 范围：

   nf4[i] = quantiles[i] / max(|quantiles|)
   

3. 最终值表（vLLM 实现）：

   [-1.0, -0.6968, -0.5251, -0.3940, -0.2844, -0.1848, -0.0910, 0.0,
    0.0799, 0.1604, 0.2495, 0.3521, 0.4781, 0.6395, 0.8654, 1.0626]
   

   值中心点在 0 附近最密集，尾部较稀疏——与正态分布概率密度匹配。

量化/反量化

// 量化
absmax = max(|W|)
scale = absmax / 1.0626   // 最大量化值
W_norm = W / scale         // 归一化到 [-1.0626, 1.0626]
index = argmin_j |nf4[j] - W_norm|   // 查找最近邻
W_q[idx] = index          // 存储为 4-bit index

// 反量化
W_deq[i] = nf4[q_idx[i]] * scale

双重量化（Double Quantization）

双重量化是对 scale 本身的二次量化。标准的逐 block 量化会产生大量 FP32 scale 值（每 block 一个），双重量化将它们进一步压缩为 INT8。

// 第一层量化
block_size = 64
for each block i:
    w_block = W[i*bs : (i+1)*bs]
    absmax = max(|w_block|)
    // 存储每个 block 的 scale (FP32)：共 OC/bs 个 scale
    scales_f32[i] = absmax / nf4_max

// 第二层量化：对 scales_f32 再做 INT8 量化
sub_block_size = 256   // 每 256 个 FP32 scales 为一组
for each sub_block j:
    s_absmax = max(|scales_f32[j*256 : (j+1)*256]|)
    s_scale = s_absmax / 127
    s_q[j] = round(scales_f32[j*256 : (j+1)*256] / s_scale)
    // 存储 s_q 为 INT8, s_scale 为 FP32

最终存储：

• 模型权重：NF4 + scale (FP32 per block)
• scale 本身也量化了（每个 sub_block 仅一个 FP32, 其余 INT8）

存储开销对比

省去了内层 32-bit scale 的大部分存储。

INT8 量化

q_i = round_clamp(x_i / scale, -127, 127)
scale = max(|x|) / 127

使用 CUDA Core 的 INT8 矩阵乘法（tcgen，类似于 int8 gemm 但使用 8-bit 整数）。

硬件要求

• NF4 推理：任何 CUDA GPU（SM70+）
• INT8 推理：Turing+ (SM75+) 的 INT8 Tensor Core 提升性能
• Double Quantization 对所有 GPU 透明
• 在 vLLM 中主要用于 LoRA 适配器场景

6. MXFP4（OCP MX 规范）

论文

• OCP Microscaling Formats (MX) Specification (Open Compute Project, 2023)
• Microscaling Data Formats for Deep Learning (Rouhani et al., 2023, arXiv:2310.10537)

概述

MXFP4 是 OCP MX 标准中定义的微观缩放格式，使用 E2M1 4-bit 元素 + E8M0 8-bit 共享 scale，每组 32 个元素共享一个 scale。

数据格式

MXFP4 元素: E2M1 (1-2-1)
┌─┬──┬─┐
│S│EE│M│
└─┴──┴─┘
S: 符号位 (1 bit)
E: 指数位 (2 bit), bias=1
M: 尾数位 (1 bit)

共享 scale: E8M0 (0-8-0)
┌──────────┐
│ EEEEEEEE │
└──────────┘
E: 8-bit 无偏指数, bias=0

值范围

分组量化

block_size = 32
for each block i:
    x_block = X[i*32 : (i+1)*32]
    // 1. 计算 shared scale
    max_val = max(|x_block|)
    e_scale = ceil(log2(max_val))   // E8M0: 8-bit 无偏指数
    scale_val = 2^(-e_scale)
    // 2. 量化每个元素
    for j in 0..31:
        mxp4_val = round_clamp(x_block[j] * scale_val, mxp4_max, 0)
    // 3. 存储: 32 × 4-bit + 1 × 8-bit = 136 bit

vLLM 实现

路径：vllm/model_executor/layers/quantization/mxfp4.py

class MXFP4Config(QuantizationConfig):
    def __init__(self, weight_block_size: Tuple[int, int] = [32, 32]):
        ...

class MXFP4LinearMethod(QuantizeMethodBase):
    def apply(self, weights, x):
        // 1. 对权重做 MXFP4 量化 (离线)
        // 2. 推理时在 kernel 内部分组反量化
        // 3. 执行 Hopper FP4 Tensor Core MMA (SM100+)

硬件与 Kernel

• SM90 (H100): 无原生 FP4 Tensor Core，通过软件模拟反量化为 FP16
• SM95/100 (B200+): 原生 FP4 Tensor Core 支持（计划中）
• vLLM 当前 SM100+ 上使用 torch.fp4_mx4_scaled_mm（尚在开发）

MX 格式族

7. Online Quantization（在线量化）

概述

Online Quantization 是在模型加载时对权重进行实时量化，不需要事先运行校准集。vLLM 支持 FP8、MXFP8、INT8 三种在线量化变体。

路径：vllm/model_executor/layers/quantization/online/

统一量化流程

def process_weights_after_loading(self, weights_dict):
    for name, w in weights_dict.items():
        if not self.should_quantize(name):
            continue
        // 在线计算 scale
        scale = self.compute_scale(w, self.quant_config)
        // 量化权重 (仅执行一次)
        w_q = quantize(w, scale, self.dtype)
        // 存储 w_q + scale 替代原始权重
        weights_dict[name] = (w_q, scale)

FP8 在线量化

// 动态计算 per-group scale
if block_size == -1:   // per-tensor
    scale = max(|W|) / 448
else:                   // per-block
    scale = [max(|W_block[i]|) / 448 for i in range(num_blocks)]

MXFP8 在线量化

使用 E8M0 scale 格式：

scale = max(|x_block|) // 2  // 取整到 2 的幂
scale_bits = pack_e8m0(scale)  // 共用 8-bit 指数

INT8 在线量化（仅 MoE）

// 对称量化到 [-127, 127]
scale_ch = 127 / max(|W_ch|)
W_i8[i] = round(W[i] * scale_ch)

硬件要求

• FP8: H100+（需要 FP8 Tensor Core）
• MXFP8: 任何 GPU（通过软件模拟）
• INT8: Turing+ (SM75+)

8. TurboQuant

论文

• (暂未正式发表；实现见 vLLM turboquant/ 目录)
• 使用 Hadamard 旋转 + Lloyd-Max MSE 最优量化

概述

TurboQuant 是 vLLM 中的 KV Cache 量化方案，通过正交变换和自适应量化极大降低量化误差。

算法流程

输入: Key 张量 K (num_tokens × head_dim)
       Value 张量 V (num_tokens × head_dim)
       group_size (默认 16)

// === 预处理 ===
// 1. 可选的 norm 校正（推导 qk 无偏输出）
K = K * (head_dim / ||K||_2)      // norm_correction=True 时

// 2. Hadamard 旋转
H_d = hadamard_matrix(head_dim)    // Sylvester 构造
K = K @ H_d.T                       // 消除异常值维度

// === 量化 ===
// Key: Lloyd-Max 量化 (per-coordinate, per-group)
for each block i (size group_size):
    stats = compute_coord_stats(K[:, coord])   // 均值, 方差
    centroids = lloyd_max(stats, num_centroids=2^bits, max_iter=20)
    K_q[:, coord] = nearest_centroid(K[:, coord], centroids)
    store(K_q, centroids)   // 4-bit index + 4-bit centroids

// Value: 统一量化 (per-group)
scale = max(|V_block|) / (2^(bits-1) - 1)
V_q = round(V_block / scale)
store(V_q, scale)

Lloyd-Max 算法

// 在已知概率密度下寻找 MSE 最优量化级
输入: 样本 x, 位数 b, 初始量化级 q_0[2^b]
输出: 最优量化级 q*[2^b], 决策边界 t*[2^b - 1]

迭代:
  while not converged:
    // 更新决策边界 (最近邻准则)
    t[k] = (q[k] + q[k+1]) / 2
    // 更新量化级 (质心条件)
    q[k] = E[x | t[k-1] < x < t[k]]
         = ∫ x·p(x) dx / ∫ p(x) dx
    检查 MSE 变化 < ε

Presets

fp8_kvcache:  // 默认
  keys: 4bit, group=16, centroid_inference=intermediate
  values: 8bit, group=256

fp8_heavy:
  keys: 4bit, group=16
  values: 8bit, group=-1 (per-tensor)

fp8_light:
  keys: 8bit, group=16
  values: 8bit, group=-1

硬件要求

• 任何 NVIDIA GPU（SM70+）
• 仅 KV Cache 算子，无需特殊 Tensor Core
• 可配合 FP8 W8A8 同时使用

9. FP8 / INT8 KV Cache

概述

对 KV Cache 在写入时量化，读取时反量化，有效减少显存占用和带宽需求。

路径：vllm/model_executor/layers/quantization/kv_cache.py

量化时机

在每个 decode step 的 attention 计算中:
  // 1. 计算 K, V (FP16/BF16)
  // 2. 在写入 kv_cache 前量化:
  kv_cache[token_pos] = quantize(K)
  // 3. 下次 decode 读取时反量化:
  K_loaded = dequantize(kv_cache[pos])

FP8 KV Cache

// Per-token, per-head, dynamic scaling
for each token t, head h:
    k_head = K[t, h, :]      // shape: [head_dim]
    absmax = max(|k_head|)
    scale[t, h] = absmax / 448  // E4M3 max
    k_q[t, h, :] = round(k_head / scale[t, h])

// 存储 K_q [num_blocks, num_heads, block_size, head_dim] (FP8)
//       scale [num_blocks, num_heads, block_size] (FP32/FP16)

反量化：

k_deq[t, h, :] = k_q[t, h, :] * scale[t, h]

INT8 KV Cache

scale[t, h] = max(|k_head|) / 127
k_q[t, h, :] = round(k_head / scale[t, h])
// INT8 KV Cache 使用 sm_75+ INT8 MMA

Scale 类型选择

class KVCacheConfig:
    SCALE_TYPES: {
        "fp8": (E4M3, dynamic per-token-head),
        "int8": (INT8, dynamic per-token-head),
    }

10. 对比总结

推荐组合（vLLM 典型部署）

说明：本文档基于 vLLM 主分支源码和对应原始论文编写。具体实现细节可能因版本迭代发生变化，建议以 vllm/model_executor/layers/quantization/ 目录下最新源码为准。