直方圖優化：Privatization、Coarsening 與 Aggregation (Histogram Optimizations)

重點總覽 (Overview)

承接 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram：basic kernel 用 atomicAdd 序列化更新，throughput 受限於 ~1/(2×memory latency)。本筆記三個技術依序「降低衝突 → 降低延遲 → 降低 commit 開銷 → 降低同位置更新次數」。

Privatization(私有化)

想法:把高度競爭的輸出資料結構複製成多份 private copies,讓一群 thread 各更新自己那份,最後再 merge 回去。

• Contention 降低倍數 ≈ 同時 active 的 block 數(across all SMs)。
• 在大規模平行系統,privatization 以 thread subset(通常一個 block) 為單位,而非每個 thread 一份(否則 merge 成本爆炸)。

                       不做 privatization                每 block 一份 private copy
   Block0  Block1 ...  ──┐                          Block0→[histo_0]  Block1→[histo_1] ...
   全部 thread 都 atomic  ├─► [ 單一 histo ] (重度衝突)        ↑只與同block衝突       ↑
                       ──┘                                   └──── commit ───► [histo_0 當公開版]

為何選「每個 thread block 一份」?

版本 A:Private copy 在 Global Memory (Fig. 9.9)

Host 端須配置 gridDim.x × NUM_BINS × 4 bytes;每 thread 加上 blockIdx.x*NUM_BINS 的 offset 指向自己 block 那份。

__global__ void histo_private_kernel(char* data, unsigned int length,
                                     unsigned int* histo) {
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < length) {
        int pos = data[i] - 'a';
        if (pos >= 0 && pos < 26)
            atomicAdd(&histo[blockIdx.x*NUM_BINS + pos/4], 1);   // 更新私有版
    }
    if (blockIdx.x > 0) {                 // block 0 那份直接當公開版
        __syncthreads();                  // 等本 block 更新完
        for (unsigned int bin = threadIdx.x; bin < NUM_BINS; bin += blockDim.x) {
            unsigned int v = histo[blockIdx.x*NUM_BINS + bin];
            if (v > 0) atomicAdd(&histo[bin], v);   // commit 到 block 0 的版本
        }
    }
}

• Commit 階段每個 bin 只有「每 block 一個 thread」競爭 → contention 極低。
• 最終結果在 histo[0 .. NUM_BINS-1](即 block 0 那份)。

版本 B:Private copy 在 Shared Memory (Fig. 9.10)

__global__ void histo_private_kernel(char* data, unsigned int length,
                                     unsigned int* histo) {
    __shared__ unsigned int histo_s[NUM_BINS];
    for (unsigned int bin = threadIdx.x; bin < NUM_BINS; bin += blockDim.x)
        histo_s[bin] = 0u;                // 平行初始化
    __syncthreads();                      // 確保全清零後才開始更新

    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < length) {
        int pos = data[i] - 'a';
        if (pos >= 0 && pos < 26)
            atomicAdd(&histo_s[pos/4], 1);   // atomic 作用在 shared memory
    }
    __syncthreads();

    for (unsigned int bin = threadIdx.x; bin < NUM_BINS; bin += blockDim.x) {
        unsigned int v = histo_s[bin];
        if (v > 0) atomicAdd(&histo[bin], v);   // commit 回 global
    }
}

Coarsening(執行緒粗化)

動機:privatization 的開銷是「每 block commit 一次」。launch 的 block 越多,commit 開銷越大。

解法:減少 block 數、讓每個 thread 處理多個輸入元素(coarsening factor = CFACTOR),把 private copy 份數降下來。分配方式有兩種。

Contiguous Partitioning (Fig. 9.12)

每個 thread 取一段連續元素:i 從 tid*CFACTOR 到 min((tid+1)*CFACTOR, length)。

    // ... shared histo_s 初始化同 Fig. 9.10 ...
    unsigned int tid = blockIdx.x*blockDim.x + threadIdx.x;
    for (unsigned int i = tid*CFACTOR;
         i < min((tid+1)*CFACTOR, length); ++i) {       // 連續區段
        int pos = data[i] - 'a';
        if (pos >= 0 && pos < 26) atomicAdd(&histo_s[pos/4], 1);
    }
    __syncthreads();
    // ... commit 同 Fig. 9.10 ...

4 thread、CFACTOR=4 (連續切):
data:  | p r o g | r a m m | i n g _ | m a s s |
          T0        T1        T2        T3
iter0:    p         r         i         m   ← warp 內位址相隔 CFACTOR → NON-coalesced
iter1:    r         a         n         a

Interleaved Partitioning (Fig. 9.14)

i 從 tid 出發,每次跳 blockDim.x*gridDim.x(全 grid 總 thread 數)。

    unsigned int tid = blockIdx.x*blockDim.x + threadIdx.x;
    for (unsigned int i = tid; i < length; i += blockDim.x*gridDim.x) {
        int pos = data[i] - 'a';
        if (pos >= 0 && pos < 26) atomicAdd(&histo_s[pos/4], 1);
    }

8 thread、交錯切:
data:  p r o g r a m m i n g _ m a s s ...
iter0: T0 T1 T2 T3 T4 T5 T6 T7  ← 相鄰位址,一次 DRAM 存取即 coalesced
iter1:                         T0 T1 ... (整體 +blockDim*gridDim)

• 第 i 輪所有 thread 共同處理一段 blockDim.x*gridDim.x 元素;位址連續 → memory coalescing,一次 DRAM burst 取回。
• 輸入長度不一定是總 thread 數倍數 → 尾端某些 thread 少跑一輪,屬正常。

Aggregation(更新聚合)

動機:某些資料局部高度集中相同值(如天空大片同色 pixel),導致對單一 bin 的重度 contention。

想法 (Merrill, 2015):每個 thread 把連續、針對同一 bin 的更新先在 register 累加,等 bin 變了才用一次 atomic flush 出去 → 大幅減少對熱門 bin 的 atomic 次數。

    // ... shared histo_s 初始化 ...
    unsigned int accumulator = 0;          // 已聚合的更新數
    int prevBinIdx = -1;                    // 上一個聚合中的 bin(-1 不會 match)
    unsigned int tid = blockIdx.x*blockDim.x + threadIdx.x;
    for (unsigned int i = tid; i < length; i += blockDim.x*gridDim.x) {
        int pos = data[i] - 'a';
        if (pos >= 0 && pos < 26) {
            int bin = pos/4;
            if (bin == prevBinIdx) {        // 同 bin → 連勝(streak),只累加
                ++accumulator;
            } else {                        // 換 bin → 結束 streak,flush 舊的
                if (accumulator > 0) atomicAdd(&histo_s[prevBinIdx], accumulator);
                accumulator = 1;
                prevBinIdx   = bin;
            }
        }
    }
    if (accumulator > 0) atomicAdd(&histo_s[prevBinIdx], accumulator); // 收尾 flush
    __syncthreads();
    // ... commit 回 global ...

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

Related Notes

• 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram
• 06-Performance-Considerations/03-Thread-Coarsening
• 06-Performance-Considerations/01-Memory-Coalescing
• 05-Memory-Architecture-And-Data-Locality/01-Memory-Access-Efficiency-and-CUDA-Memory-Types
• 10-Reduction/03-Scaling-Reduction-Hierarchical-and-Coarsening
• 15-Graph-Traversal/03-Frontiers-Privatization-and-Optimizations