本节书摘来自华章计算机《CUDA C编程权威指南》一书中的第3章，第3.2节，作者 [美] 马克斯·格罗斯曼（Max Grossman），译 颜成钢 殷建 李亮，更多章节内容可以访问云栖社区“华章计算机”公众号查看。
3.2 理解线程束执行的本质
启动内核时，从软件的角度你看到了什么？对于你来说，在内核中似乎所有的线程都是并行地运行的。在逻辑上这是正确的，但从硬件的角度来看，不是所有线程在物理上都可以同时并行地执行。本章已经提到了把32个线程划分到一个执行单元中的概念：线程束。现在从硬件的角度来介绍线程束执行，并能够获得指导内核设计的方法。
3.2.1 线程束和线程块
线程束是SM中基本的执行单元。当一个线程块的网格被启动后，网格中的线程块分布在SM中。一旦线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。一个线程束由32个连续的线程组成，在一个线程束中，所有的线程按照单指令多线程（SIMT）方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进行操作。图3-10展示了线程块的逻辑视图和硬件视图之间的关系。
然而，从硬件的角度来看，所有的线程都被组织成了一维的，线程块可以被配置为一维、二维或三维的。在一个块中，每个线程都有一个唯一的ID。对于一维的线程块，唯一的线程ID被存储在CUDA的内置变量threadIdx.x中，并且，threadIdx.x中拥有连续值的线程被分组到线程束中。例如，一个有128个线程的一维线程块被组织到4个线程束里，如下所示：
用x维度作为最内层的维度，y维度作为第二个维度，z作为最外层的维度，则二维或三维线程块的逻辑布局可以转化为一维物理布局。例如，对于一个给定的二维线程块，在一个块中每个线程的独特标识符都可以用内置变量threadIdx和blockDim来计算：
对于一个三维线程块，计算如下：
一个线程块的线程束的数量可以根据下式确定：
因此，硬件总是给一个线程块分配一定数量的线程束。线程束不会在不同的线程块之间分离。如果线程块的大小不是线程束大小的偶数倍，那么在最后的线程束里有些线程就不会活跃。图3-11是一个在x轴中有40个线程、在y轴中有2个线程的二维线程块。从应用程序的角度来看，在一个二维网格中共有80个线程。
硬件为这个线程块配置了3个线程束，使总共96个硬件线程去支持80个软件线程。注意，最后半个线程束是不活跃的。即使这些线程未被使用，它们仍然消耗SM的资源，如寄存器。
3.2.2 线程束分化
控制流是高级编程语言的基本构造中的一种。GPU支持传统的、C风格的、显式的控制流结构，例如，if…then…else、for和while。
CPU拥有复杂的硬件以执行分支预测，也就是在每个条件检查中预测应用程序的控制流会使用哪个分支。如果预测正确，CPU中的分支只需付出很小的性能代价。如果预测不正确，CPU可能会停止运行很多个周期，因为指令流水线被清空了。我们不必完全理解为什么CPU擅长处理复杂的控制流。这个解释只是作为对比的背景。
GPU是相对简单的设备，它没有复杂的分支预测机制。一个线程束中的所有线程在同一周期中必须执行相同的指令，如果一个线程执行一条指令，那么线程束中的所有线程都必须执行该指令。如果在同一线程束中的线程使用不同的路径通过同一个应用程序，这可能会产生问题。例如，思考下面的语句：
假设在一个线程束中有16个线程执行这段代码，cond为true，但对于其他16个来说cond为false。一半的线程束需要执行if语句块中的指令，而另一半需要执行else语句块中的指令。在同一线程束中的线程执行不同的指令，被称为线程束分化。我们已经知道，在一个线程束中所有线程在每个周期中必须执行相同的指令，所以线程束分化似乎会产生一个悖论。
如果一个线程束中的线程产生分化，线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。线程束分化会导致性能明显地下降。在前面的例子中可以看到，线程束中并行线程的数量减少了一半：只有16个线程同时活跃地执行，而其他16个被禁用了。条件分支越多，并行性削弱越严重。
注意，线程束分化只发生在同一个线程束中。在不同的线程束中，不同的条件值不会引起线程束分化。
图3-12显示了线程束分化。在一个线程束中所有的线程必须采用if…then两个分支来表述。如果线程的条件为true，它将执行if子句；否则，当等待执行完成时，线程停止。
为了获得最佳的性能，应该避免在同一线程束中有不同的执行路径。请记住，在一个线程块中，线程的线程束分配是确定的。因此，以这样的方式对数据进行分区是可行的（尽管不是微不足道的，但取决于算法），以确保同一个线程束中的所有线程在一个应用程序中使用同一个控制路径。
例如，假设有两个分支，下面展示了简单的算术内核示例。我们可以用一个偶数和奇数线程方法来模拟一个简单的数据分区，目的是导致线程束分化。该条件(tid%2==0)使偶数编号的线程执行if子句，奇数编号的线程执行else子句。
如果使用线程束方法（而不是线程方法）来交叉存取数据，可以避免线程束分化，并且设备的利用率可达到100%。条件(tid/warpSize)%2==0使分支粒度是线程束大小的倍数；偶数编号的线程执行if子句，奇数编号的线程执行else子句。这个核函数产生相同的输出，但是顺序不同。
现在，使用代码清单3-1中的代码可以测量这两个核函数的性能。也可以从Wrox.com中下载simpleDivergence.cu文件。因为在设备上第一次运行可能会增加间接开销，并且在此处测量的性能是非常精细的，所以，添加了一个额外的内核启动（warmingup，与mathKernel2一样）来去除这一间接开销。
使用下面的命令编译这段代码：
在Fermi M2070 GPU上运行simpleDivergence，输出报告如下。两个内核的运行时间很相近。
通过使用nvprof分析器，可以从GPU中获得指标，从而可以直接观察到线程束分化。
在这里，nvprof的branch_efficiency指标是用来计算simpleDivergence的样本执行的：
下面的结果是由nvprof报告的。
分支效率被定义为未分化的分支与全部分支之比，可以使用以下公式来计算：
奇怪的是，没有报告显示出有分支分化（即分支效率是100%）。这个奇怪的现象是CUDA编译器优化导致的结果，它将短的、有条件的代码段的断定指令取代了分支指令（导致分化的实际控制流指令）。
在分支预测中，根据条件，把每个线程中的一个断定变量设置为1或0。这两种条件流路径被完全执行，但只有断定为1的指令被执行。断定为0的指令不被执行，但相应的线程也不会停止。这和实际的分支指令之间的区别是微妙的，但理解它很重要。只有在条件语句的指令数小于某个阈值时，编译器才用断定指令替换分支指令。因此，一段很长的代码路径肯定会导致线程束分化。
如下所示，重写mathKernel1核函数，使内核代码的分支预测直接显示：
添加mathKernel3，再次编译和运行文件simpleDivergence.cu，会报告下列性能：
使用下面的命令，可以强制CUDA编译器不利用分支预测去优化内核：
如下所示，可以用nvprof再次检查没有被优化的内核分化：
结果总结如下：
另外，可以用nvprof获得分支和分化分支的事件计数器，如下所示：
结果如下：
CUDA的nvcc编译器仍然是在mathKernel1和mathKernel3上执行有限的优化，以保持分支效率在50%以上。注意，mathKernel2不报告分支分化的唯一原因是它的分支粒度是线程束大小的倍数。此外，把mathKernel1中的if. . .else语句分离为mathKernel3的多个if语句，可以使分化分支的数量翻倍。
3.2.3 资源分配
线程束的本地执行上下文主要由以下资源组成：
- 程序计数器
- 寄存器
- 共享内存
由SM处理的每个线程束的执行上下文，在整个线程束的生存期中是保存在芯片内的。因此，从一个执行上下文切换到另一个执行上下文没有损失。
每个SM都有32位的寄存器组，它存储在寄存器文件中，并且可以在线程中进行分配，同时固定数量的共享内存用来在线程块中进行分配。对于一个给定的内核，同时存在于同一个SM中的线程块和线程束的数量取决于在SM中可用的且内核所需的寄存器和共享内存的数量。
图3-13显示了若每个线程消耗的寄存器越多，则可以放在一个SM中的线程束就越少。如果可以减少内核消耗寄存器的数量，那么就可以同时处理更多的线程束。如图3-14所示，若一个线程块消耗的共享内存越多，则在一个SM中可以被同时处理的线程块就会变少。如果每个线程块使用的共享内存数量变少，那么可以同时处理更多的线程块。
资源可用性通常会限制SM中常驻线程块的数量。每个SM中寄存器和共享内存的数量因设备拥有不同的计算能力而不同。如果每个SM没有足够的寄存器或共享内存去处理至少一个块，那么内核将无法启动。一些关键的限度如表3-2所示。
当计算资源（如寄存器和共享内存）已分配给线程块时，线程块被称为活跃的块。它所包含的线程束被称为活跃的线程束。活跃的线程束可以进一步被分为以下3种类型：
- 选定的线程束
- 阻塞的线程束
- 符合条件的线程束
一个SM上的线程束调度器在每个周期都选择活跃的线程束，然后把它们调度到执行单元。活跃执行的线程束被称为选定的线程束。如果一个活跃的线程束准备执行但尚未执行，它是一个符合条件的线程束。如果一个线程束没有做好执行的准备，它是一个阻塞的线程束。如果同时满足以下两个条件则线程束符合执行条件。
- 32个CUDA核心可用于执行
- 当前指令中所有的参数都已就绪
例如，Kepler SM上活跃的线程束数量，从启动到完成在任何时候都必须小于或等于64个并发线程束的架构限度。在任何周期中，选定的线程束数量都小于或等于4。如果线程束阻塞，线程束调度器会令一个符合条件的线程束代替它去执行。由于计算资源是在线程束之间进行分配的，而且在线程束的整个生存期中都保持在芯片内，因此线程束上下文的切换是非常快的。在下面几节中，你将会认识到为了隐藏由线程束阻塞造成的延迟，需要让大量的线程束保持活跃。
在CUDA编程中需要特别关注计算资源分配：计算资源限制了活跃的线程束的数量。因此必须了解由硬件产生的限制和内核用到的资源。为了最大程度地利用GPU，需要最大化活跃的线程束数量。
3.2.4 延迟隐藏
SM依赖线程级并行，以最大化功能单元的利用率，因此，利用率与常驻线程束的数量直接相关。在指令发出和完成之间的时钟周期被定义为指令延迟。当每个时钟周期中所有的线程调度器都有一个符合条件的线程束时，可以达到计算资源的完全利用。这就可以保证，通过在其他常驻线程束中发布其他指令，可以隐藏每个指令的延迟。
与在CPU上用C语言编程相比，延迟隐藏在CUDA编程中尤为重要。CPU核心是为同时最小化延迟一个或两个线程而设计的，而GPU则是为处理大量并发和轻量级线程以最大化吞吐量而设计的。GPU的指令延迟被其他线程束的计算隐藏。
考虑到指令延迟，指令可以被分为两种基本类型：
- 算术指令
- 内存指令
算术指令延迟是一个算术操作从开始到它产生输出之间的时间。内存指令延迟是指发送出的加载或存储操作和数据到达目的地之间的时间。对于每种情况，相应的延迟大约为：
- 算术操作为10～20个周期
- 全局内存访问为400～800个周期
图3-15表示线程束0阻塞执行流水线的一个示例。线程束调度器选取其他线程束执行，当线程束0符合条件时再执行它。
对于算术运算来说，其所需的并行可以表示成隐藏算术延迟所需要的操作数量。表3-3列出了Fermi和Kepler设备所需的操作数量。示例中的算术运算是一个32位的浮点数乘加运算（a＋b×c），表示在每个SM中每个时钟周期内的操作数量。吞吐量因不同的算术指令而不同。
吞吐量由SM中每个周期内的操作数量确定，而执行一条指令的一个线程束对应32个操作。因此，为保持计算资源的充分利用，对于Fermi GPU而言，每个SM中所需的线程束数量通过计算为640÷32＝20个线程束。因此，算术运算所需的并行可以用操作的数量或线程束的数量来表示。这个简单的单位转换表明，有两种方法可以提高并行：
- 指令级并行（ILP）：一个线程中有很多独立的指令
- 线程级并行（TLP）：很多并发地符合条件的线程
对内存操作来说，其所需的并行可以表示为在每个周期内隐藏内存延迟所需的字节数。表3-4列出了Fermi和Kepler架构的指标。
因为内存吞吐量通常表示为每秒千兆字节数，所以首先需要用对应的内存频率将吞吐量转换为每周期千兆字节数。可以使用下面的命令检测设备的内存频率：
例如，Fermi的内存频率（在Tesla C2070上测量得到）是1.566 GHz。Kepler的内存频率（在Tesla K20上测量得到）是2.6 GHz。因为1 Hz被定义为每秒一个周期，所以可以把带宽从每秒千兆字节数转换为每周期千兆字节数，公式如下所示：
用内存延迟乘以每周期字节数，可以得到Fermi内存操作所需的并行，接近74KB的内存I/O运行，用以实现充分的利用。这个值是对于整个设备，而不是对于每个SM来说的，因为内存带宽是对于整个设备而言的。
利用应用程序，把这些值与线程束或线程数量关联起来。假设每个线程都把一浮点数据（4个字节）从全局内存移动到SM中用于计算，则在Fermi GPU上，总共需要18 500个线程或579个线程束来隐藏所有内存延迟，具体运算如下所示：
74 KB÷4字节/线程≌18 500个线程
18 500个线程÷32个线程/线程束≌579个线程束
Fermi架构有16个SM。因此，需要579个线程束÷16个SM＝36个线程束/SM，以隐藏所有的内存延迟。如果每个线程执行多个独立的4字节加载，隐藏内存延迟需要的线程就可以更少。
与指令延迟很像，通过在每个线程/线程束中创建更多独立的内存操作，或创建更多并发地活跃的线程/线程束，可以增加可用的并行。
延迟隐藏取决于每个SM中活跃线程束的数量，这一数量由执行配置和资源约束隐式决定（一个内核中寄存器和共享内存的使用情况）。选择一个最优执行配置的关键是在延迟隐藏和资源利用之间找到一种平衡。下一节将会更加详细地研究这个问题。
3.2.5 占用率
在每个CUDA核心里指令是顺序执行的。当一个线程束阻塞时，SM切换执行其他符合条件的线程束。理想情况下，我们想要有足够的线程束占用设备的核心。占用率是每个SM中活跃的线程束占最大线程束数量的比值。
使用下述函数，可以检测设备中每个SM的最大线程束数量：
来自设备的各种统计数据在cudaDeviceProp结构中被返回。每个SM中线程数量的最大值在以下变量中返回：
用maxThreadsPerMultiProcessor除以32，可以得到最大线程束数量。代码清单3-2展示了如何使用cudaGetDeviceProperties获得GPU的配置信息。
从Wrox.com中可以下载simpleDeviceQuery.cu文件。使用以下命令编译并运行这个示例：
Tesla M2070的输出结果显示如下。每个SM中线程数量的最大值是1 536。因此，每个SM中线程束数量的最大值是48。
CUDA工具包包含了一个电子表格，它被称为CUDA占用率计算器，有助于选择网格和块的维数以使一个内核的占用率最大化。图3-17展示了占用率计算器的一个截图。
占用率计算器包含几个部分。首先，必须提供GPU的计算能力和内核的资源使用情况的信息。
在确定GPU的计算能力后，物理限制部分的数据是自动填充的。接下来，需要输入以下内核资源信息：
- 每个块的线程（执行配置）
- 每个线程的寄存器（资源使用情况）
- 每个块的共享内存（资源使用情况）
每个线程的寄存器和每个块的共享内存资源的使用情况可以从nvcc中用以下编译器标志获得：
一旦进入这个数据，内核占用率便会显示在GPU占用率数据段。其他部分提供必要的信息，来调整执行配置和资源使用情况，以获得更好的设备占用率。
内核使用的寄存器数量会对常驻线程束数量产生显著的影响。寄存器的使用可以用下面的nvcc标志手动控制。
-maxrregcount选项告诉编译器每个线程使用的寄存器数量不能超过NUM个。使用这个编译器标志，可以得到占用率计算器推荐的寄存器数量，同时使用这个数值可以改善应用程序的性能。
为了提高占用率，还需要调整线程块配置或重新调整资源的使用情况，以允许更多的线程束同时处于活跃状态和提高计算资源的利用率。极端地操纵线程块会限制资源的利用：
- 小线程块：每个块中线程太少，会在所有资源被充分利用之前导致硬件达到每个SM的线程束数量的限制。
- 大线程块：每个块中有太多的线程，会导致在每个SM中每个线程可用的硬件资源较少。
尽管在每种情况下会遇到不同的硬件限制，但它们都会导致计算资源未被充分利用，阻碍隐藏指令和内存延迟的并行的建立。占用率唯一注重的是在每个SM中并发线程或线程束的数量。然而，充分的占用率不是性能优化的唯一目标。内核一旦达到一定级别的占用率，进一步增加占用率可能不会改进性能。为了提高性能，可以调整很多其他因素。在后续章节中将详细介绍这些内容。
3.2.6 同步
栅栏同步是一个原语，它在许多并行编程语言中都很常见。在CUDA中，同步可以在两个级别执行：
- 系统级：等待主机和设备完成所有的工作
- 块级：在设备执行过程中等待一个线程块中所有线程到达同一点
对于主机来说，由于许多CUDA API调用和所有的内核启动不是同步的，cudaDeviceSyn-chronize函数可以用来阻塞主机应用程序，直到所有的CUDA操作（复制、核函数等）完成：
这个函数可能会从先前的异步CUDA操作返回错误。
因为在一个线程块中线程束以一个未定义的顺序被执行，CUDA提供了一个使用块局部栅栏来同步它们的执行的功能。使用下述函数在内核中标记同步点：
当__syncthreads被调用时，在同一个线程块中每个线程都必须等待直至该线程块中所有其他线程都已经达到这个同步点。在栅栏之前所有线程产生的所有全局内存和共享内存访问，将会在栅栏后对线程块中所有其他的线程可见。该函数可以协调同一个块中线程之间的通信，但它强制线程束空闲，从而可能对性能产生负面影响。
线程块中的线程可以通过共享内存和寄存器来共享数据。当线程之间共享数据时，要避免竞争条件。竞争条件或危险，是指多个线程无序地访问相同的内存位置。例如，当一个位置的无序读发生在写操作之后时，写后读竞争条件发生。因为读和写之间没有顺序，所以读应该在写前还是在写后加载值是未定义的。其他竞争条件的例子有读后写或写后写。当线程块中的线程在逻辑上并行运行时，在物理上并不是所有的线程都可以在同一时间执行。如果线程A试图读取由线程B在不同的线程束中写的数据，若使用了适当的同步，只需确定线程B已经写完就可以了。否则，会出现竞争条件。第4章会更深入地研究同步问题。
在不同的块之间没有线程同步。块间同步，唯一安全的方法是在每个内核执行结束端使用全局同步点；也就是说，在全局同步之后，终止当前的核函数，开始执行新的核函数。
不同块中的线程不允许相互同步，因此GPU可以以任意顺序执行块。这使得CUDA程序在大规模并行GPU上是可扩展的。
3.2.7 可扩展性
对于任何并行应用程序而言，可扩展性是一个理想的特性。可扩展性意味着为并行应用程序提供了额外的硬件资源，相对于增加的资源，并行应用程序会产生加速。例如，若一个CUDA程序在两个SM中是可扩展的，则与在一个SM中运行相比，在两个SM中运行会使运行时间减半。一个可扩展的并行程序可以高效地使用所有的计算资源以提高性能。可扩展性意味着增加的计算核心可以提高性能。串行代码本身是不可扩展的，因为在成千上万的内核上运行一个串行单线程应用程序，对性能是没有影响的。并行代码有可扩展的潜能，但真正的可扩展性取决于算法设计和硬件特性。
能够在可变数量的计算核心上执行相同的应用程序代码的能力被称为透明可扩展性。一个透明的可扩展平台拓宽了现有应用程序的应用范围，并减少了开发人员的负担，因为它们可以避免新的或不同的硬件产生的变化。可扩展性比效率更重要。一个可扩展但效率很低的系统可以通过简单添加硬件核心来处理更大的工作负载。一个效率很高但不可扩展的系统可能很快会达到可实现性能的上限。
CUDA内核启动时，线程块分布在多个SM中。网格中的线程块以并行或连续或任意的顺序被执行。这种独立性使得CUDA程序在任意数量的计算核心间可以扩展。
图3-18展示了CUDA架构可扩展性的一个例子。左侧的GPU有两个SM，可以同时执行两个块；右侧的GPU有4个SM，可以同时执行4个块。不修改任何代码，一个应用程序可以在不同的GPU配置上运行，并且所需的执行时间根据可用的资源而改变。
