CUDA：使用具有减少在共享存储器网格跨距环(CUDA: Using grid-strided lo

我在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将被覆盖，即你需要更长的共享存储阵列或有权组织内核以另一种方式。什么是最好的（最有效）的方式来使用共享内存减少和循环跨入一起？

提前致谢。

使用网格跨入循环的减少是记录在这里。参照滑动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 ）。我们称之为网跨进循环。在这个例子中，所述threadblock本地还原操作的其余部分不被网格大小循环的影响，因而仅“前端”被示出。这种特殊的减少是做算术还原，但最大还原会简单地用类似更换操作：

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();

而在threadblock水平下降的其余部分将在一个类似的方式进行修改，以最大调查操作更换求和操作。

全平行还原CUDA示例代码可作为任何全CUDA样品安装的一部分，或在这里。