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 中的相邻线程访问(共享)内存中的相邻位置组成。