这次回顾ECE408 Lecture 15，这次介绍了Reduction Trees。

课程主页：

搬运视频：

Amdahl定律

Amdahl定律描述一个程序的可优化程度，一个程序分为并行和非并行部分，只有并行部分可以优化：

其中：

  • $k$是可并行的比例；
  • $p$是并行的资源；

所以加速比例为$\frac{1}{(1-k)+\left(\frac{k}{p}\right)}$。

Reduction无处不在，往往是并行计算的最后阶段

  • Reduction将一组输入转化为单个值；
    • Max, Min, Sum, Product；
  • 基本上，任何满足以下条件的运算符/函数都是Reduction：
    • 满足结合律和交换律；
    • 有幺元，即一个特殊的元素，sum时是0，product时是1；

Partition和Summarize

  • 数据集的常用策略
    • 如果基本操作满足结合律和交换律（即，可重新排序）
      • 将数据集分成更小的块；
      • 让每个线程处理一个块；
      • 使用Reduction Tree将每个块的结果汇总为最终答案；
  • 大数据云框架（来自谷歌和其他公司）为这种模式提供支持；

Reduction使其他技术成为可能

  • Reduction归约通常需要作为并行计算的最后阶段；
  • 例如：私有化
    • 多个线程写入输出位置；
    • 复制输出位置，使每个线程都有一个私有的输出位置；
    • 使用Reduction Tree将私有位置的值组合到原始输出位置；

实例：锦标赛使用”max” Reduction

一个Reduction的示例就是锦标赛：

顺序Reduction算法需要$O(N)$次操作

  • 将结果初始化为归约操作的特殊值$I$；
    • Max Reduction对应最小值；
    • Min Reduction对应最大值；
    • Sum Reduction对应0；
    • Production Reduction对应1；
  • 遍历输入，对结果值和当前输入值进行Recution运算；
    • 伪代码：
      • result <- I;
        • for each value X in input
          • result <- result op X;

并行Reduction需要$O(\log N)$次操作

并行算法效率很高

对于$N$个输入，总操作数为：

所示的并行算法工作效率高：需要与顺序算法相同的工作量（开销可能不同）。

但是需要大量资源

  • 对于$N$个输入值，步数为$\log (N)$，如果有足够的执行资源：
    • $N=1,000,000$需要$20$步！
  • 我们需要多少并行度？
    • 平均而言，$(N-1)/\log(N)$，在我们的示例中为$50,000$。
    • 但是峰值是$N/2$，在我们的示例中为$500,000$。

Sum Reduction例子

  • 并行实现：
    • 在每个步骤中将每个线程中的两个值相加；
    • 递归地将线程数减半；
    • 对$n$个元素执行$\log(n)$步，需要$n/2$个线程；
  • 假设使用共享内存进行in-place reduction；
    • 原始向量在设备全局内存中；
    • 共享内存用于保存部分和向量；
    • 每一步都使部分和向量更接近总和；
    • 最后的总和将在元素0中；
    • 由于使用共享内存，全局内存流量应该最小；

图示

代码

// Stride is distance to the next value being
// accumulated into the threads mapped position
// in the partialSum[] array
for (unsigned int stride = 1;
	stride <= blockDim.x; stride *= 2)
{
    __syncthreads();
    if (t % stride == 0)
    	partialSum[2*t]+= partialSum[2*t+stride];
}

如果不使用__syncthreads()，那么结果会有问题：

Block完成后的几个选项

  • 在所有reduction步骤之后，线程0
    • 从partialSum[0]写入块的总和；
    • 进入由blockIdx.x索引的全局向量；
  • 向量的长度为$N/ (2\times \mathrm{numBlocks})$；
    • 如果长度很小，将向量传输到主机并在CPU上求和；
    • 如果长度很大，再次启动内核（再一次），内核也可以使用原子操作累加到全局sum，稍后将介绍；

Segmented Reduction

执行资源分析

  • 所有线程在第一步中都处于活动状态。在所有后续步骤中，两个控制流路径：
    • 执行加法，或者什么都不做；
    • 什么都不做仍然会消耗执行资源；
  • 最多一半的线程在第一步之后执行加法：
    • （所有具有奇数索引的线程在第一步之后被禁用）；
    • 第五步之后，整个wraps什么都不做：资源利用率低，但没有divergence；
    • 活动线程束只有一个活动线程；
  • 最多五个步骤（如果限制为1024个线程）。

更好的策略

让我们试试这个方法：

  • 在每一步中；压缩部分和到partialSum数组的第一个位置；

这样做可以使活动线程保持连续。

示例

代码

for (unsigned int stride = blockDim.x; stride >= 1; stride /= 2)
{
    __syncthreads();
    if (t < stride)
    	partialSum[t] += partialSum[t+stride];
}

执行资源分析

  • 给定1024个线程；
  • 块加载2048个元素到共享内存；
  • 前六步没有分支分歧：
    • 1024、512、256、128、64和32个连续线程处于活动状态；
    • 每个warp中的线程全部处于活动状态或全部处于非活动状态；
  • 最后六个步骤有一个活动warp（最后五个步骤有分支分歧）；

完整算法

__shared__ float partialSum[2*BLOCK_SIZE];

unsigned int t = threadIdx.x;
unsigned int start = 2*blockIdx.x*blockDim.x;
partialSum[t] = input[start + t];
partialSum[blockDim.x+t] = input[start + blockDim.x + t];
for (unsigned int stride = blockDim.x; stride >= 1; stride /= 2)
{
    __syncthreads();
    if (t < stride)
    	partialSum[t] += partialSum[t+stride];
}

小结

  • 尽管“操作”的数量是$N$，但每个“操作”都涉及复杂得多的地址计算和中间结果操作。
  • 如果并行代码在单线程硬件上执行，它会比基于原始顺序算法的代码慢。

进一步改进

  • 该问题的计算强度较低
    • 每读取一个4B值一个操作；
    • 因此，请关注内存合并并避免计算资源使用不当；
  • 试验超参数：
    • 每块的线程数量；
    • 每个SM的块数量；
    • 线程粒度（3-1Reduction）；

终止状态

  • 较小的块可能看起来很有吸引力：
    • 当wrap处于活动状态时；
    • 每个SM的每个块都有一个warp；
  • 但可能有更好的方法，例如：
    • 在32个元素（或64个或128个元素）处停止减少，然后交给下一个内核；

整体流程

假设有8个SM，所以有16个block：

  1. 将整个数据集分成16块；
  2. 读取数据，直到足以填满共享内存；
  3. 计算，直到不需要某些线程为止。

  4. 然后加载更多数据！

  5. 重复直到数据用完；

  6. 然后让并行度下降。
    1. （在Host上收集16个值并进行Reduction。）