前綴和 (Prefix Sum / Scan) 與 Work Efficiency 練習題 (Practice - Scan Foundations and the Kogge-Stone Algorithm)
Related Concepts
- 11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone — Scan 基礎與 Kogge-Stone 演算法
- 11-Prefix-Sum-Scan/02-Work-Efficient-Scan-Brent-Kung-and-Coarsening — 高 Work-Efficiency 的 Scan：Brent-Kung 與 Thread Coarsening
- 11-Prefix-Sum-Scan/03-Arbitrary-Length-and-Single-Pass-Scan — 任意長度輸入與 Single-Pass Scan
| 關鍵字 / 情境 | 答案 / 重點 |
|---|---|
| Inclusive vs Exclusive | inclusive 含本身（切點）；exclusive 排除本身、首位填 identity（起點）；右移/左移一格互轉 |
| 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（P=N） |
temp + 第二個 __syncthreads() |
解 write-after-read (WAR) race；不可用 atomic（與直方圖 read-modify-write 不同） |
| Double-buffering | 兩個 shared buffer 讀寫交替 → 省掉第二個 barrier |
| Kogge-Stone control divergence | 僅第一個 warp、stride < warp_size（1,2,4,8,16） |
| Brent-Kung 兩階段 | Reduction tree（向上，N−1 ops）+ reverse distribution tree（向下，N−1−log₂N ops） |
| Brent-Kung 總操作數 | 2N − 2 − log₂N ≈ 2N − 3，O(N)，data-scalable，最多 sequential 的 2 倍 |
| Brent-Kung threads / section 上限 | SECTION_SIZE/2 threads，每 thread 載/寫 2 元素；1024 threads → 2048 元素 |
| Brent-Kung reduction 映射 | index = (threadIdx.x+1)*2*stride − 1，連續 threads 對應間隔位置以避免 divergence |
| 三階段 coarsening | P1 各 thread sequential scan → P2 段末元素 block scan → P3 加前驅段掃描值 |
| 三階段 steps 公式 | (2(N−T) + T·log₂T)/P，例 N=1024,T=64,P=32 → 72 |
| Hierarchical / segmented scan | K1 局部 scan → K2 對 S 陣列 scan → K3 加回；S 大小 = N/SECTION_SIZE |
| 三-kernel 缺點 | S 的 store/load 在獨立 kernel，延遲無法與運算重疊 |
| Single-pass / domino scan | 單 kernel，partial sum 在相鄰 block 間單向骨牌傳遞；traffic 多落 L2、可重疊 |
| Adjacent synchronization | global flags[] + scan_value[] + atomicAdd + __threadfence + __syncthreads，僅 leader |
__threadfence() 目的 |
確保 scan_value 先寫進 global memory，才設旗標（防讀到舊值） |
| Dynamic block index | atomicAdd(blockCounter,1) 取得 bid，保證 i−1 已排程 → 防 deadlock |
Question 1 - Inclusive vs Exclusive Scan [recall]
情境/題目：給定加法 operator 與輸入
[3 1 7 0 4 1 6 3]，分別寫出 inclusive scan 與 exclusive scan 的輸出，並說明兩者差別與 exclusive 第 0 個元素的值。
Inclusive → [3 4 11 11 15 16 22 25]（含本身，香腸「切點」）；Exclusive → [0 3 4 11 11 15 16 22]（排除本身，每段「起點」）。Exclusive 第 0 個元素是 operator 的 identity value，加法為 0；末位只累加到 x_{n-2}。兩者可右移/左移一格並補一格互轉。
Question 2 - Sequential Baseline 的 Work [recall]
情境/題目：循序 inclusive scan 對 N 個元素做多少次加法？其計算複雜度為何？為什麼這個數字在分析平行 scan 時很重要？
循序版 y[i] = y[i-1] + x[i] 每個輸出做 1 次加法，全程 N − 1 次加法 → O(N)。這是 work-efficiency 的最小工作量基準：任何平行 scan 的 work 都拿來和它比，越接近 N−1 越 work-efficient。
Question 3 - Naïve 平行 Scan 為何不可行 [recall]
情境/題目：若讓每個 thread 對自己的輸出位置獨立做一棵 reduction（累加
x₀…x_i），其總運算量與最長路徑步數各是多少？為何幾乎沒有 speedup？
總運算量 Σ_{i=0}^{n-1} i = n(n−1)/2 → O(N²)，比 sequential 還多；而 y_{n-1} 仍需 n 步，和 sequential 一樣長。完成時間由最慢 thread 決定 → 無 speedup，又耗更多資源。問題根源是沒有跨 reduction tree 共享 partial sums，這正是 Kogge-Stone 的突破點。
Question 4 - Kogge-Stone 不變量 [recall]
情境/題目：Kogge-Stone 是 in-place 演算法，在陣列
XY上反覆演進。請寫出「經過 k 次 iteration 後XY[i]內容」的不變量，並說明 stride 序列與總 iteration 數。
經過 k 次 iteration 後，XY[i] 含有「位置 i 及其左側共 2ᵏ 個輸入元素」之和（如 k=1 → x_{i-1}+x_i；k=2 → x_{i-3}+…+x_i）。每回合 stride 加倍（1, 2, 4, …, N/2），共 log₂N 個 iteration；threadIdx.x < stride 的 thread 已達最終值，本回合不更新。
Question 5 - Write-After-Read Race 與 temp 變數 [recall]
情境/題目：Kogge-Stone kernel 為何在迴圈內需要一個 register
temp與「第二個」__syncthreads()？這個 race 與 Chapter 9 直方圖的 race 有何不同，能否用 atomic 解決？
每個 active thread 同時讀 XY[threadIdx.x] 與 XY[threadIdx.x-stride]；若某 thread 太早把新值寫回，鄰居可能讀到新值而非舊值 → write-after-read (WAR) hazard，結果時序相依、run-to-run 不可重現。解法：先把和算進 register temp → __syncthreads()（確保大家都讀完舊值）→ 再統一寫回。不能用 atomic（atomic 解的是直方圖的 read-modify-write race）；另一替代法是 double-buffering。
Question 6 - Kogge-Stone 的 Work 與複雜度 [recall]
情境/題目：推導 Kogge-Stone kernel（單一 block）的總加法次數公式，並給出其複雜度。N=512 時大約比 sequential 多做幾倍工？
每 iteration active thread 數 = N − stride，stride ∈ {1,2,4,…,N/2}（共 log₂N 項）：work = Σ(N − stride) = N·log₂N − (N−1) → O(N log₂N)。N=512 時 log₂512 = 9 → 約比 sequential 多做 ≈ 8× 工。比 naïve 的 O(N²) 好，但仍非 work-efficient。
Question 7 - Brent-Kung 兩階段與操作數 [recall]
情境/題目：Brent-Kung scan 由哪兩個階段組成？各做多少次加法？合計總操作數與複雜度為何？為何稱為 data-scalable？
① Reduction tree（向上歸約，算出整段總和）= (N/2)+(N/4)+…+1 = N − 1 ops；② Reverse distribution tree（向下把 partial sums 推給未完成位置）= (2−1)+(4−1)+…+(N/2−1) = N − 1 − log₂N ops。合計 2N − 2 − log₂N ≈ 2N − 3 → O(N)。無論 N 多大都不超過 sequential 的 2 倍，操作數隨輸入線性成長，故稱 data-scalable algorithm。
Question 8 - Single-Pass 的 Adjacent Synchronization [recall]
情境/題目：Single-pass (domino) scan 用 adjacent synchronization 在相鄰 block 間傳遞 partial sum。請列出用到的關鍵元件，並說明
__threadfence()為何不可省略。
元件：global flags[]（producer 設旗標、consumer 輪詢構成握手）、global scan_value[]（傳遞的 partial sum 本體）、atomicAdd（原子更新旗標/counter）、__threadfence()、__syncthreads()；只由每個 block 的 leader thread (e.g. threadIdx.x==0) 執行。__threadfence() 確保 scan_value[bid+1] 先寫進 global memory，才用 atomic 設旗標；否則 weak memory model 下 consumer 可能讀到尚未寫入的舊值 → 結果錯誤。
Question 9 - 手算 Kogge-Stone [application]
情境/題目：對陣列
[4 6 7 1 2 8 5 2]執行平行 inclusive Kogge-Stone scan，寫出每個 step（stride=1,2,4）之後的中間陣列狀態（Exercise 1）。
初始：[4 6 7 1 2 8 5 2]
stride=1（XY[i]+=XY[i-1]）：[4 10 13 8 3 10 13 7]
stride=2（XY[i]+=XY[i-2]）：[4 10 17 18 16 18 16 17]
stride=4（XY[i]+=XY[i-4]）：[4 10 17 18 20 28 33 35] ← 最終 inclusive scan。
（驗證：4, 4+6, …, 4+6+7+1+2+8+5+2 = 35。）
Question 10 - 把 Inclusive Kernel 改成 Exclusive [application]
情境/題目：要把 Fig. 11.3 的 Kogge-Stone inclusive kernel 改成 exclusive scan，只需修改載入
XY的那幾行。請寫出修改方式，並說明迭代邏輯是否要改。
只改載入：XY[0] 載入 0（identity），其餘 XY[threadIdx.x] 載入 X[i-1]（整體右移一格），其餘越界補 0。對應程式碼：
if (i < N && threadIdx.x != 0) XY[threadIdx.x] = X[i-1];
else XY[threadIdx.x] = 0.0f;
迭代（for-loop）邏輯完全不變，只是頂端元素對齊不同；Brent-Kung 同理也只需調整載入。
Question 11 - 三階段 Coarsening 的步數估算 [application]
情境/題目：三階段 coarsened scan（Phase 2 用 Kogge-Stone），N=1024 元素、T=64 threads、P=32 execution units。寫出三個 phase 的 work 公式並估算總步數；和單純 Kogge-Stone 的 320 步比較。
Phase 1 = N − T = 960、Phase 2 = T·log₂T = 64·6 = 384、Phase 3 = N − T = 960。
總步數 ≈ (2(N−T) + T·log₂T)/P = (960 + 384 + 960)/32 = 72 步。
遠優於單純 Kogge-Stone 的 320 步，因為大部分工作交給最有效率的 sequential scan，Phase 2 只對少量段末元素做平行 scan。
Question 12 - Kogge-Stone vs Brent-Kung 的速度權衡 [analysis]
情境/題目：Brent-Kung 的 work-efficiency 比 Kogge-Stone 高（O(N) vs O(N log N)），但「更有效率」是否就一定「更快」？分別在「執行資源無限 (P≥N)」與「資源有限 (P 小)」兩種情況下比較，並用 1024 元素、32 execution units 的數字佐證。
不一定。關鍵在 P（執行單元數）：
- 資源無限 (P≥N)：步數主導。Brent-Kung 多了 reverse tree，約需 Kogge-Stone 的 2 倍步數 → 反而較慢。
- 資源有限 (P 小)：總操作數主導。Brent-Kung 操作數約 2N、Kogge-Stone 約 N log N，差距大 → Brent-Kung 較快。
數字（1024 元素、P=32）：Kogge-Stone ≈(1024·10)/32= 320 步 → speedup ≈ 3.2×；Brent-Kung ≈(2·1024−2−10)/32+ ~5 步 divergence ≈ 73.6 步 → speedup ≈ 14×。故資源有限時 Brent-Kung 勝；資源/延遲很大時 Kogge-Stone 反而占優。另注意 CUDA 上 Brent-Kung 的優勢被稀釋：inactive threads 因 SIMD 被綁在同一 warp 仍消耗執行資源。
Question 13 - 三-Kernel Segmented vs Single-Pass Domino [analysis]
情境/題目：對任意長度輸入，三-kernel hierarchical scan 與 single-pass (domino) scan 都能算出全域結果。請從 global memory traffic 與「延遲能否與運算重疊」的角度比較兩者效能差異，並指出 single-pass 為何不需 grid-wide barrier、又有什麼新風險。
三-kernel：K1 把各 scan block 結果寫進 global memory 的 S 陣列，K3 再讀回。由於 kernel 邊界是隱式 global barrier，這些 store/load 延遲無法與後續運算重疊，對大型輸入顯著拖慢。
Single-pass / domino：單一 kernel，partial sum 在相鄰 block 間骨牌式單向傳遞；關鍵觀察是全域 scan 步驟只需 adjacent (相鄰) 同步、不需 grid-wide barrier。其 flags/scan_value 存取多落在 L2 cache，且能與其他 block 的 phase 1 / phase 3 運算重疊 → 省去無法重疊的往返。
新風險：需自製 adjacent synchronization（flags + __threadfence + atomicAdd）；且 GPU 不保證 block 依 blockIdx 線性排程，可能 deadlock，須用 dynamic block index assignment（atomicAdd(blockCounter,1) 取得 bid，保證 i−1 已排程）來避免。
| 演算法 / 技術 | 加法次數 (Work) | 複雜度 | 步數 (P=N) | 定位 / 關鍵取捨 |
|---|---|---|---|---|
| Sequential scan | N − 1 | O(N) | N | work-efficiency 基準 |
| Naïve parallel | N(N−1)/2 | O(N²) | N（最長那條） | 又慢又耗資源，沒共享 partial sum |
| Kogge-Stone | N·log₂N − (N−1) | O(N log N) | log₂N | 快、步數少；資源充足下的中等 section（512/1024） |
| Brent-Kung | 2N − 2 − log₂N ≈ 2N−3 | O(N) | ~2×log₂N | work-efficient、data-scalable；步數較多 |
| Thread-coarsened | (N−T)+(T·log₂T)+(N−T) | O(N) | (2(N−T)+T·log₂T)/P |
串行子段 + 並行整合，硬體利用率最佳；section 受 shared memory 限 |
| Hierarchical (3-kernel) | — | — | — | 任意長度；但 S 陣列往返無法與運算重疊 |
| Single-pass (domino) | — | — | — | 單 kernel、骨牌傳遞、traffic 落 L2 可重疊；需 adjacent sync + 防 deadlock |
核心心法：scan 是「平行演算法 work 可能高於 sequential」的典型案例 —— Kogge-Stone 用多做工換少步數，Brent-Kung 用多步數換 work efficiency，coarsening 把多數工作還給最有效率的 sequential scan；任意長度則靠 hierarchical / single-pass 整合，後者用 adjacent synchronization 換取記憶體存取效率。