CUDA-MODE课程笔记 第7课: Quantization Cuda vs Triton

我的课程笔记，欢迎关注：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/tree/master/cuda-mode

CUDA-MODE课程笔记 第7课: Quantization Cuda vs Triton

适配课件详细解读

作者课件可以在这里找到：https://github.com/cuda-mode/lectures。我也下载里一份放在 https://github.com/BBuf/how-to-optim-algorithm-in-cuda/tree/master/cuda-mode/ppt 这里。

PyTorch最近一年发布了一些生成式AI模型的案例研究，这些模型运行速度极快，且代码及其精简。这些模型比如GPT-FAST，SAM-FAST都应用了量化技术，Charles很大程度上是这些量化Kernel的主要开发者。因此，这节课由Charles来分享量化技术。

这张Slides介绍了演讲者的背景和最近的研究重点，内容如下：

• Pytorch Core

  • 量化 (Quantization)
  • 剪枝 (Pruning)
  • AO (Architecture Optimization) 团队

• 最近的研究重点

  • Segment-anything-fast, gpt-fast, sdxl-fast 等项目
  • TorchAO - 提供了一个 GitHub 链接 (https://github.com/pytorch-labs/ao)
  • Int8 动态量化
  • Int8 仅权重量化
  • Int4 仅权重量化
  • i8i8->i32 vs i8i8bf16->bf16
  • bf16i8->bf16
  • bf16i4->bf16
  • GPU 量化

这张Slides介绍了三种不同的量化技术：

• 动态量化流程 (Dynamic Quantization Flow):

  • 权重和激活都从浮点开始
  • 权重在预处理阶段量化
  • 激活在运行时量化
  • 使用Int8进行乘法运算
  • 使用Int32进行累加
  • 最后rescale回浮点数

• 未量化 (Not Quantized):

  • 权重和激活都保持浮点格式
  • 所有运算（乘法和累加）都使用浮点数进行

• 仅权重量化 (Weight Only Quantization):

  • 权重在预处理阶段量化
  • 随后立即反量化回浮点数
  • 激活保持浮点格式
  • 乘法和累加都使用浮点数进行
  • 最后有一个rescale步骤

总的来说，这张Slides展示了这三种技术在处理神经网络计算时的不同流程。动态量化通过在计算过程中使用整数运算来提高效率，而仅权重量化则只对权重进行压缩，在实际计算时仍使用浮点数。未量化的方法则完全使用浮点数，可能提供最高的精度但计算效率较低。

这张Slides进一步说明了动态量化（Dynamic Quantization）的概念和流程：

• 数学表达：

  • 原始公式：Y = X.W
  • 量化后的公式：Y = (Sx*Xint).(Wint * Sw)
  • 重排后的公式：Y = Sx * (Xint.Wint) * Sw 这里，Sx 和 Sw 是缩放因子，Xint 和 Wint 是量化后的整数值。

• 动态量化流程图：

  • 开始于浮点权重（Float Weight）和浮点激活值（Float Activation）
  • 权重在预处理阶段进行量化（Quantize (preprocess)）
  • 激活值在运行时进行量化（Quantize）
  • 使用 Int8 进行乘法运算（Multiplication (Int8)）
  • 使用 Int32 进行累加运算（Accumulation (Int32)）
  • 最后将结果重新缩放（Rescale (Float)）回浮点数
  • 输出浮点激活值（Float Activation）

这张Slides展示了逐张量量化（per-tensor quantization）和逐token量化 + 逐通道量化（per-token + per-channel quantization）两种动态量化方式。性能比较（以SAM模型为例，vit_h, bsz=16）：

• 无量化：运行时间 785.313 ms，峰值内存 15.279（单位未指明，可能是GB）
• 动态量化：运行时间 731.649 ms，峰值内存 18.631 另外，这里的链接是Triton的矩阵乘法教程。

结论：动态量化可以提高计算效率，在这个例子中，运行时间减少了约7%。不同的量化策略（逐张量、逐token、逐通道）可以应用于不同的张量，以优化性能和精度。虽然动态量化提高了计算速度，但它的显存占用却更多了大概是15%-20%。

这张Slides指出，显存增加的原因是要把int8的结果累加到int32类型中，因此相比于BFloat1增加了额外的显存。

这张Slides进一步详细介绍了动态量化（Dynamic Quantization）的概念、方法和性能比较：

• 动态量化的数学表达：

  • 原始公式：Y = X.W
  • 量化公式：Y = (Sx*Xint).(Wint * Sw)
  • 重排公式：Y = Sx * (Xint.Wint) * Sw
  • 进一步优化：Sx * (XWrescaled)

其中使用了不同的数据类型：- int8：用于Xint和Wint - bf16：用于Sx和Sw - int32：用于中间计算结果XWint

• 性能比较（以SAM模型为例，vit_h, bsz=16）：

  • 无量化：运行时间 785.313 ms，峰值内存 15.279 GB
  • 动态量化：运行时间 731.649 ms，峰值内存 18.631 GB
  • 动态量化with fusion：运行时间 695.115 ms，峰值内存 14.941 GB 结论：动态量化可以显著提高计算效率，运行时间减少约7%。动态量化with fusion进一步优化了性能，运行时间比无量化减少约11.5%，同时还略微降低了内存使用。

这里展示的是要在Torch Compile中实现动态量化with fusion需要做出的努力，因为Torch Compile并不愿意融合乘法操作，所以作者不得不在Torch Compile的矩阵乘法kernel后强制添加一个乘法的epilogue（实际上这是一个编译器的PASS，需要匹配到矩阵乘法+乘法才能生效）。图片比较难看代码，这里贴一下：

# This op is a special case of the int_mm op which we use based on the pattern
# _int_mm -> mul (defined in ../fx_passes/post_grad.py) in order to prevent
# realization of the int32 _int_mm output by forcing fusion with the mul op.
# This is only used when config.force_fuse_int_mm_with_mul = True
def tuned_fused_int_mm_mul(mat1, mat2, mat3, out_dtype, *, layout=None):
    out_dtype = (
        torch.promote_types(mat3.get_dtype(), torch.int32)
        if out_dtype is None
        else out_dtype
    )
    m, n, k, layout, mat1, mat2, mat3 = mm_args(
        mat1, mat2, mat3, layout=layout, out_dtype=out_dtype
    )
    choices: List[Dict[Any, Any]] = []
    for config in int8_mm_configs(m, n, k):
        mm_template.maybe_append_choice(
            choices,
            input_nodes=(mat1, mat2, mat3),
            layout=layout,
            **dict(mm_options(config, m, n, k, layout), ACC_TYPE="tl.int32"),
            suffix_args=1,
            epilogue_fn=V.ops.mul,
        )
    return autotune_select_algorithm("int_mm", choices, [mat1, mat2, mat3], layout)

然后，Triton在实现这个需求时相比于Torch Compile会很简单，一行代码即可。

这张Slides介绍了Int8权重量化（Int8 Weight Only Quantization）的概念和流程。主要内容：

• 数学表达：

  • 原始公式：Y = X.W
  • 量化公式：Y = X.(Wint * Sw)
  • 重排公式：Y = (X.Wint) * Sw

• 权重量化流程图：

  • 从浮点权重（Float Weight）开始
  • 量化（Quantize）步骤：在预处理阶段进行
  • 反量化（Dequantize）步骤：将量化后的权重转回浮点
  • 浮点激活（Float Activation）保持不变
  • 乘法运算使用浮点（Multiplication (Float)）
  • 累加使用fp32（Accumulation (fp32)）
  • 重新缩放（Rescale (Float)）
  • 最后输出浮点激活（Float Activation）

• 特点：

  • 只对权重进行量化，而不是对激活值进行量化
  • 在实际计算前，量化的权重被反量化回浮点格式
  • 所有的计算（乘法和累加）都在浮点精度下进行

• 其它：

  • 减少模型存储空间，因为权重以Int8格式存储
  • 保持了计算精度，因为实际运算仍在浮点下进行
  • 可能比全量化方法（如动态量化）具有更高的精度
  • 适用于对精度要求较高，但仍希望减少模型大小的场景
  • 可能在某些硬件上比全量化方法更容易实现和优化

这张Slides展示了Int8权重量化（Int8 Weight Only Quantization）的性能表现，无量化: 93.08 tokens/s，int8权重量化: 40.59 tokens/s，可以看到int8权重量化反而降低了处理速度，约为无量化版本的43.6%。

在图表中，对比了Batch size 1: cublas 和 int8 weight only quantized matmul。蓝线: cublas A16W16 matmul (使用16位精度的cublas矩阵乘法)。红线: A16W8 matmul (使用16位激活和8位权重的矩阵乘法)

这张Slides讲到如果按照普通的gemm triton kernel模板，上面的Int8权重量化的性能低于预期的原因是：

• 执行了比基础matmul更多的工作，展示了一段代码，显示了额外的加载和类型转换操作，这些额外操作可能导致性能下降
• 块大小被限制为大于等于16，当前配置只执行64个块，少于A 100GPU的108个多处理器，这可能导致一些多处理器未被充分利用

然后Torch Compile通过链接里的代码解决了这个问题，贴一下：

@register_decomposition([aten.mm])
@pw_cast_for_opmath
def mm(self, input2):
    # Our matrix vector multiplies only achieve peak bandwidth with coordinate descent tuning.
    # todo: Look into why and fix it (hopefully)
    if config.coordinate_descent_tuning:
        if guard_size_oblivious(self.shape[0] == 1) or guard_size_oblivious(
            input2.shape[1] == 1
        ):
            return (self.unsqueeze(2) * input2.unsqueeze(0)).sum(dim=1)
    ...
    return NotImplemented

实际上这个操作就是让GEMV用Cuda Core而不是Tensor Core来完成计算，具体做法就是把GEMV操作等价为一个element-wise乘法加一个reduce操作。这个操作通过Torch Compile生成的Triton Kernel代码如下：

这张Slides展示了一个名为 triton_() 的函数（由Torch编译器生成），该函数实现了 Int8 权重量化的GEMV操作。完整流程为：

• xnumel 和 rnumel 都设置为 4096
• X 对应 N 维度，R 对应 K 维度
• 使用 program_id(0) 和 XBLOCK 计算偏移量
• XBLOCK 始终为 1，每个 program_id 处理输出的单个值
• 加载完整的激活张量（fp32 格式）
• 对权重的一列进行循环
• 加载权重的一列的一个chunk（可能是 int8 格式）
• 将权重列转换为 fp32 格式
• 执行矩阵乘法的核心计算
• 使用广播和累加操作
• 对结果进行掩码处理和ReduceSum求和
• 加载额外的数据（可能是偏置或缩放因子）
• 执行最后的乘法和加法操作
• 将结果存储回内存

def triton_(in_ptr0, in_ptr1, in_ptr2, in_ptr3, out_ptr1, xnumel, rnumel, XBLOCK: tl.constexpr, RBLOCK: tl.constexpr):
    xnumel = 4096
    rnumel = 4096
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    _tmp6 = tl.full([XBLOCK, RBLOCK], 0, tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r1 = rindex
        tmp0 = tl.load(in_ptr0 + (r1), None, eviction_policy='evict_last').to(tl.float32)
        tmp2 = tl.load(in_ptr1 + (r1 + (4096*x0)), xmask, eviction_policy='evict_first', other=0.0)
        tmp1 = tmp0.to(tl.float32)
        tmp3 = tmp2.to(tl.float32)
        tmp4 = tmp1 * tmp3
        tmp5 = tl.broadcast_to(tmp4, [XBLOCK, RBLOCK])
        tmp7 = _tmp6 + tmp5
        _tmp6 = tl.where(xmask, tmp7, _tmp6)
    tmp6 = tl.sum(_tmp6, 1)[:, None]
    tmp9 = tl.load(in_ptr2 + (x0), xmask, eviction_policy='evict_last').to(tl.float32)
    tmp11 = tl.load(in_ptr3 + (x0), xmask, eviction_policy='evict_last').to(tl.float32)
    tmp8 = tmp6.to(tl.float32)
    tmp10 = tmp8 * tmp9
    tmp12 = tmp10 + tmp11
    tl.store(out_ptr1 + (x0), tmp12, xmask)

这张Slides主要讲述了Int8权重量化（Int8 Weight Only Quantization）的优化过程和结果。

• 性能问题解决：通过使用torch.compile可以解决之前遇到的性能问题。

• 性能对比（LLaMA-7B模型，批次大小为1）：

  • 无量化：93.08 tokens/s
  • int8权重量化：40.59 tokens/s
  • int8权重量化优化后：135.01 tokens/s
  • 这显示优化后的int8权重量化性能显著提升，超过了无量化版本。

• 微基准测试结果：

  • 图表显示了不同权重大小下的性能比较
  • cublas A16W16 matmul（蓝线）性能最佳
  • A16W8 matmul（红线）性能较差
  • A16W8 fixed matmul（黄线）性能介于两者之间

• 优化过程中的发现：

  • 尽管性能提升明显，但仍未完全匹配默认bf16的性能
  • 这主要是由于torch.compile的开销，在端到端测试中这个差距会减小
  • 在优化过程中遇到了triton的一些限制，通过避免使用tensor cores来绕过这些限制
  • 目前仍然缺乏对批次大小大于1（bsz>1）的高性能内核

• 未来工作：

  • 需要进一步优化以完全匹配或超越bf16的性能
  • 开发支持更大批次大小的高性能内核

这里bsz=1的时候是memory bound的GEMV，如果bsz>1，这个时候就是GEMM Kernel了，很可能就是compute bound了，普通的kernel优化预计很难超越cuBLAS的性能。

从Int4 Weight Only开始，Triton开始力不从心了。要点为：

• 目前PyTorch没有原生的int4/uint4数据类型（dtype）。
• 这意味着我们需要将更大尺寸的张量拆解成多个int4类型。
• 由于Triton在类型转换和乘法操作上的限制，我们在实际操作中会失去更多性能。
• 图示展示了int4数据（4位整数）如何被打包进更大的数据类型中。

"But we can see how far we can get with just triton"（但我们可以看看仅使用triton能走多远）说明了作者打算在现有Triton框架限制下探索Int4量化的潜力。右上角显示了一个int4x2的基本结构，每个元素包含两个4位整数。下方展示了四种不同的打包/解包布局，展示了如何在更大的数据结构中组织int4数据。

Slides里面的右下角的4张图有拼写错误，注意鉴别。比如最后一张图的第一列应该是ABEF才对。

这张Slides详解了Int4权重量化（Int4 Weight Only Quantization）在矩阵乘法（matmul）中的实现策略，特别是关于数据打包和解包的选择。

• 在进行矩阵乘法时，由于这是权重，我们希望在int4x2格式中连续的信息在解包后仍然保持连续。
• 因此，我们应该使用右边两种选项之一。
• 由于矩阵乘法的实现通常让单个线程处理所有的K维度，所以选择了右下角的选项。这种选择可以避免因打包方式导致线程加载不必要的数据。

Slides里面的右下角的4张图有拼写错误，注意鉴别。比如最后一张图的第一列应该是ABEF才对。

这里提供了具体的代码来展示如何打包/解包uint8和int4：

int4[2*k,n]=(uint4x2[k,n] & 0xF) - 8
int4[2*k+1,n]=(uint4x2[k,n] >> 4) - 8

解释说选择uint8是因为triton框架对int8的位移操作存在问题。这里的uint4x2量化Kernel代码在：https://github.com/pytorch/pytorch/blob/main/torch/_inductor/kernel/unpack_mixed_mm.py

这张Slides主要讨论了Int4权重量化（Int4 Weight Only Quantization）的性能表现和一些相关观察。

• 性能数据表格（LLaMA-7B, bsz=1）：

  • 无量化：93.08 tokens/s
  • int8权重量化：40.59 tokens/s
  • int8权重量化优化版：135.01 tokens/s
  • uint4x2权重量化：43.59 tokens/s
  • Int4分组量化：187.8 tokens/s

uint4x2量化的性能(Triton实现)只有无量化情况下的1/2，而不是预期的4倍快。作者提到如果现在重新实现，会参考fast int8 kernel的方法，而不是slow int8 kernel。此外，提到Jeff Johnson（PyTorch GPU后端的开发者）使用CUDA开发了一个int4 kernel并集成到了PyTorch中，速度非常快，也就是上面表格的Int4分组量化。代码：https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/native/native_functions.yaml

这个是kernel的签名，感兴趣的读者可以自行查看代码。

从这个Int4 Weight Only的cuda量化kernel实现可以看到Triton的局限性。

这张Slides讨论了Triton的一些局限性：

• 复杂操作和非标准数据类型的问题：

  • Triton在处理复杂操作和非标准数据类型时会遇到困难。
  • 具体提到了Int4（4位整数）类型。
  • 当批处理大小大于1时，int8/int4权重量化也会遇到问题。
  • L2缓存优化在这些情况下可能会受到影响。

• 配置一致性问题：

  • 在一些测试中，启发式算法存在问题。
  • 最佳配置可能无法使用或被启发式算法错误地丢弃。

这张Slides介绍了Trito的优势：

• 擅长组合"简单"操作：

  • Triton在将简单操作组合在一起方面表现出色。
  • 提到了两个具体例子：a) Fused_int_mm_mul（融合整数矩阵乘法和乘法操作） b) SAM flash attention（Segment Anything Model中使用的快速注意力机制）

• 性能接近CUDA，但使用更简单：

  • Triton能够达到CUDA速度的约75%。
  • 最重要的是，使用Triton可以达到这种性能水平，而无需直接处理.cu文件（CUDA源代码文件）。

• 代码：

  • https://github.com/facebookre...
  • https://github.com/pytorch-la...

这里讲的就是SAM里面的Attention操作相比于标准的SelfAttention需要融合两个MASK，这个时候使用Triton实现的FlashAttention就可以非常快的实现这个需求，并且性能很好。

要复现作者的实验或者学习GPU上量化Kernel的实现可以点击这张Slides里的链接。

作者分享的Slides里面还有一些有趣的内容作为附录，这里挑选其中的一些来解读，主要是实验结果和概念部分，对torchao的使用部分的Slides有需要的读者可以自行查看。

这张Slides展示了对SAM（Segment Anything Model）模型进行不同量化和优化技术的实验结果。主要内容如下：

• 动态量化（Dynamic Quant）相比基准模型获得了约13%的速度提升。
• 仅权重量化（Weight Only Quant）对性能提升不明显，原因是模型主要受计算限制，且其内核设计并不针对大Batch进行优化。
• 所有的量化技术都只导致了很小的精度损失。图表详细展示了不同方法的性能对比：
• fp16（半精度浮点）
• compiled（编译优化）
• SDPA
• int8 weight only quant（8位整数仅权重量化）
• int8 dynamic quant（8位整数动态量化，包括权重和激活）
• 2:4 pruned cusparselt（一种稀疏化技术） 表格中比较了这些方法在以下几个方面的表现：
• 批处理大小为32的处理时间（bs 32(s)）
• 每秒处理的图像数（img/sec）
• 相对于SDPA的加速比（speedup over SDPA）
• 峰值内存使用（peak memory (GB)）
• COCO 2017验证集上的准确率（coco 2017 val accuracy）

这张Slides展示了对Llama2 7B模型进行不同量化方法的实验结果。主要内容如下：

• 使用仅权重int8和int4量化分别实现了45%和86%的加速。
• 对于int8仅权重量化，没有观察到精度下降。
• int4仅权重量化导致了小幅度的精度下降，但使用GPTQ（一种量化技术）可以恢复其中一半的精度损失。
• 动态量化（Dynamic Quantization）虽然测试过，但因为模型受内存限制，其精度和性能都不如仅权重量化，所以未列入表格。

• 表格详细展示了不同方法的性能和精度对比：

  • bf16
  • compiled（Torch 编译优化版本）
  • int8 weight only quant（8位整数仅权重量化）
  • int4g128 weight only groupwise quant（4位整数分组仅权重量化）
  • GPTQ（使用每个任务100个样本）

• 性能指标包括：

  • 每秒处理的token数（bs 1 (tok/s)）
  • 相对于compiled版本的加速比
  • hellaswag_acc_norm
  • wikitext bits_per_byte（困惑度相关指标）
  • winogrande acc
• 结果显示，int4量化提供了最大的速度提升（1.86倍），但有轻微的精度损失。int8量化在保持精度的同时也提供了显著的速度提升（1.45倍）。

这张Slides展示了对Llama2 7B模型进行模拟低精度量化的实验结果。主要内容如下：

• 实验目的：了解分组大小（groupsize）、位数（bit number）和GPTQ（量化技术）如何影响模型准确性。实验使用wikitext bits_per_byte困惑度作为评估指标。
• GPTQ效果：在大多数情况下，GPTQ能够恢复约一半的性能损失（PPL，困惑度）。特例：在G=64、2位量化的情况下，未使用GPTQ时的PPL异常地高。
• 分组大小影响：测试了不同的分组大小（G=128, 64, 32）。G=32时，4位量化的性能损失不到3%，并且GPTQ将损失进一步减少了近50%。
• 量化位数影响：比较了4位、3位和2位量化的效果。位数越低，通常性能损失越大，但GPTQ能在一定程度上缓解这种损失。
• 性能数据：表格展示了不同配置下的推理速度（tok/s）和wikitext bits_per_byte值。较低的bits_per_byte值表示更好的性能（注1）。
• 模拟低精度：3位和2位的数据是通过在4位内核上模拟得到的，方法是在量化过程中限制Qmax（注2）。
• 基准比较：结果应与bf16（bfloat16）和int8量化的基准值0.674进行比较（注1）。

这张Slides介绍了一些与量化相关的代码资源和工具：

• 量化API：

  • 量化API可在torchao仓库中找到
  • 链接：https://github.com/pytorch-labs/ao

• segment-anything-fast仓库：

  • 这个仓库展示了如何将这些API与其他技术结合使用
  • 链接：https://github.com/pytorch-labs/segment-anything-fast

• gpt-fast仓库：

  • 使用相关API来执行量化
  • 包含了int4量化和GPTQ的实现，这些实现目前在其他地方还没有
  • 链接：https://github.com/pytorch-labs/gpt-fast

后面还有几页Slides介绍了不同的量化方法。

这张Slides介绍了动态量化（Dynamic Quantization）的过程和特点：

• 动态量化流程 (Dynamic Quantization Flow):

  • 权重和激活都从浮点开始
  • 权重在预处理阶段量化
  • 激活在运行时量化
  • 使用Int8进行乘法运算
  • 使用Int32进行累加
  • 最后rescale回浮点数

• 动态量化的特点：

  • 对非平稳分布不敏感
  • 对频繁出现的异常值敏感
  • 为每个样本重新计算量化参数

• 浮点激活：

  • 用于替代非量化操作
  • 可能比允许一系列量化操作而不需要反量化的技术要慢

这张Slides介绍了Smoothquant（平滑量化）技术，并将其与动态量化进行了对比：

• 对比两种量化流程：

  • a. 权重先进行放大（Scale Up）
  • b. 激活先进行缩小（Scale Down）
  • c. 然后再进行量化
  • 左侧是动态量化流程（与前一张Slides相同）
  • 右侧是Smoothquant流程，主要区别在于：
• Smoothquant的特点：使用输入权重均衡化（Input-weight equalization）技术

• Smoothquant与LLM.int8()的结合：

  • 对激活使用逐token量化（per-token quant activations）
  • 对权重使用逐通道量化（per-channel quant weights）

• Smoothquant的优势：

  • 通过预先的缩放操作，可以更好地平衡权重和激活的数值范围
  • 有助于减少量化过程中的信息损失
  • 结合LLM.int8()技术，可以在保持精度的同时提高效率

这张Slides展示了不同量化方法在OPT-175B、BLOOM-176B和GLM-130B*模型上的性能，Smoothquant（O1、O2、O3）在大多数情况下表现接近或优于FP16和LLM.int8()。

这张Slides介绍了仅权重量化（Weight Only Quantization）为Int8的技术。

• 量化流程：

  • 浮点权重先进行量化（Quantize）
  • 然后进行反量化（DeQuantize）
  • 使用16位权重和16位激活进行乘法运算（Multiplication W16A16）
  • 最后得到浮点激活输出
• 优点：比包含激活量化的方法更精确；原因：在实践中，激活通常是更难量化的部分 计算特性：
• 混合数据类型的矩阵乘法（Mixed dtype matmul）在计算上比fp16-fp16矩阵乘法更昂贵 但在内存使用上更高效
• 实际应用：在实践中，对int8使用逐通道量化（per-channel quantization）

这张Slides介绍了仅权重量化为Int4的技术。

• 量化流程与Int8量化类似，包括权重量化、反量化和16位乘法运算。
• 实际应用中的量化策略：对int4使用分组量化（group-wise quantization）。原因是int4的精度较低，分组量化可以提高精度。
• 分组量化的具体操作：除了对每个通道进行量化外，还将通道Ci分成n个组（G0到Gn-1）。每个组有自己的量化参数。
• 右边还展示了per-token量化的激活矩阵和per-channel分组量化的权重矩阵。
• 可以考虑使用smoothquant风格的输入-权重均衡化技术。这对int4的情况可能特别有用，因为int4精度较低，更需要优化。

这张Slides介绍了GPTQ (Generative Pre-trained Transformer Quantization) 技术。

• GPTQ的核心思想：使用期望Hessian（Expected Hessian）来量化权重W，目标是最小化 argmin ||WX - ŴX||²₂，其中Ŵ是量化后的权重。
• 量化方案：可以是分组量化（group-wise）或逐通道量化（per-channel）
• GPTQ的重要性：对于获得较好的int4仅权重量化精度是必要的

• GPTQ宏观算法流程：

  • 对某一层估计多个batch的Hessian矩阵
  • 量化W的一列
  • 使用Hessian矩阵H更新未量化的W列（以保持上述方程最小化）
  • 重复步骤2和3，直到所有列都被量化

这张Slides比较了三种不同的量化技术：动态量化、静态量化和仅权重量化。特别的，对于静态量化流程（其它两种已经讲过了）：

• 权重预处理量化
• 激活量化前需要校准
• 使用Int8进行乘法运算
• Int32累加
• 重新缩放到Int8

这张Slides继续介绍了静态量化（Static Quantization）技术，并将其与动态量化进行了对比。

• 量化流程对比：左侧是动态量化流程；右侧是静态量化流程
• 静态量化的特点：通过校准集计算最佳量化参数；包含校准步骤，用于确定激活的量化参数
• 静态量化的优势：对非平稳分布更敏感（能更好地适应数据分布的变化）；对频繁出现的异常值不太敏感
• 整数激活：静态量化使用整数激活（Int8 Activation）；最适合有一系列可量化操作的情况
• 计算过程：两种方法都使用Int8乘法和Int32累加；静态量化在最后阶段重新缩放到Int8，而动态量化重新缩放到浮点数

另外，静态量化的输出也可以Rescale回Float。

总结

Lecture 7主要介绍了基于CUDA和Triton的量化技术在生成式AI模型中的应用。内容包括动态量化、仅权重量化(int8/int4)等不同量化方法的原理、实现和性能比较,以及Smoothquant、GPTQ等量化优化技术的简介。在动态量化和int8/int4 weight only量化实现讲解中，作者分析了Triton相对于CUDA在这些场景的优劣，比如对于int4 weight only 因此Triton本身的限制就不太适合来实现这个cuda kernel。up主还讨论了很多Torch Compiler针对这些量化的优化，比如decode阶段的gemv就让编译器走elementwise mul+reduce的特殊分支以提升性能。学Lecture 7能对量化，CUDA/Triton/Torch Compiler的应用有一个更好的了解，有余力的读者可以看看原视频。视频最后的QA环节也有一些有意思的问题和见解。

作者： BBuf
来源：GiantPandaCV

推荐阅读

• CUDA-MODE 第一课课后实战（上）
• 一文弄懂 LLM 结构化数据生成原理
• 属于嵌入式和移动设备的轻量级 YOLO 模型 ！
• FP8 低精度训练：Transformer Engine 简析

欢迎大家点赞留言，更多Arm技术文章动态请关注极术社区嵌入式AI专栏欢迎添加极术小姐姐微信（id:aijishu20)加入技术交流群，请备注研究方向。