CUDA 学习笔记 2

CUDA初印象

CUDA的安装

CUDA的安装比较简单，NIVIDA把驱动、Runtime、SDK、Toolkit 都打包在一起，下载按部就班地安装即可。

需要注意的是，在Windows 环境下，目前CUDA（最新版本8.0）只支持64位 OS。其次，即使在64位环境下可编译32位的应用，但CUDA对32位的支持也有限制，具体参看官方文档，总之推荐使用64位程序。

编译器尽量采用微软的，并且Visual Studio版本推荐VS2012 - VS2015， VS2017 暂不能自动集成CUDA开发环境，如果手动配置的话比较麻烦。

最后，这一切的前提当然是得有NVIDIA的显卡。

Linux下的安装基本类似，但往往需要配置编译环境，暂且不表。

Hello CUDA

安装完成后，就可以编写并运行CUDA程序了。
用VisualStudio 2015 创建一个CUDA 工程，

默认模板生成了参考代码，直接编译通过，并可在命令行运行，表明CUDA安装正确。

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

int main()
{
    cudaError_t cudaStatus;
    const int size = 5;
    const int a[size] = { 1, 2, 3, 4, 5 };
    const int b[size] = { 10, 20, 30, 40, 50 };
    int c[size] = { 0 };

    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;

    cudaStatus = cudaSetDevice(0);
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));

    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    addKernel <<< 1, size >>>(dev_c, dev_a, dev_b);
    cudaStatus = cudaDeviceSynchronize();

    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n", 
            c[0], c[1], c[2], c[3], c[4]);

    cudaStatus = cudaDeviceReset();
    return 0;
}

从这段代码可以看到，这里cuda的源代码的后缀名是cu，语法基本和C语言一致。其中有一些特殊的语法和函数与C不同。

这段代码的功能是用GPU完成两个数组的相加。核心是中间的 addKernel <<< 1, size >>>(dev_c, dev_a, dev_b)这句，作用就是在GPU 上执行add 指令。

这句在GPU上执行的函数就叫kernel，为此，在声明的时候跟普通函数略有区别，就是在最前面有__global__标示符。那么相应的，没有此标示符的函数和代码，依旧运行在CPU上。

kernel函数在GPU 中执行的时候，会在每个线程里运行一边，直到计算完成。执行时所需的线程数通过<<<...>>>来指定。每个线程在运行的时候，都有一个线程标记，这里让每个线程都计算数组的一个元素，这样5个线程运行起来之后，计算就完成了。

在这段kernel函数前后，都有一系列的cudaMalloc/cudaFree，cudaMemcpy 函数调用，作用是先分配GPU上的内存（就是显存），然后把要计算的数据拷贝到显存，kernel运行的时候，只能从显存获取数据，并把计算结果保存到显存，然后CPU再把数据从显存读回系统内存。这一切都因为GPU不能直接读写系统内存。

由于有了上述的特殊语法，在编译的过程中，c语言的编译器不可能编译这段代码，所以需要调用cuda的编译器，从下面的命令中也可以看出这一点：

CUDA 驱动程序模型

之前提到，使用CUDA可以在进行GPU并行计算的时候，避免把数据重组成图形数据的形式，但事实上我们可以认为是CUDA runtime 完成了这个转换过程，并提供简单接口给应用程序。所以本质上CUDA 还是运行在显卡驱动之上的。

Windows 下通用的驱动程序架构称为WDM （Windows Driver Model），几乎所有现代的硬件驱动程序都是基于此架构的（更新的WDF模型，或者特定设备类型的驱动模型基本都是基于此，在此不过多介绍），显卡驱动程序自然也是基于此架构编写的。只不过在Vista 之后，微软推出了新的显卡驱动模型，称为WDDM (Windows Display Driver Model), 这与WDM并不兼容，而且从Windows 7 开始，WDM的显卡驱动不再支持。

因此，在大部分情况下，CUDA的驱动程序模型自然也必须是WDDM架构的了。但在某些情况下，GPU并不是作为显示设备的，比如服务器或者远程桌面情况下，这时候GPU就单纯作为计算设备存在的，所以就无需遵守WDDM的驱动模型规范，而是把它当做一个普通设备来编写驱动程序，NIVIDIA可以随心所欲的实现GPU的功能和接口，这时候就可以采用通用的WDM 驱动模型。

这就是CUDA 驱动模型的由来，在CUDA中，这两种模型分别被称为WDDM 模型和TCC (Tesla Compute Cluster) 模型。TCC 模型因为更直接，所以性能更高。两种模型之间可以通过驱动安装包所带的nvidia-smi工具切换。但通常情况下我们都采用WDDM 模型，而且GeForce 系列的GPU 并不能支持 TTC 模式。