Enter your search text in the box above

Select a result to preview

原子操作與基本直方圖核心 (Atomic Operations and Basic Histogram Kernel)

#parallel-pattern #histogram #atomic-operations #race-condition #latency-vs-throughput

重點總覽 (Overview)

項目 重點 為何重要
Histogram pattern 統計各 value interval 內資料出現次數 輸出位置 data-dependent,無法套用 owner-computes rule
Output interference 多個 thread 可能更新同一個 histo[] bin 引入 race condition 的根源
Read-modify-write (RMW) histo[k]++ = read → +1 → write 三步驟 三步驟可被其他 thread 交錯 → 結果錯誤
Race condition 最終值隨 thread 相對時序而變 部分交錯正確、部分丟失更新 (lost update)
atomicAdd() 硬體保證 RMW 不可分割 (undividable) 序列化同一位址更新 → 正確
Atomic 不保證順序 只保證互斥,不保證誰先誰後 序列化但 thread 執行序仍任意
Throughput 瓶頸 atomic 週期 ≈ load latency + store latency 同一位址吞吐 ≈ 1/(2×latency),極低
Last-level cache (L2) atomics 在 SM 共享的 LLC 做 atomic latency 數百 → 數十 cycle,吞吐提升一個數量級
Important

前幾章的 pattern 都遵守 owner-computes rule(每個 output element 專屬一個 thread,寫入互不干擾)。Histogram 打破此規則:任何 thread 都可能更新任何 output element,因此必須處理 output interference。

直方圖背景 (Histogram Background)

循序版本 (Sequential C function, Fig. 9.2)

void histo_sequential(char *data, unsigned int length, unsigned int *histo) {
    for (unsigned int i = 0; i < length; ++i) {
        int alphabet_position = data[i] - 'a';      // ASCII 連續編碼: 'a'->0
        if (alphabet_position >= 0 && alphabet_position < 26)
            histo[alphabet_position / 4]++;          // 每 4 字母一個 interval
    }
}

原子操作與基本核心 (Atomic Operations & Basic Kernel)

平行化策略與 output interference

最直接的平行化:啟動與資料元素一樣多的 thread,每個 thread 處理一個輸入元素並遞增對應 bin。

data:   p  r  o  g  r  a  m  m  ...      (input chars)
        |  |  |  |  |  |  |  |
thread: t0 t1 t2 t3 t4 t5 t6 t7
         \  \  \________\__\____         多個 thread 寫同一 bin
          \  \         (m-p interval)    => OUTPUT INTERFERENCE
           v  v             v
histo[]: [a-d][e-h][i-l][m-p][q-t][u-x][y-z]

histo[k]++ 是一個 read-modify-write (update) 操作:read 記憶體 → +1 (modify) → 寫回 (write)。

Read-modify-write race condition

當兩個 thread 的 RMW 序列 交錯 (overlap),結果視相對時序而定:

正確交錯 (序列化)                錯誤交錯 (overlap → lost update)
T1: R(0) M(1) W(1)              T1: R(0)      M(1)      W(1)
T2:             R(1) M(2) W(2)  T2:      R(0)      M(1)      W(1)
       histo[x] = 2  (correct)         histo[x] = 1  (WRONG: T1 更新被覆蓋丟失)
這是 read-modify-write race condition(與 Ch.10 Kogge-Stone scan 的 write-after-read race 類似但不同)。最終值隨時序變動 = 存在 race condition,部分結果正確、部分錯誤。

atomicAdd 解法 (Fig. 9.6)

Atomic operation = 對某記憶體位址的 RMW 序列,保證 沒有其他 RMW 序列能與它 overlap(硬體鎖定該位址直到完成),read/modify/write 形成不可分割單位。

__global__ void histo_kernel(char *data, unsigned int length, unsigned int *histo) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;   // 取代 sequential loop
    if (i < length) {                                          // boundary check
        int alphabet_position = data[i] - 'a';
        if (alphabet_position >= 0 && alphabet_position < 26)
            atomicAdd(&(histo[alphabet_position / 4]), 1);     // 序列化更新
    }
}
Important

Atomic 只強制互斥/序列化,不強制執行順序。trailing thread 必須等 leading thread 的 atomic 完成才能開始,但「誰是 leading」是任意的。T1、T2 兩種完成順序皆合法,因為加法可交換,結果都正確。

Exercise 5 經典題:把 Partial 加到全域變數 Total 應寫 atomicAdd(&Total, Partial);(第一參數是位址 &Total,第二參數是值 Partial)。

原子操作的延遲與吞吐 (Latency & Throughput of Atomic Operations)

序列化任何部分都會拖慢大規模平行程式。GPU 靠 大量 thread 的記憶體延遲互相重疊 來達到高吞吐;但 同一位址的 atomic 無法重疊——一次只能有一個進行中。

同一位址 atomic 的時間軸 (Fig. 9.7):
atomic#1: [---- read latency ----][---- write latency ----]
atomic#2:                                                   [---- read ----][---- write ----]
atomic#3:                                                                                    [-- ...
          |<------ 一個 atomic ≈ 2 × memory latency ------>|
          每個 atomic 必須完整佔用 (load + store) latency → 吞吐受限

核心公式

量 公式 / 數值
Atomic 週期 ≈ load latency + store latency ≈ 2 × memory latency
單一位址吞吐 throughput ≈ 1 / (2 × latency)

書中範例(64-bit DDR、8 channels、1 GHz、200-cycle latency):

Peak memory BW   = 8 B × 2 transfers/clk × 1 GHz × 8 ch = 128 GB/s
Peak elements    = 128 GB/s ÷ 4 B = 32 G elements/s          (一般存取)

Atomic 同一位址 = 1 atomic / 400 cycles                       (200 read + 200 write)
                = (1/400) × 1 GHz = 2.5 M atomics/s            (慘!比 32G 低 ~4 個數量級)
一般規則「越多並行記憶體存取 → 吞吐越高」在 重度競爭同一位址的 atomic 情境下完全失效:此時延遲(而非頻寬)主導,且會 dominate kernel 執行時間。

緩解:Last-Level Cache atomics

進一步降低 latency / 競爭的手段(privatization → shared memory atomic、coarsening、aggregation)見 09-Parallel-Histogram/02-Histogram-Optimizations-Privatization-Coarsening-Aggregation。

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

情境 / 關鍵字 答案 / 技巧
為何 histogram 不能用 owner-computes rule? 輸出位置 data-dependent,多 thread 可能更新同一 bin → output interference
什麼是 read-modify-write race condition? 兩個以上同位址更新序列 交錯,最終值隨相對時序而變(可能 lost update)
atomic operation 的定義? 對某位址的 RMW 形成 不可分割單位,無其他 RMW 能與之 overlap(硬體鎖定)
atomicAdd 回傳什麼? 舊值 (old value at address)
atomic 是否保證執行順序? 否,只保證互斥/序列化;誰先誰後任意
正確 atomicAdd 把 Partial 加到 Total atomicAdd(&Total, Partial);(位址 + 值)
單一位址 atomic 最大吞吐(latency=T)? 1 / (2T);e.g. 100 ns → 1/(200ns) = 5 M atomics/s
為何 atomic 吞吐遠低於一般存取? 同位址無法 overlap,週期 ≈ 2× latency,延遲主導
LLC atomic 為何加速? latency 數百→數十 cycle → 吞吐升一個數量級
Fig. 9.6 在 global memory 做幾次 atomic? = 有效輸入元素數(每元素一次,無 privatization/coarsening)
均勻 vs 偏斜分布對吞吐影響? 均勻 → 吞吐 ≈ ×bin 數;偏斜 → boost 遠小於 bin 數(熱點競爭)