CUDA并行编程概述

CUDA

CUDA是英伟达推出的GPU架构平台,通过GPU强大的并行执行效率,为计算密集型应用加速,CUDA文件以.cu结尾,支持C++语言编写,在使用CUDA前需要下载 CUDA Toolkit

内存与显存

CPU可以访问内存,GPU可以访问显存,如果需要使用GPU进行计算,必须把数据从内存复制到显存

指向显存的指针

创建一个指向显存的指针,下面的代码可以告诉你为什么要使用 (void**)类型

使用(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计算前,需要把数据复制到显存

复制数据

1 2	cudaMemcpy(dev_a, a, length, cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, length, cudaMemcpyHostToDevice);

dev_a是显存指针,a是内存指针cudaMemcpyHostToDevice表示把长度为length的内存数据复制到显存里

计算完成后,需要把数据从显存复制到内存以供CPU计算

复制数据

1	cudaMemcpy(c, dev_c, length, cudaMemcpyDeviceToHost);

这段代码的含义是把dev_c指向的显存地址的数据复制到c指向的内存地址

在计算结束后,应该释放显存空间

释放显存

1
2
3

cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);

栅格结构

GPU的结构包含栅格(grid),块(block),线程(thread),许多线程组成一个“块”,许多个“块”组成一个“栅格”,其中grid和block都可以用三维向量表示,假设一个block有1024个线程,如果创建4个block,则总共有4096个线程同时运行

下面的代码展示了如何获取block和thread的编号

块和线程

1 2	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

1 2	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 就是一个主机函数

主机函数

1
2
3

__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流都已经执行完毕,才会继续执行下面的代码

主机同步

1
2
3

DoInKernel_1 << <1, 10 >> > ();
cudaDeviceSynchronize();
DoInKernel_2 << <1, 10 >> > ();

其中cudaMemcpy()函数必须在所有GPU任务完成后才能执行,所以它已经自带主机同步,不再需要手动阻塞