效能考量 練習題 (Practice - Performance Considerations)

DRAM cell 是一個微小電容,讀取時要靠它把微量電荷分享 (charge sharing) 給高電容的 bit line,再由 sense amplifier 偵測是否足以判定為「1」。這個過程約數十奈秒,遠慢於 sub-nanosecond 的時脈。為了在每塊晶片塞更多 bit,電容反而越做越小、越弱 → latency 不降反停滯。

每次存取一個 DRAM 位置時,硬體會把「包含該位置的一整段連續位置」一起讀出;晶片內有許多 sensor 平行運作,各偵測一個 bit,偵測完成後整段連續資料可以高速送往處理器,這段連續資料即 DRAM burst。若程式集中使用 burst 內的所有 byte,就能攤平那段昂貴的 row access latency,逼近 peak 頻寬;反之隨機存取會浪費 burst 中其它 byte。

Memory coalescing:當一個 warp 的所有 threads 在同一個 load 指令存取「連續的 global memory 位址」時(thread 0 → X、thread 1 → X+1 …),硬體會把它們合併成「對連續 DRAM 位置的單一請求」,讓 DRAM 以一個 burst 交付。warp 內 threads 因 SIMD 而同時執行同一個 load,執行排程天然一致(類比:作息相同才好共乘),所以是完美的合併候選。

以 row-major(列為主) 線性化:同一列的相鄰元素放在連續位址,M[row*Width + col];相鄰的 M0,0 與 M1,0 在記憶體中相隔 Width。口訣:stride = 索引中 threadIdx.x 的係數。係數為 1 → 連續 thread 存連續位址 → coalesced;係數為 Width → 相鄰 thread 相隔 Width → uncoalesced。

Channel = 一個 memory controller 加一條 bus,把一組 DRAM banks 連到處理器(真實系統 1–8 個 channel)。Bank = 一塊 DRAM cell array + sense amplifiers + 把 burst 送上 bus 的介面。burst 提供「單次存取內」的平行;banks 提供跨存取的延遲重疊(某 bank 等延遲時另一 bank 並行存取);channels 提供跨 bus 的頻寬平行。三層平行(bursts/banks/channels)合起來才能堆出所需頻寬。

公式:頻寬 = bus 寬度(B) × 2 × clock(GHz)。DDR (double data rate) 在每個 clock cycle 傳輸兩次——一次在 rising edge、一次在 falling edge,故乘以 2。例:64-bit(= 8 B)DDR @ 1 GHz → 8 × 2 × 1 = 16 GB/s。

單 bank 利用率 = 1/(R+1)(延遲遠大於傳輸,例 R=20 → 1/21 ≈ 4.8%)。多個 bank 可讓某 bank 等延遲時另一 bank 並行存取,重疊延遲;要跑滿 bus 至少需 R+1 個 bank。實務上更多的兩個原因:(1) 降低 bank conflict(多筆存取打到同一 bank 時延遲無法重疊)的機率;(2) 單一 cell array 容量受延遲/良率限制,需要許多 bank 才能撐起記憶體容量。

順序是 channel 優先:先填滿某 channel 某 bank 的一個 burst,再換下一個 channel;一輪 channel 用完才換到下一個 bank(toy 例 burst=2:M[0..1]→Ch0/bank0、M[2..3]→Ch1/bank0 …、M[8..9]→Ch0/bank1)。這讓連續 coalesced 存取天然被拆到不同 channel。跑滿頻寬的三條件:(1) 每筆存取 coalesced;(2) 存取平均分散到各 channel 與 bank;(3) 有夠多 thread(高 occupancy) 同時發出存取——這是 maximizing occupancy 的第二個好處(除了隱藏 pipeline 延遲,也隱藏 DRAM 延遲)。

Thread coarsening = 放棄最細粒度(每 thread 一個 work unit),改讓每個 thread 負責多個 work units,以部分串行化省下不必要的 price of parallelism(重複載入、重複計算、同步開銷)。COARSE_FACTOR = 每個 coarsened thread 負責的原始 work unit 數;coarsening loop = thread 內走訪這些 work units 的迴圈。三陷阱:(1) 沒有 price of parallelism 時硬套(如 vector add)→ 無效;(2) 過度粗化使暴露平行度不足 → 資源閒置、scalability 不再透明;(3) registers/shared memory 用量過大 → 壓垮 occupancy。

六項:(1) maximizing occupancy、(2) coalesced global memory accesses、(3) minimizing control divergence、(4) tiling of reused data、(5) privatization、(6) thread coarsening。本質:多用一種資源以減輕另一種資源的壓力;只有當「被緩解者是真正的瓶頸」且「被加重者不會反過來限制平行執行」時才有效。因此套用前必須先用 profiling tools 找出 bottleneck(瓶頸是 hardware-specific 的),否則優化可能無效甚至更糟。

M[k*Width + col]:col 的 stride = 1 → 連續 thread 存連續位址 → coalesced。M[col*Width + k]:col 被乘上 Width → 連續 thread 相隔 Width(實務上數百~數千)→ uncoalesced。修正用 corner turning:交換 threadIdx.x 與 threadIdx.y 的角色,讓連續 thread 載入同一 column 的相鄰元素(column-major 下相鄰),以 coalesced 方式搬進 shared memory,不利的存取留在 SRAM 做(SRAM 不需 coalescing,幾乎無 penalty)。

單 channel 頻寬 = 8 B × 2 × 1 GHz = 16 GB/s;所需 channel 數 = 256 / 16 = 16 channels(現代 CPU 32 GB/s 則需 2 channels)。R=20:單 bank 利用率 = 1/(R+1) = 1/21 ≈ 4.8%(16 GB/s channel 實際只送約 0.76 GB/s);跑滿 bus 至少需 R+1 = 21 個 bank(實務上更多以降低 bank conflict)。

colStart = bx * TILE_WIDTH * COARSE_FACTOR + tx(block column 跨度從 TILE_WIDTH 變成 TILE_WIDTH * COARSE_FACTOR,第 c 個元素的 col = colStart + c*TILE_WIDTH)。OP/B:simple = 0.25(2 FLOP / 8 B);tiled 32×32 = 8(≈ 0.25 × TILE_WIDTH = 0.25 × 32,資料重用 ~32 倍);tiled+coarse 4 = 12.8(每 phase 載 1 M tile + 4 N tile 算 4 個 output tile,M tile 載入從 4 塊降到 1 塊 → traffic 下降)。OP/B 越高代表越偏 compute-bound、越不受 global memory 頻寬限制。

Coalesced:warp 的存取相鄰(stride 1),硬體合併成單一 DRAM 請求,完整用滿一個 burst → global memory traffic 低。Uncoalesced:相鄰 thread 相隔 Width,硬體判定位址相距太遠無法合併 → 變成 warp 內多筆分散請求,每筆只用到 burst 中一個 byte、其餘浪費 → traffic 暴增、頻寬被癱瘓(像塞車)。Corner turning 的洞見:把「不利的存取模式」從 global memory 轉移到 shared memory——進出 global memory 那一步改成 coalesced(交換 tx/ty,沿 column 載入),不規則存取改在 SRAM 做。因為 SRAM 不需 coalescing(無 burst 限制),這幾乎沒有代價:本質上是「多用一點 shared memory + 一次重排,換掉昂貴的 uncoalesced global 存取」。

最細粒度暴露最多平行度,讓硬體自由地平行或串行,具 transparent scalability;但平行化常有 price of parallelism(如 tiled matmul 中相鄰 output tile 由不同 block 處理 → 各自重複載入同一塊 M input tile)。關鍵權衡:若這些 block 真能平行跑,重複載入值得付;但若硬體因資源不足而串行化它們,這個 price 就白付了——此時改由一個 block(每 thread 做 COARSE_FACTOR 個元素)處理多個相鄰 tile,M tile 只載入一次並重用,省下重複 traffic(OP/B 8 → 12.8)。反例:vector addition / RGB→grayscale 平行處理不同元素沒有任何 price of parallelism,套 coarsening 不會有明顯效益,還可能因 COARSE_FACTOR 過大而減少暴露平行度、壓低 occupancy → 反而變慢。結論:coarsening = 「拿暴露的平行度,換掉本來會被串行化而白付的重複成本」,只有兩個條件同時成立才划算。