GPU/CUDA: 重新排序设备内存

Question

我在设备内存中存储了一个多维数组。我想"permute"/"transpose"，即根据新的维度顺序重新排列它的元素。

例如，如果我有一个二维数组

A = [0, 1, 2
     3, 4, 5]

我想改变维度的顺序所以我得到

B = [0, 3
     1, 4
     2, 5]

这种重新排序实际上是按照 [0,1,2,3,4,5] 和 return 的顺序复制存储在内存中的元素，一个新的排序 [0,3,1,4,2,5].

我知道如何将索引从 A 映射到 B，我的问题是如何使用 cuda 在设备上高效地执行此映射？

Answer 1

你可以检查这个http://devblogs.nvidia.com/parallelforall/efficient-matrix-transpose-cuda-cc/

朴素矩阵转置：

__global__ void transposeNaive(float *odata, const float *idata)
{
  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)
    odata[x*width + (y+j)] = idata[(y+j)*width + x];
}

通过共享内存合并转置：

__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];
}

GPU/CUDA: 重新排序设备内存

GPU/CUDA: Re-ordering device memory

c++

arrays

cuda

gpu