[翻译]CUDA_C_Programming_Guide Performance Guidelines Maximize Utilization

5． GPU性能优化建议

5.1 性能优化策略概述

总体上说，性能优化主要包括以下三个策略：

• 最大化并行度以达到GPU 核的最大利用率。
• 内存优化以达到最大的内存带宽。
• 指令集优化以达到最大的指令带宽。

对于一个应用程序来讲，哪种优化策略能达到最好的性能取决于该应用程序的性能瓶颈在哪；比如说，对一段性能瓶颈在内存访问的代码进行指令使用方面的优化将不会得到很好的优化效果。所以，在性能优化之前，需要测量或者监控，识别出性能瓶颈，可以使用CUDA Profiler做性能瓶颈的分析。另外，还可以将应用程序的实际的浮点数运算吞吐量或者内存带宽与设备最大的理论值进行比较，从而估算出还有多大的优化空间。

5.2 最大利用率

为了得到设备的最大利用率，应用程序代码结构应该遵循以下原则：

• 尽可能多的线程并行。
• 并行的线程能够高效的加载到系统中的不同资源，以确保CUDA核心一直处于繁忙状态

5.2.1 应用级别优化

在应用层应该考虑最大化的并行执行，这些并行可以是在主机，GPU设备之间，主机到GPU设备传输数据时。一种可能的实现是调用asynchronous 函数和stream 机制。另外还要注意不同的处理器执行不同类型的工作，一种比较通用的做法是:串行部分在主机CPU端；并行部分在GPU设备端。

对于GPU设备中的并行部分，关注点应该在并行被终止的部分，比如线程间需要相互之间共享数据，从而在某一个时刻点执行线程同步操作。有两种情况：一种是这些线程属于同一个线程块，单个线程中使用__syncthreads()函数进行块内线程同步，从而可以同步在共享存储器中的共享数据。另一种情况是这些线程属于不同的线程块，他们通过全局存储器交互，而这种情况一般都会是在不同的kernel函数中通信，一个对共享数据写，一个对共享数据度。第二种情况因为增加了外部kernel函数请求和全局存储器访问，效率很低。可以通过尽可能的将需要线程间通信的计算部分设置在一个单独的线程块中执行，来最小化第二种情况带来的性能损耗，因为在同一个线程块中，可以通过共享存储器来进行线程同步，且损耗较小。

5.2.2 设备级别优化

GPU设备最大的优点就是核心特别多，并且能够完全并行执行。基于此特点，应用程序应该尽可能的设计成能够并行处理的结构。

5.2.3 多处理器级别优化

GPU按照warp为单位对线程进行调度，当一个warp中的所有线程准备好后，这个warp才会被执行，否则进入等待队列。GPU中驻留的warp数量直接决定了GPU的利用率。在每一个指令请求时刻，warp 调度器会选择一个已经准备好的warp执行下一条指令。一个warp从创建到能够被执行耗费的时钟周期称为”延迟”，如果所有的warp 调度器在每一个之中周期都有指令发送给一些warp执行，那么可以达到100%的利用率，也就是说，所有的延迟都被隐藏了。隐藏一个L 个时钟周期的延迟需要的指令的数量，取决于这些指令各自的吞吐量。

对于一个计算能力为3.x的设备来说，每个时钟周期发送8个指令，这8个指令分成4对分别发送4个不同的warps，每一个warp接收一对指令。

一个warp没有准备好执行的最大的可能是指令的输入操作数没有准备好，或者说是处理器无法读取到指令需要的输入数据，数据还没达到寄存器，或者指令之间的数据具有依赖性，即下一条指令需要等待上一条指令处理完才能执行。

如果指令的所有输入操作数都已经在寄存器中，那么这种情况导致延迟的原因就是数据的依赖性，比如一些输入操作数需要前面的指令重写，但是这些指令执行还未完成。这种情况称为back-to-back 寄存器依赖，它的延迟时间就等于前面指令的执行时间，在这段时间里，warp调度器必须把其它的warp调度起来。而指令执行时间的长短取决于指令本身，对于一个计算能力未3.x的设备，比较典型的执行时间是11个时钟周期，而在11个时钟周期内，warp调度器可以调度44个warp。当然这里有一个假设，就是程序具有足够大的指令级别并行度，以使得调度器总是有指令给每一个warp。

如果指令的输入操作数在片外存储器中，比如全局存储器，延迟将会达到200到400时钟周期。这种情况下，就需要有足够多的warp，来使得warp 调度器一直保持忙碌状态，需要的warp数量取决于kernel 代码和代码中指令并行的程度。这里有个名词叫做”程序的计算强度”，它表示非片外寄存器操作数的指令数量与片外寄存器操作数的指令数量之间的比值，(个人理解，这个比值就是kernel代码中计算部分需要的时钟周期与加载内存需要的时钟周期的比值)。这个比值越低，那么需要的warp数量就越大。举例说明，假设这个比值等于30，片外存储器访问的延迟是300时钟周期，如果该设备有一个时钟周期可以调度4个warp，那么就至少需要40个warp以保证warp调度器一直处于忙碌状态。

Warp未准备就绪的另一个常见原因是代码中遇到了线程同步点。线程同步点强制核处理器处于空闲状态，因为同一个块中的所有warp需要等待该块中的其它warp全部完成了同步前的所有指令。对每个多处理器分配更多的线程块有助于减轻这种影响，因为不同块之间的线程不受线程同步的影响，他们不必等待其它块内的线程。

每一个SM会分配到多少个线程块和warp，取决于kernel函数调用时线程的配置，以及SM的存储器资源(共享存储器，寄存器等)和kernel需要的资源。在编译时添加-ptxas-options=-v 选项，可以输出寄存器和共享存储器的使用情况报告。

一个线程块需要的共享存储大小等于代码静态和动态分配的共享内存。

Kerne函数中使用的寄存器的数量对SM中可以驻留的warp的数量有重大的影响。举个例子，kernel函数需要在一个计算能力时6.x的设备上运行，这个kernel函数使用了64个寄存器，同时每个块有512个线程，使用很少量的共享内存，那么一个SM上可以同时驻留两个线程块(以warp=32为单位,32个warp，1024个线程)，需要2x512x64个寄存器，这已经达到了设备的最大容量了。如果kernel函数要用65个寄存器，那么一个SM上只能运行一个线程块(以warp=32为单位，16个warp，512个线程)，因为两个 线程块需要的2x512x65个寄存器已经超出了设备的容量。通过添加maxregcount 编译选项可以控制寄存器使用，或者kernel函数加载时设置加载边界(launch bounds)

每个double和long long变量使用两个寄存器。

性能优化的效果都取决于kernel代码，需要不断的进行实验。程序中可以设置一些变量，这些变量关系到程序中对寄存器和共享内存的使用，通过调整这个变量，对不同设备进行性能优化。

每个线程块的线程数量应该都是warp大小的整数倍，这样可以避免由于warp除不断时造成计算资源的浪费。

5.2.3.1 占用量计算

CUDA提供了很多API来协议程序员确定线程块的大小，这些函数通过分析程序中需要的寄存器和共享内存资源来确定线程块大小。

• 占用量计算API

cudaOccupancyMaxActiveBlocksPerMultiprocessor，该函数基于kernel函数对线程块和共享内存的使用来预估线程占用量。返回值表示每个SM中可以同时运行的线程块的数量。

• 1. 这个值通过简单计算还可以得到其它有用的信息。乘以每个线程块的warp数量可以得到每个SM中同时运行warp数量。再将乘完之后的值除以每个SM能够同时运行的最大warp数量，就得到占用率。
• The occupancy-based launch configurator APIs,

cudaOccupancyMaxPotentialBlockSize and

cudaOccupancyMaxPotentialBlockSizeVariableSMem, heuristically calculate an execution configuration that achieves the maximum multiprocessor-level occupancy。(这一段不太理解，直接贴原文)

以下示例代码计算MyKernel函数的占用量，最后输出每个SM同时运行的warp数量与最大支持的warp数量的比值。

更多关于占用量计算的信息可以查阅CUDA的源代码：

<CUDA_Toolkit_Path>/include/cuda_occupancy.h