An Efficient Matrix Transpose in CUDA C/C++中`Coalesced Transpose Via Shared Memory`

一节讲述如何使用`shared memory`

高效地实现`matrix transpose`

：

```
__global__ void transposeCoalesced(float *odata, const float *idata)
{
__shared__ float tile[TILE_DIM][TILE_DIM];
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
__syncthreads();
x = blockIdx.y * TILE_DIM + threadIdx.x; // transpose block offset
y = blockIdx.x * TILE_DIM + threadIdx.y;
for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
}
```

（1）`idata`

和`odata`

分别是表示`1024X1024`

个`float`

元素的`matrix`

的连续内存：

（2）关于`blockIdx`

和`threadIdx`

的取值，参考下面的图：

`shared memory`

请参考下面的图：

（3）在下列代码中：

```
for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
```

每一个`block`

是`32X8`

大小，需要循环`4`

次，把一个`block`

内容`copy`

到`tile`

这个`shared memory`

中。`idata`

是按行读取的，因此是`coalesced`

。

（4）最难理解的在最后一部分：

```
x = blockIdx.y * TILE_DIM + threadIdx.x; // transpose block offset
y = blockIdx.x * TILE_DIM + threadIdx.y;
for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
```

对比从`idata`

读取数据和写数据到`odata`

：

```
......
tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
......
odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
......
```

可以看到是把`tile`

做了`transpose`

的数据（行变列，列变行）传给`odata`

。而确定需要把`tile`

放到哪里位置的代码：

```
x = blockIdx.y * TILE_DIM + threadIdx.x; // transpose block offset
y = blockIdx.x * TILE_DIM + threadIdx.y;
```

假设`blockIdx.x`

为`31`

，`blockIdx.y`

为`0`

，`threadIdx.x`

为`1`

，`threadIdx.y`

为`2`

。根据上述代码，计算`x`

和`y`

：

```
x = 0 * 32 + 1;
y = 31 * 32 + 2;
```

根据下面的图，可以看到是把东北角的内容`copy`

的西南角：