深入浅出说CUDA程序设计（一）

第一章 为什么需要并行程序

CUDA，全称是Compute Unified Device Architecture，一般翻译成中文为计算统一设备架构。笔者以为这样的名字会让人对CUDA感到很迷惑，CUDA到底是什么呢？笔者用自己的大白话来说下自己对CUDA的理解，CUDA就是一个基于GPU（Graphics processing unit）（目前是单指Nvidia公司的）的通用并行计算平台。这里有3个关键字，GPU，通用计算和并行！

关于GPU，相信它是什么，不用多说，不过关于CUDA的硬件架构后面会有一些分析，因为要写出高质量的CUDA程序，不了解底层运行机制是不够的，这就是侯捷讲的“勿在高台筑浮沙”，这里说一点题外话，笔者经常在同行老朋友面前很得意的吹嘘：“你们这些家伙写的程序都是在CPU上跑的，太土了！俺的可是在GPU上跑的哦”，所以嘛，学习CUDA是很有前途的！

第二个重要概念是通用计算，这个主要是与以前的GPU只做渲染流水线相对应的，由于笔者比较年轻，没有经历过师兄他们走过的黑暗岁月，呵呵！传统的GPU架构需要按一个标准的流水线编程，要经过vertex processor，fragment processor和pixel operation，这会使编程变得困难和不容易控制。早期的GPGPU（General Purpose Computing on GPU通用计算为目的的GPU）也需要按照那样的标准的流水线编程，而CUDA的出现对我们学软件的来讲，才标志着真正意义上GPU通用计算的到来！CUDA从技术本身来讲，仅仅是C/C++的一个小超集（这个后面会详细分析），从软件工程的角度看，降低了学习成本，加之其灵活性，必将在工程应用中发挥巨大的作用！在美国几乎所有顶级学府如哈弗，MIT等都有CUDA实验室，美国NASA的很多大型项目在多年前都已经使用CUDA技术。基本上，可以说CUDA提供了一个无论是在软件编程还是硬件处理都适合通用计算的平台架构。

在学习CUDA编程以前，我们必须深刻认识下CUDA程序中最核心的概念——并行。在操作系统的课程里面，我们会遇到另外一个名词——并发。在单核时代或者说对于普通的CPU运行程序，这两个词的意思是一模一样的！但我们应该清楚知道两者的区别，并发指复用，同一个时刻只有一个指令在处理器上执行；而并行则是在同一时刻多个指令同时运行，在GPU上处理器的称谓换成了个多处理器，每个多处理器上有8个名叫CUDA core的处理单元。在本科答辩的时候，有个老师问我，是不是可以说1个多处理器上有8个核（注：他不做CUDA的）。笔者以为作为初学者可以这样看，但如果想深入了解和掌握高级优化手段就必须明白CPU上的多核不等同于CUDA core，这个问题会在后面章节逐步说明。这里花点笔墨来说明一下并行与并发的区别，主要是让读者明白CPU并行程序和GPU并行程序在底层运行机制的区别，同时初学者往往遇到EmuDebug运行正确，而实际运行却错误的情况。这里最根本的原因还在于EmuDebug运行模式下是串行而非并行。虽然有很多牛人觉得EmuDebug没有意义，但我还是建议初学者遇到结果错误的时候先保证EmuDebug是正确的，因为绝大多数情况，保证EmuDebug是正确的说明程序逻辑没有问题（当然也有例外，后面会讲到）。

知道了并行的概念，那么为什么我们需要并行程序，根据摩尔定律：集成电路上可容纳的晶体管数目，约每隔18个月便会增加一倍，性能也将提升一倍，当价格不变时；或者说，每一美元所能买到的电脑性能，将每隔18个月翻两倍以上，如图1。他的核心思想就是硬件会越来越越快，可是随着单核CPU的时钟频率不断增加，也落入了收益递减规律里面了，其中一个主要的性能瓶颈是存储器延迟。因此我们有了新的摩尔定律：No longer get faster, just wider（未来的计算机不会更快，而是更“宽”）。新的需求也给我们这个做软件的带来了挑战和机遇：必须重新设计算法，To be aggressively parallel!笔者庆幸自己在不到两年的时光发表10来篇EI论文也得益于这样的变革机遇！在并行处理这个领域，并行体系结构、并行软件和并行算法三者缺一不可，而其中并行算法则是核心和瓶颈技术，也是我们从事软件行业的人应该肩负的使命。

在结束这个开场白的时候，我们来比较学术化的看看什么是并行算法。并行算法是指在各种并行计算机上求解问题和处理数据的算法，其本质是把多任务映射到多处理器中执行，或将现实的多维问题映射到具有特定的拓扑结构的多处理器上求解。一定要牢记的是并行算法的实现强烈的依赖于计算机硬件和软件环境，这是我们软件专业出身的人最容易忽视的，笔者个人比较特别的经历就是在系统开发时的直接合作师兄是硬件工程师，我们经常被对方弄的很崩溃。但自己现在能够写出比较高质量的CUDA程序也得益于这些对硬件问题的崩溃和思考！

 

图1 CPU chip进化速度

1.1 并行算法的目标

计算需求是永无止境的，可以说高性能计算是计算机科学研究中的“日不落”课题。并行计算是其中最有效的手段。作为软件编程人员，设计编写并行算法是最为核心的工作任务。笔者想从3个基本概念：时间重叠、资源重复和资源共享，来让读者初步直观的认识一下并行算法的总体设计目标。

首先，时间重叠是指多个处理过程在时间上相互错开，轮流重叠地使用同一套硬件设备的各个部分。这个概念可从计算复杂度的角度来理解，一个算法的复杂度可表示为空间复杂度和时间复杂度。从算法树的结构来看，通常的串行算法树“深而窄”，因为串行算法的本质是为一维问题设计的。而并行算法的目标则是尽可能减少时间复杂度，通常是通过以空间换时间的方式实现的，即增加空间复杂度。典型的时间重叠就是流水线处理。虽然CUDA平台上单个GPU暂时是不能设计流水线算法，但它也提供了异步访问以及Fermi架构的双kernel调度等时间重叠的处理方式。

除了时间重叠外，资源重复也可以实现将时间复杂度转化为空间复杂度。资源重复是指设置多个相同的处理器，同时从事处理工作，以数量取胜的方式提高处理速度。

这里给一个经典的以空间换时间的并行算法例子，寻找一个数列最大值。

我们知道在数列中寻找最大值的一般算法复杂度为O（n）但如果采用并行算法则可为O（1），假定有2n个处理单元，算法如下：

1. int i=Px     //范围0～n-1
2. int j=Py    //范围0～n-1
3.  
4. arrB[i] = 1
5.     _syncthreads()  //线程同步函数
6.         if arrA[i] < arrA[j]
7.                 arrB[i] = 0
8.    _syncthreads() 
9.                 if arrB[i] == 1
10.                 print arrA[i]

这个算法并行的思路分析：是很简单，当A[i]不是最大值时，B[i]标识为0，结合图2，相信很容易想明白。看到这个比较神奇的并行算法idea，会否激起你的挑战欲望呢！

 

图2

注意，这里给出的例子只是一个理论算法，要实现它必须假定处理系统是CRCW（Concurrent Read Concurrent Write（同时读同时写））系统而且处理器足够多（2n个），所以很难在实际中应用，当然也不是在CUDA平台上解决这一问题的理想算法。但这个经典算法体现了并行算法设计目标与传统串行程序的本质区别，同时它也揭示了并行算法的理论加速比，关于加速比在章节1.3会有详细说明。

由此，我们可以看到并行算法树采用的是与串行算法树截然不同的“浅而宽”的结构，即每个时刻可容纳的计算量相应增加，使整个算法的执行步骤数目尽可能接近问题的关键路径长度，也可以这样说，通过增加每个时刻步的算法复杂度来减少整体的时间复杂度，从而达到把时间复杂度转化为空间复杂度的目的。

最后由于数据在处理的过程中涉及到读取和写入操作，尤其是CUDA平台下存在多种访存类型，如主机端到设备端，GM（global memory全局存储器）到SM（share memory共享存储器）等等，资源共享也是并行处理技术重要概念之一，原始的定义是多个处理器按照一定的规则对同一数据资源进行访问，也包括处理器负载的均衡，即多个服务请求如何映射到多个处理器。笔者以为在CUDA平台下就是如何通过资源共享来减小访存延迟，提高存储器的带宽（见5.4黄三）。

1.2 并行算法设计的基本方法

在前面的小节里，我们已经从不同角度认识了什么是并行处理，并行程序的大体模样和设计理念。下面笔者将结合一些实例从总体上介绍下并行算法设计的基本方法。并行算法的一般设计方法大致包含以下3个途径：

（1）  检测和开发现有串行算法中固有并行性而直接将其并行化

（2）  从问题本身特征出发，设计一个新的并行算法

（3）  修改已有的并行算法使其可求解另一类相似问题

以上三个途径是最正统的说法，笔者以为凡事从事并行算法设计的同仁们都应当对它们“死去活来”。记住它们，因为它们是设计的钥匙；灵活运用，因为每一条都只是一个思想而不是一个技巧。笔者在本科毕业设计答辩时还给出了第四把钥匙：充分利用CUDA平台的各种技术特性。它不是一个设计准则而是所有设计的基础，在本书里面，笔者最强调的是是并行算法的实现强烈的依赖于计算机硬件和软件环境。为了让大家更形象化，笔者贴出自己的一页答辩PPT ，见图3。

首先，我们来看看第一个路径，是很长的标准化学术化说法。思想很简单，就是原先一个处理器一次处理一个，现在多个处理器同时处理多个。但注意，要做到这一点，需要有一个基础就是计算无关性，有很显式的，有半遮面的，也有隐式的。首先，我们还是不厌其烦的看看最简单的一个例子。

for(int i=0;i<256;++i)

              X[i]=i;

这个for循环，大家可以看到，计算X[i]与任何的X[i+n]没有任何的依赖关系，所以可以很轻易的写出并行处理程序。这里因为假定大家还没有CUDA基础，所以给出一个直观的伪代码：

//设置256线程

//让第i个线程取得值i

//让第i个线程将取得的值i写到X[i]这个位置

 

图3

上面这个例子是不是太easy了，那么稍稍加点难度，来看看下面这个例子：

X[0]=0.05;

for(i=1;i<256;i++)

   X[i]=X[i-1]+0.05;

从直观上看，与第一个例子不同，X[i]的得到是基于X[i-1]，这个说法对吗？笔者以为也对也不对，按传统的理解，这个程序的执行是相互依赖的。但稍稍想想这个犹抱琵琶半遮面的计算无关性就浮现于眼前了，并行方式如下：

//设置256线程

//让第i个线程计算（i+1）*0.05

//让第i个线程将计算结果写到X[i]这个位置

是不是也很容易呢！好吧，我们把这个例子在推广一下下，应该也就没有难度了,至少有了前面的热身不会太崩溃吧：

X[0]=b[0];

for(int i=1;i<256;++i)

{

   for(int j=0;j<i;++j)< span="">

   {

       b[i] += X[i-j-1]*a[i][j];         

   }

   X[i] = b[i];

}

这是一个线性递推公式，用数学符号表示如下：

  （1）

打了这么多铺垫，就不用笔者再说了吧！要并行化就是要将其展开，其实一点都不难，用笔划划就找到规律了，好吧，笔者来写写展开式。大家注意看看a和b的下标规律，是否想起了，咱们高一上学期数学的最后一章呢！等差数列，如果找到这个规律了，大家就会有感觉了！纵向看看我们的b，是不是阶梯形，呵呵！在横向看看b的系数a，好，我们给出两个定义式来说明下这个规律。

定义总是让人觉得学术化，但数学有时候是最直观的表达，由我们的展开式，可以得到

（2）

到了这里，这个例子的计算无关性也找到了，最直接的并行化处理（肯定不是最优，比如有大量的冗余计算可以去除等优化方法）可以这样描述如下：

//设置256线程

//让第i个线程计算公式（2）

//让第i个线程将计算结果写到X[i]这个位置

通过这三个难度等级不同的例子，可以看到挖掘现有算法中的计算无关性，虽然很直接但未必很简单，而且串行程序的固有并行性出现的方式各式各样，不一而足，需要我们的数学素养（如姜伯驹院士所讲，数学素养是国民素质的重要元素，笔者以为，然也！），同时设计经验也是我们的重要财富。

下一把钥匙是从问题本身特征出发，设计一个新的并行算法，这类情况主要是指问题本身不具备计算无关性，或者迭代或者数据相关。这里以CUDA SDK下的直方图生成算法为例。

传统的图像直方图算法是通过逐点统计实现的，基本程序算法如下：

for(int i = 0; i < BIN_COUNT; i++)

histogram[i] = 0;

for(int i = 0; i < dataNum; i++)

histogram[image[i]]++;

由于要确定的是整个图像的灰度分布情况，所以即使在并行处理环境中，不进行迭代运算也是无法生成图像的直方图的。所以最显然的算法就是分而治之，然后归并，也就是让每一个线程生成一个局部区域的直方图，然后加在一起形成全局直方图。如果不考虑性能，这样做是完全OK的，但可以很负责任的说，在CUDA平台下这样做是不好的，甚至是不行的。我的一个学生在做图像纹理分析时曾经很无奈的说，无法在CUDA平台下生成和差直方图（这个可能需要有点计算机视觉的背景才容易明白，因为对于图像的和差直方图可能会需要512个针脚）。

回到这个问题上来，问题到底出在哪里了呢？因为还没有完全介绍CUDA，这里只稍稍分析下，让某一个线程去处理其中一块形成一个局部直方图，最后归并成一个全局直方图。对于CUDA平台这个算法是非常不适宜的，因为GPU平台的存储器最优的是共享存储器而非全局存储器（见第五章黄三）。在2.0以下的CUDA设备上，共享存储器的大小为16KB，同时一个有效工作的线程块应该包含128～256个线程以获得相对高效的运行性能。那么如果采用一个线程直接对应于一个子直方图的话，一个明显的限制就是16KB平均到最大一个包含192条线程的块，每个线程只有85bytes，因此这种方法最大情况共享存储区能适合每线程子直方图可达64单字节针脚数。更为严重的是，单个执行线程所处理的数据大小也受限于字节计数器所规定的255字节，也就是说单个线程不能生成针脚数大于64的子直方图。从硬件上看，也就是受维护单个线程上下文存储器资源的限制。那么SDK下解决这个问题的方法是一个线程单干不可以，就依靠群体嘛，所谓团结就是力量，即以一个warp块（32个线程）输出一个局部区域的子直方图。多线程维护同一个直方图，在不使用原子操作的情况下（因为1.1以下设备的共享存储器不支持原子操作，笔者以为即使使用现在的高端设备，原子操作越少越好）是非常棘手的问题。我们来看看SDK的解决方案：

__device__ void addData256(volatile unsigned int *s_WarpHist,unsigned int data,unsigned int threadTag)

{

unsigned int count;

do{

count = s_WarpHist[data] & 0x07FFFFFFU;

count = threadTag | (count + 1);

s_WarpHist[data] = count;

}while(s_WarpHist[data] != count);

}

这个代码，除了关键字__device__（暂时不需要管它），其它全是普通的C代码，应该是无障碍理解吧！

算法解决访问冲突的技巧是很巧妙的，这里图像数据值的取值范围是[0，255]，每个warp线程根据输入的图像数据使数组s_WarpHist[]相应位置增加1。算法为了解决warp内的线程冲突，通过最末一个执行写入的线程将直方图标识，标识为直方图频数变量的高5位，因为warp的大小为32，所以只需要5位。这样的话，对于每一个线程读取当前直方图针脚值高5位就被替代为当前线程的标识符。如果线程间要写入的不是同一针脚位置，就没有额外的操作，但当两个或多个线程在写同一位置时硬件会执行共享存储区写结合，使线程中被标识的频数被接受而拒绝所有其它未定的线程修改。在写尝试后每个线程读取来自同一个共享内存的位置。而那个可以写回自己修改的频数的线程，则退出循环并空闲等待warp中其余的线程。Warp继续它的执行，直到所有线程退出循环。由于每个warp使用自己的子直方图并且warp线程总是同步的，所以不依赖于warp的调度顺序，这个循环执行的迭代次数不会超过32次，而且只有在warp内所有线程读到的灰度值相同时才会发生。这个解决访问冲突的算法主要包含三个步骤

第一步count = s_WarpHist[data] & 0x07FFFFFFU;是将当前该针脚的频数放入count变量中，并将高位清零。

第二步增加相应针脚频数，并写入线程标识符，这里特别说明的是threadTag = threadIdx.x << (32 - WARP_LOG_SIZE);也就是将线程号经过移位操作放到高5位作为标识符。

第三步则是将增加后的频数写回到数组s_WarpHist[]。然后判断当前写入的频数是不是最后修改频数的线程所写入的，是就接受，不是则通过循环重新执行上述操作。

在完成了子直方图计算后就是归并每一个子直方图，这个过程应该分两步首先归并各个warp的子直方图到块直方图，然后将块直方图归并为全局直方图。要理解为什么这个算法要费这么大的力气，需要掌握后面的内容，在？.？会有详细的分析。但笔者这里给出一个实例是想说明，问题本身不具备计算无关性，无法直接并行化的时候，可以针对问题本身的特征结合并行处理平台的技术特性来设计一个新的并行算法。

另外一个重要的设计原则，其实就是学会旧瓶装新酒，或者文言一点，它山之石，可以攻玉。这里就分析介绍下并行算法中最经典的例子，Fan-in算法。

所谓Fan-in算法求和就是利用多处理器实现树型求和，图1.4是一个树型求和的实例：

 

图1.4 Fan-in算法处理实例

可以看到该算法是利用多处理器的并行计算能力，将计算时间缩短，也可以很清楚的看到并行处理程序的特色“浅而宽”。这个算法如果结合CUDA平台的技术特性，会有多种变化，在后面会有详细分析。这里对算法本身做一点延伸介绍，有一种重要的并行设计策略——加速级联策略。

加速级联策略的核心思想是将一个最优但不是最快的和一个最快但不最优的级联起来。具体做法

（1）  开始使用最优算法，直到求解问题规模减少到某一阈值

（2）  接着使用快而非最优的算法，继续完成问题的求解

将Fan-in算法进行一下加速级联，算法如下：

假设，N个数a₁，a₂，…，a_n，其中N=P²

第一步：a₁+a₂，…，a_p+1+a_p+2，…,a_N-P+1+a_N-P+2

第二步：(a₁+a₂)+a₃,…，(a_p+1+a_p+2)+a_P+3,…,(a_N-P+1+a_N-P+2)+a_N-P+3

第P-1步：(a₁+a₂+…+a_P-1)+a_p,…,(a_P+1+a_P+2+…+a_2P-1)+a_2P,…,(a_N-P+1+a_N-P+2+…+a_N-1)+a_N

第P步到第P-1+log₂P步按Fan-in算法处理。

如果我们对以上两种算法做简要的理论分析，原始的Fan-in算法使用的处理线程个数为N/2，加速比为S_f和效率E_f（定义见1.3）分别为：

 

而使用加速级联之后的算法所使用的线程个数为 ，加速比S_z和效率E_z分别为

 

从上面的式子可以看到，传统Fan-in算法的加速比高于加速级联算法，但后者为常数效率的并行算法，也就是效率不随问题规模和处理器个数的改变而改变，也就是说当问题规模较小，线程切换延时也比较小的时候采用Fan-in算法，程序运行时间更快，当超过CUDA的执行函数（后面知道就是kernel）承担线程极限，函数切换延时较大时则应采用加速级联获得的执行效率更高。

这个地方简单的介绍了经典的Fan-in算法和加速级联策略，这些都是常见的并行算法设计技术。在本节中通过几个例子对并行算法设计的基本方法进行了说明，可以用一句口号来讲，一个中心，三项基本原则。

所谓一个中心，就是并行处理平台，在这里就是CUDA的各项技术特性，这是一切设计的基础。三项基本原则，挖掘计算无关性，旧瓶装新酒和独立自主创新，既有各自的技术含量，又离不开CUDA平台这个中心。下面1.3会继续并行算法设计的另一个重要方面——性能度量，接下来就会从此刻的多核时代（1.4），引出我们故事的主角CUDA（1.5）。

1.3 并行算法性能的度量

应该说在计算机科学里面，没有任何一个领域比高性能计算更加需要一个性能度量的标准了。因为它因计算需求而存在，它的价值也就体现在那一点一点的性能提升上面。对于并行程序更是这样，如果不是为了在便宜的处理器上得到高性价比的性能，何必要费弄个并行呢！所以想做CUDA的同志们一定要时刻记住，性能就是生产力！所以本节就稍稍分析下并行算法的性能度量。

对于一个给定问题，在设计完成一个并行处理算法之后，对并行处理算法性能的评价一般包括算法运行时间、算法并行度、算法成本、加速比和效率等方面。

（1）并行算法运行时间

这个标准是最直接又最简单的，首先来看看学术化定义，并行算法的运行时间是算法在并行计算机上求解一个问题所需的时间，多处理器就是最早开始到最晚结尾。在CUDA平台上一般有两种主要的时间测量方式，主机端测量和设备端测量（测量方法见5.1）。当然，还有采用CPU的计时器测量（很不精确，笔者已经将其彻底抛弃）。运行时间是我们衡量并行算法是否有用的最重要的指标，是第一生产力！

（2）算法的并行度

关于并行度，个人以为它更加是一种并行算法设计的评价。所谓算法的并行度，就是指该算法中可并行执行的操作数。若处理器资源无限，则算法的并行度可理解为可用来做并行运算的处理器个数。但在实际的算法中并行度不会是固定不变的，只能限制在一定的执行步骤或时间范围内。并行度刻画的是一个并行算法的并行程度，反映了软件并行性与硬件并行性的匹配程度。

这就引出了一个与并行度相关的主要概念是粒度，一般而言，大的粒度意味着能独立并行运行的是大任务，算法的并行度就小，反之算法的并行度就大。对于CUDA而言，由于线程是轻量级线程，所以在尽可能的情况下让线程的并行度达到GPU最大的活动量（后面会解释活动线程个数的概念），也就是达到处理器的最大吞吐量。当然有时候为了减少冗余计算和有效利用存储器带宽，也可以使用粗粒度的并行方式，毕竟我们实际的情况是资源有限。

最后给出一个形式化定义，算法的平均并行度定义为假设并行算法可以在m个并行步骤内完成，第i步时算法的并行度为DOP(i)，则算法的平均并行度为

 

这个定义可以用来评估你算法的一个整体资源使用情况，找到你并行算法的处理瓶颈之处。并行度小的地方，必然存在资源的空闲。我们后面会知道在Tesla架构下一个kernel函数占用整个GPU处理资源，如果并行度小，就会让处理资源空闲。相反并行度过大的时候，若不考虑CUDA资源的承载能力，就反而会让程序性能下降甚至执行失败，比如寄存器和共享存储器这样的稀缺资源。

（3）并行算法的成本

经济学上最经常出现的一个名词就是成本，当我们踏上并行处理这个靠效益吃饭的领域时，就不能再把自己当作象牙塔里阔佬了。我们要对节能，减排，经济，节约有深刻的认识。那么什么是并行算法的成本？下面的定义很直观，它就是并行算法的运行时间与并行算法所需的处理器个数的乘积，即

 

也就是说，成本等于最坏情况下求解某一问题时总的执行步数，其中包含了硬件和软件代价。对于我们CUDA程序而言（这里指单个GPU），每次执行占用的是处理器的全部资源，也就是说P是定值。我们的并行算法成本就是运行时间为主导，我们的目标就是在给定资源平台下让我们的时间降下来，成本下降就是效益啊！

（4）加速比，效率与加速比模型

笔者以为加速比是CUDA门外汉（比如说我们的客户）最关心的，我早年经常拿一些数据去唬人，GPU比CPU加速多少。但其实这个说法很不科学，也没有意义。由前面的并行算法成本就知道，运行时间才是评价和产生效益的关键，我的老板们都是很精明的，现在关心的都是你这个算法帧速率多少？还有没有可以优化的地方？那为什么我还要列出这个性能指标呢！因为它的理论模型对我们设计和评估并行算法还是很有借鉴意义的。

首先，加速比是衡量并行处理算法最传统的评价标准，体现了在并行计算机上运行并行算法求解实际问题所获得的效益。

加速比被定义为

 

其中T_l是最优串行算法在单处理器上的运行时间；T_p是并行算法在P个处理器上所需的时间。由这个定义就可以得到理论的最大加速比，这是我们自己给自己设定目标，以及在谈判桌上可以用到的。事实上，对于一个问题，如果按照某种条件，保持每个处理器的计算规模一定，并行算法的加速比S_p与处理器个数P成正比，则称该并行算法在该条件下，在该并行计算机上具有线性加速比。若在某些条件下，S_p＞P，则称该并行算法在该条件下，在该并行计算机上具有超线性加速比。

但上面对加速比的定义比较适合用于理论分析，一般理论分析也比较困难，所以在实际应用中往往使用下式：

 

一般的人只了解加速比这个概念，但其实与之密不可分的重要性能指标并行处理效率却容易被忽视。

首先，并行处理效率的定义是

 

其中P是处理器的个数。

我们来稍稍分析下，加速比其实是一个忽视了处理资源的评价，举个例子同一个CUDA程序在GT9600上（8个多处理器）加速40倍，在GTS 250上（16个多处理器）是不是就会80倍呢？由前面的分析可以知道，我们所期望的是处理效率高的线性加速比或超线性加速比，那么就不能脱离当前的计算资源来评价加速效果。

理解了加速比和效率这两个定义，后头去看看前面关于加速级联策略的分析，就会很清楚这些评价指标对设计的重要意义。最后再来说说两个经典的并行加速比模型，了解它们的目的在于认识到并行算法如何突破一些设计瓶颈的思路。

首先，并行加速比模型作为一个度量并行处理性能的参数，用来表示并行求解一个实际问题所获得的性能，即相对单处理器上的串行处理而言使用并行处理所获得的性能。从问题规模的角度出发可将其分为固定规模问题和可变规模问题的加速比模型。

①Amdahl模型

 

其中f是串行所占比例，N为并行所占比例

由此模型可以看出无论处理器数目如何，加速比都不能超过1/f，这就意味着在当前问题的计算需求下，无论你如何增加处理器的个数都无法继续增大加速比。例如在我们遇到一些无法展开的迭代运算时，无论你的GPU上有多少个多处理器都无能为力。但不要过于悲观，此模型忽视了问题计算规模的变化，强调的是通过并行处理来缩短求解问题的时间。再看看下一个模型之前，我们要明白，Amdahl模型告诉我们在问题规模一定的情况下，我们要尽可能的让计算负载分布到所有的计算资源上去，通过最大化并行提高加速比。

②Gustafson模型

与Amdahl模型不同，此模型说明了应该随处理器数目的增加而增加问题的规模，强调的是在同样的时间内，通过并行处理能运行多大的运算量，即通过运行时间来限制问题规模的增长程度。刚刚那个很悲观的例子，就可以采用增加运算规模来提高加速比，事实上在实时系统上这是有现实意义的。比如CUDA上PCI-Express是一个突出的性能制约因素，但通过大块数据传输可以提高处理效率；对于一些并行度很小的处理环节通过增加数据规模来提高加速性能，比如原先一个并行度小的kernel只处理一帧，现在合并处理多帧，这也可以提高处理效率和整体加速比。

虽然Amdahl模型和Gustafson模型只是是从不同角度去看待并行处理，但却告诉我们并行算法在设计和评价中的一些重要思想。比如要最大化利用计算资源，但不要盲目增加处理器；在计算规模可变的情况下，算法的串行瓶颈可以通过数据规模的增加而忽略等。

1.4 多核时代的各路诸侯

随着新摩尔定律的产生，我们已经进入到了一个多核时代。本人所在学校的外面不远处有个网吧，取名为四核时代。网吧都如此了，何况在我们这些专业的IT领域里面。本章的内容第一目的在于科普教育，让我们一起看看当今的多核技术发展，然后笔者希望通过一些比较让我们一起来看看多核GPU到底有些什么与众不同的地方。

首先，在多核处理器产生以前的并行处理主要以并行计算机和借助网络实现的大规模集群或分布式并行为主的两大派系。20世纪60年代初期，由于晶体管以及磁芯存储器的出现，处理单元的面积小型化，存储器也更小巧和廉价。这一时期出现了规模不大的共享存储多处理系统，就是我们所谓的大型机，典型代表IBM 360。与之相对应的是通常由数百数千甚至更多的处理器（机）组成的超级计算机，比如我国的天河-1A，速度全球第一，比第二名的美国国家实验室的计算机快30%，速度达到每秒2.5千万亿次运算。这两大类可以说是高性能计算领域的巨无霸。

另外一类可与之争锋的是借助网络实现的大规模集群或分布式并行计算。由于网络技术的高速发展，出现了以独立计算机连接组成的分布式并行处理系统——集群计算机。一个集群系统中的计算机节点可以是在一起的，也可以是物理上分离的，他们对于用户和应用程序来讲是透明的，如果只有一个单一的系统一样。这样就可以提供高性价比的服务，用来解决大型计算问题。随着技术的进一步发展，近年来产生了利用互联网上的计算机的 CPU的闲置处理能力来解决大型计算问题的分布式计算系统——网格计算。网格计算把一个需要非常巨大的计算能力才能解决的问题分成许多小的部分，然后把这些部分分配给许多计算机进行处理，最后把这些计算结果综合起来得到最终结果。相信大家都使用过迅雷，如果你用过离线下载，你会看到云端下载。这就是当今比较流行的云计算产物。云计算是网格计算、分布式计算、并行计算、网络存储、虚拟化等多种传统计算机技术和网络技术发展融合的产物。云计算在当今被称为是一种划时代的技术，因为它将数量庞大的廉价计算机放进资源池中，用软件容错来降低硬件成本，通过将云计算设施部署在各种能节省成本的区域，通过规模化的共享使用来提高资源利用率。国外代表性云计算平台提供商达到了惊人的10-40倍的性能价格比提升。不管有多么新潮的技术和名称，这一派系的高性能运算是以计算机网络资源为依托。

前面的两大门派，走的是高端路线，很难飞入寻常百姓家。多核处理器的出现打破了由大型机和网络集群式分布式系统的垄断，同时也使得并行算法和并行编程技术成为程序员不得不会的本领。

所谓多核处理器是指在一枚处理器中集成两个或多个完整的计算引擎。多核技术的产生源于工程师们认识到，仅仅提高单核芯片的速度会产生过多热量且无法带来相应的性能改善，同时单靠提高单核芯片速度的性价比也令人难以接受，速度稍快的处理器价格要高很多。作为处理器技术发展的先驱龙头企业，英特尔在1971推出的全球第一颗通用型微处理器4004，由2300个晶体管构成。这时，戈登摩尔提出了后来被业界奉为信条的“摩尔定律”——每过18个月，芯片上可以集成的晶体管数目将增加一倍。但到了2005年，当主频接近4GHz时，英特尔和AMD发现，速度也会遇到自己的极限：那就是单纯的主频提升，已经无法明显提升系统整体性能。以英特尔公司的奔腾系列为例，按照当时的预测，奔腾4在该架构下，最终可以把主频提高到10GHz。但由于流水线过长，使得单位频率效能低下，加上由于缓存的增加和漏电流控制不利造成功耗大幅度增加，3.6GHz奔腾4芯片在性能上反而还不如早些时推出的产品。所以，该系列只达到3.8G，就戛然而止。所以戈登摩尔本人似乎也依稀看到了“主频为王”这条路的尽头——2005年4月，他曾公开表示，引领半导体市场接近40年的“摩尔定律”，在未来10年至20年内可能失效。

多核心CPU解决方案的出现，标志着新摩尔时代的来临。事实上早在上世纪90年代末，就有众多业界人士呼吁用单芯片多处理器技术来替代复杂性较高的单线程CPU。IBM、惠普、Sun等高端服务器厂商，更是相继推出了多核服务器CPU。不过由于当时的技术并不成熟，造价也太高，并未引起大众广泛的注意。直到AMD率先推出64位处理器后，英特尔才想起利用“多核”这一武器进行“帝国反击战”。2005年4月，英特尔推出简单封装双核的奔腾D和奔腾四至尊版840。AMD在之后也发布了双核皓龙(Opteron)和速龙(Athlon) 64 X2和处理器。尽管如此，多核时代的元年应该是公元2006年。2006年7月23日，英特尔基于酷睿(Core)架构的处理器正式发布。2006年11月，又推出面向服务器、工作站和高端个人电脑的至强(Xeon)5300和酷睿双核和四核至尊版系列处理器。与上一代台式机处理器相比，酷睿2双核处理器在性能方面提高40%，功耗反而降低40%。多核CPU由于其在传统系统上重要的地位，无可争议的开启了当今的多核时代。

而与多核CPU差不多同时出现的就是本书的关键角色——多核GPU。GPU是相对于CPU的一个概念，传统的GPU是一个专门的图形的处理器，也是显卡的“心脏”。NVIDIA公司在1999年发布GeForce256图形处理芯片时首先提出GPU的概念。GPU使显卡减少了对CPU的依赖，并进行部分原本CPU的工作，尤其是在3D图形处理时。GPU所采用的核心技术有硬体T&L、立方环境材质贴图和顶点混合、纹理压缩和凹凸映射贴图、双重纹理四像素256位渲染引擎等，而硬体T&L技术可以说是GPU的标志。这也可以说是GPU脱离中央领导，自立门户的前奏。CUDA技术的前身GPGPU，即通用计算图形处理器。如笔者前面的介绍，CUDA事实上也是一种GPGPU技术，但为了与早期的GPGPU技术相区别和突出CUDA是真正意义上的的通用计算，笔者将两者视为不同。事实上，GPU的性能提升速度大大超过了CPU所遵照的摩尔定律，而且可编程性和功能都大大扩展，支持越来越复杂的运算。2006年，随着支持DirectX 10的GPU的发布，基于GPU的通用计算开始普及。Nvidia公司于2007年正式发布了CUDA这一通用计算平台。它是第一种不需要借助图形学API就可以使用以C/C++为基础的超集的通用计算的开发环境和软件体系。

从前面的介绍可以看到，在新的多核时代，多核处理器成为时代最强音。因为它们可以集成到超级计算机，集群系统与网络分布式计算系统，也可以广泛运用于个人PC，公共机等普通消费市场。这一时代的两大代表就是多核CPU和多核GPU，首先看看图1.5.

 

图1.5 CPU VS GPU

结构决定上层，有一种说法，CPU线程为重量级线程，GPU线程为轻量级线程。我们从结构对比来领会一下。首先，GPU上更多的资源放在了ALU（运算逻辑单元），这就意味着GPU拥有强大的计算能力。而在CPU上拥有更完整的控制单元和缓存空间，这就意味着CPU更适合处理复杂分支，不规则数据结构，不可预测存取，递归等情况。事实上，CUDA不支持递归程序。也正是因为CPU这样的结构，CPU线程的启动、切换、通信、同步操作开销巨大，和GPU有1000：1的关系。

总体来说，GPU 是专为计算密集型、高度并行化的计算而设计。因此，GPU 的设计能使更多晶体管用于数据处理，而非数据缓存和流控制。更具体地说，GPU 专用于解决可表示为数据并行计算的问题，即在许多数据元素上并行执行的程序，具有极高的计算密度（数学运算与存储器运算的比率）。由于所有数据元素都执行相同的程序，也就对精密流控制的要求不高；又由于有许多数据元素执行相同处理且具有较高的计算密度，因而可通过计算隐藏存储器访问延迟，而不必使用较大的数据缓存。这也是就GPU和CPU在结构设计上区别的根源。

值得注意的是，由于这两大处理器各有自己的优缺点，CPU+GPU的并行处理方式不仅可以满足更高的计算需求，节约成本和功耗，还可以将高性能计算广泛普及，引领新的程序设计方法革命。

1.5 高性能计算利剑之CUDA

首先，笔者很喜欢NVIDIA的CUDA海报——高性能计算的利器，见图1.6。2007年6月，NVIDIA推出了 CUDA。CUDA是一种将GPU作为数据并行处理设备的软硬件体系。它也是一种 GPGPU 的技术，但相对传统的GPGPU开发方法更为灵活方便，它是透过它的C语言的函数库和一些 CUDA延伸的语法来编写，因此不需用到 OpenGL 或 Direct3D，也不需要传统的图形函数库，也不会被传统的render pipeline设计束缚。

CUDA架构使GPU能够解决复杂的计算问题。它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。开发人员现在可以使用C语言来为CUDA™架构编写程序，C语言是应用最广泛的一种高级编程语言。所编写出的程序于是就可以在支持CUDA™的处理器上以超高性能运行。而且还支持其它语言，包括FORTRAN以及C++等。更为重要的是，NVIDIA CUDA技术是当今世界上第一种针对NVIDIA GPU 的C语言环境，该技术充分挖掘出NVIDIA GPU巨大的计算能力。凭借NVIDIA CUDA技术，开发人员能够利用NVIDIA GPU（图形处理器）攻克极其复杂的密集型计算难题。

 

图1.6 高性能计算之利器CUDA

作为并行计算平台的CUDA，首先是它通过采用统一处理架构，更加有效的利用了以往分布在顶点渲染器和像素渲染器的计算资源。这样既提高了GPU通用计算能力，又让开发人员无须用汇编或者高级着色器语言编写shader程序，然后再通过图形学API执行。

CUDA技术另一个重要的方面是片内共享存储器的引入，这样就使得GPU可以完成随机写入和线程间通信。在Fermi架构下，又加入了L1和L2缓存，使得GPU对数据读写的方式更为灵活，例如合并访问规则的放宽等。

从程序设计的角度来看，线程组的层次结构是一个很重要的核心概念（详见3.1.1）。它将指导程序员如何将问题分解为更小的片段，以便各个线程块通过协作的方法并行的解决。这样的分解通过允许线程在解决各子问题时协作的形式，也保持了程序语言的问题表达能力，同时保持了程序的伸缩性。各个线程块处理的子问题，安排在任何的CUDA core上进行处理。由此，编译后的 CUDA 程序可以在任何数量的处理器内核上执行（需要编译为相应的机器码，详见Nvcc），只有运行时系统需要了解物理多处理器数量。而这对程序员是透明的。

再来看看CUDA的强大计算能力，GPU的浮点运算能力是CPU的十数倍，如图1.7。有了这个数据对比，然后我们来看看CUDA到底给我们带来了什么。

 

图1.7 GPU与CPU浮点运算能力比较

第一个利用GPU的科学领域是天体物理学，在N体天体物理学研究中，天体物理学的研究人员利用GPU在一台普通的PC上模拟2百万个粒子，依靠GPU强大的并行计算能力，研究人员们做到了这一点。据性能分析报告知道，在N=16384时，GeForce 8800 GTX每秒可计算100多亿次的相互作用，每秒执行38个积分时间间隔。

在Manifold 8地理信息项目研究中，CUDA作为第一款标配应用软件用于地理信息系统（GIS）处理。该软件可制作出一幅地图，并叠加上人口信息，如该区域居民的年龄、住房类型、公路的数量等等 所有描述居住区的信息。规划人员使用GPU可以正确设计道路、房屋以及各种服务的位置，打造更加高效的城市。

在成像技术中，东京大学信息科学与技术研究生院机械信息系的Takeyoshi Dohi教授与他的同事研究了NVDIA的CUDA并行计算平台之后认为，医疗成像是CUDA这种平台非常有前途的应用领域之一。自2000年以来，这所大学的研究小组已经开发出一种系统，通过CT或MRI扫描实时获得的活体截面图被视为体纹理，这种系统不仅能够通过体绘制再现为三维图像，还可作为立体视频显示，供IV系统使用。该系统为实时、立体、活体成像带来了革命性的变化。但是，它的计算量极其庞大，仅体绘制本身就会带来极高的处理工作量，况且此后还需要进一步处理来实现立体成像。对于每一个图像帧，都有众多角度需同时显示。将此乘以视频中的帧数，您会看到令人震惊的庞大计算数量，且必须在很短的时间内高度精确地完成这样的计算。在2001年的研究中，使用了一台Pentium III 800 MHz PC来处理一些512 x 512解析度的图片，实时体绘制和立体再现要花费10秒钟以上的时间才能生成一帧。为了加速处理，研究小组起初尝试使用配备60块CPU的UltraSPARC III 900 MHz机器，这是当时性能最高的计算机。但可以得到的最佳结果也不过是每秒钟五帧。从实用的角度考虑，这样的速度还不够快。随后，研究人员使用NVIDIA GPU GeForce 8800 GTX开发了一个原型系统。在使用CUDA的GPU上运行2001年研究所用的数据集时，性能提升到每秒13至14帧。UltraSPARC系统的成本高达数千万日元，是GPU的上百倍，而GPU却交付了几乎等同于其三倍的性能，研究人员为此感到十分惊讶。同时，根据小组的研究，NVIDIA的GPU比最新的多核CPU至少要快70倍。测试显示，对于较大规模的体纹理数据，GPU的性能会更为突出。

我们通过这些不胜枚举的应用实例可以看到CUDA平台提供的强大计算能力给我们的生活带来的改变。笔者已经经历了由CUDA 1.0，1.1，2.3，3.1到现在的4.0的N多次改版，每一次改版都会引入一些新的技术特性，而且CUDA平台的功能也日趋完善，对新的硬件平台的特性也提供了良好的支持。笔者相信CUDA作为一把高性能计算的利剑，必将为解决科学研究和工程应用中的通用计算瓶颈建立卓越的功勋！多核时代，CUDA已经剑起了风尘，君等准备好了没有！

好了，到目前为止，我们已经充分的了解了为什么需要并行程序，知道了并行算法的设计目标，基本方法与性能度量指标，也清楚了我们处在一个怎样的多核时代。对我们要学习的CUDA技术也有了一个概括性的基本认识。做好准备，下一章开始，让我们完成第一个CUDA程序，You can do it！