cuda规约运算

归:递归
约:减小

对于一个矩阵做求和运算

若串行求和的话需要o(n)的复杂度
但若向下图这样，俩俩并行相加，只需要o(logn)的复杂度

__global__ void reduceNeighbored(int* g_idata, int* g_odata, unsigned int n)  //g_idata是待求和数组，g_odata存放的是每个线程块求和的结果
{
	unsigned int tid = threadIdx.x;

	if (tid >= n) return;

	int* idata = g_idata + blockIdx.x * blockDim.x;   //对每一个线程块进行求和

	for (int stride = 1; stride < blockDim.x; stride <<=1) //stride是步长，代表相加的俩个元素的跨度（坐标0和坐标1的跨度为1）。
	{
		if ((tid % (2 * stride)) == 0) //第一次的tid可以取0,2,4.... 第二次的tid可以取0,4,8... 步长如果不乘2会造成多余运算
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();    //同步一个线程块中的所有线程。等待直到同一个线程块中的所有其他线程都到达这一点
	}

	if (tid == 0) g_odata[blockIdx.x] = idata[0]; //g_odata[i]存的是线程块i的和，最后所有线程的和都存到g_odata[0]
}

还有一种更高效的规约是跨一半去相加
如下图所示

__global__ void reduceNeighbored(int* g_idata, int* g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;

	if (tid >= n) return;

	int* idata = g_idata + blockIdx.x * blockDim.x;   //对每一个线程块进行求和

	for (int stride = blockDim.x/2; stride > 0; stride >>=1)
	{
		if (tid < stride) 
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();    //同步一个线程块中的所有线程。等待直到同一个线程块中的所有其他线程都到达这一点
	}

	if (tid == 0) g_odata[blockIdx.x] = idata[0]; //g_odata[i]存的是线程块i的和，最后所有线程的和都存到g_odata[0]
}