CUDA学习笔记-1: CUDA编程概览

1.GPU编程模型及基本步骤

cuda程序的基本步骤如下：

• 在cpu中初始化数据
• 将输入transfer到GPU中
• 利用分配好的grid和block启动kernel函数
• 将计算结果transfer到CPU中
• 释放申请的内存空间

从上面的步骤可以看出，一个CUDA程序主要包含两部分，第一部分运行在CPU上，称作Host code，主要负责完成复杂的指令；第二部分运行在GPU上，称作Device code，主要负责并行地完成大量的简单指令（如数值计算）；

2.基本设施

运行在GPU中地函数称作kernel，该函数有这么几个要求：

• 声明时在返回类型前需要添加"__globol__"的标识
• 返回值只能是void

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

这就是一个合规的核函数。

除了声明时的不同，和函数的调用也是不一样的，需要以 “kernel_name <<< >>>();”的形式调用。而在尖括号中间，则是定义了启用了多少个GPU核，学习这一参数的使用，我们还需要知道下面几个概念：

• dim3：一种数据类型，包含x,y,z三个int 类型的成员，在初始化时一个dim3类型的变量时，成员值默认为1
• grid : 一个grid中包含多个block
• block: 一个block包含多个thread

我们以一种更抽象的方式来理解GPU中程序的运行方式的话，可以这么看：

GPU中的每个核可以独立的运行一个线程，那我们就使用thread来代表GPU中的核，但一个GPU中的核数量很多，就需要有更高级的结构对全部用到的核进行约束、管理，这就是block(块)，一个块中可以包含多个核，并且这些核在逻辑上的排布可以是三维的，在一个块中我们可以使用一个dim3类型的量threadIdx来表示每个核所处的位置,threadIdx.x、threadIdx.y、threadIdx.z分别表示在三个维度上的坐标；此外，每个块还带有一个dim3类型的属性blockDim，blockDim.x、blockDim.y、blockDim.z分别表示该block三个维度上各有多少个核，这个block中的总核数为blockDim.x * blockDim.y * blockDim.z；

我们一次使用的多个block，最好能使用一个容器把他们都包起来，这就是grid，类比于上文中thread和block的关系，block和grid也有相似的关系。我们使用blockIdx.x、blockIdx.y、blockIdx.z表示每个block在grid中的位置；同样，grid也具有gridDim.x、gridDim.y和gridDim.z三个属性以及三者相乘的总block数。

知道了上面这些知识后，我们可以对“kernel_name <<< >>>();”中尖括号中的参数做一个更具体的解释，它应该被定义为在GPU中执行这一核函数的所有核的组织形式，以"kernel_name <<< number_of_blocks, thread_per_block>>> (arguments)"的形式使用，一个典型的示例如下：

int nx = 16;
int ny = 4;
dim3 block(8, 2); // z默认为1
dim3 grid(nx/8, ny/2);
addKernel << <grid, block >> >(c, a, b);

这一示例中创建了一个有(2*2)个block的grid，每个block中有(8*2)个thread，下图给出了更直观的表述：

需要注意的是，对block、grid的尺寸定义并不是没有限制的，一个GPU中的核的数量同样是有限制的。对于一个block来说，总的核数不得超过1024，x、y维度都不得超过1024，z维度不得超过64，如下图

对于整个grid而言，x维度上不得有超过(2^{32}-1)个thread，注意这里是thread而不是block，在其y维度和z维度上thread数量不得超过65536.

在cuda编程中我们经常会把数组的每一个元素分别放到单独的一个核中处理，我们可以利用核的索引读取数组中的数据进行操作，但由于block、grid的存在，索引的获取需要一定的计算，在exercise2中给出了一个3D模型中取值的训练，实现如下

__global__ void print_array(int *input)
{
    int tid = (blockDim.x*blockDim.y)*threadIdx.z + blockDim.x*threadIdx.y + threadIdx.x;
    int xoffset = blockDim.x * blockDim.y * blockDim.z;
    int yoffset = blockDim.x * blockDim.y * blockDim.z * gridDim.x;
    int zoffset = blockDim.x * blockDim.y * blockDim.z * gridDim.x * gridDim.y;
    int gid = zoffset * blockIdx.z + yoffset * blockIdx.y + xoffset * blockIdx.x + tid;
    printf("blockIdx.x : %d, blockIdx.y : %d, blockIdx.z : %d,gid : %d, value: %d
", blockIdx.x, blockIdx.y, blockIdx.z, gid, input[gid]);
}

3.数据在host和device之间的迁移

我们前边提到，cuda的编程步骤是将数据移入GPU，待计算完成后将其取出，官方对可能涉及到的内存操作类的操作都给出了接口。

首先是cudaMemCpy函数，其定义为

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

该函数是将数据从CPU移入到GPU或者从GPU移出到CPU中，参数0指向目标区域的地址，参数1指向数据的源地址，参数2表示要移动的数据的字节数，最后一个参数表示数据的移动方向（cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost或cudaMemcpyDeviceToDevice）

此外，对应C语言的内存空间操作，cuda也推出了CudaMalloc, CudaMemset, CudaFree三个接口

cudaError_t  cudaMalloc ( void** devPtr, size_t size );
cudaError_t  cudaMemset ( void* devPtr, int  value, size_t count );
cudaError_t  cudaFree ( void* devPtr );

这里需要注意的一个点是cudaMalloc的第一参数的数据类型为void**，这一点怎么理解呢？

这里我们结合一个示例进行解释:

int *d_input;
cudaMalloc((void **) &d_input, bytesize);

之所以使用void，是因为这一步只管分配内存，不考虑如何解释指针，所以只需要传入待分配内存的地址，不需要传入具体的类型，其他API中的 void* 也是同理。为什么是两个*呢，这是因为我们在定义d_input时是定义了主存中的一个指针，它指向主存中的一个地址；而&d_input则是取得了存储该指针值的地址，cudaMalloc利用这一地址将在GPU中分配给该缓冲区的首地址赋值给d_input。

利用上述的几个接口函数，我们就可以实现一个基本的cuda程序的主函数：

int main()
{
    const int arraySize = 64;
    const int byteSize = arraySize * sizeof(int);

    int *h_input,*d_input;
    h_input = (int*)malloc(byteSize);
    cudaMalloc((void **)&d_input,byteSize);

    srand((unsigned)time(NULL));
    for (int i = 0; i < 64; ++i)
    {
        if(h_input[i] != NULL)h_input[i] = (int)rand()& 0xff;
    }

    cudaMemcpy(d_input, h_input, byteSize, cudaMemcpyHostToDevice);

    int nx = 4, ny = 4, nz = 4;
    dim3 block(2, 2, 2);
    dim3 grid(nx/2, ny/2, nz/2);
    print_array << < grid, block >> > (d_input);
    cudaDeviceSynchronize();

    cudaFree(d_input);
    free(h_input);

    return 0;
}

其中 cudaDeviceSynchronize(); 的作用是在此处等待GPU中计算完成后再继续执行后续的代码。

4 错误处理

在C++中，可以使用异常机制处理运行时错误，而cuda编程中由于Host和Device共同使用，难以利用异常机制，因此，cuda提供了检测运行时错误的机制。

看上面的API时会发现，每个函数的返回值类型都是 cudaError_t ，这正是cuda提供的错误检测机制，如果返回值是cudaSuccess则说明执行正确，否则就是出现了错误。可以使用 cudaGetErrorString( error )获取返回值的代表的错误的文本。前面的代码中没有使用这一机制主要是为了便于阅读，但实际的使用中这一机制是必不可少的，也会看到VS生成的demo代码中就包含着大量的错误检测代码

	cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

	cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
...
...

5 其他

1. 不同的block_size计算耗时会不同，可以多尝试后选择计算的更快的参数（学DL的调参是吧，这也搞黑盒？）；考虑GPU的计算时间时要考虑数据移入移出GPU的时间。

2. 不同的GPU有不同的性质，设备中也可能存在多个GPU，在设计程序时需要考虑这些问题，cuda也提供了访问这些信息的接口

   // 获取设备数量
   int deviceCount = 0;
   cudaGetDeviceCount(&deviceCount);
   
   //获取第一个设备的各项性质
   int devNo = 0;
   cudaDeviceProp iProp;
   cudaGetDeviceProperties(&iprop, devNo);