reduction

归约问题

并行归约问题

​ 在串行编程中，、最常见的一个问题就是一组特别多数字通过计算变成一个数字，比如加法，也就是求这一组数据的和，这种计算当具有交换律和结合律的时候，我们可以用并行归约的方法处理他们

​ 为什么叫归约，归约是一种常见的计算方式（串并行都可以），归约的归有递归的意思，约就是减少，这样就很明显了，每次迭代计算方式都是相同的（归），从一组多个数据最后得到一个数（约）

归约的方式基本包括如下几个步骤（以加法为例）：

  1. 将输入向量划分到更小的数据块中
  2. 用一个线程计算一个数据块的部分和
  3. 对每个数据块的部分和再求和得到最终的结果

​ 归约问题最常见的加法计算是把向量的数据分成对，然后用不同线程计算每一对元素，得到的结果作为输入继续分成对，迭代的进行，直到最后一个元素，成对的划分常见的方法有以下两种：

  1. 相邻配对：元素与他们相邻的元素配对
    img
  2. 交错配对：元素与一定距离的元素配对
    img

cpu 版本实现交错配对归约计算的代码：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
int recursiveReduce(int *data, int const size) {
	// terminate check
	if (size == 1)
        return data[0];
	// renew the stride
	int const stride = size / 2;
	if (size % 2 == 1) {
		for (int i = 0; i < stride; i++) {
			data[i] += data[i + stride];
		}
		data[0] += data[size - 1];
	}
	else {
		for (int i = 0; i < stride; i++) {
			data[i] += data[i + stride];
		}
	}
	// call
	return recursiveReduce(data, stride);
}

并行归约中的分化

​ 归约问题可视化

img

  • 第一步：如上图把这个一个数组分块，每一块只包含部分数据，我们假定这是线程块的全部数据
  • 第二步：每个线程要完成的任务，橙色圆圈就是每个线程做的操作
  • 第三步：将所有块得到的结果相加，就是最终结果

visualize

我们的内核代码贴出来

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
__global__ void reduceNeighbored(int * g_idata,int * g_odata,unsigned int n) {
	//set thread ID
	unsigned int tid = threadIdx.x;
	//boundary check
	if (tid >= n) return;
	//input data pointer
	int *idata = g_idata + blockIdx.x * blockDim.x;
	//in-place reduction in global memory
	for (int stride = 1; stride < blockDim.x; stride *= 2) {
		if ((tid % (2 * stride)) == 0) {
			idata[tid] += idata[tid + stride];
		}
		//synchronize within block
		__syncthreads();
	}
	//write result for this block to global mem
	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];

}

这里面唯一要注意的地方就是同步指令

1
__syncthreads();

​ 原因还是能从图上找到，我们的每一轮操作都是并行的，但是不保证所有线程能同时执行完毕，所以需要等待，执行的快的等待慢的，这样就能避免块内的线程竞争内存了

​ 理解这个线程同步：线程同步的过程可以视作：一个块内有多个核函数，__syncthreads() 就是让这一个块内多个核函数同步，等最慢的一个核函数完成 

改善并行规约的分化

​ 上面归约显然是最原始的，实际上我们有更加高效的代码

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

这个条件判断给内核造成了极大的分支，如图所示：

img

  • 第一轮 有 12\frac{1}{2} 的线程没用
  • 第二轮 有 34\frac{3}{4} 的线程没用
  • 第三轮 有 78\frac{7}{8} 的线程没用

​ 可见线程利用率是非常低的，因为这些线程在一个线程束，所以，只能等待，不能执行别的指令,对于上面的低利用率，可以想到下面这个方案来解决：

img

注意橙色圆形内的标号是线程符号，这样的计算线程的利用率是高于原始版本的，核函数如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
__global__ void reduceNeighboredLess(int * g_idata,int *g_odata,unsigned int n) {
	unsigned int tid = threadIdx.x;
	unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
	// convert global data pointer to the local point of this block
	int *idata = g_idata + blockIdx.x*blockDim.x;
	if (idx > n)
		return;
	//in-place reduction in global memory
	for (int stride = 1; stride < blockDim.x; stride *= 2) {
		//convert tid into local array index
		int index = 2 * stride * tid;
		if (index < blockDim.x) {
			idata[index] += idata[index + stride];
		}
		__syncthreads();
	}
	//write result for this block to global men
	if (tid == 0)
		g_odata[blockIdx.x] = idata[0];
}

最关键的一步就是

1
int index = 2 * stride *tid;

​ 这一步保证 index 能够向后移动到有数据要处理的内存位置，而不是简单的用 tid（thread id）对应内存地址，导致大量线程空闲