任意長度輸入與 Single-Pass Scan

重點總覽 (Overview)

前兩篇 (11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone、11-Prefix-Sum-Scan/02-Work-Efficient-Scan-Brent-Kung-and-Coarsening) 的 kernel 都只能對「一個 block 能裝下的 section」做 scan。本篇處理百萬∼數十億元素的任意長度輸入，並把 segment 結果整合成全域結果。

階層式分段 Scan (Hierarchical / Segmented Scan)

對大型資料集:先把輸入切成數個 scan block,每個剛好塞進一個 SM 的 shared memory,由單一 block 處理。執行三個邏輯步驟。

三步驟流程

輸入 X (16 元素, 切成 4 個 scan block, 每塊 4 元素)
  [2 1 3 1 | 0 4 1 2 | 0 3 1 2 | 2 6 1 2]

Step 1  各 block 局部 inclusive scan (彼此獨立, 不含跨塊貢獻)
  blk0: 2 3 6 7    blk1: 0 4 5 7    blk2: 0 3 4 6    blk3: 2 8 9 11
                ^最後一格 = 該塊總和

Step 2  蒐集各塊「最後一格」成 S 陣列, 再對 S 做一次 scan
  S  = [ 7 , 7 , 6 , 11 ]          (各塊總和)
  S' = [ 7 , 14, 20, 31 ]          (scan 後 = 全域策略點的最終結果)
        |    |    |    |
        X[3] X[7] X[11] X[15] 的最終 scan 值

Step 3  把 S'[blockIdx-1] 加回該塊每個元素 (blk0 不需加)
  blk0:        2  3  6  7
  blk1: +7  -> 7 11 12 14
  blk2: +14 -> 14 17 18 20
  blk3: +20 -> 22 28 29 31

• S 陣列維度 = N / SECTION_SIZE。S[i] 是「從 X[0] 到 scan block i 結尾」的累積和,即原問題在策略性位置 (各塊末端) 的最終 scan 值。
• 這與第二篇 thread-coarsening 的三-phase scan (11-Prefix-Sum-Scan/02-Work-Efficient-Scan-Brent-Kung-and-Coarsening) 邏輯相同,差別在 Step 2 跨越多個 thread block,故各塊末端值必須寫進 global memory 才能跨塊可見。
• 原理等同硬體加法器的 carry look-ahead(進位先行):先算各段進位,再廣播回去。

三-Kernel 實作

// K1 結尾:由 block 內最後一個 thread 把該塊總和寫進 S
__syncthreads();
if (threadIdx.x == blockDim.x - 1) {
    S[blockIdx.x] = XY[SECTION_SIZE - 1];   // = 本 scan block 的所有元素總和
}

// K2:對 S 做 scan(單一 block 即可,因元素數 = block 數,通常很小)
//     直接呼叫 Kogge_Stone_scan_kernel / Brent_Kung_scan_kernel(S, S, N/SECTION_SIZE)

// K3:把前面所有 scan block 的總和加回本塊
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
if (blockIdx.x > 0) {
    Y[i] += S[blockIdx.x - 1];              // blk0 不需加,blockIdx==0 跳過
}

Single-Pass (Domino-style) Scan

又稱 stream-based scan(注意:與 Chapter 20 的 CUDA streams 無關)。用單一 kernel 完成三個步驟,partial sum 在相鄰 block 之間透過 global memory 單向傳遞,像骨牌一樣。

骨牌傳遞 (Producer-Consumer Chain)

時間 →

phase 1 (全部 block 並行做局部 scan)
  blk0      blk1      blk2      blk3
   |         |         |         |
phase 2 (傳遞 partial sum, 序列化)
  blk0 --7--> blk1 --14--> blk2 --20--> blk3
  (blk_i 等左鄰 blk_{i-1} 的和, 加上自身 local_sum, 再傳給右鄰 blk_{i+1})
   |         |         |         |
phase 3 (收到前綴和後即可並行加回, 與其他塊重疊)
  blk0      blk1      blk2      blk3

• block i 流程:① 先對自己的 scan block 做局部 scan(與所有 block 並行)→ ② 等待左鄰 block i-1 傳來累積和 → ③ 收到後加上自己的 local sum,把新的累積和傳給右鄰 block i+1 → ④ 再把收到的值加到自己所有 partial scan 值上,產生最終輸出。
• phase 1 全並行;phase 2 傳遞被序列化;但只要和值傳得夠快,phase 3 又能在各 block 間取得大量並行。

Adjacent (Block) Synchronization

骨牌要能運作,需要 adjacent synchronization —— 一種讓相鄰 block 同步並交換資料的客製化同步,以 atomic 操作實作。只由每個 block 的 leader thread (例如 threadIdx.x == 0) 執行,其餘 thread 在最後的 __syncthreads() 等候。

__shared__ float previous_sum;
if (threadIdx.x == 0) {
    // consumer 端:反覆檢查 flag,直到 producer (block bid-1) 設好
    while (atomicAdd(&flags[bid], 0) == 0) { }   // busy-wait 直到旗標被設
    previous_sum = scan_value[bid];              // 載入前驅傳來的 partial sum
    scan_value[bid + 1] = previous_sum + local_sum;  // 累積後傳給後繼 (producer 端)
    __threadfence();                             // 確保 scan_value 先寫進 global memory
    atomicAdd(&flags[bid + 1], 1);               // 再設後繼的旗標 -> 解除其等待
}
__syncthreads();   // 其餘 thread 在此等 leader 完成

Dynamic Block Index Assignment(防 Deadlock)

解法:dynamic block index assignment —— 讓 bid 不再綁 blockIdx.x,改由原子計數器決定執行順序。

__shared__ unsigned int bid_s;
if (threadIdx.x == 0) {
    bid_s = atomicAdd(blockCounter, 1);   // 原子遞增全域計數器,取得動態索引
}
__syncthreads();
unsigned int bid = bid_s;                 // 全 block 共用此動態 bid

• leader thread 原子遞增 blockCounter,先執行 atomic 的 block 才拿到較小的 bid。
• 保證:若某 block 拿到 bid = i,則 bid = i-1 的 block 必定已被排程(因為它已執行過 atomic)→ 等待對象一定存在 → 不會死鎖,所有 scan block 邏輯上線性排程。

本章總結 (Chapter Summary)

• Scan 把「以數學遞迴描述的看似序列」運算轉成並行,削減許多應用的序列瓶頸(資源/工作分配、polynomial、recurrence、radix sort、stream compaction、tree 操作…)。
• Hierarchical scan 把上述 section 級 kernel 擴展到任意長度;但三-kernel 版本有無法重疊的冗餘 global memory 往返。
• Domino-style single-pass scan 以單 kernel + adjacent synchronization 消除往返、提升 global memory 存取效率,但需謹慎處理 atomic 旗標、__threadfence、__syncthreads,並用 dynamic block index 防死鎖。
• 進一步可用 warp-level shuffle 再優化。實務上一般人更可能直接用 Thrust 等 GPU scan library,而非自己手刻。

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone
• 11-Prefix-Sum-Scan/02-Work-Efficient-Scan-Brent-Kung-and-Coarsening
• 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram
• 10-Reduction/03-Scaling-Reduction-Hierarchical-and-Coarsening
• 04-Compute-Architecture-And-Scheduling/01-GPU-Architecture-and-Block-Scheduling
• 13-Sorting/01-Sorting-Foundations-and-Parallel-Radix-Sort