我已阅读 CUDA 编程指南,但我错过了一件事。假设我在全局内存中有 32 位 int 数组,我想通过合并访问将它复制到共享内存。 全局数组的索引从 0 到 1024,假设我有 4 个 block ,每个 block 有 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 的倍数,它不会合并?地址错位?对吗?
最佳答案
您最终想要什么取决于您的输入数据是一维数组还是二维数组,以及您的网格和 block 是一维还是二维。最简单的情况都是一维的:
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + threadIdx.x];
这是合并的。我使用的经验法则是将变化最快的坐标(threadIdx)作为偏移量添加到 block 偏移量(blockDim * blockIdx)上。最终结果是 block 中线程之间的索引步长为 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 block :
int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.x] = gmem[blockIdx.y * elementPitch +
blockIdx.x * blockDim.x + threadIdx.x];
合并的二维网格和 block :
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];
关于memory - CUDA 合并访问全局内存,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/10325244/