CUDA 是专为 GPU 设计的编程接口，由于GPU 和 CPU 有不同的硬件，编程的方式也有很大的差别。要理解CUDA 的概念，应该从硬件开始着手。下图是 CPU 和 GPU 的一个简化图，图中反映了 CPU 和 GPU 的几点不同：

1. CPU 有更大的 Local Cache，而 GPU 的 Cache 和 Local Memory 相对较少。
2. GPU 有大量的寄存器，用来支持更多的线程。相反，CPU 寄存器较少，支持的线程数目也相对较少。
3. GPU 有更多的 SIMD 单元，CPU 的 SIMD 数量较少。
4. CPU 有复杂的控制逻辑而 GPU 的控制逻辑相对简单。

CPU 和 GPU 的设计背后反映的是的两者的工作场景。CPU主要考虑的是低延迟，所以CPU的设计上主要针对的也是降低延迟。通常，CPU 都有一个强大的 ALU，以减少操作的延迟。而 Cache 容量更大能提升访存时 Cache 的命中率，由于 Cache(SRAM)比内存(DRAM)快很多， 也就支持更快速的访存操作。复杂的控制逻辑能帮助 CPU 进行分支预测、数据前送。这些设计都用来减少 CPU 的延迟以达到更快的响应速度。GPU 的 ALU 数量更多，但是每一个操作的延时都更长，它们通过管道化(pipelined)来支持大量的操作，通过创建大量的线程，在
ALU 的每个阶段都对应一个线程来充分使用这些。可以说 CPU 注重的是质，GPU 靠的是量。

很多设计良好的应用都利用了 CPU 和 GPU 的特性，用 CPU 执行顺序代码，用 GPU 来执行并行代码。

2. CUDA C

CUDA C 是 NVIDIA 开发的一套异构并行编程接口。它提供了

• 线程的组织方式
• 启动并行计算的接口
• 线程索引到数据索引的映射

CUDA 的运行模型包括了主机(CPU)执行的代码和设备(GPU)执行的代码。用于 GPU 执行的代码为叫做kernel，一个
CUDA kernel 通常会生成大量并行执行的线程，通过网格(grid)来控制线程的数量。所有的线程都执行相同的 kernel 代码。每个 线程都有其索引用来计算内存地址。一定数量的线程组成一个线程块(Thread Block).

下面以向量加法 A+B = C 来展示 CUDA C 的基本概念。

1.向量加法函数框架

1 #include<cuda.h>
2     ...
3 void vecAdd(float* A,float * B ,float * C ,int n){
4     //0.分配设备内存用于存储 A,B,C
5     //1.将主机内存中的A,B拷贝到设备内存中
6     //2.启动 kernel 函数计算 C 的值
7     //3.将计算结构拷贝回 C 中。
8     //4.释放设备内存
9 }

2.内存操作

设备内存即我们常说的显存，其容量通常能到数 GB，比如我笔记本上的 GT750M 容量为2GB，而 GTX TITAN 能到12GB(某些型号)。GPU 要想对数据进行运算，首先要把数据装载到显存中。而在装载之前首先要分配内存。CUDA 提供的分配内存 API 和 C 标准库的分配内存 API 很相像： cudaError_t
cudaMalloc(void **devPtr, size_t size); 此函数第一个参数是一个指针的地址，该指针用来保存申请到的显存的首地址。第二个参数要申请的显存大小，以字节为单位。与此对于的释放显存的函数为 cudaError_t
cudaFree(void *devPtr); 参数是指向要释放的显存首地址的指针。为设备分配A的显存的代码如下：

float * deviceA;//指向显存的指针
    //计算所需显存大小
    int size  = n * sizeof(float);

    cudaMlloc((void**)&deviceA,size);
    . . .
    cudaFree(deviceA);

注意 deviceA 是指向设备内存地址的，因此在主机代码中不可以解引用 deviceA 指针。分配完显存后需要把主机内存中的数据拷贝到显存中，CUDA 同样提供了一个和 C 标准库函数类似的API: cudaError_t
cudaMemcpy(void *dst, const void *src,size_t count, enum cudaMemcpyKind kind); 其中第一个参数是目的内存的指针，第二个参数是源内存的指针，第三个参数是要拷贝到设备内存中的数据大小，以字节为单位，第四个参数是拷贝的方式。比如把 A 中的数据拷贝到 deviceA 的代码是cudaMemcpy(deviceA,A,sizeof(float)*n,cudaMemcpyHostToDevice) 其中cudaMemcpyHostToDevice 是CUDA
中预定义的常量，表示从主机内存拷贝到设备内存。那么计算结果如何返回给主机呢？假设计算出的结果保存在 deviceC 中，那么从 deviceC中拷贝到 C 中的函数就是cudaMemcpy(C,deviceC,sizeof(float)*n,cudaMemcpyDeviceToHost)。

3.线程的组织

CUDA 中线程的组织由网格管理。每个网格包含的是线程块，它是一个线程块的数组，一个网格内所有线程块大小都一样。也就是说，每个线程块中包含相同数目的线程。

下图中，每个线程块包含256个线程，当然不是必须256个，这个值可以通过启动 kernel 函数的参数来设置，最大不超过1024，但通常是32 的倍数(因为 一个warp大小是32)。和这个参数对应的变量叫 blockDim，图中的 blockDim.x 的值就是256(线程块最大支持3维布局，所以是一个包含 x,y,z 三个维度的结构体，但是我们只用到了第一维，因此只需要访问 blockDim.x）。每个线程都有一个线程 ID ，用 threadIdx（threadIdx 同样是一个包含 x,y,z 的结构体，本例中是一维布局）
保存。在图中的例子中数据索引可以用i=blockIdx.x
* blockDim.x + threadIdx.x来获得。其中 blockDim.x 值为256,blockIdx.x 是线程块号，threaIdx.x是一个线程在线程块中的编号。本例中值从0-255，对第1号(从0 开始计数)线程块中的2号线程来说，它的 blockIdx.x 值为1，threadIdx.x 值为2。故它的索引值为258=1∗256+2,也就是说，他是所有线程中的258号，负责计数向量的第259个元素的和。

4.kernel 函数

向量加法的 kernel 函数：

1 __global__ void vecAddKernel(float * deviceA,float* deviceB,float* deviceC,int n){
2     int i = threadIdx.x + blockDim.x * blockIdx.x;
3     if(i<n) deviceC[i] = deviceA[i] + deviceB[i];
4 }

关键字__global__是CUDA
对 C 的扩展，表示此函数在设备上执行，但是只能从主机调用。最后，从主机调用这个 kernel 函数的代码为：

1 int vecAdd(float * A,float * B ,float * C, int n){
2     ...//申请内存，拷贝数据
3    ....
4    vecAddKernel<<<ceil(n/256.0),256>>>(deviceA,deviceB,deviceC,n);
5   ....//拷贝回主机内存，释放设备内存
6 }

代码的第4行就是启动 Kernel 函数的代码，<<<ceil(n/256.0),256>>>部分是对网格的配置。第一个参数要启动的线程块数量，第二个参数是每个线程块的线程数。用
ceil 函数向上取整保证了有足够多的线程块来处理向量。如果向量长度n 是1000，eil(n/256.0)=4，将会启动4个线程块，一共256*4 = 1024 个线程。从0号到999号线程都要进行计算。最后一个线程块中第1000-1023号线程跳过计算，kernel 函数中通过 if(i<n)来判断线程是否参加计算。

5.完整函数

最后，完整的向量加法函数：

 1 void vecAdd(float* A,float * B,float * C,int n){
 2     float * deviceA,deviceB,deviceC;
 3     int bytes = sizeof(float)*n;
 4     cudaMalloc((void**)&deviceA,bytes);
 5     cudaMalloc((void**)&deviceB,bytes);
 6     cudaMalloc((void**)&deviceC,bytes);
 7     cudaMemcpy(deviceA,A,bytes,cudaMemcpyHostToDevice);
 8     cudaMemcpy(deviceB,B,bytes,cudaMemcpyHostToDevice);
 9 
10     vecAddKernel<<<ceil(n/256.0),256>>>(deviceA,deviceB,deviceC,n);
11 
12     cudaMemcpy(C,deviceC,bytes,cudaMemcpyDeviceToHost);
13     cudaFree(deviceA);
14     cudaFree(deviceB);
15     cudaFree(deviceC);
16 }

代码中略去了错误处理，也不包含主机代码中读入数据的部分。此外，可以想到代码的效率不会很高，后面会对这个代码进行改进。