Enter your search text in the box above

Select a result to preview

高 Work-Efficiency 的 Scan：Brent-Kung 與 Thread Coarsening

#parallel-pattern #scan #brent-kung #thread-coarsening #work-efficiency

重點總覽 (Overview)

項目 Brent-Kung Scan Three-Phase Coarsened Scan
核心結構 Reduction tree (向上歸約) + reverse distribution tree (向下分配) Phase 1 sequential scan → Phase 2 block scan → Phase 3 加總
操作數 (N 元素) 2N - 2 - log2(N) = O(N) phase1+phase3 = 2(N-T)，phase2 = T·log2(T)
Threads 數 SECTION_SIZE/2（每 thread 載入 2 個元素） T（可遠小於元素數 N）
Section 上限 2048（1024 threads × 2） 受限於 shared memory 大小，非 thread 數
Work efficiency 最多為 sequential 的 2 倍 → data-scalable 最高（多數工作以 sequential 完成）
步數 (steps) 比 Kogge-Stone 多（兩階段樹） 居中，最有效率利用執行單元
主要缺點 步數較長；inactive threads 仍被 SIMD 綁住消耗資源 Phase 2 仍是低效率的 parallel scan
Important

核心取捨：Brent-Kung 用「略增的演算法複雜度（兩棵樹、更多步數）」換取「線性 O(N) 的 work efficiency」。當執行資源 (P) 充足時 Kogge-Stone 較快；當資源有限或重視能耗時 Brent-Kung 勝出。Coarsening 進一步把大部分工作交給最有效率的 sequential scan。

Brent-Kung 演算法 (Brent-Kung Algorithm)

延續 11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone 的問題：Kogge-Stone 雖簡單，但 work efficiency 偏低（O(N log N)）。Brent-Kung 透過兩個對稱的樹狀階段達成 O(N)。

階段一：Reduction Tree（歸約樹，向上）

16 元素 reduction tree（數字為被更新的 XY index）

stride=1 : 1  3  5  7  9 11 13 15      (8 ops)  XY[i] += XY[i-1]，i 為奇數
stride=2 :    3     7    11    15      (4 ops)  index 形如 4n-1
stride=4 :          7          15      (2 ops)  index 形如 8n-1
stride=8 :                     15      (1 op )  只剩 XY[15] = 全段總和

總操作數 = 8+4+2+1 = 15 = N-1

階段二：Reverse Distribution Tree（反向分配樹，向下）

16 元素 reverse tree（箭頭 = 被加進去的目標位置）

stride=4 : XY[7] ──► XY[11]                              (1 op)
stride=2 : XY[3]►XY[5]  XY[7]►XY[9]  XY[11]►XY[13]       (3 ops)
stride=1 : XY[1]►XY[2] XY[3]►XY[4] XY[5]►XY[6] XY[7]►XY[8]
           XY[9]►XY[10] XY[11]►XY[12] XY[13]►XY[14]      (7 ops)

XY[14] 需要 3 個 partial sums：XY[13](x12..x13)、XY[11](x8..x11)、XY[7](x0..x7)
反向階段操作數 = (2-1)+(4-1)+...+(N/2-1) = N - 1 - log2(N)

Kernel 實作 (Fig. 11.7)

__global__ void Brent_Kung_scan_kernel(float *X, float *Y, unsigned int N) {
    __shared__ float XY[SECTION_SIZE];
    unsigned int i = 2*blockIdx.x*blockDim.x + threadIdx.x;   // 每 block 處理 2*blockDim.x 個元素
    if (i < N)                XY[threadIdx.x]              = X[i];
    if (i + blockDim.x < N)   XY[threadIdx.x + blockDim.x] = X[i + blockDim.x];

    // ---- Reduction tree phase ----
    for (unsigned int stride = 1; stride <= blockDim.x; stride *= 2) {
        __syncthreads();
        unsigned int index = (threadIdx.x + 1)*2*stride - 1;  // 連續 threads → 間隔位置
        if (index < SECTION_SIZE)
            XY[index] += XY[index - stride];
    }
    // ---- Reverse distribution tree phase ----
    for (int stride = SECTION_SIZE/4; stride > 0; stride /= 2) {
        __syncthreads();
        unsigned int index = (threadIdx.x + 1)*stride*2 - 1;
        if (index + stride < SECTION_SIZE)
            XY[index + stride] += XY[index];
    }
    __syncthreads();
    if (i < N)               Y[i]              = XY[threadIdx.x];
    if (i + blockDim.x < N)  Y[i + blockDim.x] = XY[threadIdx.x + blockDim.x];
}
Tip

關鍵的 thread→data 映射 index = (threadIdx.x + 1)*2*stride - 1：讓從 thread 0 開始的一段連續 threads 對應到間隔為 stride 的資料位置。這樣每次迭代用的都是連續 thread，control divergence 只在 active threads 少於一個 warp（32）時才發生，避開了「以 2^n 為 index」那種跳號寫法的嚴重 divergence。

Warning

為何只需 SECTION_SIZE/2 個 threads？因為 reduction/reverse 兩階段中最多同時有 N/2 個位置在做加法。因此 block 用 SECTION_SIZE/2 threads，但每個 thread 負責載入 2 個 X 元素、寫回 2 個 Y 元素（見 kernel 中兩組 if）。1024 threads → section 最大 2048 元素。

Work Efficiency 分析

演算法 操作數 複雜度
Sequential N - 1 O(N)
Naïve parallel N(N-1)/2 O(N²)
Kogge-Stone N·log2(N) - (N-1) O(N log N)
Brent-Kung (N-1) + (N-1-log2(N)) = 2N - 2 - log2(N) O(N)
Warning

理論 vs CUDA 現實：Brent-Kung 的 work-efficiency 優勢在實機上被稀釋。active threads 數雖然下降得比 Kogge-Stone 快，但inactive threads 因 SIMD 被綁在同一 warp，仍會消耗執行資源直到整個 warp 結束。

Speed 比較（限定資源範例）

以 1024 元素、32 個 execution units 為例：

Kogge-Stone Brent-Kung
估計步數 (1024·10)/32 = 320 (2·1024 - 2 - 10)/32 ≈ 63.6，加 ~5 步 divergence ≈ 73.6
對 sequential 的加速 1024/320 = 3.2× 1024/73.6 ≈ 14×
Important

直覺上 Brent-Kung「更有效率卻更慢」聽起來矛盾，關鍵在 P（執行單元數）：

  • 資源有限（P 小）→ 總操作數主導，Brent-Kung 較快（如上 14× vs 3.2×）。
  • 資源無限（P ≥ N）→ 步數主導，Brent-Kung 約需 Kogge-Stone 的 2 倍步數（多了 reverse tree），反而較慢。

Thread Coarsening 的三階段 Scan (Coarsening for Work Efficiency)

scan 的平行化除了硬體未充分利用、同步開銷外，還多了一項：work efficiency 變差。若硬體實際上會把 threads 序列化執行，不如我們自己用 coarsening 序列化，把工作換成最有效率的 sequential scan。參見 06-Performance-Considerations/03-Thread-Coarsening 的通則與 10-Reduction/03-Scaling-Reduction-Hierarchical-and-Coarsening 的 reduction 版本。

16 元素、4 threads 的三階段（每 subsection 4 元素）

輸入 :  T0[2 1 3 1]  T1[0 4 1 2]  T2[0 3 1 2]  T3[1 4 2 4]

Phase 1 (各 thread 內 sequential scan，O(N) 最有效率)
        T0[2 3 6 7]  T1[0 4 5 7]  T2[0 3 4 6]  T3[1 5 7 11]
                  ▲           ▲           ▲            ▲
        每段最後元素 = 該段總和 → [7 , 7 , 6 , 11]

Phase 2 (對「各段最後元素」做 block-wide parallel scan，Kogge-Stone/Brent-Kung)
        [7 7 6 11]  ──►  [7 14 20 31]   (exclusive 前綴提供給下一段)

Phase 3 (每 thread 把「前一段的掃描結果」加到自己段，最後元素不需更新)
        T0[2 3 6 7]
        T1: +7  → [7 11 12 (14)]      ← (14) 已正確，不動
        T2: +14 → [14 17 18 (20)]
        T3: +20 → [21 25 27 (31)]

三個階段

Phase 動作 重點
1 各 thread 對自己的 subsection 做 sequential scan 段末元素 = 整段總和；最有效率
2 所有 thread 協作對「各段末元素組成的邏輯陣列」做 parallel scan 用 Kogge-Stone 或 Brent-Kung；元素間距為 stride（subsection 大小），thread→element 映射需調整
3 每 thread 把前驅段的掃描值加進自己段 各段最後元素已是正確值，不需更新

Memory Coalescing 技巧

Warning

若 Phase 1 直接從 global memory 讀各自 subsection，存取不會 coalesced（thread 0 讀 element 0、thread 1 讀 element 4…，間隔太大）。

Tip

三階段最大優勢：threads 數可遠小於元素數。Section 大小不再被 block 的 thread 上限綁住，而是受 shared memory 容量限制（整段需放得進 shared memory）。

Work 分析

對 N 元素、T threads、P 執行單元（Phase 2 用 Kogge-Stone）：

Phase Work
1 (sequential scan) N - T
2 (block scan) T·log2(T)
3 (加總) N - T
steps≈(N−T)+Tlog2⁡T+(N−T)P

範例：N=1024, T=64, P=32 →
(1024-64 + 64·6 + 1024-64)/32 = 72 steps（對照單純 Kogge-Stone 的 320 steps）。

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

情境 / 關鍵字 答案 / 技巧
「Brent-Kung 操作數？」 reduction N-1 + reverse N-1-log2(N) = 2N-2-log2(N)，O(N)
「為什麼是 O(N) 而 Kogge-Stone 是 O(N log N)？」 兩棵樹各做近 N 次加法，沒有像 Kogge-Stone 那樣每步都讓近 N 個 thread 重算
「Brent-Kung 用幾個 threads？section 上限？」 SECTION_SIZE/2；每 thread 載/寫 2 元素；1024 threads → 2048 元素
「reduction phase 的 thread→data 映射」 index = (threadIdx.x+1)*2*stride - 1，連續 threads 對應間隔位置以避免 divergence
「reverse phase 在做什麼？」 把 partial sums 往右 push：XY[index+stride] += XY[index]，stride 由 SECTION_SIZE/4 遞減到 1
「Brent-Kung 一定比 Kogge-Stone 快嗎？」 否。資源無限時 BK 約需 2× 步數較慢；資源有限時 BK 因操作數少而較快（14× vs 3.2×）
「為何 work-efficiency 優勢在 GPU 被稀釋？」 inactive threads 受 SIMD 綁定，仍佔 warp 執行資源
「三階段 coarsening 各階段做什麼？」 P1 各 thread sequential scan；P2 對段末元素 block scan；P3 加前驅段掃描值
「為何 P1 要先進 shared memory？」 直接讀 global 不 coalesced；用 coalesced 載入 shared、再做不友善樣式存取
「coarsening section 上限由什麼決定？」 shared memory 容量，不再是 thread 數上限
「三階段 steps 公式」 (2(N-T) + T·log2 T)/P，例 N=1024,T=64,P=32 → 72
「data-scalable algorithm」 操作數隨輸入線性成長者（Brent-Kung 即是）