cuda中的Grid-Stride Loops (网格跨步循环) 详解

最近在学习cuda编程的时候遇到了不少问题，其中有一个问题很费解的就是为什么cuda中循环的步长是一个网格中容纳的线程的数量。代码如下所示：

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

其中blockDim.x和gridDim.x分别为block中的线程数量和grid中的block数量，这样以来，步长就是网格中的线程数量了。这样让我很费解，后来查阅了官方文档。地址 在官方文档中是这样解释的：

Notice that the stride of the loop is blockDim.x * gridDim.x which is the total number of threads in the grid. So if there are 1280 threads in the grid, thread 0 will compute elements 0, 1280, 2560, etc. This is why I call this a grid-stride loop. By using a loop with stride equal to the grid size, we ensure that all addressing within warps is unit-stride, so we get maximum memory coalescing, just as in the monolithic version.

也就是说，这样去循环的话，假设网格中有1280个线程，那么线程0将会去计算成员0,1290,2560,等等。这就是为什么称之为Grid-Stride Loops。通过使用这种方式的循环，我们能够确保warp中的所有寻址都是单位步长，因此我们获得最大内存合并，就像在单片版本中一样。关于最大内存合并的讲解在这里。最大内存合并
当使用足以覆盖循环的所有迭代的网格启动时，网格跨步循环应该具有与单片内核中的if语句基本相同的指令成本，因为循环增量将仅在循环条件为真时评估 。

grid-stride loop 的好处

1. 可扩展性和线程重用 Scalability and thread reuse
   通过使用循环你可以计算温和大小的问题，基础超过了cuda设备所支持的大小。可以通过限制block的数量来调整性能。例如经常启动设备上多处理器(multiprocessors)倍数的blocks去平衡使用率。例子如下所示：
   例子

int numSMs;
cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, devId);
// Perform SAXPY on 1M elements
saxpy<<<32*numSMs, 256>>>(1 << 20, 2.0, x, y);

当你限制grid中blocks的数量的时候，线程将会被多次计算重用。线程重用会将线程创建和销毁成本与内核在循环之前或之后可能执行的任何其他处理（例如线程专用或共享数据初始化）进行分摊。

2. 调试 Debugging
   通过使用循环而不是单片内核，您可以通过使用一个线程启动一个块来轻松切换到串行处理。

	 saxpy<<<1,1>>>(1<<20, 2.0, x, y);

这样可以更轻松地模拟串行主机实现以验证结果，并且可以通过序列化打印顺序使printf调试更容易。 序列化计算还允许您消除由运行到运行的操作顺序的变化引起的数值变化，帮助您在调整并行版本之前验证数字是否正确。

3. 可移植性和可读性 Portability and readability
   网格跨步循环代码更像是原始顺序循环代码而不是单片内核代码，使得其他用户更清楚。 事实上，我们可以很容易地编写一个内核版本，该内核可以编译并运行在GPU上作为并行CUDA内核或作为CPU上的顺序循环运行。 Hemi库提供了一个grid_stride_range（）帮助器，它使用C ++ 11基于范围的for循环来实现这一点。

HEMI_LAUNCHABLE
void saxpy(int n, float a, float *x, float *y)
{
 for (auto i : hemi::grid_stride_range(0, n)) {
   y[i] = a * x[i] + y[i];
 }
}

我们可以使用此代码启动内核，该代码在为CUDA编译时生成内核启动，或在为CPU编译时生成函数调用。

hemi::cudaLaunch(saxpy, 1<<20, 2.0, x, y);

总结

网格跨步循环是使CUDA内核灵活，可扩展，可调试甚至可移植的好方法。 虽然本文中的示例都使用了CUDA C / C ++，但同样的概念也适用于其他CUDA语言，例如CUDA Fortran。

引用

https://devblogs.nvidia.com/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/
https://devblogs.nvidia.com/even-easier-introduction-cuda/