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个线程进行warp,所以这里可以吗？

2.

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

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

Answer 1

您想要的最终取决于您的输入数据是1D还是2D数组,以及您的网格和块是1D还是2D.最简单的情况是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];

合并的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];

Answer 2

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