大模型量化技术原理：QoQ量化及QServe推理服务系统

近年来，随着Transformer、MOE架构的提出，使得深度学习模型轻松突破上万亿规模参数，从而导致模型变得越来越大，因此，我们需要一些大模型压缩技术来降低模型部署的成本，并提升模型的推理性能。 模型压缩主要分为如下几类：

• 剪枝（Pruning）
• 知识蒸馏（Knowledge Distillation）
• 量化（Quantization）

本系列将针对一些常见大模型量化方案（GPTQ、LLM.int8()、SmoothQuant、AWQ等）进行讲述。

本系列将针对一些常见大模型量化方案（GPTQ、LLM.int8()、SmoothQuant、AWQ等）进行讲述。

• 大模型量化概述
• 量化感知训练：
• 大模型量化感知训练技术原理：LLM-QAT
• 大模型量化感知微调技术原理：QLoRA
• 训练后量化：
• 大模型量化技术原理：GPTQ、LLM.int8()
• 大模型量化技术原理：SmoothQuant
• 大模型量化技术原理：AWQ、AutoAWQ
• 大模型量化技术原理：SpQR
• 大模型量化技术原理：ZeroQuant系列
• 大模型量化技术原理：FP8
• 大模型量化技术原理：FP6
• 大模型量化技术原理：KIVI、IntactKV、KVQuant
• 大模型量化技术原理：Atom、QuaRot
• 大模型量化技术原理：QoQ量化及QServe推理服务系统
• 大模型量化技术原理：总结

之前讲述了W4A4KV4量化方案Atom和QuaRot，本文将讲述来自 MIT HAN Lab 的W4A8KV4量化方案 QoQ 及 QServe 推理服务系统。

文章较长，建议先点赞收藏，后续再慢慢观看。另外，我撰写的 大模型相关的博客及配套代码 均整理放置在Github：llm-action，有需要的朋友自取。

背景

目前，业界主要的整数量化算法可以分为三类：8位权重和8位激活（W8A8）、4位权重和16位激活（W4A16）、4位权重和4位激活（W4A4）量化。前两种方法在准确性方面几乎无损。相比之下，W4A4量化导致显著的准确性降低，尽管通过将其计算映射到高吞吐量的 INT4 Tensor Cores 上有望提供更高的吞吐量。

不幸的是，这种预期的性能提升并没有在当前的GPU上得到一致的观察。例如，最先进的W4A4服务系统 Atom 在 A100 GPU上运行Llama2-7B模型时，比TensorRT-LLM中的 W4A16 和 W8A8 性能还要低20-25%。

也就是说，社区尚未找到一种比W4A16和W8A8更优越的精度组合，用于高效的进行LLM推理服务。例如，W4A16 量化在 FP16 Tensor Cores 上执行计算，由于权重在INT4中，因此需要在 GEMM Kernel 中进行权重反量化。

另一方面，对于W4A4量化，为了保持准确性， 必须对权重和激活应用逐组（per-group）量化，在子通道基础上共享 FP16 缩放因子 。例如，最先进的W4A4量化方法 QuaRot 报告中说从逐组（per-group）量化切换到逐通道（per-channel）量化后，导致困惑度退化0.2。这种逐组（per-group）量化设计需要对局部和进行整数到浮点的反量化（因为INT4 Tensor Core 产生 INT32 局部和），它在W4A4 GEMM的顺序主循环中运行在较慢的 CUDA Core 内。在数据中心GPU上，如A100， 一个CUDA Core 操作与 50 个INT4 Tensor core 操作一样昂贵 。因此，减少CUDA Core 上的开销对于实现LLM服务的最佳吞吐量至关重要。

为了应对这一挑战，作者引入了QoQ，一种W4A8KV4量化算法。QoQ由QServe推理库实现。在QoQ算法中，引入了渐进式量化，在 W4A8 GEMM 中具有较低的反量化开销。此外，还开发了SmoothAttention 来有效减轻4位KV量化引起的准确性下降。在QServe系统中，执行计算感知的权重重排序，并利用寄存器级并行来减少反量化延迟，还使融合注意力算子保持在内存受限区域，利用KV4量化带来性能提升。

动机

权重和KV缓存量化（例如：W4、KV4）可以减少LLM服务中的内存占用。同时量化权重和激活（例如：W8A8）也可以提高峰值计算吞吐量。为LLM部署选择合适的精度是一项困难的任务。现有的解决方案如上所述可以大概分为三类：W4A16(per-group）、W8A8（per-channel权重量化+per-token激活量化）、W4A4（per-group）。

QoQ 为什么选择 W4A8KV4 精度进行量化 ？

在LLM中Attention和GEMM运算占大头，所以集中分析这2个部分。在LLM中，计算强度主要被batch_size影响，batch_size越大计算强度越大，如下图所示。

对于一个 m × n × k GEMM 问题，当n、k比m大得多时，计算强度（定义为 MACs/element）大约是m。这种情况适用于LLM解码阶段，因为m是序列数，而n、k是通道大小。

对图  3中 A100 GPU 的 Roofline 进行性能分析，当 m < 78 时，W4A16的理论吞吐量更高，而当 m > 78 时，W8A8表现更好。当输入批处理大小很小时，LLM中的GEMMs是内存受限的，内存带宽由权重流量主导。因此，W4A16的较小内存占用带来了更好的性能。然而，当m很大时，问题就变成了计算受限的。因此，W8A8由于INT8 Tensor Core具有更高吞吐量而速度更快。而作者期望 W4A8 在所有批处理大小上能结合以上两者的优势。

综上所述，在计算强度较小时，主要是memory-bound，memory带宽主要被模型权重占据；因此，占用内存更低的W4A16的吞吐量会比W8A8要高。在计算强度较大时，主要是compute-bound，这时因为能够充分利用INT8 Tensor Core，W8A8的吞吐量会更高。而W4A8可以认为兼具两者的优势，不论计算密度是高是低，它都能保持最优的计算吞吐量。

而 LLM 解码阶段Attention的计算密度很低(每个token逐步回归迭代)，这个阶段还是memory-bound，所以 KV Cache 加载越快计算吞吐就越高，因此，采用KV4可以得到KV8的两倍峰值性能。

那么为什么不选择更激进的W4A4呢 ？

当输入序列数 m 超过78时，由memory-bound变为了compute-bound，W4A4开始获得更好的理论GEMM性能，因为4位 Tensor Core 的性能是8位 Tensor Core 的两倍。然而，除了显著的准确性降低，这种理论性能提升在现有的GPU架构（Ampere和Hopper）上无法实现。如图2b所示，现有的W4A4服务系统Atom和QuaRot甚至比TensorRT-LLM中的W16A16解决方案慢得多。

虽然这种性能差距部分可以用这两个系统中的低效运行时来解释，但以前文献中忽略了将逐组（per-group）量化的 W4A4 GEMM 映射到GPU的固有困难。

最先进的系统实现 Tensor Core GEMM 如图4所示的输出固定数据流。

对于一个m × n × k GEMM 问题，每个线程块通过顺序遍历reduction维度k来计算一个 输出分片（tile）。这个顺序循环被称为主循环。主循环包含100多个迭代，并且占据了GEMM Kernel 的大部分运行时间。

在 FP16 和 W8A8 GEMM（图5a）中， 主循环完全在 Tensor cores 上执行 。

TensorRT-LLM-W4A16（图5b）和 Atom-W4A4（图5c）都 需要在主循环中进行反量化操作，这些操作在 CUDA Core 上运行 。W4A16 需要 INT4到FP16的权重转换，而 Atom-W4A4 需要INT32到FP32的局部和转换和累加。

Atom主循环中的反量化过程导致了两个显著的效率瓶颈。

• 首先，在像A100和H100这样的GPU上，FP32 CUDA Core 的峰值性能仅为 INT4 Tensor Core 的2%。也就是说，在Atom中反量化一个单独的局部和相当于 50 个 Tensor Core MACs。而主循环由慢速 CUDA Core 操作而不是快速Tensor Core 操作主导。
• 其次，Atom 创建了两组寄存器（一组用于FP32，一组用于INT32）来保存局部和。由于输出固定数据流的特性，较大的GEMM问题（例如：预填充阶段）通常在GPU上而受到寄存器限制，这导致存储局部和的寄存器消耗很高。每个warp消耗大量寄存器限制了可以在流式多处理器（SM）上同时执行的warp数量。值得注意的是，GPU 依赖于大量 in-flight warp 之间进行低成本上下文切换来隐藏延迟。因此，同时执行的warp数量减少限制了延迟隐藏的机会，进一步加剧了主循环开销。

QServe 中的 W4A8 逐组（per-group）量化 GEMM Kernel 设计如图5d。通过采用了两级渐进式分组量化方法，以确保所有计算都在INT8 Tensor Cores上执行。 选择权重反量化而不是局部和反量化，因为它的寄存器压力较低 。此外，应用4路寄存器级并行来同时解码四个INT4权重，进一步减少了主循环开销。

QoQ

为了实现 W4A8KV4 量化精度的理论吞吐量优势，同时不牺牲大语言模型的有效性。QoQ算法采用渐进式分组量化、SmoothAttention和各种通用量化优化功能。

渐进式分组量化

为了提高低比特量化的准确性，通常使用分组量化。然而，它在系统实现中的反量化开销可能会抵消这些准确性的提高。为了解决这个问题，引入了渐进式分组量化，如图6所示。

给定权重张量 ，首先，应用逐通道（per-channel）对称INT8量化：

其中， 是中间8位量化的权重张量， 是逐通道（channel-wise）量化缩放因子。

然后，进一步在中间权重张量上应用逐组（per-group）非对称INT4量化：

其中, 是无符号4位量化权重张量， 是无符号4位逐组（group-wise）量化零点， 是无符号8位逐组（group-wise）量化缩放因子。

对于W4A8 GEMM计算，4位量化权重张量 将首先根据上述方程反量化为中间8位量化权重张量 ，然后执行INT8矩阵乘法，就好像是 W8A8 逐通道（per-channel）量化一样。

a) 保护量化范围

简单地应用上述方程并不能保证中间反量化权重完全位于8位整数表示范围内。例如，经过INT8量化后，一组8位权重位于[-113, 120]。4位非对称量化将得到缩放因子 ⌈(120−(−113))/(15−0)⌉=16 和零点 ⌈0−(−113)/16⌉=7 。因此，值120被量化为 ⌈120/16+7⌉=15。它将被反量化为 (15−7)×16=128，这超出了最大8位整数127。

一个直接的解决方案是在反量化过程中的算术指令启用饱和选项。然而，简单地应用饱和将严重损害计算吞吐量，速度降低高达67%。

作者重新考虑反量化过程。

将 代入 得到：

由于 ，得到：

因此， 将INT8对称量化范围从[-127, 127]缩小到保护范围[-119, 119]，以避免去量化溢出 ，如图6顶部所示。

b) 与以前的两级量化方法比较，渐进式分组量化引入了两个层级的缩放因子： 和 。

以前的研究，如QLoRA中的VSQuant和DoubleQuant，也引入了两级缩放因子来减少组内缩放因子的内存占用。它与这里的量化流程不同，以前的方法直接使用目标精度进行组量化，然后使用组内浮点缩放因子执行逐通道（per-channel）量化，如图6底部所示：

因此，使用组内缩放因子 反量化 不能产生8位权重张量。在GPU上进行计算时，这些方法首先反量化缩放因子（scales），然后反量化权重为浮点值，这最终限制了峰值吞吐量。

DGQ也遵循VSQuant和DoubleQuant的量化方案，但对缩放因子施加限制，以确保所有计算都可以映射到INT8 Tensor Core上。然而，DGQ服务系统 将反量化Kernel与GEMM Kernel 分开 。因此，DGQ中W4A8 GEMM的端到端延迟甚至比cuBLAS中的W8A8 GEMM还要慢，未能展示4位权重量化的内存带宽优势。

相反，QoQ引入了一个保护范围，允许 将反量化操作融合到 W4A8 GEMM Kernel 中，实现全寄存器级并行 ，最小化CUDA Core开销。因此，QServe的 W4A8 逐组（per-group）GEMM 比 cuBLAS GEMM 快 1.5 倍。

SmoothAttention

如图16所示，直接将KV缓存减少到4位会显著降低LLM的准确性。图7中可视化了采样的Key和Value缓存激活的幅度分布。可以观察到：Value矩阵没有明显的异常值， 而Key矩阵在每个Attention头中都有固定的异常值通道 。

这些异常值比大多数激活值大10倍左右。虽然在以前的工作中很容易处理KV8量化，但对KV4量化来说是一个挑战，因为较低的量化层级。

受SmoothQuant的启发，作者提出了SmoothAttention，通过逐通道（per-channel）因子 缩小 Key 缓存中的异常通道：

SmoothQuant将 量化难度从激活迁移到权重 ，因此需要 通过搜索迁移强度来平衡激活和权重量化 。相反，由于 QoQ不量化Query，只需要专注于Key ，并且简单地选择 SmoothAttention 缩放因子为：

在实践中， 就足够了。

如图7所示，经过SmoothAttention处理后，Key缓存中的异常值已经大大平滑。 为了消除SmoothAttention缩放的额外Kernel调用开销，最好将缩放因子融合到前一层的权重中 。然而，现代LLM使用旋转位置嵌入（RoPE）处理Key和Query，这需要额外处理。在实践中，旋转位置嵌入将通道i与每个Attention头中的通道i + D/2配对。

因此，为了使SmoothAttention缩放在RoPE方面可交换，增加了一个硬约束，即 ，

之后，可以轻松地将SmoothAttention缩放 融合到前一层的权重中，按照 和 。

LLM量化通用优化

低比特LLM量化的一个关键挑战是每个线性层的激活异常值。作者对不同类型的线性层应用不同的优化，如下所述。

1. 块输入模块旋转

在 Transformer 块中，定义 接收块输入的组件 作为输入模块，例如：QKV投影层和第一个FFN层。如图8所示，受Quarot、Quip的启发，通过 乘以旋转矩阵来旋转块输入激活 。

为了保持线性层的数学等价性，相应地 以相反的方向旋转对应权重 。 旋转后，每个通道的激活是所有其他通道的线性组合，因此有效地抑制了异常值通道 。

此外，由于旋转是酉变换（酉变换，即酉空间V的等度量变换，复数向量空间中保持向量内积不变的线性变换。这种变换保留了向量的长度和向量之间的夹角，使其在许多数学和物理问题中变得非常有用，如在量子力学中，酉变换用来描述系统的演化，保持概率守恒；在信号处理中，酉变换用于保持信号的能量不变），可以将旋转矩阵与前一层的权重融合。这里简单地选择 缩放后的哈达玛矩阵 （哈达玛矩阵是一种方块矩阵。它的矩阵元素为1或-1。其矩阵中不同的行具备正交性质）作为旋转矩阵。

2. 块输出模块平滑

输出模块指的是 生成块输出的层 ，例如：输出投影层和第二个FFN层。如图9所示，受SmoothQuant的启发， 通过除以每个通道的平滑因子来平滑块中间激活 ，原始的SmoothQuant没有平滑块中间激活；

此外，如果这里直接用与输入模块相同的迁移强度平滑这些模块（例如：q_proj、up_proj），在Wikitext-2上Llama2-7B模型的困惑度将退化高达0.05。

在实践中，发现迁移强度 应该接近0。也就是说，平滑因子 主要由权重而不是激活决定 ，这与SmoothQuant中的观察结果非常不同。

3. 激活感知的通道重排序

AWQ和Atom都观察到，保持显著权重为FP16可以显著提高模型准确性。这些显著的权重可以通过激活分布来识别。与Atom使用的混合精度量化不同，这里提出了激活感知通道重排序，如图10所示。 使用最大(|X|)来确定通道显著性，然后重新排序通道，使得具有相似显著性的通道在同一个量化组中 。

4. 权重裁剪

权重裁剪是另一种流行的量化优化技术。它通过 和 对方程中 中的动态范围应用裁剪比率 进行裁剪。

以前的方法Quarot、GPTQ、Awq、Atom通过网格搜索裁剪比率 来最小化张量本身的量化误差（即： ）或输出均方误差（即： ）。在QServe中，最小化所有线性层（除了 q_proj 和 k_proj）的 层输出误差 ，对于 q_proj 和 k_proj，通过优化 块输出均方误差 ：

QServe

在介绍了QoQ量化算法之后，实现图3中预测的理论吞吐量优势仍然是一个挑战。因此，下面将深入探讨QServe系统设计，该设计遵循两个重要原则：

1. 减少 GEMM Kernel 中主循环的开销；
2. 使融合注意力 Kernel 保持在内存受限的范围。

QServe 系统运行时

首先介绍图11中的QServe运行时。

QServe中的 所有GEMM层都使用W4A8输入 ，在 INT8 Tensor Core 上执行计算，并生成FP16输出。所有注意力层都在CUDA Core上以FP16执行计算。因此， QServe中的每个LLM块都有FP16输入和FP16输出 。

激活量化 。为确保每个GEMM输入为INT8激活，对于QKV投影和第一个FFN层，将激活量化融合到前面的 layernorm 中；对于第二个FFN层，则融合到前面的激活 Kernel 中。此外，在注意力块的输出投影之前插入了一个单独的量化节点。

KV缓存管理 。为了避免内存碎片化，遵循vLLM和TensorRT-LLM的方法，采用分页KV缓存。与这些框架不同，它们对KV缓存执行逐层（per-tensor）静态量化（即，缩放因子离线计算），QServe由于较低的比特精度需要逐头（per-head）动态KV量化以保持准确性。因此，在每个 KV 缓存页面中量化 KV 特征之后，紧跟存储每个头的 FP16 缩放因子和零点，从而允许动态更新这些值。

此外，QServe还支持与vLLM和TensorRT-LLM相似的连续批处理。

QServe 中的 W4A8 GEMM

如前文所讨论的，主循环的开销在 量化GEMM以实现roofline模型（图3）预测的理论性能增益 方面构成了重大障碍。因此，QServe W4A8 GEMM的重点在于减少主循环开销。

具体来说，通过计算感知的权重重排序来解决指针算术操作的成本，并采用乘法后减法（subtraction after multiplicatio）计算顺序和寄存器级并行来减少反量化开销。

1. 计算感知的权重重排序：

在反量化和Tensor Core计算之前，运算对象必须从全局内存加载到L1共享内存中，每个主循环迭代期间都是如此。

如图所示，Tensor Core GEMM 本质要求在计算中为每个线程进行跨步（strided）布局。由于Tensor Core GEMM kernel要求每个线程在计算时都采用跨步布局，因此这一加载过程并不简单。

例如，线程0不是连续加载八个INT8权重，而是首先加载输入通道0-3，然后跳到输入通道16-19。也就是说，一个简单的权重加载实现将需要每个四个通道执行一次地址计算，这将导致两个效率问题。

• 首先，指针算术操作在CUDA Core上执行，其吞吐量比A100上的INT8 Tensor Core 低32倍。因此，地址计算开销变得不可忽视。
• 其次，跨步内存访问阻止了通过 打包128位加载 实现最高的HBM带宽，进一步减慢了内存流水线。

当存储和计算数据类型相同时，ldmatrix指令解决了这个问题。如图12a所示，线程i连续加载输出通道i%8的128位，ldmatrix指令自动以跨步方式分布数据，确保每个线程最终获得INT8 Tensor Core 计算所需的数据。

不幸的是，当用于存储和计算的数据类型不同时（如W4A8），ldmatrix指令将无法工作。

具体来说，在图12b中，ldmatrix确保每个线程在寄存器文件中的数据置换后获得相同数量的字节，而不是相同数量的元素。

因此，线程0获得了T0自身和线程1所需的分片（tiles），而线程1获得了线程2和线程3在随后的INT8 Tensor Core计算中所需的分片（tiles）。这在每个线程获得的数据和计算中使用的数据之间造成了不匹配。

也就是说，ldmatrix无法用于W4A8 GEMM，并且前述的指针算术开销持续存在。更糟糕的是，当我们连续加载4位权重时，内存带宽利用率进一步恶化。

因此，通过计算感知的权重重排序（图12c）解决了这个挑战。关键是以它们在计算中使用的顺序存储权重。通过将整个GEMM问题分成多个32×32分片（tiles） 。在每个分片中，线程0使用输入通道0-3和16-19用于输出通道0、8、16和24（在图12c中省略了输出通道16-31）。因此，将这32个通道连接成一个单独的128位字(word)。

线程1使用的32个通道紧跟在线程0的32个通道之后存储。 由于权重是静态的，这种重排序不会引入任何运行时开销 。

此外，它不仅将指针算术开销降低到与ldmatrix相同的水平，而且还保证了高带宽的 128位/线程 内存事务。将这种重排序应用于零点和缩放因子，以减轻反量化开销。

2. W4A8 GEMM中逐通道（per-channel）快速反量化

如图5d所示，当权重和激活使用的比特精度不同时，在主循环中反量化权重变得必要。在逐通道（per-channel）量化 W4A8 的情况下，省略了第二层级缩放因子 ，第一层级 FP16 缩放因子被有效地融合到GEMM尾部（epilogue）。因此，这里讨论的重点在于，在主循环中，将 ZINT4（即：有零点的无符号4位整数）有效地转换为SINT8。将这种转换进一步分解为两个步骤：UINT4到UINT8（权重解包）和UINT8到SINT8（零点减法）。

如图13所示，重新排列每32个UINT4权重w0、w1、...、w31为w0、w16、w1、w17、...，这使得可以利用寄存器级并行，并以仅三个逻辑运算高效地将它们解包为UINT8数字。

对于从UINT8到SINT8的转换，最直接的方法是在主循环中引入整数减法指令，将其称为 乘法前减法（subtraction before multiplication） 。虽然这种方法简单，但不可避免地给主循环带来了额外的成本，这是不可取的。相反，作者 采用乘法后减法（subtraction after multiplicatio）方法，以最小化主循环开销 。

那么，具有逐通道（per-channel）量化操作对象的GEMM层可以表示为：

其中， ( ) 是量化权重（激活）， 将大小为n（输出通道）的零点向量 扩展到k×n（k是输入通道）， , 也是从缩放向量 获得的。

将