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

# 大模型量化技术深度解析：W4A8KV4方案QoQ与QServe推理系统 这事儿还得从Transformer和MOE架构说起。自从这些架构出现，大模型的参数规模轻松突破上万亿，模型是越来越大了，但部署成本也跟着水涨船高。于是，模型压缩技术就成了刚需——核心目标很明确：降低成本，同时把推理性能拉上去。 目前主流的模型压缩技术路线主要有三条：**剪枝（Pruning）**、**知识蒸馏（Knowledge Distillation）** 和**量化（Quantization）**。本系列会逐一深入讲解常见的量化方案，包括GPTQ、LLM.int8()、SmoothQuant、AWQ等，而这篇文章，我们重点聊一聊来自MIT HAN Lab的W4A8KV4量化方案——**QoQ**，以及它背后的推理服务系统**QServe**。 ## 背景 先看一个行业现状。目前主流的整数量化大致可以分成三类： - **W8A8**：8位权重 + 8位激活 - **W4A16**：4位权重 + 16位激活 - **W4A4**：4位权重 + 4位激活 前两种方案在精度损失上基本可以忽略不计。W4A4理论上能通过INT4 Tensor Core实现更高吞吐量，但问题来了——它在实际部署中并没有展现出预期中的性能优势。举个例子，最先进的W4A4服务系统Atom，在A100上跑Llama2-7B时，反而比TensorRT-LLM中的W4A16和W8A8方案慢了20-25%。 也就是说，行业到现在还没找到一个能在所有场景下同时击败W4A16和W8A8的精度组合。 原因在哪？**W4A16**是在FP16 Tensor Core上计算的，但因为权重是INT4格式，所以GEMM Kernel里必须先做权重反量化。**W4A4**这边更麻烦——为了保住精度，必须对权重和激活做**逐组（per-group）量化**，也就是在每个子通道上都共享FP16缩放因子。像QuaRot这样的最先进方案也证实，从逐组切换到逐通道量化后，困惑度会退化0.2。逐组量化意味着在GEMM的主循环里要做INT32到FP32的反量化，而这部分操作只能在慢得多的CUDA Core上跑。在A100这类数据中心GPU上，**一个CUDA Core操作的代价相当于50个INT4 Tensor Core操作**。所以，要想让LLM服务达到最优吞吐，核心挑战就是：**把CUDA Core上的开销降下来**。 ## 动机 权重和KV缓存的量化（比如W4、KV4）可以减少内存占用，而同时量化权重和激活（如W8A8）可以提升峰值计算吞吐。但选哪个精度组合，是个难题。 那么，**为什么QoQ要选W4A8KV4？** 在LLM的推理过程中，Attention和GEMM占据了绝大部分计算量。而计算强度主要受batch_size影响——batch_size越大，计算强度就越高。 从Roofline模型分析A100 GPU的表现可以看得很清楚：当batch_size较小（比如m<78）时，GEMM是**内存受限（memory-bound）** 的，这时候W4A16因为内存占用小，吞吐量更高；当batch_size变大（m>78），就开始变成**计算受限（compute-bound）**，W8A8因为能充分利用INT8 Tensor Core，吞吐更高。那W4A8呢？这才是妙招——它在计算强度小的时候靠低内存占用取胜，在计算强度大的时候又能跑满INT8的计算单元，两边都不耽误。 至于**KV Cache为什么用4位而不是8位**？因为LLM解码阶段每个token是逐步回归迭代的，这个阶段是典型的memory-bound，KV Cache加载越快，计算吞吐越高。KV4自然能带来KV8两倍的峰值性能。 **那为什么不选更激进的W4A4？** 理论上看，当m超过78时进入compute-bound，W4A4的4位Tensor Core应该是8位Tensor Core的两倍性能。但问题出在实际硬件上——不论Ampere还是Hopper架构，这种理论优势都兑现不了。现有的W4A4系统Atom和QuaRot，连W16A16都比不过。 部分原因是运行时效率低，更根本的是，**逐组量化的W4A4 GEMM在GPU上存在天然的映射困难**。 具体来说，在FP16和W8A8的GEMM中，主循环完全是在Tensor Core上完成的。但TensorRT-LLM的W4A16和Atom的W4A4都必须在主循环里做反量化，这些反量化操作跑在CUDA Core上。W4A16要做INT4到FP16的权重转换，Atom则要做INT32到FP32的局部和转换与累加。 Atom主循环的反量化带来了两个效率瓶颈： 1. **CUDA Core太慢**：在A100和H100上，FP32 CUDA Core的峰值性能只有INT4 Tensor Core的2%。也就是说，反量化一个局部和的代价，相当于50个Tensor Core MAC操作。主循环反而被慢速的CUDA Core操作拖累了。 2. **寄存器压力过大**：Atom需要同时维护两组寄存器（FP32和INT32）来保存局部和。大GEMM问题在GPU上本身就受寄存器限制，高寄存器消耗意味着可同时运行的warp数量减少，延迟隐藏能力下降，进一步加剧了主循环开销。 QServe的做法是**选择权重反量化而不是局部和反量化**，因为权重反量化的寄存器压力更低。通过两级渐进式分组量化，让所有计算都在INT8 Tensor Core上完成，同时用4路寄存器级并行来同时解码四个INT4权重，大幅降低了主循环开销。 ## QoQ：W4A8KV4量化方案详解 要实现W4A8KV4精度在理论上的吞吐优势，同时不牺牲模型质量，QoQ算法引入了三项核心设计。 ### 渐进式分组量化 分组量化能提精度，但反量化开销会抵消精度收益。QoQ的解法是**渐进式分组量化**： 1. 先对权重张量做**逐通道（per-channel）对称INT8量化**，得到中间8位量化权重。 2. 再对中间结果做**逐组（per-group）非对称INT4量化**。 在W4A8 GEMM计算时，4位量化权重先反量化为中间8位量化权重，然后直接执行INT8矩阵乘法——就像做W8A8逐通道量化一样。 **保护量化范围**是个关键细节。简单的两级量化可能会导致反量化后的值超出8位整数范围。一个直接的解决方法是开启饱和选项，但这会导致计算吞吐量下降高达67%。 QoQ的做法更聪明：**把INT8的对称量化范围从[-127, 127]缩小到[-119, 119]**，这样就能避免反量化溢出，同时不需要牺牲性能。 与之前的两级量化方法（如QLoRA的VSQuant、DoubleQuant）不同，它们是把浮点缩放因子做逐通道量化，最终反量化后还是浮点值，没法跑在INT8 Tensor Core上。DGQ虽然对缩放因子做了限制，但它的反量化Kernel和GEMM Kernel是分开的，端到端延迟甚至比cuBLAS的W8A8还慢，4位权重的内存带宽优势根本没体现出来。 QoQ则通过引入保护范围，**把反量化操作直接融合进W4A8 GEMM Kernel**，实现了全寄存器级并行，CUDA Core开销降到最低。效果很直观：QServe的W4A8逐组GEMM，比cuBLAS的W8A8 GEMM快了1.5倍。 ### SmoothAttention：应对KV4量化的精度损失 直接把KV缓存降到4位，精度会明显下降。问题出在Key矩阵上——在每个Attention头中，Key矩阵都有固定的异常值通道，这些异常值比大多数激活值大10倍左右。对KV8量化来说还好处理，但对KV4量化，量化层级太低，是个大挑战。 受SmoothQuant启发，QoQ提出了**SmoothAttention**。思路很简单：逐通道乘以一个缩放因子s，把Key缓存中的异常通道压缩下去。QoQ不需要同时量化Query，只需要关注Key，所以缩放因子的选取很直接，在实践中s=0.5就足够了。 为了消除缩放因子带来的额外Kernel调用开销，最好把SmoothAttention的缩放因子融合到前一层的权重中。但现代LLM用旋转位置嵌入（RoPE）处理Key和Query，RoPE会把每个Attention头中的两个通道配对。为了让缩放因子在RoPE下可交换，QoQ加了一个硬约束：配对通道的缩放因子必须相等。这样，缩放因子就能无缝融合到前一层权重中。 ### 通用量化优化技巧 低比特量化的核心挑战是激活异常值。QoQ针对不同类型的线性层，用上了不同的优化手段： 1. **块输入模块旋转**：受QuaRot和Quip启发，对块输入激活乘以旋转矩阵（缩放后的哈达玛矩阵），同时逆向旋转对应权重。旋转后，每个通道的激活是所有其他通道的线性组合，异常值被有效抑制。 2. **块输出模块平滑**：对生成块输出的层（输出投影层、第二个FFN层），通过除以通道平滑因子来平滑中间激活。与SmoothQuant不同，这里的迁移强度要接近0——平滑因子主要由权重决定，而不是激活。 3. **激活感知通道重排序**：用激活分布判断通道显著性，然后重新排序通道，让相似显著性的通道进入同一个量化组。 4. **权重裁剪**：对动态范围应用裁剪比率，通过优化层输出误差或块输出均方误差来确定最佳裁剪比率。 ## QServe：把理论优势变成实际吞吐 有了QoQ量化算法，接下来要解决的就是如何实现Roofline模型预测的理论吞吐优势。QServe系统设计围绕两条原则：减少GEMM Kernel主循环开销，让融合注意力Kernel保持在内存受限区域。 ### 运行时设计 QServe里所有GEMM层都使用W4A8输入，在INT8 Tensor Core上计算，输出FP16。Attention层则在CUDA Core上用FP16计算。每个LLM块保持FP16输入和FP16输出。 激活量化方面，QKV投影和第一个FFN层的激活量化被融合进前面的layernorm，第二个FFN层则融合进前面的激活Kernel。注意力块输出投影之前单独加了一个量化节点。 KV缓存管理上，采用分页KV缓存（类似vLLM和TensorRT-LLM）来避免内存碎片化。不同之处在于，QServe因为比特精度低，需要做逐头动态KV量化来保持精度。所以在每个KV缓存页面里，量化完KV特征后，紧接着存储每个头的FP16缩放因子和零点。 ### W4A8 GEMM的核心设计 让量化GEMM达到Roofline模型预测的理论性能，关键在于减少主循环开销。QServe用了三项关键技术： **计算感知的权重重排序**：在Tensor Core GEMM中，每个线程在计算时都需要跨步（strided）布局。当存储和计算数据类型相同时，可以用ldmatrix指令解决。但在W4A8场景下，ldmatrix不适用——因为数据类型不同，加载的数据与计算所需的数据对不上。QServe的做法是把权重按照计算时的使用顺序存储，把整个GEMM问题分成多个32×32分片，每个线程使用的32个通道连续存储。权重是静态的，所以这个重排序没有任何运行时开销，不仅可以达到与ldmatrix相同的指针算术开销水平，还能保证128位/线程的高带宽内存事务。 **W4A8逐通道快速反量化**：在逐通道量化场景下，重点是把UINT4高效转换为SINT8。QoQ通过重新排列权重，充分利用寄存器级并行，三个逻辑运算就能完成解包。零点减法方面，采用了**乘法后减法（subtraction after multiplication）** 的策略，把零点减法从主循环移到尾部（epilogue）执行，在主循环里几乎消除相关开销。 **W4A8逐组快速反量化**：逐组场景下，零点定义在组基础上，不能像逐通道那样合并到尾部。核心思路是：在乘法和减法中，选乘法后减法仍然有利，因为它允许在第二层级反量化中启用寄存器级并行。通过把4位字节移位补齐到8位缩放因子的高位，可以用单个INT32 ALU操作模拟4路INT8乘法。关键是，渐进式分组量化算法保证初始乘法步骤的结果不会超出INT8范围，这样就能充分发挥RLP的性能。 ### KV4注意力Kernel优化 Attention在LLM推理中占了30-50%的时间。Roofline模型说KV4应该比KV8快2倍，但实际不是这么回事。 用TensorRT-LLM的KV8注意力Kernel做基线，直接把静态逐层8位KV缓存换成动态逐头4位，在L40S上确实有1.7倍加速，但在A100上反而比KV8基线慢1.2倍。问题出在CUDA Core上——解码阶段的注意力Kernel跑在慢速CUDA Core上，而且融合了所有算术运算和KV缓存更新后，计算强度显著增加。简单反量化一个INT4就需要5个ALU操作，包括掩码、移位、类型转换、浮点乘法和减法。 更关键的是，A100的FP32 CUDA Core的Roofline转折点只有9.8 Ops/Byte。KV反量化操作本身就已经接近饱和了，这意味着融合后的KV4注意力Kernel在A100上反而变成了计算受限。 QServe的解法分两步：首先把原始TensorRT-LLM Kernel中的所有FP32操作换成FP16，计算roof直接翻倍；其次通过位技巧，把每个元素的反量化算术强度从5个操作降到2个。再加上简化控制逻辑、预取缩放因子和零点值，最终在A100上比TensorRT-LLM的KV8 Kernel快了1.5倍。 ## 实验效果 与当前常用的后训练量化方案对比，在WikiText2上对Llama2-7B模型： - 相比W8A8 SmoothQuant和W4A16 AWQ，QoQ的困惑度最多只增加0.16 - 全面优于Atom（无论W4A4还是W4A8KV4精度） - 相比QuaRot，困惑度改进高达0.49 在五个常识任务的Zero-shot准确率上，QoQ也显著优于其他4位量化方案。在WinoGrande任务上，相比QuaRot准确率提高4.82%。对Llama-2的7B、13B和70B模型，相比FP16基线，准确率损失分别只有1.03%、0.89%和0.40%。 系统效率方面，与TensorRT-LLM（FP16、W8A8、W4A16）、Atom（W4A4）和QuaRot（W4A4）对比，在A100-80G-SXM4和L40S-48G GPU上的测试结果： - A100上：Llama1-30B实现2倍吞吐提升，Llama2系列1.2-1.4×，Mistral和Yi系列1.2×，Qwen-1.5系列2.4× - L40S上：所有评估模型提升1.47×到3.47× 值得注意的是，在L40S上，用QServe服务的34B以下的七个模型中，有五个实现了比在A100上使用TensorRT-LLM更高的吞吐量。这意味着QServe有效**将LLM服务成本降低了3倍**。 ## 结语 本文详细介绍了一种W4A8KV4量化算法QoQ，以及对应的推理系统QServe。在QoQ算法层面，渐进式分组量化实现了W4A8 GEMM中的低开销反量化；SmoothAttention有效缓解了KV4量化带来的精度损失。在QServe系统层面，计算感知的权重重排序、寄存器级并行反量化、以及KV4注意力Kernel的优化，共同将理论吞吐优势转化为了实际性能提升。 与TensorRT-LLM相比，QServe在主流模型上的吞吐提升相当可观：Llama-3-8B在A100上1.2×、在L40S上1.4×；Qwen1.5-72B在A100上2.4×、在L40S上3.5×。更关键的是，QServe在L40S上就能实现比A100上TensorRT-LLM更高的吞吐——这意味着用更便宜的硬件，也能跑出更好的性能。