在GPU成本升高，日益短缺的情况下，如何提升GPU性能的方法？
本篇根据HOW CUDA PROGRAMMING WORKS的讲解，整理下如何更好地使用GPU的一些细节，主要有三点：
- 让GPU别闲着（Use it at all）
- 把GPU中所有资源都利用起来（Use all of it）
- 高效的利用资源（Use it efficiently）
充分利用GPU资源
在我们平常的项目中，除了优化 kernel 性能，能够快速拿到收益的是整体 pipeline 的优化。
我遇到大部分项目在优化性能的时候，kernel 优化一般都是最后一步，很多项目一开始可以先从整体 pipeline 上考虑，拿到可以拿到的性能收益，最后实在扣不出来再考虑优化 kernel。
让GPU别闲着
首先明确一点，GPU是异步的，一般cpu发完命令后（也就是launch了一个 kernel 之后）就没事了。
发kernel指令这一步很快，理论上任务给了gpu之后cpu就能干自己的事儿了。有异步就会有同步，依赖gpu处理结果的cpu线程需要同步操作来等待当前gpu处理完才能进行下一步，比如处理图片：
但是上述gpu利用率并不高（看右面的曲线），这里等gpu执行完的时候其实cpu啥也没干，gpu和cpu之前的数据拷贝需要时间，然后gpu处理完了cpu会做一些后续操作，最后才会处理下一个图片，再加上一些语言层面的 overhead，GPU的利用率就很低了。
其实我们可以把所有的任务一股脑扔给GPU，因为GPU是异步的，cpu把任务传递过去可以通过同步来获取gpu处理完的数据。传递过来的kernel任务会被GPU排入队列一个一个去执行，这个时候GPU的利用率就会高些，毕竟相对没有那么闲了：
这就是最基本的让GPU别闲下来的通用办法，尽量做到：
- 减少不必要的 CPU-GPU synchronization
- 在CPU-GPU之间拷贝数据时也可以执行Kernel
- 很多任务，能通过异步提升性能就尽量异步
将 GPU 中所有资源都利用起来
当我们拿到一张显卡或者说计算卡后，对算力敏感的人一般都会比较计较这个卡的计算资源，比如A100： ( 忽略我买不起a100的事实 )
A100有108个SM（流处理单元），一共有221184个线程，每个SM都有自己的寄存器、共享显存，共享L2缓存等，我们要做的就是尽可能充分利用这些资源。如果我们某一时刻只用其中一个线程，然后一直调用（kernel一直跑），此时虽然用nvidia-smi命令查出来的GPU utilization是100%，但其实对着张显卡的使用量超级低。
比如这个单block单thread的kernel，理论上只会占用一个SM：
__global__ void simple_kernel() {
while (true) {}
}
int main() {
simple_kernel<<<1, 1>>>();
cudaDeviceSynchronize();
}但通过nvidia-smi指令得到的GPU Util是100%，这个是要注意的。
GPU如何工作
让我们先快速了解一下Cuda是如何将kernel算法分解成一块一块的，以便将它们分布到这些SM上的。
首先我们把数据分成大小相等的块，这样就可以并行独立地运行每个块。因为这些块是独立的，它们可以按任意顺序和任何时间安排。这样硬件就有尽可能大的自由度来高效地打包事物。每个块的一个保证就是它有一定数量的线程，并且它们都保证在同一个SM上同时运行。
接下来，我们要讨论的是如何使用这些 block 中的线程。每个 block 中有一个特定数量的线程，它们都被分配到同一个 SM 中。SM 中的线程会被分成一定数量的组，这些组被称为 warps。每个 warp 包含特定数量的线程（通常是 32 个），并且这些线程会在同一时钟周期内被执行。这就是为什么我们要使用大量线程的一个原因——保持 GPU 中的所有 SM 和线程都尽可能繁忙。
现在我们回到刚才说的用满。
当一个 SM 上运行的所有 block 都已经满了，就不会再有新的 block 被添加到该 SM 上，直到该 SM 上的线程开始完成它们的工作，空出一些空间。在CUDA中，每个流多处理器（SM）可以同时处理多个块，具体数量取决于块的大小，最多可以处理32个。GPU会持续调度块到各个SM上，直到所有块都被处理完毕。
这就是所谓的“occupancy”，它是一个重要的 GPU 性能指标。
我们先看一下如何编写代码来利用这些 GPU 资源。CUDA 提供了一种编程模型，称为 kernel 函数，它是在 GPU 上执行的代码。这些函数可以通过在主机上调用它们来启动 GPU 上的计算。kernel 函数的一个重要特点是，它可以被调用多次，每次使用不同的数据块。这种方式可以使得 GPU 尽可能高效地利用其并行计算资源。
为了使代码能够在 GPU 上运行，我们需要在主机代码中使用一些特殊的关键字和函数，以便让 CUDA 编译器将其转换为 GPU 可以执行的代码：
注意看相关的硬件资源和代码对应的地方，这个kernel计算多个点之间的距离来展示并行计算。程序中使用了共享内存和寄存器，共享内存允许块内的线程进行通信，而寄存器则为每个线程提供必要的工作空间。在GPU上，寄存器的用途与CPU不同，它们提供了大量立即可访问的数据空间，这对于执行复杂的数学运算尤其重要：
在GPU编程中，为每个线程分配大量寄存器是常见的，这与CPU上的情况有显著不同。首先假设当前写的kernel需要使用的资源如下（下图右侧），每个block是256个线程、每个线程使用64个寄存器，每个block的共享显存是48KB：
然后A100的一个SM的资源如下，我们写的这个kernel可以安排三个block到这个SM上，再多了共享显存不够放了：
如果我们把算法重构一下，减少shared memory的使用，比如从48K→32K，那么就可以塞4个block到这个SM上了，这个时候occupancy也就提升了：
所以为了保持GPU的忙碌状态，我们可以进一步尝试填补空闲时间。实际上，GPU可以同时运行多个程序，最多可以同时运行128个不同的程序。因此，如果我们没有完全填满GPU，就可以尝试将另一个程序“俄罗斯方块”式地放置在旁边。
比如我们想象另一个case，这次是一个大量线程、很少寄存器和没有共享内存的情况。实际上就是一个数据移动kernel的典型表示，可能是对数据进行排序或只是将数据从一个地方复制到另一个地方。
GPU很聪明，它会尝试将一个绿色网格块填充到空隙中，如果不能再放置绿色网格块，那么就会尝试另一个。
现在我们每个SM使用的线程比之前增加了三分之二，这几乎是免费的性能提升。这个“俄罗斯方块”问题又增加了一个维度，但硬件会为你解决这个问题，让内核去帮你处理就行了。
刚才聊到了如何充分利用GPU的计算资源（SM），这里我们再从另一个角度看下如何提升性能。个人水平有限，如文章有误，欢迎评论区讨论.
之前也提到，我们可以将所有要做的任务都一股脑扔给GPU，换句话说，就是可以将要计算的算法提前交到异步工作队列中（为什么是提前，就是这个任务发射出去了，kernel launch了，但实际还没有执行），这个队列一般叫做stream，如果没有设置的话，任务会处在default stream中。
GPU可以有很多stream，分任务不同使用的stream数量也不同。我们在平常使用的时候可以充分利用stream的特性：
- A stream is a queue of device work
- Kernel launches and Async CUDA API calls place work in the queue and continue
- Devices schedule work from streams to available resources
- Operations within a stream are ordered (FIFO) and cannot overlap
- Operations in different streams are unordered and can overlap
这里我们有个任务是用神经网络处理下花朵的图片，这会涉及到将一朵花或花的数据复制到GPU的内存中，然后在GPU上进行处理和处理，最后将结果复制回来。
因为处理不仅仅是一个步骤，而是多个步骤，这给了我们重叠的机会，这意味着我们可以将独立的工作分布到独立的流中，处理第二朵花实际上并不需要等待第一朵花完成。我们的copy操作（拷贝数据到gpu上），也就是数据传输，可以和计算（kernel执行）在不同的硬件路径上同时执行，这样就可以高效利用硬件的资源。
上述绿色流中的工作可以和蓝色流中的工作一起占用资源，让GPU保持最大限度的繁忙状态。其实这个很像刚才提到的俄罗斯方块，这里的方块指各种计算和内存拷贝任务，然后我们尽可能将多的任务填充到可以利用的资源当中。
目前为止，我们已经让GPU尽可能忙起来而且也尽可能利用其资源，接下来就是最难的部分：高效的利用资源（Use it efficiently）
再回顾下A100的计算资源：
看算力，将近20T的FP32 FLOPS，这还没算tensor core。我们的kernel如果能用到这些算力的80%那也是很牛逼了，不过可惜FLOPS不是最关键的，关键的是显存带宽（Memory Bandwidth），很多时候我们遇到的不是计算上限而是访存上限。
虽然A100 GPU拥有108个SM，每个SM在1410 MHz的时钟速度下每个时钟周期能请求64个字节的内存，从而使整个GPU的内存请求速率接近每秒10TB。然而，这个高速的内存请求与GPU的内存供应能力之间存在不匹配。尽管A100配备了很快的HBM2内存，每秒能提供超过1.5 TB的内存，这个供应量仍然只是SM请求量的六分之一。因此GPU的性能受到内存带宽的限制，无法完全满足所有SM的内存请求需求，所以大部分计算不是调用不了线程，而是内存性能限制。
如果gpu能提供的线程数我们没有都用到，那显然就是浪费，就是没有好好利用gpu。
以内存角度考虑，如果仅考虑数据输入速率，尽管这不完全代表真实情况（因为会有一些数据重用），但1.5 TB每秒的数据传输率是一个很好的性能指标。这相当于每秒可以处理1940亿个双精度浮点数值，除以8，即得到194 GFLOPS的峰值性能，GPU的运行速度现在大约与iPhone 11相当。这个例子说明了在进行性能评估时，应该全面考虑不仅仅是数据传输速率，还有程序实际利用这些数据的能力，以及内存带宽如何成为性能瓶颈的一个示例：
我们可以看下内存的工作原理，在DRAM存储芯片中，每一个存储单元都是一个电容，用来存储一个比特的信息——充电状态表示1，无电表示0。
这些电容通过晶体管接入电路，晶体管的开关控制着是否可以读取电容的电荷状态。这些电容和晶体管组成的存储单元被连接成一个二维矩阵，这种布局使得DRAM成为随机存取存储器（RAM），与顺序访问的磁带存储不同，DRAM可以通过行和列的索引寻址访问数据。
当进行数据读取时，首先激活选择的行，该行中所有单元的状态被复制到感应放大器中。
感应放大器将电容中微弱的电荷转换成可轻松读取的明确定义的电压信号，便于读取。然而，读取过程中会消耗掉电容中的电荷，导致原数据丢失。
接下来进行列访问，记住这是一个2D矩阵，必须先读行再读列。与从内存单元读取数据相比，它可以读取感应放大器中保存的数据。这比行更快，因为放大器产生了强有力、清晰的信号，并且读取不会破坏数据。
由于经常读取相邻的内存位置，所以也有一个连续读取模式，其中单个请求返回多个数据字。这是非常重要的，因为这意味着我不必一遍又一遍地付出列请求的代价。这也是cuda编程中一个通常的优化点之一。
这就是DRAM的突发模式，即一次请求可以返回多个数据字，这样可以减少列请求的成本。这种读取模式的设计反映了物理存储结构的影响——即DRAM单元的充放电特性。由于行变更的成本很高，优化数据访问模式，尽量从同一行读取相邻的数据值，可以显著提高读取效率。这种物理和技术的细节对于理解和优化内存访问模式至关重要，尤其是在执行内存密集型操作如CUDA编程时更是如此。
但是当需要访问另一行的时候，必须先将当前放大器中的数据重新写回到DRAM行中，以避免数据损坏（还记得当数据移动到放大器时，行已经被耗尽了，所以我们必须重写它以避免内存损坏）。
这种所谓的“switching page”操作涉及数据回写和新行加载，是一个成本较高的操作，通常切换页的代价是切换页内列的三倍：
让我们进行一个实验，我们要做的就是读取八个字节（8 byte）的值，也就是我们一直在讨论的双精度浮点型（double）。通过逐渐增加步长来读取数据。步长为8表示读取相邻的值，步长为16表示读取每隔一个值的数据，以此类推，直到步长为1000。
上图中，将 x 轴保留为字节单位，可以看到页面大小是 1K，因为当步长从 512 增加到 1024 时，我现在每页只能得到一个值，我们可以看到这里的下降。您还可以看到，仅仅通过从左侧相邻读取变为每隔一个值读取，带宽减半。但我最感兴趣的是它的最低点，因为这个最小带宽告诉我，如果不注意相邻内存访问，我可能会看到多大的影响。
实际上，当步长变长时，我们在读取随机位置，从而不断地打开和关闭页面，并发出单独的读取命令，这样会导致有效带宽下降了92%。所以，如果我实际的性能主要取决于内存性能，正如我之前所论证的那样，错误的内存访问模式会很影响性能。
A100 的宣称内存带宽是 1555 GB/s，这里只测到 1418 GB/s，因为这里读取的是 8 字节值，如果是读取16字节的话这个带宽会更高些。如果我们读取的数据块更大，那么也就会更高效。
数据访问模式确实非常重要，这个优化前提在CPU优化也很常见，但是在 GPU 上尤其严重，因为内存延迟时间更长
让我们来看一个非常常见的数据结构，一个 2D数组，也就是我们常见的图像（h,w,1）。通常以行优先格式布局，其中相邻的内存位置在 X 方向上递增。这里如果读取使用的步长每次增加都会降低性能，同样如果改为沿着一列而不是一长排进行访问，性能则会大幅下降，就和上面讲到的步长曲线中一样，直接下降了一个数量级。
那么为什么我们要谈论所有这些呢？
因为提升性能的主要因素就是高效利用GPU所有的资源，利用好可以提升10倍以上的性能，还是很值得。而我们通常使用cuda去编写kernel，这也是导致cuda编程模型成为现在这样的重要因素，所以让我们看一下cuda编程模型如何运行并行工作。
cuda将任务分解为独立block，其中每个block由大量thread一起运行。
cuda的执行基本单位是线程块，而不是线程，block实际上是并行的一种量子单位，这意味一定可以保证获得并行性，我们可以确定块中的所有线程同时运行，因此它们可以共同工作和交换数据。
cuda真正呈现线程是独立的，就像单线程程序一样，每个线程都有自己的所有变量副本和自己在代码中的程序计数器。我们怎么调用这个线程，在于两个名为“thread idx”和“block idx”的变量，它们在启动块时由硬件自动填充，以给每个线程在网格中分配索引。
唯一重要的是第一行，其中每个线程使用其线程索引和块索引来计算其正在处理的数据位。
之后，程序中的所有内容在所有线程中都是相同的。
SIMD和SIMT之间的区别是什么，GPU是SIMI模型，对于SIMT线程控制是隐含的，每个线程都有自己的状态，所有线程都是独立的，即使它们是矢量单元的一部分，它们都会保持自己的状态，包括这个索引变量。它们可以循环、分支、做条件性的事情，就像这个例子中一样，当它们一起执行相同的任务时，它们更有效率，因为当事情一起运行时更有效率，GPU按 32 个一组运行它们，我们称之为 warp。
在上述的代码中，每个线程加载两个数据点，因为它在寻找两个点 P1 和 P2 之间的距离，这由一个非常重要的线程 ID 依赖的行所索引。这实际上意味着我一次从 32 个不同的线程加载数据，偏移量由它们的线程 ID 的某个函数指示。
稍后会解释这为什么重要，但我们需要知道 float2 只是两个浮点数的结构体，每个大小为八字节的 x 和 y，它是一个二维坐标。
所以每个 warp 正在加载 256 个连续的字节数据，因为它的 32 个线程中的每一个都获取一个值，每个线程都由线程 ID 索引。线程 ID 是保证连续的，这意味着数据访问模式保证是连续的。
GPU 上的 SM 能同时管理 64 个warp，总共是每个 SM 有2048 个线程。但是你可以从 SM 的块图中看到，它实际上有四个独立的部分，所以任何时候它实际上同时运行的是四个 warp，其他的则保留在队列中，并且被切换进出。
所以我有四个 warp，每个同时加载 256 字节，四倍的 256 字节即 1024 字节，是从内存系统中同时加载的完全合并的相邻数据地址，如果你还记得，这正是充分利用我的内存系统的完美大小。
这非常重要，因为程序在各处读取数据，即使我的简单示例也是如此，它从两个不同的点数组 P1 和 P2 读取，对计算机来说，看起来我的程序在做随机内存读取，它从一个数组读取，然后在内存中完全切换页面来读取另一个。
正如我们看到的，这可能会让我损失 90% 的内存性能。这是我最关心的单一最重要的事情。但实际上，由于 warp 执行模型，因为我们同时运行四个，我的大规模并行 GPU 恰好拥有正确的数据量，恰好能够达到我的内存系统的峰值带宽，即使我的程序从各处读取数据。
如果你是一个 cuda程序员，而且你曾经想知道为什么 warp 是 32 个线程，这就是原因之一，这不是偶然的。硬件设计师精心平衡 GPU 以确切地做到这一点，而且每当有人问我他们的线程块应该有多大时，答案总是不小于 128 个线程，因为你总是想要四个 warp 一起工作，以便轻松达到内存性能的峰值。这里有一个 13 的因素在起作用，一个 warp 已经很好了，所以在你的程序中努力让所有线程在 warp 中从相邻线程读取，但你需要块中不止一个 warp 才能使一切完美工作。
总结
让我们回顾整体的内容。
对于GPU，我们需要使用异步工作提交来保持GPU的繁忙状态。我们了解到，硬件会将块尽可能广泛地分布在SM上，以最大限度地请求内存带宽。我们还学习到，硬件可以玩出相当不错的四维俄罗斯方块，这对于资源打包限制至关重要，因为资源利用率的微小变化可能会产生巨大的影响。即使你使用了GPU提供的所有线程，我们还是会遇到一个问题，即内存系统只能提供SM请求带宽的1/6，所以我们的主要限制因素是内存。
接着我们看到，内存系统极度依赖良好的访问模式，电容定律决定随机读取内存的代价要高于线性读取，如果我们弄错了，我们可能会将一切变慢，这可能导致性能降低高达92%。
但我们了解到，GPU的Warp执行系统可以通过同时从许多线程读取来节省我们的开销，但我们需要Warp中的所有线程都在处理相邻的数据，所有的线程块最好都有至少128个线程，这就是Cuda编程的方式，主要是为了让GPU保持繁忙状态，尽可能地利用它，并最重要的是小心地读写内存。
这里有很多技术细节，但实际上只有三个概念我们需要记住：
- 异步执行
- 平衡资源
- 保持内存访问的线性
几乎所有的优化内容都包含在这些主题中，如果我们能正确地处理这三个问题，就没有太多可以优化的内容了。重要的是，我们在设计算法和编写程序代码时，从一开始就考虑这些问题是最好的。CUDA的设计试图迫使你思考这些问题，因为这是最好的方法，它可以使结果产生数量级的差异。
