时间戳

首页 > 解决方案 > 在 Cuda C 上具有任意大小的矩阵转置（带共享内存）

c - 在 Cuda C 上具有任意大小的矩阵转置（带共享内存）

问题描述

我想不出一种在 CUDA C 中使用共享内存转置非平方矩阵的方法。（我是 CUDA C 和 C 的新手）

在网站上：

https://devblogs.nvidia.com/efficient-matrix-transpose-cuda-cc/

展示了如何转置矩阵（通过共享内存合并转置）的一种有效方法。但它只适用于平方矩阵。

github上还提供了代码（与博客上的相同）。

在 Stackoverflow 上有一个类似的问题。有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)-statement 来确定应该将哪个线程复制到结果矩阵（可能还有另一种方式）。就我目前的知识而言，其行为如下：在一个块中，ThreadIdx(x, y)将数据复制到图块中，而 ThreadIdx(y, x)将数据复制回结果矩阵。

我插入了另一个if语句来确定将数据复制到哪里，因为有 2（？）个可能的目的地，具体取决于 ThreadIdx。目前true插入那里，但我尝试了很多不同的东西。我能想到的最好的方法是if(threadIdx.x+1 < threadIdx.y+i)，它成功地转置了一个3x2-matrix 。

有人可以通过写回结果矩阵来解释我缺少什么吗？显然只有一个目的地是正确的。使用

matTran[n*(i_m+i) + i_n] = tile[threadIdx.x][threadIdx.y+i];

正如在博客上提到的应该是正确的，但我不知道，为什么它不适用于非平方矩阵？

标签： cmatrixcudatranspose

解决方案

我把问题复杂化了。在这里，索引并没有像我想象的那样交换。使用线程/块的 Y 坐标和 X 坐标重新计算它们。这是片段：

i_n = blockIdx.y * TILE_DIM + threadIdx.x;  
i_m = blockIdx.x * TILE_DIM + threadIdx.y

这是更正后的代码：

__global__ void transposeGPUcoalescing(double* matIn, int n, int m, double* matTran){
    __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[(i_m+i)*n + i_n];
        }
    }
    __syncthreads();

    i_n = blockIdx.y * TILE_DIM + threadIdx.x; 
    i_m = blockIdx.x * TILE_DIM + threadIdx.y;

    for (i = 0; i < TILE_DIM; i += BLOCK_ROWS){
        if(i_n < m  && (i_m+i) < n){
            matTran[(i_m+i)*m + i_n] = tile[threadIdx.x][threadIdx.y + i]; // <- multiply by m, non-squared!

        }
    }
}

感谢此评论注意到错误:)

推荐阅读