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

第二章 You can do it！

这一章将从如何搭建CUDA编程环境说起，让你一步一步完成你的第一个CUDA程序。本章会不厌其烦的把每一个环节罗列出来，虽然我不能保证未来的CUDA平台会否有所改变，但我确信可以让每一个初学者在目前的平台下完成自己的第一个CUDA程序，并且对CUDA程序的基本框架有一定认识。

对于本章的读者，有这几个必要的假设：

1，    掌握了基本的C语言编程知识，至少会写一个Hello world。

2，    完全不懂CUDA或者入门级，如果已经会了，请跳过这一章。笔者不是想忽悠无知群众，而是不想浪费您的时间。

3，    已经有了VS2005以上版本的开发平台，如果没有，请预先安装，工欲善其事，必先利其器！本书以古老的VS2005为例（不是我守旧，这是当年参加微软创新杯，送给俺们团队的，是正版的东西，也请各位不要买盗版的书，保持一下作者的写作热情哈）。

本章的内容安排2.1让你在现有IDE平台上搭建一个CUDA编程环境，然后介绍传说中的线程给你认识下，对于他们，你就是主，你对他们不认真，他们就会让你不消停！一位师兄曾说：CUDA设计是一门艺术。原因就在于线程结构设计的多样化。有了线程就如有了战士，下面就是给他们作战的空间。作为兵团的领袖，你将拥有传统程序员不曾拥有的第四空间——显存（设备存储空间），备注：笔者以为的程序员第一空间是内存空间，第二是磁盘空间，第三是网络空间（本质也是磁盘，但毕竟是远程嘛）。除了资源就是执行计划了，这就需要了解CUDA程序的核心——kernel，它可以给你带来惊喜，也可以给你带来噩梦，所以规划执行的时候一定要严谨啊！有了资源，也掌握了执行计划编写的方法，就让你创造的孩子们问候你吧！

本章最后通过一个比较的方式，看看为什么说CUDA的出现改变了GPGPU的应用前景，看看CUDA到底多了什么！

2.1 You can program on CUDA

要编写一个CUDA程序，首先你需要有一个集成开发环境（IDE），当然，对于猛人来说，也是可以没有的。我当年Unix的老师就极度讨厌windows，他说windows让软件开发的门槛下降了，因为window 95出来前一年，这位猛人狂学了命令行下的编程。本书的对象是大多数如我一般的普通人，所以我们需要一个交互性相对好的开发平台。本书用VS2005作为IDE。

关于CUDA程序执行环境的配置，首先不得不提一下硬件环境，笔者并没有将NVDIA的显卡作为必要前提，因为在理论上学习CUDA可以用Emudebug模式（相关配置见nvcc说明部分），但笔者认为既然打算学习CUDA编程，投放一点点资金在设备上还是值得的。显卡的型号和参数可以在官网上找到，网址是http://developer.nvidia.com/cuda-gpus，型号上加M表示是移动设备上的。笔者的笔记本上用的就是GeForce 9600M GT，计算能力1.1，而台式机使用的是GTX 460，计算能力2.1。显卡都不是很贵的，从学习技术的角度，使用计算能力2.0以上的设备可以学习和享用到Fermi架构带来的技术进步，但笔者认为计算能力2.0以下的设备对一些优化技巧的应用要苛刻一些，对个人能力的锻炼是有好处的，所谓由俭入奢易，由奢入俭难，在低端显卡上很苛刻的锻炼自己，必将会使自己在高端显卡使用上更得心应手！何况在某些实际应用中受功耗的限制，高端显卡会无用武之地，不过学习新技术特性又是很有必要的，不如大家学学我，配置两种，呵呵！

有了GPU（如果暂时没有，只做模拟程序，请跳过关于驱动的说明部分），下一步是安装驱动，一般来讲，只要你的电脑在正常运行，显卡驱动就已经有了，但我们是专业人士，和普通用户不同，我们应当时刻清楚一点，保持自己的显卡驱动为最新最稳定版本，理由只有一个——最新即最快。这一点，笔者可以负责任的这么认为，因为我们测试过，同样的CUDA程序，尤其是CUDA提供的一些库，更新驱动后，性能会有明显提升！

驱动可以从官方网站http://www.nvidia.cn/Download/index.aspx?lang=cn下载。笔者使用的CUDA驱动程序采用版本为270.51_winxp_32bit，注意有32/64位的区别，图2.1，2.2，2.3是安装过程的截图。笔者建议使用默认的安装路径，这样少一些配置路径的麻烦。

前面的东西其实不是CUDA特色的，因为只有你的电脑是支持CUDA的显卡，都是要装驱动的。下面要介绍的就是CUDA环境的主角SDK（SoftwareDevelopment Kit）和toolkit，在http://developer.nvidia.com/cuda-downloads 可以下载。这里笔者采用的是Toolkit4.0和SDK4.0。同样安装时建议使用CUDA的默认路径，这样减少配置头文件和库文件的一些麻烦。安装过程比安装驱动还简单，一路确定就OK了。下面就来讲讲在VS2005里的相关配置，如果你采用的是默认路径，就可以直接按我的说明来做，否则就注意修改下路径。

图2.1 安装选项，一般来讲，我选自定义，把所有东西都装上，反正现在的磁盘那么大，呵呵！

1，添加包含文件

在VS2005的工具->选项->项目与解决方案->VC++目录->包含文件中添加CUDA目录中的include和NVIDIA CUDASDK\common中的inc，如下：

C:\Documents andSettings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK4.0\C\common\inc

C:\ProgramFiles\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include

2，添加库文件

在VS2005的工具->选项->项目与解决方案->VC++目录->库文件中添加CUDA中的lib和NVIDIA CUDA SDK\common中的lib，如下：

C:\Documents andSettings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK4.0\C\common\lib\Win32

C:\ProgramFiles\NVIDIA GPU Computing Toolkit\CUDA\v4.0\lib\Win32

3，添加可执行文件

在VS2005的工具->选项->项目与解决方案->VC++目录->可执行文件中添加以下两个bin目录：

C:\Program Files\NVIDIA GPUComputing Toolkit\CUDA\v4.0\bin

C:\Documents andSettings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK4.0\C\common\bin

除此之外，还要重编译生成C:\Documents andSettings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK4.0\C\common的cutil_vs2005工程，debug和release各编译一次，然后检查C:\Documents andSettings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK4.0\C\common\lib\Win32中cutil32.lib和cutil32D.lib两个文件是否生成成功。环境配置完成。

友情提示：

1，要使SDK中的范例在VS2005下正常工作，就必须重编译生成OpenCL\common里的工程。

2，在VS2005的工具->选项->文本编辑器->文件扩展名中选择VC++，添加.cu扩展名，可使CUDA程序如普通C++程序一样，高亮显示。

图2.2 安装ing

图2.3 安装完毕，一般我是最后都装完了，再重启，呵呵！

好了，完成了以上步骤的准备，我们就可以开始正式进入到CUDA程序编写的阶段了！Are you ready？Let’s go！

2.2 You can createsmany threads

由第一章，我们已经清楚明白的知道，CUDA是一个并行处理的平台，如果将这个环境比喻为战场，那么你是集团军总司令，要解决的问题就是你面对的敌人，线程就是最可爱的士兵。

在操作系统课程里面两个重要的概念，进程与线程。个人认为两者的定义在实践上说，区别并不是那么的重要，都可以看成是一次执行的实体，只是单位和量级不一样！这里不详细分析，单说线程。CUDA程序设计中首先需要考虑的就是线程的设置，你可以把它当作你想派出去战斗的部队规模，每位士兵干的工作都基本相同（同样的控制指令，不同的数据而已）。那么如何来设置呢！我们还是用“臭名昭著”的例子hello world作为例子，我们先不管这个函数怎么实现的，假定别人已经为我们写好了函数哈！下面的这个语句相信大家都知道，一个标准的C风格函数调用：

Hello ();

同样的函数在CUDA程序里面就变成了下面这个样子：

int threadNum =128;//线程数

Hello<<<1, threadNum >>>();

多了什么呢？一个奇怪的<<<>>>，它就是我们首先要介绍的，CUDA函数执行配置（注：如果你使用驱动API，可以通过标准的C语言方式进行执行配置，而不需要有这个新的奇怪的配置语法）。完整的执行配置最多有四项，这里把它先列出来，在后面的章节逐个解释哈。

<<<gridDim，blockDim，sizeofSM，streamId>>>

gridDim：线程网格的维度设置信息，可以是一维或二维，暂时不支持三维, 详解见3.1.1

blockDim：线程块的维度设置信息，可以是一维、二维或者三维，详解见3.1.1

sizeofSM：动态使用共享存储器（share memory，SM）的字节数（byte为单位），默认为0，详解见3.3

streamId：设置流的索引，也就是设置工作流，默认为0，表示只有一个工作流。

如果没有CUDA基础，就暂时不用理解上面四个配置项的含义，也会不用去管怎么设置。这里只设置了1个一维的线程快，里面包含了128个线程，就OK了。

简而言之，要想拥有你的战斗部队，只要在普通函数前设置这个<<<>>>就可以了，而我们的第一个例子就用这个设置了包含128个线程的小部队吧！

进入下一节之前，再来看看这个东东<<<>>>，有点奇怪的符合，但很重要，有了它，C函数的名称改头换面为kernel；有了它，你就可以建立你的高性能战斗部队！呵呵！

2.3 You can havethe device memory

由第一章，我们已经知道CUDA平台为我们开辟了一个新的战场——显存空间，我把它称为第四空间。关于存储空间的详解将放在下一章，这一节主要想让读者初步了解如何分配和释放最基本的显存空间，也就是全局存储器。

在正式引入第一个CUDA库函数之前，想稍稍说明一下，CUDA提供的API（应用程序接口）主要分成两大类驱动API和运行时API。本书后面会对驱动API有一些介绍，但本书的主要目标是面对工程实践，驱动API过于繁琐，一般的应用不需要使用，所以本书大部分时候都只提供运行时API版本的程序实例，而不给出对应的驱动API版本。

这里首先给出一个CUDA库函数：

cudaError_t cudaMalloc (void ** devPtr, size_t size)

功能说明：分配一个线性的显存空间，与C语言的malloc功能完全一样，除了执行的对象由内存变成了显存，准确地讲是全局存储器（global memory）。

参数说明：

devPtr：指向分配空间的指针

size： 要分配的空间大小，以字节为单位

错误产生的可能原因：分配未释放空间

使用范例：

float *d_fData;

cudaError_t enCudaRet =cudaMalloc((void**)&d_fData, sizeof(float)*256);

if(cudaSuccess != enCudaRet)

{     

       returnfalse;

}

这是我们第一次使用CUDA的库函数，这里的异常检测不是最常用的写法，一般使用CUDA_SAFE_CALL()或者其它类似的方式，这里写出来主要是为了让大家看看函数的返回情况，后面就不会这样写了。

事实上，因为CUDA提供了多种不同的显存使用方式，分配显存的库函数很多，我们会在第三章详解全局存储器的时候予以更多的说明。

对于指针的使用，一个最好的编写程序的习惯就是写好了分配，马上就写释放，所以第二个出场的库函数就是cudaFree。函数原型如下：

cudaError_t cudaFree (void *devPtr)

功能说明：与C语言的free一样，释放已分配的线性空间。它既可以释放由cudaMalloc分配的线性空间，还可以释放由cudaMallocPitch和cudaMalloc3D分配的线性空间。

错误产生的可能原因：释放尚未分配的空间或重复释放，两者本质是一个意思。我们只要记住最简单的原则，一分，一放，交替进行就好了。

这一节推出了两个最基本的CUDA运行时API，目的不只在于让你知道如何使用它们，更重要的是是让你了解我们必须要了如指掌的一个资源空间，因为往后它会是我们的重要家底，下一节我们将来编写被称为kernel的CUDA函数，也就是要为线程们安排工作任务了。

2.4 You can writea kernel

前面在讲如何创建线程的时候已经提到了kernel函数，这里继续用最简单的Hello来介绍如何编写一个CUDA执行函数。先给出一个最简单的示例，不绕弯子，先给出函数的定义：

__global__ void Hello()

{

       constunsigned int tid = threadIdx.x + 1;

       constunsigned int bid = blockIdx.x + 1;

       cuPrintf("NO.%din group %d says hello to you, sir!\n", bid,tid);

      

}

让我们一行一行的来分析，这个函数的第一行比起一般C/C++函数，就多了一个__global__关键字，首先特别说明一下，CUDA上的关键字基本都是用__×××__这样的形式，习惯了就好了，呵呵！

所有的kernel函数都必须使用__global__这个函数类型限定符，这就是kernel函数的特别标志。它只能在主机端调用，也就是说kernel不能调用kernel（这个就意味着不能递归哈），使用时要配置2.2中讲的配置参数。

__global__定义的函数的返回类型必须为void，我想这个问题可以很好的直观理解，kernel是在GPU上运行的多线程函数，如果返回值如int，主机端如何接收多个值，何况GPU处理的所有数据均在显存上，CPU根本无法触及数据本身。所以任何的kernel计算结果都要先放在显存上，然后通过device to host或者zero-copy方式（后面会有详细解释）返回给主存。这是kernel的一个特色。

然后__global__定义的函数是异步返回的，也就是说它不会等设备完成所有工作就返回，但对于每个工作流，kernel的启动是自动同步的。笔者认为在这里展开讨论异步与同步是时候未到，所以这里大家有个印象，kernel函数的返回是异步的。这里用表2.1总结下我们的第一个关键字。

跳过花括号，就来到了函数的主体部分，前两句基本是一样的，而且是以后写kernel最常见的语句，也许初看有点奇怪，结构体threadIdx和blockIdx没有任何的声明和定义，怎么不报编译错误呢？！

Table 2.1关于__global__

关键字	作用	用法说明
__global__	定义函数为kernel函数	1，返回类型必须为void 2，只能在主机端调用 3，调用时要进行执行配置 4，异步返回

这是因为它们是CUDA上的内建变量。内建变量的详细解释会在下一章讲解线程结构时给出。这里你可以把它们看成是已有的变量。

threadIdx.x表示了当前线程在块内x维度上的索引号，从0开始

blockIdx.x表示了当前线程在grid内x维度上的索引号，从0开始

有了内建变量，我们就可以为每个线程建立一个标识，这是我们使用它们的基石。这里加1，只是想从1开始编号而已，没有什么特殊的用意哈！

cuPrintf的作用和参数用法与C语言的printf基本一致，这里用它作为我们的第一个例子，当然使用它要遵循一些限制，所以要让这个程序跑起来还需要一些额外的工作，这个会在下一节全部给出。

总结一下，kernel函数是在GPU上执行的多线程函数，声明和定义必须使用__global__关键字予以限定，不能返回任何东西（所以为void）。

CUDA已经为我们提供了内建变量，我们可以很轻易的通过建立索引来指挥我们的线程工作。

2.5 You can be greeted by so many threads

由2.2我们已经知道如何设置多线程（也就是执行配置），从2.4我们又知道了如何定义kernel，所以现在让我们的线程战士们来问候阁下你吧！

第一步，建一个工程，做法很多。最经济实用的就是使用由赵开勇编写的CUDA_VS_Wizard，网上搜一下就有了，而且很方便！装好后在VS2005里面会多出下面这个部分：

图2.4 使用CUDA_VS_Wizard建工程

然后取好名字确定，选择一些基本的东西就OK。大部分的时候，我都这样干，谢谢勇哥这样的先驱。

但是考虑到大家有可能会在已有的工程中添加CUDA，所以我还是把最一般的建工程方法给出，以Win32 console工程为例，MFC的配置与之类似，但要稍稍多一点点工作，在本书的第四章nvcc时会有所说明。

言归正传，首先用VS建一个空的Win32 console工程，添加一个cpp源文件，这样就如下图：

图2.5 空的空的Win32 console工程

然后将后缀由cpp改为cu，然后右击工程，进入自定义生成规则，选择你所装的CUDA rules（如果是CUDA1.1的开发包，还需要你自己找到cuda.rules文件添加），建议你与时俱进，如果你和我一样，是CUDA4.0就会看到图2.6。选择run time的就好了，退出时的提示，直接确定。

图2.6 CUDA生成规则选择

到此为止，工程建立好了，在现有工程中添加CUDA文件的做法和相关设定与这个类似。当然不同工程会遇到一些些特殊的问题，笔者将尽力在本书中把自己遇到过的配置编译问题的解决方案给出，但毕竟水平有限，很难面面俱到。

建好了工程，然后就是编写代码，第一个部分是相关头文件和库的引用。关于库的引用，你可以在VS工程属性->链接器->输入里面的添加附加依赖项添加基本的lib，cudart.lib和cutil32d.lib是最基本的。当然笔者建议你用预编译指令的方式，如下：

#pragma comment(lib,"cudart.lib")

#if defined(DEBUG) || defined(_DEBUG)

#pragma comment(lib,"cutil32d.lib")

#else

#pragma comment(lib,"cutil32.lib")

#endif

#include <cuda_runtime.h>

#include <cutil.h>

#include "cuPrintf.cu.h"

注意最后一个文件引用，这里要使用cuPrintf，它不是CUDA的库函数，你需要自己添加，首先在SDK的工程simplePrintf里（默认路径为C:\Documents and Settings\All Users\ApplicationData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\src\simplePrintf）找到cuPrintf.cu和cuPrintf.cuh文件，将其拷贝到自己的工程目录下。然后将cuPrintf.cu文件更名为cuPrintf.cu.h，这是为了去除重复定义的链接错误。当然你可以有其它做法，这里是笔者给出的一种解决方案。

接下来，就是kernel函数，如下：

/************************************************************************/

/* Hello, Sir                                                              */

/***********************************************************************/

__global__ void Hello()

{

constunsigned int tid = threadIdx.x + 1;

constunsigned int bid = blockIdx.x + 1;

cuPrintf("NO.%din group %d says hello to you, sir!\n", bid,tid);

 

}

Main函数部分如下：

int main(int argc, char* argv[])

{     

 

       cudaPrintfInit();

 

       Hello<<<1,32, 0>>>();

 

       cudaPrintfDisplay(stdout,true);

       cudaPrintfEnd();

 

       return0;

}

里面除了Hello这个kernel外的三个函数是使用cuPrintf必须的，分别是初始化，显示和结束，这个不用去深究它。只要知道cuPrintf使用的两个最限制：

1，必须是1.1以上的GPU

2，必须使用CUDA 2.3以上的开发包

你通过改变配置参数的值，就可以看到不同的线程组织方式，图2.7给出了两种不同的配置输出结果。

<<<1, 32, 0>>>

<<<4, 8, 0>>>

图2.7 Threads say hello to you (I)

到现在我们已经完成了第一次CUDA程序Hello，但笔者认为由于cuPrint的使用，我们的第一个CUDA程序少了一些环节，想一想2.3节，是的，我们还没有自己动手使用显存空间呢！所以，笔者这里给出一个不使用cuPrintf的Hello示例，来让大家认识下一个完成CUDA程序的基本框架流程——Hello_v2.0哈！

首先，我们来看看新的kernel是怎么样的：

__global__ static void Hello_v2(char*result, int N)

{

       constunsigned int tid = threadIdx.x;

       constunsigned int bid = blockIdx.x;

       constunsigned int index = bid * blockDim.x + tid;

       inti;

       charhello[] = "says hello to you, sir!";

       result[index* N] = (char)(index + 1);

       result[index* N + 1] = ' ';

       for(i= 2; i < N; ++i)

       {

              result[index* N + i] = hello[i - 2];                      

       }

}

程序是比1.0复杂了一点点，但也没有什么太多很神秘的。参数表多了两个参数，一个是用于指向存放结果buffer的指针，另外一个参数是我们要得到的字串长度。

通过const unsigned int index = bid * blockDim.x+ tid，我们可以得到每个线程属于自己的一个索引号。

这里的result[index * N] = (char)(index + 1)，是我想线程在报到时从1开始数起。

后面的部分就是最普通的字符串赋值，相信大家应该可以明白这个kernel在做什么吧！

对于结果buffer的写入result[index * N + i]，大家也不要觉得太奇怪！这里分配的空间是一个一维空间，每个线程都需要写入N个char，所以通过index * N就可以让每个线程知道自己写入的起始地址。参考图2.8，不难理解吧！

图2.8 Hello_v2.0线程访存示意图

下面来看看新的main函数：

int main(int argc, char* argv[])

{     

       char*d_result;

       char*h_result;

       intblocks = 1;

       intthreads = 32;

       intN = 25;

       intsize = N * threads * blocks;

       //分配主机端线程

       h_result= new char[size];

       //分配显存

       cudaMalloc((void**)&d_result, sizeof(char) * size);

 

       Hello_v2<<<blocks,threads, 0>>>(d_result, N);

       //将结果写回主机端

       cudaMemcpy(h_result,d_result, sizeof(char) * size, cudaMemcpyDeviceToHost);

       //打印到控制台显示

       printf("%d",h_result[0]);

       for(inti = 1; i < size; ++i)

       {

              if((i% N) == 0)

              {

                     printf("\n");

                     printf("%d",h_result[i]);

                     continue;

              }

              printf("%c",h_result[i]);

       }

       printf("\n");

       //显存空间释放

       cudaFree(d_result);

       //主机端内存释放

       delete[]h_result;

 

       return0;

}

2.0是不是突然比1.0多了好多好多，呵呵，事实上1.0也包含了这么多步骤，只不过别人帮你写好了而已。好，现在我们来一步一步看看这个CUDA程序。

笔者在给buffer命名时喜欢用d和h的前缀来区分设备端和主机端的空间，主机端与一般的C/C++相同，自不必多讲。而设备端的空间分配在2.3已经说明，显存的分配采用与C语言一样的形式，因为每个线程需要N个char空间，size就表示了所有线程所需要的结果空间大小。

接下来是Kernel的调用执行，如一般的C函数一样传参，然后加上2.4里已经说明的执行配置。这里需要说明下，对于kernel，可以直接参数数值，它会被作为线程的私有变量放入寄存器（也可能是局部存储器）中，而不需要预先放入显存。但GPU只能对显存进行寻址（即使零拷贝方式，也要经过一个地址映射的步骤），所以传入kernel的指针必须是显存空间。

Kernel完成之后，结果仍然在设备端，而我们想在主机端的控制台进行显示。所以就要把结果拷贝回来（当然，如果你是做图像处理后的显示，可以采用D3D或者OpenGL与CUDA的互操作直接显示，而不需要这样的拷贝）。所以这里我们要介绍另外一个非常常用的库函数cudaMemcpy。

cudaError_t cudaMemcpy (void _ dst, const void _src, size_t count, enum cudaMemcpyKind kind)

功能说明：将源地址src的count字节大小的数据拷贝到目标地址dst，必须指定是何种拷贝类型kind

参数说明：

dst：拷贝的目的地址

srt：用于此次拷贝的数据源地址

count：要拷贝的数据字节数

kind：拷贝类型，共有4种拷贝类型：

cudaMemcpyHostToHost：主机端到主机端，与C语言的memcpy差不多

cudaMemcpyHostToDevice：主机端到设备端，一般用于数据输入

cudaMemcpyDeviceToHost：设备端到主机端，一般用于结果返回

cudaMemcpyDeviceToDevice：设备端到设备端，在显存上搬移。

由于这四种数据搬移对程序性能有很大的影响，后面章节会有更全面深入的分析。这里希望读者对CUDA程序中的这点特色有个初步的认识，数据会有主机端和设备端两种物理上不同的存放空间。

错误产生的可能原因：

1，拷贝的源和目的地址有重叠，这个一般只会发生在主机端和设备端的内部搬移过程中，零拷贝的地址映射除外，所以准确的说是拷贝的源和目的物理地址有重叠。

2，输入的源和目的地址与所选择的拷贝类型不匹配，例如选择主机端到设备端cudaMemcpyHostToDevice，dst却为主机端内存地址，或者srt为显存地址。

事实上cudaMemcpy是一个函数族的普通一员，还有cudaMemcpy2D，cudaMemcpy3D，cudaMemcpyToArray等等各种用于不同场合的数据拷贝函数，这些函数虽然参数表有一些不同，但使用的核心思想是完全一致，只是用在CUDA提供的各种特殊存储器和存储方式上而已。后面章节里面也会有更多的一些介绍。

经过这一步的拷贝后，数据就回到了主机端。这里用控制台进行输出显示。最后释放掉这里用到的主机端和设备端空间。程序运算结果见图2.9。

现在是不是觉得这个2.0 Hello程序看起来比1.0也没有多什么东西，的确如此。不过我们稍稍总结一下，就可以从这个简单的例子中看到CUDA程序的一般性框架。

如图2.10所示，CUDA程序框架与一般的C程序相比，主要多了显存分配，数据传输，GPU执行和显存释放这几个环节。作为GPU执行的程序，CUDA平台为程序员开辟了第四空间，并且由此引出了多种具有不同性能特性的存储器，如常量存储器，纹理存储器，表面存储器等。显存的管理是CUDA程序员首先需要了解并掌握的，任何一个CUDA程序都需要先进行显存的分配，这是GPU处理所必需的环节。

对于多数应用程序，输入数据一般存放在主机端，这就需要通过主机端到设备端的拷贝为GPU提供输入，当然也有不需要这样输入的，如本章的例子。

有了空间和数据后，GPU就可以执行相关多线程任务，如kernel函数或者已经开发好的CUDA函数库（cufft，cublas等见4.3）

GPU的结果输出，与输入类似，可通过拷贝传回主机端，本章例子就是如此，也可通过其它方式输出，如D3D等见4.2。

与一般C程序的内存空间管理一样，显存空间同样需要合理的进行释放。

图2.9 Hello_v2.0运行结果

图2.10 CUDA程序的一般性框架

2.6 CUDA比C多的这么一点点

经过本章啰哩啰嗦，应该还算细致的分析介绍，希望并相信每一个读者都已经知道了CUDA环境的配置，CUDA程序的基本组成，并完成了自己的第一个CUDA程序了吧！

CUDA程序是C语言程序的一个超集，目前应该还不是C++的超集哈！笔者以对比和总结的方式来谈谈CUDA程序本身多了些什么，以此作为本章的结束。

1，函数的声明和调用形式

CUDA程序的主要作用是实现GPU上的并行处理，因此，它引入了一个名为kernel的多线程函数。Kernel函数相比较于传统的C函数在函数声明和定义时需要使用__global__关键字来告诉编译器，这是一个会在设备上生成多线程执行的函数。线程的设置机制则是通过引入<<<>>>这样的执行配置完成的。同时，这个多线程函数也不能像一般的C函数一样返回运行结果，返回类型必须为void，也不能递归调用kernel函数。

上面所说的这些kernel函数与传统C函数的差异，其实都非常直观！它需要有一个标识来告诉编译器（后面第四章会讲CUDA的编译过程），它是一个在调用时需要配置执行参数的多线程设备函数。

执行配置的语法样子有些奇怪，但它本身却是很自然的，目前的四个配置项都有自己很重要的用途。最基本的前两项恰恰是多线程函数所必须的线程结构配置项。

GPU多线程函数不能返回值，这个做法也是完全合情合理的。因为这里的多线程是设备端的多线程，主机端的线程没有办法直接访问显存，也就不可能得到返回结果。

因为kernel的调用是需要执行配置的，而执行配置里面的项只能在主机端完成。也就是说GPU线程不能在创建新的GPU线程，这个是和目前的GPU工作机制一致的（SPMD单线程多数据），所以Kernel不能递归调用。

除了这些之外，kernel的语法与一般的C函数并无差别。

2，内置变量

要编写一个真正的CUDA程序，必须使用CUDA提供的几个基本的内置变量。这些变量的主要的目的就是标识线程，你作为程序的设计者，如战地指挥官一样，需要给你的每一个战士编个号，这样才能形成战斗序列，安排任务。

CUDA程序最基本最核心的地方就是用这些内置变量让不同的线程去处理不同的数据，如此而已！在语法与C相比，没有任何新东西，不过就是使用几个内置的结构体嘛！

3，显存空间

传统的C程序在CPU上运行，我们需要管理的是内存空间。而在GPU上运行的CUDA程序，多的就是一个显存空间（当然它的内容很丰富，见下一章），在程序设计的思想上是完全一致的。

对于CUDA程序员来讲，显存空间的引入，就需要增加对显存空间使用和管理的意识，同时也要洞悉数据在程序中存在的物理位置，看起来多了一些东西，但理解和使用起来，其实和一般的C程序没有太大的区别！

笔者前面特别说CUDA目前还不是C++的超集，这是因为nvcc主机端只支持部分的C++ ISO/IEC 14882:2003，而设备端不支持运行时类型（RTTI），异常检测和C++标准库。使用时也要遵循一些限制（见CUDA编程指南附录D2）

 

好了，本章非常肤浅的讨论了如何完成一个CUDA程序的编写，很多地方只是把一些最基本的概念解释了下，目的在于让读者快速完成自己的第一个CUDA程序！本书的第三章会展开讨论CUDA程序设计的最核心部分，应该是每个CUDA程序员所不得不知道和掌握的部分，第四章则是讨论一些CUDA技术相对周边的东西，这些并非不重要，而是不同的应用需求以及不同CUDA程序员的不同的要求，并不一定都需要知道的。