CUDA：使用网格跨步循环减少共享内存答案

【问题标题】：CUDA: Using grid-strided loop with reduction in shared memoryCUDA：使用网格跨步循环减少共享内存
【发布时间】：2014-10-01 11:36:05
【问题描述】：

我有以下关于在 CUDA 内核中一起使用共享内存中的网格跨步循环和优化缩减算法的问题。想象一下，您有一个一维数组，其元素数量多于网格中的线程数 (BLOCK_SIZE * GRID_SIZE)。在这种情况下，您将编写这种内核：

#define BLOCK_SIZE (8)
#define GRID_SIZE (8)
#define N (2000)

// ...

__global__ void gridStridedLoop_kernel(double *global_1D_array)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int i;

    // N is a total number of elements in the global_1D_array array
    for (i = idx; i < N; i += blockDim.x * gridDim.x)
    {
        // Do smth...
    }
}

现在你想通过减少共享内存来寻找global_1D_array中的最大元素，上面的内核将如下所示：

#define BLOCK_SIZE (8)
#define GRID_SIZE (8)
#define N (2000)

// ...

__global__ void gridStridedLoop_kernel(double *global_1D_array)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int i;

    // Initialize shared memory array for the each block
    __shared__ double data[BLOCK_SIZE];

    // N is a total number of elements in the global_1D_array array
    for (i = idx; i < N; i += blockDim.x * gridDim.x)
    {
        // Load data from global to shared memory
        data[threadIdx.x] = global_1D_array[i];
        __syncthreads();

        // Do reduction in shared memory ...
    }

    // Copy MAX value for each block into global memory
}

很明显data 中的某些值将被覆盖，即您需要更长的共享内存数组或必须以另一种方式组织内核。一起使用减少共享内存和跨步循环的最佳（最有效）方法是什么？

提前致谢。

【问题讨论】：

你研究过cuda parallel reduction sample and PDF吗？它涵盖了所有这些，最终展示了使用网格跨步循环并行减少共享内存。
是的，我看到了这个演示文稿，目前使用其中讨论的算法之一进行归约，但本文中的步幅意味着归约级别的步幅，并且隐含地认为您在网格中有更多线程比输入全局数组中的元素。在这篇文章中，情况正好相反，我想使用网格步幅而不是步幅来降低水平。
你应该再读一遍。 “据认为，网格中的线程多于输入全局数组中的元素。”这不是真的。本文也涵盖了相反的情况（在本文末尾）。如果您的陈述是正确的，则不需要网格跨步循环。
例如，看一下the pdf 的幻灯片 32。那是一个网格跨步循环。网格的大小（在线程中）小于输入全局数组中的数据元素数量。

标签： c memory cuda shared

【解决方案1】：

记录了使用网格跨步循环的减少here。参考幻灯片 32，网格跨步循环如下所示：

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockSize*2) + threadIdx.x;
unsigned int gridSize = blockSize*2*gridDim.x;
sdata[tid] = 0;
while (i < n){
  sdata[tid] += g_idata[i] + g_idata[i+blockSize];
  i += gridSize;
  }
__syncthreads();

请注意，while 循环的每次迭代都会将索引增加gridSize，并且此 while 循环将继续进行，直到索引 (i) 超过（全局）数据大小 (n)。我们称之为网格跨步循环。在这个例子中，线程块局部缩减操作的其余部分不受网格大小循环的影响，因此只显示了“前端”。这种特殊的归约是在进行总和归约，但最大归约会简单地将操作替换为以下内容：

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockSize + threadIdx.x;
unsigned int gridSize = blockSize*gridDim.x;
sdata[tid] = 0;
while (i < n){
  sdata[tid] = (sdata[tid] < g_idata[i]) ? + g_idata[i]:sdata[tid];
  i += gridSize;
  }
__syncthreads();

线程块级别减少的其余部分必须以类似的方式进行修改，将求和操作替换为最大查找操作。

完整的并行缩减 CUDA 示例代码可作为任何完整 cuda 示例安装的一部分，或 here。

【讨论】：