CUDA学习——Chapter 2(4)内存空间布局对核函数性能的影响(3)

本文探讨了在CUDA编程中,使用二维网格和一维块对矩阵求和的影响。通过实例展示了如何构建这种结构,并分析了不同内存空间布局对核函数性能的效应,指出传统布局可能并非最优,改变内存布局可以提升性能。实验结果显示,不佳的内存布局可能导致GPU性能低于CPU。

第二章

3.使用二维网格和一维块对矩阵求和

那么刚才使用了一维网格一维块的结构来对矩阵进行求和,现在我们来尝试着用二维网格一维块的结构来进行矩阵求和的运算。
二维网格一维块的结构如下:
二维网格一维块的结构示意图
那么也就是一个线程处理一个对应的矩阵元素。
线程的内存索引依然满足:idx=iy*nx+ix;

依葫芦画瓢嘛,我们可以写出使用二维网格一维块的核函数

__global__ void sumMatrixOnGPUMix(float *MatA,float *MatB,float *MatC,int nx,int ny)
{
    unsigned int ix=threadIdx.x+blockIdx.x*blockDim.x;
    unsigned int iy=blockIdx.y;
    unsigned int idx=iy*nx+ix;
    if(ix<nx&&iy<ny)
        MatC[idx]=MatA[idx]+MatB[idx];
}

因为线程本身没有y方向上的量,所以iy并不会加上threadIdx.y,而块是一维的,所以blockDim.y缺省为1,可以省略。

设置块和网格的大小

dim3 block(32);//x方向上有32个线程的块
dim3 grid((nx+block.x-1)/block.x,ny);//x方向上创建可以容纳nx个线程的k个块,其中k=nx/block.x的向上取整。y方向上创建可以容纳ny的线程的ny个块。

还是用上次的代码,跑一下这个结果
该代码稍微有点变动,使用CPU时钟频率来获得更精确的时间花销计算。
二维网格一维块的矩阵求和计算
使用这个修改过的代码来修改2-7,2-8,得到二维网格二维块,一维网格一维块,二维网格一维块的运行时间(Release模式):
二维网格二维块
一维网格一维块
二维网格一维块
由此,我们得出结论:
1.传统的(也就是朴素的)核函数内存分割布局方式不一定能获得最佳性能。
2.改变传入的global memory布局对核函数的性能有不同的影响。
3.多试几种布局方式可能能获得更好的性能。

附:CPU计算该矩阵求和的时间:
CPU单独计算矩阵求和的时间
可以看到,布局不好的时候,GPU甚至会比CPU慢上许多。

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值