本篇根据HOW CUDA PROGRAMMING WORKS的讲解，整理下如何更好地使用GPU的一些细节，主要有三点：

• 让GPU别闲着（Use it at all）
• 把GPU中所有资源都利用起来（Use all of it）
• 高效的利用资源（Use it efficiently）

在我们平常的项目中，除了优化 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有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使用的线程比之前增加了三分之二，这几乎是免费的性能提升。这个“俄罗斯方块”问题又增加了一个维度，但硬件会为你解决这个问题，让内核去帮你处理就行了。

公众号后台回复“数据集”获取100+深度学习各方向资源整理