8

我已经阅读了 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，每个被 32 个线程在 warp 中复制，所以可以吗？

2.

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

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

4

4 回答 4

16

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

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

合并的 2D 网格和块：

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];
于 2012-04-26T05:29:46.480 回答
1

您在 1 处的索引是错误的（或者故意如此奇怪，看起来是错误的），某些块在每个线程中访问相同的元素，因此无法在这些块中进行合并访问。

证明：

例子：

Grid = dim(2,2,0)

t(blockIdx.x, blockIdx.y)

//complete block reads at 0
t(0,0) -> sData[threadIdx.x] = gData[0];
//complete block reads at 2
t(0,1) -> sData[threadIdx.x] = gData[2];
//definetly coalesced
t(1,0) -> sData[threadIdx.x] = gData[threadIdx.x];
//not coalesced since 2 is no multiple of a half of the warp size = 16
t(1,1) -> sData[threadIdx.x] = gData[threadIdx.x + 2];

所以如果一个块被合并，它就是一个“运气”游戏，所以一般来说没有

但是合并内存读取规则对较新的 cuda 版本没有以前那么严格。
但是对于兼容性问题，如果可能的话，您应该尝试针对最低 cuda 版本优化内核。

这是一些不错的来源：

http://mc.stanford.edu/cgi-bin/images/0/0a/M02_4.pdf

于 2012-04-26T03:17:53.340 回答
0

可以合并访问的规则有些复杂，并且随着时间的推移而变化。每个新的 CUDA 架构在合并方面都更加灵活。我会说一开始不要担心。相反，以最方便的方式进行内存访问，然后查看 CUDA 分析器所说的内容。

于 2012-04-25T23:53:57.973 回答
-1

如果您打算使用一维网格和螺纹几何，您的示例是正确的。我认为您打算使用的索引是[blockIdx.x*blockDim.x + threadIdx.x].

我相信，对于 #1，warp 中的 32 个线程“同时”执行该指令，因此它们的请求是顺序的并与 128B (32 x 4) 对齐，在特斯拉和费米架构中合并。

对于＃2，它有点模糊。如果someIndex是 1，那么它不会在一个扭曲中合并所有 32 个请求，但它可能会进行部分合并。我相信 Fermi 设备会将线程 1-31 的访问合并为一个 128B 连续内存段的一部分（并且浪费了前 4B，不需要线程）。我认为特斯拉架构设备会由于未对齐而使其成为未合并的访问，但我不确定。

someIndex比如说 8，特斯拉将有 32B 个对齐的地址，费米可能会将它们分组为 32B、64B 和 32B 。但底线是，根据价值someIndex和架构，发生的事情是模糊的，不一定是可怕的。

于 2012-04-26T03:00:35.087 回答