# Day 22｜CUDA 是什麼？GPU 平行運算基礎

17th鐵人賽 cuda

Lontoone

團隊スターバースト・ストリーム

2025-10-07 18:14:00

86 瀏覽

分享至

先給一句話：CUDA 是 NVIDIA 推出的「用 C/C++/Python 等語言操控顯示卡做平行計算」的工具與生態系。
它讓你把原本在 CPU 上慢慢做的工作，搬到擁有成千上萬「小工人」的 GPU 上，同時做完。

1）為什麼需要 CUDA？用「便當店」想像 CPU vs GPU

CPU：像一位多才多藝、會處理複雜客訴的店長，核心（Core）很少、單核很強，適合流程多變、條件分支多的任務。
GPU：像一間擁有上百上千名「包飯糰小幫手」的中央廚房，核心數超多、每個小核心擅長做同樣的動作很多次。

當你的任務是：

「對 1 億個數字各自做相同運算」
「對一張 8K 影像每個像素做濾鏡」
「讓 100 萬個粒子同時更新位置」
→ GPU 超拿手，CUDA 就是把這類任務丟給 GPU 的橋樑。

2）三個最重要的字：Host、Device、Kernel

Host：主機端，指 CPU 跟你的主程式（C/C++/Python）。
Device：裝置端，指 GPU（NVIDIA 顯示卡）。
Kernel：你的「GPU 函式」。你會啟動它，讓海量執行緒同時跑這支函式。

流程長這樣（概念示意）：

CPU(Host) 準備資料 → 傳到 GPU(Device)
            ↓
        啟動 Kernel（一次啟動許多執行緒）
            ↓
GPU 算完 → 把結果傳回 CPU

關鍵心法：Host 安排工作、Device 真的動手做。

3）SIMD？SIMT？GPU 的運作哲學

CPU 常用 SIMD（Single Instruction, Multiple Data）：一個指令同時處理多筆資料（像 AVX）。
CUDA 讓你寫的程式邏輯像是每個執行緒在跑自己的指令，但底層實作是 SIMT（Single Instruction, Multiple Threads）：
- 許多執行緒被**「成群結隊」**一起走相同的指令流程。
- 在 NVIDIA GPU，這個小隊常被稱為 Warp（一組 32 個執行緒，常見設定）。
- 同一個 Warp 內如果有人走 if 的「左邊路」，有人走「右邊路」，就會 分歧（Divergence），造成隊伍分開排隊，速度下降。

口訣：讓同一個 Warp 內的執行緒走相同路線，就會跑得快。

4）執行緒是怎麼排隊的？Thread → Block → Grid

當你啟動一個 Kernel，會一次啟動非常多的執行緒。CUDA 用三層來管理：

Grid（整個任務）
└─ 多個 Block（中隊）
   └─ Block 內有很多 Thread（士兵）

Thread（執行緒）：真正跑你 Kernel 函式的最小單位。
Block：一群執行緒的集合。同一個 Block 的執行緒可以彼此協作（共享記憶體、同步）。
Grid：這次呼叫 Kernel 的全部 Block。

示意圖（1D 版本）：

Grid
 ├─ Block 0: [Thread 0 ... Thread 255]
 ├─ Block 1: [Thread 256 ... Thread 511]
 └─ Block 2: [Thread 512 ... Thread 767]

核心觀念：Block 是協作單位；同一 Block 內才能共享某些快取資源與同步。不同 Block 之間不能彼此直接同步。

5）SM、Warp 與排程：誰在安排誰上場？

SM（Streaming Multiprocessor）：GPU 上的「小舞台」。每個 SM 可以同時容納多個 Block，並切換很多 Warp 上下場。
Warp：上面說過，一隊 32 個執行緒。SM 會輪流發派 Warp 執行，隱藏記憶體延遲。

當有些 Warp 在等資料（例如等記憶體讀取），SM 就讓其他 Warp 先上，這叫「延遲隱藏（Latency Hiding）」。
為了隱藏延遲，你需要足夠多的活躍執行緒（稱為 Occupancy 載入度）來填滿 SM。

6）記憶體層級：為什麼讀資料常常是瓶頸？

GPU 有多層記憶體，速度與容量差很多：

暫存器（Registers）：超快，每個執行緒私有。
Shared Memory（共享記憶體）：同一個 Block 內的執行緒共享，速度快、容量小。
Global Memory（全域記憶體）：整張卡都能存取，容量大、最慢（相對）。
另外還有 Constant/Texture Cache 等特殊用途。

心法：

先想如何少讀 Global Memory。

能把重複用的資料放進 Shared Memory 就放。

讀取時盡量讓連續執行緒讀連續位址（合併存取 / Coalesced Access），速度會差很多。

7）同步與通訊：Block 裡能一起等，Block 間不要等

同一個 Block 的執行緒可以用「同步」指令（如 __syncthreads() 的概念）一起等齊再繼續，並共享 Shared Memory。
不同 Block 之間無法直接同步；要同步通常是結束這輪 Kernel，回到 Host 安排下一輪。

設計演算法時，要把「需要頻繁交換資料的工作」盡量安排在同一個 Block 裡完成。

8）資料搬運：PCIe 與統一記憶體（Unified Memory）

PCIe 傳輸：Host（CPU）與 Device（GPU）之間的資料要走 PCIe，這步可能很慢。
- 常見策略：少搬運、多計算；把資料一次搬過去，多做幾步再搬回。
Unified Memory（統一記憶體）：讓你不必手動管理「拷貝到 GPU / 拷回 CPU」這件事，由系統在需要時搬運。
- 好上手，但效能與可控性不一定比手動拷貝好。入門可以用，進階多半會回到手控。

9）常見應用：什麼工作適合 CUDA？

影像處理：濾鏡、去雜訊、邊緣偵測、超解析度。
科學計算 / 模擬：流體、粒子、有限元素、蒙地卡羅。
機器學習 / 深度學習：張量運算、訓練與推論（PyTorch、TensorFlow 底層大量用 CUDA）。
金融工程：期權定價、風險模擬。
圖形與光線追蹤：路徑追蹤、後處理。

共通特徵：大量、規則、可平行的工作，拆成很多小任務同時做。

10）效能三件事：資料、分工、分歧

資料：
- 減少 Global Memory 存取次數；
- 看能不能把重用資料放 Shared Memory；
- Coalesced：連續執行緒讀連續位址。
分工：
- Block / Thread 數量要合理（通常很多）；
- 讓 SM 有足夠 Warp 輪替（高 Occupancy）以隱藏延遲；
- 但不要讓每個執行緒用太多暫存器/Shared Memory，會擠壓可同時上場的 Warp 數。
分歧（Divergence）：
- 同一 Warp 內如果 if 走不同路，會排隊分時執行，變慢；
- 設計上盡量讓同一 Warp 走一樣的分支，或重排資料讓相似工作排在一起。

11）Debug 與思維：先正確，再快速

先小資料集跑通：先用小 N 讓結果可人工驗證。
與 CPU 版對照：做一個「慢但正確」的 CPU 版本作為「標準答案」。
逐步上效能：量測（profiling）後再針對瓶頸（多半是記憶體存取）優化。
觀念上要接受：GPU 很快，但要「給得出足夠多且規律的工作」才快。

12）生態圈：CUDA 不只是一個 API

CUDA Runtime / Driver API：C/C++ 介面。
CUDA Libraries：像 cuBLAS（矩陣）、cuFFT（傅立葉）、cuDNN（深度學習），可直接用。
高階語言綁定：Python（Numba、CuPy、PyTorch）、Julia、MATLAB 等都能調用 CUDA。
對應他牌：AMD 有 HIP/ROCm（概念相似）；跨家平台也有 OpenCL、近年 SYCL。入門先把 CUDA 概念學好，遷移會更容易。

13）你現在該有的「CUDA 心智圖」

任務是否大量、可平行、重複？
  ├─ 否：CPU（或混合）
  └─ 是：
      Host(準備資料/啟動Kernel) → Device(GPU執行)
         ├─ 設計 Grid/Block/Thread 分工
         ├─ 最小化 PCIe 搬運
         ├─ Shared Memory 作快取
         ├─ 合併存取（Coalesced）
         └─ 避免 Warp 分歧、提升 Occupancy