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。