AI infra：性能模型与逐元素优化
围绕 GPU 程序性能优化中的核心问题展开，从计算与访存瓶颈出发，介绍 Roofline 性能模型如何用于判断程序是受算力限制还是受内存带宽限制。随后以向量加法为例，分析算术强度、全局内存访问、内存对齐、内存合并以及向量化访存对性能的影响，并进一步讨论 FP16、half2 等逐元素优化手段，帮助理解低算术强度算子在 CUDA 中的优化思路。
性能能否进一步提升？
之前提到 HtoD 和 DtoH 的耗时远大于核函数的时间，这说明瓶颈主要出现在 CPU 和 GPU 通信时间上，即传统带宽相较于算力更加是瓶颈。
GPU 内部的带宽其实也会出现这样的情况，从 NVIDIA A100 的案例来看，A100 的 FP32 算力达到了 19.5 TFLOPS，即每秒可以进行 次计算，内存带宽仅仅为 2039 GB/s，这说明每从显存里搬运 1 个字节的数据，平均至少要做大约 9.6 次浮点运算，GPU 才有机会把 FP32 算力跑满，显存的速度相对于算力而言是有限的。
根据摩尔定律，算力每 18~24 个月翻倍，但是 DRAM 每年只提升大约 10%，现有的计算机大多是冯·诺依曼架构，数据必须在处理器和内存之间频繁传输，当数据量大或者计算密度低时，就会遇到内存瓶颈。
对于一个 Float32 加法而言，每次要进行 12 字节的内存访问（2 次读+1 次写），以及一次加法操作，本身需要的内存访问较多，而计算量很少，理论上加法的性能瓶颈在内存访问上，那又该怎么量化这个性能瓶颈呢？
Roofline 模型
Roofline 模型是一种用于判断程序性能瓶颈在哪里的分析模型，它把程序的性能上限用一张图表示出来，用于表示当前程序是受算力限制还是内存带宽限制：
纵轴通常是性能，如 FLOP/s，横轴通常是算术强度，即 ，表示每搬运 1 字节数据，做了多少次浮点运算。
图中的拐点也称为脊点，表示可达到该硬件理论最高性能的最低计算强度。在拐点以左，表示当前的算术强度下主要受访存影响，计算任务是访存密集型，而在拐点以右，表示当前的算术强度下主要受计算影响，计算任务是计算密集型。
仍以 Float32 为例：，而在 A100 上之前计算得到的 ，由此可见 Float 加法确实是一个访存密集型算子。
Nsight Compute
NVIDIA 还提供了一个内核级性能分析工具 Nsight Compute，可以通过使用 ncu cli 来对 cuda 程序进行性能分析：
也可以导出文件后使用 ncu GUI 查看：
每个核函数一行，且包含函数名、GPU 运行时间、计算吞吐和内存吞吐等信息。
还能够查看 roofline 图：
可以看到加法算子的访存时间要比计算时间多出许多：
再一次证实了加法算子是访存密集型的，那既然加法的性能瓶颈在内存访问上，为了提高整体的效率，就该提升其访存效率，要增大内存带宽的使用，就要一次多传输数据，那么应该怎么做呢？
向量化
答案就是把标量操作转化为向量操作，通过一次同时操作多个标量数据来达到一次多传输数据的目的，以提升访存效率。
向量化的本质就是 SIMD（Single Instruction Multiple Data），通过同时对一组数据中的每一个分别执行相同的操作从而实现空间上的并行，这与之前的 SIMT 听起来很像，但是实际上二者仍有区别：
使用向量化访存最简单的方式是内置向量化访存类型：float1 、float2 、float3、float4。
向量化访存并行加法
可以使用一个模板 add() 函数来封装不同类型的加法操作：
其中的 __device__ 表示这是一个仅在 GPU 上运行的函数，make_float2()、make_float4() 为 float2 和 float4 的创建函数。向量化访存后的结果如下：
完成向量化访存后，float2 、float4 时间均有效提升，而 float3 却显著得比另外两个慢。这就涉及到一个内存对齐与内存合并的问题。
GPU 全局内存访问以内存事务为单位（由总线宽度决定），这说明即便只需要 1 字节的数据，也必须搬运整个数据块，全局内存的基础事务单位是 32 bytes，缓存行为 128 bytes，所谓的内存对齐，即事务的起始地址需要是 32 的整数倍。
之前提到 CUDA 的线程结构为 Grid、Block、Thread，但是实际上调度的时候并不是每个线程单独调度的，而是以线程束（warp）为最小单位进行调度的，CUDA 的一个 warp 是 32 个线程，即一个周期内，一个 warp 中的线程统一执行同一条指令。内存合并就是将同一个 warp 内的多个内存访问合并为少数事务，最大化内存带宽利用率。
要实现内存合并需要的条件有 warp 中的线程需要访问连续的内存地址，并且访问的起始地址必须满足架构对齐要求。
因为 float2 和 float4 分别是 8 字节和 16 字节，满足 的内存对齐要求，能够完美触发 64-bit 或 128-bit 的单指令向量加载，性能能够达到最佳。float3 则是 12 字节，不是 2 的幂次方，因此无法合并。
向量化访存性能提高
从以上可以看出，使用向量化访存确实会提升原本访存受限的并行加法性能。用更直观的方式来看，之前的编译流程中提到设备端编译流程中间每个虚拟架构会编译出一个 .ptx 文件，这里面包含着该文件对该虚拟架构的 PTX （Parallel Thread Execution）代码，这是一种低级的并行线程执行的虚拟指令集。
通过对 PTX 的代码查看可以看到 float3 确实没有能够合并访存，而是通过单纯的三个 float 的操作来执行的。
再一步提升？
对于向量加法这种低算术强度、强带宽需求的程序，想要进一步提升性能，本质上是在同样的算术强度下提高更多的带宽，但是如果已经做过一些访存优化，例如合并访存，向量化访存，那么继续往上提升的空间是有限的。
想让性能点进一步提升，不如让算术强度变高，让每搬运一字节数据做更多计算。想要提升算术强度，根据算术强度的公式：
主要有三种方式：
- 提高 ：增加计算量，可以通过算子融合的方式，例如原本可能有很多个小的算子，每个算子都要从显存读取数据、做一点计算、把中间结果写回显存，这样会导致 很大，而 不高，如果把多个算子融合成为一个 kernel，那么同一批数据读进来以后，可以在寄存器或 shared memory 中连续做更多操作，最后统一写回，这样计算量增加了而中间回写减少了，就能提高 AI。
- 降低 ：减少访存压力，传输更少的数据。例如减少不必要的全局内存读写、提高数据复用、用 shared memory/register 暂存、避免把中间结果频繁写回显存，改善 cache 命中以及压缩数据类型，例如将 FP32 改成 FP16/BF16。
- 提高 W 的同时降低 Q，综合以上两种方法。
针对向量加法，就可以通过传输更少的数据来进行，使用 FP16 来进行计算。
float 是 IEEE754 binary32,单精度浮点数，总共 32 位，由 1 位符号位、8 位指数位、23 位位数组成。half 则是 IEEE754 binary 16，半精度浮点数，总共 16 位，由 1 位符号位、5 位指数位、10 位尾数位组成。当任务能够接受更小的数值范围和更低的精度时，就可以把单精度浮点数换为半精度浮点数。在 CUDA 中使用 half 需要 include<cuda_fp16.h> 以使用 CUDA 关于 fp16/half 的相关支持。然后使用 _hadd() 函数进行两个半精度数加法。此外还能够使用 __half2float() / __float2half() 在 half 与 float 之间进行转换。
一些类型转换函数会有不同的舍入模式版本，分别是 .rd 、.ru、.rn、.rz：
- RD：Round Down，往更小的方向舍入，比如 -2.5 舍入为 -3.0。
- RU：Round Up，往更大的方向舍入，比如 -2.5 舍入为-2。
- RN：Round-to-Nearest-Even，四舍五入，在中间时会舍入到最近的偶数，比如 -2.5 舍入为-2。
- RZ：Round Toward Zero，向零方向舍入，如-2.5 会舍入为-2.0。
使用半精度后，可以看到 half 确实是要比 float 要更快一些：
此外 half 在 CUDA 中支持内置半精度类型 half2，其是一个 32 位的数据结构，内部紧凑地打包了两个独立的 half 数值。 为了进一步加快速度，可以加法可以使用 __hadd2()，这个函数能够在一个时钟周期内用一条硬件指令同时完成两对半精度浮点数的加法，而之前的 __hadd() 只能够处理单个 half 类型的数据。
