教材:10710周志遠教授平行程式
https://www.youtube.com/playlist?list=PLS0SUwlYe8cxqw70UHOE5n4Lm-mXFXbZT

20250910 筆記
內容可能有錯僅供參考

18B 平行記憶體 / GPU的架構
19A 記憶體層級 / CUDA的優化

今日大綱
1. GPU 硬體架構與 CUDA 映射
2. 執行緒管理與 Warp 概念
3. GPU 硬體限制與排程
4. 分支預測與 Warp Divergence
5. 迴圈展開 (Loop Unrolling)
6. Atomic Operations
7. GPU 記憶體層級 (Memory Hierarchy)
8. 變數與記憶體的對應關係

### GPU 硬體架構與 CUDA 映射

* 接下來將從 **GPU 的硬體架構 (hardware architecture)** 出發，回顧 CUDA 的行為或函數定義，並深入解釋 CUDA 這樣設計的原因是基於硬體線路考量。
* 主要探討記憶體與執行兩大塊，尤其是執行部分的硬體和記憶體如何運作。
* 了解 **local variable (區域變數)**、**shared variable (共用變數)** 和 **global variable (全域變數)** 在 GPU 上是如何被管理 (managed)。
* **CUDA 軟體與 GPU 硬體之間的映射關係 (mapping relationship)**：
* 一個 **thread (執行緒)** 對應到一個 C 語言執行的 ALU (算術邏輯單元)。
* 一個 **thread block (執行緒區塊)** 對應到一個 **Streaming Multiprocessor (SM)**，也稱為 SM 處理器。
* 一個 **grid (網格)**，也就是一個 **kernel (核心)**，對應到一整個 **device (裝置)**。

### 執行緒管理與 Warp 概念

* GPU 在管理執行緒時，以 **thread block 為基本單位**。
* 一個完整的 block 會被排程 (schedule) 到其中一個 SM 處理器上執行。
* 由於 SM 處理器上有 **shared memory** 的硬體晶片，同一 block 內的執行緒可以透過 shared memory 進行溝通。
* 一個 block 內的所有執行緒必須執行完畢後，才會讓下一個 block 執行。
* 然而，同一 block 內的所有執行緒不一定會同時執行，不同 block 之間由於可能被放到不同的 SM 上，其執行順序也無保證。
* 一個 block 被映射到 SM 後，其執行緒會被切割成所謂的 **warp**。
* **Warp 才是真正的 C 語言處理器 (physical processor) 執行單位**。
* **同一個 warp 裡的執行緒會實體上同時 (physically) 被執行**。
* 同一個 block 裡的執行緒只是邏輯上 (logically) 排在同一處理器上執行，但執行順序可能不同。
* **一個 warp 通常包含 32 個執行緒**。

### GPU 硬體限制與排程

* GPU 硬體在平行執行不同 warp 或 block 時，會有一些基本限制。
* 以較早的 Fermi 架構為例：
* 每個 SM 最多只能有 **8 個 block**。
* 同一個 SM 上，**活躍 (active) 的 warp 數量最多只有 48 個**。
* 如果一個 block 超過這些限制，它將不會被排程到該 SM 上執行。
* 每個 SM 最多可以同時執行的執行緒數量計算方式：48 個 warp * 32 個執行緒/warp = **1536 個執行緒**。
* 這個數量決定了 GPU 規格中「最大平行執行緒數量」的限制。
* GPU 的內部排程器 (scheduler) 在符合平行度要求且資源充足的情況下，會將 block 排程執行。
* 儘管一個 SM 可以有 48 個 active warp，但這不代表它有 48 個實體 C 語言處理器。
* 一個 warp 在執行過程中可能因為 **memory stall (記憶體停滯)** 而無法使用 GPU 的 C 語言處理器。
* **active warp 的數量通常大於實體 C 語言處理器的數量**，目的是最大化 GPU C 語言處理器的利用率，避免浪費時間等待。
* **內部機制**：每個處理器會對應到多個 warp。當一個 warp 中的執行緒需要進行記憶體存取時，處理器會立即進行 **context switch** 到另一個 active warp。
* 這種 context switch 是 **物理性的且速度非常快**，只需重新指向暫存器 (register) 即可。
* 只要程式碼中有記憶體存取動作，GPU 就會在 active warp 之間不斷切換，確保計算資源不被閒置。
* 在一個 warp 中，只要有一個執行緒需要進行記憶體存取，**整個 warp 的所有執行緒都必須停下來**，讓出處理器給其他 warp 的執行緒執行。
* GPU 的架構設計主要目的就是盡量 **隱藏記憶體存取時間**，以達到最高的計算能力.
* GPU 擁有大量的核心，且所有執行緒通常執行相同的指令。如果任何一個執行緒進行記憶體存取，會導致整個 warp 的執行緒停滯。因此，設計上會讓更多的執行緒可用於替換，以便隨時使用計算資源。

### 分支預測與 Warp Divergence

* 由於 GPU 本質上是 SIMD (Single Instruction, Multiple Data) 處理器，特別是同一個 warp 內的執行緒，**在 kernel 程式碼中應避免使用 branch (分支)**。
* 如果同一個 warp 內的不同執行緒執行不同的分支 (例如根據 `thread ID` 進行 `if-else` 判斷)，實際結果是這些執行緒會 **同時執行所有分支的路徑 (做 A 又做 B)**，然後再根據判斷結果合併兩者，這會使工作量變成兩倍，嚴重拖慢執行速度。
* **最佳化分支的方法**：
* 盡量避免或取代分支。
* 如果必須使用分支，應確保 **同一個 warp 內的所有執行緒都遵循相同的分支結果**。
* 例如，將條件判斷基於 `warp size` 為單位 (如 `threadIdx.x / warp_size`)，確保一個 warp 內的所有執行緒做出相同的判斷，這樣即使有分支也不會產生效能問題。
* **迴圈發散 (Loop Divergence)**：
* 如果每個執行緒都跑一個 for loop，且迴圈次數因 `index` 位置而異 (例如 `thread 0` 跑 10 次，`thread 1` 跑 12 次)。
* 結果將是 **整個 warp 會以迴圈次數最多的那個執行緒為準**，所有執行緒都必須執行那麼多次，導致許多執行緒進行無謂的計算。
* **一般建議**：
* 編寫程式碼時，需留意每個執行緒所做的事情是否完全相同、指令是否相同、次數是否相同。
* 如果不同，需判斷是否在同一個 warp 內。若在同一個 warp 內做不同事情，則會產生問題；若在不同 warp 內則影響較小。
* GPU 優化很多時候是 `case by case` 的，但要留意這些狀況。
* 需留意數據模式 (data pattern)，例如迴圈的 `index` 值可能導致計算量不同。
* 可能會使用 **decomposition (分解)** 或 **flattening (扁平化)** 的方式，強制讓每個 kernel 和每個 warp 內的計算量相同。例如，將計算分解成多個 kernel，讓大部分執行緒執行相同次數，零散的部分交由另一個較少執行緒的 kernel 處理。

### 迴圈展開 (Loop Unrolling)

* 在 GPU 上，任何分支或 for loop 都對效能不利，因為 GPU 需要判斷是否為你重複執行不同的指令。
* **一個非常基本的 GPU 優化是將 for loop 進行 loop unrolling**。
* **目的**：大幅減少執行時 `if condition` 的判斷次數，提高執行效率。
* **原理**：將一個執行 n 次的 for loop 展開成多個重複的指令。例如，`unroll` 三次意味著每次迭代處理三份原始工作。
* **考量**：
* `unroll` 的次數 (展開層數) 很重要。
* 通常假設迭代次數是 `unroll` 次數的倍數。如果不是，最後一次迭代可能會產生多餘的計算量。
* 選擇合適的 `unroll` 程度，可避免多餘的計算量。
* **優點**：減少無謂的指令，避免 GPU 誤以為需要執行各種可能性，產生更高效的程式碼。
* **實作方式**：
* 手動展開 (繁瑣)。
* 使用編譯器支援的 **`#pragma unroll` 指令**，例如 `#pragma unroll 5`，編譯器會自動為你展開迴圈。

### Atomic Operations

* **應用場景**：用於在 `shared memory` 上，當多個 block 內的執行緒共用變數 (例如 `C++`)，並對其進行修改時 (例如 `C = C + 1`)。
* 這會產生 **同步問題 (synchronization problem)**，類似於需要 `lock/unlock` 來保護。
* **建議**：在演算法設計上，應盡量 **避免使用 `lock/unlock` 或原子操作**。因為鎖定會導致執行緒執行率極低，嚴重降低平行度，GPU 本來就不是為序列化運算設計的。
* **特殊情況下的應用**：對於簡單的聚合 (aggregation) 或總結 (summarize) 類操作 (例如計算總數、建構直方圖 `histogram`)，GPU 內部支援 `atomic instruction`。
* **支援的指令**：加法 (addition)、減法 (subtraction)、最大值 (max) 等基本原子指令。
* **範例**：使用 `atomicAdd` 指令來建構直方圖，即使多個執行緒同時修改同一位置，GPU 會進行保護。
* **效能影響**：**效能會非常差，會變慢很多**。因為原子操作會強制執行緒序列化 (serialize)。例如，一個 warp 有 32 個執行緒，`atomicAdd` 指令實際上會執行 32 次 (在 warp 內序列化)，而不是同時執行。
* **結論**：雖然程式碼易於編寫且 GPU 支援，但應盡量避免。如果序列化嚴重，效能甚至可能不如 CPU。
* 接下來將介紹記憶體層級 (memory hierarchy)。

### GPU 記憶體層級 (Memory Hierarchy)

* 記憶體層級是理解 `local variable`、`shared variable` 和 `global variable` 效能與大小差異的關鍵。
* 即使是 `local variable` 也可能存儲在不同的硬體記憶體晶片上，導致效能差異。
* **GPU 上的四種主要記憶體晶片/類型**：

1. **Register (暫存器)**
* **範圍 (Scope)**：每個執行緒私有 (per thread)。
* **速度**：**最快**，延遲極低 (latency 非常小)，頻寬極大 (bandwidth 非常大)。讀寫大約 24 個週期 (cycle)，相較於 Global Memory 的 200 個週期快很多。
* **位置**：最接近每個物理 GPU 核心 (physical GPU core)。每個 GPU 核心都有自己的 register set。
* **Context Switch**：允許非常快速的上下文切換，因為只需重新指向不同的 register set。
* **限制**：
* **大小有限**：每個 SM 的 register 總量有限制，影響能有多少 active block 或 thread。一個 block 可用 register 大小是固定的 (例如 Compute Capability 6.1 下每個 block 最多 64KB)，一個 thread 的 register 限制是 255 個。
* **不可索引 (Non-indexable)**：不支援陣列 (array)。`register` 是一對一的映射，無法使用 `index` 存取多個值。
* **因此，`local variable` 如果是陣列，就無法存在 `register` 中，一定會存到 `local memory`**。只有 `int i, int j, int x` 這類簡單變數才能放在 `register`。
* **Register Spilling (暫存器溢出)**：
* 當宣告太多 `local variable` 或啟動過多執行緒時，`register` 空間不足，會導致 `register spilling`。
* 溢出的變數會被放到 **`local memory`** (實際上是 off-chip 的 `global memory`)。
* 在同一個 SM 內，一個 block 的 `register file` 空間是固定的，所以當 block 內的執行緒數量越多，每個執行緒能分配到的 `register` 空間就越少。
* **過高的平行度有時反而會降低效能**：因為增加執行緒數量可能導致 `register spilling`，使得每個執行緒的執行變慢。最佳設定是計算與記憶體存取時間達到平衡。
* **編譯時決定**：編譯器 (compiler) 在編譯程式碼時 (會產生 PTX code)，會決定每個 kernel 需要多少 `register` 空間，以及有多少 `local variable` 會放在 `register` 或溢出到 `local memory`。這些資訊在程式執行前就已確定。

2. **Shared Memory (共用記憶體)**
* **範圍**：**整個 block 共用**。
* **位置**：位於 SM 處理器內部，屬於 **on-chip memory**。
* **速度**：效能與 `register` 差異不大，因為都在同一個晶片上。
* **功能**：資料可在同一 block 內的執行緒之間共用，而 `register` 無法。
* **分配方式**：可以動態 (dynamically) 或靜態 (statically) 分配。
* **與 L1 Cache 的關係**：
* 較舊的架構 (例如 Fermi) 中，`shared memory` 與 L1 Cache 共用同一塊物理空間，可能產生競爭。
* 較新的架構 (例如 Volta, Ampere) 中，`shared memory` 與 L1 Cache 已被分開，擁有獨立的專屬空間，不再競爭，空間更大，效能更好。

3. **Global Memory (全域記憶體)**
* **範圍**：所有執行緒 (thread) 和任何 `stream processor` 都可以存取。
* **位置**：**off-chip memory** (在 SM 晶片之外)，所以容量非常大。
* **速度**：延遲較高 (latency 較高)，頻寬較低 (bandwidth 較低)。
* **與 Local Memory 的關係**：`global memory` 與 `local memory` 指的是同一塊物理記憶體空間，只是用途不同。
* **快取 (Cache)**：**不能被快取 (cannot be cached)**，因此存取速度永遠較慢。

4. **Local Memory (區域記憶體)**
* **定義**：它不是指一個獨立的物理記憶體，而是指 **`global memory` 中用於儲存 `local variable` 的空間**。
* **用途**：專門用來存放那些無法存放在 `register` 中的 `local variable` (例如陣列或 `register spilling` 溢出的變數)。
* **範圍**：是該執行緒的私有區域 (private area)，其他執行緒無法存取。因此，沒有平行存取問題。
* **位置**：實際上位於 **off-chip 的 `global memory` 上面**。
* **速度**：
* 其存取效能與 `global memory` 一樣慢。
* **關鍵優點：可被快取 (can be cached)**。它可以在 `per-block memory` (即 `shared memory` 的位置) 或 GPU 的 L1 cache 中被快取。
* **快取機制**：第一次讀取會很慢，但如果能被快取到 L1 cache 中，第二次讀取的速度就會大幅提升，甚至接近 `shared memory` 的速度。
* **編譯時決定**：編譯器在編譯時就決定哪些 `local variable` 會被放在 `global memory` (即 `local memory` 區域)，並產生其記憶體位址。

5. **Constant Memory (常數記憶體)**
* **位置**：與 `global memory` 位於不同的區塊。
* **功能**：**只能讀取 (read-only)**。
* **速度**：效能與 `global memory` 相似，但因為可被快取，所以實際存取速度會更快。

### 變數與記憶體的對應關係

* 理解軟體中宣告的變數如何映射到物理記憶體晶片對效能至關重要。
* **Local Variable (區域變數)**：
* **宣告方式**：在 `kernel function` 內部宣告，無特殊前綴。
* **可見範圍**：只對宣告它的執行緒可見。
* **物理映射**：
* 理想情況下映射到 **Register** (非常快)。
* 如果空間不足 (spilling) 或變數是陣列，則會映射到 **Local Memory** (較慢，但可快取)。
* **注意**：`local variable` 不一定很快，取決於它是否存於 `register` 或 `local memory`。

* **Shared Variable (共用變數)**：
* **宣告方式**：在 `kernel function` 內部宣告，前綴為 `__shared__`。
* **可見範圍**：同一個 block 內的執行緒可共用。
* **物理映射**：一定映射到 **Shared Memory** (快，on-chip)。

* **Global Variable (全域變數)**：
* **宣告方式**：
* 在 `kernel function` 外部宣告，前綴為 `__device__`。
* 或作為 `kernel function` 的輸入引數 (input argument) 指標傳入。
* **可見範圍**：所有執行緒可存取。
* **物理映射**：
* 映射到 **Global Memory** (最慢，off-chip)。
* 如果變數是 `read-only` 且需要快取，可映射到 **Constant Memory** (速度較快，因為可快取)。

---
其他課程連結

[平行程式1C~2B Introduction parallel programming](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/Syxh3H7Kxe)
[平行程式3A~3D The Latest Developments and Applications Using Parallel Programming](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/HJh7QFVKle)
[平行程式4A~4B IO Parallel IO and Program Analysis](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/HJLMsuHFgg)
[平行程式5A~5B The Latest Developments and Applications Using Parallel Programming](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/SJh57hIFle)
[平行程式6A~6B Communication Routines and Parallel Function Code](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/r1X9kX_Fle)
[平行程式 6C~6D Communication Routines and Parallel Function Code](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/S1DPjoYFlx)
[平行程式 7A~8A Pthread:Synchronization Problem & Tools](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/HJu-_0tKge)
[平行程式 8B~8D Synchronization Tools & Open Multi-Processing(OpenMP)](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/H1ki4E2Fee)
[平行程式 9A~9B Synchronization Construct](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/BJTYMrpKlx)
[平行程式 10A~10B Synchronization Tools & Open Multi-Processing Synchronization Construct](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/B1cY6M1qee)
[平行程式 10C~10D Synchronization Tools & Open Multi-Processing Synchronization Construct](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/BkgFaNg5gg)
[平行程式 11A~11B Parallel Work Pool and Termination / Parallel Sorting](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/H1hfOw-5xl)
[平行程式 12A~12B Parallel Sorting and Pipelined Computations](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/Symo-zQ9eg)
[平行程式 12C~12D Parallel Sorting and Pipelined Computations](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/BJYNKDVceg)
[平行程式 13A-13B Sychronous Parallelism](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/HJ2UJ2Bqex)
[平行程式 14A~14B Heterogeneous Computing](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/BksS4yP5eg)
[平行程式 14C~14D Heterogeneous Computing](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/BJrfTUd9xx)
[平行程式 15A~15B Parallel Programming Model on GPU](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/ByWnl-t5gg)
[平行程式 16A~16B What is Compute Unified Device Architecture(CUDA)?](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/HyYpsjcqgl)
[平行程式 17A~18A 平行運算的CUDA](https://hackmd.io/@6FOC2dvARe-Vz0kVSyajew/H1dUeBT5lg)
[平行程式 18B~19A 記憶體層級 / CUDA的優化](https://hackmd.io/@JuitingChen/HyF44e1jge)
[平行程式 19B~19D 記憶體層級 / CUDA的優化 ](https://hackmd.io/@JuitingChen/ryPEu4lieg)
[平行程式 20A~20B CUDA優化全域和區域記憶體/共享記憶體](https://hackmd.io/@JuitingChen/r1X659Zoxl)
[平行程式 21A~21B Parallel Reduction / Distributed Computing Framework](https://hackmd.io/@JuitingChen/HyiOpozjxl)
[平行程式 NTHU-PP-Chap10-Big Data-Part1 ](https://hackmd.io/@JuitingChen/Hyc-e3Golx)
[平行程式 NTHU-PP-Chap10-Big Data-Part2 ](https://hackmd.io/@JuitingChen/ryC_QTXoxl)
[平行程式 NTHU-PP-Chap11-MapReduce](https://hackmd.io/@JuitingChen/HJgBXJOsge)
[平行程式 NTHU-PP-Chap12-Distributed Training-Part1](https://hackmd.io/@JuitingChen/ryh5hBtsge)
[平行程式 NTHU-PP-Chap12-Distributed Training-Part2](https://hackmd.io/@JuitingChen/rJ2G7kdjxg)
[平行程式 NTHU-PP-Chap12-Distributed Training-Part3](https://hackmd.io/@JuitingChen/HkA471dilx)
[平行程式 NTHU-PP-Chap13-UCX-Part1](https://hackmd.io/@JuitingChen/rJbq103ieg)
[平行程式 NTHU-PP-Chap13-UCX-Part2](https://hackmd.io/@JuitingChen/SJpNmk_ixl)
[平行程式 NTHU-PP-Chap13-UCX-Part3](https://hackmd.io/@JuitingChen/HkIUYa13xe)