記憶體存取效率與 CUDA 記憶體類型 (Memory Access Efficiency & CUDA Memory Types)

重點總覽 (Overview)

記憶體存取效率的重要性 (Importance of Memory Access Efficiency)

Compute-to-Global-Memory-Access Ratio

以 Ch.3 矩陣乘法 kernel 最熱的內層 dot-product 迴圈為例:

// 內層迴圈每一次迭代:
for (int k = 0; k < Width; ++k) {
    Pvalue += M[row*Width + k] * N[k*Width + col];
    //         ^^^^^^^^^^^^^^ 1 次 global load (4B)
    //                          ^^^^^^^^^^^^^^ 1 次 global load (4B)
    // 1 次 FP multiply + 1 次 FP add = 2 FLOP
}

• 每次迭代:2 FLOP 對 8 B global memory 存取 → ratio = 2/8 = 0.25 FLOP/B (OP/B)。
• 這個比值就是 compute to global memory access ratio,文獻中也叫 arithmetic intensity 或 computational intensity。

要餵滿運算單元需要的 ratio:

代表每讀 1 個 4-byte 浮點值,必須做約 50 個浮點運算才能完全發揮 A100。能否達到取決於計算本身的 data reuse 程度。

Memory-bound vs Compute-bound

Roofline Model

評估「應用程式相對於硬體上限」效率的視覺模型。

 GFLOPS (throughput, y軸, log)
   ^
   |                 ______________________  <- 水平線: peak compute (GFLOPS)
   |                /        A3 (compute-bound)
   |               /  • A3
   |   斜線:       /
   |   peak BW   /•A1   <- A1 貼著斜線: 高效使用頻寬
   |   (slope)  /
   |          / • A2    <- A2 遠低於斜線: 頻寬利用率差,有優化空間
   |         /
   |        /
   +-------+------------------------------> arithmetic intensity (OP/B, x軸 log)
        交點 = memory-bound ↔ compute-bound 分界

• x 軸:arithmetic intensity (OP/B) = 每 byte 載入做多少 work。
• y 軸:實測 computational throughput (GFLOPS)。
• 斜線(從原點出發,正斜率)= peak memory bandwidth 上限。
• 水平線 = peak compute throughput 上限。
• 任何 app 的點必在兩線之下(不可能超過硬體 peak)。
• 點離線越近 = 資源用得越有效率;交點橫座標 = memory↔compute bound 的分界 intensity。

CUDA 記憶體類型 (CUDA Memory Types)

CUDA device 提供多種記憶體幫助提高 compute-to-memory ratio。底層回到 von Neumann model:global memory 對應 Memory box(off-chip DRAM,高延遲低頻寬),register file 在 processor chip 上(極短延遲、極高頻寬)。

            CUDA Device (晶片邊界 = processor box)
  +-------------------------------------------------+
  |  SM                                             |
  |   +------------+   +------------------------+   |  on-chip:
  |   | Registers  |   |   Shared Memory        |   |  超高速、平行
  |   | (per-thread)|  |   (per-block, scratchpad)| |
  |   +------------+   +------------------------+   |
  |        ^                    ^                   |
  +--------|--------------------|-------------------+
           |  (R/W, no load     |  (需 load 指令,
           |   instruction)     |   但 on-chip 低延遲)
           v                    v
  +-------------------------------------------------+
  |  Global Memory (DRAM, off-chip)  | Constant Mem |  off-chip:
  |  R/W by host & device            | R-only, 有 cache|  高延遲 (數百 cycle)
  |  (Local memory 也實際放這裡)      | host W / dev R |
  +-------------------------------------------------+

Registers vs Shared vs Global 三者差異

Shared Memory 的角色

• 在電腦架構術語裡是一種 scratchpad memory(程式員手動管理,非自動 cache)。
• 因為 on-chip,延遲遠低於 global;但因仍需 load 操作,延遲/頻寬不如 register。
• 關鍵差異:shared 變數 block 內所有 thread 可見,而 register 是 thread 私有 → shared memory 是 thread 間高頻寬協作的手段(這是 tiling 的基礎)。

宣告語法、Scope 與 Lifetime (Table 5.1)

// 各記憶體類型宣告範例
__global__ void kernel(...) {
    int idx = threadIdx.x;          // 自動純量 -> Register, scope=thread
    float tmp[4];                   // 自動陣列 -> Local memory (見警告)
    __shared__ float tile[16][16];  // Shared, scope=block, 也可寫 __device__ __shared__
    // ...
}
__device__ int   GlobalVar;          // Global, 宣告在任何 function 之外
__constant__ int ConstVar[256];      // Constant, 必須在 function 之外宣告

Scope:能存取此變數的 thread 集合。

• thread scope:每個 thread 有私有副本。launch 一百萬 thread → 產生一百萬份。
• block scope:每個 block 一份,block 內 thread 共用。
• grid scope:所有 grid 所有 thread 共用同一份。

Lifetime:變數在程式執行期間可用的時段。

• Grid lifetime:必須宣告在 kernel function body 內;kernel 多次呼叫間「不保留」值,每次都要重新初始化。
• Application lifetime:必須宣告在任何 function body 之外;值在整個 application 期間維持,所有 kernel 可見。

Constant 與 Global 變數細節

Pointer 用法

kernel/device function 中 pointer 指向 global memory 的兩種典型方式:

1. host 用 cudaMalloc 配置物件,pointer 當 kernel 參數傳入(Ch.2 / Ch.3 的作法)。
2. 把 global 變數的位址指派給 pointer,例如 kernel 內 float* ptr = &GlobalVar;。

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

Related Notes

• 05-Memory-Architecture-And-Data-Locality/02-Tiling-and-Tiled-Matrix-Multiplication
• 05-Memory-Architecture-And-Data-Locality/03-Boundary-Checks-and-Memory-Occupancy
• 04-Compute-Architecture-And-Scheduling/03-Resource-Partitioning-and-Occupancy
• 06-Performance-Considerations/01-Memory-Coalescing
• 07-Convolution/02-Constant-Memory-and-Caching
• 22-Advanced-Practices-And-Future-Evolution/03-Memory-Bandwidth-and-Compute-Throughput