cuda编程_CUDA编程4-执行模型（上）

1 概述

执行模型会提供一个操作视图，说明如何在特定的计算架构上执行指令。CUDA执行模型解释了GPU并行架构的抽象视图，有助于编写指令吞吐量和内存访问高效的代码。

1.1 GPU架构概述

GPU架构围绕一个流式多处理器（SM）的可扩展阵列搭建，通过复制这种架构的构件来实现硬件并行。每个GPU通常有多个SM，每个SM都能支持数百个线程并发执行。当启动一个内核网格时，它的线程被分布在可用的SM上执行。线程块一旦被调度到一个SM上，其中的线程只会在那个指定的SM上并发执行；多个线程块可能被分配到同一个SM上，且根据SM资源的可用性进行调度；同一线程中的指令利用指令级并行性进行流水化。

CUDA采用单指令多线程（SIMT）架构来管理和执行线程。每32个线程为一组，被称为线程束（warp）。线程束中所有线程同时执行相同的指令，每个线程都有自己的指令地址计数器和寄存器状态，利用自身数据执行当前指令。每个SM都将分配给它的线程划分到线程束中，然后在可用的硬件资源上调度执行。32在CUDA程序里是一个神奇的数字，它来自于硬件系统，对软件的性能有重要影响。它是SM用SIMD方式所同时处理的工作粒度。优化工作负载以适应线程束的边界，可以更有效地利用GPU资源。

SIMT架构与SIMD架构相似，两者都将相同的指令广播给多个执行单元来实现并行。一个关键区别是，SIMD要求同一向量的所有元素在一个同步组中一起执行，而SIMT则允许同一线程束的多个线程独立执行。尽管一个线程束中所有线程在相同的程序地址同时开始执行，但是单独的线程仍有可能有不同的行为。SIMT确保可以编写独立的线程级并行代码、标量线程、以及用于协调线程的数据并行代码。SIMT模型包含3个SIMD所不具备的关键特征：

• 每个线程都有自己的指令地址计数器；
• 每个线程都有自己的寄存器状态；
• 每个线程都可以有一个独立的执行路径。

一个线程块只能在一个SM上被调度，一旦线程块在一个SM上被调度，就会保存在该SM上直到执行完成。下图从逻辑视图和硬件视图的角度描述了CUDA编程对应的组件。在SM中，共享内存和寄存器是非常重要的资源。共享内存被分配在SM上的常驻线程块中，寄存器在线程中被分配。线程块中的线程通过这些资源可以进行相互合作和通信。尽管线程块里的所有线程在逻辑上都可以并行运行，但在物理层面并非都能同时执行，因此线程块里的线程可能会以不同的速度前进。

在并行线程中共享数据会引起竞争：多个线程以不定的顺序访问同一数据，会导致不可预测的程序行为。CUDA提供了一种方法来同步线程块里的线程，以保证所有线程在进一步动作之前都达到执行过程中的一个特定点，但没有提供块间同步的原语。

尽管线程块里的线程束可以任意顺序调度，但活跃线程束的数量还是会由SM的资源所限。当线程束因任何原因闲置时，SM可以从同一SM上的常驻线程块中调度其他可用线程束。在并发的线程束间切换没有开销，因为硬件资源已经被分配到了SM上的所有线程和块中，所以最新被调度的线程束的状态已经被存储在SM上。

SM是GPU架构的核心，寄存器和共享内存是SM中的稀缺资源。CUDA将这些资源分配到SM中的所有常驻线程里，因此这些有限的资源限制了在SM上活跃的线程束数量，活跃的线程束数量对应于SM上的并行量。了解一些SM硬件组成的基本知识，有助于组织线程和配置内核执行以获得最佳性能。下面介绍两种GPU架构，分别是Fermi架构和Kepler架构。

1.2 线程架构示例

Fermi架构是第一个完整的GPU计算架构，其逻辑图如下所示。Fermi的特征是多达512个用于加速计算的CUDA核心，每个CUDA核心都有一个全流水线整数算数逻辑单元（ALU）和一个浮点运算单元（FPU），在这里每个时钟周期执行一个整数或浮点数指令。CUDA核心被组织到16个SM中，每个SM含有32个CUDA核心。Fermi架构有6个384位的GDDR5 DRAM存储器接口，支持多达6GB的全局 内存，这是许多应用程序关键的计算资源。主机接口通过PCIe总线与CPU相连。GigaThread引擎是一个全局调度器，用来分配线程块到SM线程束调度器上。

Fermi架构包含一个耦合的768KB的二级缓存，被16个SM共享。上图中的每个垂直矩形条表示一个SM，包括执行单元（CUDA核心）、线程束调度器、指令分派单元、共享内存/一级缓存、寄存器文件、加载/存储单元、特殊功能单元等。如下图所示，每个多处理器有16个加载/存储单元，允许每个时钟周期有16个线程（线程束的一半）计算源地址和目的地址。特殊功能单元（SFU）执行固有指令，如正弦、余弦、平方根和插值。每个SFU在每个时钟周期内的每个线程上执行一个固有指令。

每个SM有两个线程束调度器和两个指令分派单元。当一个线程块被指定给一个SM时，线程块中的所有线程被分成线程束。两个线程束调度器选择两个线程束，再把一个指令从线程束中发送到一个组上，组里有16个CUDA核心、16个加载/存储单元和4个特殊功能单元，如下图所示。Fermi架构的计算性能为2.X，能在每个SM上同时处理48个线程束，即可在一个SM上同时常驻1536个线程。

Fermi架构的一个关键特征是有一个64KB的片内可配置存储器，它在共享内存与一级缓存之间进行分配。对于许多高性能应用程序，共享内存是影响性能的一个关键因素。共享内存允许一个块上的线程相互合作，这有利于芯片内数据的广泛重用，并大大降低了片外的通信量。根据给定内核中共享内存或缓存的使用，修改片内存储器的配置，可以提高性能。

Fermi架构也支持并发内核执行：在相同GPU上执行相同应用程序的上下文中，同时启动多个内核。并发内核执行允许执行一些小的内核程序来充分利用GPU，如下图所示。Fermi架构允许同时多达16个内核同时在设备上运行。从程序员角度看，并发内核执行使得GPU表现得更像MIMD架构。

Kepler架构是Fermi之后的一代架构，其特点是使得混合计算更加容易理解，包含了强化的SM、动态并行、一级Hyper-Q技术三项重要创新。下图展示了Kepler K20X芯片的框图，包含了15个SM和6个64位内存控制器，提供了超过1TFlop的峰值双精度计算能力。

如下图所示，Kepler架构的关键部分是新的SM单元，包含了一些结构创新，以提高编程和功耗效率。每个Kepler的SM包含192个单精度核心、64个双精度单元、32个特殊功能单元、以及32个加载/存储单元；此外还包括4个线程束调度器和8个指令调度器，以确保在单一SM上同时发送和执行4个线程束；Kepler架构的计算能力是3.5，每个SM上可同时调度64个线程束，即可常驻2048个线程；其寄存器文件容量达到64KB；同时还允许片内存储器在共享内存和一级缓存间有更多的分区。

动态并行是Kepler的另一新特性，它允许GPU动态启动新的网格，这样任一内核都能启动其他内核，并管理任何核间需要的依赖关系。这一特点可以更容易地创建和优化递归以及数据相关的执行模式。下图展示了没有动态并行时主机在GPU上启动每一个内核的情况、以及有动态并行GPU能启动嵌套内核从而消除与GPU的通信需求。

Hyper-Q增加了CPU与GPU之间的硬件连接，使CPU核心能够在GPU上同时运行更多任务，在增加GPU利用率的同时，减少CPU闲置时间。Fermi GPU依赖单一硬件工作队列来从CPU到GPU传输任务，这可能会导致某个任务阻塞队列中其后的所有任务。Kepler Hyper-Q消除了这种限制，它在主机与GPU之间提供了32个硬件工作队列，保证了GPU上有更多并发执行，如下图所示，最大限度提高了GPU的利用。

2 线程束执行本质

当启动内核时，从逻辑上看，内核中所有线程都并行运行；但从硬件上看，并非所有线程都在物理上可以同时并行执行。

2.1 线程束和线程块

线程束是SM的基本执行单元。当一个网格启动后，网格中的线程块分布在SM中；一旦线程块被调度到一个SM上，线程块中的线程束会被进一步划分为线程束；每个线程束由32个连续线程组成，一个线程束中的所有线程按单指令多线程（SIMT）方式执行：所有线程执行相同指令，每个线程在私有数据上执行操作。下图展示了线程块的逻辑视图和硬件视图之间的关系：

线程块在逻辑上可以被配置为一维、二维或三维，但从硬件角度看，所有线程都被组织成一维形式。在一个块中每个线程都有唯一ID，如果将x作为最内层维度，y作为第二个维度，z作为最外层维度，则二维或三维逻辑块布局可以转化为一维物理布局，且拥有连续值的线程被分组到线程束中。

• 对于一维线程块，唯一的线程ID被存储在CUDA内置变量threadIdx.x中；
• 对于二维线程块，线程ID可根据公式threadIdx.y*blockDim.x+threadIdx.x计算；
• 对于三维线程块，则可以根据threadIdx.z*blockDim.y*blockDim.x+ threadIdx.y*blockDim.x+threadIdx.x计算。

一个线程块中线程数的数量可以根据ceil(num_threads/warp_size)获得，因此硬件总是给一个线程块分配一定数量的线程束。线程束不会在不同线程块之间分离，如果线程块大小不是线程束的偶数倍，则在最后的线程束中有些线程就不会活跃。下图是一个在x为40、y轴为2的二维线程块，整个块中有80个线程。硬件为这个线程配置了3个线程束，使总共96个硬件线程去支持80个软件线程，最后半个线程束不活跃。但即使这些线程未被使用，它们仍然消耗SM资源，如寄存器。

总之，从逻辑角度来看，线程块是线程集合，可以被组织为一维、二维或三维布局；但从硬件角度看，线程块是一维线程束的集合，在线程块中线程被组织成一维布局，每32个连续线程组成一个线程束。

2.2 线程束分化

控制流是高级语言的一种基本构造，GPU支持传统、C风格、显式的控制流结构，如if...then...else、for和while。CPU拥有复杂的硬件结构来执行分支预测，即在每个条件检查中预测控制流会使用哪个分支，若预测正确，则只需付出很小的性能代价；若预测不正确，则可能会停止运行很多周期，因为指令流水线被清空了。

GPU则没有复杂的分支预测机制，一个线程束中的所有线程在同一周期中必须执行相同的指令，如果一个线程执行一条指令，则线程束中的所有线程都必须执行该指令。若同一线程束中的线程使用不同的路径通过同一个应用程序，则有可能产生问题。比如下面的语句：

if(cond)
{ ... }
else
{ ... }

若一个线程束有16个线程执行这段代码时cond为True，另外16个为False，则有一半的线程束需要执行if语句块中的指令，另一半需要执行else语句块的指令。在同一线程束中的线程执行不同的指令，被称为线程束分化。前面已经说过，一个线程束中所有线程在每个周期必须执行相同的指令，所以线程束分化会产生一个悖论。

如果一个线程束中的线程产生分化，线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。因此线程束分化会导致性能明显地下降， 在前面的例子中线程束中并行线程数量减少了一半。而条件分支越多，并行性削弱越严重，下图显示了线程束分化。需要注意的是，线程束分化只发生在同一线程束中，不同线程束中不同的条件值不会引起线程束分化。

为了获得最佳性能，应该避免在同一线程束中有不同的执行路径。需要记住在一个线程块中，线程的线程束分配是确定的，因此以这样的方式对数据区分是可行的，以确保同一个线程束中的所有线程子在一个应用程序中使用同一控制路径。

下面的代码中，mathKernel1使用一个奇数和偶数线程方法来模拟简单的数据分区，目的是导致线程束分化；mathKernel2使用线程束方法来交叉存储数据来避免线程束分化，这个核函数产生相同的输出，但顺序不同。

__global__ void mathKernel1(float *c)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a = 0.0f, b = 0.0f;
    if(tid % 2 == 0)
        a = 100.f;
    else
        b = 200.f;
    c[tid] = a + b;
}

__global__ void mathKernel2(float *c)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a = 0.0f, b = 0.0f;
    if((tid/warpSize) % 2 == 0)
        a = 100.f;
    else
        b = 200.f;
    c[tid] = a + b;
}

2.3 资源分配

线程束的本地上下文主要由程序计数器、寄存器、共享内存这几个资源组成，由SM处理的每个线程束的执行上下文，在整个线程束的生存期中是保存在芯片内的，因此从一个执行上下文切换到另一个执行上下文没有损失。

每个SM都有32位的寄存器，它存储在寄存器文件中，可以在线程中进行分配，同时有固定数量的共享内存在线程块中进行分配。对于一个给定的内核，同时存在于同一个SM中的线程块和线程束的数量，取决于在SM中可用的且内核所需的寄存器和共享内存数量。

如上图所示，若每个线程消耗的寄存器越多，则可以放在一个SM中的线程束就越少；若可以减少内核消耗寄存器的数量，就可以同时处理更多的线程束。如下图所示，若一个线程块消耗的共享内存越多，则在一个SM中可同时处理的线程块就会越少。如果每个线程块使用的共享内存数量变少，则可以同时处理更多的线程块。

资源可用性通常会限制SM中常驻线程块的数量，每个SM中寄存器和共享内存的数量因设备的计算能力不同而不同。如果SM中没有足够的寄存器或共享内存去处理至少一块，则内核无法启动。当计算资源已分配给线程块时，线程块被称为活跃块，它所包含的线程束被称为活跃线程束。活跃线程束可以进一步分为下面三种类型：

• 选定线程束：一个SM上的线程束调度器在每个周期都选择活跃线程束将它们调度到执行单元，活跃执行的线程束被称为选定线程束；
• 符合条件线程束：如果一个活跃线程束准备执行但尚未执行，则是一个符合条件线程束；
• 阻塞线程束：如果一个线程束没有做好执行准备，则是一个阻塞线程束，如果同时满足这两个条件：有32个CUDA核心可用于执行、当前指令中的所有参数都已就绪，则线程束符合执行条件。

例如，在Kepler SM上活跃的线程束数量，从启动到完成在任何时候都必须小于或等于64个并发线程束的架构限度；在任何周期，选定线程数量都小于等于4；如果线程束阻塞，线程束调度器会令一个符合条件线程束去代替它执行；由于计算资源在线程束之间进行分配，而且线程束在整个生存期中都保持在芯片内，因此线程束上下文切换非常快。

总之，在CUDA编程中需要特别关注计算资源分配：计算资源限制了活跃线程束数量，因此必须了解由硬件产生的限制和内核用到的资源。为最大程度地利用GPU，需要最大化活跃线程束数量。

2.4 延迟隐藏

SM依赖线程级并行来最大化功能单元的利用率，因此利用率与常驻线程的数量直接相关。在指令发出与完成之间的时钟周期被定义为指令延迟。当每个时钟周期中所有线程调度器都有一个符合条件线程束时，可以达到计算资源的完全利用。这样就能保证，通过在其他常驻线程束中发布其他指令，可以隐藏每个指令的延迟。

延迟隐藏在CUDA编程中十分重要，GPU是为处理大量并发和轻量级线程以最大化吞吐量而设计的，GPU的指令延迟被其他线程束的计算隐藏。按照指令延迟，指令可分为两种基本类型：

• 算数指令：指令延迟是一个算数操作从开始到产生输出之间的时间，大约10～20个周期；
• 内存指令：指令延迟是指发出的加载或存储操作和数据到达目的地之间的时间，全局访存为400～800个周期。

下图表示线程束0阻塞执行流水线的示例，线程束调度器选取其他线程束执行，当线程束0符合条件时再执行：

利特尔法则（Little's Law）提供了一个隐藏延迟所需活跃线程数量的合理近似值，它源于队列理论中的定理，也可应用于GPU中：

，下图形象说明了利特尔法则：

需要注意吞吐量和带宽两个概念，它们经常被混淆，根据实际情况它们可被交换使用，都是用来度量性能的速度指标。带宽通常是理论峰值，用来描述单位时间内最大可能的数据传输量；而吞吐量是指已达到的值，用来描述单位时间内任何形式的信息或操作的执行速度，如每个周期完成了多少指令。

2.5 占用率

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

。

函数cudaGetDeviceProperties 中，来自设备的统计数据在cudaDeviceProp结构中返回，每个SM的最大线程数量在maxThreadsPerMultiProcessor中，将其数值除以32，就可以得到最大线程束数量。

每个线程的存储器和每个块的共享内存资源的使用情况可以从nvcc中用--ptxas-options=-v标志获得。

2.6 同步

栅栏同步是一个原语，在许多并行编程语言中都十分常见。在CUDA中，同步可以在两个级别执行：

• 系统级：等待主机和设备完成所有工作；
• 块级：在设备执行过程中等待一个线程中的所有。

对于主机来说，许多CUDA API调用和所有内核启动都是异步的，cudaDeviceSynchronize函数可以用来阻塞主机应用程序，直到所有CUDA操作完成。这个函数可能会从先前异步CUDA操作返回错误。因为在一个线程块中线程束以一个未定义的顺序执行，CUDA提供了一个块局部栅栏来同步它们的执行功能，可以使用__device__ void __syncthreads(void)函数在内核中标记同步点。当__syncthreads被调用时，在同一个线程块中每个线程都必须等待直至该线程块中所有其他线程都达到这个同步点。在栅栏之前所有线程产生的所有全局内存和共享内存访问，将会在栅栏之后对线程块中所有其他线程可见。该函数可以协调同一个块中线程之间的通信，但它强制线程束空闲，因此可能对性能产生负面影响。线程块中的线程可以通过共享内存和寄存器来共享数据，当线程之间共享数据时，要避免竞争条件。

在不同块之间没有没有线程同步，块间同步唯一安全的方法是在每个内核执行结束端使用全局同步点，即在全局同步后，终止当前核函数，开始执行新核函数。不同块中线程不允许相互同步，因此GPU可以任意顺序执行块，这使得CUDA程序在大规模并行GPU上是可扩展的。

2.7 可扩展性

对任何并行程序而言，可扩展性是一个理想特性。可扩展性意味着为并行应用程序提供了额外的硬件资源，相对于增加的资源，并行应用程序会产生加速。一个可扩展的并行程序可以高效地使用所有的计算资源以提高性能。可扩展性意味着增加的计算核心可以提高性能，串行代码本身是不可扩展的，并行代码有可扩展的潜能，但真正的可扩展性取决于算法设计和硬件特性。

能够在可变数量的计算核心上执行相同应用程序代码的能力被称为透明可扩展性。一个透明可扩展平台拓宽了现有应用程序的应用范围，并减少了开发人员的负担，因为它们可以避免新的或不同的硬件产生的变化。可扩展性比效率重要，一个可扩展但效率很低的系统可以通过简单添加硬件核心来处理更大的工作负载。一个效率很高但不可扩展的系统可能很快会达到可实现性能的上限。

CUDA内核启动时，线程块分布在多个SM中，网格中的线程块以并行或连续或任意的顺序被执行，这种独立性使得CUDA程序在任意数量的计算核心间可以扩展。下图展示了CUDA架构可扩展性的一个例子。