1. 基本概念
-
Thrust 是 NVIDIA CUDA 提供的類似 C++ STL 的并行算法庫。
-
Scan (前綴和):給定數組
[a0, a1, a2, ...]
,產生前綴和序列。 -
Exclusive Scan (排他前綴和):
- 輸出位置
i
存放的是輸入數組中 0 到 i-1 的累積結果。 - 換句話說,結果向右錯位一格,首位放初始值。
- 輸出位置
公式:
output[0] = init
output[i] = input[0] ⊕ input[1] ⊕ ... ⊕ input[i-1], (i > 0)
其中 ⊕
是二元操作符(默認是加法)。
和 inclusive_scan 的區別:
- inclusive:包含當前位置 →
output[i] = input[0] + ... + input[i]
- exclusive:不包含當前位置 →
output[i] = input[0] + ... + input[i-1]
2. 函數原型
// 版本 1:默認加法
template <typename InputIterator, typename OutputIterator>
OutputIterator exclusive_scan(InputIterator first, InputIterator last, OutputIterator result);// 版本 2:指定初始值
template <typename InputIterator, typename OutputIterator, typename T>
OutputIterator exclusive_scan(InputIterator first, InputIterator last, OutputIterator result, T init);// 版本 3:指定二元運算符
template <typename InputIterator, typename OutputIterator, typename T, typename BinaryFunction>
OutputIterator exclusive_scan(InputIterator first, InputIterator last, OutputIterator result, T init, BinaryFunction binary_op);
參數說明:
first, last
:輸入區間result
:輸出區間init
:初始值(默認 0)binary_op
:二元操作(默認thrust::plus<T>()
)
返回值:輸出迭代器(即 result + (last - first)
)
3. 示例代碼
🔹 默認加法 + 默認初始值
#include <thrust/scan.h>
#include <thrust/device_vector.h>
#include <iostream>int main() {thrust::device_vector<int> d_in{1, 2, 3, 4};thrust::device_vector<int> d_out(4);thrust::exclusive_scan(d_in.begin(), d_in.end(), d_out.begin());// 輸出:0 1 3 6for (int x : d_out) std::cout << x << " ";
}
解釋:
- 輸入
[1, 2, 3, 4]
- 輸出
[0, 1, 3, 6]
指定初始值
thrust::exclusive_scan(d_in.begin(), d_in.end(), d_out.begin(), 10);
輸入 [1, 2, 3, 4]
,結果 [10, 11, 13, 16]
自定義運算符(乘法)
struct multiply
{__host__ __device__ int operator()(const int &a, const int &b) const {return a * b;}
};thrust::exclusive_scan(d_in.begin(), d_in.end(), d_out.begin(), 1, multiply());
輸入 [1, 2, 3, 4]
,結果 [1, 1, 2, 6]
4. 內部實現原理
-
exclusive_scan
在 GPU 上通常使用 并行前綴和算法(Blelloch Scan):- 上行階段(reduce):樹形歸約,計算部分和。
- 下行階段(down-sweep):回傳累積和,得到 exclusive 結果。
-
復雜度:
O(n)
-
并行化效率:在 CUDA 中可利用 warp shuffle / shared memory 優化。
5. 常見應用
- 數組索引計算:比如稀疏矩陣非零元素位置。
- 并行 compact/filter:布爾標記數組 → 前綴和 → 計算輸出位置。
- 動態內存分配:統計每個線程寫入位置。
- 圖算法:CSR 格式構建、鄰接表索引。
6. 易錯點
exclusive vs inclusive:結果差一個位置,常見 bug。
初始值:沒設置時默認 0,很多人誤以為會報錯。
in-place 使用:輸入和輸出可以是同一個 vector(安全)。
自定義運算符:必須滿足結合律(associative),否則結果不穩定。
7. 性能優化
exclusive_scan
已由 Thrust 內部優化,適合中大規模數組。- 小規模數組時,CPU
std::exclusive_scan
可能更快。 - 多次調用時建議 預分配 device_vector,避免反復分配內存。
8、 exclusive_scan
vs inclusive_scan
對比表
特性 | exclusive_scan | inclusive_scan |
---|---|---|
定義 | 輸出位置 i 是輸入 [0..i-1] 的累積和,當前位置 不包含 | 輸出位置 i 是輸入 [0..i] 的累積和,當前位置 包含 |
首元素 | 結果 output[0] = init (默認 0) | 結果 output[0] = input[0] |
數學表達式 | out[i] = init ⊕ in[0] ⊕ ... ⊕ in[i-1] | out[i] = in[0] ⊕ in[1] ⊕ ... ⊕ in[i] |
輸入 [1, 2, 3, 4] ,init=0 | 輸出 [0, 1, 3, 6] | 輸出 [1, 3, 6, 10] |
典型用途 | 計算 索引偏移、數組 compaction、CSR 圖結構 | 直接得到前綴累計量,例如直方圖累計、累積分布函數 (CDF) |
9、典型應用場景代碼
1?、數組索引偏移(exclusive_scan)
常見于 稀疏矩陣 / 圖的 CSR 格式構建
#include <thrust/scan.h>
#include <thrust/device_vector.h>
#include <iostream>int main() {thrust::device_vector<int> flags{1, 0, 1, 1, 0};thrust::device_vector<int> indices(flags.size());// exclusive_scan:計算每個位置在輸出數組的起始索引thrust::exclusive_scan(flags.begin(), flags.end(), indices.begin());// 輸出: 0 1 1 2 3for (int x : indices) std::cout << x << " ";
}
解釋:
flags
表示某位置是否有效exclusive_scan
得到每個有效元素在輸出數組中的索引位置
2?、數組 compaction(exclusive_scan + scatter)
保留布爾條件為真的元素,常用于 過濾數據
#include <thrust/scan.h>
#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <iostream>int main() {thrust::device_vector<int> data{3, 7, 0, 4, 9};thrust::device_vector<int> flags{1, 0, 1, 0, 1};thrust::device_vector<int> indices(flags.size());thrust::device_vector<int> result(3); // 預估有效數量thrust::exclusive_scan(flags.begin(), flags.end(), indices.begin());// 根據 flags 把有效元素寫入結果數組for (int i = 0; i < flags.size(); i++) {if (flags[i]) result[indices[i]] = data[i];}// 輸出: 3 0 9for (int x : result) std::cout << x << " ";
}
3?、累積分布函數 CDF(inclusive_scan)
常用于 直方圖后處理
#include <thrust/scan.h>
#include <thrust/device_vector.h>
#include <iostream>int main() {thrust::device_vector<int> hist{2, 3, 5, 4};thrust::device_vector<int> cdf(hist.size());thrust::inclusive_scan(hist.begin(), hist.end(), cdf.begin());// 輸出: 2 5 10 14for (int x : cdf) std::cout << x << " ";
}
解釋:
- 輸入直方圖
[2,3,5,4]
inclusive_scan
得到累積分布[2,5,10,14]
4?、并行前綴乘積(inclusive_scan with multiply)
用于 概率鏈式計算
struct multiply {__host__ __device__ float operator()(float a, float b) const {return a * b;}
};thrust::device_vector<float> probs{0.9, 0.8, 0.7, 0.6};
thrust::device_vector<float> cumprod(probs.size());thrust::inclusive_scan(probs.begin(), probs.end(), cumprod.begin(), multiply());// 輸入: 0.9 0.8 0.7 0.6
// 輸出: 0.9 0.72 0.504 0.3024
10、 總結
- exclusive_scan:計算“排他前綴和”,結果右移一位,首位為初始值。
- 三種重載:默認 / 指定 init / 指定 init + 運算符。
- 應用廣泛:稀疏矩陣、圖算法、并行 compaction。
- 注意事項:exclusive vs inclusive,init 值,自定義運算符必須結合律。