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

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

16A~16B What is Compute Unified Device Architecture(CUDA)?

今日大綱
1.指標與記憶體管理
2.NVCC 編譯流程與 GPU 程式結構
3.GPU Kernel 執行與執行緒管理
4.GPU 效能優化
5. 演算法應用範例：Floyd-Warshall

### GPU 程式設計基礎概念與應用

#### 1. 指標與記憶體管理

* **裝置指標 (Device Pointer)**：指向 GPU 全域記憶體 (global memory) 上的位置。
* **函數回傳值限制**：在 GPU 上執行的函數（例如 Kernel）通常不能直接回傳複雜的數值。若要「回傳」數值，必須透過**指標參數 (pointer argument)** 的方式來儲存結果。這些裝置指標的記憶體位置只要不被釋放，就可以在多個 Kernel 呼叫之間持續存在。
* **記憶體複製 (cudaMemcpy)**：
* 用於在 **Host (CPU)** 和 **Device (GPU)** 之間進行資料傳輸，支援**雙向**複製 (Host to Device, Device to Host)。
* 也可以在 **Device 之間**進行資料複製 (Device to Device)。
* `cudaMemcpy` 函數需要指定：目標位址 (`destination`)、來源位址 (`source`)、資料大小 (`size_in_bytes`) 以及複製方向 (`cudaMemcpyKind`)。
* **關鍵注意事項**：`destination` 和 `source` 的指標類型（是裝置指標還是主機指標）決定了 `cudaMemcpyKind` 參數的正確選擇。若為裝置指標，則需使用 `cudaMemcpyDeviceToDevice` 或適當的跨記憶體類型複製；若為主機指標，則通常是跨記憶體類型複製。這部分需要非常留意指標的正確性。

#### 2. NVCC 編譯流程與 GPU 程式結構

* **函數類型修飾符 (Specifiers)**：`__global__`、`__device__` 和 `__host__` 用於指定函數在哪裡執行以及可從哪裡呼叫。這些修飾符告知編譯器該函數應該生成哪種程式碼，若使用不當會導致編譯錯誤。
* **NVCC 編譯器特性**：
* **兩階段編譯**：NVCC 編譯器並非一次性將程式碼編譯成最終的二進位碼。它會先將程式碼編譯成**PTX (Parallel Thread eXecution)** 這種中間表示語言。
* **SASS 轉換**：PTX 碼再由另一個編譯器轉換成特定 GPU 架構的 **SASS (assembler code)**。
* **兩階段原因**：這是因為 GPU 加速器（尤其是 GPU）的硬體架構和指令集變化非常快速。PTX 作為一個**中介且與硬體架構無關**的表示層，允許底層的 SASS 編譯器根據不同的 GPU 型號生成優化的程式碼，而無需每次都重新編譯整個原始程式碼。這類似於 Java 的位元組碼 (bytecode) 概念。
* **產生兩種程式碼**：NVCC 編譯器會生成一份在 CPU 上執行的程式碼，以及一份在 GPU 上執行的程式碼。
* **GPU 程式執行流程**：程式碼經由 NVCC 編譯後，產生的二進位碼會被 GPU 驅動程式載入到主機記憶體中，然後驅動程式再將其載入到 GPU 的記憶體中執行。

#### 3. GPU Kernel 執行與執行緒管理

* **基本概念**：Kernel 是在 GPU 上執行的函數。早期的 `__device__` 變數只能在裝置記憶體中，但現在的 `__global__` 函數也能使用裝置變數。
* **Kernel 函數參數**：在 `__global__` 函數中，若要傳遞資料到 GPU 上進行處理，其參數必須是**裝置指標**。
* **正確的 GPU 程式流程範例（陣列加法）**：
1. **宣告主機與裝置陣列**：同時定義主機端 (`h_A`, `h_B`, `h_C`) 和裝置端 (`d_A`, `d_B`, `d_C`) 的陣列。
2. **裝置記憶體配置**：使用 `cudaMalloc` 在 GPU 裝置記憶體上為 `d_A`, `d_B`, `d_C` 分配空間，並取得裝置指標。
3. **資料傳輸（主機到裝置）**：使用 `cudaMemcpy(d_A, h_A, ..., cudaMemcpyHostToDevice)` 等，將主機資料複製到裝置上。
4. **Kernel 啟動**：呼叫 `__global__` 函數，並將裝置指標 (`d_A`, `d_B`, `d_C`) 作為參數傳入。
5. **資料傳輸（裝置到主機）**：Kernel 執行完畢後，使用 `cudaMemcpy(h_C, d_C, ..., cudaMemcpyDeviceToHost)` 將結果從裝置複製回主機。
6. **裝置記憶體釋放**：使用 `cudaFree` 釋放裝置記憶體。
* **多元素處理與平行化**：為充分利用 GPU 的平行處理能力，Kernel 通常會處理多個資料元素，例如陣列的相加。
* **索引計算 (Indexing)**：
* GPU 上的每個執行緒 (`threadIdx.x`) 結合其所屬的區塊 (`blockIdx.x`)，可以計算出一個全域索引 (global index) 來識別要處理的資料元素。
* 常見的 1D 索引計算公式為：`global_idx = blockIdx.x * blockDim.x + threadIdx.x`。
* **索引方式對效能的影響**：不同的索引計算方式（例如 `blockIdx.x * blockDim.x + threadIdx.x` 與 `threadIdx.x * gridDim.x + blockIdx.x`）會導致**不同的記憶體存取模式**。
* **連續記憶體存取**對 GPU 效能至關重要，因為它能促進記憶體合併 (memory coalescing)。不連續的存取模式會導致更多 I/O 操作，顯著降低效能。
* **區塊限制 (Block Limits)**：GPU 對於每個區塊可包含的執行緒數量以及每個網格 (grid) 可包含的區塊數量都有硬體限制。如果處理的資料量非常大，可能需要將其拆分成多個區塊來處理。
* **處理無法整除的資料大小**：
* 當總資料量 (`N`) 無法被執行緒或區塊數量完全整除時，一些執行緒可能會嘗試存取**超出有效範圍**的記憶體。
* **解決方法**：將實際的有效資料大小 `N` 作為參數傳遞給 Kernel。在 Kernel 內部，每個執行緒在執行計算前應**檢查其全域索引是否小於 `N`** (`if (global_idx < N)`)，以避免越界存取。

#### 4. GPU 效能優化

* **分支 (Branching) 對 GPU 效能的影響**：
* GPU 的基本架構是 **SIMD (Single Instruction, Multiple Data)** 或 **SIMT (Single Instruction, Multiple Threads)**。相同 Warp (通常是一組 32 個執行緒) 中的執行緒會同步執行相同的指令。
* **Warp 分歧 (Warp Divergence)**：如果同一個 Warp 內的執行緒由於 `if/else` 條件判斷而採取了不同的執行路徑（例如，部分執行緒進入 `if` 區塊，部分進入 `else` 區塊），GPU 處理器會讓**所有執行緒都執行這兩個分支的所有指令**。對於那些不符合條件的執行緒，其計算結果會被捨棄，這實際上**浪費了大量的計算資源**。
* **優化建議**：應盡量**避免在 Kernel 中使用分支語句**，或者將計算重寫為無分支的形式。例如，透過修正索引或使用其他數學技巧，即使可能需要執行更多的算術操作，但只要能確保 Warp 內執行緒的指令流一致，效能通常會更好。這是 GPU 編程優化中的一個重要考慮點。

#### 5. 演算法應用範例：Floyd-Warshall

* **演算法概述**：Floyd-Warshall 演算法是一種用於尋找**所有節點對之間最短路徑**的演算法。它使用鄰接矩陣來表示圖中節點之間的權重。
* **演算法特性**：
* 透過迭代的方式不斷更新最短路徑距離，直到沒有任何路徑可以進一步縮短為止。
* 其核心思想是，對於任意兩個節點 `i` 和 `j`，嘗試透過一個中間節點 `k` 來尋找更短的路徑 (`dist[i][j] = min(dist[i][j], dist[i][k] + dist[k][j])`)。
* **適合 GPU 平行化**：
* Floyd-Warshall 演算法非常**適合在 GPU 上進行平行化**。
* 其計算結構包含三層巢狀迴圈 (`k`、`i`、`j`)。在最內層，每個 `(i,j)` 對都執行相同的比較和更新操作，這是一個典型的**資料平行 (data-parallel)** 模式，其迭代的性質和每個迭代中固定的計算模式，非常符合 GPU 的 SIMD/SIMT 架構。
* 儘管 Floyd-Warshall 可能會執行一些冗餘計算，但由於其在每次迭代中**所有執行緒的存取模式都非常一致**，因此能夠實現高效的平行處理。
* **與 Bellman-Ford 的對比**：相較於 Bellman-Ford 等演算法，Floyd-Warshall 更適合 GPU。Bellman-Ford 的不規則更新和不確定的新增/移除操作，使得其記憶體存取模式更難以預測和平行化。

---
其他課程連結

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