平行直方圖 (Parallel Histogram) 練習題 (Practice - Atomic Operations and the Basic Histogram Kernel)

因為每個 thread 的 輸出位置是 data-dependent(由輸入值決定 histo[(data[i]-'a')/4]),多個 thread 可能更新 同一個 bin,無法把每個 output element 專屬給某個 thread。這種多 thread 寫同一輸出位置稱為 output interference,是 race condition 的根源。

它讀取 address 指向的 32-bit word(global 或 shared memory),加上 val,把結果寫回同一位址,整個 read-modify-write 形成 不可分割 (undividable) 單位,並 回傳該位址的舊值 (old value)。atomicAdd 是 intrinsic function,編譯成單一硬體 atomic 指令,程式中沒有真正的函式呼叫。

當兩 thread 的 RMW overlap:T1 R(0)→T2 R(0)(此時 T1 還沒寫回)→T1 W(1)→T2 W(1),最終 histo[x]=1,T1 的更新被覆蓋(lost update)。問題本質:histo[x]++ 是 read→+1→write 三步驟,可被其他 thread 交錯;最終值隨相對時序變動即為 read-modify-write race condition。

同一位址的 atomic 無法 overlap:trailing thread 的 RMW 必須等 leading thread 的 RMW 完成才能開始(Fig. 9.7),一次只有一個進行中。每個 atomic 約佔 load latency + store latency ≈ 2 × memory latency,故單一位址吞吐 ≈ 1 / (2 × latency)。此時 延遲(非頻寬)主導,且會 dominate kernel 執行時間。

因為 throughput ≈ 1/(2×latency),降低 latency 直接提升吞吐。L2 存取約 數十 cycle(vs DRAM 數百 cycle)。高度競爭的變數被多 thread 頻繁存取,一旦載入便 常駐 L2(命中即直接在 cache 更新),使 atomic 吞吐相較早期 GPU 至少提升一個數量級。

把高度競爭的 histogram 複製成多份 private copies,讓一群 thread 各更新自己那份;通常以 每個 thread block 一份 為單位。Contention 約降低 同時 active 的 block 數 倍。Global 版須配置 gridDim.x × NUM_BINS × 4 bytes,thread 以 offset blockIdx.x * NUM_BINS 指向自己 block 那份;最後每 block 把私有值 commit(atomicAdd)回 block 0 的版本。

Per-block 的好處:(1) block 內可用 __syncthreads() 等彼此更新完再 commit(若跨 block 共用就得另開 kernel 或更複雜手法);(2) 若 bins 夠少,private copy 可宣告在 shared memory(block 間看不到彼此的 shared memory,所以必須 per-block)。Shared memory latency 僅 數 cycle,直接轉成更高的 atomic throughput。前提是 bins 數夠小 能塞進 shared memory,否則退回 global 版。

每個 thread 從 i = tid(global thread index)開始,每輪 i += blockDim.x * gridDim.x(全 grid 總 thread 數)。同一輪內,相鄰 thread 存取 相鄰位址,一次 DRAM burst 即可 coalesced 取回,充分利用記憶體頻寬;contiguous 則讓 warp 內位址相隔 CFACTOR,non-coalesced。

同一位址的 atomic 不能 overlap,故每 100 ns 最多完成一個:throughput = 1 / 100 ns = 10 M atomics/s。
(若題目把 100 ns 視為單次 load 或 store 的 latency,則一個 atomic ≈ read+write = 200 ns → 1/(2×latency) = 5 M atomics/s;按 Ch.9 公式 throughput ≈ 1/(2×latency)。)

(d) atomicAdd(&Total, Partial);。第一參數是要更新位置的 位址(&Total),第二參數是要加上的 值(Partial,非位址)。

(a) 524,288(每元素一次 global atomic,無 privatization)。
(b) blocks = 524288/1024 = 512;histogram 階段 atomic 在 shared memory,只有 commit 階段打 global,每 block 至多 128 bins → 512 × 128 = 65,536。
(c) blocks = 524288/(1024×4) = 128 → 128 × 128 = 16,384。

CPU:thread 少、每個 cache 只服務少量 thread,contiguous 的循序存取能重用 cache line 且互不干擾 → 最佳。GPU:SM 內大量同時 active thread 互相干擾 cache,單一 thread 的連續資料無法保留在 cache;且 contiguous 讓 warp 內位址相隔 CFACTOR → non-coalesced,浪費 DRAM 頻寬。Interleaved 讓 warp 相鄰 thread 取相鄰位址 → memory coalescing,一次 burst 服務整個 warp,故為 GPU 首選。

Aggregation 把連續、同一 bin 的更新先在 register 累加,bin 變了才 atomic flush → 減少對熱門 bin 的 atomic 次數。高 contention(skewed,如天空大片同色) 時顯著加速;低 contention 時因多了變數與指令,可能比 simple kernel 更慢。if 的 divergence 在兩個極端都小(全部 streak 或全部 flush);中間情況的 divergence 通常被 降低的 contention 補償 掉。