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

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

19B~19D 記憶體層級 / CUDA的優化

今日大綱
1.記憶體類別與變數處理
2.共用記憶體
3.全域記憶體與常數記憶體
4.記憶體階層與效能
5.CUDA 優化技巧

### **記憶體類別與變數處理**

1. **局部變數 (Local Variables)**
* **類型與儲存**:
* 在 C 語言中，像 `M`, `N`, `BAC` 這樣的變數，如果沒有特別宣告，通常是 **局部變數 (local variable)**。
* 它們會儲存在 **stack 記憶體**中。
* **動態配置的限制**:
* CUDA 中的局部變數不允許動態配置 (dynamic allocation)。這表示你無法在執行時根據變數大小（例如 `n` 的數量）來配置陣列大小。
* 這是因為編譯器在編譯時就需要知道變數的大小。
* 因此，局部變數的 **size 必須是固定的**。
* **陣列處理**:
* 如果局部變數是陣列 (array)，即使它是局部變數，它也會放在 **local memory** 中，而不是 stack。
* 對局部變數的陣列進行動態配置是不可能的。

#### **補充**
- **局部變數 (Local Variables)**：在 Hopper/Ada 架構中，局部變數若溢出暫存器，仍會置於本地記憶體（等同全域記憶體，延遲高達 400-800 週期）。新優化：使用 CUDA 13.0 的共享記憶體暫存器溢出（Shared Memory Register Spilling），可優先將溢出變數置於共享記憶體，降低延遲並減輕 L2 快取壓力。效能提升可達 5-10%。
- **陣列處理**：動態配置仍受限，但 Hopper 的 Tensor Memory Accelerator (TMA) 支援非同步資料移動，可隱藏陣列載入延遲，尤其適合大型陣列。

### **共用記憶體 (Shared Memory)**

1. **定義與特性**
* **快取 (Cache)**: 共用記憶體的定義上就是一種 **可程式化的快取**。
* **非永久性**: 它是為了加速計算而存在，一旦 block 執行完畢，共用記憶體中的資料就會被清除，不會永久儲存。
* **存取速度**: 比全域記憶體 (Global Memory) 快很多，因為它位於 GPU 的 on-chip memory。
* **管理**: 由程式設計師自己來管理，不像許多硬體快取是自動管理的。
* **大小限制**: 每個 block 能使用的共用記憶體大小是固定的且有上限，例如 48KB。
* 在較舊的 CUDA 版本中，48KB 的共用記憶體空間可能還要與 L1 快取共用。
* 在 Fermi 或 Kepler 架構中，可以透過 API 設定共用記憶體與 L1 快取的比例，例如 1:1 或 1:2，通常預設是 1:2。
2. **使用模式**
* 共用記憶體的主要用途是當資料需要 **重複使用** 時。
* **流程**:
1. 資料最初通常在 **全域記憶體 (Global Memory)** 中。
2. 每個 block 負責特定的資料區塊。
3. 將 block 負責的資料區塊 **複製** 到該 block 的共用記憶體中。
4. block 中的 thread 開始對共用記憶體中的資料執行計算。
5. 計算過程中，所有 thread 頻繁存取共用記憶體，因為速度快。
6. 計算完成後，將結果從共用記憶體寫回全域記憶體。
* **`__syncthreads()` 的重要性**
* 在將資料複製到共用記憶體之後，以及開始計算之前，必須使用 `__syncthreads()`。
* 這是因為所有 thread 都是平行執行的，如果沒有同步，某些 thread 可能會存取到其他 thread 尚未複製完成的資料，導致錯誤的結果。
3. **效能提升**
* 透過將資料從全域記憶體搬到共用記憶體，可以大幅減少全域記憶體的存取次數。
* 例如，一個計算可能將全域記憶體存取次數從六次減少到兩次，速度可提升三倍或更多。
* 這是因為共用記憶體的存取速度比全域記憶體快至少數十倍。
4. **動態共用記憶體 (Dynamic Shared Memory Allocation)**
* **宣告方式**: 使用 `extern __shared__` 宣告，例如 `extern __shared__ int s_array[];`。
* **大小指定**: 大小不是在程式碼中 hard code ，而是在啟動 kernel (launch kernel) 時透過 `<<<gridDim, blockDim, shared_memory_size>>>` 的 **第三個參數** 來動態指定。這個參數是 `shared_memory_size`。
* **每個 block 只能有一個**: 每個 block 只允許 **一個** 動態共用變數。
* 如果你宣告了多個 `extern __shared__` 變數 (例如 `A` 和 `B`)，它們實際上會指向 **相同的記憶體位置**。
* 若要模擬多個動態共用陣列，你需要宣告一個大的單一陣列，然後自己透過索引來劃分邏輯區塊 (logical block)，例如 `s_array[size_A]` 用於 A，`s_array[size_A + size_B]` 用於 B。
* **陣列型別**: 動態共用記憶體只能接受 **一維陣列 (one-dimensional array)**。
* 若需使用多維陣列，程式設計師必須手動進行位址映射 (address mapping)。
* **資料搬移**: 每個 thread 負責搬移它自己需要的一個資料元素到共用記憶體中，而不是整個陣列。所有 thread 完成搬移後，整個資料區塊才算進入共用記憶體。
5. **靜態共用記憶體 (Static Shared Memory Allocation)**
* **宣告方式**: 在 kernel 程式碼內部宣告，並直接指定其固定大小，例如 `__shared__ int s_array;`。
* **限制**: 大小必須是編譯時已知的 **常數 (constant)**，不能是變數。
* **適用情境**: 當你知道資料大小固定，或者空間足夠且願意浪費時使用。
* **優點**: 相對動態配置更簡單。

#### **補充**
- **大小限制與配置**：現代 GPU 已大幅增加共享記憶體容量。
- Ampere (A100)：每個 SM (Streaming Multiprocessor) 可配置高達 164KB（共享與 L1 快取共用）。
- Ada Lovelace (RTX 4090)：每個 SM 128KB。
- Hopper (H100/H800)：每個 SM 256KB（比 Ampere 增加 56%），並可透過 API 動態調整比例（預設 1:1 或 1:3）。Blackwell (B100) 預計進一步擴大至 300KB+。
- 舊版（如 Fermi 48KB）限制已解除；現在每個 block 上限為 64KB，但總容量取決於 SM 資源。
- **新功能**：
- **Distributed Shared Memory (DSM)**：Hopper 獨有，允許跨 SM 通信（inter-SM），延遲僅 181 週期（比全域記憶體低 6.1 倍）。適用於叢集大小 2-8 的圖形或 AI 工作負載，輸送量高達 3.28 TB/s（環狀存取模式下）。
- **Shared Memory Register Spilling**：CUDA 13.0 新增（需在 kernel 定義中加入 `.pragma enable_smem_spilling`）。當暫存器不足時，優先溢出至共享記憶體而非本地記憶體。益處：降低延遲、減少 L2 快取壓力；基準測試顯示 kernel 持續時間縮短 7.76%、SM 活躍週期減少 9.03%。
- **非同步複製 (Asynchronous Copy)**：從 CUDA 11.0 強化，在 Hopper 上更高效。允許從全域記憶體非同步複製至共享記憶體（無需中間暫存器），元素大小 8-16 位元組時效能最佳（比同步複製快 1.5-2 倍）。不需複製計數為 4 的倍數，即可達峰值。
- **效能提升**：在 Hopper 上，使用共享記憶體可將全域存取次數減至 1/10，結合 TMA 可使矩陣乘法 (GEMM) 加速 1.5 倍（大型矩陣 n>7000 時達 600 TFLOPS）。
- **動態/靜態配置**：無變化，但 Hopper 支援 warpgroup-level 操作（128 執行緒），需使用 wgmma 指令確保記憶體一致性（fence 操作）。

### **全域記憶體 (Global Memory) 與常數記憶體 (Constant Memory)**

1. **全域記憶體 (Global Memory)**
* **作用**: 用於主機 (Host) 與設備 (Device) 之間傳輸資料，是 GPU 上最基本且必須使用的記憶體。
* **配置與操作**:
* 透過 `cudaMalloc` 在 GPU 設備上配置記憶體，返回一個 device pointer。
* 透過 `cudaMemcpy` 在 Host 記憶體和 Global Memory 之間複製資料。
* 透過 `cudaFree` 釋放記憶體。
* **生命週期**: 配置的 Global Memory 在 kernel 執行完畢後不會被清除，可以在不同的 kernel 呼叫之間持續存在和存取。
* **存取速度**: 相對於共享記憶體和暫存器較慢。
2. **常數記憶體 (Constant Memory)**
* **作用**: 用於儲存不會改變的常數資料。
* **宣告方式**: 使用 `__constant__` 關鍵字宣告，例如 `__constant__ int const_array[SIZE];`。
* **特點**:
* 它是 **沒有動態配置 (no dynamic allocation)** 的。
* 資料會被快取 (cache)，第二次存取會更快。
* 存取模式類似於將其值複製到局部變數或暫存器中進行計算，以達到更快的速度。
* **效能**: 如果能夠被快取，常數記憶體的存取速度會很快，接近暫存器或共享記憶體。但如果無法快取，第一次存取仍會較慢。

#### **補充**
- **全域記憶體**：
- **存取速度與頻寬**：Hopper 延遲約 656 週期（比 Ampere 高 2.1 倍，但頻寬提升）。HBM3 記憶體提供 3 TB/s 頻寬（Ampere HBM2e 為 2 TB/s）。新優化：L2 快取管理（CUDA 11.0+），支援持續存取 (persisting) vs. 串流存取 (streaming)，適用計算能力 8.0+（Ampere/Ada/Hopper）。
- **合併存取 (Coalesced Access)**：計算能力 6.0+ 簡化為 32 位元組交易（比 Fermi/Kepler 更高效）。Hopper 支援 FP8 精準度，加速 AI 工作負載。
- **常數記憶體**：
- **大小與特點**：仍為 64KB（無變化），唯讀、快取式。快取命中時延遲接近暫存器（1-4 週期），但 warp 內不同位址存取會序列化。
- **效能**：在 Hopper/Ada 上，若 warp 內執行緒存取相同位置，可達暫存器速度。無重大更新，但 Rust CUDA (2025 年 5 月更新) 支援自動置放靜態變數至常數空間，提升編譯效率。基準測試顯示，與全域記憶體相比，無明顯改善時可切換回全域（RTX 3090 測試）。
- **紋理記憶體 (Texture Memory) 相關**：2024-2025 討論中，紋理記憶體快取仍相關於 2D 空間局部性，但全域記憶體快取已優化為線性存取，適合 warp 連續讀取。

### **記憶體階層與效能**

1. **記憶體階層**
* 從快到慢依序為: **暫存器 (Register)** -> **共享記憶體 (Shared Memory)** -> **常數記憶體 (Constant Memory)** (如果命中快取) -> **局部記憶體 (Local Memory)** -> **全域記憶體 (Global Memory)**。
* **暫存器** 和 **共享記憶體** 的存取速度非常接近。
* **全域記憶體** 和 **局部記憶體** 則慢得多。
2. **GPU 效能瓶頸**
* GPU 的計算能力非常強大，但 **記憶體頻寬 (memory bandwidth)** 是其主要的限制，或稱為 **瓶頸 (bottleneck)**。
* **計算與記憶體存取的速度差異**:
* 一個舊的架構數據顯示，GPU 每秒可以進行 4000 個計算操作，但內部存取全域記憶體每秒只能達到 250 個操作，換算成浮點數傳輸約為每秒 63 個操作。
* 這意味著每次存取全域記憶體，你最好能進行約 60 次計算，才能維持效率。
* **時脈週期 (Clock Cycles) 差異**:
* 計算操作可能只需要 1 到 4 個週期。
* 共享記憶體存取可能需要 1 到 20 個週期。
* 全域記憶體存取可能需要 400 到 800 個週期，甚至高達 100 倍的差異。
* **結論**: 不管從哪個角度看，全域記憶體存取都非常慢，因此必須 **盡量減少對全域記憶體的存取**。
3. **效能優化的關鍵**
* **資料重複使用 (Data Reuse)**: 程式設計師需要管理資料流，確保資料在被載入到快速記憶體 (如共享記憶體) 後能被最大限度地重複使用。
* **記憶體存取與計算的重疊 (Memory Access & Computation Overlap)**: 盡量讓記憶體操作和計算操作同時進行，以隱藏記憶體存取的延遲。

#### **補充**
- **記憶體階層**：從快到慢：暫存器 ≈ 共享記憶體 (1-20 週期) > 常數記憶體 (快取命中) > L1 快取 (32-33 週期) > L2 快取 (Hopper: 264-502 週期) > 全域/本地記憶體 (400-800 週期)。
- **更新**：Hopper L2 快取增至 50MB（Ampere 40MB 的 1.25 倍），輸送量 2.6 倍 Ampere。L1 與共享記憶體效能維持前代，但 Hopper 分區設計（雙分區）影響遠端命中延遲。
- **瓶頸**：記憶體頻寬仍是限制（Hopper 每秒計算操作達萬億級，但全域存取僅數百）。建議計算/存取比至少 60:1。
- **GPU 效能瓶頸**：Hopper 引入 DPX 指令，加速動態規劃（16 位元操作提升 13 倍），隱藏延遲。FP8 tensor cores 提供 FP16 的雙倍效能。

### **優化技巧 (Optimization Techniques)**

1. **一般原則**
* **減少記憶體存取大小**: 盡量只存取必要的資料。
* **最大化資料重複使用**: 將資料搬到更快的記憶體後，集中進行計算，減少重複從慢速記憶體載入。
2. **硬體優化 (Hardware Optimization)**
* GPU 架構不斷更新，透過增加記憶體大小、頻寬，以及優化記憶體存取演算法 (例如 L1 快取) 來提升效能。
3. **軟體優化 (Software Optimization) - 程式設計師的角色**
* **從瓶頸 (bottleneck) 開始優化**: 找出最慢的部分進行優化，通常是外層或最慢的記憶體存取。
* **優化層級**:
* **Host-Device 記憶體傳輸優化 (外層，最慢)**
* **Pinned Memory (Page-locked Memory)**
* **概念**: 將主機記憶體鎖定 (pin)，使其不能被作業系統交換 (swapped out) 到磁碟。
* **優點**: 允許 GPU 直接存取這塊主機記憶體，不需經過作業系統的地址轉換，從而加速 `cudaMemcpy`。傳輸速度可能提升 7-10%。
* **缺點**: 鎖定過多記憶體會減少主機的動態記憶體使用量，可能導致主機程式執行變慢。建議不要鎖定過多記憶體。
* **CPU 與 GPU 計算重疊 (Overlap CPU/GPU Computation)**
* **概念**: CPU 和 GPU 是兩個獨立的處理單元，如果它們的計算之間沒有資料依賴 (data dependency)，應盡量讓它們平行工作。
* **預設行為**: 大部分 CUDA 操作（例如 `cudaMemcpy`、kernel launch）預設是非同步的 (asynchronous)，CPU 在發出指令後會立即回傳並繼續執行自己的工作。**唯一的例外是 `cudaMemcpy` 從 Device 到 Host 預設是同步的 (synchronous)**。
* **避免同步**:
* **1. 獨立計算**: 如果 CPU 和 GPU 的計算完全獨立，可以直接讓 CPU 執行自己的程式，GPU 執行 Kernel，最後再同步結果。
* **2. 使用 `cudaMemcpyAsync`**: 對於 Device 到 Host 的資料複製，可以將預設的同步 `cudaMemcpy` 改為非同步的 `cudaMemcpyAsync`。
* **前提**: `cudaMemcpyAsync` 需要配合 **Pinned Memory** 使用。這是因為非同步傳輸時，GPU 需要直接存取主機記憶體，而 Pinned Memory 確保了記憶體位址的穩定性。
* **額外參數**: `cudaMemcpyAsync` 會多一個 `cudaStream_t` 類型的參數，用於指定操作所屬的 CUDA Stream。
* **CUDA Streams**
* **概念**: Stream 是 GPU 上一系列操作的序列，這些操作會依照發布的順序執行。不同 Stream 中的操作可以平行執行，只要 GPU 有可用的資源。
* **目的**: 實現 GPU 內部的計算與資料傳輸重疊，以及多個 kernel 的平行執行。
* **使用方式**:
1. **創建/銷毀 Stream**: 透過 `cudaStreamCreate` 創建，`cudaStreamDestroy` 銷毀。
2. **指定 Stream ID**: 在 `cudaMemcpyAsync` 和 kernel launch (`<<<gridDim, blockDim, shared_memory_size, stream_id>>>`) 時，指定其所屬的 Stream ID 作為最後一個參數。如果沒有指定，則預設是 Stream 0。
3. **Pipelining (管線化)**: 將大型資料或複雜計算分解為多個小塊。
* **流程**: 先傳輸第一塊資料到 Device，立即啟動第一個 Kernel 處理它。同時，開始傳輸第二塊資料，然後啟動第二個 Kernel。如此循環，實現資料傳輸和計算的重疊。
* **好處**: 大幅減少總執行時間。
* **同步**:
* **`cudaStreamSynchronize(stream)`**: 等待指定 Stream 中的所有操作完成。
* **`cudaEventRecord` 與 `cudaStreamWaitEvent`**: 允許更精細的跨 Stream 同步，例如一個 Stream 等待另一個 Stream 中的特定事件完成。
* **範例**: 複製完第一塊資料後，記錄一個 event。在第二個 Stream 的 Kernel 啟動前，等待這個 event，確保資料已準備就緒。
* **非同步傳輸**: Stream 配合 `cudaMemcpyAsync` 要求 Host 記憶體必須是 Pinned Memory。
* **Global Memory / Local Memory 優化 (GPU 內部)**
* **Coalesced Memory Access (合併記憶體存取)**: 確保 thread 存取全域記憶體時，是連續且對齊的，以便 GPU 硬體能將多個較小的存取合併為一個大的、高效的存取。
* **Padding / Bank Conflicts**: 透過在資料結構中加入填充 (padding) 或重新組織資料，來避免共享記憶體 bank conflicts，以提升效能。
* **Shared Memory 優化 (GPU Block 內部)**
* **Optimal Shared Memory Size**: 選擇每個 block 或 thread block 最適合的共享記憶體大小，以平衡資源利用和平行度。
* **Address Translation**
* **概念**: 多維陣列在硬體層面最終會被線性化 (linearized) 成一維。
* **影響**: 不同的線性化方式 (例如行主序 vs. 列主序) 會影響資料在記憶體中的擺放位置，進而影響記憶體存取模式是否能符合 coalesced access 的要求，或是否會產生 bank conflicts。
* **Padding**: 也可以透過 padding 來改變記憶體佈局，避免 bank conflicts。

#### **補充**
- **一般原則**：強調非同步操作與資料重用。Hopper TMA 支援非同步資料移動，隱藏傳輸延遲。
- **硬體優化**：Blackwell (2025 預計) 將引入更高頻寬 HBM3e，預計頻寬達 5 TB/s。
- **軟體優化**：
- **Host-Device 傳輸**：Pinned Memory 加速 7-10%，結合 cudaMemcpyAsync 與 Streams（Hopper 支援多達 32 Streams 平行）。Pipelining 將資料分塊，重疊傳輸與計算，總時間減半。
- **Global/Local Memory 優化**：確保 coalesced access；Hopper DSM 避免 bank conflicts（環狀模式最佳）。Padding 仍有效。
- **Shared Memory 優化**：啟用 register spilling；warpgroup 操作（wgmma）需 fence 同步。
- **新技巧**：
- **L2 Cache Access Window**：計算能力 8.0+，調整持續存取，提升重複讀取頻寬。
- **Sparse Operations**：Hopper 稀疏 wgmma 延遲 144 週期（比稠密多 16 週期），適合 AI 模型壓縮。
- **精準度優化**：使用 FP8/INT8 加速（Hopper tensor cores），但需注意精度損失。
- **從瓶頸開始**：使用 Nsight Compute 剖析，聚焦高延遲存取。迭代測試 block 大小（128-256 執行緒最佳）。

---
其他課程連結

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