Reduction 基礎與簡單 Kernel

重點總覽 (Overview)

背景:Reduction 的定義 (Background)

• Reduction = 從一個 array of values 導出單一值(sum / max / min / product / ...),輸入型別可為 int、float、double、half、char,計算結構皆相同。
• 數學上:給定一個有 identity value 的 binary operator 即可定義 reduction。identity 指「v ⊕ identity = v」的那個值。
• 與 histogram 一樣,reduction 是把大量資料壓成摘要的重要 pattern;平行版需要 threads 互相協調 (coordination),是展示效能瓶頸與其解法的好教材。

Sequential reduction(O(N),N 次迭代,N 步):

// Fig. 10.1 — sum 版本
float sum = 0.0f;                       // 初始化為 identity
for (unsigned int i = 0; i < N; ++i)
    sum += input[i];                    // 逐一累加

// Fig. 10.2 — 通用形式(可換成 max / min / product)
acc = IDENTITY;
for (unsigned int i = 0; i < N; ++i)
    acc = Operator(acc, input[i]);      // max 回傳較大、min 回傳較小 ...

Reduction Tree 平行模型 (Reduction Trees)

把 sequential 的一條長鏈,重新加括號成一棵樹:每個 time step 同時對所有「pair」做運算,逐層折半收斂到 root。葉子是原始輸入,root 是最終結果(注意:這是概念上的樹,edge 只代表 time step 間的資訊流,不是指標型 tree 資料結構)。

8 個輸入的 sum reduction tree(time 由上往下)
index:    0    1    2    3    4    5    6    7
         [3]  [1]  [7]  [0]  [4]  [1]  [6]  [3]
           \  /      \  /      \  /      \  /
step 1:    [4]       [7]       [5]       [9]      ← N/2 = 4 次運算 (peak)
             \       /           \       /
step 2:       [ 11 ]              [ 14 ]          ← N/4 = 2 次運算
                 \               /
step 3:             [   25   ]                    ← 1 次運算 = root

順序被改變了:sequential 是 ((((((3 max 1) max 7) max 0)...) 一路左結合;tree 則是 ((3⊕1)⊕(7⊕0)) ⊕ ((4⊕1)⊕(6⊕3))——同一份運算清單,只是括號位置不同。

兩個必要性質

• max / min / sum / product 皆同時 associative 且 commutative。
• 把 sequential → tree(Fig 10.7)只需 associative;若要像 10-Reduction/02-Optimizing-Single-Block-Reduction-Kernel 那樣重排 operand(讓 active threads 靠在一起),則還需 commutative。

工作量、步數與平行度

Simple Reduction Kernel (A Simple Reduction Kernel)

因為 reduction tree 需要所有 threads 協作,而跨整個 grid 無法做 barrier,所以第一版只在單一 block 內做。對 N 個元素,launch 一個 block、N/2 threads。

• 一個 block 最多 1024 threads → simple kernel 最多處理 2048 個 elements(§10.8 才解除此限制)。
• 採 "owner computes":每個資料位置由唯一一個 thread「擁有」,且只有 owner 會寫入它。thread k 的 owner 位置是 i = 2*threadIdx.x(即所有偶數索引)。
• stride 從 1 開始每輪 ×2(1,2,4,...);第 n 輪只有 threadIdx.x % stride == 0(即 index 為 2ⁿ 倍數)的 thread 為 active,逐輪減半,最後只剩 thread 0。

// Fig. 10.6 — 單一 block 的 simple sum reduction kernel
__global__ void SimpleSumReductionKernel(float* input, float* output) {
    unsigned int i = 2 * threadIdx.x;                      // owner = 偶數索引
    for (unsigned int stride = 1; stride <= blockDim.x; stride *= 2) {
        if (threadIdx.x % stride == 0) {                   // active = 2^n 的倍數
            input[i] += input[i + stride];                 // 讀 owner + 讀 stride 外 + 寫 owner
        }
        __syncthreads();                                   // 等全部 partial sum 寫完才進下一輪
    }
    if (threadIdx.x == 0) {
        *output = input[0];                                // thread 0 輸出最終結果
    }
}

Thread → Data 映射與執行(N=8, 4 threads)

thread:   T0       T1       T2       T3          (T_k owns input[2k])
index:    0   1    2   3    4   5    6   7
         [a] [b]  [c] [d]  [e] [f]  [g] [h]
stride=1  T0:0+=1  T1:2+=3  T2:4+=5  T3:6+=7      ← 4 threads 全 active,相鄰 pair
         [ab]     [cd]     [ef]     [gh]          (寫回 idx 0,2,4,6)
            __syncthreads()
stride=2  T0:0+=2           T2:4+=6               ← 只 T0,T2 active;T1,T3 idle
         [abcd]            [efgh]                 (寫回 idx 0,4)
            __syncthreads()
stride=4  T0:0+=4                                 ← 只 T0 active
         [abcdefgh]                               (寫回 idx 0)
            __syncthreads()
output = input[0]            (T0)

• 每個 active thread 每輪做 2 reads + 1 write:讀自己的 owner、讀距離 stride 外的元素、寫回 owner。
• __syncthreads() 確保「本輪所有 partial sum 都已寫進 input array」後,下一輪的 active threads 才會去讀(避免 race)。詳見 04-Compute-Architecture-And-Scheduling/01-GPU-Architecture-and-Block-Scheduling 的 barrier。

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

Related Notes

• 10-Reduction/02-Optimizing-Single-Block-Reduction-Kernel
• 10-Reduction/03-Scaling-Reduction-Hierarchical-and-Coarsening
• 04-Compute-Architecture-And-Scheduling/01-GPU-Architecture-and-Block-Scheduling
• 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence
• 11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone
• 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram