CUDA实例系列二:矩阵转置优化

矩阵转置优化

矩阵转置是CUDA初学者们刚开始面临的一个比较绕的问题, 他的"绕"在于其中坐标的转换, 以及使用Shared Memory来解决合并访存的过程.

在很多实际项目或应用中, 矩阵转置多数都是其中一个步骤.接下来我就来解释下如何利用CUDA优化矩阵转置问题.

首先, 解释下矩阵转置矩阵转置

按照矩阵转置的公式，我们设定（按照下图所示）

输入矩阵为: A[16][16] M=16
输出矩阵为: B[16][16]
保证：A[y][x] = B[x][y]

接下来, 话不多说直接上代码(不要划走, 解释在代码之后):

__global__ void transpose(int A[M][M], int B[M][M])
{
    __shared__ int rafa[TILE_SIZE][TILE_SIZE + 1]; //tell me why?
	
    int x = threadIdx.x + blockDim.x * blockIdx.x;
    int y = threadIdx.y + blockDim.y * blockIdx.y;
    if (x < M && y < M)
    {
	rafa[threadIdx.y][threadIdx.x] = A[y][x];
    }
    __syncthreads();
	
    int y2 = threadIdx.y + blockDim.x * blockIdx.x;
    int x2 = threadIdx.x + blockDim.y * blockIdx.y;
    if (x2 < M && y2 < M)
    {
	B[y2][x2] = rafa[threadIdx.x][threadIdx.y];
    }
}

1. 先说明一下实验环境

---

---

2. 将数据从global Memory中读取到shared memory, 注意这里申请shared memory的时候TILE_SIZE+1是为了避免Shared memory的冲突

---

---

3. 将对应的结果写入global memory, 具体说明写在图片中

---

---

---

---

这样我们就完成了矩阵的转置, 最麻烦的部分是坐标问题, 大家一定要看仔细