大模型量化技术原理：FP6

【点击】 加入大模型技术交流群

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

• 剪枝（Pruning）

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

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

• 大模型量化概述
• 量化感知训练：
• 大模型量化感知训练技术原理：LLM-QAT
• 大模型量化感知微调技术原理：QLoRA
• 训练后量化：
• 大模型量化技术原理：GPTQ、LLM.int8()
• 大模型量化技术原理：SmoothQuant
• 大模型量化技术原理：AWQ、AutoAWQ
• 大模型量化技术原理：SpQR
• 大模型量化技术原理：ZeroQuant系列
• 大模型量化技术原理：FP8
• 大模型量化技术原理：FP6
• 大模型量化技术原理：FP4
• 大模型量化技术原理：总结

之前对FP8进行了基本介绍，同时，讲述了FP8在不同推理框架中的应用，本文将围绕微软研究人员在FP6量化上相关的一些工作（ ZeroQuant(4+2) : Redefining LLMs Quantization with a New FP6-Centric Strategy for Diverse Generative Tasks、 FP6-LLM : Efficiently Serving Large Language Models Through FP6-Centric Algorithm-System Co-Design）讲述。

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

ZeroQuant(4+2)

背景

近年来，LLMs 的广泛应用受到了其需要很大计算和内存需求的挑战，为了缓解这些挑战，业界诞生了很多量化技术来压缩模型的大小。但模型的压缩通常会带来模型质量下降。目前，一些先进的算法设计（例如：GPTQ  和 LoRC ）使质量下降大大减少。但是是针对较大的模型大小（大于 13B），对于较小的模型大小（例如：1B）会导致模型质量下降很多。同时，它们只关注zero-shot的评估。而在生产环境中，保持原始模型的不同任务的性能至关重要，模型质量的任何损失都是一个主要问题。现有方法虽然具有创新性，但并不能完全满足在实际应用中部署LLMs的实际要求。

为了应对这些挑战，论文主要的工作包括：

• 扩大模型量化评估范围 。研究表明，现有的定量（如：GPTQ）方法往往会过度拟合校准数据集。本文扩大了 LLMs 中 4 位量化的评估范围，以包括zero-shot之外的任务，例如：代码生成和摘要。文中发现 INT4 量化通常表现不佳，尤其是在较小的模型中，甚至在参数高达 130 亿的情况下，例如：LLaMA-13b。
• 证明了 FP6 量化的卓越性能 。本文说明了 FP6 采用基本的舍入到最近 (RTN) 算法和粗粒度量化方法，始终能够达到与全精度模型相当的精度，并在广泛的生成任务中证明了其高度有效。使用 FP6 量化的 StarCoder-13B 模型在代码生成任务中的性能，与 StarCoder-13B FP16 模型相匹配。而对于 406M 等较小模型在总结任务中，与基线结果接近。这超出了 INT4 量化的能力。
• 创新的4+2 FP6设计 。本文为 FP6 引入了创新的 4+2 设计，克服了 AI 硬件上的集成和加速问题。该设计实现了与最先进的 INT4 细粒度量化类似的延迟，使 FP6 成为 LLMs 中现有 4 位量化方法的可行替代方案。

LLM量化需要进行全量的评估

本文通过Zero-Shot任务(困惑度、准确性)、生成任务（ROUGE或Pass@1）、摘要任务来评估模型的性能。全面比较了 RTN 和 GPTQ 方法的 INT4 精度的细粒度量化（FGQ）和粗粒度量化（CGQ），得出的结论是 位精度和性能之间没有明确的联系 。

同时发现虽然 GPTQ 在PTQ量化方面具有创新性，但它往往会过度拟合特定的数据集，尤其在其细粒度量化时。对于较大的模型（从 1B 到 13B 或 65B），过拟合现象不太严重。

LLMs的核心优势在于其生成序列的能力，通过评估大模型的代码生成能力发现，与标准基准测试相比， INT4 的性能存在显著差异 。比如：

• 在 JavaScript 代码中， CodeLLaMA-34B 模型的性能从 45.05 (FP16) 下降到 43.45 (INT4, CGQ) 或 43.22 (INT4, FGQ)，分别下降了 1.6 和 1.83 个点。
• 对于较小的模型和 JavaScript 代码中，虽然 INT4 上的 FGQ 比 CGQ 有相当大的改进，但与 FP16 相比仍然存在差距。
• 但在 Python 代码中，FGQ 上的 INT4 CodeLLaMA-34B 得分为 46.88，超过了其基准，而 FGQ 上的 INT4 CodeGeeX2-6B 得分仅为 29.8，甚至落后于其 INT4-CGQ 性能。这凸显了 INT4 的不一致性 。

因此，需要更深入探索 INT4 在复杂生成任务中的有效性。

补充：pass@k 评估指标

使用pass@k指标评估功能正确性，即每个问题生成k个代码样本，如果有任何样本通过单元测试，则认为问题已解决，并报告问题解决的总比例。然而，以这种方式计算pass@k会有很高的方差性。相反，为了评估pass@k，我们为每个任务生成n≥k的样本，计算通过单元测试的正确样本 的数量，并计算出无偏估计。

最佳解决方案：FP6

由于 INT4 量化在代码生成和摘要任务中的不稳定性以及之前的研究表明 FP8 在激活过程中的应用比 INT8 的使用具有显著的改进。受此启发，提高位精度（例如：提高到 5 位或 6 位）能否在生成任务中提供更稳定、更稳健的结果？本论文主要是探讨 FP6（FP5）的有效性程度及其对不同量化算法的适应性，为之前 INT4 量化挑战带来的困境提供潜在的解决方案。

之前在FP8量化一文中介绍过标准的浮点数表示格式，主要由三部分组成：符号位、指数位和尾数位。

其中，S 表示符号 (±1)，E 表示指数( ，e 是指数位数），b 是指数的偏置（通常为 ），M 表示尾数 ( )。

因此，对于 而言，可实现的最大/最小值为 。

FP16（或BF16）权重矩阵进行如下量化：

其中， 是原始全精度权重，Quant(·)表示量化函数， 是缩放因子， 是 数。比例因子 的计算公式为： ，从而确保在不影响精度的情况下，最佳地使用 的范围。

为什么不用 INT6 而不是 FP6？

选择 FP6 而不是 INT6 是由两个关键因素驱动的：

• FP 格式简化了转换过程，因为最终计算通常使用 FP16 或 BF16 执行。
• ZeroQuant-FP的研究表明这些格式之间的准确性没有观察到差异，从而无需额外的实验验证。

FP6 性能评估

通过代码生成、摘要和Zero-Shot任务上进行了全面测试。总的来说，FP6 量化方法，特别是使用 CGQ 在不同任务、模型甚至量化算法（RTN 和 GPTQ）之间得到了高性能和鲁棒性的平衡，比 FP5 和 INT4 量化都有显著改进。同时，并没有发现位精度和性能之间存在明确的联系。

关于 FP6 系统支持探讨

4+2 格式用于 FP6

为了解决使用非标准 6 位数字格式的挑战，该格式偏离了传统的 2 次幂数字格式，本文提出了一种新颖的方法，此方法与两种通常考虑的策略不同：

• 第一种方法：直接将 6 位格式转换为 8 位浮点 (FP8) 格式。虽然这是一个简单的解决方案，但它抹杀了 6 位格式的主要优点：节省内存。
• 第二种方法：将多个 6 位数字组合在一个连续的内存块中，并使用 32 位整数 (INT32) 或 32 位浮点 (FP32) 格式表示它们。该方法保持了节省内存的优势，但增加了反量化过程的复杂性。

本文侧重于将 6 位数字划分为两个不同的子数字：第一个子数字代表最初的 4 位，第二个子数字代表剩余的 2 位。本文提出的“4+2”方法可以看作是第二种标准方法的高级变体。在此基础上，将 6 位数字分为两个部分：

• 第一部分由最初的 4 位组成，处理符号位和 3 位指数等元素。
• 第二部分包含剩余的 2 位，专用于 2 位尾数。

这种划分为 4+2 位的方式 有利于这些子数的同时加载和反量化 ，生成最终的 16 位浮点 (FP16) 权重。

该方法创新地平衡了减少内存占用的需求与反量化的实用性，特别是在解决跨分段数字内存访问的方面。

偏置移位

在 GPU 上运行时将 FP6 反量化为 FP16 可能会占用大量资源，这主要是由于操作指数字段所涉及的复杂性。指数的偏置项通常由指数位确定，对于 FP6 为 3，对于 FP16 为 15。数学上，将FP6反量化为FP16（不包括符号）的过程表示为：

其中，上标 FP16/FP6 表示各自的格式。

注意：对于细粒度（子行）量化方案，缩放因子反量化可以在矩阵乘法之后、累加之前进行；对于粗粒度（按行）量化方案，缩放因子反量化可以在累加之后进行。

虽然填充（padding）可以轻松调整尾数，但由于偏置的差异，对齐指数需要更多的努力。可以通过将 INT4 数字转换回对称 INT8 格式来进行类比：如果 INT4 采用对称格式（对于尾数），则零填充就足够了。然而，在非对称格式中，单独填充是不够的，还需要额外的步骤。

为了解决这个问题，论文定制了 FP6 格式，使用非标准指数偏置为 15。此修改不会影响精度或准确度，因为：

这意味着偏置位移可以无缝集成到缩放因子中。至关重要的是，由于 小于 1，因此将其与 相乘仍然可以通过简单的指数位移以 FP16 格式准确表示，从而避免数值错误。

下图 a 概述了原始对每个 FP6 权重进行反量化的两步过程。

• 第一步涉及将 转换为 。
• 第二步需要乘以量化比例 。

此过程中要求最高的部分是重新计算 的指数，其中涉及从 中提取 ，加上 12，然后将其合并回 . 此外，对subnormal FP6 数进行反量化的过程进一步增加了运行时反量化的复杂性。

然而，利用本文的偏置移位策略，如下图 b 所示，指数调整变成了一个简单的bit-level填充过程。最初在步骤 1 中需要添加常量整数 12，现在可以推迟到步骤 2，从而消除任何运行时开销。这是可行的，因为量化比例与常数整数的乘法可以在模型量化之后和运行时之前静态地执行。此外，这种简化的方法还有效地适应了subnormal数的反量化。

系统评估

通过对比对仅权重量化 GPU kernels 的性能，FP6 kernel 通过 Bias-Shift 增强，速度比 cuBLAS 最高快 2.37 倍，平均快 1.92 倍。鉴于 LLM 推理通常受到 GPU DRAM 的限制，本文的方法通过最小化模型权重内存访问，有效地缓解了这一瓶颈。此外，FP6 kernel  在速度上优于 TensorRT-LLM 中最先进的细粒度 INT4 实现，FFN1平均快 1.06 倍，FFN2平均快 2.05 倍。另外，带有 Bias-Shift 的 FP6 kernel 比没有 Bias-Shift 的相同 FP6 kernel 平均快 1.36 倍，这 凸显了 Bias-Shift 的关键作用 。

小结

在ZeroQuant(4+2)中，探索了INT4量化技术(如GPTQ算法) 在LLMs中的表现能力。虽然这些技术可以减小模型大小和参数存储量，但由于过拟合问题, 它们在更一般的许多任务中往往表现不佳，包括代码生成和摘要等更多生成任务。因此，当前迫切需要新的方法来提高LLMs的效率和有效性。

通过对不同量化方法的探索，发现FP6精度格式在各种任务的性能和灵活性方面均表现出色。比如：使用FP6量化的模型，如StarCoder-15B，在代码生成方面达到了与FP16模型相当的结果，而较小的模型(如BART-406M)在摘要方面达到了标准FP16性能水平。但是FP6数据格式在当前AI硬件的高效支持中存在挑战。为了提高FP6在当前主流AI硬件上的执行效率，提出了一种新颖的 4+2  FP6 GPU kernel 方案，以实现与最先进的 INT4 细粒度量化相似的延迟。这一创新使FP6成为提高LLMs效率的有效途径。

FP6-LLM

背景

当今 部署 LLMs 具有挑战性，一方面，它需要大量的 GPU 内存（FP16 中的 GPT-3 需要 326 GB）才能容纳模型权重，而 A100/H100 GPU 仅具有最多 80 GB 内存。另一方面，LLM推理在token生成过程中面临严重的"内存墙"问题，其中LLM推理的速度主要受到从 GPU DRAM 读取模型权重的时间限制。它使 LLM 推理内存受到限制，无法充分利用 GPU 的计算能力。最近的研究表明， LLM 部署时，FP6 是推理成本和模型质量之间的良好权衡（推理成本低于 8 位量化， 模型质量比 4 位量化更好）。FP6 可以有效地减小 LLMs 的大小，并在不同的应用程序中保持一致的模型质量。

然而现有系统不提供对 FP6 量化的 Tensor Core 支持，并且在 LLM 推理期间难以实现性能改进。在 GPU 上支持 FP6 量化具有以下挑战：

• (1) 位宽不规则的模型权重对内存访问不友好。
• (2) 权重反量化的运行时开销较高。

为了解决这些问题，本文提出了TC-FPx，这是第一个全栈GPU KERNEL设计方案，具有统一的Tensor Core支持各种量化位宽的浮点权重。同时，将 TC-FPx kernel 集成到现有的推理系统中，为量化 LLM 推理提供新的端到端支持（FP6-LLM）。实验表明，FP6-LLM 仅使用单个 GPU 即可实现 LLaMA-70b 的推理，实现比 FP16 基线高 1.69× - 2.65× 的推理吞吐量。

补充：CUDA Cores 与 Tensor Cores

SIMT cores (CUDA Cores) 负责 GPU 中的通用处理任务，处理各种指令，包括整数运算、浮点运算、加载/存储操作等。SIMT cores 执行单个（或矢量）数据元素。

Tensor Cores 是为加速矩阵乘法而设计的专用硬件。在 A100 /H100 GPU 上，Tensor Cores 的 FLOPS 比 SIMT Cores 高 16.0 倍/ 14.8 倍。此外，Tensor Cores 以粗粒度工作，例如：使用单个 mma（矩阵乘法和累加）指令在形状为 16 × 16 和 16 × 8 的两个 FP16 矩阵之间执行矩阵乘法。

FP6 Kernel 设计的选择与挑战

设计选择

量化的主要目标是线性层的权重（即矩阵乘法），占总体 LLM 权重的99%以上。而现有对线性层的支持主要是针对位宽为2的指数的数据类型（例如：4位、8位和16位）而设计的。本文之前尚不清楚如何在现代 GPU 上有效支持 FP6。因此，本文提出了两个重要的设计选择。

启用 Tensor Cores ：本文发现在执行量化 LLMs 推理时支持 Tensor Core 很有必要。通过评估 AWQ 纯 SIMT Cores 执行在各种批量大小上的性能，以测试其可扩展性。如下图所示，随着推理批量大小的增加，不带 Tensor Core 支持的线性层 (AWQ_W4A16_SIMT) 的运行时性能变得极低。主要有两点原因：

• 一方面，对于线性层执行，传统 SIMT Core 比 Tensor Core 慢一个数量级。
• 另一方面，SIMT Core 计算能力的很大一部分将用于在运行时对模型权重进行反量化，这进一步降低了 SIMT Core 用于计算矩阵乘法的可用计算能力。

因此，本文启用 Tensor Core 来进行矩阵乘法的密集计算，同时利用多功能 SIMT Core 进行权重反量化。

统一kernel而不是使用双kernel ：WxA16量化的独特之处在于激活矩阵使用FP16，但权重矩阵以较窄的位宽存储。但是 Tensor Core 要求将权重矩阵和激活矩阵存储在同一数据类型中（例如：FP16/INT8/INT4）。简单的解决方案（即双kernel解决方案）添加一个额外的 GPU kernel，在调用普通 FP16 kernel之前将权重反量化为 FP16。然而，这样的推理速度会比没有量化的模型还要慢。如下图左边所示，将启动两个 GPU kernel用于线性层执行，反量化的 FP16 权重将再由第二个 GPU Kernel 读取之前写入 GPU DRAM，从而导致 2 倍 DRAM 访问。

将反量化和矩阵乘法过程融合到单个 GPU kernel 中将更加高效，从而消除了反量化权重的读/写（FP16 中的 W'），如下图右边所示。

设计挑战

在现代 GPU 上设计支持 FP6×FP16 矩阵乘法的统一 GPU Kernel 具有挑战性。

• 硬件不友好的内存访问 ：现代 GPU 内存系统天然不支持不规则位宽（不是 2 的指数），因为 GPU 全局/共享内存的最小访问大小是每个线程 8/32 位，并且要访问的内存地址必须对齐。并且 Tensor Core 复杂的数据布局要求使得不规则位宽更具挑战性。
• 反量化计算开销高 ：反量化计算成本昂贵，因为它需要大量复杂的bit-level操作（按位操作）。因此，如何将反量化融合到线性层计算中而不影响整体性能也很重要。

设计方法论

本论文为了解决不友好的内存访问的挑战，提出了运行前 Bit-level 预包装。为了应对反量化高计算开销的挑战，实现了运行时 SIMT 高效 GPU 计算。同时，设计了高效流水线，使 SIMT Cores、Tensor Cores和 GPU 内存层次结构以最佳性能协同工作。

下图对比了本文设计的 TC-FPx（仅 x 位权重量化线性层 Kernel）与传统的通用矩阵乘法 (GEMM) ，其中两个输入矩阵均采用 FP16。

在 TC-FPx 中，模型权重的存储位数有所减少。同时，在寄存器层级引入了额外的反量化阶段 (Dequant W)，其中，使用 SIMT cores 在每个线程内将 FP6 权重本地反量化为 FP16。

注意 ：FP16 权重不会写回共享内存，而是存储在寄存器中以供将来使用，从而消除了对共享内存不必要的往返访问。

另一个区别是TC-FPx使用细粒度的lds指令将x位权重从共享内存加载到寄存器，而不是使用粗粒度的固有ldmatrix（加载矩阵），后者具有严格的布局要求并且灵活性更小。

运行前比特层级（Bit-level）预包装

为了解决对于现代 GPU 内存层次结构对于不规则位宽的权重的内存访问不友好的问题。本文提出可以每 32 个 x 位权重的内存读取合并，从而产生 x 请求（每个 GPU 线程 4 字节字）。在这种情况下，所有内存访问都将以 32 位字的粒度对齐，而不是不规则的位宽。

然而，由于 Tensor Core 严格的数据布局要求，将权重的内存读取合并起来并非易事，因为 每个 GPU 线程所需的权重并不存储在连续的内存空间中 。

为了解决这个问题，本文通过 重新排序每个权重矩阵中的权重并提前预包装权重来优化运行时内存访问 。由于模型权重是在模型训练和量化后静态确定的，因此可以 提前对权重应用复杂的内存布局转换，从而不会引入运行时开销 。

由于只需要预包装一次权重，因此预打包权重的开销可以通过每次推理服务有效摊销，并且可以忽略不计 。

一般来说，权重预封装包括两个步骤。

第一步， 收集每个 GPU 线程所需的所有权重，并在本地组合这些权重 。鉴于每个 GPU 线程所需的权重最初并不位于每个权重矩阵内的连续位置，因此我们必须仔细选择每个 GPU 线程的权重。然后，为每个线程选择的权重在本地以相对时间顺序组合，因为它们在运行时被 Tensor Core 消费。

第二步， 将整个GPU WARP（由32个GPU线程组成）所需的所有权重组合到一个统一的线性内存空间中，权重将在运行时按顺序存储在GPU DRAM中 。为了完全消除共享内存库冲突，以锯齿状顺序组合每个线程的 32 位字。

注意 ：这里讨论的所有技术都独立于模型权重的实际位宽（始终使用 x 表示）。因此，我们的权重预封装可以自然地应用于任何位宽。

步骤一：逐线程权重收集

下图展示了T0(Thread#0)选取的权重以及组合它们的顺序。

假设 WARP 层级切片大小为 64 × 64 ，这意味着每个权重矩阵被分为 64 × 64 数据块，并以每个 WARP 的粒度加载到 GPU 的共享内存。然后，每个权重块进一步分为四个片，因为权重是从共享内存加载的，并逐片用于 Tensor Core 计算。每个切片被分为四个 16 × 16 块，因为 Tensor Core 在每条指令中处理 16 × 16 数据块。

在每个 16 × 16 块中，为 T0 选择四对 FPx 权重并组合。

如图所示，经过步骤1，我们得到了32组（即WARP大小）FPx权重。权重在每组中合并连续存储，并且每组权重将被某个GPU线程消耗。

综上所述， 每个 64 × 64 权重块最终分配给 32 个线程（一个 WARP） ，每个线程将消耗 128 个 x 位权重。

步骤二：按 WARP 进行 Bit-level 封装

在步骤二中，将不同组的所有权重组装成统一的内存空间。

在这个 Bit-level 预包装过程中，将组合权重作为连续的数据进行复制，暂时忽略每个位的含义。具体来说，x 位的 128 个条目被视为 32 位的 4x 个条目。

本文建议按照图中的锯齿状顺序组装所有组的权重 。

首先，将每个线程的第一个 32 位条目连接在一起。之后，每个线程的第二个 32 位条目被连接并附加到之前的结果中。通过重复这个过程，所有权重可以连续存储在线性内存空间中，并且对齐良好（128字节对齐）。这样，所有权重就可以以 128 字节块的粒度简单地从 DRAM 复制到共享内存，无需任何更改，轻松实现最佳 DRAM 访问。

此外，这些权重可以在运行时以最佳性能从共享内存加载。

具体来说，线程的 WARP 将为每个内存请求读取共享内存中的连续 32 位条目，完全避免了存储体冲突。

运行时高效的SIMT计算

并行反量化 ：

为了减少 FP-x 权重反量化的运行时开销，使用优化的按位 SIMT core 指令实现了 FP-x 反量化。此外，建议 并行对多个 FPx 权重进行反量化 ，通过利用每个 32 位寄存器内的 bit-level 并行，进一步将 SIMT 开销减少 4 倍。

（1）优化的按位运算：当将 FPx 转换为等效的 FP16 时，FP16 的指数应为 。为了简化这个过程，我们采用了Zeroquant(4+2)中的数学变换，用 来计算FP16的指数。为了保持正确性，将 FP16 的结果与 FP16 常数