如何理解此 CUDA 矩阵复制代码中的合并访问？答案

【问题标题】：How to understand the coalesced access in this CUDA matrix copy code?如何理解此 CUDA 矩阵复制代码中的合并访问？
【发布时间】：2021-12-30 02:38:53
【问题描述】：

__global__ void Matrixcopy(float *odata, const float *idata)
{
  // threadblock size = (TILE_DIM, BLOCK_ROWS) = (32, 8)
  // each block copies a 32 * 32 tile
  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[(y+j)*width + x] = idata[(y+j)*width + x];
}

我对多维度数组的合并访问的概念感到很困惑。合并的全局内存访问的定义是

顺序内存访问是相邻的

根据文献Learn CUDA Programming。对于一维数组，很容易理解线程被索引为threadIdx.x + blockDim.x * blockIdx.x，它可以很容易地映射到真正的一维数组：warp 中的相邻线程访问一维数组的相邻物理地址。

然而对于 2-dim 数组或矩阵和 2-dim 线程块，如上面的代码，我不确定我是否理解正确：warp 中的相邻线程位于同一行，即相同的 y值，不同的 x 值。如果 y = 0，则连续线程为 x = [1, 2, 3, 4, 5, 6...]，并且它们访问连续地址 [1, 2, 3, 4, 5, 6...] 如果j = 0。所以这段代码是合并访问的。我理解正确吗？而这只是一个简单的cuda代码，如果我们有一个复杂的cuda内核，如何快速判断一个访问是否被合并？

【问题讨论】：

标签： c++ arrays memory cuda

【解决方案1】：

因此，此代码具有合并访问权限。我的理解正确吗？

是的，差不多。我会说线程是 x = [0, 1, 2, 3, 4, 5, 6...]，它们访问连续的地址 [0, 1, 2, 3, 4, 5, 6...] 但基本上我们是一致的。

如果我们有一个复杂的 cuda 内核，我们如何快速确定一个访问是否被合并？

您可以查看任何索引构造并使用以下测试：如果threadIdx.x 变量作为加法因子包含在索引中，并且没有乘法因子，则访问将在典型用法中合并（你有方形线程块的地方）。任何可以表示为的索引：

idx = f + threadIdx.x

其中f 是任意的，但不包括threadIdx.x 将导致合并访问。 “in x”中的相邻线程将访问内存中的相邻位置。对于“非方形”线程块，您可以使用threadIdx.y 开发类似的规则。例如，尺寸为 (1,32) 的线程块将要求将 threadIdx.y 作为仅加法因子包含在内。

【讨论】：

感谢您的明确答复。