第2.3节演示了如何让标准C函数在设备端运行。通过将限定符__global__添加到函数并调用它使用特殊的尖括号语法,我们在GPU上执行了该函数。不过,这是一个启动在GPU上串行运行的核函数,这个核函数非常简单且效率很低。
前面说过,GPU最适合重复性和高度并行的计算任务。在本节中,我们将以向量加法为例讲解如何启动一个并行设备的核函数,以提升程序的效率。向量加法就是两个长度为 N 的向量相加,如图3-9所示,Vector A中的每个元素和Vector B中对应位置的元素相加,其结果保存为Vector C中对应位置的元素。
图3-9 向量加法示意图
用传统的CPU编程完成向量相加的逻辑如图3-10所示。
图3-10 用传统的CPU编程完成向量相加的逻辑
完整的向量相加纯CPU编程代码见示例代码3-1。
注:
● float*a=(float*)malloc(size)表示分配一段内存,并使用指针a指向它。
● for循环产生一些随机数,并放在分配的内存里。
● cpuVectorAdd(float*A,float*B,float*C,int n)表示要输入指向3段内存的指针名,也就是a、b、c。
● 函数gettimeofday用于得到精确时间,其精度可以达到微秒级,是C标准库的函数。
● 最后用函数free把申请的3段内存释放掉。
将上面的完整代码保存到cpuVectorAdd.cpp,然后用g++编译。
用MXMACA异构编程完成向量相加的逻辑如图3-11所示,这也是典型的MXMACA程序。该程序通常包括以下几个步骤。
(1)分配主机端的系统内存,并进行数据初始化。
(2)从主机端申请设备端全局内存(通常也叫显存),把要复制的内容从主机端内存复制到申请的设备端全局内存里。
(3)设备端的核函数对复制的内容进行计算,得到运算结果。
(4)把运算结果从设备端全局内存复制到申请的主机端内存里,并释放设备端的显存和主机端的系统内存。
(5)释放设备端分配的显存和主机端上分配的系统内存。
上述程序逻辑中最重要的一步是调用设备端的核函数来执行并行计算。核函数是在GPU线程中并行执行的函数,用__global__来声明,在调用时需要用<<<grid,block>>>来指定核函数要执行的线程数量。在MXMACA编程中,每个GPU线程都要执行核函数,并且每个线程会分配唯一的线程号(Thread ID),这个线程号可以通过核函数的内置变量threadIdx来获得。
下面,我们进行进一步细化这些内容。
● 设备端代码:读写线程寄存器;读写线程网格中全局内存;读写线程块中共享内存。
图3-11 用MXMACA异构编程完成向量相加的逻辑
● 主机端代码:显存、内存的申请(内存是插在主板内存插槽上的内存条,而显存是独立显卡上封装在GPU中的HBM);线程网格中全局内存的复制转移(显存、内存相互复制);内存、显存的释放。
● 申请显存的函数mcMalloc:在主机端完成显存的申请,得到相应的指针。
● 内存和显存之间相互复制的函数mcMemcpy:其参数包括终点的指针、起点的指针、复制的大小和模式(主机端到设备端、设备端到主机端、设备端之间的复制)。
● 释放显存的函数mcFree:释放指向显存的指针。
完整的向量相加MXMACA异构程序代码见示例代码3-2。
注:
● 首先要用__global__来进行修饰。
● gpuVectorAddKernel(float*A_d,float*B_d,float*C_d,int n)用于输入指向3段显存的指针名,也就是A_d、B_d、C_d。
● float*da=NULL;用于定义空指针。
● mcMalloc((void**)&da,size);用于申请显存,da指向申请的显存,注意mcmalloc函数传入指针的指针(指向申请得到的显存的指针)。
● mcMemcpy(da,a,size,mcMemcpyHostToDevice)用于把内存的内容复制到显存,也就是把a、b、c里的内容复制到d_a、d_b、d_c中。
● int threadPerBlock=256;int blockPerGrid=(n+threadPerBlock-1)/threadPerBlock;用于计算线程块和线程网格的数量。
● vecAddKernel<<<blockPerGrid,threadPerBlock>>>(da,db,dc,n);用来调用核函数。
● 函数gettimeofday用于得到精确的时间,其精度可以达到微秒级,是C标准库的函数。
● 最后的free函数用于把申请的3段内存释放掉。
将代码保存到文件gpuVectorAdd.cpp,然后用mxcc编译它。
编译得到可执行文件cpuVectorAdd.cpp和gpuVectorAdd.cpp之后,我们可以执行程序并比较在以下两种情形下的运行时间(注意要在Linux系统下运行)。
(1)GPU可以加速。在CPU上,执行程序10 9 次需要运行约1.03秒。
在GPU上,执行程序10 9 次只需要运行约3×10 -3 秒。
经测试验证,对于计算简单且并行度高的计算,GPU可以大幅提速。
(2)GPU不能加速。在CPU上,执行程序10 4 次需要运行1.2×10 -5 秒。
在GPU上,执行程序10 4 次却需要运行1.36×10 -4 秒。
经测试验证,对于少量计算,GPU的效率反倒不如CPU。
(1)在gpuVectorAdd.cpp中固定输入n的值为2048,将threadPerBlock的值分别修改为1/16/32/64/128/256/512,然后编译并运行,看看结果有哪些变化。如果threadPerBlock=1024,编译并运行会有什么现象?
(2)在第3.4节中,已经有了可以并行的线程块,为什么还要使用并行线程?并行线程相对于并行的线程块,有哪些不一样的地方?在gpuVectorAdd.cpp中固定输入n的值为256,分别修改threadPerBlock的值为1和256,然后编译并运行,看看结果有什么不同。
(3)在第3.6.3节中,对于少量计算,为什么GPU的效率不如CPU?
(4)请编写一个矩阵乘法程序。