CUDA优化冷知识23|如何执行配置优化以及对性能调优的影响
CUDA优化冷知识23|如何执行配置优化以及对性能调优的影响
这一系列文章面向CUDA开发者来解读《CUDA C Best Practices Guide》 （CUDA C最佳实践指南）
我们今天主要进行<CUDA Best Practices Guide>的章节10的剩余内容https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#occupancy，
也就是接上一篇的occupancy后面，继续说说寄存器的延迟掩盖，blocks形状和使用，shared memory的使用，以及，concurrent kernels和CUDA Context等方面，对性能调优的影响。
首先我们从寄存器的延迟掩盖开始。本小结首先讲述了，当需要使用寄存器中的数据，而该数据没有准备好的时候，从而无法取得数据喂给SM中的执行单元，从而可能导致执行的线程被卡住(stall)而不能就绪执行的状态。小结只讲述了常见的A = XXX; 这种形式的寄存器上的结果计算延迟。并用volta举例常规的计算有4个周期的延迟，在此期间内，立刻使用结果数据是不可以的，需要等待4个周期才可以。并讲述了可以临时切换到其他warps中的指令继续执行来掩盖的方式。本小结是乐观的，认为这一般不构成对性能的影响。
但是实际上，随着现在nsight compute的流行，long/short scoreboard的stall reason之类的分析指标的公开，很多操作对寄存器的结果写入，可能要超过这例子中的4个周期不少。
我们这里只额外说一下，使用s_xxx[idx] = d_xxx[idx]形式的，从global memory看似'一步到位'写入到shared memory的做法。实际上会被编译成中间的分步的tmp = d_xxx[idx]; s_xxx[idx] = tmp; 的经过寄存器(tmp)的分解过程，导致中间第二次写入的时候有一次对寄存器的依赖。使用8.6和8.7计算能力的人们，建议考虑新版的cuda::memcpy_async的载入方式，这种可以直接越过寄存器。
这是今天的第一小节。
第二小节讨论了block和grid的形状对性能的影响问题。这个是个喜闻乐见的讨论，在我们夏令营和冬令营的活动中，被人讨论了无数次了。小节首先澄清了，grid和block的1D还是2D还是3D的形状，从本质上并不影响性能，影响性能的只是无论1D还是到3D时候的，计算出来的每个block里的线程总数量，和blocks的总数量。
小节同时说明了，这些线程和blocks的数量（和其他资源），影响了在SM上的active warps的数量。能达到的active warps数量，才是之前的occupancy之类的很重要的原因。而active warps的数量，往往决定了延迟掩盖，和对SM各个单元的利用程度。这样性能就取决于这些单元的利用率情况，因为一旦我们买回来了一张卡，硬件的SM数量，和SM里面的执行单元配置是固定死的了，硬件本身乘以利用率，才会影响最终的性能发挥。
然后小节往下说了，该如何调整kernel启动时候的方括号里的第一个和第二个参数。大部分情况下，调优kernel，需要同时(in tandem)试验性的调整这两个参数。但每个参数也有他们自己的调整策略：
对于第一个参数(blocks数量): 基本的策略是要足够多，至少每个SM上得有1个block。同时，考虑到了1个SM上如果只有1个block的话，一旦该block中的线程们，执行了__syncthreads()进行等待同步的话，很可能导致SM上warps大部分都处于等待状态了，降低该SM的使用率。所以这个至少的1个block还需要调更多。手册的建议是，亲这边应该至少上几千个blocks每张卡。理由很简单：考虑到现在的8.6的3090的卡，有82个SM。每个SM上可以上到多达16个blocks，这样82 * 16等于差不多1000。几千个差不多能将一张卡上个几批次。手册说到，我们要面向未来考虑，将来的卡更强。所以数量不能保守。
阅读到这里，我们应当结合实际一点。因为随着block对资源的使用不同（例如shared memory), 一个批次能上多少个blocks，对于固定的卡，随着kernel的不同是不同的。建议读者使用nsight compute, 观察里面特定kernel的waves数量指标，该指标说明了某kernel的blocks需要分成几个wave（批次），才能上完。
以及，对于某些因为算法的角度的限制，不能有效扩大blocks数量的情况下，针对本章节讨论到的，因为__syncthreads()而导致1给block中的warps在SM上整体stall的问题。可以考虑使用细粒度的部分同步手段。也就是使用cuda::barrier（需要计算能力7.0+），进行1个block中的部分线程进行同步。这样当部分线程在wait()或者arrive_and_wait()进行同步的话。该block中的其他不参与barrier同步的线程依然有机会执行，继续利用SM上的执行单元。
以及，新版本的上一部分手册(CUDA Programming Guide), 现在已经正式引入了很多C++风格的东西了。上一段说到的asynchronous barrier, 在当年我们阅读编程指南的时候，没有涉及。建议读者重新阅读相关章节。
然后继续回到<<<>>>的第二个参数，也就是block中的线程数量的优化考虑。手册这里主要考虑了你不能用过小的blocks，例如只有32个线程的block. 因为SM往往还有例如16个block/SM的硬限制。使用过小的block往往会导致SM上去的总warps数量不足，可能会影响性能。手册这里建议的方式是，至少上64个线程的block，然后逐步调整block中的线程数量, 找到特定kernel的最佳性能点。这个逐步调整，可以从128或者256个线程起步。
手册继续说，调整到适可而止就行了，没必要追求极限。例如通过调整前两个参数，让SM能上到66%的occupancy，和能上到100%的occupancy，可能并不会对性能起到太显著的影响。因为调整的目的是追求性能，而不是单纯追求指标。为了得到过高的occupancy，有的时候你只能降低寄存器数量之类的，从而导致使用了过多的local memory, 反而影响性能。
而另外一方面，因为除了我们之前说过的TLP（例如依靠切换warps）来充分利用硬件的执行单元，还存在ILP的方式，也就是线程内部的前后指令本身的并行性，来提高效率。手册这里指出了，只要内部的ILP程度足够，哪怕较低的occupancy也是足够的。对于这个问题，我们建议读者继续扩展阅读经典文章：《Better performance at lower occupancy》（链接： http://dmacssite.github.io/materials/volkov10-GTC.pdf ），该文章描述了哪怕很低的occupancy，也可以通过ILP取得优异性能的方式。虽然这个文章较老，但是依然非常经典。
另外的，我们夏天搞夏令营活动的时候，客串出场的樊博士，也在他的实践中(GPUMD项目），指出了这点，例如在他的《Efficient molecular dynamics simulations with many-body potentials on GPU》中，老樊写道：“哪怕使用float的时候只有50%的occupancy；或者使用double的时候只能到25%的occupancy。性能也相当不错"。（arvix: https://arxiv.org/abs/1610.03343 ), 感兴趣的读者也可以扩展阅读。
这两篇文章都分别有12年和5年的历史了，但是里面的思想，是正确和不过时的。
（这里推荐一下樊博士写的CUD编程书籍，也是NVIDIA CUDA夏令营/冬令营推荐参考书籍）
此外，追求调整occupancy的时候，如果是寄存器受限，可以考虑调整-maxrregcount参数来调整常规寄存器数量（CUDA通用的predicate register，和7.5+的标量的uniform register是固定的8个和64个，不可调）。如果在实践中，发现单一.cu文件中存放了多个kernel, 不能统一用maxrregcount参数调整的话，也可以上__launch_bounds__针对特定kernel单独调整。
下一小节手册谈论了shared memory的使用对性能的影响问题。主要提及了，shared memory有助于global memory上的合并访存、消除global的重复访问、和block内部的数据交换等方面的用途。并指出了虽然这些用途很有用，但有的时候需要做出取舍，一个劲的使用shared memory不一定总是能有正面效果。
对于这里的shared memory对global memory的读写合并访存上的帮助，我们搞夏令营活动的时候，已经给大家演示过了嵌入式jetson设备上，消除读取或者写入时候的不合并情况，对性能带来了有效的提升。但是在现在的逐代更新的台式卡，随着各级cache的扩大，这种效应在递减，例如我们的老樊在他的github上的链接（ github.com/brucefan1983 ）指出，较新卡在进行矩阵转置的例子的时候，哪怕不合并的读取或者写入，因为cache的效应，哪怕不使用shared memory, 很多情况下问题也不太大。所以，如果当优化的时候，shared memory的资源使用，成为了限制因素的话，该情况下Shared memory也可以减少使用。
下面老樊的图: （较新代数的卡上的不合并访存的效果弱化演示）
然后对于block内部的数据交换，读者如果能够将范围细化到每个warp内部级别的话，可以考虑上shuffle操作。该操作可以将数据缓存在寄存器内，从而减少了了对Shared memory的使用率。感兴趣的读者可以参考次链接进行进一步的扩展阅读：NVIDIA: 《Register Cache: Caching for Warp-Centric CUDA Programs》（ Register Cache: Caching for Warp-Centric CUDA Programs | NVIDIA Technical Blog ）。这些都有助于你的性能优化。
以及，如果在特定的计算能力的卡（8.6）上的话，因为shared memory会强制的被自动额外占用1KB/block, 这些结合block/线程形状+shared使用量这两个小结的整体内容，你可能在8.6上不能使用过小的block，一面导致无辜的额外资源占用。
此外，本小结还提出了，很多的思路往往喜欢1个线程对应1个Shared memory中的数据，这样如果shared memory种有一个32x32的矩阵，上一个(32,32)的block还勉强凑合。但是如果shared memory中有(64,64)的矩阵的话，上(64,64)的block将不可能（超过了1024个每个block中的线程数量限制）。此时可以考虑每个线程计算多个数据。这样不仅仅shared memory和block中的线程形状这两点结合了起来，额外每个线程计算的多个数据还有助于ILP的进行，这样shared, 线程数量，ILP三者就结合了起来了。
手册继续将讨论每个GPU上多个小的并发kernels，和考虑有无MPS存在的情况下，上多个CUDA Contexts对性能的影响。
首先手册讨论了多个并发kernels。这点在我们阅读本手册最开头的APOD原则的时候，也就是对现有的CPU项目逐个热点的发掘，并移植到GPU上加速的过程中，往往很有帮助。
一个需要加速的老程序的多个方面，往往在应用该原则被改成到GPU的过程中，不同的代码片段往往会被实现成为多个kernel。这些kernel本身，如果单一来看压榨不出来足够的并行性。则可以考虑通过concurrent kernels特性，使用多个无关的流，来并发的启动他们，规避单一kernel无法充分利用GPU硬件的情况。这样，不仅仅多流对于我们之前说过的计算--传输的并发上有帮助，在计算---计算的并发上，也对性能有帮助，这也是优化的过程中需要考虑的一点。
关于这点，和下面即将谈论到的multiple cuda contexts，我建议读者阅读《Characterizing Concurrency Mechanisms for NVIDIA GPUs under Deep Learning Workloads》（arvix: https://arxiv.org/abs/2110.00459 )。不要被这篇文章的标题所迷惑（深度学习），这里谈论到streams和contexts的各种情况，以及对MPS的性能测试分析。
好了，回到今天的最后一小节，手册讨论了到多CUDA Context的使用。这里主要有两点用途，一点是CUDA Context在Driver API和Runtime API混合调用时候的帮助。我们知道runtime api是没有context这个概念的，而driver api有。同时runtime api稍微易用点，而driver api稍微难用点。而很多代码，例如NV的Video Codec SDK的例子中，很多代码使用的driver api进行的。则本小节指出了，可以通过特殊的primary context的概念，来进行和runtime api的交互。注意，这点并不能直接提高性能，但是交互操作，能让你切换到使用runtime api，从而节省了你的时间。你节省的时间可以用来优化成本，或者花费到优化代码的其他方面，来提升性能。例如这里举例的Video Code SDK Samples代码，就可以直接方面的改成简单版本的runtime api版的，能节省很多的开发时间。
此外, primary context不是必须的，同时和常规context比较起来比较奇怪（例如只能用引用计数方式自动被创建和使用、销毁，而不能手工创建销毁）。但是实际上你也可以创建普通的CUDA Context来在Driver和Runtime API之间交互，这并没有问题。这个只是看起来比较奇怪而已，实际上依然是一个普通context。
最后，今天手册谈论了MPS的情况。指出了多个CUDA Context的并行问题。例如上一节的多流的concurrent kernels, 只能在1个Context内部真正并行。如果有kernels同时存在，在不同的contexts中，则他们不能真正并行，必须通过MPS才可以，否则只能一个context中的kernels暂时切换到显存保存状态，然后另外一个context中的kernels再从显存切换回来继续执行，手册这里称为time-slicing.
而在有MPS存在的情况下，则这些context会被合并成1个真正的context，消灭了时间片轮换的代价。所以我们建议优化的时候，如果可能，尽量使用单一CUDA Context, 如果不能避免使用多个CUDA Context(例如你在调用一个第三方的库，你不能安心的将你的context交给他，万一他内部有BUG，可能你的代码会被连累）。此时可以可以考虑上MPS来提高性能。关于MPS，Streams的各点评测，可以看上面的文章。此外，上面的文章也少见的提到了Priority Streams, 不同优先级的流中的任务调度对性能的影响（手册今天没涉及），感兴趣的读者也可以看一下。