cuda矩阵转置

一直以为经典的cuda矩阵转置只能用于矩阵的宽高都能被线程块大小整除的情况，也是很奇怪，不知道怎么形成这个概念的，然后这次又要用到，本想着大干一番，把宽高不能被线程块整除的矩阵转置攻克了，可是没想到一测试，人本来就可以实现，这就尴尬了, 所以在此记录下来，纠正自己的这个错误。

代码：

#define BLOCK_DIM 16

__global__ void myTranspose(float *indata,float *odata,int height,int width){__shared__ float block[BLOCK_DIM][BLOCK_DIM + 1];unsigned int xindex = blockIdx.x * BLOCK_DIM + threadIdx.x;unsigned int yindex = blockIdx.y * BLOCK_DIM + threadIdx.y;/*int height1 = height - height%BLOCK_DIM;int width1 = width - width%BLOCK_DIM;*/if(xindex < width && yindex < height){unsigned int index_in = yindex * width + xindex;block[threadIdx.y][threadIdx.x] = indata[index_in];}__syncthreads();xindex = blockIdx.y * BLOCK_DIM + threadIdx.x;yindex = blockIdx.x * BLOCK_DIM + threadIdx.y;if(xindex < height && yindex < width){unsigned int index_out = yindex * height + xindex;odata[index_out] = block[threadIdx.x][threadIdx.y];}}

再解释一下开辟的共享内存块大小为什么是[BLOCK_DIM][BLOCK_DIM+1]吧。

这是为了共享内存不产生bank conflict。

共享内存有32个bank，32位连续的内存映射到连续的bank上，一个bank一个时钟周期的带宽是32位。

这里我们的BLOCK_DIM是16，如果共享内存块的大小是[BLOCK_DIM][BLOCK_DIM],那么共享内存的bank如左下图：

我用绿色标出了bank1的冲突，可以看到同时是有这么多线程访问bank1的不同地址的，由于一个bank一个时钟周期的带宽只有32位，那么这些线程的操作就会变成串行的，顺序执行。右边的图表示共享内存块的大小是[BLOCK_DIM][BLOCK_DIM+1], 可以看到没有线程同时访问同一bank的不同地址，所以不会发生bank conflict。