首页 > 代码库 > 4.1Reduction模型

4.1Reduction模型

1. Reduction

Reduction是一种广泛使用的计算模型,特别是在并行计算领域。简单地来说,Reduction就是一系列的划分(Partition)和汇总(Summarize)操作的集合:对输入数据分块,对每一个分块汇总,然后再将汇总后的数据视为新的输入数据,重复分块和汇总,直到得到最终结果,可以想象为一个倒置的树。Google和Hadoop的Map/Reduce中的Reduce计算就是一个很好的例子。

Reduction模型满足联合性和交换性:联合性:(a + b) + c = a + (b + c),交换性:a + b = b + c。换言之,Reduction模型对输入数据的顺序没有特定的要求。模型满足这两个性质,则里面的算符也必须满足这两个性质。常见的Reduction算符,也可以说是Reduction操作,包括了Max,Min,Sum,Product等,用户也可以自定义算符,但一定要注意,自定义的算符首先需要满足上述两个性质,其次需要定义单位值(Identity Value)。单位值是指一种特殊的值,它的存在对计算结果没有影响,比如,对Sum计算,0就是单位值。单位值常常被用来填充数据,使得数据长度为2的幂方值或某一特定值的整数倍。

一个Max算符的Reduction树如下图所示:

技术分享

 

2. Reduction树的性能考虑

这里以一个求向量元素之和的计算任务为例来说明如何提高Reduction模型的性能。

向量求和可以如下图表示:

技术分享

显而易见,这个过程包含了O(llogN)步,一共有O(N)的算符操作,每一步都会让活动的线程数减半。

现在来考虑性能。假设一个Warp大小为2个Thread(仅仅是假设,大多数情况下Warp大小为32T),那么一共需要2个Warp。在STEP1时,所有的4个Thread都处于活动状态,而到了STEP2,就只有T0和T2处于活动状态。但是,前面有提到过,Warp是调度的最小单位,所以,只要在Warp中有任何的活动线程,整个Warp就处于活动状态,其中的非活动状态的线程仍然会占用计算资源。因此,在STEP2时,尽管只有2个活动线程,但是仍然会占用4个线程的计算开销。

通过改变Reduction树计算方式可以解决这个问题。注意到上面的计算中,随着计算步骤增加,待计算的两个数据之间的距离是增加的。比如STEP1中3+1=4,3和1间距离为1,而STEP2中4+7=11,4和7间的距离为2。活动线程间的距离是数据距离的2倍,在它们之间是非活动线程。这样在计算中间过程中总有活动线程被非活动线程所分隔,就总会有多余1个的Warp同时包含了活动和非活动线程。

现在将上面的计算步骤反过来,先从距离为4的数据开始计算,即STEP1中T0会计算3+4=7而不是3+1=4。可以验证,这样的计算步骤所有的活动线程都集中在前部,所有的非活动线程都在后部。因此,至多只会有1个Warp同时包含了活动和非活动线程。

从上述例子可以看出,尽管Reduction模型对数据顺序没有要求,不同的计算顺序对性能是有影响的。

3. 代码示例

// just the kernel functionvoid sumOfVec(float *input) {    // take advantage of shared memory to speed up calculation    // each thread take 2 elements from input to shared memory    __shared__ float partialSum[2 * BLOCK_SIZE]    int t = threadIdx.x;    int start = 2 * blockIdx.x * blockDim.x;    partialSum[t] = input[start + t];    partialSum[t + blockDim.x] = input[start + t + blockDim.x];    // reduction begin    for (int i = blockDim.x; i > 0; i /= 2) {        __syncthreads();        if (t < i) {            partialSum[t] = partialSum[t + i];        }    }}            

 

4.1Reduction模型