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，每一个由在经32个线程复制，因此在这里它确定？

2.

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

如果someIndex不是32多个它不聚结？错位的地址？那是对的吗？

Answer 1

访问可以合并的规则有点复杂，并且随着时间的推移而变化。每个新的CUDA架构在可以合并时都更加灵活。起初我会说不要担心。相反，以任何最方便的方式访问内存，然后查看CUDA分析器的说明。

Answer 2

如果您打算使用1D网格和线程几何，则您的示例正确。我认为你打算使用的索引是[blockIdx.x*blockDim.x + threadIdx.x]。我相信，在＃1中，warp中的32个线程同时执行该指令，因此他们的请求顺序排列并与128B（32 x 4）对齐，并且在Tesla和Fermi体系结构中都合并在一起。

＃2，它有点模糊。如果someIndex为1，则它不会合并所有32个请求中的一个warp，但它可能会进行部分合并。我相信费米设备将会像128B连续内存段（以及没有线程需要的第一个4B被浪费掉）的一部分一样，在线程中将线程1-31的访问合并在一起。我认为特斯拉架构设备会因错位导致不合并访问，但我不确定。

与someIndex一样，如8，特斯拉将有32B对准的地址，费米可能将它们分组为32B，64B和32B。但底线是，取决于someIndex的价值和架构，发生的事情是模糊的，并不一定是可怕的。

Answer 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

Answer 4

你想要什么，最终取决于你的输入数据是否是一维或二维阵列，以及是否网格和块是一维或二维。最简单的情况都是1D：

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

Coa lesced，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];