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

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

14C~14D Heterogeneous Computing

今日大綱
GPU 硬體與效能指標
Compute Capability (計算能力)
CUDA 介紹
CUDA 執行流程 (Execution Flow)
CUDA 程式設計模型 - Kernel 與執行模型
軟體與硬體架構的對應關係 (Software-Hardware Mapping)

#### **1. GPU 硬體與效能指標**

* **記憶體頻寬 (Memory Bandwidth)**:
* **GPU Global Memory Bandwidth**: 2025 年最新旗艦 GPU 如 NVIDIA Blackwell 系列 (例如 B200) 已達到高達 8 TB/s 的 HBM3e 記憶體頻寬。
* **Inconnect Bandwidth (對外)**: 指 GPU 與主機記憶體或 CPU 之間的連接頻寬。
* 2025 年，PCIe 5.0 提供高達 128 GB/s 的頻寬，但仍遠低於 GPU 內部記憶體頻寬
* 為突破限制，引入 **NVLink**: 另一種互連頻寬技術，**只能支援 GPU 之間**。速度非常快，從 300GB/s 提升到 2025 年的 NVLink 5.0 版本，提供高達 1.8 TB/s 的雙向頻寬。價格也更高。
* **計算能力 (FLOPS)**: 衡量 GPU 的浮點運算能力。
* 主要關注 **單精度 (single precision)** 和 **雙精度 (double precision)** 浮點運算。
* 不同精度的表現會有落差。
* 例如，GeForce 系列 (如 RTX 5090) 在單精度表現上可達約 90-100 TFLOPS，與數據中心 GPU (如 H100) 差異不大，但雙精度則相差非常多 (GeForce 雙精度約 1-2 TFLOPS，而 H100 可達 34 TFLOPS)。
* **低精準度計算 (如深度學習)**: 傾向使用一般顯示卡，因為性價比高 (CP值好、價格低、效能差異不大)。
* **高精準度 HPC (高性能計算)**: 則必須使用數據中心系列 GPU (如 H100 或 Blackwell 系列)，因其對精準度要求很高。
* **GPU 架構演進與技術**:
* 每代架構都有很大變化，持續改進。
* **內部記憶體** 越來越大: 受製程奈米數越來越小 (如4奈米甚至更小) 的影響。
* **3D 堆疊記憶體 (3D stacking memory)**: 在 Kepler (cal) 之後出現。不僅使記憶體變大，更重要的是**大幅提升記憶體頻寬**，因為距離變短，不需線性化。目前在 Blackwell 架構中繼續優化，使用 HBM3e 技術。
* **NVLink**: 從 Pascal (PC) 時代引入，現已發展到 2025 年的 5.0 版本，並引入 NVLink Fusion 技術，允許與第三方 CPU 或加速器整合，提供更高系統級頻寬 (單一 NVLink Spine 可達 130 TB/s)。
* **Tensor Core**: 針對深度學習優化，從 Volta (VTA) 架構開始引入。使用 Tensor Core 可以大幅提升深度學習的效能。在 Blackwell 架構中已演進到第五代，提供更高效率的矩陣運算。
* **MIG (Multi-Instance GPU)**: Ampere 架構引入的最強功能。可以將一張 GPU 卡**切分成七張獨立運作的卡**。旨在解決計算資源閒置問題，讓多個程式同時運行，提高資源利用率。在 Hopper 和 Blackwell 架構中繼續支援並優化。

#### **2. Compute Capability (計算能力)**

* **定義**: 是一個定義集，說明每個 Compute Capability 所支援的功能 (feature)。
* **設定上限**: 包含許多設定，例如最大值等上限設定。
* **軟硬體中介**: 介於軟體和硬體之間。
* 硬體架構會說明它支援的 Compute Capability。
* CUDA 軟體版本也會說明它支援的 Compute Capability。
* **重要性**: 對於寫作業和最佳化設定非常重要，因為需要根據這些上限來設定最佳參數。
* **查詢方式**: 可以透過 `deviceQuery` 等指令查詢，或上網查詢，或詢問助教。
* **版本更新**: CUDA 的 Compute Capability 會不斷上升，以適應新的硬體架構。在建立環境時需留意版本和 Compute Capability。目前最新為 Compute Capability 10.0 (Blackwell 架構)，支援 CUDA 13.0 以上的版本。

#### **3. CUDA 介紹**

* **什麼是 CUDA?**:
* 是 "Compute Unified Device Architecture" 的縮寫。
* **不只是一種程式語言**: 它是一個完整的軟體開發套件 (SDK)，包含工具包 (Tool Kits)、函式庫 (Library)、編譯器 (Compiler) 等所有與程式設計相關的組件。
* 在電腦上安裝 CUDA 環境時，通常需要安裝相關的工具、函式庫、驅動程式和編譯器。
* 它是一種 **API 擴展 (API extense)**。
* 支援 C 和 C++: 程式設計師**不需要學習新的程式語言**，只需學習 C/C++ 中的 CUDA API 即可。
* **程式設計基礎**:
* 基於 **SIMD (Single Instruction, Multiple Data)** 或 **資料平行 (data parallelism)** 模型。
* 從硬體角度看是 SIMD 架構，從軟體角度看是資料平行。
* **版本更新與挑戰**:
* CUDA 版本更新頻繁，目前最新為 CUDA 13.0 (2025 年 8 月發布)。
* 每個版本的 API 和功能都有很大變化。
* 這些變化可能對底層開發 (例如 API 攔截、驅動程式運作) 造成麻煩或帶來便利。
* 硬體的不斷演進 (GPU 一直在變) 導致軟體端也有許多變革，因為許多機制與硬體緊密綁定。
* 例如，CUDA 11.0 的連結方式不同，可能導致原有的 API 攔截方法失效。
* 支援 GPU 分割 (partition) 或統一記憶體 (unified memory) 等機制，也可能在不同版本中有重大變化。
* **持續關注更新**: 如果從事 GPU 相關開發，需要實時關注新的功能和 API，才能大幅提升效能。
* **函式庫的重要性 (Software Stack)**:
* 計算硬體提供運算能力，但其**軟體堆疊 (software stack) 非常重要**。
* NVIDIA 提供了非常多的函式庫 (e.g., 深度學習函式庫、通訊函式庫 PTT CE)。
* 許多時候是**透過這些函式庫來使用 GPU**，而非直接手寫 GPU 程式碼。
* 只有在做底層開發、優化，或理解函式庫內部運作時，才需要更深入的 CUDA 知識。
* 如果需要開發支援 GPU 的軟體或進行程式移植 (porting)，就需要學習如何撰寫 CUDA 程式。

#### **4. CUDA 執行流程 (Execution Flow)**

CUDA 程式遵循異質計算 (heterogeneous computing) 模型，假設主機 (host) 和設備 (device) 完全獨立運作。

1. **資料複製到設備 (Copy data to device)**:
* 必須將計算任務所需的記憶體資料從主機記憶體 (host memory, CPU) 複製或移動到 GPU 設備上的記憶體 (GPU memory / global memory)。
* 這是讓 GPU 能夠存取資料的第一步。
2. **啟動 Kernel (Launch Kernel)**:
* 啟動一個在 GPU 上執行的函式，稱為 **Kernel**。
* 啟動 Kernel 時，會將 Kernel 的程式碼及其相關的上下文 (context)，包括記憶體資料位置、執行配置 (configuration)、中繼資料 (metadata) 等，全部傳送到 GPU 上，準備執行條件。
* Kernel 呼叫是**非同步 (asynchronous)** 或非阻塞 (non-blocking) 的: CPU 呼叫後會立即回傳，可以繼續執行其他任務。
3. **設備上執行 Kernel (Execute Kernel on device)**:
* Kernel 函式在 GPU 上執行時，CPU 無法干預或中斷。
* Kernel 會在 GPU 上自行完成所有運算。
* GPU 通常會等到整個 Kernel 函式執行完畢後，才會通知主機任務已完成。主機與設備之間的同步化方式比較複雜，是非同步溝通。
4. **資料複製回主機 (Copy data back to host)**:
* Kernel 執行完畢後，還需要將結果資料從 GPU 記憶體複製回主機記憶體。
* **優化**: 上述是基本流程，但存在優化機制，例如讓 GPU 不需複製資料直接存取主機記憶體，但這些優化通常有缺點，且不屬於基本的執行流程。
* **重要性**: 了解資料在 host 和 device 之間傳輸的重要性。

#### **5. CUDA 程式設計模型 - Kernel 與執行模型**

* **Kernel**:
* 是一個被卸載 (offload) 到 GPU 上執行的函式。
* 它本身就是一個**平行程式**，由 GPU 的多個執行緒 (threads) 執行。
* 程式碼中，Kernel 的宣告會包含 `__global__` 關鍵字。
* **程式碼結構**:
* Kernel 函式在 `main` 函式中被呼叫時，會使用三層括號 `<<<>>>` 定義執行配置。
* 例如：`kernel_function_name<<<n_blocks, block_size>>>(arguments)`。
* `n_blocks` 和 `block_size` 是兩個數字，相乘等於執行該 Kernel 函式所需的總執行緒數量。
* Kernel 函式定義時，其簽名 (declaration) 前面會加上 `__global__` 關鍵字。
* `__global__` 表示這個函式碼要被編譯成 GPU 指令，並可從 CPU 呼叫到 GPU 執行。
* 還有其他關鍵字 (如 `__device__`, `__host__`)，用於定義不同執行位置的函式。
* **資料平行模型**:
* 所有 GPU 執行緒都執行**相同的程式碼 (Single Program, Multiple Data - SPMD)**。
* 每個執行緒有唯一的 **執行緒 ID (thread ID)**。類似於 MPI process 的 rank ID。
* 執行緒透過其 ID 來區分處理不同的資料。
* 如果多個執行緒由同一個物理處理器執行，它們可能同時執行相同的指令。但若跨不同處理器，執行時間點可能不同。
* **分歧 (Divergence)**:
* 執行緒可以透過分支 (branch) 執行不同的指令。
* 但在 GPU 上，這樣做的代價非常高，會導致**效能下降 (diversion problem)**，甚至引入多餘的計算。
* **執行緒排程與切換 (Context Switching)**:
* GPU 上的平行執行緒數量通常遠多於其核心數。
* 當任何一個執行緒遇到記憶體存取延遲 (memory stall) 時，該執行緒就會被切換出 (context switch)，換到另一組沒有記憶體延遲的執行緒，以持續利用 GPU 的計算資源。
* 記憶體存取時間遠高於核心計算週期，因此執行緒切換非常頻繁，需要大量過度訂閱 (oversubscribe) 執行緒數量 (可能達10倍)。
* GPU 核心通常有各自的**暫存器檔案 (register file)**: 不像 CPU OS 需要在切換時複製記憶體資料，GPU 核心的切換速度快很多。
* 這類執行緒通常稱為**物理執行緒 (physical threads)**。
* **執行緒階層 (Thread Hierarchy)**:
* 程式設計師在啟動 Kernel 時，可決定平行度。
* CUDA 採用**兩層階層式 (2-level hierarchical)** 描述。
* 一個 Kernel 執行時，被稱為一個 **Grid (網格)**:
* Grid 由多個 **Thread Blocks (執行緒塊)** 組成。
* 每個 Thread Block 由多個 **Threads (執行緒)** 組成。
* 在啟動 Kernel 時，`<<<n_blocks, block_size>>>` 中的 `n_blocks` 描述 Grid 中的 Block 數量，`block_size` 描述每個 Block 中的 Thread 數量。
* 每個 Block 有其唯一的 **Block ID**。
* 描述一個 Thread 時，需要同時指明其所在的 Block ID 和 Thread ID。
* **同步化 (Synchronization)**:
* 不同 Block 中的執行緒**無法保證同時執行**，執行時間可能錯開，甚至可能被序列化 (sterilized)。
* GPU 雖然是 SIMD 處理器，但**跨 Block 的執行緒無法使用 API 進行同步**。
* 程式設計師**只能做到讓同一 Block 內的執行緒同步**，透過 `__syncthreads()` API 實現。
* 這在程式設計上需要調整和技巧，是優化中會介紹的內容。

#### **6. 軟體與硬體架構的對應關係 (Software-Hardware Mapping)**

軟體模型與硬體架構之間有明確的對應關係。

* **Grid (網格)**:
* 對應到一個 **Device (設備/GPU 卡)**。
* 一個 Kernel 或 Grid 只能被丟到一個 Device 執行，這是一對一的關係。
* 一個 Device 可以運行多個 Kernel，但很少這麼做，因為 GPU 缺乏像 OS 一樣複雜的記憶體管理，多個 Kernel 共用 GPU 記憶體需要使用者程式自行處理。
* **Block (執行緒塊)**:
* 對應到一個 **Streaming Multiprocessor (SM) / SD 處理器**。
* 一個 SM 可以運行多個 Blocks，這不是一對一的關係。
* 不同的 Block 可能會被派送到不同的 SM (或不同的 SD 處理器) 上。
* **Thread (執行緒)**:
* 對應到一個 **Core (核心)**。
* 一個 Core 可以同時運行多個 Threads (透過頻繁的執行緒切換)。
* **重要性**: 程式設計模型抽象化了硬體架構，因此**理解硬體架構**將有助於更容易理解和想像軟體程式如何被執行，並進行優化。

---
其他課程連結

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