5.1 简介

英伟达为它的硬件调度方式选择了一种比较有趣的模型，即SPMD(单程序多数据Single Program，Multiple Data)，属于SIMD(单指令多数据)的一种变体。从某些方面来说，这种调度方式的选择是基于英伟达自身底层硬件的实现。并行编程的核心是线程的概念，一个线程就是程序中的一个单一的执行流，就像一件衣服上的一块棉，一块块棉交织在一起织成了衣服，同样，一个个线程组合在一起就形成了并行程序。CUDA的编程模型将线程组合在一起形成了线程束、线程块以及线程网格。本章，就让我们一起来详细了解这些概念。

5.2 线程

线程是并行程序的基本构建块。对大多数做过多核程序设计的C程序员而言，这个概念并不陌生。即使一个程序员从来没有在代码中发起过一个线程，起码也执行过一个线程，因为任何序列化的代码都是以单线程的方式执行的。

随着双核、四核、十六核甚至更多核的处理器的出现，我们将更多的注意力放到了程序员如何充分利用这些硬件上。除了近十年，在过去的很多年里，大多数程序员编写的程序都是单线程的，因为当时运行程序的CPU也是单核的。当然，你可以利用更多的硬件设备或者成千上万的商用服务器来取代少数的实力强劲的机器，从而通过集群式计算机和超级计算机的方式来尝试进行一个更高层次的并行编程。然而，这些大多仅限于一些大学和大型机构使用，一般无法提供给大众使用。

实现多线程并行很难，但实现一次执行一个任务却要简单得多。在当时，每隔几年，串行处理速度就会提升一倍，因此基本没有什么需要去进行困难的并行编程。串行编程语言C/C++就在这样一个时代应运而生。直至大约十年前，这种情况发生了变化。现在，不管你接不接受，要想提高程序的速度，就必须考虑并行设计。

5.2.1 问题分解

CPU领域的并行化是向着一个CPU上执行不止一个(单一线程)程序的方向发展。但这只是我们之前所提到的任务级的并行。一般而言，程序拥有比较密集的数据集，例如，视频编码。对这种程序我们可以采用数据并行模型，将任务分解成N个部分，每个部分单独处理，其中，N代表可供使用的CPU核数。例如，你可以让每个CPU核计算一帧的数据，帧与帧之间没有相互的关联。又或者，你可以选择将每一帧分成N个片段，将每个片段的计算分配到每个独立的核中。

在GPU的领域，恰好能看到这些选择方案，当我们尝试加速渲染3D游戏中的现实场景时，就会采用多GPU的方式。你可以交替发送完整的帧数据到每个GPU(如图5-1)。此外你也可以让一个 GPU 渲染屏幕不同的部分。

 

然而，这里有一点需要我们来权衡。如果数据集是独立的，通过向GPU(或者CPU)提供需要计算的子数据集，我们可以只需要很小的内存来传递少量的数据。如果使用分割帧渲染模式(SFR)，那么对于渲染地面的GPU3来说，就没有必要去知道渲染天空的GPUO里面的数据内容。然而，在实际场景渲染中，有时地面上可能会有一些飞行物的影子，地面的照明度也会随着一天中时间的不同而发生变化。如果出现这种情况，由于共享数据的存在，此时使用 (Alternate Frame Rendering，AFR)可能对渲染更有益。

此处所说的 SFR是根据粗粒度的并行度来划分的，就是以某种方式将大块的数据分配到N个强劲的设备中，在数据处理之后又将它们重构成一整块。当我们为一个并行环境设计应用程序时，这一步的选择非常重要，它将严重影响到程序的性能。通常，最好的选择与所使用的设备密切相关。在后文中，你将看到多个贯穿全书的应用程序。

当只有数量较少的强劲设备时，例如在CPU上，我们的中心议题是解决平均分配工作量的问题。当然，这个问题很好解决，因为毕竟设备的数量较少。但如果像GPU那样拥有大量较小设备时，尽管也能很好的平均工作量，但我们却需要花大量的精力在同步和协调上。

世界经济有宏观(大规模)和微观(小规模)经济，相应地，并行也有粗粒度的和细粒度的并行。然而，只有在那些支持大量线程的设备上才能真正实现细粒度的并行，例如，GPU。相比之下，CPU同样支持线程，但伴随着大量的开销，因此它只适合解决粗粒度的并行问题。CPU与GPU不同，它遵从多指令多数据(MIMD)模型，即它可以支持多个独立的指令流。这是一种更加灵活的方式，但由于这种方式是获取多个独立的指令流，而不是平摊多个处理器的单指令流，因此它会带来额外的开销。

在此背景下，让我们来考虑用一个图像校正函数来增强数码照片的亮度。如果在GPU上，你可能会为照片上的每个像素点分配一个线程。但如果是在一个四核的CPU上，你可能会为每个 CPU 核分配照片的1/4图像的数据进行处理。

5.2.2 CPU与GPU的不同

GPU和CPU设备的架构是迴异的。CPU的设计是用来运行少量比较复杂的任务。GPU的设计则是用来运行大量比较简单的任务。CPU的设计主要是针对执行大量离散而不相关任务的系统。而GPU的设计主要是针对解决那些可以分解成成千上万个小块并可独立运行的问题。因此，CPU适合运行操作系统和应用程序软件，即便有大量的各种各样的任务，它也能够在任何时刻妥善处理。

    CPU与GPU支持线程的方式不同。CPU的每个核只有少量的寄存器，每个寄存器都将在执行任何已分配的任务中被用到。为了能执行不同的任务，CPU将在任务与任务之间进行快速的上下文切换。从时间的角度来看，CPU上下文切换的代价是非常昂贵的，因为每一次上下文切换都要将寄存器组里的数据保存到 RAM 中，等到重新执行这个任务时，又从 RAM中恢复。相比之下，GPU同样用到上下文切换这个概念，但它拥有多个寄存器组而不是单个寄存器组。因此，一次上下文切换只需要设置一个寄存器组调度者，用于将当前寄存器组里的内容换进、换出，它的速度比将数据保存到RAM 中要快好几个数量级。

CPU和GPU都需要处理失速状态。这种现象通常是由I0操作和内存获取引起的。CPU在上下文切换的时候会出现这种现象。假定此时有足够多的任务，线程的运行时间也较长，那么它将正常地运转。但如果没有足够多的程序使CPU处于忙碌状态，它就会闲置。如果此时有很多小任务，每一个都会在一小段时间后阻塞，那么CPU将花费大量的时间在上下文切换上，而只有少部分时间在做有用的工作。CPU的调度策略是基于时间分片，将时间平均分配给每个线程。一旦线程的数量增加，上下文切换的时间百分比就会增加，那么效率就会急剧的下降。

GPU就是专门设计用来处理这种失速状态，并且预计这种现象会经常发生。GPU采用的是数据并行的模式，它需要成千上万的线程，从而实现高效的工作。它利用有效的工作池来保证一直有事可做，不会出现闲置状态。因此，当GPU遇到内存获取操作或在等待计算结果时，流处理器就会切换到另一个指令流，而在之后再执行之前被阻塞的指令。

CPU和GPU的一个主要差别就是每台设备上处理器数量的巨大差异。CPU是典型的双核或者四核设备。也就是说它有一定数量的执行核可供程序运行。而目前费米架构的GPU拥有16个SM(流多处理器)，每个SM可看作是CPU的一个核。CPU通常运行的是单线程的程序，即它的每个核的每次迭代仅计算一个数据。然而，GPU默认就是并行的模式，它的 SM 每次可同时计算 32个数而不是像CPU那样只计算一个数，因此，相对于一个四核的CPU来说，GPU的核数目就是其4倍，数据的吞吐量则是其32倍。当然，你可能会说CPU也可以使用所有的可供使用的计算核，以及像MMX、SSE和 AVX那样的指令扩展集，但问题是又有多少CPU 程序使用了这种扩展集呢。  

GPU 为每个 SM 提供了唯一并且高速的存储器，即共享内存。从某些方面来说，共享内存使用了连接机和cell 处理器的设计原理，它为设备提供了在标准寄存器文件之外的本地工作区。自此，程序员可以安心地将数据留在内存中，不必担心由于上下文切换操作需要将数据移出去。另外，共享内存也为线程之间的通讯提供了重要机制。

5.2.3 任务执行模式

任务执行的模式主要有两种。一种基于锁步(lock-step)思想，执行N个SP(流处理器)组，每个SP都执行数据不同的相同程序。另一种则是利用巨大的寄存器文件，使线程的切换高效并且达到零负载。GPU能支持大量的线程就是按照这种方式设计的。

所谓的锁步原则到底是什么?指令队列中的每条指令都会分配到SM的每个SP中。每个SM就相当于嵌人了N个计算核心(SP)的处理器。

传统的 CPU会将一个单独的指令流分配到每个CPU核心中，而GPU所用的SPMD模式是将同一条指令送到 N个逻辑执行单元，也就是说 GPU 只需要相对于传统的处理器 1/N的指令内存带宽。这与许多高端的超级计算机中的向量处理器或单指令多数据处理器很相似。

然而，这样做并不就意味着没有开销。通过后面的学习我们将看到，当个线程执行相同的控制流，如果程序未遵循整齐的执行流，对于每一个分支而言，将会增加额外的执行周期。

图5-3 锁步指令分配

5.2.4 GPU线程

现在，我们再回过头来看看线程。首先，来看一段代码，从编程的角度看看它有什么意义。

这段代码很简单。它让数组b和数组c中下标相同的元素进行相乘，然后将所得的结果保存到相同下标的数组a中。串行代码需要128次for循环(从0~127)。而在CUDA 中我们可以将这段代码直接转换成用128个线程，每个线程都执行下面这段代码:

a[i]= b[i]*c[i];

由于循环中每一轮计算与下一轮计算之间没有依赖，因此将这段代码转换成并行程序非常简单。这种并行转换叫做循环并行化。这种并行化是另一种流行的并行语言扩展openMP 的基础。

在一个四核的CPU上，你可以将此计算任务平均分成四部分，让CPU的第一个核计算数组下标为0~31的元素，第二个核计算下标为32~63的元素，第三个核计算下标为64~95的元素，第四个核计算下标为96~127的元素。有些编译器自动就可以做这样的并行划分，而有些则需要程序员在程序中指出哪些循环需要并行。其中，Intel的编译器就非常擅长此道。这种编译器可以按照这种方式，而不是通过增加线程数量的方式，产生嵌人式SSE 指令以使循环矢量化。GPU的并行模式与我们所说的这两种并行模式相差并不太多。

在CUDA中，你可以通过创建一个内核函数的方式将循环并行化。所谓的内核函数就是一个只能在 GPU上执行而不能直接在 CPU上执行的函数。按照CUDA的编程模式CPU将主要处理它所擅长的串行代码。当遇到密集计算的代码块时，CPU则将任务交给GPU，让 GPU利用它超强的计算能力来完成密集计算。应该还有人记得CPU曾经搭载浮点协处理器的那段时光吧，应用程序在装有浮点协处理器的机器进行大量的浮点计算异常的快。而 GPU也是如此，它们就是用来加速程序中运算密集的模块的。

    从概念上看，GPU的内核函数和循环体是一样，只不过将循环的结构移除了。下面这段代码就是一个内核函数:

    仔细观察你会发现循环结构没有了，循环控制变量i也没有了。除此之外，在的函数前面还多了一个__global__的前缀。__global__前缀是告诉编译器在编译这个函数的时候生成的是GPU代码而不是CPU代码，并且这段GPU代码在CPU上是全局可见的。

CPU和GPU有各自独立的内存空间，因此在GPU代码中，不可以直接访问CPU端的参数，反过来在CPU代码中，也不可以直接访问GPU端的参数。稍后，我们将介绍一种特殊的方法来解决这个问题。现在，我们只需要知道它们是在不同的存储空间。因此，我们之前申明的全局数组a，b，c全是在CPU端的内存中，GPU端的代码是无法直接访问的，所以我们必须在GPU端的内存中也声明这几个数组，然后将数据从CPU端复制到GPU端以GPU内存指针的方式传递给GPU的内存空间进行读写操作，在计算完毕之后，再将计算的结果复制回CPU端。这些步骤我们会在之后的章节中做-一详解。

    下-个问题是，i不再是循环控制变量，而是用来标识当前所运行的线程的一个变量。在此，我们将以线程的形式创建128个该函数的实例，而CUDA则提供了一个特殊的变量，它在每个线程中的值都不一样，使得它可以标识每一个线程。这就是线程的索引，即线程ID。我们可以直接将这个线程标号用作数组的下标对数组进行访问。这和MPI中获取程序优先级很相似。

    线程的信息是由一个结构体存储的。在这个例子中，我们只用到了这个结构体中的一个元素，因此，我们将它保存到一个名为thread_idx变量中，以避免每次都访问这个结构体。具体代码如下:

    注意，有些人可能会使用idx或tid 来保存线程的标号，因为这样更加简短方便如此一来，线程0中的threadidx值为0，线程1的为1，依此类推，线程127中的thread idx值为127。每个线程都进行了两次读内存操作，一次乘法操作，一次存储操作，然后结束。我们注意到，每个线程执行的代码是一样的，但是数据却不相同。这就是CUDA的核心--SPMD模型。

在openMP和MPI中，你可能找到与这相似的代码块。对一个给定的循环迭代，将线程标号或线程优先级提取出来并分配给每一个线程，然后在数据集中作为下标使用。

5.2.5 硬件初窥

现在我们知道每个SM中有N个核，那么我们该如何运行128个线程?与CPU 很相似，GPU的每个线程组被送到SM中，然后N个SP开始执行代码。在得到每个线程的标号之后的第一件事就是从数组b和数组c中各取一个数然后进行乘法运算。不幸的是，这不是立即发生的。实际上，当从存储子系统取得所需要的数之后，已经过去了400~600个GPU时钟周期。在这期间，这一组中的N个线程都将挂起。

事实上，线程都是以每32个一组，当所有32个线程都在等待诸如内存读取这样的操作时，它们就会被挂起。术语上，这些线程组叫做线程束(32个线程)或半个线程束(16个线时，程)，这个概念在后面的内容中将会介绍。

因此，我们可以将这128个线程分成4组，每组32个线程。首先让所有的线程提取线程标号，计算得到数组地址，然后发出一条内存获取的指令(如图5-4所示)。接着下一条指令是做乘法，但这必须是在从内存读取数据之后。由于读取内存的时间很长，因此线程会挂起。当这组中的32个线程全部挂起，硬件就会切换到另一个线程束。

图5-4 周期0

在图5-5中我们可以看到，当线程束0由于内存读取操作而挂起时，线程束1就成为了正在执行的线程束。GPU一直以此种方式运行直到所有的线程束到成为挂起状态(如图5-6所示)。

当连续的线程发出读取内存的指令时，读取操作会被合并或组合在一起执行。由于硬件在管理请求时会产生一定的开销，因此这样做将减少延迟(响应请求的时间)。由于合并，内存读取会返回整组线程所需要的数据，一般可以返回整个线程束所需要的数据。在完成内存读取之后，这些线程将再次置成就绪状态，当再次遇到阻塞操作时，例如另一个线程束进行内存读取，GPU 可能将这个线程束用作另一块内存的读取。当所有的线程束(每组32个线程)都在等待内存读取操作完成时，GPU将会闲置。但到达某个时间点之后，GPU将从存储子系统返回一个内存块序列，并且这个序列的顺序通常与发出请求的顺序是一致的。

假设数组下标为0~31的元素在同一时间返回，线程束0进入就绪队列。如果当前没有任何线程束正在执行，则线程束0将自动进入执行状态(如图5-4所示)。渐渐地其他所有挂起的线程束也都完成了内存读取操作，紧接着它们也会返回到就绪队列。一旦线程束0的乘法指令执行完毕，它就只剩下一条指令需要执行，即将计算得到的结果写入相同下标的数组a中。由于再没有依赖该操作的其他指令，线程束0全部执行完毕然后消亡。其他的线程束也像这样，最终发出一条写数据的请求，完成之后便消亡。当所有的线程束都消亡之后，整个内核函数也就结束了，最终将控制返回到CPU端。

图5-5 周期 1

图5-6  周期 8

图5-7 周期9

5.2.6 CUDA内核

现在，我们来仔细介绍一下如何调用一个内核。CUDA专门定义了一个C语言的扩展用以调用内核。牢记，一个内核仅仅是一个运行在GPU上的函数。调用内核时必须按照以下语法:

kernel function<<<num blocks, num threads>>>(paraml, param2,...)

一个内核函数中可以传递很多参数，至于如何传递，我们稍后再进行详细介绍。现在，我们来看看另外两个比较重要的参数，numblocks与numthreads。它们可以是实参也可以是形参。在这里，建议使用变量，因为在之后进行性能调优时用起来更加方便。

参数num_blocks现在还没涉及，在下一节中我们将进行详细的介绍。现在我们只需保证至少有一个线程块。

参数num_threads表示执行内核函数的线程数量。在这个例子中，线程数目即循环迭代的次数。然而，由于受到硬件的限制，早期的一些设备在一个线程块中最多支持512个线程，而在后期出现的一些设备中则最多可支持1024个线程。本例中，我们无须担心这个问题，但对任何现实的项目而言，这个问题必须注意。在接下来的小节中我们将介绍如何来解决这个问题。

内核调用的下一部分是参数的传递。我们可以通过寄存器或常量内存来进行参数传递,而具体是哪一种方式则视编译器而定。如果使用寄存器传参，每个线程用一个寄存器来传递一个参数。如果现在有128个线程，每个线程传递3个参数，那么就需要3 x 128 = 384个寄存器。这听起来很多，但其实在每个SM(流处理器)中至少有8192个寄存器，而且随着后续硬件的发展，可能会更多。因此，如果在一个SM 上只运行一个线程块，每个线程块中只有128个线程，那么每个线程就可以使用64个寄存器(8192个寄存器÷128个线程)。

尽管每个线程能使用64个寄存器，但一个SM上只运行一个含有128个线程的线程块并不是一个好方案。只要我们访问内存，SM 就会闲置。因此，只有在很少数的情况下，运算强度很强的时候，我们才会考虑选择这种用64个寄存器进行计算的方案。在实际的编程过程中，我们都会尽量避免SM闲置状态的出现。