因为std::inclusive_scan是CPU端顺序/线程并行实现，不生成GPU可执行kernel，也无法利用warp-level原语（如__shfl_down_sync），故无法直接在GPU上运行。
为什么不能直接用 std::inclusive_scan 上 GPU？
因为 std::inclusive_scan 是 CPU 端顺序/线程并行实现，不生成 GPU 可执行的 kernel，也不能利用 warp-level 原语（比如 __shfl_down_sync）。直接丢数组到 CUDA kernel 里手写串行前缀和，会彻底失去并行性——每个线程等前一个算完，O(n) 时间，毫无意义。
GPU 上高效 Scan 必须分层：先做 block 内快速 scan（用 shared memory + warp shuffle），再对各 block 的末尾元素做全局 scan（reduce-then-scan），最后用 block scan 结果修正本块输出。这是经典两阶段法（Hillis-Steele 或 Blelloch 变种），CUDA Toolkit 自带的 thrust::exclusive_scan 就是这么干的。
用 thrust::device_vector + thrust::exclusive_scan 最简可行路径
90% 场景下，别自己重写。Thrust 已针对不同数据规模、GPU 架构（Turing/Ampere）自动选择最优策略（比如小数组用单 block，大数组切分+多级 reduce）。
- 输入必须是 device-side 连续内存：
thrust::device_vector<int></int>，不是std::vector或裸指针（否则触发隐式 host-device 拷贝，性能归零） - 注意 inclusive vs exclusive：如果需要 [a₀, a₀+a₁, a₀+a₁+a₂]，用
thrust::inclusive_scan；若要 [0, a₀, a₀+a₁]，才用exclusive_scan - 自定义二元操作符（如 max-scan）需确保满足结合律，且 CUDA 支持该 functor 的 device 编译（加
__device__ __host__）
thrust::device_vector<float> d_in = {...};
thrust::device_vector<float> d_out(d_in.size());
thrust::inclusive_scan(d_in.begin(), d_in.end(), d_out.begin());手动写 block-level scan 时，__shfl_down_sync 的坑在哪？
Warp shuffle 是 block 内 scan 的核心加速点，但极易出错：
立即学习“C++免费学习笔记（深入）”；
- 必须用 mask（如
0xffffffff），不能传 0 —— 否则在部分 warp 活跃线程数 - 偏移量
delta必须是 2 的幂（1,2,4,8,16），且threadIdx.x % 32 >= delta的线程才参与 shuffle，否则读到旧值或越界 - shared memory 的 bank conflict：若按线程号顺序存
sdata[tid]，当 tid=0,16,32… 时会撞同一 bank，拖慢 5–10 倍；改用sdata[tid ^ 16]或 padding 可缓解
典型错误现象：cuda-memcheck 报 invalid address，或结果每 32 个元素出现重复/跳变——基本就是 shuffle mask 或 bank conflict 搞错了。
大规模数据（>1GB）Scan 的吞吐瓶颈常不在算法，而在内存访问模式
即便用了最优两阶段 scan，实测发现瓶颈常是 global memory 的 coalescing 效率：
- block scan 阶段，每个 warp 读
d_in[base + tid], d_in[base + tid + 1], ...是连续的，没问题 - 但 global scan 阶段，各 block 只取末尾一个元素（如
d_in[block_id * N - 1]），这些地址跨距极大，变成随机访存，带宽利用率掉到 30% 以下 - 解决办法：预处理阶段把所有 block 尾部元素 gather 到连续数组（用
thrust::gather+ stride iterator），再 scan 这个紧凑数组，最后 scatter 回去
这个细节在论文里常被忽略，但实测对 2GB 浮点数组能提升 1.7× 总耗时——尤其在 A100/H100 上，global memory 带宽吃不满比算力不足更常见。
