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