資料平行性與 CUDA C 程式結構 (Data Parallelism and CUDA C Program Structure)
重點總覽 (Overview)
| 項目 | 重點 | 為什麼重要 |
|---|---|---|
| Data parallelism | 對資料不同部分的計算彼此獨立,可同時進行 | 是平行程式可擴展性 (scalability) 的主要來源 |
| Task parallelism | 將應用拆成多個獨立的任務 (task),例如 I/O、資料傳輸、不同函式 | 補充角色;數量有限,擴展性不如 data parallelism |
| Host / Device | host = CPU(序列碼),device = GPU(大量平行 kernel) | CUDA 是異質 (heterogeneous) 模型,單一原始檔可混合兩者 |
| Kernel | 以 __global__ 標記、由所有 thread 平行執行的函式 |
一次 kernel call 觸發一個 thread grid |
| Grid → Thread | 一次 launch 產生的所有 thread 合稱一個 grid | thread 是平行執行的基本載具 |
| SPMD | 所有 thread 跑同一份程式碼、處理不同資料 | CUDA 的核心程式設計風格(≠ SIMD) |
本篇聚焦 2.1 (為什麼用 data parallelism) 與 2.2 (host/device 程式結構) 的「為什麼」與「全貌」。
實際的 vecAdd 主機碼、記憶體 API、kernel 索引計算與 launch 語法分別留給 02-Heterogeneous-Data-Parallel-Computing/02-Vector-Addition-Host-Code-Device-Memory-and-Data-Transfer、02-Heterogeneous-Data-Parallel-Computing/03-Kernel-Functions-and-Threading、02-Heterogeneous-Data-Parallel-Computing/04-Calling-Kernels-Compilation-and-Summary。
資料平行性 (Data Parallelism)
當現代軟體跑得慢,問題通常是資料太多 (too much data):影像有數百萬到數兆 pixel、流體模擬有數十億 grid points、分子動力學有數千到數十億 atoms。關鍵觀察是:這些資料元素大多可以彼此獨立地處理。
Data parallelism:對資料集不同部分要做的計算,彼此獨立 (independent),因此可平行完成。撰寫 data parallel 程式 = 圍繞「資料」重新組織計算,讓得到的獨立計算可平行執行,更快完成整體工作。
範例:彩色轉灰階 (Color-to-Grayscale)
每個 pixel 存成 (r, g, b) tuple(0=暗,1=全強度),灰階亮度 (luminance) 由加權和算出:
L = r × 0.21 + g × 0.72 + b × 0.07
關鍵:O[i] 只依賴 I[i],不依賴任何其他 pixel → 每個 pixel 可獨立計算。
輸入 I (RGB array) 輸出 O (luminance array)
┌──────┬──────┬──────┬────┐ ┌────┬────┬────┬────┐
│ I[0] │ I[1] │ I[2] │ … │ │O[0]│O[1]│O[2]│ … │
└──┬───┴──┬───┴──┬───┴────┘ └─▲──┴─▲──┴─▲──┴────┘
│ │ │ │ │ │
t0 t1 t2 … 每個 thread 算一個 pixel(1:1 對應,互不干擾)
L=.21r+.72g+.07b (所有 thread 同時進行)
「每個元素獨立」是 data parallelism 的理想情況(如灰階、vector add)。
例外:blur 需要鄰域 pixel、reduction / 求平均亮度 是「全域」操作 —— 這些仍含豐富 data parallelism,但需拆解成可獨立執行的小計算(見 10-Reduction/01-Reduction-Fundamentals-and-Simple-Kernel)。找出並利用這種平行性正是全書的核心訓練。
Data Parallelism vs. Task Parallelism
| 面向 | Data Parallelism | Task Parallelism |
|---|---|---|
| 拆解依據 | 依資料 (data decomposition) | 依任務 (task decomposition) |
| 典型例子 | 對陣列每個元素做同一運算 | vector add 與 matrix-vector mult 為兩個獨立 task;I/O 與資料傳輸 |
| 平行度規模 | 大(隨資料量成長) | 通常較小(task 數量有限) |
| 對擴展性 | 主要來源 (main source of scalability) | 輔助角色,仍可助達成效能目標 |
| 在本書 | 從本章起貫穿全書 | 之後介紹 streams 時再談 |
資料集越大,往往就有越多可利用的 data parallelism,足以餵飽大規模平行處理器;於是應用效能可隨每一代擁有更多執行資源的硬體一起成長(同一份程式,硬體越大跑越快)。task parallelism 因 task 數量有限,無法提供這種隨硬體放大的擴展。
CUDA C 程式結構 (CUDA C Program Structure)
CUDA C 以最少的新語法與函式庫擴充 ANSI C,讓程式能同時瞄準含 CPU 核心與大規模平行 GPU 的異質運算系統。
| 概念 | 說明 |
|---|---|
| Host | CPU 及其記憶體;執行序列 (serial) 程式碼 |
| Device | 一個或多個 GPU 及其記憶體;以 data-parallel 方式執行 kernel |
| 原始檔 | 單一 .cu 檔可混合 host code 與 device code;device code 用特殊 CUDA 關鍵字明確標記 |
| 預設 | 任何傳統 C 程式都是只含 host code 的合法 CUDA 程式(可逐步加入 device code) |
| Kernel | 以 data-parallel 方式被執行的 device 函式 |
CUDA C 刻意讓介面盡量貼近原本的 C runtime library(例如 cudaMalloc 與 malloc 極為相似),把 C 程式設計師重新學習的成本降到最低。
執行流程 (Execution Flow)
執行從 host 序列碼開始;當呼叫一個 kernel function 時,device 上會 launch 大量 thread 來執行該 kernel。一次 kernel call launch 的所有 thread 合稱一個 grid。grid 內所有 thread 完成後 grid 結束,控制權回到 host,直到下一個 grid 被 launch。
Host (CPU serial) Device (GPU parallel)
───────────────────────────────────────────────
serial code
│
kernel A <<<…>>> ──► Grid A: ┃┃┃┃┃ … ┃┃┃┃┃ (大量 thread 同時跑)
│ ◄── (全部完成後 grid 結束)
serial code
│
kernel B <<<…>>> ──► Grid B: ┃┃┃┃┃ … ┃┃┃┃┃
│ ◄──
serial code
上圖假設 CPU 與 GPU 的執行不重疊 —— 這是教學用的簡化。實務上許多異質應用會重疊 (overlap) CPU/GPU 執行以同時壓榨兩者(屬 task parallelism,之後以 streams 處理,見 20-Heterogeneous-Computing-Cluster/03-Overlapping-Computation-and-Communication)。
Threads、Grid 與 SPMD
- Thread:對「處理器如何執行序列程式」的簡化抽象 —— 包含程式碼、目前執行位置、變數/資料結構的值。從使用者角度看,單一 thread 的執行是序列的。
- CUDA 以呼叫 kernel 來啟動平行:runtime 機制 launch 一個 thread grid,各 thread 平行處理資料的不同部分。
- thread 在 GPU 上極輕量:CUDA 程式設計師可假設產生與排程一個 thread 只需極少 clock cycles(硬體支援);對比傳統 CPU thread 通常需數千 clock cycles。因此灰階轉換可為每個 pixel 開一個 thread(thread 數 = pixel 數)。
kernel 規定「所有 thread 要執行的同一份程式碼」,因此 CUDA 屬 SPMD (Single-Program Multiple-Data) 風格:所有處理單元跑同一支程式、處理不同資料。
與 SIMD 的差別:SIMD 要求所有處理單元在任一瞬間執行同一條指令;SPMD 不要求,各單元不必同時跑到同一條指令。(GPU 硬體底層以 warp 實作 SIMD,見 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence。)
考試/面試重點 (Exam / Test Patterns)
| 情境 / 關鍵字 | 答案 / 技巧 |
|---|---|
| 「平行擴展性的主要來源是什麼?」 | Data parallelism(資料越大、可平行度越大,效能隨硬體成長)。task parallelism 是輔助。 |
| 「data vs task parallelism 差別」 | data = 依資料拆解、平行度大;task = 依任務拆解、平行度有限。task 例子:I/O、資料傳輸、不同函式。 |
| 「灰階亮度公式」 | L = 0.21r + 0.72g + 0.07b;每個 O[i] 只依賴 I[i] → 完全獨立。 |
| 「SPMD 是不是 SIMD?」 | 不是。SPMD 同程式不同資料、不要求同時執行同指令;SIMD 要求每瞬間同指令。 |
| 「kernel call 產生什麼?」 | 一個 thread grid;grid 完成後才回到 host。一個 source file 可混 host/device code。 |
| 「GPU thread 與 CPU thread 開銷」 | GPU thread 產生/排程僅需極少 cycles;CPU thread 需數千 cycles。 |
| 「為何能為每個 pixel 開一個 thread?」 | 各 pixel 計算獨立 + GPU thread 極輕量,thread 數可 = 資料元素數。 |
| 「哪種運算難直接 data-parallel?」 | 全域/相依運算(求平均、reduction、blur 鄰域)—— 仍可拆成獨立小計算。 |
Related Notes
- 02-Heterogeneous-Data-Parallel-Computing/02-Vector-Addition-Host-Code-Device-Memory-and-Data-Transfer
- 02-Heterogeneous-Data-Parallel-Computing/03-Kernel-Functions-and-Threading
- 02-Heterogeneous-Data-Parallel-Computing/04-Calling-Kernels-Compilation-and-Summary
- 01-Introduction/01-Heterogeneous-Parallel-Computing-and-the-Demand-for-Speed
- 19-Parallel-Programming-And-Computational-Thinking/03-Problem-Decomposition-Output-vs-Input-Centric