1、CUDA并行编程的内容概要
在上一篇推文中，讨论了如何安装CUDA并使用它编写程序。尽管示例并不令人印象深刻，但它证明了使用CUDA是非常容易的。
在本次推文和下一次推文中，继续以这个概念为基础，讨论一下如何使用CUDA为GPU编写高级程序。从变量加法程序开始，然后逐步构建CUDA C中的复杂向量操作示例，同时也会介绍内核如何工作以及如何在CUDA程序中使用设备属性。本部分还会讨论在CUDA程序中向量是如何运算的，以及与CPU处理相比，CUDA如何能加速向量运算。除此之外，还会介绍与CUDA编程相关的术语。
本部分将讨论以下主题（紫色部分为本次推文的内容）：
1、内核调用的概念
2、在CUDA中创建内核函数并向其传递参数
3、配置CUDA程序的内核参数和内存分配
4、CUDA程序中的线程执行
5、在CUDA程序访问GPU设备属性
6、在CUDA程序中处理向量
7、并行通信模型
2、CUDA程序结构
上一篇推文中介绍了一个非常简单的“Hello,CUDA!”程序，其中展示了一些与CUDA程序相关的重要概念。CUDA程序是在主机或GPU设备上执行的函数的组合。不显示并行性的函数在CPU上执行，显示数据并行性的函数在GPU上执行。GPU编译器在编译期间隔离这些函数。如前一篇所示，在设备上执行的函数是使用__global__关键字定义的，由NVCC编译器编译，而普通的C主机代码是由C编译器编译的。CUDA代码基本上与ANSI C代码相同，只是添加了一些开发数据并行性所需的关键字。
因此，本次将用一个简单的双变量加法程序来解释与CUDA编程相关的重要概念，如内核调用、从主机到设备传递参数到内核函数、内核参数的配置、利用数据并行性需要的CUDA API，以及发生在主机和设备上的内存分配。
话不多说，直接上代码：
#include <iostream> #include <cuda.h> #include <cuda_runtime.h> #include <stdio.h> //定义两数相加的 kernel cuda 函数 __global__ void gpuAdd(int d_a, int d_b, int *d_c) { /*gpuAdd函数与ANSI C中的一个普通add函数非常相似。它以两个整数变量 d_a和d_b作为输入，并将加法存储在第三个整数指针d_c所指示的内存位置。 设备函数的返回值为void，因为它将结果存储在设备指针指向的内存位置中， 而不显式地返回任何值。*/ *d_c = d_a + d_b; } //主函数 int main(void) { //定义CPU计算时的指针 int h_c; //定义GPU指针 int *d_c; //使用cudaMalloc函数在GPU设备上分配d_c的内存；其作用类似于C语言中的Malloc函数 cudaMalloc((void**)&d_c, sizeof(int)); //Kernel call by passing 1 and 4 as inputs and storing answer in d_c //<< <1,1> >> = 一块设备和一个线程执行gpuAdd函数 //(6, 18, d_c)代表给gpuAdd函数传的参数 gpuAdd << <1, 1 >> > (6, 18, d_c); //将GPU运算的结果复制到主机设备上，也就是复制到CPU上进行接下来的操作 cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost); printf("6 + 18 = %d\n", h_c); //释放GPU设备上的内存，防止资源占用 cudaFree(d_c); return 0; }
2.1、内核调用
使用ANSI C关键字和CUDA扩展关键字编写的设备代码称为内核。它是主机代码（Host Code）通过内核调用的方式来启动的。简单地说，内核调用的含义是我们从主机代码启动设备代码。内核调用通常会生成大量的块（Block）和线程（Thread）来在GPU上并行地处理数据。内核代码非常类似于普通的C函数，只是这段代码是由多个线程并行执行的。
它以我们想要启动的内核的名称开始。你应该确保这个内核是使用__global__关键字定义的。然后，它具有<<<>>>内核启动配置，该配置包含内核的配置参数。它可以包含三个用逗号分隔的参数。第一个参数表示希望执行的块数，第二个参数表示每个块将具有的线程数。因此，内核启动所启动的线程总数就是这两个数字的乘积。第三个参数是可选的，它指定内核使用的共享内存的大小。在变量相加程序中，内核启动语法如下：
gpuAdd << <1,1> >> (6, 18, d_c)
在这里，gpuAdd是想要启动的内核的名称，<<<1,1>>>表示想用每个块一个线程启动一个块，这意味着只启动一个线程。圆括号中的三个参数是传递给内核的参数。这里，传递了两个常数，6和18。第三个参数是指向d_c设备显存的指针。它指向设备显存中的位置，内核将在那里存储相加后的结果。
程序员必须记住的一件事是，作为参数传递给内核的指针应该仅指向设备显存。如果它指向主机内存，会导致程序崩溃。内核执行完成后，设备指针指向的结果可以复制回主机内存，以供进一步使用。只启动一个线程在设备上执行不是设备资源的最佳使用。
2.2、配置内核参数
为了在设备上并行启动多个线程，必须在内核调用中配置参数，内核调用是在内核启动配置中编写的。它们指定了Grid中块的数量，和每个块中线程的数量。可以并行启动很多个块，而每个块内又有很多个线程。通常，每个块有512或1024个线程。每个块在流多处理器上运行，一个块中的线程可以通过共享内存（Shared Memory）彼此通信。程序员无法选定哪个流多处理器将执行特定的块，也无法选定块和线程以何种顺序执行。
假设要并行启动500个线程，你可以对前面解释的内核启动语法进行哪些修改？一种选择是通过以下语法启动一个包含500个线程的块：
gpuAdd << <1,500> >> (6, 18, d_c)
程序员必须注意，每个块的线程数量不能超过GPU设备所支持的最大限制。
如果需要要处理一个图像，你可以启动一个16×16的块网格，所有的块都包含16×16个线程。语法如下：
mykernel << < dim3(16, 16), dim3(16, 16) > >> ()
总之，在启动内核时，块数量和线程数量的配置非常重要。根据正在开发的应用程序和GPU资源的不同，应该谨慎地选择。
2.3、CUDA API函数
在变量加法程序中，会遇到了一些常规C或C++程序员不熟悉的函数或关键字。这些关键字和函数包括：
__global__
cudaMalloc
cudaMemcpy
cudaFree
__global__：它与__device__和__host__一起是三个限定符关键字。这个关键字表示一个函数被声明为一个设备函数，当从主机调用时将在设备上执行。应该记住，这个函数只能从主机调用。如果要在设备上执行函数并从设备函数调用函数，则必须使用__device__关键字。__host__关键字用于定义只能从其他主机函数调用的主机函数。这类似于普通的C函数。默认情况下，程序中的所有函数都是主机函数。__host__和__device__都可以同时用于定义任何类型函数。它生成同一个函数的两个副本。一个将在主机上执行，另一个将在设备上执行。
cudaMalloc：它类似于C中用于动态内存分配的Malloc函数。此函数用于在设备上分配特定大小的内存块。
cudaMemcpy：这个函数类似于C中的Memcpy函数，用于将一个内存区域复制到主机或设备上的其他区域。
cudaFree：类似于C中的free函数；
CUDA除了现有的ANSI C函数之外，还有许多其他的关键字和函数。我们会经常使用这三个函数，因此对它们进行了讨论。要了解更多细节，你可以看一下CUDA的编程指南。
3、在GPU设备上执行线程
#include <iostream> #include <stdio.h> __global__ void myfirstkernel(void) { //blockIdx.x gives the block number of current kernel printf("Hello!!! I'm thread in block: %d\n", blockIdx.x); } int main(void) { //kernel函数有16个并行块，每个块只有一个线程 myfirstkernel << <16, 1 >> > (); /*Function used for waiting for all kernels to finish 因为启动内核是一个异步操作，只要发布了内核启动命令， 不等内核执行完成，控制权就会立刻返回给调用内核的CPU线程。*/ cudaDeviceSynchronize(); printf("All threads are finished!\n"); return 0; }
从代码中可以看出，正在启动一个内核，它有16个并行块，每个块只有一个线程。在每个执行该段内核代码的线程里，我们打印出来它们各自获取到的块ID。可以认为，并行启动了16个执行相同myfirstkernel代码的线程副本。每个副本线程将拥有一个属于自己的块ID和线程ID。本例中，前者可以通过blockIdx.x的CUDA C的内置变量读取到。后者则可以通过threadIdx.x内置变量读取到。这两个ID将告诉我们正在执行内核的是具体哪个块和其中的哪个线程副本。当你多次运行程序时会发现，每次运行，线程块都是以不同的顺序执行的。一个样本输出如图：
这个程序还含有额外的一个CUDA函数调用：cudaDeviceSynchronize()。为何要加这句？这是因为启动内核是一个异步操作，只要发布了内核启动命令，不等内核执行完成，控制权就会立刻返回给调用内核的CPU线程。在上述的代码中，CPU线程返回，继续执行的下一句是printf()。而再之后，在内核完成之前，进程就会结束，终止控制台窗口。所以，如果不加上这句同步函数，你就看不到任何的内核执行结果输出。在程序退出后内核生成的输出结果，将没有地方可去，你没法看到它们，因此，如果我们不包含这个指令，你将不会看到任何内核执行的printf语句的输出结果。要能看到内核生成的输出结果，我们必须包含这句同步函数。这样，内核的结果将通过可用的标准输出显示，而应用程序则会在内核执行完成之后才退出。
