CUDA学习——Chapter 3（7）归约问题（2）相邻归约及交错归约

第三章

并行归约的分化及相邻归约的优化

在上一篇的并行归约里面，我们要注意这样一条语句：

if ((tid % (2*stride))==0)

我们之前曾经讲过线程束分化的问题——当一个线程束里面并不是全部线程都是活跃的时候，我们称这个线程束是分化的。而在上述并行归约中，在第一次计算的时候只有1/2（甚至更少）的线程是活跃的，第二次只有1/4，第三次只有1/8，以此类推。

要改变这种情况，就必须让tid小的线程先被利用，这样可以让一些线程束是完全空闲的而不是分化的，并且提高SM中的占用率。

其思路如图3-23所示：

还是老样子，我们先把改进后的设备端代码贴出来，再一点点分析这个代码：

__global__ void reduceNeighoredLess(int *g_idata,int *g_odata,unsigned int n)
{
   
	unsigned int tid=threadIdx.x;
	unsigned int idx=blockIdx.x*blockDim.x+threadIdx.x;