我无法找到在 CUDA C 中使用共享内存转置非方矩阵的方法。(我是 CUDA C 和 C 的新手)
In 这篇博文 https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/展示了如何转置矩阵的有效方法(通过共享内存合并转置)。但它只适用于方阵。
还提供了代码github https://github.com/NVIDIA-developer-blog/code-samples/blob/master/series/cuda-cpp/transpose/transpose.cu(与博客上的相同)。
StackOverflow 上有一个类似的question https://stackoverflow.com/questions/13208982/non-square-matrix-transpose-with-shared-mem-in-cuda. There TILE_DIM = 16
已设置。但通过该实现,每个线程只需将矩阵的一个元素复制到结果矩阵。
这是我当前的实现:
__global__ void transpose(double* matIn, double* matTran, int n, int m){
__shared__ double tile[TILE_DIM][TILE_DIM];
int i_n = blockIdx.x*TILE_DIM + threadIdx.x;
int i_m = blockIdx.y*TILE_DIM + threadIdx.y; // <- threadIdx.y only between 0 and 7
// Load matrix into tile
// Every Thread loads in this case 4 elements into tile.
int i;
for (i = 0; i < TILE_DIM; i += BLOCK_ROWS){
if(i_n < n && (i_m+i) < m){
tile[threadIdx.y+i][threadIdx.x] = matIn[n*(i_m+i) + i_n];
} else {
tile[threadIdx.y+i][threadIdx.x] = -1;
}
}
__syncthreads();
for (i = 0; i < TILE_DIM; i += BLOCK_ROWS){
if(tile[threadIdx.x][threadIdx.y+i] != -1){ // <- is there a better way?
if(true){ // <- what should be checked here?
matTran[n*(i_m+i) + i_n] = tile[threadIdx.x][threadIdx.y+i];
} else {
matTran[m*i_n + (i_m+i)] = tile[threadIdx.x][threadIdx.y+i];
}
}
}
}
其中 4 个元素从线程复制到图块中。此外,图块中的四个元素也被复制回结果矩阵中。
这里是内核配置<<<a, b>>>
:
where a: (ceil(n/TILE_DIM), ceil(n/TILE_DIM)) (-> is casted to doubles) and
b: (TILE_DIM, BLOCK_ROWS) (-> (32, 8))
我目前正在使用if(tile[threadIdx.x][threadIdx.y+i] != -1)
- 确定哪个线程应该复制到结果矩阵的语句(可能还有另一种方法)。就我目前的知识而言,其行为如下:在一个块中,线程索引(x, y)
将数据复制到图块和线程索引中(y, x)
将数据复制回结果矩阵。
我插入了另一个if
- 确定将数据复制到哪里的语句,因为有 2(?) 个可能的目的地,具体取决于线程索引。现在true
插入那里,但我尝试了很多不同的东西。我能想到的最好的办法是if(threadIdx.x+1 < threadIdx.y+i)
,它转置3x2
- 矩阵成功。
有人可以解释一下,我通过写回到结果矩阵中缺少什么吗?显然只有一个目的地是正确的。使用
matTran[n*(i_m+i) + i_n] = tile[threadIdx.x][threadIdx.y+i];
正如博客中提到的应该是正确的,但我不明白,为什么它不适用于非平方矩阵?