这次回顾ECE408 Lecture 17，这次介绍了Parallel Scan第二部分。

课程主页：

搬运视频：

目标

  • 掌握并行扫描（前缀和）算法
    • 工作效率与延迟；
    • Brent-Kung树算法；
    • 分层算法；

Kogge-Stone并行Scan算法

回顾

回顾Kogge-Stone并行Scan算法：

提高效率

  • 一种常见的并行算法模式：
    • 平衡树；
    • 在输入数据上构建一个平衡二叉树，并将其从root开始遍历；
    • 树不是一种实际的数据结构，而是一种概念模式；
  • 对于scan：
    • 从叶子节点向上遍历到根，在树的内部节点处构建部分和；
      • root是所有叶子节点的总和；

Brent-Kung并行scan算法

Brent-Kung并行scan步骤

第一步先计算一些特殊的部分和，称为Scan Step：

第二步，根据第一步的结果再次求和，称为Post-Scan Step：

数据流：

Scan Step Kernel

// float T[2*BLOCK_SIZE] is in shared memory
// for previous slide, BLOCK_SIZE is 8
int stride = 1;
while(stride < 2*BLOCK_SIZE) {
    __syncthreads();
    // 2n - 1
    int index = (threadIdx.x+1)*stride*2 - 1;
    if(index < 2*BLOCK_SIZE && (index-stride) >= 0)
    	T[index] += T[index-stride];
    stride = stride*2;
}

在之前的例子中：

// threadIdx.x+1 = 1, 2, 3, 4, 5, 6, 7, 8
// stride = 1, index = 1, 3, 5, 7, 9, 11, 13, 15

分析：

  • 在Scan Step中，索引为$i=2^k(2s + 1) - 1$的位置的和为$\sum_{j=i-2^{k} +1}^i a_j$；
  • 在Post Step中，每一步是使得上面求和的下界$j$逐渐减少到$0$；
    • 这是直观理解，完全理解需要仔细分析；

Post Scan Step Kernel

int stride = BLOCK_SIZE/2;
while(stride > 0) {
    __syncthreads();
    int index = (threadIdx.x+1)*stride*2 - 1;
    if ((index+stride) < 2*BLOCK_SIZE)
    	T[index+stride] += T[index];
    stride = stride / 2;
}
// In our example,
// BLOCK_SIZE=8 stride=4, 2, 1
// for first iteration, active thread = 0 index = 7, +stride = 11

分析

  • 并行scan执行$2\log(n)$次并行迭代；
    • Scan和Post Scan分别使用$\log n$；
    • 一共有$n/2, n/4,\ldots, 1, (2-1), \ldots, (n/4-1), (n/2-1)$次加法；
    • 在我们的示例中，$n = 16$，加法次数为$16/2 + 16/4 + 16/8 + 16/16 + (16/8-1) + (16/4-1) + (16/ 2-1)$；
    • 总加法次数：$(n-1) + (n-2)-(\log(n) -1) = 2\times (n-1) – \log(n)=O(n)$；
  • 加法的总数不超过高效顺序算法的两倍；
    • 在硬件充足的情况下，并行可以轻松克服两倍的工作量；

Kogge-Stone vs. Brent-Kung

  • 与Kogge-Stone相比，Brent-Kung使用的线程数只有一半；
    • 每个线程应该将两个元素加载到共享内存中；
  • Brent-Kung的步数是Kogge-Stone的两倍；
    • Kogge-Stone在GPU中使用block进行并行scan更受欢迎；
    • 块中的线程被“占用”，直到整个块完成；

完整Scan的整体流程，分层方法

在CUDA中使用全局内存

  • 一个线程块的寄存器和共享内存中的数据对其他块不可见；
  • 要使数据可见，必须将数据写入全局内存；
  • 但是，写入全局内存的任何数据在内存fence之前都是不可见的，这通常是通过终止内核执行来完成的；
  • 启动另一个内核以继续执行，由终止的内核完成的全局内存写入对所有线程块都是可见的。

处理任意长度输入的Scan

  • 基于Scan内核构建，最多可处理来自Brent-Kung的$2\times \mathrm{blockDim}.x$个元素；
    • 对于Kogge-Stone，将$\mathrm{blockDim}.x$元素的每个部分分配给一个块；
  • 让每个块使用其$\mathrm{blockIdx}.x$作为索引将其部分和写入Sum数组
    ；
  • 对Sum数组运行并行Scan；
    • 如果对一个块来说太大，可能需要将Sum分解成多个部分；
  • 将扫描的Sum数组值添加到对应的元素部分；

Inclusive Scan定义

定义：

扫描操作采用二元结合运算符$\oplus$和$n$个元素的数组：

并返回前缀和数组：

例子：

如果$\oplus$是加法，那么在数组$[3,1,7,0,4,1,6,3]$上执行exclusive scan操作的结果会返回$[0, 3,4,11,11,15,16,22]$。

为什么要Exclusive Scan

  • 使用缓冲区的起始地址；
  • Inclusive和Exclusive Scan可以互相转换；

例子：

一个简单的Exclusive Scan内核

  • 适配inclusive的Kogge-Stone扫描内核
    • Block 0：
      • 线程0将0加载到（共享）$XY[0]$；
      • 其他线程将（全局）$X[\mathrm{threadIdx}.x-1]$加载到$XY[\mathrm{threadIdx}.x]
        $；
    • 所有其他块：
      • 所有线程将$X[\mathrm{blockIdx}.x\times \mathrm{blockDim}.x+\mathrm{threadIdx}.x-1]$加载到$XY[\mathrm{threadIdex}.x]$；
  • Brent-Kung内核的类似适配，但注意每个线程加载两个元素；
    • 只加载一个零；
    • 所有元素都应该只移动一个位置；