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 模式。