第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实现:
__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):
__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结果
// 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实现
#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进行数组压缩
#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比较。