第5章 性能指南

5.1 指令性能

要处理一个warp的线程的指令，多处理器必须：
 读取warp的每个线程的指令操作数，
 执行指令，
 写入warp的每个线程的结果。
因此，有效的指令吞吐量取决于名义指令吞吐量以及内存延迟和带宽。它通过下列方式最大化：
 最小化具有低吞吐量的指令的使用（参见5.1.1），
 最大化每种内存的可用内存带宽（参见5.1.2），
 允许线程调度器尽可能地将内存事务与数学计算重叠，这需要：
 由线程执行的程序具有高的算术密度，也就是说，每个内存操作具有高的算术操作数；
 每个多处理器具有许多活动线程，详见5.2。

5.1.1 指令吞吐量

5.1.1.1 算术指令

要执行warp的一个指令，多处理需要：
 4个时钟周期，用于浮点加、浮点乘、浮点乘-加、整数加、位操作、比较、求最小、求最大、类型转换指令；
 16个时钟周期，用于倒数、倒数平方根、__log(x)（参见表B-2）。
32位整数乘法使用16个时钟周期，但__mul24和__umul24（参见附录B）提供了4个时钟周期的有符号和无符号24位整数乘法。但是，在将来的架构中，__[u]mul24将比32位整数乘法慢，所以我们建议提供两个内核由应用程序相应地调用，其中一个使用__[u]mul24，另一个使用一般的32位整数乘法。

整数除法和模数操作特别昂贵，如果可能的话应该尽可能地避免，或者尽量替换为位操作：如果n是2的幂，则(i/n)等于(i>>log2(n))，(i%n)等于(i&(n-1))；如果n是文本，则编译器将执行这些转换。
其他函数使用更多时钟周期，因为它们实现为多个指令的组合。
浮点平方根实现为倒数平方根与求倒，而非倒数平方根与乘法，所以它对于0和无穷大获得正确的结果。因此，它对于warp使用32个时钟周期。
浮点除法使用36个时钟周期，但__fdividef(x, y)提供了更快的版本，即20个时钟周期（参见附录B）。
__sin(x)、__cos(x)、__exp(x)使用32个时钟周期。
有时候，编译器必须插入转换指令，从而引入附加的执行周期。这种情况包括：
 操作在其操作数通常需要转换为int的char或short上的函数，
 用作单精度浮点计算输入的双精度浮点常量（不使用任何类型后缀定义），
 用作表B-1中定义的数学函数的双精度版本的输入参数的单精度浮点变量。
最后两种情况可以通过下列方式避免：
 单精度浮点常量，使用f后缀定义，比如3.141592653589793f、1.0f、0.5f，
 数学函数的单精度版本，也使用f后缀定义，比如sinf()、logf()、expf()。
 对于单精度代码，我们强烈建议使用浮点类型和单精度数学函数。当在不支持原生双精度的设备（比如计算能力1.x的设备）上编译时，双精度类型默认降级为浮点数，双精度数学函数映射为其单精度对应值。但是，如果将来这些设备将支持双精度，则这些函数将映射为双精度实现。

5.1.1.2 控制流指令

任何流控制指令(if, switch, do, for, while)通过导致同一warp的线程分散，也就是说，按照不同的执行路径执行，可以显著影响有效的指令吞吐量。如果这种情况发生，则不同的执行路径必须序列化，增加此warp执行的指令总数。当所有不同的执行路径都已完成时，线程将集中到同一执行路径。
当控制流取决于线程ID时，要获得最佳性能，就应该写入控制条件，以便最小化分散的warp数。这可能是因为warp在块中的分布是确定性的，参见3.2。一个小示例是当控制条件仅取决于(threadIdx / WSIZE)，其中WSIZE是warp大小。在这种情况下，没有任何warp会分散，因为控制条件与warp已完美对齐。

有时，编译器可以通过使用分支预测展开循环或可以优化if 或switch语句，详细说明如下。在这些情况下， warp绝不会分散。程序员也可以使用#pragma unroll指令控制循环展开（参见4.2.5.2）。
当使用分支预测时，其执行取决于控制条件的指令没有一个将跳过。相反，每个指令都与基于控制条件设置为真或假的每线程条件代码或谓词相关，虽然其中每个指令都会调度执行，但只有具有真谓词的指令将实际执行。具有假谓词的指令不写入结果，而且不求地址或读取操作数。
仅当由分支条件控制的指令数小于或等于特定临界值时，编译器才将分支指令替换为谓词指令；如果编译器确定条件可能生成许多分散的warp，则此临界值是7，否则是4。

5.1.1.3 内存指令

内存指令包括从共享或全局内存中读取或写入的任何指令。多处理器使用4个时钟周期来执行warp的一个内存指令。此外，当访问全局内存时，还有400到600个时钟周期的内存延迟。
例如，下列示例代码中的赋值操作符

使用4个时钟周期执行从全局内存中的读取，使用4个时钟周期执行到共享内存的写入，但最重要的是使用400到600个时钟周期从全局内存中读取浮点数。
如果在等待全局内存访问完成时，可以执行足够的独立算术指令，则此全局内存延迟的大部分可以由线程调度器隐藏。

5.1.1.4 同步指令

如果没有任何线程必须等待其他任何线程，则__syncthreads将使用4个时钟周期执行warp。

5.1.2 内存带宽

每个内存空间的有效带宽主要取决于内存访问模式，详见下列小节。
因为设备内存与芯片上内存相比具有较高的延迟和较低的带宽，所以设备内存访问必须最小化。典型的编程模式是将来自设备内存的数据存储到共享内存中；换句话说，就是让块的每个线程：
 将设备内存中的数据加载到共享内存中，
 与块的所有其他线程同步，以便每个线程可以安全读取由不同线程写入的共享内存位置，

 处理共享内存中的数据，
 如果必要的话，重新同步以确保共享内存已经由结果更新，
 将结果写回到设备内存中。

5.1.2.1 全局内存

全局内存空间没有高速缓存，所以最重要的是按照正确的访问模式获得最大的内存带宽，尤其是已知对设备内存的访问有多昂贵时。
首先，设备能够在单个指令中将32位、64位或128位字从全局内存读取到寄存器中。要将如下赋值：

编译到单个加载指令中，type必须使得sizeof(type)等于4、8或16，且类型为type的变量必须对齐为sizeof(type)个字节（也就是说，让其地址是sizeof(type)的倍数）。
对于4.3.1.1一节中介绍的内置类型，比如float2或float4，对齐要求将自动完成。
对于结构体，大小和对齐要求可以由编译器使用对齐指定符__align__(或__align__(16)来强制执行，比如

更准确地说，在每个半warp中，半warp中的线程号N应访问地址

其中，HalfWarpBaseAddress具有类型type*，type满足上述的大小和对齐要求。此外，HalfWarpBaseAddress应对齐为16*sizeof(type)个字节（比如，是16*sizeof(type)的倍数）。驻留在全局内存中或由D.5或E.8中的内存分配例程之一返回的变量的任何地址BaseAddress始终对其为至少256个字节，所以为了满足内存对齐约束，HalfWarpBaseAddress-BaseAddress应是16*sizeof(type)的倍数。
注意，如果半warp满足上述所有要求，即使半warp的一些线程不实际访问内存，每线程内存访问也将合并。
与仅分别履行其每个半warp的合并要求相反，我们建议履行整个warp的合并要求，因为将来的设备将使其成为适当合并的必需操作。
图5-1显示了已合并内存访问的一些示例，而图5-2和图5-3显示了未合并内存访问的一些示例。
已合并64位访问提供了比已合并32位访问稍低的带宽，已合并128位访问提供了比已合并32位访问低很多的带宽。然而，当访问是32位时，尽管未合并访问的带宽比已合并访问的带宽低大约一个数量级，但当访问是64位时，仅低大约4倍，当访问是128时，仅低大约2倍。

左：已合并的float内存访问。
右：已合并的float内存访问（分散warp）。
图5-1. 已合并全局内存访问模式的示例

左：非顺序的float内存访问。
右：未对齐的开始地址。
图5-2. 未合并全局内存访问模式的示例

左：不相邻的float内存访问。
右：未合并的float3内存访问。
图5-3. 未合并全局内存访问模式的示例

常见的全局内存访问模式是当线程ID为tid的每个线程访问位于类型为type*的地址BaseAddress上的数组的一个元素时，使用下列地址：

要获得内存合并，type必须满足上述大小和对齐要求。特别地，这意味着，如果type是大于16个字节的结构体，则应分割为满足这些要求的多个结构体，而且数据应在内存中排列为这些结构体的多个数组，而非类型为type*的单个数组。
另一个常见的全局内存访问模式是当索引为(tx,ty)的每个线程访问位于类型为type*、宽度为width的地址BaseAddress上的2D数组的一个元素时，使用下列地址

在这种情况下，仅当满足下列条件，用户才能获得线程块的所有半warp的内存合并：
 线程块的带宽是半个warp大小的倍数；
 width是16的倍数。
特别地，这意味着，如果宽度不是16的倍数的数组实际使用向上取整为最接近的16的倍数进行分配，且其行相应地进行填补，则此数组将获得非常有效的访问。cudaMallocPitch()和cuMemAllocPitch()函数及其相关的内存复制函数（参见D.5和E.8）允许程序员编写不依赖于硬件的代码来分配符合这些约束的数组。

5.1.2.2 常量内存

常量内存空间具有高速缓存，所以从常量内存中的读取仅在高速缓存缺失时，耗费从设备内存中的一个内存读取，否则它仅耗费从常量高速缓存中的一个读取。
对于半warp的所有线程，只要所有线程读取同一地址，则从常量内存中读取与从寄存器中读取一样快。成本随所有线程读取的不同地址数线性扩展。与仅让每个半warp中的所有线程读取同一地址相反，我们建议让整个warp的所有线程读取同一地址，因为将来的设备将需要此操作来实现完全的快速读取。

5.1.2.3 纹理内存

纹理内存空间具有高速缓存，所以纹理拾取仅在高速缓存缺失时，耗费从设备内存中的一个内存读取，否则它仅耗费从纹理高速缓存中的一个读取。纹理高速缓存针对2D空间局部性进行了优化，所以读取紧密相邻的纹理地址的同一warp的线程将达到最佳性能。此外，它还设计用于流水化具有恒定延迟的拾取。比如，高速缓存命中降低了DRAM带宽需求，但没有降低拾取延迟。
通过纹理拾取读取设备内容可能是从全局或常量内存中读取的设备内存的有利备选方案，详见5.4。

5.1.2.4 共享内存

因为位于芯片上，所以共享内存要比本地和全局内存空间快得多。实际上，对于warp的所有线程，访问共享内存与访问寄存器一样快，只要在线程之间没有任何库冲突，详见下文。
要获得高内存带宽，请将共享内存划分为同样大小的内存模块，命名为库，可以同时访问这些库。因此，由属于n个显式内存库的n个地址组成的任何内存读取或写入请求都可以同时获得服务，最后可收益n倍的有效带宽与单个模块的带宽一样高。
但是如果两个内存请求地址组成同一个内存库，则会导致库冲突和访问必须序列化。硬件将带有库冲突的内存请求按需分成许多单独的无冲突的请求，通过一个因子来减少有效带宽使其与单独内存请求的数目相等。如果单独内存请求的数目为n，则最初的内存请求据说会导致n种方式的库冲突。
要获得最高性能，因此了解如何将内存地址映射到内存库变得非常重要，这样做的目的是调度内存请求，以便最小化库冲突。
在共享内存空间这种情况下，将库组织为：将连续的32位字分配到连续的库中，且每两个时钟周期每个库都有一个32位带宽。
对于计算能力1.x的设备，warp大小为32，库为16（参见5.1）；将warp的共享内存请求划分为第一半warp的一个请求和第二半warp的一个请求。因此，属于第一半warp的线程和属于同一warp的第二半warp的线程间不会发生任何库冲突。
一个常见的情况是每个线程从按线程ID tid索引的数据中使用某个跨度s来访问一个32位字：

这种情况下，只要s*n是库数m的倍数，或者同等地，只要n是m/d的倍数，其中d是m和s最大公约数，则线程tid和tid+n将访问同一库。因此，仅当半warp大小小于或等于m/d时，才不会发生库冲突。对于计算能力1.x的设备，仅当d等于1时，或者换句话说，因为m是2的倍数，所以仅当s是奇数时，才不会发生任何库冲突。
图5-4和图5-5显示了一些无冲突内存访问的示例，而图5-6显示了一些导致库冲突的内存访问示例。
值得一提的其他情况是当每个线程访问大小小于或大于32位的元素时。例如，如果按下列方式访问char的数组，则会发生库冲突。

例如，因为shared[0]、shared[1]、shared[2]和shared[3]属于同一个库。但是，如果按下列方式访问同一数组，则不会发生任何库冲突：

将导致下列结果：
 如果type定义如下，则结果为三个单独的无库冲突的内存读取

因为每个成员是使用三个32位字的跨度来访问。
 如果type定义如下，则结果为两个单独的无库冲突的内存读取

因为每个成员是使用三个32位字的跨度来访问。
 如果type定义如下，则结果为两个单独的无库冲突的内存读取

因为每个成员是使用五个字节的跨度来访问。
最后，共享内存还具有广播机制，当服务一个内存读取请求时，可以读取一个32位字并同时广播到多个线程。当半warp的多个线程从同一32位字内的地址读取时，这将减少库冲突的数目。更精确地说，由多个地址组成的内存读取请求随时间由多个步来服务——每两个时钟周期一步——每步服务一个这些地址的无冲突子集，直到所有地址都已服务完毕；在每一步，子集从尚未服务的剩余地址中构建，过程如下：
 选择由剩余地址指向的其中一个字作为广播字，
 将下列内容包括在子集中：
 位于广播字内的所有地址，
 由剩余地址指向的每个库的一个地址。
选择哪个字作为广播字以及在每个周期为每个库选择哪个地址均未指定。
常见的无冲突情况是当半warp的所有线程从同一32位字内的地址中读取时。
图5-7显示了一些涉及广播机制的内存读取访问的示例。

左：跨度为一个32位字的线性寻址
右：随机排列
图5-4. 无库冲突的共享内存访问模式示例

跨度为三个32位字的线性寻址。
图5-5. 无库冲突的共享内存访问模式示例

左：跨度为两个32位字的线性寻址将导致2路库冲突。
右：跨度为八个32位字的线性寻址将导致8路库冲突。
图5-6. 有库冲突的共享内存访问模式示例

左：因为所有线程从同一32位字中的地址读取，所以此访问模式是无冲突的。
右：如果在第一步期间选择库5中的字作为广播字，则此访问模式不会导致任何库冲突，否则会导致2路库冲突。
图5-7. 有广播的共享内存读取访问模式示例

5.1.2.5 寄存器

通常，访问寄存器对于每条指令需要零个额外时钟周期，但是由于寄存器读后写依赖关系和寄存器内存库冲突，可能会发生延迟。
由读后写依赖关系导致的延迟可以忽略，只要每个多处理器至少有192个活动线程使其隐藏。
编译器和线程调度器调度指令尽可能最佳，以避免寄存器内存库冲突。当每块中的线程数是64的倍数时，可以获得最佳结果。除了遵循此规则之外，应用程序对这些库冲突没有任何直接的控制。特别地，无需将数据打包为float4或int4类型。

5.2 每块的线程数

给定每网格的线程总数，选择每块的线程数或等同的块数时应该最大化可用计算资源的利用率。这意味着块的数目应该至少与设备中的多处理器的数目一样多。
此外，当每个多处理器仅运行一块时，如果每块没有足够的线程来覆盖加载延迟，则在线程同步期间，以及设备内存读取期间，将强制每个多处理器进入空闲状态。因此，最好的方法是每个多处理器上允许存在两个或多个活动块，以允许在等待的块和可以运行的块之间出现重叠。要让这种情况发生，不仅块的数目至少应该是设备中多处理器数目的两倍，而且每块分配的共享内存量至多应该是每个多处理器可用共享内存总量的一半（参见3.2）。更多线程块以管线方式在设备中分流，并在更大程度上分摊开销。
有了足够大数目的块，每块线程的数目应选择为warp大小的倍数，以避免使用未充满的warp而浪费计算资源，或者更好地，选择为64的倍数，究其原因，参见5.1.2.5。为每块分配更多线程有利于有效的时间分片，但是每块的线程越多，每线程可用的寄存器就越少。如果内核编译的线程数大于执行配置所允许的数目，这可能会阻止内核调用继续。当使用--ptxas-options=-v选项编译时，内核编译的寄存器数目（还有本地、共享和常量内存使用）由编译器报告。
对于计算能力1.x的设备，每线程可用的寄存器数等于：

其中，R是每个多处理器的寄存器总数(参见附录A），B是每个多处理器的活动块数，T是每块的线程数，ceil(T, 32)是T向上取整为32的最近倍数。

每块64个线程是最小的，并且仅当每个多处理器有多个活动块时才有意义。每块有192或256个线程比较好，而且通常允许有足够的寄存器进行编译。
如果用户想将其扩展到将来的设备，则每个网格的块数应该至少是100；1000个块将扩展到几代。
每个多处理器的活动warp数与活动warp（参见附录A）的最大数目的比率称作多处理器占有率。为了最大化占有率，编译器应尝试最小化寄存器使用，而且程序员需要小心选择执行配置。CUDA软件开发工具包提供了一个电子表格以帮助程序员基于共享内存和寄存器要求来选择线程块大小。

5.3 主机和设备之间的数据传送

设备和设备内存之间的带宽比设备内存和主机内存之间的带宽高得多。因此，用户应该争取最小化主机和和设备之间的数据传送，例如，将更多代码从主机移动到设备，即使这意味着要使用低并行计算来运行内核。中间数据结构可以在设备内存中创建，由设备操作，销毁，而且永远不会由主机映射，或复制到主机内存。
另外，由于每次传送都会有开销，所以将许多小的传送分批为一次大的传送要比单独执行每一个传送要好得多。
最后，使用页面锁定内存时，可以在主机和设备之间获得较高带宽，详见4.5.1.2。

5.4 纹理拾取与全局或常量内存读取

与从全局或常量内存中读取相比，通过纹理拾取进行设备内存读取具有下列几个优点：
 高速缓存，如果在纹理拾取中有位置，则可以潜在地展示较高带宽；
 不受内存访问模式的约束，此约束即全局或常量内存读取必须尊重以获得好的性能（参见5.1.2.1和5.1.2.2）；
 寻址计算的延迟隐藏得更好，可能会改善执行随机访问数据的应用程序的性能；
 打包的数据可以在单个操作中广播到多个独立变量中。
 8位和16位整数输入数据可以有选择地转化为[0.0,1.0]或[-1.0,1.0]范围内的32位浮点值（参见4.3.4.1）。
 如果纹理是CUDA数组（参见4.3.4.2），则硬件提供了可能对不同应用程序有用的其他能力，尤其是图像处理：

然而，在同一内核调用中，纹理高速缓存与全局内存写不保持一致，从而对已经在同一内核中通过全局写而写入的某个地址的纹理拾取将返回未定义的数据。换句话说，仅当此内存位置已经由先前的内核调用或内存复制更新时，而不是已经由同一内核调用中的同一个或另一个线程更新时，线程才可以通过纹理安全地读取某个内存位置。仅当作为内核从线性内存中的拾取无论如何不能写入到CUDA数组中时，这才相关。
5.5 整体性能优化策略
性能优化围绕三个基本策略：
 最大化并行执行；
 优化内存使用以获得最大内存带宽；
 优化指令使用以获得最大指令吞吐量。
最大化并行执行首先应通过暴露尽可能多的数据并行来结构化算法。在算法中，因为一些线程需要同步以便互相共享数据，而破坏了并行性的情况有两种：这些线程属于同一块，这种情况下，线程应该使用__syncthreads()，并通过同一内核调用中的共享内存来共享数据，或者这些线程属于不同块，这种情况下，必须使用两个单独的内核调用通过全局内存来共享数据，一个内核调用写入全局内存，另一个从全局内存读取。
暴露算法的并行之后，则需要尽可能有效地将其映射到硬件。通过仔细选择每个内核调用的执行配置来完成此操作，详见5.2。
应用程序还可以通过显式暴露通过流在设备上的并发执行，如4.5.1.5所述，在更高的水平上最大化并行执行，以及最大化主机和设备之间的并发执行。
最优化内存使用首先应最小化具有低带宽的数据传送。这意味着最小化主机和设备之间的数据传送，详见5.3，因为这要比在设备和全局内存之间的数据传送的带宽低得多。这也意味着通过最大化设备上共享内存的使用来最小化设备和全局内存之间的数据传送，详见5.1.2。有时候，最好的优化甚至可能是通过简单地重新计算数据来避免任何数据传送，而不管是否需要这样做。

详见5.1.2.1、5.1.2.2、5.1.2.3和5.1.2.4，取决于每种内存类型的访问模式，有效带宽可能有一个数量级的变化。因此，优化内存使用的下一步是基于最佳的内存访问模式，尽量优化地组织内存访问。此优化对于全局内存访问尤其重要，因为全局内存访问的带宽很低，且其延迟是数百个时钟周期（参见5.1.1.3）。
另一方面，通常仅当共享内存访问具有高度的库冲突时才值得优化。
对于优化指令使用，应该最小化具有低吞吐量的算术指令的使用。这包括在不影响最终结果时用精度换速度，比如使用固有函数，而不使用常规函数（固有函数在表B-2中列出），或使用单精度而不使用双精度。由于设备的SIMD本质，所以要特别注意控制流指令，详见5.1.1.2。