推理加速之-GPU加速

模型推理加速一般可以从两方面考虑。一方面是模型本身的剪枝及蒸馏，另一方面就是充分利用GPU进行推理加速了。这篇文章主要讲解了GPU加速概述，原理以及CUDA概述。

先看一下 NVIDIA GPU 显卡架构发展历程：

• Tesla（特斯拉）2008年，应用于早期的 CUDA 系列显卡芯片中，并不是真正意义上的 GPU 芯片。

• Fermi（费米）2010年，是第一个完整的 GPU 计算架构。首款可支持与共享存储结合纯 cache 层次的 GPU 架构，支持 ECC(Error Correcting Code) 的 GPU 架构。

• Kepler（开普勒）2012年，Fermi 的优化版。

• Maxwell（麦克斯韦）2014年，首次支持实时的动态全局光照效果，

• Pascal（帕斯卡）2016年，GPU 将处理器和数据集成在同一个程序包内，以实现更高的计算效率。

• Volta（伏打）2017年，首次将一个 CUDA 内核拆分为FP32 和 INT32 两部分，首次支持混合精度运算，提高了5倍于 Pascal 计算速度，还增加了专用于深度学习的 Tensor Core 张量单元。

• Turing（图灵）2018年，增加了 RT Core 专用光线追踪处理器，将实时光线追踪运算加速至上一代架构的 25 倍，并能以高出 CPU 30 多倍的速度进行电影效果的最终帧渲染。去掉了对 FP64 计算的支持。

• Ampere（安培）2020年，重新支持 FP64，新增异步拷贝指令能够从 global memory 中将数据直接加载到 SM shared memory，降低中间寄存器堆（RF）的需求。新增 BF16 数据类型，专为深度学习优化。

• Hopper架构 2022年3月 这一全新架构以美国计算机领域的先驱科学家 Grace Hopper 的名字命名，将取代两年前推出的 NVIDIA Ampere 架构。

1. GPU加速概述

• 存取效率：GPU和显存之间的数据交换效率，其中全局内存具有最大的容量和最慢的访问效率，且对是否对齐和连续访问很敏感；共享内存访问速度快，且对是否对齐和连续访问不敏感，但是对Bank Conflict非常敏感；寄存器具有最快的访问速度，只对每个线程可见，线程内可多使用寄存器，但是需注意，一个SM内的寄存器数量有限，当单个线程的寄存器数量超过限制，会影响线程的实际占用率，从而影响加速效果；其他存储包括纹理内存常量内存等。具有的一些特殊特性使其有时候在特定的情况下被使用以获得更高的效率，比如纹理内存带有的纹理缓存具备硬件插值特性，可以实现最邻近插值和线性插值，且针对二维空间的局部性访问进行了优化，所以通过纹理缓存访问二维矩阵的邻域会获得加速，这个特性使得纹理内存在一些图像处理算法中具有一定的优势。

• 计算效率：指出去内存交换过程以外的算法计算部分的效率。GPU有三类基础运算：整数运算、单精度浮点运算、双精度浮点运算。其中单精度浮点运算速度最快而双精度浮点运算速度最慢，FLOPS（floating-point operations per second, 每秒执行的浮点运算次数）也是衡量GPU运算性能的关键指标。

优化要点：

• 要展现足够的并行性；

为了最大程度的利用GPU多线程的优势，应该在GPU上安排尽量多的并发任务，以使指令带宽和内存带宽都达到饱和，在一个SM(流处理器)中保证有足够多的并发线程束，这不单单是要为GPU每个线程都安排任务，还需要检查SM资源占用率的限制因素（共享内存、寄存器以及计算周期等）以找到达到最佳性能的平衡点，因为GPU的内存资源是有限的，为每个线程分配的资源也是有限的，如果算法设计者在一个线程中使用了过多的共享内存或者寄存器，那么并发运行的线程数必然会减少，使得SM资源的实际占用率小于理论占用率；另一方面可以为每个线程/线程束分配更多独立的工作。

• 优化内存访问；

大多数GPU算法的性能瓶颈都在于内存访问速度，由于显存访问的高延迟和低效率，内存访问模式对内核性能有着显著的影响。内存访问优化的目标是最大限度的提高内存带宽的利用率，重点在于优化内存访问模式和保证充足的并发内存访问。

在GPU中，线程是以线程束为单位执行的，一个线程束包含32个线程，所以一方面最好将并发线程数设置为32的倍数，另一方面当一个线程束发送内存请求(加载或存储)时，都是32个线程一起访问一个设备内存块，因此对于全局内存来说，最好的访问模式就是对齐和合并访问，对齐内存访问需要所需的设备内存的第一个地址是32字节的倍数，合并内存访问指的是通过线程束中的32个线程来访问一个连续的内存块。

这表示在算法设计中，一定要尽量为一个线程束的现成分配连续的内存块，比如0~31号线程(同一个线程束)访问影像中连续存储的31个像素，而不是访问不连续的31个像素，由于合并访问对内存访问效率影响非常大，因此需要严格遵守该要求。

共享内存由于是片上内存，所以比本地和设备的全局内存具有更高的带宽和更低的延迟，使用共享内存有两个原因：①. 减少全局内存的访问次数；②. 通过重新安排数据布局，避免未合并的全局内存的访问。

在物理角度上，共享内存通过一种线性方式排列，通过32个存储体（bank）进行访问。Fermi和Kepler架构各有不同的默认存储体模式：4字节存储体模式和8字节存储体模式，共享内存地址到存储体的映射关系随着访问模式的不同而不同，当线程束中的多个线程在同一存储体中访问不同字节时，会发生存储体冲突（Bank Conflict）,由于共享内存重复请求，所以多路存储体冲突可能要付出很大的代价，应该尽量避免存储体冲突，每个存储体（Bank）每个周期只能指向一次操作（一个32bit 的整数或者一个单精度的浮点型数据），一次读或者一次写，也就是说每个存储体（Bank）的带宽为每周期 32bit，比如一个32*32的二维单精度浮点数组，每一列属于一个Bank，如果一个线程束里的不同线程访问该数组里同一列的不同数据，则会发生Bank Conflict，解决或减少存储体冲突的一个非常简单有效的方法是填充数组，在合适的位置添加填充字，可以使其跨不同存储体进行访问，从而减少延迟并提高了吞吐量。

寄存器是GPU上最快的存储机制，但是数量非常有限，如果一个线程使用过多的寄存器，会导致SM能够同时启动的线程数变少，实际上很多情况下寄存器都成为了资源占用率无法达到100%的主要限制条件，所以往往要注意监控寄存器的数量，当数量没有超标时，适当的增加数量可以提升性能，而一旦数量超标，最好还是将寄存器的数量减少以保证100%的资源占用率，这可以通过重新排列代码的顺序来实现，比如当变量的赋值和使用靠的很近时，编译器会重复使用少量寄存器以达到减少寄存器数量的目的。

• 优化指令执行；

GPU属于单指令多数据流架构，每个线程束中的所有线程在每一步都执行相同的指令，如果每个指令都能够得到对结果有效的运算值，就能够避免线程的浪费，而如果由于条件分支造成线程束内有不同的控制流路径，则线程运行可能出现分化，这时线程束必须顺序执行每个分支路径，并禁用不在此执行路径上的线程，而如果算法的大部分时间都耗在分支代码中，必然显著的影响内核性能，所以尽量避免使用分支是很关键的，或者尽量使分支有非常大的概率执行对结果有效的哪一个路径。

2. GPU框架与异构并行计算
一个的异构并行架构，包括一个CPU及其内存和一个GPU机器内存，GPUI设备端通过PCle总线与基于CPU的主机端进行交互，一个异构并行应用包括主机代码和设备代码，分别运行在主机端和设备端。应用由CPU初始化，在设备端进行数据运算前，CPU负责管理设备端的环境，代码和数据。称host为CPU及其内存，device为GPU及其内存。

CPU计算适合处理控制密集型任务，GPU计算适合处理包含数据并行的计算密集型任务。在CPU上执行串行部分或任务并行部分，在GPU上执行数据密集型并行部分，这种异构并行架构使得计算能力可以充分被利用。

3. CUDA编程模型

3.1 cuda软件体系

• cuda驱动API(cuda driver api)是一种基于句柄的底层接口，大多数对象通过句柄被引用，其函数前缀均为cu，在调用driver api前，必须进行初始化，再创建cuda上下文。该上下文关联到特定设备并成为主机线程的当前上下文，通过加载PTX汇编形式，或二进制对象形式的内核，然后启动内核计算。driver api可以通过直接操作硬件执行一些复杂的功能，但其变成较为复杂，难度较大；

• cuda运行时api(cuda runtime api)runtime api 对 driver api 进行了一定的封装，隐藏了部分实现细节，因此使用起来更加方便。runtime api没有专门的初始化函数，他将在第一次调用运行函数时，自动完成初始化。使用时，通常需要包含头文件cuda_runtime.h，其函数前缀均为cuda。

3.2 cuda函数库

• CUFFT：利用cuda进行傅里叶变换的函数库

• CUBLAS：利用cuda进行加速的完整标准矩阵与向量的运算库

• CUDPP：并行操作函数库

• CUDNN：利用cuda进行深度卷积神经网络

3.3 cuda应用程序

• 分配 host 内存，并进行数据初始化（CPU初始化）

• 分配 device 内存，并从 host 将数据拷贝到 device 上（GPU初始化）

• 调用 CUDA 的核函数在 device 上完成指定的运算（GPU并行运算）

• 将 device上的运算结果拷贝到 host 上（将GPU结果传回CPU）

• 释放 device 和 host 上分配的内存（初始化清空）

3.4 cuda硬件结构

• SP(Streaming Processor)也成为了CUDA core，是最基本的处理单元，最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。

• SM(Streaming Multiprocessor)多个SP加上其他资源组成一个SM，也叫GPU大核，其他资源如包括warp scheduler，register，shared memory 等。SM可以看做GPU的心脏（类似 CPU 核心）。每个 SM 都拥有 register 和 shared memory，CUDA 将这些资源分配给所有驻留在 SM 中的线程，但资源非常有限，SM 结构如下图所示。

每个 SM 包含的 SP 数量依据 GPU 架构而不同，如 Fermi 架构 GF100 是 32 个，GF10X 是 48 个，Kepler 架构都是 192 个，Maxwell 都是128 个。

4. kernel，thread，block，grid与warp

4.1 cuda线程模型

cuda的线程模型从小到大依次是：

• Thread，线程，并行的基本单位；

• Block，线程块，互相合作的线程组，线程块有如下几个特点：

• 以1维、2维或3维组织

• 允许彼此同步

• 可以通过共享内存快读交换数据

• Grid，网格，由一组Block组成

• 共享全局内存

• 以1维、2维组织

4.1.1 kernel

4.1.2 warp

对应下图：

从硬件角度来看，所有的 thread 以一维形式组织，每个 thread 都有个唯一的 ID，于是作为补全整数倍的 thread 在所在的 warp 中为 inactive 状态，会额外消耗 SM 资源，所以要设定 block 中的 thread 一般为32的倍数。
从硬件角度和软件角度解释 CUDA 的线程模型:

4.1.3 线程索引

确定线程的唯一索引，以2D grid和2D block的情况为例。

若需要计算的数值矩阵在内存中是row-major(行主序)线性存储的，如下：

将 thread 和 block 索引映射到矩阵坐标
ix = threadIdx.x + blockIdx.x * blockDim.x
iy = threadIdx.y + blockIdx.y * blockDim.y
idx = iy * nx + ix

以下为block和thread索引，矩阵坐标以及线性地址之间的关系：

在实践应用中，常常会多一维 grid， 那就是三维情况的索引，设 (gridDim.x, gridDim.y) = (2, 3), (blockDim.x, blockDim.y) = (4, 2)，以 thread_id(3,1) block_id(0,1) 为例：(grid>block>thread)

可以得到：

ix = threadIdx.x + blockIdx.x * blockDim.x = 3 + 0 * 4 = 3
iy = threadIdx.y + blockIdx.y * blockDim.y = 1 + 1 * 2 = 3
coordinate(3,3)
global index: idx = iy * blockDim.x * gridDim.x + ix = 3 * 4 * 2 + 3 = 27

以上内容均为学习过程中笔记以及增加了一些自己的思考和计算，如有什么问题请各位批评指正。后面会记录一些关于CUDA C编程加速的内容。

【参考】

1. https://zhuanlan.zhihu.com/p/384826488

2. https://zhuanlan.zhihu.com/p/683016265

3. 知乎：丽台科技

4. https://zhuanlan.zhihu.com/p/674217175