Scan 基礎與 Kogge-Stone 演算法
重點總覽 (Overview)
Scan (prefix sum) 把「每一項都依賴前一項」的 sequential recurrence 轉成平行運算，是 radix sort、stream compaction、polynomial evaluation 等演算法的 primitive。本筆記涵蓋定義、循序基準、naïve 平行版本，以及 Kogge-Stone kernel 與其 work-efficiency 分析。
| 演算法 (Algorithm) | 運算量 (Work) | 步數 (Steps, P=N) | Work-efficient? | 備註 |
|---|---|---|---|---|
| Sequential scan | N − 1 加法, O(N) | N | ✅ (基準) | 單一迴圈累加 |
| Naïve parallel (每個輸出各做一棵 reduction) | N(N−1)/2, O(N²) | N (最長那條) | ❌ | 沒共享 partial sum，又慢又耗資源 |
| Kogge-Stone | N·log₂N − (N−1), O(N log N) | log₂N | ❌ (但遠優於 naïve) | 共享 partial sum；步數少、resource 充足時最快 |
Scan 是「平行演算法的 work complexity 可能高於 sequential」的典型例子。Kogge-Stone 用「多做工 (O(N log N))」換「少步數 (log₂N)」——只有在硬體 execution unit 夠多時才划算。
Scan 的定義：Inclusive vs Exclusive
給定 binary associative operator ⊕ 與輸入 [x₀, x₁, …, x_{n-1}]：
| 類型 | 輸出 | 第 0 個元素 | 直覺 (香腸範例) |
|---|---|---|---|
| Inclusive scan | [x₀, (x₀⊕x₁), …, (x₀⊕…⊕x_{n-1})] |
x₀ |
切點 (cut points)：每段結束的位置 |
| Exclusive scan | [i, x₀, (x₀⊕x₁), …, (x₀⊕…⊕x_{n-2})] |
identity i |
起點 (beginning points)：每段開始的位置 |
- Inclusive：每個輸出包含對應輸入元素的貢獻。
- Exclusive：每個輸出排除對應輸入元素；首位填 operator 的 identity value（加法為
0），末位只累加到x_{n-2}。
範例（⊕ = 加法，輸入 [3 1 7 0 4 1 6 3]）：
- Inclusive →
[3 4 11 11 15 16 22 25]（香腸切點） - Exclusive →
[0 3 4 11 11 15 16 22]（每段起點，用於 memory allocation 回傳起始指標）
- Inclusive → Exclusive：整體右移一格，第 0 位填 identity。
- Exclusive → Inclusive：整體左移一格，末位填「舊末位 ⊕ 末項輸入」。
因此書中只實作 inclusive；exclusive 只需改載入方式（見下）。
循序基準 (Sequential O(N) Baseline)
void sequential_scan(float *x, float *y, unsigned int N) {
y[0] = x[0];
for (unsigned int i = 1; i < N; ++i) {
y[i] = y[i - 1] + x[i]; // 把前一個輸出累加一個新輸入
}
}
- 每個輸出只做 1 次加法 → 全程 N − 1 次加法 → O(N)。
- 這就是 work-efficiency 的最小工作量基準：任何平行版本的 work 都拿來和它比。
Naïve 平行 Scan：為什麼不可行
直覺：讓每個 thread 對自己的輸出位置 y_i 獨立做一棵 reduction（累加 x₀…x_i）。
- 輸出
y_{n-1}需要 n 步，和 sequential 一樣長 → 沒有 speedup（完成時間由最慢的 thread 決定）。 - 總運算量：
Σ_{i=0}^{n-1} i = n(n−1)/2→ O(N²)，比 sequential 還多。 - 改用 reduction tree 雖能降單點步數，但若不共享 partial sum，每棵樹仍是 i 次加法 → 整體仍 O(N²)。
Naïve 平行 scan 既不 work-efficient 又不快，還要更多 execution resource。Kogge-Stone 的關鍵突破就是跨 reduction tree 共享 partial sums。
Kogge-Stone 平行 Scan (Kogge-Stone Algorithm)
源自 1970 年代高速加法器電路設計 (Kogge & Stone, 1973)。是 in-place 演算法，在陣列 XY 上反覆演進：
不變量 (invariant)：經過 k 次 iteration 後，
XY[i]含有「位置 i 及其左側共 2ᵏ 個輸入元素」的和。
演進 (8-element 範例，stride 1→2→4)
初始: x0 x1 x2 x3 x4 x5 x6 x7
| | | | | | | |
stride=1: | +←| +←| +←| +←| +←| +←| +←| XY[i] += XY[i-1]
x0 x0:1 x1:2 x2:3 x3:4 x4:5 x5:6 x6:7
| | | | | | | |
stride=2: | | +←┘ +←┘ +←┘ +←┘ +←┘ +←┘ XY[i] += XY[i-2]
x0 x0:1 x0:2 x0:3 x1:4 x2:5 x3:6 x4:7
| | | | | | | |
stride=4: | | | | +←──┘ +←──┘ +←──┘+←┘ XY[i] += XY[i-4]
x0 x0:1 x0:2 x0:3 x0:4 x0:5 x0:6 x0:7 ← 最終 inclusive scan
(xa:b 表示 x_a + … + x_b；+← 表示「加上左方 stride 距離的舊值」)
- 每個 iteration
stride加倍（1, 2, 4, …, N/2），共 log₂N 個 iteration。 threadIdx.x < stride的 thread 已得最終值，本回合不更新（造成輕微 control divergence，見下）。
Kernel (Fig. 11.3)
__global__ void Kogge_Stone_scan_kernel(float *X, float *Y, unsigned int N) {
__shared__ float XY[SECTION_SIZE]; // block size == SECTION_SIZE
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
if (i < N) XY[threadIdx.x] = X[i]; // 合併載入 global → shared
else XY[threadIdx.x] = 0.0f; // 越界補 identity
for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
__syncthreads(); // (1) 確保上一回合寫入完成
float temp; // partial sum 暫存到 register
if (threadIdx.x >= stride)
temp = XY[threadIdx.x] + XY[threadIdx.x - stride]; // 先「讀」
__syncthreads(); // (2) 確保所有 thread 都讀完舊值
if (threadIdx.x >= stride)
XY[threadIdx.x] = temp; // 才「寫」回
}
if (i < N) Y[i] = XY[threadIdx.x];
}
每個 active thread 同時讀 XY[threadIdx.x] 與 XY[threadIdx.x - stride]。若 thread i 太早把新值寫回 XY[i]，thread i+stride 可能讀到新值而非舊值 → 結果錯誤（時序相依、run-to-run 不可重現）。
- 這是 write-after-read (WAR) hazard，不能用 atomic 解決（直方圖那種是 read-modify-write）。
- 解法：先把和寫入 register
temp→__syncthreads()→ 再統一寫回XY。需要第二個 barrier。
用兩個 shared memory buffer，每回合「讀 buffer A、寫 buffer B」交替。讀寫位置不同 → 沒有 WAR race → 每回合只需 1 個 __syncthreads()（書中 exercise）。
對比 Chapter 10 Reduction：為何 reduction 不需要 temp？
| Reduction (Ch.10) | Kogge-Stone Scan | |
|---|---|---|
| 讀取位置 | input[tid] 與 input[tid+stride] |
XY[tid] 與 XY[tid-stride] |
| 被寫入位置是否被他人讀 | 否（stride 距離位置不被任何 active thread 寫） | 是（共享 partial sum 導致他人會讀） |
| 需要第二個 barrier? | 否 | 是（或 double-buffer） |
Exclusive 版本
只改載入 4 行（取代 Fig. 11.3 的 04–08）：XY[0] 載入 0（identity），其餘 XY[threadIdx.x] 載入 X[i-1]（右移一格），iteration 邏輯完全不變。
小位置的 thread 比大位置早結束（if (threadIdx.x >= stride)）。divergence 只發生在第一個 warp、且只在 stride < warp_size 的回合（stride 1,2,4,8,16）。large block 時影響很小。
速度與 Work Efficiency 分析 (Speed & Work Efficiency)
Work efficiency = 演算法實際做的工作量，距離「最小必要工作量」有多近。最小 = sequential 的 N − 1 次加法。
Kogge-Stone work 推導
每個 iteration 的 active thread 數 = N − stride（inactive = stride）：
Work = Σ (N − stride), stride ∈ {1, 2, 4, …, N/2} ← 共 log₂N 項
= N·log₂N − (1 + 2 + 4 + … + N/2) ← 後段為等比級數 = N−1
= N·log₂N − (N − 1)
→ O(N log₂N)：比 naïve 的 O(N²) 好，但仍比 sequential 的 O(N) 多做工。
用了 N 個 thread；即使許多 thread 停止參與，它們仍佔住整個 warp 的 execution resource 直到 warp 完成。實際消耗接近 N·log₂N（連那 −(N−1) 都省不掉）。
量化範例
| 指標 | 公式 | N=512 | N=1024 |
|---|---|---|---|
| Sequential 步數 | N | 512 | 1024 |
| 比 sequential 多做幾倍工 | ≈ log₂N | ≈ 8× | ≈ 10× |
| 步數 (P=N，資源無限) | log₂N | 9 | 10 |
| 理想步數縮減 | N / log₂N | ≈ 56.9× | ≈ 102× |
| 步數 (P=32 execution units) | (N·log₂N)/P | — | (1024·10)/32 = 320 |
| 實際 speedup (P=32) | N / steps | — | 1024/320 = 3.2× |
- 代價 1：硬體利用率低；若 P 太小，平行版甚至比 sequential 還慢。
- 代價 2：多做的工 = 多耗能源，不適合 mobile / power-constrained。
- 強項：resource 充足時速度極快、control divergence 極少、可用 warp-level shuffle 指令加速。適合處理 512 / 1024 的中等 section，是現代高速 scan 的重要組件。
考試/面試重點 (Exam / Test Patterns)
| 情境 / 關鍵字 | 答案 / 技巧 |
|---|---|
| Inclusive vs exclusive 差別 | inclusive 含本身（切點）；exclusive 排除本身、首位填 identity（起點）。右移/左移一格互轉 |
| Exclusive scan 第 0 個元素是什麼 | operator 的 identity value（加法 → 0） |
| Sequential scan 的 work / 複雜度 | N − 1 次加法，O(N)（work-efficiency 基準） |
| Naïve 平行 scan 為何爛 | 每輸出獨立 reduction → O(N²)、y_{n-1} 仍要 N 步、無 speedup |
| Kogge-Stone 不變量 | k 個 iteration 後 XY[i] = 位置 i 及左側 2ᵏ 個元素的和 |
| Kogge-Stone work / 步數 | work = N·log₂N − (N−1) → O(N log N)；步數 = log₂N |
為何要 temp + 第二個 __syncthreads() |
解決 write-after-read race；先全讀進 register，barrier 後再寫回 |
| Kogge-Stone 的 race 與直方圖的差別 | 此處 WAR（不可用 atomic）；直方圖是 read-modify-write（可用 atomic） |
| 避免第二個 barrier 的方法 | double-buffering（兩個 buffer，讀寫交替） |
| Kogge-Stone control divergence 範圍 | 僅第一個 warp、stride < warp_size（1,2,4,8,16） |
| 何時選 Kogge-Stone | execution resource 充足、section 中等（512/1024）、要最少步數時 |
| 512 元素比 sequential 多做幾倍工 | 約 8 倍（log₂512 = 9） |
Related Notes
- 11-Prefix-Sum-Scan/02-Work-Efficient-Scan-Brent-Kung-and-Coarsening
- 11-Prefix-Sum-Scan/03-Arbitrary-Length-and-Single-Pass-Scan
- 10-Reduction/02-Optimizing-Single-Block-Reduction-Kernel
- 10-Reduction/03-Scaling-Reduction-Hierarchical-and-Coarsening
- 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram
- 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence