CUDA全局内存的合并访问

Question

我已经阅读了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的倍数，它就不是合并的吗？地址不对齐？这正确吗？

- Hlavson

---

除了网格中的第一个块，这些块都无法合并。线程按列主序编号。 - talonmies

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

合并、二维网格、一维块：

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

- harrism

---

增加了更多的严谨性和示例。 - harrism

Answer 2

你的索引从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

- djmj

Answer 3

哪些访问可以合并的规则有些复杂，而且随着时间的推移它们也在不断变化。每个新的CUDA架构在可合并性方面都更加灵活。我建议一开始不要过于担心这个问题。相反，以最方便的方式进行内存访问，然后查看CUDA分析器的结果。

- Roger Dahl

Answer 4

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

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

对于#2，有点模糊。如果someIndex为1，则它不会合并warp中的所有32个请求，但它可能会进行部分合并。我相信Fermi设备将把warp中线程1-31的访问作为128B顺序存储器段的一部分合并（第一个4B没有用），而由于不对齐，Tesla架构设备将使其成为不合并的访问，但我不确定。

假设someIndex为8，特斯拉将具有32B对齐的地址，而费米可能将它们分组为32B、64B和32B。但归根结底，取决于someIndex的值和架构，发生的情况是模糊的，并不一定会很糟糕。

- Vanwaril

---

无法确定，因为他的索引方式错误或非常奇怪，请参考我的答案。 - djmj

嗯，你说得对，发现得好。@Hlavson，根据你的问题，我假设你有一个一维网格和一维线程几何结构。因此，你需要使用[blockIdx.x*blockDim.x + threadIdx.x]进行索引。 - Vanwaril

很抱歉，这个答案完全是错误的。线程编号在块内是按列主序排列的，每个线程都乘以一个步长（blockIdx.x）。在第一种情况下，第一个块将完全协同，但之后的块则不会。第二种情况与第一种情况相同，只是有一个偏移量。 - talonmies

抱歉，这不是。对于情况＃1，如果您有一个1D块，则第一个块具有1个字的读取跨度，这将被合并。第二个块的读取跨度为2，无法合并，第三个块的跨度为3，依此类推。在具有1D块的情况下，情况＃1的等效公式为threadIdx.x * blockIdx.x + gridDim.x。那永远不会完全合并。情况＃2只是具有额外偏移量的情况＃1。 - talonmies

对不起，我不知道你在说什么。在任何块中，两个线程之间唯一的区别是threadIdx.x的差异；因此，在一个warp内，如果它开始对齐，它就会合并，如果它没有对齐，它就会出现奇怪的情况。我同意他问题中的索引是错误的 - 我在我的评论中解决了这个问题。但这并不是忽略手头的问题的理由，这个问题是关于何时内存访问合并以及何时不合并的。 - Vanwaril