[C++知识库]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任务完成后才能执行，所以它已经自带主机同步，不再需要手动阻塞