在開發 CUDA 應用程式時，有效的記憶體使用 對於發揮 GPU 的全部潛力至關重要。在眾多最佳化策略中，合併記憶體存取（Coalesced Memory Access） 在降低記憶體延遲與最大化頻寬使用率方面扮演關鍵角色。本文將探討此概念的核心原理、其重要性，以及如何在 CUDA 程式中實作。

什麼是合併記憶體存取？

在 CUDA 中，全域記憶體（Global Memory） 相較於 共享記憶體（Shared Memory） 來說速度較慢。當一個 warp（32 個執行緒） 存取全域記憶體時，GPU 會嘗試以單一記憶體交易（memory transaction）讀取或寫入資料。若要高效執行，所有執行緒的記憶體存取應該是合併的，也就是存取連續的記憶體位址。如果存取模式是非合併的，GPU 會將該操作拆分為多個較小的交易，進而顯著增加記憶體延遲。

為何合併記憶體存取很重要？

合併與非合併記憶體存取的效能差異可能極為顯著。例如，當執行緒按照合併模式存取記憶體時，CUDA 核心（Kernel）的執行速度可能是非合併存取模式的 兩倍以上。以下是一個簡單的 CUDA 核心的效能比較：

  • 合併存取：232 微秒
  • 非合併存取：540 微秒

非合併存取速度幾乎是合併存取的 2.3 倍，這凸顯了適當對齊記憶體存取模式的必要性。

合併記憶體存取的技巧

為了在 CUDA 核心中實作合併記憶體存取模式，可以考慮以下策略：

1. 對齊執行緒與記憶體布局

確保執行緒索引（thread ID）對應到記憶體中的連續位置。例如，執行緒 i 應該存取陣列的第 i 個元素：

@cuda.jit
def coalesced_access(a, b, out):
    i = cuda.grid(1)
    out[i] = a[i] + b[i]  # 合併存取

2. 使用共享記憶體（Shared Memory）

共享記憶體是一種快取，位於 GPU 晶片上，由區塊內的執行緒共享。透過共享記憶體，我們可以在不規則的存取模式下實現合併存取：

@cuda.jit
def shared_memory_example(a, out):
    tile = cuda.shared.array((32, 32), dtype=numba.types.float32)
    i, j = cuda.grid(2)
    tile[cuda.threadIdx.y, cuda.threadIdx.x] = a[i, j]  # 合併讀取
    cuda.syncthreads()
    out[j, i] = tile[cuda.threadIdx.x, cuda.threadIdx.y]  # 合併寫入

3. 最佳化 2D 和 3D 格狀結構

當處理 二維（2D）或三維（3D）資料 時，應當合理設計 CUDA 的網格（Grid）和區塊（Block），確保執行緒與記憶體布局對齊，以減少非合併存取的發生。

共享記憶體與 Bank Conflict（記憶體銀行衝突）

儘管共享記憶體能夠帶來顯著的效能提升，但不當的使用方式可能導致記憶體銀行衝突（Bank Conflict）。CUDA 的共享記憶體由多個記憶體銀行組成，若同一個 warp 中的多個執行緒同時存取相同的記憶體銀行，這些存取將會序列化，導致效能下降。

解決方案：增加記憶體填充（Padding），確保每個執行緒存取不同的記憶體銀行。例如：

tile = cuda.shared.array((32, 33), dtype=numba.types.float32)  # 增加填充

這樣做可以確保連續的執行緒存取不同的記憶體銀行，避免衝突。

案例研究：矩陣轉置（Matrix Transpose）最佳化

考慮矩陣轉置（Matrix Transpose）這一運算，若使用合併讀寫模式，效能將顯著提升。以下是不同方法的效能比較：

  1. 天真方法（Naive Kernel）：合併讀取，但寫入不合併。
  2. 共享記憶體方法（Shared Memory Kernel）：透過共享記憶體實現合併讀取與寫入。
  3. 最佳化方法（Optimized Kernel）：使用共享記憶體並解決記憶體銀行衝突。

效能比較：

  • 天真方法：1.61 毫秒
  • 共享記憶體方法：1.1 毫秒
  • 最佳化方法：0.79 毫秒

重要結論

  • 合併記憶體存取 可以降低延遲、提高頻寬利用率，是 CUDA 最佳化的重要技術。
  • 共享記憶體 可幫助實現合併存取，但需注意 記憶體銀行衝突。
  • 優化記憶體存取模式 往往只需少量代碼更改，但可獲得 顯著效能提升。

透過掌握合併記憶體存取與共享記憶體技術，你可以撰寫高效能的 CUDA 核心，最大化 GPU 的運算能力。此外，別忘了使用 CUDA Profiler 來分析效能瓶頸，驗證你的最佳化策略！