量化与多精度算子:INT8 / FP8、反量化与 per-channel
低精度是大模型推理和训练的标配,原因有两个,都和本系列的主线对应:Tensor Core 篇 说过低精度的算力屋顶更高(FP16 比 FP32 快约 4.5 倍),访存篇 的逻辑则是低精度的数据更小、HBM 流量更少。这一篇讲量化算子的工程细节:量化怎么做、反量化在哪做、不同粒度的权衡、各种精度的硬件路径,以及容易踩的精度坑。
一、低精度的两个收益
收益一:减少访存(memory-bound 算子)
数据从 FP32(4 字节)换成 FP16(2 字节)或 INT8(1 字节),HBM 搬运量直接减半、减到四分之一。对 memory-bound 算子,耗时几乎正比于字节数。
实测一个逐元素算子在 FP32 和 FP16 下的耗时(RTX 3060 Ti,\(n=2^{26}\)):
| 精度 | 耗时 | 加速 |
|---|---|---|
| FP32(4 字节/元素) | 1.363 ms | 1.00× |
| FP16(2 字节/元素) | 0.752 ms | 1.81× |
FP16 快 1.81 倍,接近字节减半带来的 2 倍上限。这就是为什么大模型推理几乎都跑在 FP16/BF16/INT8 上——很多算子卡在带宽,精度减半就接近翻倍提速。
收益二:提高算力(compute-bound 算子)
GEMM 这类 compute-bound 算子,低精度走 Tensor Core 的更高算力屋顶(第 11 篇 实测 FP16 约 73 TFLOP/s vs FP32 16 TFLOP/s)。INT8 在 Tensor Core 上吞吐更高,是推理 GEMM 的常见选择。
二、量化的基本形式
量化把浮点值映射到低位宽整数。最常见的线性量化:
\[ q = \text{round}(x / s) + z, \qquad x \approx s \cdot (q - z) \]
其中 \(s\) 是 scale(缩放因子),\(z\) 是 zero-point(零点)。两种形式:
- 对称量化(\(z=0\)):\(q = \text{round}(x/s)\),映射到 \([-127, 127]\)。实现简单,反量化只需一次乘法,适合权重(通常零对称分布)。
- 非对称量化(\(z \ne 0\)):能更好覆盖非对称分布(如 ReLU 后的激活,全非负),但反量化和矩阵乘的展开更复杂。
INT8 GEMM 的核心:\(x \approx s_x(q_x - z_x)\)、\(w \approx s_w(q_w - z_w)\),于是 \(x \cdot w\) 展开成整数点积 \(\sum q_x q_w\) 加上几个修正项。整数点积在硬件上跑,最后用 \(s_x s_w\) 反量化回浮点。
三、量化粒度:per-tensor 还是 per-channel
scale 用多细的粒度,直接影响精度:
- per-tensor:整个张量一个 scale。最省、最快,但一个离群值(outlier)会撑大 scale,让其他值的有效位宽变低,精度损失大。
- per-channel(per-row/per-column):每个输出通道一个 scale。权重量化常用这个——不同通道的数值范围差异大,逐通道的 scale 能保住精度。代价是反量化时要按通道乘不同的 scale。
- per-group / per-block:每若干个元素一个 scale(如权重每 128 个一组),是 LLM 权重量化(GPTQ、AWQ 等)常用的折中,精度接近 per-channel,开销可控。
粒度越细,精度越好,但 scale 的存储和反量化的索引开销越大。算子实现要把”按哪个维度取 scale”和数据布局对齐,避免反量化时的非合并访问(第 05 篇)。
四、反量化在哪做:算子设计的关键
反量化(dequant)把整数结果转回浮点。它放在哪一步,决定了算子的访存和精度:
- GEMM 内累加用高精度:INT8 输入,但累加器用 INT32 或 FP32,避免累加溢出/精度损失。Tensor Core 的 INT8 MMA 正是 INT8 输入、INT32 累加。
- 反量化融进 epilogue:GEMM 算出 INT32 结果后,在 epilogue(第 16 篇)里乘 \(s_x s_w\)、加 bias、过激活,趁结果还在寄存器时一次完成,不写回 INT32 再读。
- 权重在线 vs 离线反量化:weight-only 量化(如 LLM 推理)里,权重以 INT4/INT8 存(省显存带宽),加载后在 kernel 内反量化成 FP16 再做 FP16 GEMM。这把瓶颈从权重带宽转移到反量化计算,是 LLM 解码阶段(memory-bound)的常见手法。
设计量化算子时,“什么时候转精度”要服从一个目标:让大块数据以低精度搬运,反量化尽量晚、尽量在片上做。
五、各精度的硬件路径
| 精度 | 硬件路径(本卡 / 通用) | 说明 |
|---|---|---|
| FP16 / BF16 | Tensor Core HMMA | 训练/推理主力,本卡支持 |
| TF32 | Tensor Core | FP32 范围、降低尾数,训练用,本卡支持 |
| INT8 | Tensor Core IMMA;CUDA core __dp4a |
__dp4a 一条指令做 4 路 INT8 点积累加到
INT32 |
| INT4 | Tensor Core(部分架构)/ 解包后算 | LLM weight-only 量化常用 |
| FP8 (E4M3/E5M2) | Hopper sm_90+ Tensor Core | 本卡(Ampere)不支持,需 H100 及以后 |
__dp4a(a, b, c) 是 CUDA core 上的 INT8
点积内建:把两个打包成 int32 的 4×INT8 向量点积,累加到
INT32。它适合没有 IMMA 的场景或小规模 INT8 计算;大规模 INT8
GEMM 还是走 Tensor Core 的 IMMA,吞吐高得多。FP8 是当前 LLM
训练/推理的前沿精度,但需要 Hopper
及更新的硬件,本系列测试卡跑不了,相关讨论见 第 21 篇。
六、精度对齐的工程坑
量化算子最容易出问题的是数值正确性,几个常见坑:
- 累加精度不足:INT8 点积若用 INT16 累加会溢出,必须 INT32;FP16 累加长序列会损失精度,常需 FP32 累加。
- round 方式不一致:训练时的量化模拟(fake quant)和推理 kernel 的 round 行为(round-half-to-even vs round-half-away)不一致,会导致训练-推理结果对不上。
- scale 的布局与广播:per-channel scale 要按正确维度广播,索引错了结果整体偏移。
- clamp / 饱和:量化后要 clamp 到 \([-127,127]\) 或 \([0,255]\),漏掉会溢出回绕。
- 离群值:per-tensor 下的离群值会毁掉精度,LLM 激活量化常需要专门处理离群通道(如保留部分 FP16)。
这些都要靠和高精度参考实现逐元素对齐来验证,方法见 调试与数值正确性篇。
七、小结与下一步
- 低精度有两个收益:减少访存(实测 FP16 逐元素比 FP32 快 1.81 倍)和提高算力(Tensor Core 低精度屋顶更高)。
- 线性量化用 scale + zero-point,对称适合权重、非对称适合激活;粒度从 per-tensor 到 per-group 在精度和开销间权衡。
- 反量化应尽量晚、在 epilogue 片上完成;累加用高精度(INT32/FP32)。
- INT8 走 Tensor Core IMMA 或 CUDA core
__dp4a,FP8 需 Hopper(本卡不支持)。 - 精度对齐(累加位宽、round、scale 布局、clamp、离群值)是量化算子的主要正确性风险。
单卡算子讲完。下一篇跨出单卡,看多卡场景下 通信与计算重叠:NCCL collective 与 kernel overlap。
同主题继续阅读
把当前热点继续串成多页阅读,而不是停在单篇消费。
【GPU 算子工程】Tensor Core 与 MMA:wmma、mma.sync 与数据布局
Tensor Core 把矩阵乘做进专用硬件。实测 RTX 3060 Ti 的 FP16 Tensor 吞吐达 72.8 TFLOP/s,约 FP32 峰值的 4.5 倍。讲清 MMA 指令、wmma fragment API、数据布局与精度要求,以及为什么喂数据才是真正的瓶颈。
【GPU 算子工程】趋势:TMA、Blackwell、ThunderKittens 与编译器协同
算子工程的前沿方向:Hopper 的 TMA 异步搬运与 wgmma、Blackwell 的更低精度、ThunderKittens 等 tile 级库降低门槛、Triton/MLIR 的编译器自动生成算子。本系列测试卡为 Ampere,相关特性为引用与前瞻,明确标注。
【GPU 算子工程】全景:算子工程在 AI 计算栈的位置
从框架一行 matmul 到 PTX/SASS,拆开 AI 计算栈的分层:框架算子、算子库、手写 kernel、编译器生成。回答工程师什么时候才需要自己写或调 kernel,以及本系列的实验环境与方法。
【GPU 算子工程】GPU 执行模型:SM、warp、线程层次与 occupancy
讲清 grid/block/warp 如何映射到 SM,SIMT 执行与 32 线程 warp 的本质,分支发散为何昂贵(实测 1.7 倍),以及 occupancy 的含义。建立一切 GPU 性能优化的硬件直觉。