CUDA矩阵转置(共享内存 tile)

Udacity的CUDA编程课程中介绍了CUDA实现矩阵转置的六种方式,本文介绍其中的一种方式

如果矩阵为N*N的方阵。该方式让每个线程处理一个矩阵元素,总共需要N*N个线程。首先,声明两个常量并配置blocks,threads:

const int N=1024;
const int K=32;
dim3 blocks(N/K,N/K); 
dim3 threads(K,K);	

内核函数:

__global__ void 
transpose_parallel_per_element_tiled(float in[], float out[])
{
	// (i,j) locations of the tile corners for input & output matrices:
	int in_corner_i  = blockIdx.x * blockDim.x, in_corner_j  = blockIdx.y * blockDim.y;
	int out_corner_i = in_corner_j, out_corner_j = in_corner_i;

	int x = threadIdx.x, y = threadIdx.y;

	__shared__ float tile[K][K];

	// coalesced read from global mem, TRANSPOSED write into shared mem:
	tile[y][x] = in[(in_corner_i + x) + (in_corner_j + y)*N];
	__syncthreads();
	// read from shared mem, coalesced write to global mem:
	out[(out_corner_i + x) + (out_corner_j + y)*N] = tile[x][y];
}

内核涉及两个输入参数,in代表输入矩

评论 5
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值