这次回顾ECE408 Lecture 16，这次介绍了Parallel Scan。

课程主页：

搬运视频：

Inclusive Scan定义

定义：

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

并返回前缀和数组：

例子：

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

Scan应用例子

  • 假设我们有一块$100$厘米的木头；
  • 我们需要以下长度（以厘米为单位）的木块；
    • $[3,5,2,7,28,4,3,0,8,1]$；
  • 我们如何快速切割这块木头？会剩下多少？
  • 方法一：
    • 按顺序切割：首先是3，然后是5，接着是2，以此类推；
  • 方法二：
    • 计算前缀和数组；
    • $[3, 8, 10, 17, 45, 49, 52, 52, 60, 61]$（剩下39厘米）；

Scan的经典应用

Scan是并行算法的简单和实用的构建块：

将序列算法：

for(j=1; j<n; j++) out[j] = out[j-1] + f(j);

转换为并行算法：

forall(j) { temp[j] = f(j) };
scan(out, temp);

在很多并行算法中都有用：

  • Radix sort；
  • Quick sort；
  • String comparison；
  • Lexical analysis；
  • Stream compaction；
  • Polynomial evaluation；
  • Solving recurrences；
  • Tree operations；
  • Histograms；
  • Memory buffer allocation；

顺序扫描

给定序列$[x_0, x_1, x_2, \ldots,]$，计算输出$[y_0, y_1, y_2, \ldots]$，使得：

顺序算法的C实现

y[0] = x[0];
for (i = 1; i < max_i; i++)
	y[i] = y[i-1] + x[i];

计算效率：$N$个元素需要$N$次加法，时间复杂度为$O(N)$。

朴素的i并行扫描

  • 分配一个线程来计算每个$y$元素；
  • 让每个线程将$y$元素所需的所有$x$元素相；
    • “只要您不关心性能，并行编程就很容易。”

使用Reduction Tree的并行扫描

  • 将每个输出元素计算为所有先前元素的reduction
    • 一些reduction部分和将在输出元素的计算中共享；
    • 基于IBM的Peter Kogge和Harold Stone在1970年代的硬件附加设计，Kogge-Stone Trees；

Kogge-Stone并行扫描算法展示

Kogge-Stone并行扫描算法说明

每个线程将输入（全局内存）数组中的一个值加载到共享内存数组$T$中，假设$T$的大小是2的幂：

  • 对id为stride到$n-1$的线程进行操作，$n - stride$个活跃的线程；
  • 线程$j$添加元素$T[j]$和$T[j-stride]$并将结果写入元素$T[j]$
  • 每次迭代需要两个同步线程；
    • syncthreads(); // make sure that input is in place
      • 确保输入到位；
    • float temp = T[j] + T[j-stride];
    • syncthreads(); // make sure that previous output has been consumed
      • 确保所有输入元素都已使用；
    • T[j] = temp;

部分实现

__global__
void Kogge_Stone_scan_kernel(float *X, float *Y, int InputSize)
{
    __shared__ float XY[SECTION_SIZE];
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < InputSize) XY[threadIdx.x] = X[i];
    
    for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
        __syncthreads();
        if (threadIdx.x >= stride)
            // This code has a data race condition
            XY[threadIdx.x] += XY[threadIdx.x-stride];
    }
    Y[i] = XY[threadIdx.x];
}

双缓冲

  • 使用两份数据$T_0$和$T_1$；
  • 首先使用$T_0$作为输入，$T_1$作为输出；
  • 每次迭代后切换输入/输出角色：
    • 迭代 0：$T_0$作为输入，$T_1$作为输出；
    • 迭代 1：$T_1$作为输入，$T_0$和输出；
    • 迭代 2：$T_0$作为输入，$T_1$作为输出；
  • 这通常使用两个指针（源和目标）实现，它们将它们的内容从一个迭代交换到下一个迭代；
  • 这消除了对第二个__syncthreads()调用；

示意图

算法：

  • 每次迭代只需要一次syncthreads()
    • syncthreads(); // make sure that input is in place
    • float destination[j] = source[j] + source[j-stride];
    • temp = destination; destination = source; source = temp;
  • 循环后，将目标内容写入全局内存；

工作效率分析

  • Kogge-Stone scan内核执$\log(n)$次并行迭代
    • 每个步骤执行$(n-1), (n-2), (n-4),\ldots,(n- n/2)$次加操作；
    • 添加操作的总数：$n\log n - (n - 1), O(n\log n)$；
  • 这种扫描算法工作效率不高
    • 顺序扫描算法使用$n$次加法；
    • $\log(n)$导致1,000,000个元素时候增加20倍加；
    • 所以通常在$n\le 1024$的块中使用；
  • 由于工作效率低下导致执行资源饱和时，并行算法可能会变慢；