CUDA programming

cuda c 的后缀是 .cu, .cu的后缀是为了告诉VS（visual studio）该文件是要传送给NVCC 处理的。 cuda 程序部分是用 NVCC compiler 编译的。 这个编译器的作用就是将 the cuda code  变成 PTX（CUDA的汇编语言）， 剩下的普通的 C++ code 被 C++ 编译器编译。

注意， PTX code 是嵌入在 EXE中， 然后just在该程序运行之前， 被我们的device driver（指代驱动GPU的程序）编译。 这样PTX 可以编译成适合你的GPU的类型的 machine code。

1.1 Header and Libraries

使用cuda c， 相应的头文件和链接库是必不可少的。

#include <cuda.h>

link 库：

cudart.lib

完成了对VS的这些配置， 我们就可以使用cuda 为我们提供的API程序了。

 1.2  host and device

host 和 device ， Kernel使我们常常会碰到的术语。 

简而言之， Host 指的是 CPU， 控制值自己的system RAM, disks, 以及外围器件（peripheral device）

Device 指的是GPU， 也拥有自己的RAM ， 但是无法访问系统的RAM， 无法使用系统的 disks, 以及无法控制 外围器件。

Kernel 就是一个函数， 只不过这个函数不是运行在CPU， 而是在GPU上运行的。 通常一次性的（同时）被GPU的100个或者1000个threads 运行着， 一个线程运行kernel的copy。 不难看出， kernel是用cuda c 语言写的。  CPU的功能就是只需通过特殊的语法命令 告诉GPU启动多少个 thread（线程）去执行kernel。

cuda c 和 c 语言很像， 只不过 cuda c 中没有递归函数。 kernel 返回值总是void, 而且kernel 无法使用变量个数的 参数（variable numbers of arguments）， 并且无法使用 system memory。

kernel 既然也是函数， 那么如何区分他和普通的函数的区别呢？？

这就用到了 function qualifiers。 具体有__global__， __device__，  __host__。

（1）__global__(注意下划线是因为输入下两个下划线)

表明该函数是有host(即 CPU) 调用， 运行在 device（即GPU）上面， 该函数是kernel

（2）__device__

该函数被device (即GPU)调用（call）， 运行在device上面， 也是kernel， 又称为 kernel helper functions

(3)__host__

或者函数没有任何的qualifier（默认为__host__）， 或者有 __host__， 都表示  该函数是普通的函数， 不是kernel， 将会在CPU上运行的。

NVCC compiler 会将这部分的程序返还给 VS的compiler 进行编译。 

2 cuda 内存管理函数

由于GPU无法访问 系统的RAM， 所以数据必须拷贝到GPU的RAM中去， 处理完后， GPU（device）还要将处理后的数据返还给CPU(host)， 这样就会出现反反复复的复制， 影响运行的速度和效率。 解决办法就是使用cuda 的API函数， 用于GPU的内存管理， 接口如下：

2.1 device 上的内存分配

cudaMalloc(void** devPtr, size_t sizeInBytes);

此函数在GPU上的global memory 预定（reserve）memory， 预定的内存使用 GPU 的 指针指向。 预定的内存大小用第二个参数指定。

值得注意的是GPU的内存（如1G）远远小于CPU的系统内存（system RAM, 例如4G）。 anyway , 此函数是在device 上分配的内存。

2.2 device 上的内存释放

cuda 释放由cudaMalloc分配的内存的函数是：

cudaFree(void **devPtr)

该函数的目的就是释放device上被占用的， 但是现在不用了的内存。 由于device具有的内存有限， 所以never forget to free memory once you are done。

2.3 device 和host 之间的整块内存中数据的复制接口函数

cudaMemcpy(void* dest, void* src, size_t sizeInBytes, enum direction);

该函数的作用就是实现device 和 host memory之间数据的复制。 dest 是复制的目的地， src 是复制的源头， 即把src 指针指向的数据复制到dest 指针指向的内存中去。 方向有如下两个：

cudaMemcpyDeviceToHost
cudaMemcpyHostToDevice

有了上述的基本知识， 我们接下来编写一个程序， 用GPU实现两个整数的相加， 得到结果。 当然这只是一个toy example， 我们直接可以使用CPU实现两个数的和， 而且速度是GPU的几百甚至几千倍了， anyway， 这只是一个例子， 帮助我们了解GPU编程的一些流程。

为了在GPU上实现这个功能， 我们遵循如下流程：

(1) create 两个integers for host(即CPU)

(2) 在device上为两个integers 分配 device 的memory

(3) 将host上的两个整数复制到device 的 memory上

(4) 调用kernel 实现device这两个整数的相加

(5)将计算的结果复制到host的memory

(6)CPU打印出结果

(7) 释放我们在device 上分配的内存。

最后注意， 启动kernel的syntax 如下：

someKernel<<<1, 1>>>(parameters)

<<<n, k>>>的意思是启动n 个 blocks，  每个blocks k 个threads。 所以总共启动了 n * k 个 线程。

上面的意思是我们启动了1 block ， 1 thread per block， 也就是一个线程。

#include <iosteam>
#include <cuda.h>

using namespace std;

__global__ void AddIntsCUDA(int* a, int* b) {
    a[0] += b[0];
} 
int main() {
    int a= 5;
    int b = 4;
    int *d_a;
    int *d_b;
    
    cudaMalloc(&d_a, sizeof(int));
    cudaMalloc(&d_b, sizeof(int));
    
    cudaMemcpy(d_a, &a, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &b, sizeof(int), cudaMemcpyHostToDevice);
    
    AddIntsCUDA<<<1, 1>>>(d_a, d_b);
    
    cudaMemcpy(&a, d_a, sizeof(int), cudaMemcpyDeviceToHost);
    cout << "The answer is " << a << endl;
    
    cudaFree(d_a);
    cudaFree(d_b);
    return 0;
}

命名为 myFirst.cu， VS编译运行。