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

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

16C~16D What is Compute Unified Device Architecture(CUDA)?

今日大綱
1.GPU程式碼轉換與最佳化
2.GPU與CPU的效能特性與考量
3.非同步執行與同步化機制
4.效能測量與時間計算

### **GPU程式碼轉換與最佳化**

1. **目標**：將包含多層 `for` 迴圈（例如三個）的程式碼放到GPU上執行，目的是利用不同的執行緒同時進行迴圈內的計算，以達成加速。
2. **K迴圈的特殊考量**：
* K的意義在於「**刷新最後的結果**」，確保每一次更新後，資料能從任何一個點遍歷到其他所有點。
* 在圖論演算法（如Floyd-Warshall）中，K代表的更新次數應至少為N（點數），因為N次迭代後，無論是誰，最後都能夠到達目的地。
* **K迴圈的循序性 (Sequential Nature)**：K的迭代本質上是循序的。我們需要確保第一輪的所有更新（例如針對第一個K值）都完全完成後，才能開始計算第二輪（針對第二個K值）的結果。
* **將K迴圈移入GPU內核 (Kernel) 的問題**：
* 如果將K迴圈移動到內核內部，可能導致錯誤，無法保證結果的正確性。
* 原因在於，若在第一輪的更新尚未完全完成時，就開始進行下一輪的計算，可能會取到舊的、尚未被新資料更新的值，這會使得最短路徑的更新無法被保證。也就是說，你必須「刷過一頁」的結果後，才能進行下一頁的刷新。
* **解決方案 (`__syncthreads()`)**：
* 在GPU內核內部使用 `__syncthreads()` 可以解決這個問題。它會強制在同一個**區塊 (block)** 內的所有執行緒，必須等待當前K值的計算完成後，才能繼續進行下一個K值的計算。
* **限制**：`__syncthreads()` 的同步範圍僅限於**單一區塊內**。它無法用於不同區塊之間的同步，也無法解決跨多個GPU設備的問題。

### **GPU程式碼重構策略**

1. **策略一：展開I與J的維度**
* 將I和J兩個 `for` 迴圈的維度展開，使每個執行緒負責更新矩陣中的一個元素。
* 這種方法透過讓每個執行緒ID對應到一個特定的元素來加速計算。
2. **策略二：嘗試將K移入內核並展開（錯誤策略）**
* 這種做法是將K值移入GPU內核，並嘗試與I、J一同展開。
* **問題**：這並沒有實質上增加演算法的平行度，因為它仍然只有兩個維度的執行緒（例如 `gridDim.x` 和 `gridDim.y`）在處理，而不是三個維度。同時，如前所述，它會導致計算結果不正確。

### **GPU與CPU的效能特性與考量**

1. **GPU專長**：
* GPU專門用於執行高度平行化的運算，例如加法和乘法。這些運算在GPU上執行速度非常快。
* 在GPU上，許多運算（例如開根號、除法、取餘數）可能非常慢，應盡量避免。程式碼的改寫常常是為了避開這些低效率的運算。
2. **記憶體存取**：
* GPU在存取連續記憶體時表現極佳。
3. **主機 (Host) 與設備 (Device) 通訊**：
* 盡量減少主機（CPU）與設備（GPU）之間的資料傳輸，因為這是影響效能的瓶頸之一。
4. **資料量考量**：
* GPU通常用於處理**大量資料**和**大規模計算**。
* 若資料量非常小，GPU啟動和資料傳輸的開銷可能使其比CPU更慢。在這種情況下，直接使用CPU進行運算反而可能更有效率。

### **非同步執行與同步化機制**

1. **CUDA函式的非同步性**：
* 幾乎所有的CUDA核心啟動 (kernel launch) 函式對主機 (Host) 而言都是**非同步 (asynchronous)** 的。這意味著主機發送GPU任務後，會立即返回並繼續執行自己的其他任務，而GPU則在背景平行地執行計算。
* 這種特性使得CPU和GPU可以同時工作，提高整體系統利用率。
2. **`cudaMemcpy` 的同步性例外**：
* `cudaMemcpy`（資料複製）函式是一個重要的例外，它預設是**同步 (synchronous)** 的。
* 當主機執行 `cudaMemcpy` 時，它會等待資料傳輸完成後才返回。這是因為對於GPU而言，`cudaMemcpy` 本身也是一個必須完成的執行任務，它必須先完成，後續的GPU核心才能正確啟動。
3. **Host-Device 同步方法**：
* 有**三種**主要的同步方式（外加Event提供更多彈性）。
* **`cudaDeviceSynchronize()`**：
* 此函式會阻塞主機，直到**所有設備上**（或通過 `cudaSetDevice()` 指定的設備上）所有先前發布的CUDA任務都完成為止。
* 若有多張GPU卡，`cudaDeviceSynchronize()` 會根據當前設定的設備（使用 `cudaSetDevice()`）來同步。例如，如果設定為設備0，則只會等待設備0上的任務完成，不會阻塞到設備1。
* 目的：確保所有GPU工作已完成，防止GPU過載或資料不一致。
* **`cudaStreamSynchronize(stream)`**：
* 此函式會阻塞主機，直到**指定串流 (stream)** 上所有先前發布的CUDA任務都完成為止。
* 它允許更細粒度的同步，因為主機可以等待特定串流完成，而其他串流可能繼續執行，實現部分任務的平行執行。
* **`cudaEventSynchronize(event)`**：
* 這是最具彈性的同步方式。
* 首先，使用 `cudaEventRecord(event, stream)` 在GPU執行流程中的**任何指定點**記錄一個事件。
* 接著，主機可以呼叫 `cudaEventSynchronize(event)`。這個呼叫會阻塞主機，直到該特定的事件在GPU上被記錄完成為止。
* `cudaEventRecord()` 本身是異步的，只有 `cudaEventSynchronize()` 會導致主機阻塞。
* 用途：**精確控制同步點**，並常用於測量GPU特定任務的執行時間。
* **`cudaSetDevice(deviceId)`**：
* 這個函式用來指定當前所有後續CUDA操作將會在哪一個GPU設備上執行。

### **效能測量與時間計算**

1. **使用CUDA事件進行計時**：
* **最佳方法**：使用CUDA事件 (events) 是測量GPU執行時間的最準確方式。
* **步驟**：
1. 創建兩個CUDA事件：`start` 和 `stop`。
2. 在要測量的GPU工作（例如核心啟動）之前記錄 `start` 事件：`cudaEventRecord(start, 0)`。
3. 在GPU工作之後記錄 `stop` 事件：`cudaEventRecord(stop, 0)`。
4. **關鍵步驟**：呼叫 `cudaEventSynchronize(stop)` 來**阻塞主機**，直到 `stop` 事件在GPU上被記錄完成為止。這確保了所有被測量的GPU工作都已完成。
5. 最後，使用 `cudaEventElapsedTime(&elapsed_time, start, stop)` 計算 `start` 和 `stop` 事件之間經過的時間。
* **重要注意事項**：
* 由於GPU核心啟動和事件記錄本身都是非同步的，**必須**在計算時間之前加入 `cudaEventSynchronize()`（或 `cudaDeviceSynchronize()`）。
* 如果沒有同步，主機可能會在GPU任務尚未完成（甚至尚未開始）時就計算時間，導致測量結果不準確（例如得到接近零的時間）。

### **GPU內部的同步**

1. **GPU處理器內部的任務排程**：
* 在GPU內部，當多個任務提交給它時，它會強制這些任務依序執行，除非透過特定的複雜策略（如使用串流）來實現重疊執行。
* 這意味著在預設情況下，GPU會自動處理任務的序列性，開發者通常無需擔心不同核心任務在GPU內部會不小心「重疊」執行，除非有明確的優化需求並使用串流來實現。