模型量化综述

前言 模型量化作为模型压缩优化的一个常见方法,现有的工具链都非常成熟易用。但很少有人了解其原理与实现。 本文总结了一下常见的量化方法,量化算法,量化如何加速,以及量化算法的实现。 量化为什么能加速 首先普遍认为量化的加速原理是,数据规模变小了,所以计算就加速了。亦或者是整数运算会比浮点运算更快。 但是实际上现代硬件对于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,