DecDEC: A Systems Approach to Advancing Low-Bit LLM Quantization

Yeonhong Park, Jake Hyun, Hojoon Kim, and Jae W. Lee, Seoul National University

主要贡献

本文针对大型语言模型（LLM）的量化问题，特别是低比特设置（如3比特和4比特精度）下模型质量下降的挑战，提出了一种推理方案DecDEC。该方案旨在改善低比特LLM的质量，同时保留量化的关键益处：GPU内存节省和延迟减少。核心问题是：在给定内存预算下，使用最佳努力配置的量化LLM，是否存在一种方法来恢复量化引起的质量损失？研究目标是利用外部CPU内存存储残差矩阵（全精度权重与量化权重之间的差异），并动态获取仅针对少量权重的残差，这些权重对应于激活异常值标记的显著通道，从而纠正这些通道中的量化错误。显著通道在每个解码步骤通过分析输入激活动态识别，以适应激活分布的动态性质，从而最大化错误补偿的有效性。创新点包括：对LLM推理中激活异常值的动态性质进行深入分析；提出DecDEC推理方案，通过动态识别显著通道并补偿其量化错误来增强量化LLM；引入DecDEC调优器，用于推荐满足目标延迟界限的系统参数；在五个不同消费级GPU上评估DecDEC，展示显著质量改进，同时内存和延迟开销最小。

背景知识/关键Observation/设计原则

LLM推理

现代LLM架构基于Transformer【[59], Attention is all you need, 2017, NeurIPS】，由多个Transformer解码器块组成，每个块包含多个线性层，这些层占据了大部分推理时间，此外还有自注意力层和归一化层等组件。

LLM推理涉及两个阶段：预填充和解码。在预填充阶段，所有输入标记（如‘I’、‘am’、‘a’）并行处理以生成单个输出标记（如‘computer’）。随后开始解码阶段，其中前一步的输出标记反馈到模型中生成下一个标记，此过程重复直到序列结束。解码阶段的顺序性质使其成为主要延迟瓶颈。

解码阶段特别受内存限制，因为一次只处理一个标记，将线性层简化为GEMV操作。在数据中心设置中，可以通过批处理多个查询来缓解此问题【[32], Efficient memory management for large language model serving with pagedattention, 2023, SOSP】【[66], Orca: A distributed serving system for Transformer-Based generative models, 2022, OSDI】。然而，这通常无法应用于设备端推理，其中LLM仅服务于单个用户。

LLM量化

LLM量化是一种流行的压缩技术，可减少内存使用和推理延迟，可分为两类【[28], An empirical study of llama3 quantization: From llms to mllms, 2024】【[69], A survey on model compression for large language models, 2024】：权重-激活量化和仅权重量化【[10], Llm.int8(): 8-bit matrix multiplication for transformers at scale, 2024, NeurIPS】【[34], Enhancing computation efficiency in large language models through weight and activation quantization, 2023, EMNLP】【[37], LLM-QAT: Data-free quantization aware training for large language models, 2024, ACL】【[49], Omniquant: Omnidirectionally calibrated quantization for large language models, 2023】【[64], SmoothQuant: Accurate and efficient post-training quantization for large language models, 2023, ICML】【[65], Zeroquant: Efficient and affordable post-training quantization for large-scale transformers, 2022, NeurIPS】。权重-激活量化主要用于数据中心设置，其中需要最小化内存和计算成本以提高吞吐量。通过量化权重和激活，可以有效利用现代GPU上的低精度算术单元（如INT4、INT8、FP8）【[1], NVIDIA H100 Tensor Core GPU, 2024】。相反，对于设备端推理，仅权重量化是首选方法【[31], SqueezeLLM: Dense-and-sparse quantization, 2024, ICML】【[69], A survey on model compression for large language models, 2024】。在此方法中，量化权重从内存加载并在与全精度激活相乘之前即时反量化到全精度（即FP16）【[23], Lutgemm: Quantized matrix multiplication based on luts for efficient inference in large-scale generative language models, 2024, ICLR】。虽然它仅减少内存流量，但这足以加速设备端推理，其中内存是瓶颈，如第2.1节所述。

仅权重量化进一步分为两个子类别：量化感知训练（QAT）【[11], QLoRA: Efficient finetuning of quantized LLMs, 2023, NeurIPS】【[30], Memory-efficient fine-tuning of compressed large language models via sub-4-bit integer quantization, 2023, NeurIPS】【[37], LLM-QAT: Data-free quantization aware training for large language models, 2024, ACL】和后训练量化（PTQ）【[9], QuIP: 2-bit quantization of large language models with guarantees, 2023, NeurIPS】【[12], SpQR: A sparse-quantized representation for near-lossless LLM weight compression, 2024, ICLR】【[17], Extreme compression of large language models via additive quantization, 2024, ICML】【[19], OPTQ: Accurate quantization for generative pre-trained transformers, 2023, ICLR】【[31], SqueezeLLM: Dense-and-sparse quantization, 2024, ICML】【[33], Owq: Outlier-aware weight quantization for efficient fine-tuning and inference of large language models, 2024, AAAI】【[36], Awq: Activation-aware weight quantization for on-device llm compression and acceleration, 2024, MLSys】【[57], QuIP#: Even better LLM quantization with hadamard incoherence and lattice codebooks, 2024, ICML】。虽然QAT通过重新训练减少量化错误从而产生更好结果，但其成本使其对许多终端用户不切实际【[31], SqueezeLLM: Dense-and-sparse quantization, 2024, ICML】【[69], A survey on model compression for large language models, 2024】。因此，不需要重新训练的PTQ已成为设备端LLM推理的首选方法。因此，本文特别关注设备端推理的仅权重PTQ。

使用CPU内存增强量化LLM

概念。我们旨在利用CPU内存来改善量化LLM质量，而不增加GPU内存成本。量化在模型质量和更小内存足迹之间权衡。在受限设置中，从业者通过选择统一比特宽度或应用细粒度策略（如逐层【[8], Zeroq: A novel zero shot quantization framework, 2020, CVPR】【[14], Hawq-v2: Hessian aware trace-weighted quantization of neural networks, 2020, NeurIPS】【[15], Hawq: Hessian aware quantization of neural networks with mixed-precision, 2019, ICCV】或逐通道【[58], ExLlamaV2】【[61], Rdo-q: Extremely fine-grained channel-wise quantization via rate-distortion optimization, 2022, ECCV】分配）在固定GPU内存预算内优化此权衡。我们的目标是通过利用CPU内存事后进一步改善质量，而不增加GPU内存使用。

基本机制。图2说明了利用CPU内存增强量化LLM的概念。我们主要针对桌面或笔记本平台，其中GPU通过PCIe互连连接到CPU。与传统推理系统一样，量化权重参数（W）和激活（x）保存在GPU内存中。不同之处在于，R（原始全精度权重与量化权重之间的残差）存储在CPU内存中。在解码阶段，从CPU获取残差以帮助补偿量化错误，从而潜在改善模型质量。由于PCIe的有限带宽通常比GPU内存带宽低一个数量级（例如32 GB/s vs. 1 TB/s），获取整个残差矩阵将导致禁止性的延迟瓶颈。因此，只应以选择性方式获取一小部分残差。简而言之，此过程将量化LLM的每个线性操作从Wx增强到(W + R ⊙ M)x，其中M是一个二进制掩码，用于稀疏化R。

关键研究问题。设计有效CPU增强量化LLM推理系统的关键研究问题是确定残差子集或掩码M。一个好的掩码M应：1）在带宽约束内选择对改善模型质量贡献最大的残差部分，2）保持结构化形式，以最小化索引开销并确保有效残差传输，同时启用GPU上的高效处理。

机会：并非所有残差都同等重要。一些残差比其他更重要——这是一个可利用的机会来确定M。此机会源于LLM推理中激活异常值（即具有大幅值的激活值）的存在，这是一个众所周知的现象【[6], Quantizable transformers: Removing outliers by helping attention heads do nothing, 2023, NeurIPS】【[10], Llm.int8(): 8-bit matrix multiplication for transformers at scale, 2024, NeurIPS】【[25], Rethinking channel dimensions to isolate outliers for low-bit weight quantization of large language models, 2024, ICLR】【[27], Slim-llm: Salience-driven mixed-precision quantization for large language models, 2024】【[33], Owq: Outlier-aware weight quantization for efficient fine-tuning and inference of large language models, 2024, AAAI】【[36], Awq: Activation-aware weight quantization for on-device llm compression and acceleration, 2024, MLSys】。当某些激活值明显较大时，即使相应权重通道中的小量化错误也会被乘以并放大，导致输出中相当大的扰动。图3说明了此问题。在此示例中，权重矩阵的第三个输入通道（第三行）乘以激活异常值-3.0。我们将此类通道称为显著通道。根据输入激活的幅度在输入通道粒度上构建M可以满足有效掩码的两个关键条件：选择残差的有影响部分并保持结构化形式。

事实上，使用相应残差补偿显著通道中的错误非常有效。图4说明了量化错误（定义为使用FP16权重（Wx）和量化权重（Wx）的计算结果之间的均方误差）如何通过依次用相应FP16值替换量化权重的输入通道来减少。对于此分析，我们评估LLaMA-3-8B-Instruct模型【[16], The llama 3 herd of models, 2024】的3比特和4比特版本，使用最先进方法AWQ【[36], Awq: Activation-aware weight quantization for on-device llm compression and acceleration, 2024, MLSys】量化，并以C4数据集【[47], Exploring the limits of transfer learning with a unified text-to-text transformer, 2020, JMLR】的文本样本作为输入提示。评估包括第8、16和24解码器块中的所有四个线性层。对于3比特和4比特模型，当我们按激活幅度降序逐步补偿通道时，量化错误急剧下降（实线红色和蓝色线）。此趋势密切遵循激活幅度分布（实线黑色线），表示按降序排序的激活幅度。相反，当按随机顺序补偿输入通道时，量化错误减少显著较慢，如虚线所示。这突显了基于激活幅度优先考虑显著通道的重要性。

图4: 在Llama-3-8B-Instruct的第8、16和24解码器块中，按排序顺序（实线）和随机顺序（虚线）依次用FP16值替换量化权重的输入通道时观察到的量化错误减少趋势。还显示了按排序顺序的激活幅度分布（黑色线）

挑战：激活异常值的动态性质。

图5: (a) 激活异常值（前5%）分布和(b) 基于静态分析的异常值识别对真实前1%和5%异常值的召回率，跨越100个解码步骤。使用Llama-3-8B-Instruct模型的第8、16和24解码器块的下投影层进行剖析

识别显著通道具有挑战性，因为激活异常值的分布本质上动态变化。虽然可以通过在小校准集上静态分析激活值统计来推断显著通道【[10], Llm.int8(): 8-bit matrix multiplication for transformers at scale, 2024, NeurIPS】【[27], Slim-llm: Salience-driven mixed-precision quantization for large language models, 2024】【[33], Owq: Outlier-aware weight quantization for efficient fine-tuning and inference of large language models, 2024, AAAI】【[36], Awq: Activation-aware weight quantization for on-device llm compression and acceleration, 2024, MLSys】，但此类静态方法次优。图5(a)显示了Llama-3-8B-Instruct模型第8、16和24解码器块的下投影层中激活异常值（定义为前5%幅度的激活）的分布，跨越100个解码步骤，使用C4数据集【[47], Exploring the limits of transfer learning with a unified text-to-text transformer, 2020, JMLR】的文本样本作为输入提示。为可见性，仅显示前512个通道。虽然一些通道（例如第24块中的通道306，由箭头突出）一致显示高激活幅度并保持持久异常值，但异常值分布通常在解码步骤中显示显著不规则性。

为了量化激活异常值的动态性质，我们计算使用校准集静态分析识别的前1%和前5%异常值的召回率，与每个解码步骤中观察到的真实前1%和前5%异常值（ground truth）相比。我们使用Pile数据集【[21], The pile: An 800gb dataset of diverse text for language modeling, 2020】的子集进行校准，遵循先前工作【[36], Awq: Activation-aware weight quantization for on-device llm compression and acceleration, 2024, MLSys】。具体而言，我们剖析每个激活值的均方平均值，并使用此作为识别异常值的指标。图5(b)呈现结果，显示召回率保持较低（约20%）对于前1%和前5%异常值。这突显了静态分析的明确局限性，因为它在运行时遗漏了大多数异常值，强调了动态识别显著通道的必要性。

方法细节

概述

在本节中，基于第3.2节概述的机会并解决第3.3节描述的挑战，我们提出DecDEC，一种用于量化LLM的CPU增强推理系统，进行带有动态错误补偿的解码。图6呈现DecDEC概述。在解码阶段，DecDEC通过动态错误补偿增强每个线性层，本质上是GEMV操作。为了产生最终输出o，DecDEC将错误补偿项odec添加到基GEMV结果ob = Wx。odec通过将输入向量与从CPU选择性获取的权重残差子集相乘来计算。此选择动态进行，完全考虑输入向量的可变性。为了在PCIe带宽约束下最大化获取的残差值数量，DecDEC存储和检索残差的量化版本（R），包括量化值Qr(R)和相关元数据，而不是全精度残差（R）。这里，Qr是一个量化器，将全精度残差映射到低比特形式，并不同于用于权重的基量化器Qb。

动态错误补偿过程包括四个顺序步骤。1 首先，通过调查输入激活向量，DecDEC创建sc_indices，一个显著通道索引列表。要补偿的显著通道数量k是一个预配置参数。此步骤本质上是TopK操作，选择输入激活向量中幅度最大的值。2 接下来，从CPU通过PCIe获取对应于显著通道的量化残差部分Qr(R)[sc_indices, :]（连同必要的量化元数据）。3 然后，将获取的残差乘以稀疏化激活向量(x[sc_indices])，产生odec。4 最后，将结果odec添加到基GEMV结果ob，产生最终输出o。所有步骤与基GEMV并行运行在不同的GPU流上，并且必须高度高效，以保持隐藏在基GEMV运行时内。

以下各节提供DecDEC每个组件的细节。第4.2节解释DecDEC如何进行残差量化。第4.3节详细说明动态错误补偿的GPU实现。第4.4节说明DecDEC如何配置系统参数，包括k，要补偿的通道数量。

残差量化

图7描绘了残差的量化方案。DecDEC对残差的每个输出通道（即列）采用4比特量化。为了最小化元数据，使用对称均匀量化。此方法仅需要每个输出通道的一个标量尺度因子作为元数据。第i个输出通道的残差量化器(Qr,i)定义为：

Si，尺度因子，通过网格搜索确定为最小化原始和量化权重之间均方误差的值。

使用此量化器，每个浮点残差值r投影到-7和7之间的整数。在运行时，从CPU获取选定输入通道的量化残差（图7中突出显示）以及所有尺度因子。量化残差的每个输入通道以及尺度都连续存储在CPU内存中，实现合并数据传输。

动态错误补偿的高效实现

实现动态错误补偿的首要优先级是确保低延迟，允许其执行保持隐藏在基GEMV执行时间内。为了在此短时间窗口内获取足够数量的残差通道，DecDEC引入三种关键软件优化策略：1) 零拷贝残差获取，2) 用于通道选择的快速近似Top-K，3) 内核融合。

零拷贝残差获取。DecDEC利用CUDA零拷贝【[42], CUDA C++ Best Practices Guide, 2024】，而不是常用API函数如cudaMemcpy()或cudaMemcpyAsync()，来从CPU获取残差。这些API依赖直接内存访问（DMA）引擎进行数据传输，对于大数据块传输高效，但对于较小传输次优，由于DMA设置开销。然而，获取残差属于后者。残差获取的粒度发生在量化残差矩阵的行级。使用4比特量化和典型行长几千到几万，每个数据块传输仅几KB。为了最佳PCIe带宽利用，数据块大小理想上至少几百KB（例如256 KB）【[41], Large graph convolutional network training with gpu-oriented data communication architecture, 2021, VLDB】【[46], Evaluating characteristics of cuda communication primitives on high-bandwidth interconnects, 2019, ICPE】。在零拷贝访问中，GPU直接发送缓存线大小的内存请求，使其适合细粒度数据访问。虽然零拷贝访问有占用GPU核心生成内存请求的缺点——潜在减慢其他并发运行内核——但这对DecDEC不是主要问题。我们案例中的并发内核，基GEMV，通常受内存限制，因此为该内核使用更少核心不太可能对其执行时间有显著影响。

用于通道选择的快速近似Top-K。对于通道选择，DecDEC不是使用精确Top-K操作，而是采用快速且GPU友好的近似Top-K方法，同时保持精度。图8使用从4096维激活向量（din = 4096, k = 128）选择128个元素的示例说明该方法。如图8(a)所示，DecDEC不是单一全局选择，而是将输入分区成四个连续的1024维块，并在每个块内执行局部Top-kchunk选择。在此示例中，kchunk = k/4 = 32。然后将局部选择的元素连接以形成最终结果。虽然这引入近似，但通过避免全局同步显著减少延迟——每个局部选择由线程块独立处理。更大的块大小可以改善效率但可能增加近似错误；我们将块大小设置为1024以有效平衡此权衡。

对于每个局部选择，DecDEC使用基于桶的Top-K算法的变体【[2], Fast k-selection algorithms for graphics processing units, 2012, ACM JEA】。图8(b)说明此过程。1 首先，根据幅度将块中的1024个元素分散到桶中，桶边界为(bk0, bk1, ..., bk30)。桶数量设置为32，与warp中的线程数匹配，允许高效线程级并行，以便每个线程处理不同桶。2 接下来，从桶0开始收集元素，直到总数达到kchunk。3 如果当前桶中的元素数量超过kchunk的剩余位置，如图8(b)中的桶9，使用随机选择填充剩余位置。此随机选择添加额外近似层，但通过避免精确排序显著减少延迟。

确定适当桶边界对于最小化随机选择引入的近似错误至关重要（图8(b)中的3）。DecDEC使用小校准集剖析激活值分布，并旨在设置平衡Top-K准确性和处理更广值范围的能力的边界值。在预期第k个最大值周围放置更细粒度桶通常改善准确性，但限制系统处理分布外值的能力。为了解决此问题，DecDEC使用离线分析确定两个关键边界，bk0和bk15，从中推断其他边界（如图9所示）。让校准集的激活值分布为X ∈ RN×din，其中N是集的大小。边界bk设置为abs(X)中所有向量中第k个最大值的最大值。0和bk15之间的间隔均匀分为16个桶(bk15, bk16, . . . , bk30)，专注于第k个最大值最可能发生的范围。为了处理可能导致选择精度显著下降的分布外情况，DecDEC为超出bk15的值分配额外16个桶。具体而言，bk0设置为所有abs(X)中的最大值，bk0和bk15之间的范围也均匀分为剩余16个桶。

内核融合。DecDEC广泛地将所有动态错误补偿操作融合到一个单一内核中。图10可视化融合内核的执行流程，使用处理大小为4096×6144权重矩阵的两个线程块示例。在此示例中，每个四个块选择一个通道（即kchunk = 1, k = 4）。最初，每个线程块顺序处理两个块，选择总共两个通道（图6(b)中的步骤1）。选定通道的索引(sc_indices)以及相应激活值(x[sc_indices])存储在GPU内存中。然后使用协作组的网格级同步特征同步线程块【[43], CUDA C++ Programming Guide, 2024】。此同步是必需的，因为每个线程块处理所有选定通道的段——不是不相交子集——当从CPU获取量化残差并执行GEMV时（图6(b)中的步骤2和3）。例如，线程块0处理Qr(R)[sc_indices][: 3072] 而不是Qr(R)[sc_indices[: 2]][:]。线程块级同步允许所有线程块访问完整的sc_indices和x[sc_indices]。此分区方案允许残差GEMV的有效减少，而不需要广泛全局同步。残差GEMV的结果使用原子原语直接添加到基GEMV的结果(ob)，产生最终输出(o)（图6(b)中的步骤4）。GPU内存开销。融合内核中sc_indices和x[sc_indices]的缓冲区是DecDEC唯一的额外GPU内存使用。从CPU的零拷贝不消耗GPU内存。桶边界值不存储在GPU上；相反，仅bk0和bk15作为参数传递给内核。如果根据最大k调整大小，则单个缓冲区可重用于所有线性层。在Llama-3-8B中跨越所有层获取10%通道的极端情况下，最大k将是下投影层的1433。这需要8.6 KB缓冲区(1433×(4+2))，假设3比特精度——本质上是可忽略的开销，小于模型大小的0.0003%。

参数调优器

参数调优器的必要性。DecDEC的有效使用需要仔细调优两个关键参数：

ntb: 指定用于动态错误补偿的线程块数量。由于DecDEC与基GEMV并行运行，为动态错误补偿分配太多线程块可能减慢基GEMV计算，而分配太少可能未充分利用PCIe带宽，因为零拷贝传输需要GPU核心发出内存请求。
kchunk: 指定每个块（1024通道）补偿的通道数量。更大的kchunk改善模型质量但可能增加推理延迟。

选择ntb和kchunk具有挑战性，由于大设计空间。每个线性层类型的ntb值(nqkvtb , notb, ngutb , ndtb)有多个可行候选，受相应权重矩阵维度约束。例如，在Llama-3-8B中，nqkvtb有9个可能候选(1, 2, 3, 4, 5, 6, 8, 12, 24)，其他层也有多个选项。选择kchunk呈现更大挑战，由于其更广的可能性范围。每个线性层类型的kchunk值(kqkvchunk, kochunk, kguchunk, kdchunk)可以是任何在内存中的值。虽然依赖于平台，但此上限较大（例如，使用48 KB每块共享内存的367），导致广阔搜索空间。本节末尾呈现给定模型和平台识别kchunk和ntb候选值的的技术细节。

为了解决此问题，我们提供DecDEC调优器，根据目标减慢率建议ntb和kchunk值。调优器最大化kchunk，同时保持总执行时间——包括所有线性层的基GEMV和动态错误补偿——在相对于基线（即无补偿）的指定减慢内。此调优是给定模型-设备对的一次性过程。

参数调优过程。图11(a)说明调优过程，包括两个阶段：阶段1确定ntb值，阶段2确定kchunk值。

在阶段1中，调优器通过将其替换为搜索单个元参数nmaxtb来简化每个层的ntb搜索。nmaxtb表示用于动态错误补偿的线程块数量上限。然后每个层的ntb设置为低于nmaxtb的最大候选。确定nmaxtb涉及测试值直到总SM计数的一半，以减少搜索空间。在图11所示示例中，GPU有56个SM，因此测试直到28。对于每个测试nmaxtb，粗粒度kchunk搜索检查所有层中kchunk的均匀增量可以应用多少，而不超出目标减慢率。图11(b)显示nmaxtb = 24（即(nqkvtb , notb, ngutb , ndtb) = (24, 16, 23, 14)）的粗粒度kchunk搜索示例，产生19个有效步骤。如果对于任何nmaxtb值无法进行步骤，调优器将具有最小权重矩阵的层的kchunk固定为0并重复过程，因为较小矩阵通常对kchunk增加最敏感。

阶段1后，调优器选择具有最多步骤的nmaxtb并在阶段2中进行细粒度kchunk搜索。图11(c)显示选定nmaxtb = 24的细粒度k搜索。在此阶段，不是所有kchunk值一起增加。在每个步骤，调优器尽可能为更多层递增kchunk，优先考虑执行时间增加较小的层。例如，在步骤1，首先递增kguchunk，其次是kdchunk。步骤1在此点停止，因为进一步增加kqkvchunk和kochunk将超出目标；因此设置它们的最终值（即kqkvchunk, kochunk = 19）。此过程重复直到无法为任何层进一步递增。技术细节：ntb和kchunk候选。参数调优期间考虑的ntb值是那些对内核执行的两个部分之一有意义影响的值：近似Top-K选择和残差获取。在近似Top-K选择部分中，每个线程块的最小处理粒度是一个块。因此，增加ntb超出块数量不会对性能有额外影响。因此，与此部分相关的ntb值集是：

对于残差获取部分，4比特残差以256值（128字节）的合并段通过PCIe传输，导致总共s = dout /256个段。这些段分布在ntb线程块中，每个块处理⌈s/ntb⌉个段。如果多个ntb值导致每个块相同段数（⌈s/ntb⌉），则仅考虑最小此类值，其余是冗余的。排除这些情况，与此部分相关的候选集是：

ntb的最终候选集是两者的并集（即N = A ∪ B）。

同时，kchunk受共享内存限制约束。在内核的近似Top-K选择部分中，共享内存使用随kchunk值增加。具体而言，此部分的共享内存使用是：

这里，128字节用于跟踪落入32个桶中每个的元素数量的整数计数器；128 × kchunk字节用于临时存储分配到每个桶的元素索引；2 × 1024字节考虑块中的输入激活值。总数必须保持低于每块共享内存限制（例如49,152字节），这约束了最大允许kchunk值。

实验环境

数据集包括：WikiText【[39], Pointer sentinel mixture models, 2016】用于困惑度评估，规模为测试集；BIG-Bench Hard (BBH)【[54], Challenging big-bench tasks and whether chain-of-thought can solve them, 2022】包含23个任务，用于问题解决能力评估；MT-Bench【[68], Judging llm-as-a-judge with mt-bench and chatbot arena, 2024, NeurIPS】包含80个响应，用于多轮对话评估，使用GPT-4o作为评判器。校准数据集包括C4【[47], Exploring the limits of transfer learning with a unified text-to-text transformer, 2020, JMLR】和Pile【[21], The pile: An 800gb dataset of diverse text for language modeling, 2020】子集，用于静态分析和桶边界确定。

模型架构：Llama-3-8B-Instruct【[16], The llama 3 herd of models, 2024】（8B参数）和Phi-3-medium-4k-instruct【[40], Phi-3 technical report: A highly capable language model locally on your phone, 2024】（14B参数），均为指令调整LLM。量化方法：AWQ【[36], Awq: Activation-aware weight quantization for on-device llm compression and acceleration, 2024, MLSys】（均匀量化）和SqueezeLLM【[31], SqueezeLLM: Dense-and-sparse quantization, 2024, ICML】（非均匀量化）。比特宽度：3比特、3.5比特（块级分配，半块3比特、半块4比特，使用KL散度敏感性【[8], Zeroq: A novel zero shot quantization framework, 2020, CVPR】）和4比特。残差量化：4比特对称均匀量化。

硬件配置：消费级GPU包括RTX 4050 Mobile（6GB GDDR6, 16 SMs, 128 GB/s内存带宽, 16 GB/s PCIe带宽）、RTX 4070 Mobile、RTX 4070 Super、RTX 4080 Super、RTX 4090（24GB GDDR6X, 60 SMs, 1 TB/s内存带宽, 32 GB/s PCIe带宽）；服务器级GPU包括H100 SMX5和GH200（两者3.36 TB/s内存带宽，H100 64 GB/s PCIe，GH200 450 GB/s NVLink-C2C）；跨代GPU包括RTX 3080、4080S、5080。平台为桌面/笔记本，CPU-GPU通过PCIe连接。

软件配置：基于PyTorch的推理管道，使用torch.compile优化【[5], Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation, 2024, ASPLOS】；基GEMV内核为LUTGEMM【[23], Lutgemm: Quantized matrix multiplication based on luts for efficient inference in large-scale generative language models, 2024, ICLR】（均匀量化）和Any-Precision LLM【[45], Any-precision llm: Low-cost deployment of multiple, different-sized llms, 2024, ICML】（非均匀量化）；CUDA用于零拷贝和内核融合；NVIDIA Nsight Systems用于内核时间测量；操作系统未指定，但兼容CUDA。

实验结果

GPU内核基准：评估DecDEC GPU内核在RTX 4070 Super、RTX 4090和RTX 4050 Mobile上，针对Llama-3-8B-Instruct的输出、门/上和下投影层的GEMV操作，假设3比特宽度。基GEMV使用LUTGEMM【[23], Lutgemm: Quantized matrix multiplication based on luts for efficient inference in large-scale generative language models, 2024, ICLR】。结果显示执行时间随kchunk呈分段线性函数：小kchunk下延迟接近基GEMV时间（补偿隐藏）；超过膝点后线性增加。膝点敏感于ntb和Rbw（内存带宽与PCIe带宽比率），较低Rbw（如RTX 4050M）支持更大kchunk。更大权重尺寸允许更高kchunk而最小延迟开销。图12显示了变化ntb和kchunk下的归一化执行时间，理论膝点标记为虚线（源自图12）。

图12: 基GEMV + DecDEC的执行时间，随kchunk和ntb变化，归一化到基GEMV执行时间

模型质量影响：在3比特、3.5比特和4比特Llama-3和Phi-3上评估DecDEC，基量化AWQ和SqueezeLLM。kchunk统一设置为8、16、32、64、128。WikiText困惑度（图13）：困惑度随kchunk增加一致下降；3比特模型在kchunk=8时显著改善（AWQ Llama-3从10.15降到9.63，Phi-3从5.96到5.53）；4比特影响较小。BBH准确率（图14）：类似趋势，使用CoT【[62], Chain-of-thought prompting elicits reasoning in large language models, 2024, NeurIPS】。MT-Bench分数（图15）：当基线接近FP16时分数不变，否则小kchunk显著改善，但进一步增加不总是明显。DecDEC通道选择有效性（图16）：与随机、静态（Hessian排名【[33], Owq: Outlier-aware weight quantization for efficient fine-tuning and inference of large language models, 2024, AAAI】）和精确Top-K比较；DecDEC优于静态（召回率~80% vs. ~30%），接近精确。残差比特宽度影响（表2）：4比特在等效数据传输下最佳或接近最佳。

图13: WikiText上的困惑度。x标记对应无DecDEC基线（kchunk=0）。越低越好

图14: BBH上的准确率。x标记对应无DecDEC基线（kchunk=0）。越高越好

图15: MT-Bench分数。x标记对应无DecDEC基线（kchunk=0）。越高越好

图16: 与随机、静态和精确通道选择比较。困惑度在上行（越低越好），召回率在下行（越高越好）。黑色x标记对应无DecDEC基线（kchunk=0）

端到端评估：在RTX 4090、4080 Super、4070 Super、4070 Mobile和4050 Mobile上评估，目标减慢率2.5%、5%、10%、20%。调优器配置列于表3，实际减慢低于目标。困惑度 vs. 每标记时间（图17）：DecDEC提供Pareto最优权衡；2.5%和5%目标特别令人印象深刻；在高PCIe/GPU带宽比率平台上，3比特DecDEC有时优于3.5比特基线（例如AWQ Llama-3在4050M上从10.15降到9.12，1.7%减慢）。

图17: 各种NVIDIA GPU上的困惑度对每标记时间。x标记表示基线值，无DecDEC（kchunk=0）。后续标记显示DecDEC在目标减慢率2.5%、5%、10%和20%上的结果

表3: 四个目标减慢率（2.5%、5%、10%、20%）的调优器结果(nmaxtb /(kqkvchunk, kochunk, kguchunk, kdchunk))以及3比特Llama-3和Phi-3的相应实际减慢率。Phi-3在4050M GPU上内存不足(OOM)

跨GPU代评估：在RTX 3080、4080S和5080上评估（表4），Rbw不变或降低。结果（图18(a)）显示DecDEC改进在三代中相当，确认鲁棒性。

图18: 困惑度 vs. 每标记时间：(a) 跨不同GPU代和(b) 在服务器级GPU上

服务器级GPU适用性：在H100 SMX5和GH200上评估AWQ Llama-3-70B-Instruct。结果（图18(b)）显示困惑度改进，最小延迟开销；GH200优势小于Rbw差距建议，因为量化GEMV受L1限制而非DRAM。

结论

DecDEC是一种低比特LLM推理方案，通过选择性检索存储在CPU内存中的残差来纠正量化错误，从而改善模型质量。通过在每个解码步骤动态识别显著通道，DecDEC在有限传输量内最大化错误补偿。DecDEC以最小开销显著改善量化LLM的质量。未来工作可包括优化服务器级GPU的量化GEMV内核以缓解L1瓶颈，进一步解锁收益。