優化單一 Block 的 Reduction Kernel (Control/Memory Divergence 與 Shared Memory)

重點總覽 (Overview)

本筆記針對 simple sum reduction kernel (Fig. 10.6) 做三個連續優化，每一步只改幾行卻有顯著效能差異。基準輸入皆以 N = 256 (即 128 threads = 4 warps) 為例。

降低 Control Divergence (Minimizing Control Divergence, 10.4)

問題：active thread 隨時間越離越遠

Simple kernel 把 thread 對應到 i = 2*threadIdx.x (偶數位置)，條件是 threadIdx.x % stride == 0。隨 stride 遞增，留下的 active thread 在 warp 內越分越散，使每個 warp 同時混有 active / inactive thread → divergence。

Simple (Fig.10.7): stride 遞增, active thread 越離越遠
in[]: 0  1  2  3  4  5  6  7  ...
T:    0     1     2     3        (i = 2*tid)
s=1:  +1    +1    +1    +1   每個 active thread 與鄰居間隔大
s=2:  T0          T2           (只剩 tid 為 2 的倍數 → warp 內穿插 divergence)
s=4:  T0                       (只剩 tid 為 4 的倍數)

解法：讓 active thread 保持「相鄰」(stride 遞減)

改成 i = threadIdx.x、stride 從 blockDim.x 遞減，條件 threadIdx.x < stride。如此 active thread 永遠是 thread 0..stride-1 的連續區段，整段 warp 要嘛全 active、要嘛全 inactive。

Convergent (Fig.10.8, N=16, 8 threads): stride 遞減, active thread 永遠相鄰
in[]:  0  1  2  3  4  5  6  7 | 8  9 10 11 12 13 14 15
s=8:   T0 T1 T2 T3 T4 T5 T6 T7  in[t]+=in[t+8]   (全部 active)
s=4:   T0 T1 T2 T3 .  .  .  .   in[t]+=in[t+4]
s=2:   T0 T1 .  .                in[t]+=in[t+2]
s=1:   T0 .                      in[0]+=in[1]  → 結果在 in[0]

// Fig.10.9  Convergent kernel — 低 control divergence + coalesced
__global__ void ConvergentSumReductionKernel(float* input, float* output) {
    unsigned int i = threadIdx.x;                 // 改 1: 不再乘 2，owner 位置相鄰
    for (unsigned int stride = blockDim.x; stride >= 1; stride /= 2) {  // 改 2: 遞減
        if (threadIdx.x < stride) {               // 改 3: 連續區段的 thread 才 active
            input[i] += input[i + stride];
        }
        __syncthreads();
    }
    if (threadIdx.x == 0) *output = input[0];
}

量化：execution resource utilization efficiency

分母關鍵：一個 warp 只要有任何一個 active thread，就吃掉整 32 個 thread 的 resource。

• N/64 = 啟動的 warp 數 (N/2 threads ÷ 32)。Simple 前 5 次迭代所有 warp 都活著 (warp 內 ≥1 active)，故 ×5。
• Convergent：整個 warp 會一起退場，前段每次迭代 active warp 減半，直到剩 1 個 warp 撐完最後 5 次迭代 (+5·1)。

降低 Memory Divergence (Minimizing Memory Divergence, 10.5)

Memory divergence = warp 內相鄰 thread 沒有存取相鄰位址 → 無法 coalesce，一次 warp 存取被拆成多個 global memory requests。

每次迭代每個 active thread 做 2 reads + 1 write (讀 owner、讀 stride 外、寫回 owner)。

Simple kernel: 相鄰 thread 的 owner 位址相隔 2 (= 2*tid) → 非 coalesced
warp 第一次 read:  in[0] in[2] in[4] ...  位址間隔 2
  → 觸發 2 個 memory requests，回傳資料有一半沒用到

Convergent kernel: 相鄰 thread 存取相鄰位址 (i = tid) → 全 coalesced
warp read:  in[0] in[1] in[2] ... in[31]  連續
  → 每個 read/write 只觸發 1 個 memory request

量化：global memory requests (N=256，×3 代表 2讀+1寫)

• Simple 的 ·2：前 5 次迭代每個 warp 因非 coalesced 而觸發 2 個 request。
• 比值 141/36 ≈ 3.9×。在 N=2048 時為 1149 vs 204 ≈ 5.6× (warp 越多、初期非 coalesced 浪費越大)。

降低 Global Memory 存取 (Minimizing Global Memory Accesses, 10.6)

觀察：convergent kernel 每次迭代仍把 partial sum 寫回 global memory，下一次又讀回來。由於 shared memory 延遲更低、頻寬更高，可把 partial sum 全程留在 shared memory。

Fig.10.10: 第一次加法直接在 load 時於 global memory 完成，
之後所有迭代都只在 shared memory input_s[] 進行：

global in[]: [.... 2*BLOCK_DIM 個原始元素 ....]
   每個 thread t:  input_s[t] = in[t] + in[t + BLOCK_DIM]   (coalesced, 一次做完第一層)
                          │
                          ▼  (以下只動 shared memory)
   shared input_s[]:  reduction tree (stride 從 blockDim/2 遞減)
                          │
                          ▼
   thread 0:  output = input_s[0]   (唯一寫回 global)

// Fig.10.11  Shared-memory kernel
#define BLOCK_DIM 1024
__global__ void SharedMemorySumReductionKernel(float* input, float* output) {
    __shared__ float input_s[BLOCK_DIM];
    unsigned int t = threadIdx.x;
    input_s[t] = input[t] + input[t + BLOCK_DIM];   // 第一層在 global 上做掉, 兩個 read 皆 coalesced
    for (unsigned int stride = blockDim.x/2; stride >= 1; stride /= 2) { // 從 blockDim/2 起
        __syncthreads();                            // 移到迴圈開頭: 同步 shared 寫入與下一輪讀取
        if (threadIdx.x < stride) {
            input_s[t] += input_s[t + stride];      // 只動 shared memory
        }
    }
    if (threadIdx.x == 0) *output = input_s[0];      // 唯一寫回 global
}

量化：global memory accesses / requests

• 對比 convergent 的 36 → 9，再 ~4× 改善。
• __syncthreads() 必須移到迴圈開頭：確保第一層 (迴圈外) 的 shared memory 寫入，與第一次迭代的讀取之間有 barrier synchronization。

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

Related Notes

• 10-Reduction/01-Reduction-Fundamentals-and-Simple-Kernel
• 10-Reduction/03-Scaling-Reduction-Hierarchical-and-Coarsening
• 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence
• 06-Performance-Considerations/01-Memory-Coalescing
• 05-Memory-Architecture-And-Data-Locality/02-Tiling-and-Tiled-Matrix-Multiplication
• 06-Performance-Considerations/02-Hiding-Memory-Latency