第9章并行模式：前缀和 - Scan算法 (Parallel Patterns: Prefix Sum - Scan Algorithms)

一、导读

9.1 本章学习目标

通过本章的学习，读者将能够：

深入理解前缀和（Prefix Sum/Scan）操作的定义、分类及其在并行计算中的基础地位

掌握Hillis-Steele和Blelloch两种经典并行扫描算法的原理与差异

能够在CUDA中实现高效的inclusive和exclusive scan kernel

理解scan操作与归约的关系，以及如何从scan派生出其他并行原语

掌握基于共享内存的分块scan策略，处理大规模数据

学会使用scan解决实际问题，如数组压缩、基数排序、稀疏矩阵转换等

能够分析scan算法的计算复杂度和内存访问模式，识别优化机会

了解现代GPU库（如CUB、Thrust）中scan的实现特点和最佳实践

9.2 为什么本章如此重要

前缀和（也称为Scan）是并行计算中最基础且最强大的原语之一。它将一个二元结合运算符应用于数组的前缀子序列，生成每个位置的前缀累积结果。虽然概念简单，但scan是许多复杂并行算法的核心构建块。

Scan的重要性体现在以下几个方面：

首先，scan是众多并行算法的基础。基数排序使用scan来计算直方图的前缀和以确定元素位置；数组压缩（filtering）使用scan来分配输出空间；稀疏矩阵格式转换（如CSR到COO）依赖scan来计算行偏移；甚至一些图算法和动态规划也可以通过scan来表达。掌握了scan，就等于掌握了一整套并行算法的钥匙。

其次，scan展示了并行算法设计的精妙之处。直观上，前缀和似乎是一个 inherently serial 的操作，因为每个元素依赖于前面所有元素的结果。然而，通过巧妙的算法设计，我们可以在O(log N)的并行步数内完成scan，实现了显著的并行加速。这种从串行思维到并行思维的转变，是学习并行编程的关键。

第三，scan的实现涉及多种并行优化技术：树形遍历、上下扫描（up-sweep/down-sweep）、共享内存管理、多阶段协调等。通过深入理解scan的实现，读者可以建立起系统的并行算法设计能力。

在实际应用中，高效的scan实现对于性能至关重要。例如，在数据库查询处理中，scan用于快速计算累积分布；在计算机图形学中，scan用于粒子系统的紧凑存储；在深度学习中，scan用于序列模型的某些操作。因此，掌握scan不仅是理论学习的需要，也是工程实践的要求。

本章将从scan的基本定义出发，详细介绍两种经典的并行scan算法，然后通过CUDA代码展示如何实现和优化scan，最后讨论scan在各种应用中的使用方法。

9.3 核心问题引导

什么是前缀和？Inclusive scan和exclusive scan有什么区别？

Hillis-Steele算法的工作原理是什么？它的优缺点分别是什么？

Blelloch算法如何通过up-sweep和down-sweep实现work-efficient的scan？

Scan与归约有什么关系？能否用scan来实现归约？

如何处理超过单个Block处理能力的大规模数据的scan？

Scan在数组压缩和基数排序中是如何应用的？

对于非交换运算符（如字符串连接），scan是否仍然有效？

现代GPU库中的scan实现有哪些优化技巧？

二、核心概念详解

9.1 前缀和的定义与分类

前缀和的定义：

给定一个数组[x₀, x₁, x₂, ..., xₙ₋₁]和一个二元结合运算符⊕，前缀和（scan）产生一个新数组[y₀, y₁, y₂, ..., yₙ₋₁]，其中每个元素是输入数组前缀的累积运算结果。

根据是否包含当前元素，scan分为两类：

Inclusive Scan：

$$y_i = x_0 \oplus x_1 \oplus ... \oplus x_i$$

即每个输出元素包含对应位置的输入元素。

示例（加法）：

Input:  [1, 2, 3, 4, 5]
Output: [1, 3, 6, 10, 15]

Exclusive Scan：

$$y_0 = identity$$

$$y_i = x_0 \oplus x_1 \oplus ... \oplus x_{i-1} \quad \text{for } i > 0$$

即每个输出元素不包含对应位置的输入元素，第一个元素为运算符的单位元。

示例（加法，单位元为0）：

Input:  [1, 2, 3, 4, 5]
Output: [0, 1, 3, 6, 10]

常见运算符的单位元：

运算符	单位元
加法 (+)	0
乘法 (×)	1
最大值 (max)	-∞
最小值 (min)	+∞
逻辑与 (&&)	true
逻辑或 (\	\	)	false

Inclusive与Exclusive的转换：

Inclusive scan和exclusive scan可以相互转换：

Exclusive → Inclusive：将exclusive结果向右移位，并在末尾添加最后一个元素的完整前缀和
Inclusive → Exclusive：将inclusive结果向左移位，并在开头填入单位元

在实际实现中，exclusive scan更为常用，因为它直接给出了"前面所有元素的累积"，这在分配索引时非常有用。

9.2 Hillis-Steele并行Scan算法

Hillis-Steele算法由Guy Steele和Daniel Hillis在1986年提出，是一种直观的并行scan方法。

算法原理：

Hillis-Steele算法通过log₂N次迭代完成inclusive scan。在第k次迭代中，每个元素加上距离为2^k的前面元素的值。

Iteration 0 (stride=1):
  [1, 2, 3, 4, 5]
       ↓
  [1, 3, 5, 7, 9]    // Each element adds element at distance 1

Iteration 1 (stride=2):
  [1, 3, 5, 7, 9]
          ↓
  [1, 3, 6, 10, 14]  // Each element adds element at distance 2

Iteration 2 (stride=4):
  [1, 3, 6, 10, 14]
              ↓
  [1, 3, 6, 10, 15]  // Each element adds element at distance 4

伪代码：

for k = 0 to log2(N) - 1:
    for each element i in parallel:
        if i >= 2^k:
            output[i] = output[i] + output[i - 2^k]

复杂度分析：

并行步数：O(log N)
总工作量：O(N log N) — 这不是work-efficient的
处理器需求：每步需要N个处理器

Hillis-Steele算法的优点是实现简单，并行度高。缺点是总工作量比串行的O(N)多出log N倍，这意味着它做了大量冗余计算。对于GPU这样拥有大量处理器的架构，这种冗余是可以接受的，但在处理器数量有限的情况下，这可能不是最优选择。

CUDA实现：

cpp

__global__ void hillisSteeleInclusiveScan(float* data, int N) {
    extern __shared__ float temp[];
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Load into shared memory
    if (idx < N) {
        temp[tid] = data[idx];
    }
    __syncthreads();
    
    // Hillis-Steele iterations
    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        __syncthreads();
        if (tid >= stride && idx < N) {
            temp[tid] += temp[tid - stride];
        }
    }
    __syncthreads();
    
    // Write back
    if (idx < N) {
        data[idx] = temp[tid];
    }
}

注意这个实现仅限于Block内的scan。对于跨Block的scan，需要额外的处理。

9.3 Blelloch Work-Efficient Scan算法

Blelloch算法由Guy Blelloch在1990年提出，是一种work-efficient的并行scan算法，总工作量仅为O(N)。

算法原理：

Blelloch算法分为两个阶段：Up-Sweep（归约阶段）和Down-Sweep（分发阶段）。

Up-Sweep阶段：

Up-Sweep本质上是一个归约操作，但它保留了中间结果。我们将数组视为一棵二叉树，自底向上计算部分和。

Up-Sweep (addition example with 8 elements):

Level 0: [1, 2, 3, 4, 5, 6, 7, 8]
              ↓       ↓       ↓       ↓
Level 1: [3, X, 7, X, 11, X, 15, X]   // Pairs summed
                    ↓           ↓
Level 2: [X, X, 10, X, X, X, 26, X]   // Pairs of pairs summed
                                ↓
Level 3: [X, X, X, X, X, X, X, 36]    // Total sum

在每一步中，我们将距离为stride的元素对相加，并将结果存储在较高位置。

Down-Sweep阶段：

Down-Sweep自顶向下传播信息，计算最终的scan结果。关键思想是：对于每个节点，左子树的前缀和等于父节点的值，右子树的前缀和等于父节点的值加上左子树的总和。

Down-Sweep (starting from up-sweep result):

Level 3: [X, X, X, X, X, X, X, 36]
         Set last to 0 for exclusive scan
                                ↓
Level 2: [X, X, 10, X, X, X, 26, 0]
         Left gets parent, right gets parent + left
                    ↓           ↓
Level 1: [0, X, 10, X, 26, X, 26, 0]
              ↓       ↓       ↓       ↓
Level 0: [0, 1, 3, 6, 10, 15, 21, 28]  // Final exclusive scan

伪代码：

// Up-Sweep
for d = 0 to log2(N) - 1:
    for each pair (i) in parallel:
        offset = 2^(d+1) - 1
        index = (i + 1) * offset - 1
        data[index] += data[index - 2^d]

// Set last element to identity (for exclusive scan)
data[N-1] = identity

// Down-Sweep
for d = log2(N) - 1 down to 0:
    for each pair (i) in parallel:
        offset = 2^(d+1) - 1
        index = (i + 1) * offset - 1
        temp = data[index - 2^d]
        data[index - 2^d] = data[index]
        data[index] += temp

复杂度分析：

并行步数：O(log N)
总工作量：O(N) — work-efficient
处理器需求：每步需要N/2个处理器

Blelloch算法的总工作量与串行算法相同，只是将这些工作分布在log N个并行步中。这使得它在处理器数量有限的情况下更加高效。

CUDA实现（Block内exclusive scan）：

cpp

__global__ void blellochExclusiveScan(const float* input, float* output, int N) {
    extern __shared__ float temp[];
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Load input into shared memory
    if (idx < N) {
        temp[tid] = input[idx];
    }
    __syncthreads();
    
    // Up-Sweep
    int offset = 1;
    for (int d = N / 2; d > 0; d >>= 1) {
        __syncthreads();
        if (tid < d) {
            int ai = offset * (2 * tid + 1) - 1;
            int bi = offset * (2 * tid + 2) - 1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }
    
    // Clear the last element for exclusive scan
    if (tid == 0) {
        temp[N - 1] = 0.0f;
    }
    
    // Down-Sweep
    for (int d = 1; d < N; d *= 2) {
        offset >>= 1;
        __syncthreads();
        if (tid < d) {
            int ai = offset * (2 * tid + 1) - 1;
            int bi = offset * (2 * tid + 2) - 1;
            
            float t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }
    __syncthreads();
    
    // Write output
    if (idx < N) {
        output[idx] = temp[tid];
    }
}

这个实现假设N是2的幂。对于一般情况，需要将数组填充到最近的2的幂。

9.4 Scan与归约的关系

Scan和归约有着密切的关系：

Scan包含归约：

Inclusive scan的最后一个元素就是整个数组的归约结果。因此，scan可以看作是"保留所有中间结果的归约"。

归约可以用于分块scan：

在处理大规模数据时，我们可以：

对每个Block进行局部scan

提取每个Block的最后一个元素（即该Block的归约结果）

对这些归约结果进行scan

将步骤3的结果加回到步骤1的每个Block

这种方法将全局scan分解为多个局部scan和一个小的全局scan，非常适合GPU的层次结构。

算法对比：

特性	Hillis-Steele	Blelloch
总工作量	O(N log N)	O(N)
并行步数	O(log N)	O(log N)
实现复杂度	简单	较复杂
适用场景	GPU（大量处理器）	CPU/GPU通用

9.5 大规模数据的分块Scan

当数据规模超过单个Block的处理能力时，需要进行分块scan。

三阶段策略：

局部Scan：每个Block对自己的数据段进行exclusive scan

归约与Scan：提取每个Block的最后一个元素（归约结果），对这些值进行scan

合并：将阶段2的结果加回到每个Block的局部scan结果

cpp

// Phase 1: Local exclusive scan per block
__global__ void localScan(const float* input, float* output, float* blockSums, int N) {
    extern __shared__ float temp[];
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int blockSize = blockDim.x;
    
    // Load and pad to power of 2 if needed
    if (idx < N) {
        temp[tid] = input[idx];
    } else {
        temp[tid] = 0.0f;
    }
    __syncthreads();
    
    // Perform Blelloch exclusive scan (simplified, assumes blockSize is power of 2)
    int offset = 1;
    for (int d = blockSize / 2; d > 0; d >>= 1) {
        __syncthreads();
        if (tid < d) {
            int ai = offset * (2 * tid + 1) - 1;
            int bi = offset * (2 * tid + 2) - 1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }
    
    if (tid == 0) temp[blockSize - 1] = 0.0f;
    
    for (int d = 1; d < blockSize; d *= 2) {
        offset >>= 1;
        __syncthreads();
        if (tid < d) {
            int ai = offset * (2 * tid + 1) - 1;
            int bi = offset * (2 * tid + 2) - 1;
            float t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }
    __syncthreads();
    
    if (idx < N) {
        output[idx] = temp[tid];
    }
    
    // Store block sum (the inclusive sum of this block's input)
    if (tid == 0 && idx < N) {
        // Need to compute inclusive sum separately or save before clearing
        blockSums[blockIdx.x] = /* block's total sum */;
    }
}

// Phase 2: Scan the block sums (can use same kernel if few blocks)
// Phase 3: Add phase 2 results to each block's output
__global__ void addBlockOffsets(float* output, const float* blockOffsets, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N && blockIdx.x > 0) {
        output[idx] += blockOffsets[blockIdx.x - 1];
    }
}

9.6 Scan的应用

数组压缩（Filtering）：

从数组中筛选出满足条件的元素，形成紧凑的输出数组。

Input:    [1, -2, 3, -4, 5]
Predicate: x > 0
Flags:    [1,  0, 1,  0, 1]
Scan:     [0,  1, 1,  2, 2]  (exclusive scan of flags)
Output:   [1, 3, 5]  (scatter using scan results as indices)

基数排序（Radix Sort）：

基数排序的每一位排序需要使用scan来计算直方图的前缀和，确定每个桶的起始位置。

稀疏矩阵转换：

将稀疏矩阵从Coordinate Format (COO) 转换为 Compressed Sparse Row (CSR) 格式，需要对行索引进行scan来计算行偏移。

流压缩：

在粒子系统中，移除已死亡的粒子，压缩存活粒子的数组。

9.7 性能优化技巧

共享内存银行冲突：

Blelloch算法的访问模式可能导致bank conflict。使用填充或重新索引可以缓解。

寄存器压力：

Down-Sweep阶段需要临时变量，可能增加寄存器使用。注意监控寄存器溢出。

Warp级优化：

在最后几次迭代中，可以使用warp shuffle代替共享内存，减少同步开销。

使用库函数：

CUB和Thrust提供了高度优化的scan实现，在生产代码中应优先使用。

三、程序员的视角

3.1 在实际项目中的应用

数据库查询：在OLAP查询中，scan用于计算累积总和、运行平均值等窗口函数。

图形渲染：在粒子系统渲染中，scan用于压缩可见粒子列表，避免渲染不可见粒子。

文本处理：在并行字符串处理中，scan用于计算字符偏移，支持快速的子串提取。

机器学习：在某些序列模型中，scan用于累积隐藏状态。

3.2 常见陷阱与误区

忘记单位元：Exclusive scan的第一个元素必须是单位元，否则结果错误。

非2的幂处理：Blelloch算法假设输入大小是2的幂。对于一般情况，需要填充或特殊处理。

溢出问题：对于整数scan，累积和可能溢出。使用更大的数据类型或检查溢出。

跨Block同步：不要试图在一个kernel内进行跨Block的同步。使用多kernel启动。

3.3 与其他层级的联系

与归约的联系：Scan可以看作扩展的归约。许多归约优化技术也适用于scan。

与排序的联系：基数排序依赖scan来计算直方图前缀和。

与前缀积的联系：Scan的概念可以推广到任何结合运算，包括矩阵乘法等。

四、代码实战

4.1 完整的Blelloch Scan实现

cpp

#include <cuda_runtime.h>
#include <stdio.h>

#define BLOCK_SIZE 256

// Blelloch exclusive scan for a single block (assumes power of 2 size)
__global__ void blellochScan(const float* input, float* output, int N) {
    extern __shared__ float temp[];
    
    int tid = threadIdx.x;
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Load input
    if (globalIdx < N) {
        temp[tid] = input[globalIdx];
    } else {
        temp[tid] = 0.0f;
    }
    __syncthreads();
    
    // Up-Sweep
    int offset = 1;
    for (int d = blockDim.x / 2; d > 0; d >>= 1) {
        __syncthreads();
        if (tid < d) {
            int ai = offset * (2 * tid + 1) - 1;
            int bi = offset * (2 * tid + 2) - 1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }
    
    // Set last element to 0 for exclusive scan
    if (tid == 0) {
        temp[blockDim.x - 1] = 0.0f;
    }
    
    // Down-Sweep
    for (int d = 1; d < blockDim.x; d *= 2) {
        offset >>= 1;
        __syncthreads();
        if (tid < d) {
            int ai = offset * (2 * tid + 1) - 1;
            int bi = offset * (2 * tid + 2) - 1;
            
            float t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }
    __syncthreads();
    
    // Write output
    if (globalIdx < N) {
        output[globalIdx] = temp[tid];
    }
}

// Multi-block scan
void multiBlockScan(const float* d_input, float* d_output, int N) {
    int blockSize = BLOCK_SIZE;
    int gridSize = (N + blockSize - 1) / blockSize;
    
    float *d_localOutput, *d_blockSums;
    cudaMalloc((void**)&d_localOutput, N * sizeof(float));
    cudaMalloc((void**)&d_blockSums, gridSize * sizeof(float));
    
    // Phase 1: Local scan
    blellochScan<<<gridSize, blockSize, blockSize * sizeof(float)>>>(
        d_input, d_localOutput, N);
    
    // Extract block sums (last element of each block's inclusive scan)
    // This requires a separate inclusive scan or modification of the kernel
    
    // Phase 2: Scan block sums (assuming gridSize is small)
    // ... (omitted for brevity, use same kernel or CUB)
    
    // Phase 3: Add offsets
    // ... (omitted for brevity)
    
    cudaMemcpy(d_output, d_localOutput, N * sizeof(float), 
               cudaMemcpyDeviceToDevice);
    
    cudaFree(d_localOutput);
    cudaFree(d_blockSums);
}

int main() {
    int N = 8;
    float h_input[] = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f};
    float h_output[N];
    
    float *d_input, *d_output;
    cudaMalloc((void**)&d_input, N * sizeof(float));
    cudaMalloc((void**)&d_output, N * sizeof(float));
    
    cudaMemcpy(d_input, h_input, N * sizeof(float), cudaMemcpyHostToDevice);
    
    blellochScan<<<1, N, N * sizeof(float)>>>(d_input, d_output, N);
    
    cudaMemcpy(h_output, d_output, N * sizeof(float), cudaMemcpyDeviceToHost);
    
    printf("Input:  ");
    for (int i = 0; i < N; i++) printf("%.0f ", h_input[i]);
    printf("\nOutput: ");
    for (int i = 0; i < N; i++) printf("%.0f ", h_output[i]);
    printf("\nExpected (exclusive): 0 1 3 6 10 15 21 28\n");
    
    cudaFree(d_input);
    cudaFree(d_output);
    
    return 0;
}

4.2 使用Scan进行数组压缩

cpp

#include <cuda_runtime.h>
#include <stdio.h>

#define N 1000
#define BLOCK_SIZE 256

// Mark elements that satisfy predicate
__global__ void markElements(const float* input, int* flags, int N, float threshold) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        flags[idx] = (input[idx] > threshold) ? 1 : 0;
    }
}

// Compact array using scan results
__global__ void compactArray(const float* input, const int* scanResult, 
                              float* output, int* flags, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N && flags[idx]) {
        output[scanResult[idx]] = input[idx];
    }
}

int main() {
    float h_input[N];
    for (int i = 0; i < N; i++) {
        h_input[i] = (float)(i % 10 - 5);  // Mix of positive and negative
    }
    
    float *d_input, *d_output;
    int *d_flags, *d_scanResult;
    cudaMalloc((void**)&d_input, N * sizeof(float));
    cudaMalloc((void**)&d_output, N * sizeof(float));
    cudaMalloc((void**)&d_flags, N * sizeof(int));
    cudaMalloc((void**)&d_scanResult, N * sizeof(int));
    
    cudaMemcpy(d_input, h_input, N * sizeof(float), cudaMemcpyHostToDevice);
    
    // Step 1: Mark
    markElements<<<(N + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE>>>(
        d_input, d_flags, N, 0.0f);
    
    // Step 2: Exclusive scan of flags (using CUB or custom implementation)
    // For simplicity, assume we have a scan function
    // scanExclusive(d_flags, d_scanResult, N);
    
    // Step 3: Compact
    // compactArray<<<...>>>(d_input, d_scanResult, d_output, d_flags, N);
    
    printf("Array compression using scan demonstrated.\n");
    
    cudaFree(d_input);
    cudaFree(d_output);
    cudaFree(d_flags);
    cudaFree(d_scanResult);
    
    return 0;
}

五、深入思考

5.1 历史背景

并行scan的研究始于1980年代。Hillis和Steele的1986年论文首次提出了完全并行的scan算法。Blelloch在1990年的工作中指出了work-efficiency的重要性，提出了up-sweep/down-sweep方法。

在GPU领域，Mark Harris的系列文章和NVIDIA SDK示例使scan成为CUDA程序员的标准知识。CUB和Thrust库提供了生产质量的scan实现。

5.2 前沿方向

硬件加速Scan：一些新兴架构正在探索硬件级别的scan支持，如专用指令或网络原语。

异步Scan：在异步并行模型中，scan如何与非阻塞操作结合是一个活跃的研究领域。

近似Scan：对于某些应用，近似的scan结果已经足够。随机化算法可以减少通信和计算开销。

5.3 延伸阅读

Guy Blelloch, "Prefix Sums and Their Applications", 1990

Mark Harris, "Parallel Prefix Sum (Scan) with CUDA", NVIDIA

CUB Library Documentation - DeviceScan

Thrust Library Documentation - exclusive_scan

六、实验笔记

实验1：Hillis-Steele vs Blelloch性能对比

目的：比较两种scan算法的实际性能。

步骤：

实现两种算法的CUDA版本

测试不同数组大小下的执行时间

分析总工作量差异

预期：对于小数组，两者差异不大。对于大数组，Blelloch应更高效，因为总工作量更少。

实验2：Scan在基数排序中的应用

目的：验证scan作为排序子程序的效果。

步骤：

实现基数排序的一位排序步骤

使用scan计算直方图前缀和

测量sort的整体性能

预期：Scan的质量直接影响排序性能。高效的scan可以显著提升排序吞吐量。

七、本章小结

前缀和（Scan）是并行计算中的基础原语，它将二元结合运算符应用于数组的前缀子序列。本章介绍了两种经典的并行scan算法：Hillis-Steele和Blelloch。

Hillis-Steele算法实现简单，并行度高，但总工作量为O(N log N)，不是work-efficient的。Blelloch算法通过up-sweep和down-sweep两阶段处理，实现了O(N)的总工作量，是work-efficient的。

Scan与归约有密切关系：scan的最后一个元素就是归约结果。利用这一关系，我们可以设计分块scan策略，处理大规模数据。

Scan在众多应用中发挥核心作用：数组压缩、基数排序、稀疏矩阵转换等。掌握scan是实现这些高级算法的前提。

在现代GPU编程中，推荐使用CUB或Thrust等库提供的scan实现，它们经过充分优化，可以达到接近硬件峰值的性能。

八、习题

Scan定义：给定数组[3, 1, 4, 1, 5, 9, 2, 6]，计算其inclusive和exclusive scan（加法）。

Hillis-Steele追踪：手动追踪Hillis-Steele算法对4元素数组[1, 2, 3, 4]的执行过程，展示每次迭代后的数组状态。

Blelloch正确性：证明Blelloch算法的down-sweep阶段正确计算了exclusive scan。提示：使用归纳法。

单位元的作用：为什么exclusive scan的第一个元素必须是单位元？如果运算符是乘法，单位元是什么？

非2的幂处理：如果输入大小不是2的幂，如何修改Blelloch算法？填充法的优缺点是什么？

Scan变体：实现一个inclusive-to-exclusive转换kernel。

应用设计：设计一个使用scan实现的并行直方图算法。说明如何使用scan计算bin偏移。

性能建模：估算Blelloch scan的共享内存访问次数。与Hillis-Steele相比，哪个的内存效率更高？

多级Scan：对于10^9元素的数组，设计一个三级scan策略（Block内、Grid内、多Grid）。

综合实现：使用scan实现一个完整的基数排序（至少4位）。测量排序性能并与thrust::sort比较。

第9章 并行模式：前缀和 - Scan算法 (Parallel Patterns: Prefix Sum - Scan Algorithms)