内容出处:
1. 前言
这是一份简单的CUDA编程入门,主要参考英伟达的官方文档进行学习,本人也是刚开始学习,如有表述错误,还请指出。官方文档链接如下:
An Even Easier Introduction to CUDA | NVIDIA Technical Blog
developer.nvidia.com/blog/even-easier-introduction-cuda/
本文先从一份简单的C++代码开始,然后逐步介绍如何将C++代码转换为CUDA代码,以及对转换前后程序的运行时间进行对比
本文所使用的CPU为i7-4790,GPU为GTX 1080,那就开始吧。
2. 一份简单的C++代码
首先是一份简单的C++代码,主要的运行函数为add函数,该函数实现功能为30M次的for循环,每次循环进行一次加法。
编译以及运行代码:
不出意外的话,你应该得到下面的结果:
第一行表示add函数的运行时间,第二行表示每个for循环里的计算是否与预期结果一致。
这个简单的C++代码在CPU端运行,运行时间为85ms,接下来介绍如何将主要运算的add函数迁移至GPU端。
3. 把C++代码改成CUDA代码
将C++代码改为CUDA代码,目的是将add函数的计算过程迁移至GPU端,利用GPU的并行性加速运算,需要修改的地方主要有3处:
- 首先需要做的是将add函数变为GPU可运行函数,在CUDA中称为kernel,为此,仅需将变量声明符添加到函数中,告诉 CUDA C++ 编译器这是一个在 GPU 上运行并且可以从 CPU 代码中调用的函数。
那么修改后的add函数的调用也比较简单,仅需要在add函数名后面加上三角括号语法
指定CUDA内核启动即可,称为执行配置(execution configuration),用于配置程序运行时的线程,后续会讲到,目前先将其设置为
:
2. 那么为了在GPU进行计算,需要在GPU上分配可访问的内存。CUDA中通过Unified Memory(统一内存)机制来提供可同时供GPU和CPU访问的内存,使用cudaMallocManaged()函数进行分配:
同时,在程序最后使用cudaFree()进行内存释放:
其实就相当于C++中的new跟delete。
3. add函数在GPU端运行之后,CPU需要等待cuda上的代码运行完毕,才能对数据进行读取,因为CUDA内核启动时并未对CPU的线程进行固定,需要使用cudaDeviceSynchronize()函数进行同步。
4. 整体的程序如下所示:
使用nvcc对程序进行编译并运行:
或者使用nvprof进行速度测试:
不出意外的话,你会得到以下输出:
框出来的就是add函数在GPU端的运行时间,为4s。没错,就是比CPU端85ms还要慢,那还学个锤子。
4. 使用CUDA代码并行运算
好的回过头看看,问题出现在这个执行配置 <<<i,j>>> 上。不急,先看一下一个简单的GPU结构示意图,按照层次从大到小可将GPU按照 grid -> block -> thread划分,其中最小单元是thread,并行的本质就是将程序的计算模块拆分成多个小模块扔给每个thread并行计算。
再看一下前面执行配置 `<<<i,j>>>` 的含义,`<<<i,j>>>` 应该写成 `<<<numBlocks, blockSize>>>` ,即表示函数运行时使用的block数量以及每个block的大小,前面我们将其设置为`<<<1,1>>>` ,说明程序是单线程运行的,那当然慢了~~。下面我们以单个block为例,将其改为`<<<1,256>>>`,add函数也需要适当修改:
修改的部分也比较好理解,不赘述了,接下来运行看看结果:
你看,开始加速了吧,4s加速到了77ms。
那么,`<<<numBlocks, blockSize>>>` 的两个参数应该怎么设置好呢。首先,CUDA GPU 使用大小为 32 的倍数的线程块运行内核,因此 `blockSize` 的大小应该设置为32的倍数,例如128、256、512等。确定 `blockSize` 之后,可以根据for循环的总个数`N`确定 `numBlock` 的大小(注意四舍五入的误差):
当然因为变成了多个`block`,所以此时add函数需要再改一下:
这里index跟stride的计算可以参考上面GPU结构图以及下面的图(图取自An Even Easier Introduction to CUDA | NVIDIA Technical Blog),自行推算,较好理解。
搞定之后再编译运行一下:
看看,又加速了不是,通过提升并行度而加速,相比于CPU端(85ms)加速了接近一倍左右。
5. 结论
以上仅是一份简单的CUDA入门代码,看起来还算比较简单,不过继续深入肯定有更多的坑,期待后面有时间继续学习。