CUDA 联合访问全局内存

我已经阅读了 CUDA 编程指南，但我错过了一件事。假设我在全局内存中有 32 位 int 数组，我想通过合并访问将其复制到共享内存。 全局数组的索引从 0 到 1024，假设我有 4 个块，每个块有 256 个线程。

__shared__ int sData[256];

何时执行合并访问？

1.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y];

全局内存中的地址从 0 复制到 255，每个地址由 warp 中的 32 个线程复制，所以这里可以吗？

2.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y + someIndex];

如果 someIndex 不是 32 的倍数，那么它不会合并？地址错位？那是对的吗？

您想要什么最终取决于您的输入数据是一维还是二维数组，以及您的网格和块是一维还是二维。最简单的情况都是一维的：

shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + threadIdx.x];

这是合并的。我使用的经验法则是，将变化最快的坐标（threadIdx）添加为块偏移量（blockDim * blockIdx）的偏移量。最终结果是块中线程之间的索引步长为 1。如果步长变大，则会失去合并。

简单的规则（在 Fermi 和更高版本的 GPU 上）是，如果 warp 中所有线程的地址落入相同的对齐 128 字节范围，则将产生单个内存事务（假设为负载启用了缓存，这是默认）。如果它们落入两个对齐的 128 字节范围，则会产生两个内存事务，依此类推。

在 GT2xx 和更早的 GPU 上，情况变得更加复杂。但您可以在编程指南中找到详细信息。

其他示例：

未合并：

shmem[threadIdx.x] = gmem[blockDim.x + blockIdx.x * threadIdx.x];

没有合并，但在 GT200 及更高版本上还不错：

stride = 2;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

根本没有合并：

stride = 32;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

合并、2D 网格、1D 块：

int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.x] = gmem[blockIdx.y * elementPitch + 
                          blockIdx.x * blockDim.x + threadIdx.x]; 

合并的二维网格和块：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.y * blockDim.x + threadIdx.x] = gmem[y * elementPitch + x];