CUDA C编程（四）理解线程束执行的本质

  对于我们来说，在内核中似乎所有的线程都是并行地运行的。这在逻辑上是正确的，但从硬件的角度来看，不是所有线程在物理上都是可以同时并行地执行。前面提到，线程束的概念是把32个线程划分到一个执行单元中，接下来从硬件的角度来介绍线程束执行，并能够获得指导内核设计的方法。

线 程 束 和 线 程 块

  线程束是SM中基本的执行单元。当一个线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。一个线程束由32个连续的线程组成，在一个线程束中，所有的县城按照单指令多线程(SIMT)方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进行操作。下图展示了线程块的逻辑视图和硬件视图之间的关系。

  然而，从硬件的角度来看，所有的线程都被组织成了一维的，线程块可以被配置成一维、二维或三维的。在一个块中，每个线程块都有一个唯一的ID。对于一维的线程块，唯一的线程ID被存储在CUDA的内置变量threadIdx.x中，并且，threadIdx.x中拥有连续值得线程被分组到线程束中。
  用x维度作为最内层的维度，y维度作为第二个维度，z作为最外层的维度，则二维或三维线程块的逻辑布局可以转化为一维物理布局。例如，对于一个给定的二维线程块，在一个块中的每个线程的独特标识符都可以用内置变量threadIdx和blockDim来计算：threadIdx.y * blockDim.x + threadIdx.x;对于一个三维线程块，计算如下：threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;一个线程块的线程束的数量可以根据下式确定：

  因此，硬件总是给一个线程块分配一定数量的线程束。线程束不会在不同的线程块之间分离。如果线程块的大小不是线程束大小的偶数倍，那么在最后的线程束里有些线程就不会活跃。从逻辑角度来看，线程块是线程的集合，它们可以被组织成一维、二维或三维布局。从硬件角度来看，线程块是一维线程束的集合。在线程块中线程被组织成一维布局，每32个连续线程组成一个线程束。

线 程 束 分 化

  控制流是高级编程语言的基本构造中的一种。GPU支持传统的、C风格的、显式的控制流，例如，if…then…else、for和while。CPU拥有复杂的硬件以执行分支预测，也就是在每个条件检测中预测应用程序的控制流会使用哪些分支。如果预测正确，CPU中的分支只需要付出很小的性能代价。如果预测不正确，CPU可能会停止运行很多个周期，因为指令流水线被清空了。GPU是相对简单的设备，他没有复杂的分支预测机制。一个线程束的所有线程在同一周期中必须执行相同的指令，如果一个线程执行一条指令，那么线程束中的所有线程都必须执行该指令。如果同一线程束中的线程使用不同的路径通过同一个应用程序，这可能会产生问题。如下：

  假设在一个线程束中有16个线程执行这段代码，cond为true，但对于其他16个来说cond为false。一半的线程束需要执行if语句块中的指令，而另一半需要执行else语句块中的指令。在同一线程束中的线程执行不同的指令，被称为线程束分化。我们已经知道，在一个线程束所有线程在每个周期中必须执行相同的指令，所以线程束分化似乎会产生一个悖论。
  如果一个线程束中的线程产生分化，线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。线程束分化会导致性能明显的下降。在前面的例子中可以看到，线程束中并行线程的数量减少了一半：只有16个线程同时活跃地执行，而其他16个被禁用了。条件分支越多，并行性削弱越严重。注意，线程束分化只发生在同一线程束中。在不同的线程束中，不同的条件值不会引起线程束分化。
  线程束分化如下图所示，在一个线程束中所有的线程必须采用if…then两个分支来表述。如果线程的条件为true,它将执行if子句；否则，当等待执行完成时，线程结束。

  为了获得最佳的性能，应该避免在同一线程束中有不同的执行路径。请记住，在一个线程块中，线程的线程束分配是确定的。因为，以这样的方式对数据进行分区是可行的，以确保在同一线程束中的所有线程在一个应用程序中使用同一控制路径。
  分支效率被定义为未分化的分支与全部分支之比，可以使用以下公式来计算，通过使用nvprof分析器，可以从GPU中获取指标，从而可以直接观察到线程束分化。其中，nvprof的branch_efficiency指标是用来计算的分支率的。

  线程束分化总结：当一个分化的线程采取不同的代码路径时，会产生线程束分化；不同的if-then-else分支会连续执行；尝试调整分支粒度以适应线程束大小的倍数，避免线程束分化；不同的分化可以执行不同的代码且无需以牺牲性能为代码。

资源分配

  线程束的本地执行上下文主要由以下资源组成：1.程序计数器；2.寄存器；3.共享内存。由SM处理的每个线程束的执行上下文，在整个线程束的生存期中是保存在芯片内的，因此，从一个执行上下文切换到另一个执行上下文没有损失。
  每个SM都有32位的寄存器，它存储在寄存器文件中，并且可以在线程中进行分配，同时固定数量的共享内存用来在线程块中进行分配。对于一个给定的内核，同时存在同一个SM中的线程块和线程束的数量取决于在SM中可用的且内核所需的寄存器和共享内存的数量。
  下图显示了若每个线程消耗的寄存器越多，则可以放在一个SM中的线程束越少，如果可以减少内核消耗寄存器的数量，那么就可以同时处理更多的线程束。若一个线程块消耗的共享内存越多，则在一个SM中可以被同时处理的线程块就会变少。如果每个线程块使用的共享内存数量变少，那么可以同时处理更多的线程块。


  资源可用性通常会限制SM中常驻线程块的数量。每个SM中寄存器和共享内存的数量因设备拥有不同的计算能力而不同。如果每个SM没有足够的寄存器或共享内存去处理至少一个块，那么内核将无法启动。
  当计算资源（如寄存器和共享内存）已分配给线程块时，线程块被称为活跃的块。它所包含的线程束被称为活跃的线程束。活跃的线程束可以进一步被分为以下3种类型：选定的线程块；阻塞的线程块；符合条件的线程块。一个SM上的线程束调度器在每个周期都选择活跃的线程束，然后把它们调度到执行单元。活跃执行的线程束被称为选定的线程束。如果一个活跃的线程束准备执行但尚未执行，它是一个符合条件的线程束。如果一个线程束没有做好执行的准备，它是一个阻塞的线程束。如果同时满足以下两个条件则线程束符合执行条件：32个CUDA核心可用于执行；当前指令中的所有参数都已就绪。
  例如，Kepler SM上的线程束数量，从启动到完成在任何时候都必须小于或等于64个并发线程束的架构限度。在任何周期中，选定的线程束数量都小于或等于4。如果线程束阻塞，线程束调度器会令一个符合条件的线程束代替它去执行。由于计算资源是在线程束之间进行分配的，而且在线程束的整个生存期中都保持在芯片中，因此线程束上下文的切换是非常快的。
  资源分配：在CUDA编程中需要特别关注计算资源分配：计算资源限制了活跃的线程束的数量。因此必须了解由硬件产生的限制和内核用到的资源。为了最大化利用GPU，需要最大化活跃的线程束数量。

延迟隐藏

  SM依赖于线程级并发，以最大化功能单元的利用率，因此，利用率与常驻线程束的数量直接相关。在指令发出和完成之间的时钟周期被定义为指令延迟。当每个时钟周期中所有的线程调度器都有一个符合条件的线程束时，可以达到计算资源的完全利用。这就可以保证，通过在其他常驻线程束中发布其他指令，可以隐藏每个指令的延迟。
  与在CPU上用C语言编程相比，延迟隐藏在CUDA编程中尤为重要。CPU核心是为了同时最小化延迟一个或两个线程而设计的，而GPU则是为处理大量并发和轻量级线程以最大化吞吐量而设计的。GPU的指令延迟被其他线程束的计算隐藏。
  考虑到指令延迟，指令可以被分为两种基本类型：算术指令；内存指令。算术指令延迟是一个算术操作从开始到它产生输出之间的时间。内存指令延迟是指发送出的加载或存储操作和数据到达目的地之间的时间。对于每种情况，相应的延迟为：算术延迟为10 ~ 20个周期；全局内存延迟为400~800个周期。下图表示线程束0阻塞执行流水线的一个示例。线程束调度器选取其他线程束执行，当线程束0符合条件时再执行它。

  我们一般利用特尔法则提供的一个合理的近似值来估算延迟隐藏所需要的活跃线程束的数量。它起源于队列理论中的一个定理，它也可以应用于GPU中：所需线程束数量 = 延迟 × 吞吐量。下图形象的说明了利特尔法则。假设在内核里一条指令的平均延迟是5个周期。为了保持在每个周期内执行6个线程束的吞吐量，则至少需要30个未完成的线程束。

  (PS:吞吐量和带宽的区别：都是用来度量性能的速度指标。带宽通常是理论峰值，而吞吐量是指以达到的值。带宽通常是用来描述单位时间内最大可能的数据传输量，而吞吐量是用来描述单位时间内任何形式的信息或操作的执行速度，例如，每个周期完成多少个指令)
  对于算术运算来说，其所需的并行可以表示成隐藏算术延迟所需要的操作数量。下表列出了Fermi和Kepler设备所需的操作数量。实例中的算术运算是一个32位的浮点数乘加运算(a+b×c)，表示在每个SM中每个时钟周期内的操作数量，吞吐量因不同的算术指令而不同。

  吞吐量由SM中每个周期内的操作数量确定，而执行一条指令的一个线程束对应32个操作，因此，为保持计算资源的充分利用，对于Fermi GPU而言，每个SM中所需的线程束数量通过计算为640÷32=20个线程数。因此，算术运算所需的并行可以用操作的数量或者线程束的数量来表示。这个简单的单位转换表明，有两种方法可以提高并行：指令级并行（ILP）:一个线程中有很多独立的指令；线程级并行（TLP）:很多并发地符合条件的线程。对于内存操作来说，其所需的并行可以表示为在每个周期内隐藏内存延迟所需的字节数。下表列出了Fermi和Kepler架构的指标。

  因为内存吞吐量通常是每秒千兆字节数，所以需要用对应的内存频率将吞吐量转换为每周期千兆字节数。公式如下：吞吐量÷内存频率=每周期千兆字节数。接下来用内存延迟乘以每周期字节数，可以得到内存操作所需的并行。其中，这个值是对于整个设备，而不是对于每个SM来说的，因为内存带宽是对于整个设备而言的。
  利用应用程序，把这些值与线程束或线程数量关联起来。假设每个线程都把一浮点数据(4个字节)从全局内存移动到SM中用于计算，则在Fermi GPU上，需要18500个线程或579个线程束来隐藏所有内存延迟，具体运算如下：
 74KB÷4字节/线程≈18500个线程
 18500个线程÷32个线程/线程束≈579个线程束
  Fermi架构有16个SM。因此，需要579个线程束÷16个SM=36个线程束/SM，以隐藏所有的内存延迟。如果每个线程执行多个独立的4字节加载，隐藏内存延迟需要的线程就可以更少。
  与指令延迟很像，通过在每个线程/线程束中创建更多独立的内存操作，或创建更多并发地活跃的线程/线程束，可以增加可用的并行。
  延迟隐藏取决于每个SM中活跃线程束的数量，这一数量由执行配置和资源约束隐式决定（一个内核中寄存器和共享内存的使用情况）。选择一个最优执行配置的关键是在延迟隐藏和资源利用之间找到一个平衡。
  因为GPU在线程间分配计算资源并在并发线程束之间切换的消耗（在一个或两个周期命令上）很小，所以所需的状态可以在芯片内获得。如果有足够的并发活跃线程，那么可以让GPU在每个周期内的每一个流水线阶段中忙碌。在这种情况下，一个线程束延迟可以被其他线程束的执行隐藏。因此，向SM显示足够的并行对性能是有利的。
  计算所需并行的一个简单公式是，用每个SM核心的数量乘以在该SM上一条算术指令的延迟。例如，Fermi有32个单精度浮点流水线线路，一个算术指令的延迟是20个周期，所以，每个SM至少需要有32×20=640个线程使设备处于忙碌状态，然而，这只是一个下边界。

占用率

  在每个CUDA核心里指令是顺序执行的。当一个线程束阻塞时，SM切换执行其他符合条件的线程束。理想情况下，我们想要有足够的线程束占用设备的核心。占用率是每个SM中活跃的线程束占最大线程束数量的比值。

  cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop,int device);函数可以检测设备中每个SM的最大线程束数量。来自设备的各种统计信息在cudaDeviceProp结构中被返回。每个SM中线程数量的最大值在maxThreadsPerMultiProcessor变量中返回，用其除以32，可以得到最大线程束数量。
  CUDA工具包包含了一个电子表格，他被称为CUDA占用率计算器，有助于选择网格和块的维数以使一个内核的占用率最大化。如下图所示：

  占用计算率包含几个部分。首先，必须提供GPU的计算能力和内核的资源使用情况的信息。在确定GPU的计算能力后，物理限制部分的数据都是自动填充的。接下来需要输入以下内核资源信息：每个块的线程（执行配置）；每个线程的寄存器（资源使用情况）；每个块的共享内存（资源使用情况）。
  为了提高占用率，还需要调整线程块配置或重新调整资源的使用情况，以允许更多的线程束同时处于活跃状态和提高计算资源的利用率。极端的操纵线程块会限制资源的利用：
  >小线程块：每个块中线程太少，会在所有资源被充分利用之前导致硬件达到每个SM的线程束数量的限制。
  >大线程块：每个块中有太多的线程，会导致在每个SM中每个线程可用的硬件资源太少。
  网格和线程块大小的准则：保证每个块中线程数量是线程束大小的倍数；避免块太小：每个块至少有128或者256个线程；根据内核资源调整块大小；块的数量要远多于SM的数量，从而在设备中可以显示有足够的并行；通过实验得到最佳执行配置和资源使用情况。

同步

  在CUDA中，同步可以在两个级别执行：系统级：等待主机和设备完成所有的工作；块级：在设备执行过程中等待一个线程块中所有线程到达同一点。对于主机而言，由于许多CUDA API调用和所有的内核启动不是同步的，cudaDeviceSynchronize函数可以用来阻塞主机应用程序，直到所有的CUDA操作完成。cudaError_t cudaDeviceSynchronize(void)；这个函数可能会从先前的异步CUDA操作返回错误。
  因为在一个线程块中线程束以一个未定义的顺序被执行，CUDA提供了一个块局部栅栏来同步他们的执行的功能。使用__device__ void _syncthreads(void);这个函数可以在内核中标记同步点。当_syncthreads被调用时，在同一个线程块中的每个线程都必须等待直至该线程块中所有其他线程都已经达到这个同步点。在栅栏之前所有线程产生的所有全局内存和共享内存访问，将会在栅栏后对线程块中所有其他的线程可见。该函数可以协调同一个块中线程之间的通信，但它强制线程束空闲，从而可能对性能产生负面的影响。
  线程块中的线程可以通过共享内存和寄存器来共享数据。当线程之间共享数据时，要避免竞争条件。竞争或危险，是指多个线程无序地访问相同地内存位置。在不同的块之间没有线程同步，块间同步，唯一安全的办法是在每个内核执行结束端使用全局同步点零零也就是说，在全局同步之后，终止当前的核函数，开始执行新的核函数。不同块中的线程不允许相互同步，因此GPU可以以任意顺序执行块。这使得CUDA程序在大规模并行GPU是不可行的。

可扩展性

  对于任何并行应用程序而言，可扩展性是一个理想的特性。可扩展性意味着并行应用程序提供了额外的硬件资源，相对于增加的资源，并行应用程序会产生加速。例如，若一个CUDA程序在两个SM中是可扩展的，则与在一个SM中运行相比，在两个SM中运行会使运行时间减半。一个可扩展的并行程序可以高效地使用所有的计算资源以提高性能。可扩展性意味着增加的计算核心可以提高性能。串行代码本身是不可扩展地，因为在成千上万地内核上运行一个串行单线程应用程序，对性能是没有影响的。并行代码有可扩展地潜能，但真正的可扩展性取决于算法设计和硬件特性。
  能够在可变数量的计算核心上执行相同的应用程序的能力被称为透明可扩展性。一个透明的可扩展平台拓宽了现有应用程序的应用范围，并减少了开发人员的负担，因为它们可以避免新的或不同的硬件产生的变化。可扩展性比效率更重要。一个可扩展但效率很忙的系统可以通过简单添加硬件核心来处理更大的工作负载。一个效率很高但不可扩展的系统可能很快会达到可实现性能的上限。
  CUDA内核启动时，线程块分布在多个SM中。网络中的线程块以并行或连续或任意的顺序被执行。这种独立性使得CUDA程序在任意数量的计算核心间可以扩展。
  下图展示了CUDA架构可扩展性的一个例子。左侧的GPU有两个SM，可以同时执行两个块；右侧的GPU有4个SM，可以同时执行4个SM，可以同时执行4个块。不修改任何代码，一个应用程序可以在不同的GPU配置上运行，并且所需的执行时间根据可用的资源而改变。