1、CUDA程序结构
CUDA程序是在主机或者设备上执行的函数的组合。不显示并行性的函数在CPU上执行，显示数据并行性的函数在GPU上执行，GPU在编译期间要隔离这些函数。CUDA代码基本上与C代码相同，只是添加了一些开发数据并行性所需的关键字。
２、CUDA C中的双变量加法程序
编写一个将两个变量相加的内核（kernel）函数
#include<iostream> #include<cuda.h> #include<cuda_runtime.h> __global__ void gpuAdd(int d_a, int d_b, int* d_c) { *d_c = d_a + d_b; }
main函数代码
int main(void) { //定义主机变量，统一以h_开头 int h_c; //定义设备变量，统一以d_开头，作为参数传递给内核的指针应该仅指向设备显存 int* d_c; //在设备上分配d_c变量的内存，cudaMalloc类似于C中的malloc函数
//分配了一个整数变量大小的内存块，并返回指向该内存位置的指针d_c
cudaMalloc((void**)&d_c, sizeof(int)); //调用gpuAdd，其中1和4是两个输入变量，d_c是一个作为输出指针变量的设备显存指针
// gpuAdd << <1, 1 >> > (1, 4, d_c); //将结果从设备内存复制到主机内存中 cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost); //释放设备上通过cudaMalloc使用的所有内存 cudaFree(d_c); return 0; }
需要注意的是，作为参数传递给内核的指针应该仅指向设备显存，如果它指向主机内存，会导致程序崩溃。其次，当通过对调用的内核传递参数时i，所有的指针仅指向设备显存。
３、内核调用
使用ASNSI C关键字和CUDA扩展关键字编写的设备代码称为内核。它是主机代码（Host Code）通过内核调用的方式来启动的。简单地说，内核调用的含义是我们从主机代码启动设备代码。内核调用通常会产生大量的块（Block）和线程（Thread）来在GPU上并行地处理数据。内核使用__global__关键字定义，使用kernel<< <希望执行的块数， 每个块将具有的线程数， 内核使用的共享内存大小>> >来启动内核。在变量相加程序中，内核启动的语法如下：
gpuAdd << <1, 1 >> > (1, 4, d_c)
4、CUDA中的向量运算
上面的双变量加法没有利用到GPU设备的并行处理能力，下面利用GPU的并行结构修改CPU上的向量加法程序，首先在CPU上编写向量加法程序
4.1、CPU向量加法
#include "stdio.h" #include<iostream> //定义向量元素个数 #define N 5 //定义CPU向量加法 void cpuAdd(int *h_a, int *h_b, int *h_c) { int tid = 0; while (tid < N) { h_c[tid] = h_a[tid] + h_b[tid]; tid += 1; } } int main(void) { int h_a[N], h_b[N], h_c[N]; //向量初始化 for (int i = 0; i < N; i++) { h_a[i] = 2 * i*i; h_b[i] = i; } //调用CPU向量加法函数 cpuAdd (h_a, h_b, h_c); //输出结果 printf("Vector addition on CPU\n"); for (int i = 0; i < N; i++) { printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]); } return 0; }
4.2、GPU向量加法
在GPU上添加修改后的内核函数：
#include "stdio.h" #include<iostream> #include <cuda.h> #include <cuda_runtime.h> //定义向量元素个数 #define N 5 //定义向量加法内核函数 __global__ void gpuAdd(int *d_a, int *d_b, int *d_c) { //获取当前内核的块ID int tid = blockIdx.x; //在当前块ID来初始化tid
//因为是同时启动的，每个块必须直到自己要处理的数据是什么
if (tid < N) d_c[tid] = d_a[tid] + d_b[tid]; } int main(void) { //定义主机变量 int h_a[N], h_b[N], h_c[N]; //定义设备指针，指向cudaMalloc分配的显存 int *d_a, *d_b, *d_c; //给设备变量分配内存 cudaMalloc((void**)&d_a, N * sizeof(int)); cudaMalloc((void**)&d_b, N * sizeof(int)); cudaMalloc((void**)&d_c, N * sizeof(int)); //主机向量初始化 for (int i = 0; i < N; i++) { h_a[i] = 2*i*i; h_b[i] = i ; } //将主机向量复制到设备向量中 cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice); //使用内核启动符号<<<>>>调用内核函数，启动N个块（和向量长度一样，使用不同的块并行处理向量N个位置的加法），每个块1个线程，传递参数为设备指针 gpuAdd << <N, 1 >> >(d_a, d_b, d_c); //将结果从设备内存复制到主机内存中 cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost); printf("Vector addition on GPU \n"); //在控制台中输出结果 for (int i = 0; i < N; i++) { printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]); } //释放内存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; }
在main函数中调用内核函数时，并行启动了N个块，这意味着我们同时启动了N个执行该内核代码的线程副本。因为N个块是同时启动的，只有每个块分工合作效率才会高，main函数调用gpuAdd内核函数时，启动的块是随机的，但可以通过blockIdx.x查看启动的块的ID，用该ID作为块处理的数据对应的ID，这样就可以让N个块分别处理向量中N个数据的加法。GPU的并行计算的方式比CPU的串行计算提高了吞吐量。
4.3、GPU对向量每个元素进行平方
上面是并行启动N个块，每个块只有一个线程。也可以在1个块里面并行启动N个线程。
#include "stdio.h" #include<iostream> #include <cuda.h> #include <cuda_runtime.h> //定义向量元素个数 #define N 5 //平方内核函数 __global__ void gpuSquare(float *d_in, float *d_out) { //获取当前调用内核函数启动的线程ID int tid = threadIdx.x; //使用内核函数计算线程ID对应的向量中的元素 float temp = d_in[tid]; d_out[tid] = temp*temp; } int main(void) { //定义主机向量 float h_in[N], h_out[N]; //定义设备指针 float *d_in, *d_out; //在CPU上分配内存 cudaMalloc((void**)&d_in, N * sizeof(float)); cudaMalloc((void**)&d_out, N * sizeof(float)); //Initializing Array for (int i = 0; i < N; i++) { h_in[i] = i; } //将主机上的数组复制到设备上 cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice); //调用1个块，N个线程的内核函数 gpuSquare << <1, N >> >(d_in, d_out); //将设备上的结果复制到主机内存上 cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost); //控制台输出结果 printf("Square of Number on GPU \n"); for (int i = 0; i < N; i++) { printf("The square of %f is %f\n", h_in[i], h_out[i]); } //释放内存 cudaFree(d_in); cudaFree(d_out); return 0; }
当使用这种方式启动N个线程时，每个块的最大线程数不超过512或1024，如果N=2000，而GPU的线程最大数量为512，那么不能写成<<<1, 2000>>>，而应该用<<<4, 500>>>。应当选择合适的块和线程的数量。
5、并行通信模式
当多个线程并行执行时，他们遵循一定的通信模式，直到它们在显存哪里输入，哪里输出，通讯模式包括映射、收集、分散式、蒙版、转置。
5.1、映射（Map）
单一输入，单一输出。Map的代码模式看起来如下：
d_out[i] = d_in[i] * 2;
5.2、收集（Gather）
多个输入，单个输出。Gather模式的代码看起来如下：
out[i] = (in[i - 1] + in[i] + in[i + 1]) / 3;
5.3、分散式（Scatter）
单一输入，多个输出。Scatter模式的代码看起来如下：
out[i - 1] += 2 * in[i] and out[i + 1] += 3 * in[i];
5.4、蒙版（Stencil）
从数组中读取固定形状的相邻元素时，称为stencil模式。比如3X3或者5X5卷积就是Stencil，它是Gather操作的一种模式，代码和Gather相似。
5.5、转置（Transpose）
对矩阵进行转换。