模型量化综述

前言 模型量化作为模型压缩优化的一个常见方法,现有的工具链都非常成熟易用。但很少有人了解其原理与实现。 本文总结了一下常见的量化方法,量化算法,量化如何加速,以及量化算法的实现。 量化为什么能加速 首先普遍认为量化的加速原理是,数据规模变小了,所以计算就加速了。亦或者是整数运算会比浮点运算更快。 但是实际上现代硬件对于fp16和fp32,int32的计算速度是差不多的。 真正影响的是速度的主要原因是两个: 1. GPU的访存非常昂贵,量化会让计算数据变小,而相应的访存请求也降低了。 2. 量化结束之后,因为数据大小降低了,因此可以把多条指令进行合并一条进行处理(向量化计算)。在指令层降低了延迟。(包括取值,译码,执行,全都压缩了) 这里稍微回顾一下指令执行的大概流程。分别是,取指,译码,访存,执行,写回。可以看到基本每一个流程都被优化了。 量化方法 对称量化,非对称量化,整数量化 tensor量化,通道量化 待完善 量化算法 minmax kl entropy 待完善…

深度学习数据精度总结

前言 fp32 & fp16 & bf16 & tf32 fp32和fp16都是国际标准IEEE的单精度和半精度,相信计算机专业的都比较熟悉。 bf16是由google brain提出开发的,全称 brain floating point,主要为了加速机器学习的一种格式。 而tf32是由nvidia提出开发的,全称 Tensor float32,也是为了加速机器学习的一种格式。tf32目前只有在ampere架构以上才支持。是为了专门在 tensor core 上加速的一种数据结构。 他们的格式表示是完全一样的,都是指数小数形式,区别只有指数的位数和小数的位数。 在看格式表示之前,以float32为例,回顾一下浮点数的通用表示法。 float32,有1个符号位,8个指数位(exponent),23个小数位(fraction)。它的计算公式见下图: * 符号位: 0代表正数,1代表负数 * 指数位: 其数值范围为 [0, 255], 但是全0和全1有特殊作用,另外指数需要减去一个固定大小的偏置,因此指数的实际范围为 [-126, 127] *…

LLM 推理加速

前言 本文主要是个人对 商汤技术分享 的摘抄,并加上了一些个人理解,以及详解了一些视频中被快速省略的部分。 视频说的比较详细,如果想了解的更加细致建议移步视频。 本人仅供个人技术学习记录与快速分享使用。 LLM 的特性与不同点 如果你使用过chatgpt或者其他LLM模型,你就会知道,在你输入一段话的时候,他会有非常短的延迟,然后会逐字逐句的输出。在这个实际场景下,LLM模型有一个区别于传统语言模型的极大区别。就是他只有一次encoding,但会有N次的decoding。在提交问题到gpt回答的这个时间,其实是encoding的时间,而逐字生成答案的时间,都是decoding的时间。 这个设计给模型放松了很大的latency限制,也很符合人类本身对对话的实际情况。 在LLM中,encoding和decoding的基本上一致的,他们的结构见下图。 其中只有结构中的self attention,在具体实现上,encoding和decoding不相同。 在LLM中 encoding 的结构发生了变化,在论文中都称之为 prefill。这个变化源自…

两个简洁高效的cuda reduce写法

前言 cuda reduce 操作是非常常见的一个操作。在自己写kernel的时候也经常会有需要。 其中以 ReduceMax 和 ReduceSum 最为常见,其他操作原理也是完全一致的。 reduce kernel是很多cuda新手经常会遇到的课题,基本都是拿这个课题进行一步一步优化的。 那么实际工业界如何写一个简洁高效的kernel呢?这里总结2个写法。 不过这里需要注意的是,标题说的是两个,其实原理都是一样的,只是写法不一样。 这里用的都是warp原语写法。 以warp原语有很自然的好处: 1. 这个kernel基本上就是一个warp级kernel,再在外层包一层使其成为一个block级kernel。是一种很规范的kernel写法。 2. warp级kernel能更方便的规避bank conflict。 1. __shfl_down_sync warp kernel 第一种是pytorch的写法 template __inline__ __device__ T WarpReduceSum(T val)

浅谈SIMD和SIMT

之前对这个的概念不够,被问倒了,今天学习补充一下。 个人理解 基于个人的学习经历来说,他们的目的非常明确且一致,都是为了在加快运算速度。 但是两者有一个显著的差别,simd最开始是为了从单指令单数据 拓展到 多指令多数据 后,想要追求更加高效的并行方式。在此前提下在编码层把多指令砍到单指令,在硬件层面上砍掉了单指令以外的所有指令硬件,包括取值,译码,发射等硬件和流程。 通过指令集的方式,取代了砍掉的硬件。 这尽管会带来高效,但也同时带来了编程困难。用并行的思维去解决同步的问题尚且不提,更重要的是,他对高级语言支持不好,你敲代码方式几乎完全变了,包括你的所有变量,所有常用函数,并且会迎来大量的库不相通,不支持的问题。 举个例子: #include #include // 包含SIMD指令集的头文件 int main() { const int size = 8; // 向量大小 alignas(32) float a[size] = {1.0f, 2.0f,

C++原子操作总结

关于乱序 c++程序中,没有依赖关系的语句的执行顺序是无法预测的。 乱序的原因 1. 编译器优化。编译器优化假设程序是单线程环境时,如果编译器出于某种原因会对程序执行顺序并且不会改变其结果,此时是可以乱序执行的。编译器一般不会将函数之后的指令移到函数前面,因为编译器一般不知道函数体内是否含有memory barrier 2. 处理器优化。处理器允许指令乱序,这样避免指令等待资源而暂时hang住,处于闲置状态。eg. cache miss。如果当前指令cache miss,下一条命中,则可能会乱序执行。 3. 存储系统。在每个cpu核中都存在一个store buffer(可以理解成一些特殊寄存器)。在一些写指令执行的时候,写的数据先存在各自的store buffer中,只有自己感知得到,这之后会以先进先出的方式往l2 cache等逐级向下到内存中。而存往store buffer时,cpu就已经认为这条指令执行结束了,因此可能产生乱序。 直面乱序 1. 乱序执行不可避免 2. 如果读写指令涉及的不是共享变量,则不会有坏影响 3. 否则,应该用锁或者原子变…

关于右值

右值 右值定义 右值与左值其实并没有详细官方的定义,官方的其实是一个类似映射表(这个是左值,这个是右值)。但通常来说,下面这个图能覆盖大多数情况。 纯右值(pure right value) 对于大多数临时变量皆为右值。这里的临时变量是指,不能被程序员读写,不存在于内存,其可能只会在寄存器中存在数秒,或者直接被优化掉的变量。 * literal (不含string literal): 42 true nullptr * 返回值不是引用类型的函数 * this指针 * lambda xvalue(expiring lvalue) 即原本为左值,但是程序员确定已经不会再对其读写,可以使用std::move强制让它变为右值。 因此,xvalue必定在一个变量的读写生命周期末尾。 左值 与右值相对,有内存,可以被程序员读写的为左值 * string literal: "2132" * 普通变量: string s * 通过指针访问对象或者对象的数据成员 * 返回是lvalue的函数 * 通过名字访问lvalue对象的数据成…

cnn从入门到入土

cnn全称卷积神经网络(Convolutional Neural Networks),常用于图像分类,语音,机器翻译等。其中最为常见的还是图像分类。 LeNet 总结 各层主要作用 * 卷积层:提取特征 * 池化层:降低计算量 * 全连接:联系各个区域,在分类中结合softmax起到分类器的作用 AlexNet 2012 主要工作与改进点 * 引入ReLU激活函数: ReLU的全称是修正线性单元(Rectified Linear Unit), 其函数表达式为 max(0, x) * 多GPU训练: 加快训练速度 * 引入局部响应归一化(LRN): 实际证明没什么用,在此不对其做解释 * 使用重叠池化:一种池化类型,普通池化是 stride == kernel size,而重叠池化 stride < kernel size * 使用数据增强:一个trick,在数据过少的时候,通过翻转,…