CUDA, 用于大量数据的超级运算：第二节

本文来自：http://www.ddj.com/hpc-high-performance-computing/207402986

第一个内核
Rob Farber 是西北太平洋国家实验室（Pacific Northwest National Laboratory）的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究，并且是几个新创企业的合伙人。大家可以发邮件到[e-mail]与他沟通和交流。

在该系列文章的第一节，我展示了第一个简单的CUDA（Compute Unified Device Architecture，计算统一设备架构之首字母缩写）程序——moveArrays.cu，让您熟悉下用于构建和执行程序的CUDA工具。对于C程序员而言，这个程序只是调用了CUDA API将数据移入和移出CUDA设备。并没有添加新内容，以免在学习如何使用工具构建和运行CUDA程序时发生混淆。
这一节是在第一个示例的基础上添加了几行代码，以便在CUDA设备上进行简单的计算——特别是在浮点数组中以1为增量增加每个元素。令人惊喜的是，该样例已经提供了使用CUDA解决很多问题的基本框架（“将数据移动到CUDA启动的设备、进行计算并获取结果”）！
在涉及更多高级的话题之前，您需要首先了解：
什么是内核？内核是一个函数，可以从主机调用，并可以在CUDA设备上执行－由多个线程平行执行；
主机如何调用内核？这涉及确定内核名称和执行配置。就该专栏文章来说，执行配置仅仅意味着定义在运行CUDA设备内核时，组中的平行线程数和组的数量。这实际上是个很重要的话题，我们将在以后的专栏里详细介绍。
如何同步内核和主机代码。
在列表1的最上面（incrementArrays.cu），我们可以看到主机例程的例子，incrementArrayOnHost 和我们的第一个内核，incrementArraysOnDevice。 
主机函数incrementArrayOnHost只是对数组元素数的简单循环，以1为增量增加每个数组元素。此函数用于在此代码末尾进行比较，以验证内核在CUDA设备上进行了正确的计算。
在列表1稍下面的位置是我们的第一个CUDA内核，incrementArrayOnDevice。CUDA提供了几个对C语言的扩展。该函数类型限定符__global__将函数声明为CUDA设备上的可执行内核，只能从主机调用。所有内核必须声明返回类型为void。
内核incrementArrayOnDevice与incrementArrayOnHost进行相同的计算。仔细查看incrementArrayOnDevice会发现里面没有循环！这是因为该函数是由CUDA设备上的一组线程同时执行的。但是，每个线程都具有一个唯一的ID，可以用于计算不同的数组索引或制定控制决策（比如，如果数组索引超过数组大小则不进行任何操作）。这使得incrementArrayOnDevice的计算变得非常简单，如同计算寄存器变量idx中的唯一ID一样，然后使用该变量唯一地引用数组中的每个元素并以1为增量递增。因为线程数可能超过数组大小，因此先将idx与N相比较（N是向内核中传递的一个参数，用于指定数组中的元素数），看一下是否需要进行一些操作。
那么内核是如何调用，执行配置又是如何指定的呢？控制按顺序流过源代码，从main开始到包含列表1中Part 2 of 2语句的注释下面。

 

 

view plaincopy to clipboardprint?

1. // incrementArray.cu  
2. #include <stdio.h>  
3. #include <assert.h>  
4. #include <cuda.h>  
5. void incrementArrayOnHost(float *a, int N)  
6. {  
7.   int i;  
8.   for (i=0; i < N; i++) a[i] = a[i]+1.f;  
9. }  
10. __global__ void incrementArrayOnDevice(float *a, int N)  
11. {  
12.   int idx = blockIdx.x*blockDim.x + threadIdx.x;  
13.   if (idx<N) a[idx] = a[idx]+1.f;  
14. }  
15. int main(void)  
16. {  
17.   float *a_h, *b_h;           // pointers to host memory  
18.   float *a_d;                 // pointer to device memory  
19.   int i, N = 10;  
20.   size_t size = N*sizeof(float);  
21.   // allocate arrays on host  
22.   a_h = (float *)malloc(size);  
23.   b_h = (float *)malloc(size);  
24.   // allocate array on device   
25.   cudaMalloc((void **) &a_d, size);  
26.   // initialization of host data  
27.   for (i=0; i<N; i++) a_h[i] = (float)i;  
28.   // copy data from host to device  
29.   cudaMemcpy(a_d, a_h, sizeof(float)*N, cudaMemcpyHostToDevice);  
30.   // do calculation on host  
31.   incrementArrayOnHost(a_h, N);  
32.   // do calculation on device:  
33.   // Part 1 of 2. Compute execution configuration  
34.   int blockSize = 4;  
35.   int nBlocks = N/blockSize + (N%blockSize == 0?0:1);  
36.   // Part 2 of 2. Call incrementArrayOnDevice kernel   
37.   incrementArrayOnDevice <<< nBlocks, blockSize >>> (a_d, N);  
38.   // Retrieve result from device and store in b_h  
39.   cudaMemcpy(b_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);  
40.   // check results  
41.   for (i=0; i<N; i++) assert(a_h[i] == b_h[i]);  
42.   // cleanup  
43.   free(a_h); free(b_h); cudaFree(a_d);  

 

表1： incrementArrays.cu.

这将排队启动支持CUDA的设备上的incrementArrayOnDevice，并说明添加到C语言的另一个CUDA，即对CUDA内核的异步调用。该调用指定了内核的名称和封闭在三角括号"<<<" and ">>>"之间的执行配置。注意指定执行配置的两个参数：nBlocks和blockSize，将在下面对它们进行讨论。任何对内核调用的参数都通过标准C语言参数列表提供，该列表包含以标准C语言样式"(" and ")"分界的函数。在本示例中，指向设备全局内存的指针a_d（它包含数组元素）和N（数组元素数）都被传递到内核。
因为CUDA设备是空闲的，内核立即开始根据执行配置和函数参数运行。同时，内核启动后，主机继续执行代码的下一行。此时，CUDA设备和主机同时运行他们各自的程序。在incrementArrays.cu中，主机立即调用cudaMemcpy，它会等待设备上的所有线程完成（例如，从incrementArrayOnDevice返回），之后它将修改后的数组拉回主机。该程序在主机系统进行完串行比较后完成，以验证我们在平行CUDA设备上通过incrementArrayOnDevice得到的结果与在主机上通过串行版incrementArrayOnHost得到的结果相同。
在内核启动时，通过执行配置（在本样例中，通过包含在三角括号 "<<<" and ">>>"之间的变量nBlocks和blockSize），确定几个变量。nBlocks和blockSize之后的思路非常精妙，程序开发人员可以应对硬件限制，而无需重新编译应用程序－这是使用CUDA开发商业软件的本质特征。
在以后的文章里，我将检查块中的线程是否有能力互相通信和同步。这是个非常棒的软件特征，但是从硬件的角度来说，花费比较大。相对于比较便宜（也比较旧式）的设备，这需要更为昂贵的（在未来）设备来支持每块更多的线程数。创建网格抽象，这样程序开发人员可以考虑（无需重新编译）区分硬件的能力，而不用考虑价位和年代。事实上，网格将具有同样维度和大小的块对同一内核的调用汇成一批处理，然后有效地乘以一个因子nBlocks（该因子是可以在单个内核调用中启动的线程数）。性能稍查的设备可能只能同时运行一个或几个线程块，而性能较强大的（如较贵，未来开发出来的）设备可能同时运行很多个线程块。使用网格抽象设计软件要求在同时运行的诸多独立线程之间进行平衡，并且要求块内有更多的线程，能够彼此合作。请注意两种线程的成本。当然，不同的算法会有不同的要求，但是如有可能，尽量使用更多的线程块。
在CUDA启动的设备上的内核中，有几个可用的内置变量，它们是通过内核调用的执行配置设置的。它们是：
blockIdx包含网格内的块索引。
threadIdx包含块内的线程索引。
blockDim包含块内的线程数。
这些变量是包含整数变量组件的结构。例如，块有x-、y-和z-整数组件，因为它们是三维的。而网格只有 x-和y-组件，因为它们是二维的。本样例只使用了这些变量的x-组件，因为我们移动到CUDA设备的数组是一维的（下面的专栏文章将介绍二维和三维配置能力的功效，以及如何利用这种功效）。
我们的示例内核使用这些内置的变量，通过下面的语句来计算线程索引idx：
int idx = blockIdx.x * blockDim.x + threadIdx.x;

变量nBlocks和blockSize分别是网格中块的数量和每个块中的线程数。在本样例中，它们就在主机代码的内核调用前初始化： 
int blockSize = 4;
int nBlocks = N/blockSize + (N%blockSize == 0?0:1);

当N不能被blockSize整除时，nBlocks计算中的最后一项会加上一个额外的块，这意味着在有些情况下，块中的某些线程将不会进行任何有用的工作。
显然，本样例被故意简化了，因为它假定数组大小小于能够被包含在4个线程块中的线程数。这显然是过于简单，但它让我们能够通过简单的代码了解对incrementArrayOnDevice的内核调用。
还有很重要的一点需要强调，每个线程都能够访问设备上的整个数组a_d。在内核启动时没有固有的数据分区。这由程序员在编写内核时根据要识别和利用的计算的数据平行情况来决定。
表1说明了如何计算idx，及如何引用数组a_d。（如果前面的文本有任何不清楚的地方，我建议向incrementArrayOnDevice添加一个printf语句，以便将idx和用于计算它的相关变量打印出来。为仿真器编译程序，"make emu=1"，运行它看看会发生什么。一定要指定到仿真器可执行程序的正确路径来查看printf输出。）

 

同样，内核调用是不同步的——在内核启动后，控制立即返回到主机CPU。之前所有CUDA调用结束后，内核将在CUDA设备上运行。不同步的内核调用是重叠主机和设备上的计算的极好方式。在本例中，对incrementArrayOnHost的调用可以放在对incrementArrayOnDevice的调用之后，以重叠主机和设备上的计算来获得更好的性能。主机和设备可以同时计算，这取决于内核完成计算需要的时间量。
在继续阅读下一篇专栏文章之前，我建议：
尝试改变N和nBlocks的值。看一下当它们超过设备能力时会发生什么。
想想如何引入循环来处理任意大小的数组。
区分不同类型的支持CUDA的设备的内存（比如，全局内存、寄存器、共享内存和持久内存）。看一下CUDA占用率计算器，以及nvcc选项 - cubin或 --ptxas-options=-v，来决定内核中使用的寄存器的数量。