CUDA：使用网格划分循环减少共享内存

关于在CUDA内核中共享内存中网格跨越循环和优化缩减算法的使用，我有以下问题。 想象一下，你有1D数组，其元素数多于网格中的线程数（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中的某些值将被覆盖，即您需要更长的共享内存arrays或者必须以另一种方式组织内核。 将共享内存和跨步循环一起使用减少的最佳（最有效）方法是什么？

提前致谢。

Answer 1

这里记录了使用网格跨越环的减少。 参考幻灯片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样本安装的一部分提供，或在此处提供 。