CUDA——用于大量数据的超级计算：第二部分（Rob Farber专栏）

本文在第一个示例的基础上添加了几行代码，以便在CUDA设备上进行简单的计算——特别是在浮点数组中以1为增量增加每个元素。令人惊喜的是，该示例已经提供了使用CUDA解决很多问题的基本框架（“将数据移动到支持CUDA的设备、进行计算并获取结果”）！

在本系列文章的第一部分，我展示了第一个简单的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语句的注释下面。 // incrementArray.cu #include <stdio.h> #include <assert.h> #include <cuda.h> void incrementArrayOnHost(float *a, int N) { int i; for (i=0; i < N; i++) a[i] = a[i]+1.f; } __global__ void incrementArrayOnDevice(float *a, int N) { int idx = blockIdx.x*blockDim.x + threadIdx.x; if (idx<N) a[idx] = a[idx]+1.f; } int main(void) { float *a_h, *b_h; // pointers to host memory float *a_d; // pointer to device memory int i, N = 10; size_t size = N*sizeof(float); // allocate arrays on host a_h = (float *)malloc(size); b_h = (float *)malloc(size); // allocate array on device cudaMalloc((void **) &a_d, size); // initialization of host data for (i=0; i<N; i++) a_h[i] = (float)i; // copy data from host to device cudaMemcpy(a_d, a_h, sizeof(float)*N, cudaMemcpyHostToDevice); // do calculation on host incrementArrayOnHost(a_h, N); // do calculation on device: // Part 1 of 2. Compute execution configuration int blockSize = 4; int nBlocks = N/blockSize + (N%blockSize == 0?0:1); // Part 2 of 2. Call incrementArrayOnDevice kernel incrementArrayOnDevice <<< nBlocks, blockSize >>> (a_d, N); // Retrieve result from device and store in b_h cudaMemcpy(b_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); // check results for (i=0; i<N; i++) assert(a_h[i] == b_h[i]); // cleanup 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输出。）图1 同样，内核调用是不同步的——在内核启动后，控制立即返回到主机CPU。之前所有CUDA调用结束后，内核将在CUDA设备上运行。不同步的内核调用是重叠主机和设备上的计算的极好方式。在本例中，对incrementArrayOnHost的调用可以放在对incrementArrayOnDevice的调用之后，以重叠主机和设备上的计算来获得更好的性能。主机和设备可以同时计算，这取决于内核完成计算需要的时间量。在继续阅读下一篇专栏文章之前，我建议： • 尝试改变N和nBlocks的值。看一下当它们超过设备能力时会发生什么。 • 想想如何引入循环来处理任意大小的数组。 • 区分不同类型的支持CUDA的设备的内存（比如，全局内存、寄存器、共享内存和持久内存）。看一下CUDA占用率计算器，以及nvcc选项 - cubin或 --ptxas-options=-v，来决定内核中使用的寄存器的数量。更多信息有关本系列文章的后续部分，请参见CUDA, Supercomputing for the Masses: Part 3。有关本系列文章的前一部分，请参见CUDA, Supercomputing for the Masses: Part 1。单击这里CUDA，获得有关CUDA的更多信息，单击这里NVIDIA，获得有关NVIDIA的更多信息。