目录

背景

总体性能优化策略

最大化使用率

应用层面

设备层面

多处理器层面

占有率计算器

最大化内存吞吐量

主机和设备间的数据迁移

设备内存访问

全局内存

尺寸和对齐要求

二维数组

局部内存

共享内存

常量内存

纹理和表面内存

最大化指令吞吐量

算术指令

单精度浮点除法

单精度浮点倒数开方

单精度浮点开方

正余弦

整型运算

半精度运算

类型转换

控制流指令

同步指令

结语

背景

本文来进行对CUDA10.0官方文档的第五章，也是正文最后一章——性能指南的翻译

总体性能优化策略

性能优化可以被概括为以下三个基本策略：

• 最大化并行度以达到最大使用率；
• 优化内存使用以达到最大内存吞吐量；
•  优化指令使用以达到最大指令吞吐量。

一个应用的某部分使用哪个策略能达到最好的性能取决于那部分的性能掣肘因素；例如，对主要由内存访问限制的核进行指令优化不会得到任何的显著性能提升。因此，优化工作应该持续朝着衡量和键值优化掣肘因素，比如可以使用cuda优化器来进行。而且，比较指定核函数浮点操作的吞吐量或内存吞吐量与对应设备的理论巅峰吞吐量的差距可以指示这个核还有多大的优化空间。

最大化使用率

为了最大化使用率，我们应该以暴露尽可能大的并行度、高效地把并行度映射到系统组件中以让这些组件在大多数时候保持忙碌的方式组织应用。

应用层面

在高层面，应用应该通过使用前文中异步并发执行一节中描述的异步函数和流来最大化主机、设备和连接它们的总线的并行执行度。应该为每个处理器分配它们最适合执行的任务：主机执行串行任务，设备执行并行任务

对于并行任务，算法存在着一些因为一些线程保持同步来互相共享数据而打破并行的时刻，这有两种情况：要么这些线程处于同一个线程块，此时它们应该在一个核函数调用中使用__syncthreads()，再通过共享内存来共享数据；要么它们属于不同的线程块，此时它们必须通过使用两个独立的核函数调用——一个从全局内存中读、一个往全局内存中里写的方式来通过全局内存共享数据。第二种情况更耗性能，因为它增加了由额外的核函数调用和全局内存通信产生的开销。因此这种情况应该尽量避免出现，通过让要求线程间通信的计算尽可能在单一线程块中执行的方式来把算法映射到cuda编程模型中

设备层面

在低层面，应用应该最大化设备中多处理器的并行执行度。多个核函数可以在设备上并发执行，所以最大化使用率也可以通过使用流的方式来让足够多的核函数并发执行，如前文中异步并发执行一节中所述

多处理器层面

在更低的层面上，应用应该最大化一个多处理器内部多个功能单元的并行执行度。

如前文中硬件多线程一节中所记载，一个GPU多处理器依赖于线程级别的并行来最大化其功能单元的利用率，因此使用率直接和共存的伪线程数相关。在每个指令发起的时刻，一个伪线程调度器选择已经准备好执行下一条指令的伪线程(如果存在的话)，再把指令分发给伪线程中的活跃线程。伪线程用以为下一条指令的执行做好准备所消耗的时钟周期数被称之为延迟，当在延迟期间的每个时钟周期内，所有的伪线程调度器都有一些要分发给某些伪线程的指令时，处理器的完全利用也就达到了，换句话说，此时延迟周期被彻底地“隐藏”了。隐藏长度为L的时钟周期所需的指令条数取决于这些指令各自的吞吐量。例如，所有指令的最大吞吐量在计算能力为3.X的设备上为8L，因为多处理器在每个时钟周期一次性为四个伪线程中的每一个都分发一对指令。因此对于计算能力为3.X的设备，每个周期发起的8条指令是发给四个不同伪线程的四对，每一对都属于同一个伪线程。

伪线程没有准备好执行下一条指令的最常见原因是指令的输入操作数还不可用。如果所有的输入操作数都在寄存器中，造成延迟的原因就是寄存器依赖，比如某些输入操作数是被之前还没有执行完的指令写入的。在背靠背寄存器依赖的（例如，一些输入操作数由前一条指令写入）情况下，延迟等于上一条指令的执行时间，在这段时间超线程调度器必须为别的超线程调度指令。不同的指令有不同的执行时间，但在计算能力3.X的设备上大多都是11个时钟周期左右，这段时间可以让调度器调度44个伪线程（假设伪线程执行吞吐量最大的指令，否则调度的伪线程数会下降；以及指令级的并行度足够，这样调度器总是能为每个伪线程分配一对指令）。

如果一些输入操作数存在于主机内存，那延迟就高得多了：对于计算能力为3.X的设备来说，可能有200~400个时钟周期。在这么长的延迟周期中让伪线程调度器保持忙碌所需的伪线程数取决于核函数的实现以及核函数的指令级并行度。总体上说，如果不需要主机内存操作数的指令数和需要主机内存操作数的指令数的比值越低(这个比值通常称为程序的算术强度)，就需要更多的伪线程来让伪线程调度器保持忙碌。例如，假设这个比值是30，计算能力为3.X的设备上的延迟为300个时钟周期，那么所需的伪线程数量大约为40（其余假设和上一段最后相同）。

伪线程没有准备好执行下一条指令的另一个原因是它在等待一些内存栅栏或者同步点。同步点可以强迫多处理器空闲，因为越来越多的伪线程在等待同一线程块中其他伪线程执行完同步点之前的指令。在处理器上使用多个线程块可以缓解这种情况下的空闲，因为不同线程块的伪线程不必在同步点彼此等待。

对于一个给定的核函数，每个多处理器上能够共存的线程块和伪线程数取决于调用它时的执行配置，而多处理器的内存资源和核函数的资源需求的关系在硬件多线程中有所记载。通过在编译时使用-ptxas-options=-v选项可以让编译器报告寄存器和共享内存的使用情况

一个核函数使用的寄存器数量对共存的伪线程数量有着重大影响。例如在计算能力为6.x的设备上，如果一个核函数使用了64个寄存器，每个线程块有512个线程，而且需要很少量的共享内存，那么两个块(总共32个伪线程)可以在一个多处理器上共存，因为它们总共需要2 * 512 * 64个寄存器(64K)。但是，一旦核函数使用了更多的寄存器，那就只有一个线程块(总共16个伪线程)可以共存于一个多处理器上，因为2个块就至少需要2 * 512 * 65 = 65K > 64K个寄存器了。因此编译器会在把寄存器溢出和指令数量保持到最少的同时，尝试减少寄存器的使用量。寄存器的使用可以通过maxrregcount编译选项或者启动边界来进行控制。

每个双精度变量和长长整型变量会使用两个寄存器。

对于一个指定的核函数调用而言，执行配置对性能的影响基本是取决于核函数代码，因此推荐去做实验。应用也可以基于寄存器文件大小和共享内存数量来参数化执行配置，其中共享内存数取决于设备的计算能力、多处理器数以及设备上的内存带宽，这些都可以使用运行时API查询

每个线程块的线程数应该是伪线程容量的整数倍，以尽量避免线程较少的伪线程对计算资源的浪费

占有率计算器

存在着一些帮助程序员基于寄存器和共享内存数选择线程块大小的API函数。

占有率计算API——cudaOccupancyMaxActiveBlocksPerMultiprocessor()可以提供基于某个核的线程块大小和共享内存使用计算出来的占有率预测值，这个函数会针对每个处理器上并发线程块数来汇报占有率。注意这个值可能被转换成其他指标，它乘以每个块的伪线程数会得到处理器上并发的伪线程数；再除以处理器支持的最大伪线程数，会得到占有率的百分比。

基于占有率的启动配置API(cudaOccupancyMaxPotentialBlockSize()和cudaOccupancyMaxPotentitalBlockSizeVariableSMem())会启发式地计算达到多处理器层面最大的禅忧虑所需的执行配置。

下面的代码计算了核函数MyKernel()的占有率，汇报了并发伪线程数和处理器上最大伪线程数的比值：

#include "stdio.h"
#include "cuda_runtime.h"__global__ void MyKernel(int* a, int* b, int* c) {int idx = blockIdx.x * blockDim.x + threadIdx.x;c[idx] = a[idx] + b[idx];
}int main() {int numBlocks;int blockSize = 32;int device = 0, activeWraps, maxWarps;cudaDeviceProp prop;cudaGetDevice(&device);cudaGetDeviceProperties(&prop, device);cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, MyKernel, blockSize, 0);activeWraps = numBlocks * blockSize / prop.warpSize;maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;printf("Occupancy: %.3f\n", activeWraps / (double ) maxWarps);return 0;
}

    运行结果如下：

下面的代码根据用户输入配置了MyKernel的基于占有率的启动参数：

#include "stdio.h"
#include "cuda_runtime.h"__global__ void MyKernel2(int* a, int size) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < size) {a[idx] = a[idx] + 3;}
}int main() {int blockSize, minGridSize, gridSize;int size;scanf("%d", &size);int *h_data = (int *) malloc(size * sizeof(int));int *d_data;for (int i = 0; i < size; i++) {h_data[i] = i * i;}cudaMalloc(&d_data, size * sizeof(int ));cudaMemcpy(d_data, h_data, size * sizeof(int ), cudaMemcpyHostToDevice);cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, (void *)MyKernel2, 0, size);gridSize = (size + blockSize - 1) / blockSize;MyKernel2<<<gridSize, blockSize>>>(d_data, size);cudaMemcpy(h_data, d_data, size * sizeof(int ), cudaMemcpyDeviceToHost);for (int i = 0; i < size; i++) {printf("%d\n", h_data[i]);}cudaFree(d_data);free(h_data);return 0;
}

输出结果如下：

cuda工具包也提供了一个自注释、单独的占有率计算器和启动配置器的实现：/usr/local/cuda-10.0/include/cuda_occupancy.h，用来在不依赖cuda软件栈的情况下使用。还提供了一个电子表格版本的占有率计算器，这个电子表格计算器作为一个将影响占有率的参数改变时的影响可视化学习工具特别有用。

最大化内存吞吐量

最大化应用的内存吞吐量是减少低带宽的数据传输。这意味着要减少主机和设备间的数据传输，因为这种数据传输的带宽比设备和全局内存间数据传输的要低得多。这也意味着要通过使用片上内存(共享内存、纹理和常量内存，以及计算能力2.X及以上设备可用的L1和L2缓存)来减少设备和全局内存间的数据传输。

共享内存等同于用户管理的缓存：应用要显式地分配和访问它，一个典型的编程模板是把来自设备内存的数据缓存到共享内存中，换句话说，让一个线程块中的每个线程从设备内存中把数据加载到共享内存、和块内其他线程同步以让每个线程可以安全读取被不同线程共享的内存地址、在共享内存中处理数据、如有必要再次同步以确保共享内存已经更新了结果、把结果写回到设备内存中。

对于某些应用(例如全局内存访问模板是数据依赖的应用)，出纳痛的硬件管理缓存更适合用来利用数据本地性。对于计算能力3.x和7.x的设备来说，L1和共享内存使用的相同的片上内存数，并且与共享内存相比，有多少内存是专门用于L1缓存的都是可以为每个核函数调用配置的，

对于每种类型内存的不同访问模式，核函数访问的内存吞吐量可能有数量级的差异。最大化内存吞吐量的下一步因此就是基于最优内存访问模板尽可能优化内存访问的组织结构。这种优化对于全局内存尤为重要，因为全局内存的带宽很低，所以未优化的全局内存访问对性能的掣肘更为显著。

主机和设备间的数据迁移

应用应该尽力减少主机和设备间的数据迁移，一个方法是把更多的代码从主机端移到设备端，即便这样会让核函数的运行有比较低的计算并行度。中间数据结构可以在设备端创建、操作和销毁，避免任何与主机内存进行任何的映射或复制。

而且，因为每次迁移都有代价，把很多小批量的迁移合并成一个大批量的迁移总是会得到更好的性能。在使用前端总线的系统上，主机和设备之间的数据迁移可以通过使用页锁主机内存(参见前文)来达到更好的性能。

另外，当使用映射页锁内存时，无需分配任何设备内存，也不需要进行显式地主机与设备内存间的数据复制，每次核函数访问映射内存时会隐式地执行数据迁移。为了更好的性能，这些内存访问应该和全局内存访问合并。假设合并已经完成，并且映射内存只被读或写一次，使用映射页锁内存会比在主机和设备内存间进行显式复制有着更好的性能表现。

在设备内存和主机内存都是同一块物理内存的集成系统上，主机和设备内存间的任何复制都是多余的，我们应该转而使用页锁内存。应用可以对设备是否集成(intergrated字段是否为1)进行查询，请参见前文中设备枚举一节

设备内存访问

访问可寻址内存(全局、局部、共享、常量或纹理内存)的指令可能根据伪线程中的多个线程的内存地址分布被重新发起多次，这种情况下分布是如何影响指令吞吐量取决于不同的内存类型。比如，对于全局内存而言，基本的规则是地址分布越零散，吞吐量下降得越快

全局内存

全局内存存在于设备内存中，设备内存是通过32、64或128字节的内存转换来访问的。这些内存转换必须自然对齐：只有和各自尺寸对齐的32、64和128字节的设备内存段才可以被内存转换进行读写。

当伪线程执行访问全局内存的指令时，它把伪线程内线程对内存的访问合并成一个或多个这种内存转换，其数量取决于每个线程访问的字数量以及线程间的内存地址分布。通常来说，需要的转换数越多，除了被线程访问的字外被转换的无用字数也就越多，从而降低了指令吞吐量。例如，如果为每个线程的4字节访问生成一个32字节的内存转换的话，吞吐量会降至原来的1/8。

需要进行多少转换和最终影响多少吞吐量取决于设备的计算能力。为了最大化全局内存的吞吐量，因此有必要最大化合并访问，方式有：设备计算能力为3.x、5.x、6.x和7.x时要使用最优的访问模板；使用满足尺寸和对齐要求的数据类型；在某些情况下填充数据(比如访问二维数组时)

尺寸和对齐要求

全局内存指令支持对大小等于1、2、4、8、16字节的数据进行读写，当且仅当数据类型的尺寸为1、2、4、8、16字节并且数据自然对齐(比如其地址是尺寸的整数倍)时，任何对存在于全局内存的数据的访问(通过变量或指针)会被编译成单个全局内存指令。

如果尺寸和对齐要求没有被满足，数据访问会被编译成多条访问模板不同的指令，这些模板阻止指令的充分合并。因此推荐对存在于全局内存的数据使用满足这些要求的数据类型。

自动满足对齐要求的内置数据类型有char、short、int、long、longlong、float、double、float2和float4。对于结构体而言，对齐要求可以通过使用__align__(8|16)来被编译器强行满足，比如下面的写法：

struct __align__(8) {int x;int y;
} A;struct __align__(16) {int x;int y;int z;
} B;

由驱动或运行时API的内存分配函数返回的或者任何存在于全局内存中的变量的地址总是至少256字节对齐的。读取非8字节或16字节对齐的字会产生错误的结果(几个字的偏差)，所以我们要特别刘依保持这些类型的数组或变量的起始地址。一个容易被忽视的典型情况是当使用一些自定义全局内存分配机制时，被分成多个数组的单个大块内存分配取代了多次调用cudaMalloc()或cuMemAlloc()函数进行的多次数组分配，这是每个数组的起始地址就应该是块起始地址的偏移。

二维数组

一个常见的全局内存访问模板是当每个索引为(tx, ty)的线程使用下面的取址方式来访问宽高均为width的二维数组中的某个元素，这个二维数组的起始地址为BaseAddress，类型为type*(符合尺寸和对齐要求中的要求)：

BaseAddress + width * ty + tx

对于这些要被充分合并的访问，线程块的宽度和数组宽度必须是伪线程容量的整数倍。特别地，这意味着宽度不是伪线程容量整数倍的数组会被高效地访问，如果它被分配了一个最近且是伪线程容量整数倍的宽度，并且其行据此进行了填充。cudaMallocPitch()函数和cuMemAllocPitch()函数以及相关的内存复制函数可以让程序员写出硬件独立的代码以依照这些限制来分配数组。

局部内存

局部内存访问只发生在某些自动类型的变量上，可能被编译器放到局部内存中的自动变量包括：无法确定是否使用常量进行索引的数组；消耗太多寄存器空间的结构体或数组；使用超过可用量的寄存器的任何变量(也就是会发生寄存器溢出的变量)。查看ptx汇编代码(通过使用-ptx或-keep选项进行编译)会看到一个变量是否在编译第一阶段被放到局部内存中，因为被放到局部内存的变量会使用.local助记符声明，访问时会使用ld.local和st.local助记符进行读写。即便第一阶段没有把变量放到局部内存里，随后的编译阶段也有可能这么做，如果编译器发现变量消耗了目标架构的太多寄存器空间：使用cuobjdump查看cubin对象会告诉我我们是不是这种情况。而且，当使用--ptxas-options=-v编译时，编译器会报告每个核函数的局部内存总用量lmem。注意，一些数学函数有着可能访问局部内存的实现。

局部内存存在于设备内存中，因此局部内存访问拥有和全局内存访问一样的高延迟低带宽，也要服从设备内存访问中和全局内存一样的限制与要求。但是，局部内存是由被连续的线程id访问的连续4字节的字组成的，因此只要一个伪线程内的所有线程访问相同的相对地址(比如一个数组中的相同索引、结构体中的相同变量)，这种访问就是完全合并的。

在一些计算能力为3.x的设备上，局部内存总是用和全局内存相同的方式缓存到L1和L2上。在计算能力为5.X和6.X的设备上，部内存总是用和全局内存相同的方式缓存到L2上。

共享内存

因为共享内存是片上内存，它就有着比局部内存和全局内存更高的带宽和更低的延迟。为了达到高带宽，共享内存被分成几个等量的、可以被同时访问的内存模块，称之为银行。任何由n个地址(每个各属于一个独立的银行)组成的的内存读写请求都可以被同时服务，从而得到的整体带块是每个单独模块带宽的n倍。

然而，如果一个内存请求中的两个地址落到了同一个内存银行中，就会产生银行冲突，对应的访问就得串行。硬件会把带银行冲突的内存请求分成很多的没有冲突的请求，这样会减少吞吐量，减少的幅度等于独立请求数。如果独立请求数为n，那么我们就说最初的内存请求造成了n步银行冲突。

为了达到最好的表现，理解内存地址如何映射成内存银行以调度内存请求来减少银行冲突就显得很重要。

常量内存

常量内存存在于设备内存中，被缓存在常量缓存里。如果初始请求中放了不同的内存地址，它就会被分成与内存地址数等量的独立请求，并会减少吞吐量，减少的幅度等于独立请求数。

如果产生缓存命中，就是用常量缓存的吞吐量来服务产生的独立请求；否则就得使用全局内存。

纹理和表面内存

纹理和表面内存存在于设备内存中，缓存于纹理缓存里，所以一个纹理获取或表面读只会在缓存不命中的情况下读取一次设备内存，命中时就读取纹理缓存。纹理缓存对二维空间局部性有优化，所以访问纹理或表面地址且在二维空间上挨得更近、处于同一伪线程内的线程会达到最好的表现。而且，它也是为带有不变延迟的流获取而设计的，缓存命中会减少带宽消耗，但不会降低获取延迟

通过纹理或表面获取来读取设备内存有以下几个好处，从而让这种方式是从全局或常量内存读取设备内存的一种有效替代手段：

• 如果内存读取没有依照全局或常量内存读要得到好的性能而必须依照的模版的话，只要在纹理获取或表面读取中存在本地性，就可以得到更高的带宽；
• 取址计算是在核函数外被专用的单元执行的；
• 包装好的数据可以在单个操作中的独立变量间广播；
• 单字节或双字节的整型输入数据可以被转换成四字节的浮点数，值域为[-1.0, 1.0]或[0.0, 1.0]，可参见前文中纹理内存一节

最大化指令吞吐量

为了最大化指令吞吐量，应用应该：

• 尽量减少低吞吐量算术指令的使用；包括当不影响最终结果时改变精度换取速度，比如使用指令级别而非常规函数、单精度而非双精度、把非正规化的数清零等；
• 尽量避免因为控制流指令产生的线程块分散；
• 通过优化同步点或者使用受限指针来减少指令数量

在这一节，吞吐量以每个多处理器每个时钟周期进行的操作数量表示。对于容量为32的伪线程，一条指令对应32个操作，所以如果N是每时钟周期进行的操作数量，那么每时钟周期的指令吞吐量就是N / 32。所有的吞吐量都针对的是一个多处理器，如果要针对整台设备，就得把这个数乘以设备上多处理器的数量。

算术指令

下表给出了各种计算能力的设备硬件本地支持的运算指令的吞吐量(每个多处理器上每时钟周期产生的结果数)

其中多指令表示这个算术指令会被编译成多条指令，剩下的上标的含义如下表所示：

其他指令和函数都是在本地指令之上实现的，不同的计算能力的设备实现可能不一样，编译过后本地指令的数量也随编译器版本的不同而有所变化。对于复杂的函数，可能针对不同的输入有多条代码路径。cuobjdump可以被用来检测一个cubin对象里的具体实现

一些函数的视线在cuda头文件里是可见的。整体上，使用-ftz=true(清零非正规化数字)编译的代码要比-ftz=false编译的有着更好的性能，同样，使用-prec-div=false(低精度除法)和-prec-sqrt=false(低精度开方)也都比各自对立面有着更好的性能。

单精度浮点除法

__fdividef(x, y)提供了比除法操作符更快的单精度浮点除法

单精度浮点倒数开方

为了保留IEEE-754语义，仅当求倒和开方都是近似计算(例如-prec-div=false而且-prec-sqrt=false)时，编译器可能把1.0/sqrtf()优化成rsqrtf()，因此建议需要时直接调用rsqrtf()

单精度浮点开方

单精度浮点开方是通过先求倒数平方根再求倒数实现的，而不是先求倒数平方根再用乘法，所以它会为0和∞得到正确的结果

正余弦

sinf(x)、cosf(x)、tanf(x)、sincosf(x)及其对应的双精度指令的开销很大，如果x增大一个数量级，那开销会更大。更准确地说，参数归约代码由一条快的代码路径和一条慢的代码路径组成。快路径由少量的乘加混合操作组成，用来计算数量级较小的参数结果；慢路径由较长的计算组成，以在全参数范围得到一个正确的结果，因此慢路径是为数量级较大的参数服务的。目前，三角函数的参数归约代码会为参数数量级小于105615.0f的单精度函数和小于2147483648.0的双精度函数选择快路径。由于慢路径比快路径需要更多的寄存器，通过把一些中间变量存到因为延迟和带宽影响性能(请参见前文中设备内存访问一节)的局部内存中来减少慢路径的寄存器压力的工作已经在进行，目前，单精度函数会使用28字节的局部内存，双精度函数则会使用44字节，但准确数量不会不变。

因为慢路径中计算链的长度个和局部内存的使用，当快慢两条路径的参数归约比较时，使用前者时这些三角函数的吞吐量就会降一个数量级。

整型运算

整数除法和取模开销很大，因为它们会被编译成20条指令。某些情况下，它们可以被位运算代替：如果n是2的整数次幂，那么i / n = i >> log2(n)，i % n = i & (n - 1)，当n是字面量时，编译器会帮我们做这些转换。

 __brev和__popc映射成一条指令，__brevll和__popcll则会映射成多条指令；

__[u]mul24是没有任何理由再去使用的指令级函数。

半精度运算

为了让半精度浮点值的加、乘、乘加达到好的吞吐量，我们建议使用half2数据类型。向量指令(__hadd2、__hsub2、__hmul2、__hfma2等)就可以被用来在一条指令里做两个操作了。使用half2代替half的两次调用也可以帮助提升其他指令的性能，比如伪线程混洗。

指令__halves2half2被提供用来把两个半精度的值转换成half2数据类型

类型转换

有时，编译器必须插入转换指令，引入加法执行环，这种情况有：

• 操作char或short类型变量，但这些变量的操作数一般要被转换成int的函数；
• 为单精度浮点计算(交给标准C/C++执行)传入双精度浮点常量(比如没有使用任何类型后缀声明的常量)作为输入时

后一种情况可以通过使用单精度浮点常量来避免，这种浮点常量用后缀f声明，比如3.141592653589793f、1.0f、0.5f等

控制流指令

任何控制流指令(if、switch、do、while、for)都可以通过让同一伪线程内的线程分散(执行不同的代码逻辑)来显著影响有效指令吞吐量。如果这种情况发生，不同的执行路径不得不串行，这样就增加了这个伪线程要执行的指令总数。

为了在控制流取决于线程id的情况下得到最好的表现，控制条件应该有利于减少伪线程的分散度，这是可能的，因为线程块间的伪线程分布是确定的，如SIMT架构中所提及。当控制条件只取决于threadIdx / warpSize(32)时，不会有伪线程内的分散，因为控制条件和伪线程完美对齐。

有时编译器可能会展开循环或者通过使用分支预测把短的if或switch代码块优化取代，这种情况下也没有伪线程能够分散，程序员可以通过使用#pragma unroll来直接控制循环的展开。

当使用分支预测时，取决于控制条件执行的指令都不会被跳过。相反，他们当中每一条都和一个针对每个线程的、根据控制条件被设置成真或假的判断代码或者判断谓语，尽管这些指令都会在执行中被调度，但只有判断谓语为真的指令才会被实际执行，判断谓语为假的指令就不会写结果、计算地址或者读取操作数。

同步指令

__syncthreads()函数在不同计算能力设备上的吞吐量(每时钟周期执行的操作数量)如下表所示，注意这个函数可以通过让多处理器强制空闲来影响性能，详情请参见前文中设备内存访问一节

 

计算能力	3.X	5.X	6.0	6.1	6.2	7.x
每时钟周期的操作数量	128	64	32	64	64	16

(正文完)

结语

CUDA10.0的正文部分翻译到此结束，最后的两篇文章我将会翻译附录CUDA动态并行和合作组部分，原文中最后一部分统一内存我没有翻译，因为文章说统一内存不能用在嵌入式设备上，而我最终就是想在嵌入式终端上利用CUDA进行异构编程，所以就没有翻译统一内存部分。