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

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

15A~15B Parallel Programming Model on GPU

今日大綱
1. Block 排程與執行
2. Warp 是什麼?
3. block 大小限制與多 block 使用
4. Memory Hierarchy
5. 記憶體分配與資源管理
6. CUDA 程式設計基本術語
7. CUDA 語言擴充與 API

#### 1.Block 排程與執行

* **區塊分配方式**：區塊 (blocks) 會被任意地分配到串流多處理器 (SM) 上。只要有可用的 SM，或 SM 擁有足夠的核心 (core) 來執行該區塊，該區塊就可能被派送到該 SM 執行。
* **區塊獨立性**：每一個 block 的情形是完全獨立的，這是其非常重要的限制。SM 之間不應有任何溝通或同步化的行為發生。
* **獨立性原因**：因為一個 SM 中包含許多核心和執行緒 (threads)。若兩個 SM 之間的執行緒有任何同步化點 (synchronization point)，將會停止這兩個 SM 的所有執行，只為了解決兩個執行緒或區塊的同步，這會浪費許多其他正在執行的計算資源。
* **SM 執行能力**：每個 SM 有很多核心，並依據區塊大小和所需執行緒數量有其上限。在達到上限之前，可以丟給它任意數量的區塊。
* **GPU 卡差異**：不同 GPU 卡的 SM 數量不同，但由於區塊的獨立性，程式行為不會因此改變。SM 數量越多，平行度越高，計算吞吐量也越高，但程式本身沒有差異。
* **區塊一旦派送，不可換 SM**：一個區塊一旦被派送到某個 SM，就必須留在該 SM 上，直到執行完畢。
* **核心 (Kernel) 執行**：
* 一個 Kernel 可以跨多個 SM 執行，因為一個 Kernel 包含多個區塊，這些區塊可以執行在不同的 SM 上。
* Kernel 不一定需要多個 SM 才能執行多個區塊；即使只有一個 SM，也可以依序執行任意數量的區塊。
* 同一區塊內的執行緒不能在不同的 SM 上執行，因為區塊是以單一單位被派送給 SM 的。

#### 2.Warp 是什麼?

* **定義**：一個 Warp 包含 32 個執行緒 (threads)。
* **執行單位**：一個 Warp 裡的所有 32 個執行緒會組成一個群組，並被執行在一個執行緒處理器 (thread processor) 上，由硬體排程器決定。
* **Context Switch**：以 Warp 為單位進行 Context Switch，即 32 個執行緒一批一批地執行指令。
* **同步行為**：若 Warp 中的任何一個執行緒有記憶體停滯 (memory stall)，整個 Warp 的執行緒都會被暫停，然後換另一個 Warp 執行。
* **實體平行與邏輯平行**：
* 只有在同一個 Warp 裡面的執行緒才是真正的**實體平行 (physically parallel)**，保證同時執行相同的指令。
* 跨 Warp 或跨區塊的執行緒只是**邏輯平行 (logically parallel)**，它們的執行可能會有時間上的差異。

#### 3.block 大小限制與多 blocks 使用

* **區塊大小限制**：區塊有其大小限制 (例如執行緒數量)，這是因為一個區塊無法跨 SM 執行，它受限於單一 SM 的核心數量。
* **強制使用多區塊**：由於存在這些上限，在需要高平行度時，有時會被迫使用多個區塊，而不是單一一個大區塊。
* **區塊間同步問題**：使用多個區塊可能導致區塊之間需要同步化。
* **Compute Capability**：這些區塊大小限制等資訊可以在 GPU 的 Compute Capability 文件中查詢到。
* **Kernel 的區塊數量限制**：一個 Kernel 可以擁有的區塊數量上限通常遠大於單一區塊的執行緒數量上限 (例如一個區塊上限約 1024 個執行緒，一個 Kernel 的區塊上限可達 2^31-1 或更高)。由於區塊是獨立執行，多一點區塊通常問題較小。
* **程式設計考量**：
* 過多的執行緒不一定能加速，因為它仍受限於實體硬體資源。
* 需要考量硬體層面的平行度，例如一個 GPU 最多可以同時執行多少個 Kernel，每個 SM 最多可以開啟多少執行緒，或每個區塊的執行緒上限等。
* 超越硬體限制意味著無法實現平行。

#### 4. Memory Hierarchy

GPU 記憶體分為多個層次：

1. **Global Memory (全域記憶體)**：
* **最外層**的記憶體，所有執行緒 (包括跨 Kernel) 都可以讀取。
* **生命週期 (Lifetime)**：從分配後持續存在，直到整個使用者程式結束或透過 API 提前釋放。它與應用程式或 CPU 程序的生命週期相同。
* **用途**：由於其生命週期長，可以在不同 Kernel 之間共享資料，避免 CPU 和 GPU 之間頻繁的資料搬移。
* **大小限制**：有可使用上限，例如 32GB。
* **分配方式**：通常透過 API 函數（如 `cudaMalloc`）進行分配。

2. **Shared Memory (共享記憶體)**：
* 每個區塊專用的記憶體。每個區塊有其獨立的共享記憶體。
* **生命週期**：僅限於其所屬區塊的生命週期或 Kernel 的生命週期。當 Kernel 執行完畢，區塊和其記憶體資源會被釋放。
* **範圍**：僅限於同一個區塊內的執行緒可以存取。不同區塊的執行緒即使在同一個 SM 上，也無法存取彼此的共享記憶體。
* **大小限制**：以區塊為單位定義最大可使用大小。一個區塊內所有執行緒的共享記憶體用量總和不能超過這個上限。
* **分配方式**：可以在 Kernel 程式碼中宣告 (Compile-time 決定)，或在 Kernel 啟動 (launch kernel) 時指定大小 (Run-time 決定，但必須在 Kernel 啟動前決定)。

3. **Local Memory / Registers (區域記憶體/暫存器)**：
* 每個執行緒專用的記憶體。
* **生命週期**：僅限於該執行緒的生命週期。當執行緒或其所屬區塊/Kernel 結束時釋放。
* **範圍**：只有該執行緒本身能看到和存取，即使是同一個區塊內的執行緒也無法看到彼此的暫存器或區域記憶體資料。
* **大小限制**：
* **暫存器 (Registers)**：以區塊為單位有上限 (例如 65536 個暫存器)。開啟的執行緒越多，每個執行緒能用到的暫存器就越少。
* **區域記憶體 (Local Memory)**：來源資料未明確說明其上限，但應與暫存器類似，受限於 SM 資源。

4. **Constant Memory (常數記憶體)**：
* 與 Global Memory 位於同一塊空間。
* **特性**：**唯讀 (read-only)** 記憶體。
* **優勢**：因為是唯讀，資料被存取一次後可以被快取 (cache) 在 SM 的記憶體中，之後的存取會非常快。這避免了處理資料同步和一致性的問題。

#### 5.記憶體分配與資源管理

* **多數限制以區塊為單位**：許多記憶體和計算資源的限制都是以區塊為單位。當程式需要更高的平行度或更多記憶體資源時，通常需要使用多個區塊。
* **資料分割**：由於這些上限的存在，必須將資料進行分割以適應可用的空間。
* **編譯時檢查 (Compile-time Check)**：若程式碼超過這些限制，通常會在**編譯時 (compile time)** 被檢測出來並報錯，而不是在執行時 (run time) 崩潰。
* **GPU 資源管理特性**：
* GPU 是硬體，沒有像 CPU 那樣的作業系統 (OS) 或複雜的動態記憶體管理輔助。
* 大部分資源配置是在編譯時預先決定和靜態分區 (static partition) 的。
* 若在執行時才進行動態記憶體分配，GPU 可能沒有足夠的空間，也無法處理，導致崩潰。因此，許多設定在編譯時就已決定。
* **最佳實踐**：建議查閱 GPU 的規格 (compute capability) 並仔細計算，盡量使用接近上限的資源以提高平行效率。

#### 6.CUDA 程式設計基本術語

* **Host (主機)**：指 CPU。
* **Device (裝置)**：指 GPU。
* **Kernel (核心函數)**：在 GPU 上執行的函數。
* **Thread (執行緒)**：GPU 上最基本的執行單位。
* **Block (區塊)**：一組執行緒，會在某一個 SM 上執行。
* **Grid (網格)**：由許多區塊組成，代表一個 Kernel 的總體執行範圍。

#### 7.CUDA 語言擴充與 API

1. **Kernel 啟動參數 (Launch Kernel Arguments)**：
* 使用三箭頭 `<<<...>>>` 語法啟動 Kernel 函數。
* 共有四個參數，但後兩個是可選的，有預設值。
* **第一個參數 (`num_blocks`)**：指定網格中的區塊數量 (number of blocks)。必須填寫。
* **第二個參數 (`num_threads_per_block`)**：指定每個區塊中的執行緒數量 (number of threads per block)。必須填寫。
* **第三個參數 (`shared_memory_size`)**：可選。指定每個區塊動態分配的共享記憶體大小，以位元組 (bytes) 為單位。這個值必須是**具體的數字** (compile-time constant)，不能是變數。
* 原因：GPU 記憶體管理在 Kernel 執行期間不允許動態變動，最晚在 Kernel 啟動瞬間決定。
* 編譯器會檢查共享記憶體用量是否超過上限。
* **第四個參數 (`stream_ID`)**：可選。指定 Kernel 將在哪個串流 (stream) 中執行。預設值為 0。
* 用途：用於管線化 (pipeline) 優化技巧，讓 GPU 在載入時實現管線化，並可以重疊記憶體傳輸和 Kernel 執行時間。

2. **內建變數 (Built-in Variables)**：
* **`threadIdx`**：執行緒在區塊內的索引 (thread index)。
* **`blockIdx`**：區塊在網格內的索引 (block index)。
* **`blockDim`**：區塊的維度 (block dimension)，即每個區塊有多少執行緒。
* **`gridDim`**：網格的維度 (grid dimension)，即網格有多少區塊。
* **特性**：
* 這些變數**無需宣告**，直接使用即可。
* 它們只在 Kernel 函數（GPU 函數）內部有定義和被賦值。
* `blockDim` 和 `gridDim` 的值對應於 Kernel 啟動時指定的 `num_threads_per_block` 和 `num_blocks`。

3. **多維度索引 (Multi-dimensional Indexing)**：
* 上述內建變數 (`threadIdx`, `blockIdx`, `blockDim`, `gridDim`) 實際上是 `dim3` 結構體，可以支援一維、二維或三維的索引。
* **設計原因**：方便處理二維或三維的資料 (如矩陣)，因為 GPU 最初設計用於圖形處理。
* **限制**：最多支援三維。
* **宣告方式**：若要使用多維度，Kernel 啟動參數也需使用 `dim3` 型別來宣告，例如 `dim3 grid(2, 3, 1);`。
* **存取分量**：使用 `.x`, `.y`, `.z` 來存取各個維度的索引，例如 `threadIdx.x`。
* **資料索引計算**：當 Kernel 中的所有執行緒操作一個共享陣列時，會使用 `blockDim`、`blockIdx` 和 `threadIdx` 等變數來計算資料元素的實際索引。

4. **內聯函數 (Intrinsic Functions)**：
* 以雙底線 `__` 開頭的函數，例如 `__syncthreads()`。
* **執行位置**：只在 GPU 上執行，只能在 Kernel 函數內部呼叫。
* **`__syncthreads()` (區塊內同步)**：
* 作用：在一個區塊內進行執行緒同步。
* 必要性：如果區塊大小超過 32 個執行緒 (即包含多個 Warp)，那麼跨 Warp 的執行緒之間是不同步的，需要呼叫 `__syncthreads()` 才能保證同步。
* 限制：**無法支援跨區塊的同步**。因為跨區塊同步會導致 SM 之間引入同步點，大幅降低 GPU 執行效率，甚至可能比 CPU 還慢，所以硬體層面不提供此功能。

5. **函數修飾符 (Function Qualifiers)**：
* **`__device__`**：
* 在 GPU 上執行，只能由其他 GPU 函數呼叫（不能從 CPU 或 Kernel 函數直接呼叫）。
* 通常用於 Kernel 函數內部的**輔助函數**。
* **`__global__`**：
* Kernel 函數，在 GPU 上執行，只能由 CPU (host) 呼叫。
* 一個 `__global__` 函數不能從另一個 `__global__` 函數內部呼叫（動態平行 (dynamic parallelism) 是一個例外，它允許遞迴呼叫其他 Kernel 函數）。
* **`__host__`**：
* 在 CPU (host) 上執行。
* 如果函數沒有指定任何修飾符，則預設就是 `__host__`。通常不需顯式寫出。
* **`__host__ __device__`**：
* 表示該函數會有兩個版本，一個用於 CPU，一個用於 GPU。
* 這種情況較少見。

6. **變數修飾符 (Variable Qualifiers)**：
* 用於指定變數被分配到何種類型的硬體記憶體上。
* **`__device__`**：
* 表示變數分配在 **Global Memory** 上。
* 在 Kernel 程式碼中宣告，分配在編譯時決定。
* **`__constant__`**：
* 表示變數分配在 **Constant Memory** 上。
* 唯讀，可快取，用於加速存取。
* **`__shared__`**：
* 表示變數分配在 **Shared Memory** 上。
* 分配方式有兩種：在 Kernel 程式碼中宣告 (Compile-time 決定) 或在 Kernel 啟動時指定 (Run-time 決定，但在 Kernel 啟動前確定)。

---
其他課程連結

[平行程式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)