CUDA基本语法

示例：向量加的CUDA实现

这里通过向量加计算这个例子介绍CUDA的基本知识，主要包括内存操作、核函数（kernel function）以及线程配置等。

向量加问题描述

向量a，b， c为有N个元素的一维向量，计算c=a+b.
由于c中每个元素的计算与其他的元素间没有依赖关系，所以该问题适合于并行计算。见下图：

向量加操作示意图

CUDA是对现有编程语言的扩展，比如CUDA C是对C语言的扩展，大部分为C语言语法。下面我们分别来看下该问题的CPU代码实现以及CUDA代码实现，然后针对CUDA代码中新出现的语法作一一说明。

CPU 代码实现

CPU向量加的函数实现：

void vecAdd (int n, int *a, 
             int *b, int *c)
{
    for(int i = 0; i < n; i++)
    {
        c[i] = a[i] + b[i];
    }
} 

CPU代码的main函数：

void main(){
    int size = N * sizeof(int);
    int *a, *b, *c;
    a = (int *)malloc(size);
    b = (int *)malloc(size);
    c = (int *)malloc(size);
    memset(c, 0, size);
    init_rand_f(a, N);
    init_rand_f(b, N);

    vecAdd(N, a, b, c);
}

这一部分非常简单，没有需要解释的地方，下面我们来看一下CUDA代码的实现有什么不同之处。

CUDA代码实现

CUDA向量加的函数实现：

__global__ void
vectoradd (int *a, int *b, 
            int *c, int n)
{
    int i = blockDim.x * 
             blockIdx.x +
             threadIdx.x;

    if (i < n)
    {
        c[i] = a[i] + b[i];
    }
}

CUDA代码main函数实现：

int main(void) {
    size_t size = N * sizeof(int);
    int *h_a, *h_b; int *d_a, *d_b, *d_c;
    h_a = (int *)malloc(size);
    h_b = (int *)malloc(size);
    ...
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);

    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    vectoradd<<<Grid, block>>>(d_a, d_b, d_c, N);
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b);
    return 0;
}

代码对比与分析

对比CPU代码以及CUDA代码，可以发现两者的部分代码是相似的，并且实现流程一致。我们这里主要关心的是两者的不同之处：

1. 向量加函数的实现处不同。CUDA代码实现中函数最前面多出一个标识符__global__，该标识符表明这段函数会在GPU端执行，我们通常称这样的函数为核函数（Kernel function）。在函数的内部int i = blockDim.x * blockIdx.x + threadIdx.x;表示的是在GPU上执行时的线程号。另外CPU代码使用for循环对向量中的每个元素依次进行操作，而在CUDA代码中则省去了该循环操作，取而代之的是所有的元素同时操作，每个CUDA线程对应一个向量元素，这也就是并行计算。
2. 内存的分配与释放。观察两者main函数的实现，总共可以发现两处不同之处，一是内存的分配与释放，另外一个是向量加实现函数的调用。前面我们也提到CPU与GPU分别具有各自的内存系统，数据可以在CPU内存以及GPU内存之间传递，要想在GPU上进行计算，我们首先要把待处理的数据放到GPU端，GPU完成计算后再把结果拷贝回CPU端。GPU端内存的分配与释放分别有相应的CUDA Runtime API来完成。
3. 实现函数的调用。与CPU代码的另外一个明显的不同就在于实现函数的调用方式不同。从CUDA代码可以看出，其调用方式为kernel_name<<<grid, block>>>(…)，其中<<<grid, block>>>为线程的配置信息，它表示核函数在GPU上执行时会有多少线程。
   上面简要说明了CUDA代码相对于CPU代码的不同之处，后面将针对这几个部分作详细的说明。