擴展 Reduction：Hierarchical Reduction 與 Thread Coarsening

重點總覽 (Overview)

階層式分段多 Block Reduction (Hierarchical Segmented Multiblock Reduction)

為何單 block 不夠

• 之前所有 kernel 都假設 launch 一個 block,因為用 __syncthreads() 當 barrier。
• __syncthreads() 只能同步同一個 block 內的 threads → 平行度上限 = 1024 threads = 2048 元素。
• 百萬/十億級輸入需要更多 thread → 必須讓不同 block 獨立執行(因為跨 block 沒有好的 barrier)。

分段策略

• 把輸入切成多個 segment,每段大小 = 2 * blockDim.x(一個 block 處理的量)。
• 每個 block 在自己的 segment 上跑「整棵」reduction tree,就好像 segment 是整個輸入。
• 每個 block 最終的 partial sum (input_s[0]) 用 atomicAdd 累加到單一 output。

segment size = 2 * blockDim.x   (blockDim.x = 1024 → 2048)

global input:  [ ---- block 0 ---- ][ ---- block 1 ---- ][ ---- block 2 ---- ] ...
segment start:  0                    2048                 4096
                 \  reduce tree /     \  reduce tree /     \  reduce tree /
                  input_s[0]            input_s[0]            input_s[0]
                       \                    |                    /
                        \____ atomicAdd ___ output ____________/

Kernel (Fig 10.13)

__global__ void SegmentedSumReductionKernel(float* input, float* output) {
    __shared__ float input_s[BLOCK_DIM];               // BLOCK_DIM == blockDim.x
    unsigned int segment = 2 * blockDim.x * blockIdx.x; // 此 block 的段起點 (line 03)
    unsigned int i = segment + threadIdx.x;             // thread 在全域陣列的擁有位置 (line 04)
    unsigned int t = threadIdx.x;                       // thread 在 shared 陣列的位置
    input_s[t] = input[i] + input[i + BLOCK_DIM];       // 第一步:用全域 index i 載入並相加 (line 06)
    for (unsigned int stride = blockDim.x/2; stride >= 1; stride /= 2) {
        __syncthreads();
        if (t < stride) {
            input_s[t] += input_s[t + stride];          // reduction tree (與 Fig 10.11 完全相同)
        }
    }
    if (t == 0) {
        atomicAdd(output, input_s[0]);                  // 把段的 partial sum 原子累加 (line 14)
    }
}

• 關鍵差異 vs 單 block 版:(1) 用 i (全域、含 segment 位移) 讀 global memory;(2) tree 迴圈內仍用 t (因每個 block 有自己私有的 input_s);(3) 結尾改成 atomicAdd 而非直接寫。
• i 是「全域擁有位置」,t 是「shared 內擁有位置」── 兩者分工是這支 kernel 的核心。

執行緒粗化降低 Overhead (Thread Coarsening for Reduced Overhead)

動機:我們為平行付了「重稅」

• Reduction tree 每往上一層,就有更多 warp 閒置、最後一個 warp control divergence 更嚴重 → hardware underutilization。
• 這個「逐漸閒置」的階段,每一個 launch 的 block 都要付一次。
• 若硬體資源不足,多餘的 block 本來就會被硬體序列化執行(跑完一個換下一個)。
• 既然要被序列化,不如自己用更有效率的方式序列化 → 這就是 thread coarsening(把工作塞進較少 thread 以減少平行化 overhead)。

做法:每個 block 多吃資料,前段序列相加

• 把 block 粗化 COARSE_FACTOR 倍 → 每個 block 拿 COARSE_FACTOR * 2 * blockDim.x 個元素。
• 例:原本每 block 16 元素(每 thread 2 個);coarsen ×2 → 每 block 32 元素(每 thread 4 個)。
• 每個 thread 獨立、序列地把它負責的 COARSE_FACTOR*2 個元素加進 register sum:
  • 此階段全部 thread 都活躍(硬體全利用)。
  • 不需要 __syncthreads()、不需要寫 shared memory(thread 彼此獨立)。
• 加總完才把 sum 寫進 input_s[t],接著進入和之前完全相同的 reduction tree。

COARSE_FACTOR = 2, blockDim.x = 8, 32 elements/block, 每 thread 4 元素

序列粗化階段 (全 thread 活躍, 無 sync, 無 shared):
 step 1   sum  = input[i]
 step 2   sum += input[i + 1*BLOCK_DIM]
 step 3   sum += input[i + 2*BLOCK_DIM]
          sum += input[i + 3*BLOCK_DIM]   (寫入 input_s[t])
            |  全 8 threads 活躍 → 硬體全利用
            v
reduction tree 階段 (每步半數 thread drop, 需 sync + shared):
 step 4   stride=4  ▓▓▓▓░░░░
 step 5   stride=2  ▓▓░░
 step 6   stride=1  ▓░

Kernel (Fig 10.15)

__global__ void CoarsenedSumReductionKernel(float* input, float* output) {
    __shared__ float input_s[BLOCK_DIM];
    // 段更大 COARSE_FACTOR 倍 (line 03)
    unsigned int segment = COARSE_FACTOR * 2 * blockDim.x * blockIdx.x;
    unsigned int i = segment + threadIdx.x;
    unsigned int t = threadIdx.x;
    float sum = input[i];                               // 粗化迴圈:序列相加 (lines 06-09)
    for (unsigned int tile = 1; tile < COARSE_FACTOR * 2; ++tile) {
        sum += input[i + tile * BLOCK_DIM];             // 全 thread 活躍、無 __syncthreads()
    }
    input_s[t] = sum;                                   // 加總完才入 shared memory
    for (unsigned int stride = blockDim.x/2; stride >= 1; stride /= 2) {
        __syncthreads();
        if (t < stride) {
            input_s[t] += input_s[t + stride];
        }
    }
    if (t == 0) {
        atomicAdd(output, input_s[0]);
    }
}

為何省:2 序列 block vs 1 coarsened block (Fig 10.16)

• coarsening 同時砍掉三種 overhead:hardware underutilization、barrier synchronization、shared memory access。
• 全利用階段在 register 上做、且無需同步 → 比 tree 階段便宜得多。

本章總結 (Chapter Summary)

平行 reduction 看似簡單(sequential 版只是一個 for-loop),但要對大型輸入達到高效能,需要層層疊加的技巧:

• Reduction 是許多資料彙整應用的核心,也是 prefix-sum (scan) 的重要基礎 → 延伸至 11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone。

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

Related Notes

• 10-Reduction/01-Reduction-Fundamentals-and-Simple-Kernel
• 10-Reduction/02-Optimizing-Single-Block-Reduction-Kernel
• 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram
• 09-Parallel-Histogram/02-Histogram-Optimizations-Privatization-Coarsening-Aggregation
• 06-Performance-Considerations/03-Thread-Coarsening
• 11-Prefix-Sum-Scan/03-Arbitrary-Length-and-Single-Pass-Scan