CUDA 合并访问全局内存答案

【问题标题】：CUDA coalesced access to global memoryCUDA 合并访问全局内存
【发布时间】：2012-05-06 16:52:32
【问题描述】：

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

__shared__ int sData[256];

何时执行合并访问？

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

全局内存中的地址从 0 复制到 255，每个被 32 个线程在 warp 中复制，这样就可以了？

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

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

【问题讨论】：

这些都不能合并，除了网格中的第一个块。线程按列主要顺序编号。

标签： memory cuda copy coalescing

【解决方案1】：

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

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

【讨论】：

添加了更多的严谨和示例。

【解决方案2】：

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

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

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

将someIndex 设为 8，Tesla 将拥有 32B 个对齐的地址，Fermi 可能会将它们分组为 32B、64B 和 32B。但归根结底，根据someIndex 的值和架构，发生的事情是模糊的，不一定很糟糕。

【讨论】：

不能说，因为他的索引是错误的或者很奇怪，看我的回答
嗯，你说得对，不错。 @Hlavson，根据你的问题，我假设你有一个一维网格和一维线程几何。因此，您需要使用 [blockIdx.x*blockDim.x + threadIdx.x] 进行索引。
恐怕这个答案是完全错误的。线程编号是块内的主要列，并且都具有 threadIdx.x 乘以步幅 (blockIdx.x)。在第一种情况下，第一个块将发生完全合并，但在那之后不会发生。第二种情况与第一种情况相同，但有一个偏移量。
抱歉不是。对于案例 #1。如果你有一个 1D 块，那么 first 块的读取步长为 1 个单词，这将被合并。第二个块的读取步长为 2，不会合并，第三个块的步长为 3，依此类推。具有一维块的案例#1 的等效公式是threadIdx.x * blockIdx.x+gridDim.x。那永远不会完全融合。案例#2 只是案例#1 有一个额外的偏移量。
对不起，我不知道你在说什么。在 any 块中，两个线程之间的唯一区别是 threadIdx.x 的区别；所以在一个扭曲中，如果它开始对齐，它就会合并，如果没有，它会做一些奇怪的事情。我同意他问题中的索引是错误的——我在评论中提到了这一点。但这没有理由忽略手头的问题，即内存访问何时合并，何时不合并。

【解决方案3】：

您在 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

【讨论】：

【解决方案4】：

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

【讨论】：