資料平行性與 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-Transfer02-Heterogeneous-Data-Parallel-Computing/03-Kernel-Functions-and-Threading02-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 同時進行)
一般規則 vs 例外

「每個元素獨立」是 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 帶來擴展性

資料集越大,往往就有越多可利用的 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(例如 cudaMallocmalloc 極為相似),把 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

SPMD ≠ SIMD

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 鄰域)—— 仍可拆成獨立小計算。