CUDA编程初体验

前言

真的比较难蚌，谁能想到有朝一日我终于走上了CUDA编程的道路，要不是这个该死的论文非要用什么“弹性可分离卷积”，还真的就不能使用torch直接写出来，必须使用自定义的算子，必须要接触到CUDA编程。

无心插柳柳成荫吧，毕竟，搞深度学习的人不接触一点高性能计算，不懂得一点C++和CUDA编程，说出去给人的感觉也太不专业了，学一点，总归是好的。

CUDA编程基础

了解CUDA，可参考如下博客，必须要对GPU的硬件架构方式有所了解，对于英伟达的GPU，可使用以下示意图说明：
CPU与GPU的硬件架构示意图
不同于CPU的冯诺依曼架构，GPU中的控制器与缓存器几乎同时要对应极大数量的低精度算术逻辑单元，这一批算术逻辑单元共享同一个缓存单元，全局共享内存位于DRAM上，DRAM通过PCIE接口与CPU的外接PCIE接口，GPU由于其本身极差的单核处理性能，其不能单独作为一个计算机系统中央处理单元使用，因此只能作为CPU的一个外部设备device，因此CPU与GPU的关系是协同处理关系。早期的图形处理单元作为外设的普通一员，使用与其他外设共用的南桥接口，受CPU节制；随着图像处理单元的工艺升级和深度学习的应用，现代GPU使用的PCIE接口与CPU内部直联，客观上提高了CPU与GPU的协同处理能力。

超大规模的线程调用是使用GPU的关键步骤，因此一个重要概念是GPU线程。GPU中的线程呈现出空间三维的关系，其示意图如下所示：
线程示意图
两个重要的线程组织单元是线程块Block和线程网格Grid，线程块是一个线程的集合，线程网格是一个线程块的集合。线程块中的线程可以共享同一块共享内存，线程网格中的线程块可以共享同一块全局内存。线程块和线程网格的大小是可以在程序运行时动态调整的，但是线程块的大小是有限制的，线程块的大小不能超过1024个线程，在使用GPU运行计算程序时，需要指定线程时，应当注意使用所有线程唯一的网序号gridIdx和块序号threadIdx，这两个变量将发挥类似于CPU中循环变量i的作用，需要特别牢记。

gridIdx和threadIdx都是结构体变量，同时可以取x、y和z三个成员变量，共计可取三个维度的序号。

核函数程序

CUDA编程语法是C++的扩展，整体上与C++完全兼容（也就是老黄所谓的护城河），因此使用起来极为方便。CUDA程序主要运行在GPU上，如果该程序的指令来自于主机CPU，则该程序又被称为核函数程序，是需要编码的主要程序。核函数程序最为主要的一个标志是：使用修饰符__global__装饰函数，返回值必定为void。例如：

__global__ void kernel(float *d_out, float *d_in){
    int idx = threadIdx.x;  
    d_out[idx] = d_in[idx] * d_in[idx];
}

该核函数完成的是一个向量的点乘功能，在主机中如果调用该程序，则与调用普通的函数程序相同，但是需要在函数名前增加模板修饰符<<<grid_size, block_size>>>，该模板第一个变量是网格数，第二个变量是块数，这两个变量均可以使用CUDA内建的dim3变量（即三维变量结构体）作为输入，也可以使用整型变量等。

线程块大小最好是线束大小的整数倍，在大于32时，通常最好取为32的整数倍。

一个完整的CUDA程序在编译完毕后，由主机端自动完成配置，CUDA架构内部完成对该调用指令的分发与下达。一个可能的调用上述函数的例子是：

int main(){
    float *d_in, *d_out;
    float *h_in, *h_out;
    int size = 1024 * sizeof(float);
    h_in = (float *)malloc(size);
    h_out = (float *)malloc(size);
    for(int i = 0; i < 1024; i++){
        h_in[i] = i;
    }
    cudaMalloc((void **)&d_in, size);
    cudaMalloc((void **)&d_out, size);
    cudaMemcpy(d_in, h_in, size, cudaMemcpyHostToDevice);
    kernel<<<1, 1024>>>(d_out, d_in);
    cudaMemcpy(h_out, d_out, size, cudaMemcpyDeviceToHost);
    for(int i = 0; i < 1024; i++){
        printf("%f", h_out[i]);
    }
    free(h_in);
    free(h_out);
    cudaFree(d_in);
    cudaFree(d_out);
    return 0;
}

本例中基本上已经陈列出了全部常用的CUDA函数，包括设备指定、内存空间分配与释放、内存的复制等，除了更了复杂（为了取得进一步的高性能运算）的目的以外，加上C语言的编程知识已经能够应付绝大多数情况了。

核函数中不支持C++的iostream头文件，核函数不可成为一个成员函数

核函数中最好不使用CUDA函数，也最好不使用自定义的数据结构

核函数是直接在GPU内部运行的，因此其功能应当尽量减少封装，而CUDA函数往往是主机函数，不能够被核函数调用（除非是设备函数），而且相互调用也容易造成错误的内存访问产生报错，不推荐使用。因此最好让核函数的数据类型和函数回调尽量简单且底层，可以在核函数的封装函数中进一步完善其调用接口。

通常采用的方式是将核函数经包装函数warpper封装后，将封装函数定义为类的成员，在Pytorch的CUDA自定义算子接口中，就使用到了这样的方法。

内建修饰符

除了__global__修饰符以外，CUDA编程中还提供了另外两种修饰符以说明程序运行的位置，如：

__device__修饰符，用于修饰设备函数，即在GPU上运行的函数，该函数可以被其他设备函数调用，但是不能被主机函数调用。
__host__修饰符，用于修饰主机函数，即在CPU上运行的函数，该函数可以被其他主机函数调用，但是不能被设备函数调用。

注意，这两个修饰符是可以叠加的，叠加后的函数即可以在主机上运行也可以在设备上运行；但是，这两个修饰符都不能与__global__修饰符共同修饰一个函数。另外，修饰符与void声明的顺序可以任意交换。

CUDA程序调试

以下介绍一下部分通用的调试、编译和测试工具。

调试思路

之前一直忽略了这一点，反正这才是最重要的调试环节，尤其是在多卡服务器上进行调试，稍不注意就造成内存非法访问，那个错误已经见惯不惊了。需要明确的调试思路是：

先语法后程序
先单卡后多卡

第一点基本上所有的软件工程都会涉及到，第二点的先单卡后多卡真的很重要！很多错误往往就是在多卡的情况下才发生的。切记不要因为自己的程序只在单卡上运行就忽略了显式指定单卡的步骤！

请使用环境变量CUDA_VISIBLE_DEVICES显式指定使用的GPU设备。

在调试程序运行前，请强制显式指定该变量，保证调试程序运行在单张GPU显卡上，例如：

1	export CUDA_VISIBLE_DEVICES=0

该变量将保证所有CUDA程序仅在0卡上运行，事先排除了可能的显卡内存错误分配等设备配置问题。在单卡上完成一系列测试后，再设置该变量为多卡：

1	export CUDA_VISIBLE_DEVICES=0,1,2,3

然后完成调试，事实上，在程序运行中也可以摆脱直接指定设备的用法（如Pytorch的to函数，因此有的代码中完全不含有指定设备的函数，一定程度上也起到了简化代码增强可读性的作用），但是如果使用本方法，必须要在文档中予以指出，避免复现者产生疑惑。

编译工具

编译CUDA程序必须使用英伟达开发的编译工具nvcc，编译有两种方式，分别是即时编译just-in-time compilation，即著名的JIT编译，另一种则是和普通的C++程序一样，先编译成中间代码，再链接成可执行文件。编译过程中输入至nvcc命令的参数列表其中有一项非常关键，即-gencode arch=compute_XY, code=sm_ZW/compute_ZW。其中表示先指定一个虚拟架构compute以确定代码中使用到的CUDA功能模块，然后使用一个真实的架构以确定可执行文件能够使用的GPU；通常为了保证代码在不同的GPU设备上也能运行，需要增加多个架构选项，也即简单地将上述参数重复多次，而初学者可以使用简单的一个参数完成大部分情况的应付：-arch=sm_XY参数。

对于工程项目文件，使用nvcc编译就与使用gcc编译一样麻烦，因此使用cmake完成编译文件的生成和配置就非常关键，在使用cmake文件时，在CMakeLists.txt文件中需要增加如下的配置：

1	project(<PROGRAM_NAME> LANGUAGES CXX CUDA)

该配置将告诉make程序使能CUDA编译，则C++部分的代码由gcc编译，而CUDA部分的代码由nvcc编译。通常生成的make文件自动就会配置好适合当前编译环境的编译文件，然后就可以如同编译一个C++程序一样编译CUDA程序了，参考自博客。

内存错误检查工具

CUDA编程出bug乃兵家常事，反复debug就是编译的正道，英伟达官方提供了内存错误检查工具cuda-memcheck，该工具指令共计包括memcheck、racecheck、initcheck和synccheck四个子指令，分别用于检查内存错误、检查内存竞争、检查内存初始化和检查同步错误。使用方法示例如下，这里仅以第一个内存检查为例：