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<nx && iy<ny){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<nx && iy<ny){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:
接着我们就开始用共享内存进行操作了。
使用共享内存的矩阵转置
完整内容 https://face2ai.com/CUDA-F-5-4-合并的全局内存访问/