CUDA 循环中的空间局部性

如何解决CUDA 循环中的空间局部性

我正在阅读《更简单的 CUDA 简介》，我在想这样的例子：

__global__
void add(int n,float *x,float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

其中每个线程跨越数组。在普通的 cpu 计算中，人们宁愿将数组拆分为连续的子数组，这些子数组在线程之间拆分，以便每个子数组都能更好地利用空间局部性。

这个概念是否也适用于 CUDA 的统一内存？我想了解在这种情况下最有效的方法是什么。

解决方法

grid-stride loop 有利于内存访问的原因是它提升了 "coalesced" access to global memory。简而言之，合并访问意味着 warp 中的相邻线程正在访问内存中的相邻位置，在任何给定的读取或写入周期/操作中，被认为是 Warp-wide。

网格步幅循环在整个扭曲中排列索引以促进这种模式。

这与内存是使用“普通”设备分配器（例如 cudaMalloc）还是“统一”分配器（例如 cudaMallocManaged）分配的正交。无论哪种情况，设备代码访问此类分配的最佳方式是使用合并访问。

您没有问过它，但是 CUDA shared memory 也有它的一种“最佳访问模式”，它由 warp 中的相邻线程访问（共享）内存中的相邻位置组成。