避免bank conflict的使用共享内存的矩阵转置。及其循环展开。转置前后都是行主序，中间共享内存中是列主序

在GPU并行编程中，共享内存被划分为多个bank，每个bank独立访问。当多个线程同时访问同一个bank的不同字节时，就会产生bank conflict，这会导致访问延迟和性能下降。那么，我们为什么不直接减少使用的线程数目，使之少于共享内存中每行数据的数量，反而要在每行中添加空数据这么麻烦呢？

这是因为：如果直接让线程数少于共享内存中每行的数据数量，虽然也能避免bank conflict，但会导致共享内存的利用率变低，影响性能。因此，为了充分利用共享内存的性能优势，通常会在共享内存中填充空数据来避免bank conflict。

为了避免bank conflict，可以使用padding的方式将数组中的每一行各自占据一个bank，这样每个线程读取不同行的数据时就不会发生bank conflict。padding的方式是在每行数据后面添加一定数量的空数据，使得每行的字节数恰好是bank的整数倍，这样就能保证每行数据各自占据一个bank。

具体来说，当一个线程访问共享内存中的一个元素时，GPU会根据该元素在共享内存中的地址计算出其所在的bank，然后将该元素从该bank中读取出来。如果该元素所在的bank同时被其他线程访问，就会产生bank conflict。

通过填充空数据，可以使得每个线程访问的数据都位于不同的bank中，从而避免bank conflict，提高访问效率和程序性能。

向原宽度数据填充一列空数据的过程只是对声明的共享内存宽度多加了一列。这个加了一列的宽度在映射到硬件上是不改变实际宽度的，如下图所示：

矩形共享内存的静态填充，填充前：

__global__ void setRowReadColRect(int * out)
{
    __shared__ int tile[BDIMY_RECT][BDIMX_RECT];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[icol][irow];
}

填充后：

__global__ void setRowReadColRectPad(int * out)
{
    __shared__ int tile[BDIMY_RECT][BDIMX_RECT+IPAD];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[icol][irow];
}

这里有个要注意的地方：输出矩阵的out[idx]依然是行主序，就像输入的矩阵（输入和输出都是来自全局内存，本函数是从全局内存传到共享内存中再传回全局内存）也是行主序[idx]，只是中间的tile[icol][irow]是列主序。

这个理解是从谭升的这个图里看到的： 

更详细的如下，填充前：

__global__ void transformSmem(float * in,float* out,int nx,int ny)
{
	__shared__ float tile[BDIMY][BDIMX];
	unsigned int ix,iy,transform_in_idx,transform_out_idx;
// 1.计算当前块中的线程的全局坐标（相对于整个网格），计算对应的一维线性内存的位置
	ix=threadIdx.x+blockDim.x*blockIdx.x;
    iy=threadIdx.y+blockDim.y*blockIdx.y;
	transform_in_idx=iy*nx+ix;
// 2.bidx表示block idx也就是在这个（block块？）块中的线程的坐标的线性位置（把块中的二维线程位置按照逐行排布的原则，转换成一维的），然后进行转置，也就是改成逐列排布的方式，计算出新的二维坐标，逐行到逐列排布的映射就是转置的映射，这只完成了很多块中的一块，而关键的是我们把这块放回到哪
	unsigned int bidx,irow,icol;
	bidx=threadIdx.y*blockDim.x+threadIdx.x;
	irow=bidx/blockDim.y;
	icol=bidx%blockDim.y;
// 3.计算出转置后的二维全局线程的目标坐标，注意这里的转置前的行位置是计算出来的是转置后的列的位置，这就是转置的第二步。
	ix=blockIdx.y*blockDim.y+icol;
	iy=blockIdx.x*blockDim.x+irow;
// 4.计算出转置后的二维坐标对应的全局内存的一维位置
	transform_out_idx=iy*ny+ix;
//5.读取全局内存，写入共享内存，然后按照转置后的位置写入
	if(ix<nx&& iy<ny)
	{
		tile[threadIdx.y][threadIdx.x]=in[transform_in_idx];
		__syncthreads();
		out[transform_out_idx]=tile[icol][irow];

	}
}

填充后：

__global__ void transformSmemPad(float * in,float* out,int nx,int ny)
{
	__shared__ float tile[BDIMY][BDIMX+IPAD];
	unsigned int ix,iy,transform_in_idx,transform_out_idx;
	ix=threadIdx.x+blockDim.x*blockIdx.x;
    iy=threadIdx.y+blockDim.y*blockIdx.y;
	transform_in_idx=iy*nx+ix;

	unsigned int bidx,irow,icol;
	bidx=threadIdx.y*blockDim.x+threadIdx.x;
	irow=bidx/blockDim.y;
	icol=bidx%blockDim.y;


	ix=blockIdx.y*blockDim.y+icol;
	iy=blockIdx.x*blockDim.x+irow;


	transform_out_idx=iy*ny+ix;

	if(ix<nx&& iy<ny)
	{
		tile[threadIdx.y][threadIdx.x]=in[transform_in_idx];
		__syncthreads();
		out[transform_out_idx]=tile[icol][irow];

	}

}

使用循环展开来进行进一步加速：

__global__ void transformSmemUnrollPad(float * in,float* out,int nx,int ny)
{
	__shared__ float tile[BDIMY*(BDIMX*2+IPAD)];
//1.
	unsigned int ix,iy,transform_in_idx,transform_out_idx;
	ix=threadIdx.x+blockDim.x*blockIdx.x*2;
    iy=threadIdx.y+blockDim.y*blockIdx.y;
	transform_in_idx=iy*nx+ix;
//2.
	unsigned int bidx,irow,icol;
	bidx=threadIdx.y*blockDim.x+threadIdx.x;
	irow=bidx/blockDim.y;
	icol=bidx%blockDim.y;
//3.
	unsigned int ix2=blockIdx.y*blockDim.y+icol;
	unsigned int iy2=blockIdx.x*blockDim.x*2+irow;
//4.
	transform_out_idx=iy2*ny+ix2;
	if(ix+blockDim.x<nx&& iy<ny)
	{
		unsigned int row_idx=threadIdx.y*(blockDim.x*2+IPAD)+threadIdx.x;
		tile[row_idx]=in[transform_in_idx];
		tile[row_idx+BDIMX]=in[transform_in_idx+BDIMX];
//5
		__syncthreads();
		unsigned int col_idx=icol*(blockDim.x*2+IPAD)+irow;
        out[transform_out_idx]=tile[col_idx];
		out[transform_out_idx+ny*BDIMX]=tile[col_idx+BDIMX];

	}
}

1. 计算当前块中的线程的全局坐标——一维线性内存的位置
2. 在共享内存内进行转置
3. 计算出转置后的二维全局线程的目标坐标，注意这里的转置前的行位置是计算出来的是转置后的列的位置，这就是转置的第二步。
4. 计算出转置后的二维坐标对应的全局内存的一维位置，注意这里不是一次计算一个块，而是计算两个块，换个理解方法，我们把原来的块x方向扩大一倍，然后再对这个大块分成两个小块（，B），每个小块中的对应位置差BDIMX，然后对其中A，B中数据按行写入共享内存。
5. 将4中读取到共享内存中的数据，按照转置后的位置写入全局内存

矩形共享内存动态版：

填充前：

__global__ void setRowReadColRectDyn(int * out)
{
    extern __shared__ int tile[];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    unsigned int col_idx=icol*blockDim.x+irow;
    tile[idx]=idx;
    __syncthreads();
    out[idx]=tile[col_idx];
}

填充后：

__global__ void setRowReadColRectDynPad(int * out)
{
    extern __shared__ int tile[];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    unsigned int row_idx=threadIdx.y*(IPAD+blockDim.x)+threadIdx.x;
    unsigned int col_idx=icol*(IPAD+blockDim.x)+irow;
    tile[row_idx]=idx;
    __syncthreads();
    out[idx]=tile[col_idx];
}

方形共享内存：

静态共享内存版，填充前：

//存储操作是无冲突的，但是加载操作显示有16路冲突
__global__ void setRowReadCol(int *out)
{
   __shared__ int tile[BDIMY]BDIMX];
   unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
   tile[threadIdx.y][threadIdx.x] = idx;
   __synthreads();
   out[idx] = tile[threadIdx.x][threadIdx.y];
}

填充后：

__global__ void setRowReadColIpad(int * out)
{
    __shared__ int tile[BDIMY][BDIMX+IPAD];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.x][threadIdx.y];
}

所以填充内存不用改其他地方的代码，只要改一下对共享内存宽度的声明就可以了。

动态共享内存版，填充前：

//写操作无冲突，读操作会报告一个16路径冲突
__global__ void setRowReadColDyn(int *out)
{
   extern __shared__ int tile[];
   unsigned int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
   unsigned int col_idx = threadIdx.x * blockDim.y + threadIdx.y;
   tile[row_idx] = row_idx;
   __synthreads();
   out[row_idx] = tile[col_idx];
}

填充后：

__global__ void setRowReadColDynPad(int *out)
{
   extern __shared__ int tile[];
   unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;
   unsigned int col_idx = threadIdx.x * (blockDim.x + IPAD) + threadIdx.y;
   unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;
   tile[row_idx] = g_idx;
   __synthreads();
   out[g_idx] = tile[col_idx];
}