目录
背景
Nvidia的cuda 示例代码transpose_kernel.cu 提供了两种矩阵转置的实现 - transpose_naive和transpose,其中后者针对BANK Conflict 和 Memory Coalescing做了专门的优化,本文旨在结合代码说明transpose的算法和优化
BANK Conflict
百度bank conflict可以得到很多的说明,简单的来说cuda kernel的实现算法中如果出现bank conflict,那么就会导致share memory的访问效率降低
举个例子🌰 :假设硬件的Share Memory的Bank数目=4,WARP的线程数=4,这4个线程要访问Share Memory中一个4X4的矩阵的同一列,如下所示就会出现bank conflict(性能下降);而优化的方案很简单,就是给矩阵增加一列空白数据,就可以让4个线程错开访问不同bank,避免发生conflict(性能提升)
![](https://i-blog.csdnimg.cn/blog_migrate/3dffe03ccb1e9b5742e4af37f1d60f2f.gif)
Memory Coalescing
这个概念很直白,由于DRAM的物理特性决定了访问连续地址数据的时候可以做到1次访问多块数据(性能提升),而访问不连续地址数据的时候就变成多次访问多块数据(性能下降),参考下图示意
![](https://i-blog.csdnimg.cn/blog_migrate/8504cad01e74ff278be653c9be3c3839.gif)
transpose算法解析
__global__ void transpose(float *odata, float *idata, int width, int height)
{
//每个ThreadBlock分配了一个同样大小的share memory块,其中增加了一列空白数据避免bank conflict
__shared__ float block[BLOCK_DIM][BLOCK_DIM+1];
//根据ThreadBlock坐标和Thread坐标按照行主序的次序读取矩阵元素到share memory块中
//此时ThreadBlock坐标和Thread坐标都没有发生转置
unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;
unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;
//没有发生转置,所以xIndex的边界是width,yIndex的边界是height
if((xIndex < width) && (yIndex < height))
{
unsigned int index_in = yIndex * width + xIndex;
block[threadIdx.y][threadIdx.x] = idata[index_in];
}
//等待ThreadBlock中的所有线程完成上面的步骤
__syncthreads();
//这一步非常重要且最难理解:将ThreadBlock坐标系进行转置,而Thread坐标系保持不变
//在这样的情况下blockIdx.y 和 threadIdx.x都是x轴方向,blockIdx.x和threadIdx.y都是y轴方向
xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;
yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;
//在上面的坐标系转置条件下,xIndex的边界是height,yIndex的边界是width了
if((xIndex < height) && (yIndex < width))
{
//由于threadIdx.x是变化快的方向,所以xIndex连续变化可以确保index_out访问连续地址空间
//从而通过memory coalescing提升性能
unsigned int index_out = yIndex * height + xIndex;
//这一步也很重要且也难以理解:访问share memory块的行列和输出矩阵的行列要颠倒
//这样做的目的是将share memory块做转置后写入输出矩阵
//从而实现了ThreadBlock和Block内的元素都发生转置,整个输入矩阵都转置的目标
odata[index_out] = block[threadIdx.x][threadIdx.y];
}
}
算法的核心注释如上,短短几行代码但内容颇为丰富,并且还精妙的囊括了对Memory BANK Conflict和Coalescing的优化提升性能
文字总是苍白的,举个实例:假设输入矩阵A为6x4,ThreadBlock为2x2,则Grid为3x2,下面的算法示意视频展示了一个ThreadBlock执行算法以及所有ThreadBlock执行算法的详细过程
Matrix-Transpose算法示意