CUDA学习——Chapter 2(4)内存空间布局对核函数性能的影响(2)
第二章
3.使用一维网格和一维块对矩阵求和
前面我们使用了二维网格和二维块对矩阵求和,这种分割是非常直观的。那么现在我们就要使用抽象一点的概念,以稍复杂的编程来获得可能能取得的更高的性能。这就需要我们对映射有比较好的深入了解。
首先,我们再复习一下一维网格和一维块的结构:
其中,nx是x方向上最大的线程数,ny是一个线程需要处理的数据元素个数(因为这个块是一维的,照理来说是不应该有ny的)。所以这里只有ix是对线程的真正索引,iy是线程内部数据的索引(这个时候要把线程看成一个大的单位,里面是由ny个子线程组成的,每个子线程依次处理一个数据。但一定要记住,这个子线程实际上并不存在,是并行里面的串行)。
那么每个数据的idx就依然满足idx=iy*nx+ix;其中iy是从0迭代到ny的。
所以我们写出核函数:
__global__ void sumMatrixOnGPU1D(float *MatA,float *MatB,float *MatC,int nx,int ny)
{
unsigned int ix=threadIdx.x+blockIdx.x*blockDim.x;//获得x方向上的网格坐标
if(ix<nx)//防止越界
{
//从这里开始,就已经是线程里面的串行了
for(int iy=0;iy<ny;i++)
{
int idx=iy*nx+ix;//得到计算矩阵的坐标idx
C[idx]=A[idx]+B[idx];
}
}
}
接下来布置网格和块,我们题目说的是一维网格一维块,那就是:
dim3 block(32,1);//y方向上是1也就是没有
dim3 grid((nx+block.x-1)/block.x,1)//y方向上是1也就是没有
然后调用的时候按照二维网格和二维块的代码替换就可以了。
二维网格二维块和一维网格一维块的运行速度对比:
(单位错了,这里是以毫秒为单位的)
可以看到,这个运行速度是明显优于二维网格二维块的。
我们粗略地分析一下两者之间的不同,二维网格二维块是完全使用了GPU进行计算的。而一维网格一维块是采用了GPU+CPU的混合计算的,笔者使用的CPU是i7-7700HQ 8核,可能在CPU内部自己也存在并行计算,或是时钟频率高,从而导致性能远远优于仅使用GPU计算。
我们可以看到,用GPU纯计算的时间是238ms,而用CPU纯计算的时间是241ms。
如果我们用nvprof来观察时间,我们得到:
我们可以看到,核函数的运行时间是38ms,和用CPU时钟计算出来的时间相当。而cudaMemcpy才是其中最耗时间的部分。如果算上复制数据的时间,GPU在小规模的并行计算的性能可能还要差于CPU计算(当然,这是要看型号的,这里只是提供了一种分析方法而已)。