CUDA并行归约reduce核函数调用：为何grid尺寸设置为grid.x/2？

你的核函数本身采用了单block处理2倍blockSize长度输入的设计，因此启动的grid数量需要减半，具体逻辑如下：

• 核函数的设计隐含了单block处理双倍数据的规则：
  1. 你在核函数中定义全局索引的逻辑是 unsigned int idx = blockDim.x*blockIdx.x * 2 + threadIdx.x，同时每个block对应的输入数据偏移是 int *idata = input + blockIdx.x*blockDim.x * 2，这两个参数的步进都是2*blockDim.x，天然就要求每个block负责两段各blockDim.x长度的输入数据
  2. 核函数在归约循环前还做了一次预合并：if (idx + blockDim.x<N) input[idx] += input[idx + blockDim.x];，直接把每个block负责的后半段数据加到前半段对应位置，后续归约只需要在单段blockDim.x长度内完成即可
• 启动时传grid.x/2的必要性：
  你一开始定义的dim3 grid(N / blocksize + 1);是按照每个block仅处理1倍blockSize长度数据计算出来的所需总block数，现在每个block可以处理2倍长度的数据，只需要原来一半的block就可以覆盖全部N个输入元素，所以启动时传入grid.x/2才是匹配核函数设计的参数
• 直接传grid.x出错的原因：
  如果传入完整的grid.x启动核函数，后半部分的block的起始偏移会超出输入数组的有效范围，出现非法内存访问，同时前半部分数据会被重复计算合并，最终归约结果自然错误。

内容的提问来源于stack exchange，提问作者Qiang Zhang