cuda合并访问的要求_【CUDA 基础】5.4 合并的全局内存访问

title: 【CUDA 基础】5.4 合并的全局内存访问

categories:

- CUDA

- Freshman

tags:

- 合并

- 转置

toc: true

date: 2018-06-04 21:34:22

Abstract: 本文介绍使用共享内存进行矩阵转置以减少内存的交叉访问

Keywords: 合并，转置

开篇废话

没废话，看以前的废话感觉自己像个傻瓜。。就像以后看我正在写的文字一样。

还记得我们矩阵转置的例子么，在全局内存部分介绍的：4.4核函数可达到的带宽

在4.4中我们当时只有共享内存这一种工具可以使用，为了达到最高效率，我们要配合一级缓存，二级缓存进行编程，来提高转置的效率，因为转置只能在行读取列写入或者列读取行写入之间选择一个，这样就必然会引发非合并的访问，虽然我们利用一级缓存的性质可以提高性能，但是我们今天会介绍我们的新工具共享内存，在共享内存中完成转置后写入全局内存，这样就可以避免交叉访问了。

基准转置内核

在介绍我们的神奇共享内存之前，我们最好先研究出来一下我们的问题的极限在哪，换句话说，我们需要清楚的知道我们最慢的情况(最简单的方式能达到的速度)以及最快的理论速度，理论速度可能会达不到，但是可以接近，最慢速度肯定可以超越，你永远可以写出更慢的程序，所以我们用最简单的方法作为下界，而用正行读取，然后不经变换的写入来作为上限，这一招我们在前面使用过，就是在4.4中，那次我们突破极限了(哈哈，很有可能是计时有问题)，但是正常来讲，极限是最好的参考值。

完整的代码在github:https://github.com/Tony-Tan/CUDA_Freshman(欢迎随手star? )

上限：

__global__ void copyRow(float * in,float * out,int nx,int ny)

{

int ix=threadIdx.x+blockDim.x*blockIdx.x;

int iy=threadIdx.y+blockDim.y*blockIdx.y;

int idx=ix+iy*nx;

if (ix

{

out[idx]=in[idx];

}

}

下限是我们的too young too naive版本，就是最常规的方法：

__global__ void transformNaiveRow(float * in,float * out,int nx,int ny)

{

int ix=threadIdx.x+blockDim.x*blockIdx.x;

int iy=threadIdx.y+blockDim.y*blockIdx.y;

int idx_row=ix+iy*nx;

int idx_col=ix*ny+iy;

if (ix

{

out[idx_col]=in[idx_row];

}

}

这两段代码中第一段并没有转置的功能，只是为了测试上限，第二段是naive的转置，前面也讲过，这里就直接贴结果了

copyRow的cpu计时和nvprof结果：

transformNaiveRow的cpu计时和nvprof结果：

我们可以得到下表：

核函数

CPU计时

nvprof计时

copyRow

0.001442 s

1.4859 ms

transformNaiveRow

0.003964 s

3.9640 ms

可以看出计cpu计时还是比较准的，在数据量比较大情况下，我们现在的矩阵大小是 212×2122^{12}\times 2^{12}212×212 的大小。

然后是加载和存储全局内存请求的平均事务数(越少越好)

copyRow:

transformNaiveRow:

接着我们就开始用共享内存进行操作了。

使用共享内存的矩阵转置