Harris 缩减优化第 4 阶段的网格大小

Question

我正在学习展开循环以优化内核计算。

这是一个代码snippet from the book Professional CUDA C Programming:

if (idx + 4 * blockDim.x <= n)
{
    int a1 = g_idata[idx];
    int a2 = g_idata[idx + blockDim.x];
    int a3 = g_idata[idx + 2 * blockDim.x];
    int a4 = g_idata[idx + 3 * blockDim.x];
    tmpSum = a1 + a2 + a3 + a4;
}

据我了解，每个线程处理 4 个数据块并处理每个数据块中的单个元素。所以，当我们启动内核时，相对于内核w/o展开grid.x，配置改为 reduceSmemUnroll<<<grid.x / 4, block>>>.

然后我有一个关于 Mark Harris presentation on parallel reduction 第 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();

我的问题是启动内核时如何确定网格的大小？相比配置w/o多负载应该是grid.x/2？

Answer 1

是的，应该是块数的一半；它在幻灯片上这样说，第一次出现您在 Mark 的演示文稿中引用的代码片段 - 已经在幻灯片 18 上：

Halve the number of blocks, and replace single load:

[code snippet]

with two loads and [the] first add of the reduction

当然，您需要注意尺寸。为简单起见，该演示文稿假设您的总长度是 2 的幂，因此当还有多个元素时，您始终可以安全地除以 2。在现实生活中情况并非如此，因此您可能需要留出余量（例如“网格大小的一半加一，如果是奇数”）。

Harris 缩减优化第 4 阶段的网格大小

Grid size in phase #4 of Harris' reduction optimization

parallel-processing

cuda

reduction