欢迎您访问程序员文章站本站旨在为大家提供分享程序员计算机编程知识!
您现在的位置是: 首页

CUDA 学习笔记 2

程序员文章站 2024-02-27 12:40:39
...

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 学习笔记 2

默认模板生成了参考代码,直接编译通过,并可在命令行运行,表明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/cudaFreecudaMemcpy 函数调用,作用是先分配GPU上的内存(就是显存),然后把要计算的数据拷贝到显存,kernel运行的时候,只能从显存获取数据,并把计算结果保存到显存,然后CPU再把数据从显存读回系统内存。这一切都因为GPU不能直接读写系统内存。

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


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