Harris减缩优化阶段#4中的网格大小 [英] Grid size in phase #4 of Harris' reduction optimization

问题描述

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

这是书Professional CUDA C Programming中的代码snippet：

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个数据块，并处理每个数据挡路中的单个元素。 所以当我们启动内核时，与没有展开的内核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();

我的问题是在启动内核时如何确定网格的大小？是否应将其grid.x/2与没有多重加载的配置进行比较？

推荐答案

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

将块数减半，替换单一负载：

[代码片段]

两次加载并[第一次]添加减法

当然，您需要注意尺寸。为简单起见，演示文稿假定您的总长度是2的幂，因此在剩下多个元素的情况下，您始终可以安全地被2整除。在现实生活中并非如此，因此您可能需要考虑到网格大小的松弛(例如，网格大小的一半加上1(如果是奇数))。

这篇关于Harris减缩优化阶段#4中的网格大小的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！