CUDA 循环中的空间局部性答案

【问题标题】：Spatial locality in CUDA loopsCUDA 循环中的空间局部性
【发布时间】：2021-08-08 22:52:48
【问题描述】：

我正在阅读 Even Easier Introduction to 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 的统一内存？我想了解在这种情况下最有效的方法是什么。

【问题讨论】：

虽然 NVIDIA 称它们为“线程”，但这是用词不当。您应该将 CUDA 线程视为 SIMD 矢量处理单元的通道。那么这个模式就更有意义了。话虽如此，这个类比并不适用于所有层面，因此他们称之为 SIMT。
而且 CUDA 也支持每个线程中的像 float4 这样的 SIMD 类型，作为较小的并行化方法（这可以减少内存繁重内核的内存事务数量）。

标签： c++ performance cuda unified-memory

【解决方案1】：

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

grid-stride 循环将索引排列在 warp 上以促进这种模式。

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

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

【讨论】：

这个答案中一个非常重要的概念是 32 个线程的 warp，可以认为是同时执行。