神经网络INT8量化~部署

这里详解量化的基本概念以及卷积量化的实际操作。

刚开始接触神经网络量化是2年前那会,用NCNN和TVM在树莓派上部署一个简单的SSD网络。那个时候使用的量化脚本是参考于TensorRT和NCNN的PTQ量化(训练后量化)模式,使用交叉熵的方式对模型进行量化,最终在树莓派3B+上部署一个简单的分类模型(识别剪刀石头布静态手势)。

转眼间过了这么久啦,神经网络量化应用已经完全实现大面积落地了、相比之前成熟多了!

我工作的时候虽然也简单接触过量化,但感觉还远远不够,趁着最近项目需要,重新再学习一下,也打算把重新学习的路线写成一篇系列文,分享给大家。

本篇系列文的主要内容计划从头开始梳理一遍量化的基础知识以及代码实践。因为老潘对TensorRT比较熟悉,会主要以TensorRT的量化方式进行描述以及讲解。不过TensorRT由于是闭源工具,内部的实现看不到,咱们也不能两眼一抹黑。所以也打算参考Pytorch、NCNN、TVM、TFLITE的量化op的现象方式学习和实践一下。

当然这只是学习计划,之后可能也会变动。对于量化我也是学习者,既然要用到这个技术,必须要先理解其内部原理。而且接触了挺长时间量化,感觉这里面学问还是不少。好记性不如烂笔头,写点东西记录下,也希望这系列文章在能够帮助大家的同时,抛砖引玉,一起讨论、共同进步。

参考了以下关于量化的一些优秀文章,不完全统计列了一些,推荐感兴趣的同学阅读:

  • 神经网络量化入门--基本原理(https://zhuanlan.zhihu.com/p/149659607)

  • 从TensorRT与ncnn看卷积网络int8量化(https://zhuanlan.zhihu.com/p/387072703)

  • 模型压缩:模型量化打怪升级之路 - 1 工具篇(https://zhuanlan.zhihu.com/p/355598250)

  • NCNN Conv量化详解(一)(https://zhuanlan.zhihu.com/p/71881443)

当然在学习途中,也认识了很多在量化领域经验丰富的大佬(田子宸、JermmyXu等等),嗯,这样前进路上也就不孤单了。

OK,废话不多说开始吧。

Why量化

我们都知道,训练好的模型的权重一般来说都是FP32也就是单精度浮点型,在深度学习训练和推理的过程中,最常用的精度就是FP32。当然也会有FP64、FP16、BF16、TF32等更多的精度:

FP32 是单精度浮点数,用8bit 表示指数,23bit 表示小数;FP16半精度浮点数,用5bit 表示指数,10bit 表示小数;BF16是对FP32单精度浮点数截断数据,即用8bit 表示指数,7bit 表示小数。TF32 是一种截短的 Float32 数据格式,将 FP32 中 23 个尾数位截短为 10 bits,而指数位仍为 8 bits,总长度为 19 (=1 + 8 + 10) bits。

对于浮点数来说,指数位表示该精度可达的动态范围,而尾数位表示精度。之前老潘的一篇文章中提到,FP16的普遍精度是~5.96e−8 (6.10e−5) … 65504,而我们模型中的FP32权重有部分数值是1e-10级别。这样从FP32->FP16会导致部分精度丢失,从而模型的精度也会下降一些。

其实从FP32->FP16也是一种量化,只不过因为FP32->FP16几乎是无损的(CUDA中使用__float2half直接进行转换),不需要calibrator去校正、更不需要retrain

而且FP16的精度下降对于大部分任务影响不是很大,甚至有些任务会提升。NVIDIA对于FP16有专门的Tensor Cores可以进行矩阵运算,相比FP32来说吞吐量提升一倍。

实际点来说,量化就是将我们训练好的模型,不论是权重、还是计算op,都转换为低精度去计算。因为FP16的量化很简单,所以实际中我们谈论的量化更多的是INT8的量化,当然也有3-bit、4-bit的量化,不过目前来说比较常见比较实用的,也就是INT8量化了,之后老潘的重点也是INT8量化。

那么经过INT8量化后的模型:

  • 模型容量变小了,这个很好理解,FP32的权重变成INT8,大小直接缩了4倍

  • 模型运行速度可以提升,实际卷积计算的op是INT8类型,在特定硬件下可以利用INT8的指令集去实现高吞吐,不论是GPU还是INTEL、ARM等平台都有INT8的指令集优化

  • 对于某些设备,使用INT8的模型耗电量更少,对于嵌入式侧端设备来说提升是巨大的

所以说,随着我们模型越来越大,需求越来越高,模型的量化自然是少不了的一项技术。

如果你担心INT8量化对于精度的影响,我们可以看下NVIDIA量化研究的一些结论:

出自《INTEGER QUANTIZATION FOR DEEP LEARNING INFERENCE: PRINCIPLES AND EMPIRICAL EVALUATION》。

量化现状

量化技术已经广泛应用于实际生产环境了,也有很多大厂开源了其量化方法。不过比较遗憾的是目前这些方法比较琐碎,没有一套比较成熟比较完善的量化方案,使用起来稍微有点难度。不过我们仍可以从这些框架中学习到很多。

Google

谷歌是比较早进行量化尝试的大厂了,感兴趣的可以看下Google的白皮书Quantizing deep convolutional networks for efficient inference: A whitepaper以及Quantization and Training of Neural Networks for Efficient Integer-Arithmetic-Only Inference

TensorFlow很早就支持了量化训练,而TFLite也很早就支持了后训练量化,感兴趣的可以看下TFLite的量化规范 (https://www.tensorflow.org/lite/performance/quantization_spec) ,目前TensorRT支持TensorFlow训练后量化的导出的模型。

TensorRT

TensorRT在2017年公布了自己的后训练量化方法,不过没有开源,NCNN按照这个思想实现了一个,也特别好用。不过目前TensorRT8也支持直接导入通过ONNX导出的QTA好的模型,使用上方便了不少,之后老潘会重点讲下。

NVIDIA自家也推出了针对Pytorch的量化工具(为什么没有TensorFlow,因为TF已经有挺好用的官方工具了),支持PTQ以及QTA,称为Pytorch Quantization,之后也会提到。

TVM

TVM有自己的INT8量化操作,可以跑量化,我们也可以添加自己的算子。不过TVM目前只支持PTQ,可以通过交叉熵或者percentile的方式进行校准。不过如果动手能力强的话,应该可以拿自己计算出来的scale值传入TVM去跑,应该也有人这样做过了。

比较有参考意义的一篇:

  • ViT-int8 on TVM:提速4.6倍,比TRT快1.5倍(https://zhuanlan.zhihu.com/p/365686106)

当然还有很多优秀的量化框架,想看详细的可以看这篇(https://zhuanlan.zhihu.com/p/355598250),后续如果涉及到具体知识点老潘也会再提到。

量化基本知识

进入主题前需要提两个概念,也就是量化的两个重要过程,一个是量化(Quantize),另一个是反量化(Dequantize):

  • 量化就是将浮点型实数量化为整型数(FP32->INT8)

  • 反量化就是将整型数转换为浮点型实数(INT8->FP32)

量化和反量化操作在最终的模型推理中都会用到,接下来我们就具体说下。

之后实数就代表我们的FP32浮点数,而整数就代表INT8整型数。

量化操作

基于线性量化的对称量化和非对称量化

对不对的关键在于我们是否是采用对称量化,什么是对称量化呢?这里的对称指的是以0为中心进行量化(还有另一种说法,这里老潘先略过),然后0两边的动态范围都是一样的。

可以看上图,左边是非对称量化,右边是对称量化(也称为Affine quantization和Scale quantization)。可以观察到: 

需要说明一点,不论是非对称还是对称量化,是基于线性量化(也可以称作均匀量化)的一种。线性量化将FP32映射到INT8数据类型,每个间隔是相等的,而不相等的就称为非线性量化。非线性量化因为对部署并不是很友好,虽然能够更好地捕捉到权重分布的密集点,但感觉用的并不多,这里也就先不多说了。

关于详细的非对称量化,对称量化对比可以参考这篇文章:

  • Affine Quantization vs Scale Quantization(https://liq.opengenus.org/affine-quantization-vs-scale-quantization/)

对称量化

接下来的重点是对称量化,也就是TensorRT中使用的量化方式,这里的范围也就是[-127,127],因为只比[-128,127]少了一个范围,所以实际量化中并没有太大的影响。

 

那么实际操作过程中,scale系数是怎么用呢?或者说这个量化系数是怎么作用于所有的输入、所有的权重呢?

一般量化过程中,有pre-tensorpre-channel两种方式,pre-tensor显而易见,就是对于同一块输入(比如某个卷积前的输入tensor)我们采用一个scale,该层所有的输入数据共享一个scale值;而pre-channel呢一般是作用于权重,比如一个卷积的权重维度是[64,3,3,3](输入通道为3输出通道为64,卷积核为3x3),pre-channel就是会产生64个scale值,分别作用于该卷积权重参数的64个通道。

为什么权重不能是pre-tensor呢?这个对精度的影响太大了,所以一般不用。输入就可以pre-tensor?当然可以,也经过测试了,对精度的影响不是很大,完全可以用。

那为什么权重必须是pre-channel呢?不能是每个权重值都有自己的scale么?呃,这个问题嘛,首先可以想到,这个计算量,应该挺大,其次嘛,让我们分析一下。

卷积操作量化

铺垫了这么多,那么接下来说下量化最核心的操作吧,量化过程中最核心的操作当然是卷积量化

我们都知道卷积操作可以拆分为im2col+sgemm,而大部分的计算都在矩阵运算也就是sgemm中,我们量化的重点也就是这个操作。以前是FP32计算,而现在变成INT8去计算,这是怎么转换的呢?

那么pre-channel又是怎么来的呢?

还记得之前说过的 im2col+sgemm 操作吗(如果不记得强烈建议去看看),其中的sgemm是这样的,需要注意,下图左边的kernel矩阵,每一行代表一个输出通道的kernel集合(这里因为输入图像是三通道的,因此kernel有三个,不同颜色代表一个kernel):

这就是pre-channel或者详细点就是per-output-channel也就是卷积输出通道,我们对每一个卷积权重的输出通道那一维进行量化,然后共享一个scale,这也就呼应了上述的公式!

后记

到此我们已经讲述了量化的基本概念以及卷积量化的实际操作是什么样的,当然想说的还有很多...就是现在实在写不动了,关于非对称量化的公式以及为什么非对称量化计算量比较大,之后再谈.

 whaosoft aiot http://143ai.com 

  • 1
    点赞
  • 9
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
以下是一个简单的使用CUDA C编写卷积操作的示例代码: ```c #include <stdio.h> #include <stdlib.h> #include <cuda_runtime.h> #define BLOCK_SIZE 16 __global__ void convolve(float *input, float *output, float *kernel, int input_width, int input_height, int kernel_size) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; float sum = 0.0f; for (int i = 0; i < kernel_size; ++i) { for (int j = 0; j < kernel_size; ++j) { int input_row = row + i - kernel_size / 2; int input_col = col + j - kernel_size / 2; if (input_row >= 0 && input_row < input_height && input_col >= 0 && input_col < input_width) { sum += input[input_row * input_width + input_col] * kernel[i * kernel_size + j]; } } } output[row * input_width + col] = sum; } int main() { int input_width = 512; int input_height = 512; int kernel_size = 5; float *input = (float*)malloc(input_width * input_height * sizeof(float)); for (int i = 0; i < input_width * input_height; ++i) { input[i] = rand() % 256 / 255.0f; } float *kernel = (float*)malloc(kernel_size * kernel_size * sizeof(float)); for (int i = 0; i < kernel_size * kernel_size; ++i) { kernel[i] = rand() % 256 / 255.0f; } float *output = (float*)malloc(input_width * input_height * sizeof(float)); float *d_input, *d_kernel, *d_output; cudaMalloc(&d_input, input_width * input_height * sizeof(float)); cudaMalloc(&d_kernel, kernel_size * kernel_size * sizeof(float)); cudaMalloc(&d_output, input_width * input_height * sizeof(float)); cudaMemcpy(d_input, input, input_width * input_height * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(d_kernel, kernel, kernel_size * kernel_size * sizeof(float), cudaMemcpyHostToDevice); dim3 block_size(BLOCK_SIZE, BLOCK_SIZE); dim3 grid_size((input_width + BLOCK_SIZE - 1) / BLOCK_SIZE, (input_height + BLOCK_SIZE - 1) / BLOCK_SIZE); convolve<<<grid_size, block_size>>>(d_input, d_output, d_kernel, input_width, input_height, kernel_size); cudaMemcpy(output, d_output, input_width * input_height * sizeof(float), cudaMemcpyDeviceToHost); cudaFree(d_input); cudaFree(d_kernel); cudaFree(d_output); free(input); free(kernel); free(output); return 0; } ``` 该示例代码中,`convolve()`函数是卷积操作的核函数,它接收输入图像、输出图像和卷积核作为输入参数,以及输入图像尺寸和卷积核尺寸。在核函数中,每个线程负责计算输出图像中的一个像素值,它将卷积核与输入图像中对应像素的值相乘,并将结果累加到一个变量中。最后,输出图像中对应像素的值被赋为累加的结果。 在主函数中,我们首先生成随机的输入图像和卷积核,然后在GPU上分配内存,将输入图像和卷积核从主机内存复制到设备内存中,调用卷积核函数进行卷积操作,最后将输出图像从设备内存复制到主机内存中,并释放分配的内存。 在实际使用中,您需要根据自己的需求修改输入图像、卷积核和卷积核函数,以适应不同的场景。
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值