CUDA并行编程概述
CUDA
CUDA是英伟达推出的GPU架构平台，通过GPU强大的并行执行效率，为计算密集型应用加速，CUDA文件以.cu结尾，支持C++语言编写，在使用CUDA前需要下载 CUDA Toolkit
内存与显存
CPU可以访问内存，GPU可以访问显存，如果需要使用GPU进行计算，必须把数据从内存复制到显存
指向显存的指针
创建一个指向显存的指针，下面的代码可以告诉你为什么要使用 (void**)类型
int* p; // 这是一个指向int变量的内存指针
function(p); // 如果直接把指针传入函数，那么它会以参数的形式被带入计算，函数中的操作无法修改p的值
function(&p); // 只要把p的地址传入函数，函数就可以通过地址修改指针的值
void* v; // 但是指针类型很多，为了统一，通常会使用无类型指针
function(&v); // 所以我们应该传入一个指向无类型指针地址的指针
(void*)p; // 这样可以把 p 变成无类型指针，但是我们需要的是指向 p 的地址的无类型指针
(void**)&p; // 这样我们就得到了指向 p 的地址的无类型指针
function((void**)&p); // 这样function函数就可以直接修改p的数值
void* p;
function(&p); // 如果你的 p 已经是无类型指针，那么可以直接使用取址符在GPU中申请显存，并获得指向显存的指针
int length = size * sizeof(int);
int a[size];
int b[size];
int c[size];
int* dev_a;
int* dev_b;
int* dev_c;
cudaMalloc((void**)&dev_a, length);
cudaMalloc((void**)&dev_b, length);
cudaMalloc((void**)&dev_c, length);此时的dev_a, dev_b, dev_c已经指向显存地址，空间大小为 length
内存与显存的数据交换
在使用GPU计算前，需要把数据复制到显存
cudaMemcpy(dev_a, a, length, cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, length, cudaMemcpyHostToDevice);dev_a是显存指针，a是内存指针cudaMemcpyHostToDevice表示把长度为length的内存数据复制到显存里
计算完成后，需要把数据从显存复制到内存以供CPU计算
cudaMemcpy(c, dev_c, length, cudaMemcpyDeviceToHost);这段代码的含义是把dev_c指向的显存地址的数据复制到c指向的内存地址
在计算结束后，应该释放显存空间
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);栅格结构
GPU的结构包含栅格(grid)，块(block)，线程(thread)，许多线程组成一个“块”，许多个“块”组成一个“栅格”，其中grid和block都可以用三维向量表示，假设一个block有1024个线程，如果创建4个block，则总共有4096个线程同时运行
下面的代码展示了如何获取block和thread的编号
int i = threadIdx.x;
int j = blockIdx.x;合理使用这些编号，可以帮助你定位变量位置
__global__ void DoInKernel(int* a, int* b, int* c) {
int i = blockIdx.x * 1024 + threadIdx.x;
c[i] = a[i] + b[i];
}函数限定词
核函数
核函数使用 __global__ 修饰，它在CPU上调用，在GPU上执行
__global__ void DoInKernel(int* a, int* b, int* c) {
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
int main(){
DoInKernel << <1, size >> > (dev_a, dev_b, dev_c);
}其中 <<< >>>运算符决定了执行核函数的方式，第一个参数是block的数量，即一个grid里有几个block，它实际上是一个dim3类型的变量，在处理多维数组时它可以让你的代码编写更加方便，但是这里不做演示
dim3 dg(10, 10, 10);
DoInKernel << <dg, size >> > (dev_a, dev_b, dev_c);第二个参数是thread的数量，即一个block里有几个线程，它同样是dim3类型的变量，如果输入的是int，则默认y和z都是1
后面还有两个可选参数，分别用来表示共享内存大小和流，共享内存大小限制了可以动态分配的共享内存的最大值，流指定使用哪个IO通道在内存和显存之间复制数据，使用不同的流可以防止阻塞
内联函数
内联函数使用 __device__ 修饰，它必须在GPU上调用，只能在GPU上执行
__device__ int add(int a, int b) {
return a + b;
}
__global__ void DoInKernel(int* a, int* b, int* c) {
int i = threadIdx.x;
c[i] = add(a[i], b[i]);
}主机函数
所有不加修饰的函数都是主机函数，它也可以使用 __host__ 修饰，主机函数只能在CPU上调用和执行，例如 main 就是一个主机函数
__host__ int main(){
return 0;
}异常处理
CUDA代码极难调试，因此最好在每一步都检查一次错误，一旦发生错误，立即转到错误处理
int main()
{
//无关代码
if (cudaMalloc((void**)&dev_a, length) != cudaSuccess) {
goto OnError;
}
if (cudaMalloc((void**)&dev_b, length) != cudaSuccess) {
goto OnError;
}
if (cudaMalloc((void**)&dev_c, length) != cudaSuccess) {
goto OnError;
}
//无关代码
return 0;
OnError:
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
}主机同步
由于GPU的代码是异步执行的，如果两个核函数被写在一起，那么它们很可能会被同时执行，使用cudaDeviceSynchronize()阻塞主机线程，可以确保所有的核函数或者IO流都已经执行完毕，才会继续执行下面的代码
DoInKernel_1 << <1, 10 >> > ();
cudaDeviceSynchronize();
DoInKernel_2 << <1, 10 >> > ();其中cudaMemcpy()函数必须在所有GPU任务完成后才能执行，所以它已经自带主机同步，不再需要手动阻塞