CUDA学习——Chapter 3（7）归约问题（1）相邻归约及并行归约

最新推荐文章于 2024-06-03 10:18:50 发布

KarK_Li

最新推荐文章于 2024-06-03 10:18:50 发布

阅读量836

点赞数

分类专栏： CUDA 文章标签： CUDA C 并行计算

本文链接：https://blog.csdn.net/weixin_40427089/article/details/90301870

版权

本文详细介绍了CUDA中的归约问题，包括相邻归约和交错配对两种基本单位划分方法，以及如何实现并行归约算法。通过线程的同步和设备端、主机端的代码分析，阐述了如何在CUDA中进行高效的数据加总。

摘要由CSDN通过智能技术生成

第三章

归约问题

首先来介绍一个并行计算的案例：
有一个长度为n的数组L，求他们的和。
如果按顺序叠加，那么时间复杂度就是O(n)。
伪代码就是：

count←0
for i←1 to n
	count+=L[i]
return count

那还有一种在串行是O(n)，在并行上可以是 $\left \lceil log_2n \right \rceil$ 的算法，称为并行归约算法。这种算法是基于二分法的思想来做的，其基本思路就是，把输入的数组划分成更小的基本单位，直到计算每一个基本单位的和都只需要一个加法指令，然后自底向上的合并这些结果，称为归约。而并行计算的优势在于并行，因此最底层的基本单位的和的计算就可以使用并行来处理，同时，得到的基本单位的和组成的集合可以就地保存在输入向量中，不需要有额外的空间开销。

那么这些基本单位的划分可以分为两种：相邻配对和交错配对。

还是拿上面的L来说，相邻配对指的是相邻的两个元素组成一个基本单位，比如L[0]和L[1]，L[2]和L[3]，如果数组有奇数个元素，那么最后就剩一个单独的L[n-1]做基本单位。

交错配对指的是，以n/2为分界线，左右两边各自重新分序号，序号相同的组在一起。比如L[0]和L[4]，L[1]和L[5]（如果L的长度为8），直到L[3]和L[7]。

下列的示意图应该可以更好地帮助大家理解（3-19为相邻配对，3-20为交错配对）：

如何实现并行归约？

首先，思路就是，每次都是线程全部并行计算完之后，再开始下一步的线程并行计算，直至活跃的线程只有一个。这个思路既可以是循环的，也可以是递归的。在Chapter 3（5）里面我曾经介绍过一个设备端的API：__syncthreads，可以使用这个API来使得线程块内的活跃线程全部到达同一个点。

特别需要注意的一点就是：线程块之间是无法等待的，所以需要把每个线程块得到的结果通过cudaMemcpy拷贝回主机端进行串行计算。

其设备端代码如下：

__global void reduceNeighbored(int *g_idata,int *g_odata,unsigned int n)
{
   
	//set thread ID
	unsigned int tid=threadIdx.x;
	unsigned int idx=blockIdx.x*blockDim.x+threadIdx.x;
	int *idata=g_idata+blockIdx.x*blockDim.x;
	//boundary check
	if (idx>=n) return;
	//in-place reduction in global memory
	for(int stride=1;