昨天先把第一個 CUDA 程式跑起來之後，今天想把「GPU 到底怎麼執行這些程式」補起來。
一開始看 CUDA Programming Guide，很容易被 kernel、thread、block、grid、SM、warp、shared memory 這些名詞轟炸。
其實可以先拆成三件事：
把這三條線串起來，後面學 CUDA optimization 會容易很多。
CUDA 預設的是 heterogeneous system（異質運算系統）。
簡化來看：
CPU / Host
│
│ PCIe / NVLink
▼
GPU / Device
一般桌機中：
System DRAM：平常講的系統 RAM，例如 DDR5GPU DRAM：GPU 使用的記憶體，也就是常說的 VRAM，例如 GDDR6XPCIe / NVLink：CPU、GPU 或 GPU 彼此之間的高速連線CUDA 中：
例如：
__global__ void add(float *a, float *b, float *c) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
c[i] = a[i] + b[i];
}
啟動：
add<<<gridSize, blockSize>>>(a, b, c);
它不是單純讓 GPU 執行一次 add()，而是：
Kernel Launch
↓
啟動大量 GPU Threads
↓
每個 Thread 執行同一份 Kernel Code
每個 thread 再透過自己的 index，決定自己要處理哪一筆資料。
從 CUDA programming model 的角度，可以把 GPU 想成很多個：
SM（Streaming Multiprocessor）
簡化後：
GPU
├── GPC
│ ├── SM
│ ├── SM
│ └── ...
├── GPC
│ ├── SM
│ └── ...
└── GPU Memory
其中：
每個 SM 裡會有：
初學時可以先記：
CUDA 把工作切成很多 Thread Blocks，再把 Blocks 排到各個 SM 上執行。
一次 kernel launch 通常會建立大量 threads。
CUDA 會把它們分組：
Grid
├── Thread Block
│ ├── Thread
│ ├── Thread
│ └── ...
├── Thread Block
└── ...
最基本的層級：
Grid
↓
Thread Block
↓
Thread
主要是讓 thread index 更容易對應資料的形狀。
例如：
處理圖片時：
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
就可以直接理解成：
Thread(x, y)
↓
Pixel(x, y)
這些維度是方便 programmer 使用的 programming abstraction，不代表 GPU 實體真的排成二維或三維。
CUDA 提供幾個 built-in variables：
| 變數 | 功能 |
|---|---|
threadIdx |
Thread 在 Block 裡的位置 |
blockIdx |
Block 在 Grid 裡的位置 |
blockDim |
一個 Block 的大小 |
gridDim |
整個 Grid 的大小 |
最常見的一維 global index：
int i = blockIdx.x * blockDim.x + threadIdx.x;
如果一個 block 有 256 threads：
Block 0, Thread 0 → i = 0
Block 0, Thread 255 → i = 255
Block 1, Thread 0 → i = 256
Block 1, Thread 255 → i = 511
因此每個 thread 都能取得自己的 global identity，再決定自己負責哪筆資料。
這是 CUDA 很重要的一條規則：
同一個 Thread Block 的所有 Threads 都會在同一個 SM 上執行。
原因是同一個 block 裡的 threads 需要能快速：
概念上：
SM 0
├── Block A
│ ├── Thread
│ ├── Thread
│ └── ...
├── Block B
└── ...
注意兩件事：
實際可以同時放多少 blocks，會受到：
影響。
Grid 可能有幾十萬甚至幾百萬個 blocks，但 GPU 可能只有幾十或幾百個 SM。
例如：
1,000,000 Blocks
↓
Scheduler
↓
100 個 SM
GPU 會：
第一批 Blocks
↓
在 SM 上執行
↓
完成、釋放資源
↓
下一批 Blocks
而且 CUDA 不保證不同 Blocks 的執行順序。
所以：
Block 0
Block 1
Block 2
Block 3
實際上可能：
Block 2 → Block 0 → Block 3 → Block 1
也可能同時平行執行。
因此 CUDA 的基本原則是：
不同 Thread Blocks 應該盡量彼此獨立，不依賴其他 Block 的執行順序。
這也是為什麼同一份 CUDA kernel，可以從少量 SM 的 GPU 一路 scale 到大量 SM 的 GPU。
一般情況下，不同 blocks 應該彼此獨立。
但在 Compute Capability 9.0 以上的 GPU，CUDA 還提供一個額外層級：
Thread Block Cluster
可以把它想成：
Grid
├── Cluster
│ ├── Block
│ ├── Block
│ └── Block
└── Cluster
同一個 Cluster 裡的 Blocks：
這算是前面「Block 彼此獨立」的一個進階例外。
初學 CUDA 先知道有這個概念即可。
雖然我們寫程式時看到的是一個個 thread，但 GPU 不會單獨排程每一個 thread。
在一個 Thread Block 中，threads 會被分成：
Warp = 32 Threads
例如：
Block：256 Threads
Warp 0 → Thread 0 ~ 31
Warp 1 → Thread 32 ~ 63
Warp 2 → Thread 64 ~ 95
...
每個 thread 在 warp 裡的位置稱為 warp lane：
lane 0 ~ lane 31
CUDA 採用的是：
SIMT（Single Instruction, Multiple Threads）
可以先口語理解成：
一組 32 個 threads，大部分時間一起執行同一條 instruction，但每個 thread 還是保有自己的資料與控制流程。
假設：
if (threadIdx.x % 2 == 0) {
doSomething();
}
一個 warp 裡：
Thread 0 → 執行
Thread 1 → 不執行
Thread 2 → 執行
Thread 3 → 不執行
...
因為同一個 warp 原本希望一起執行相同 instruction，所以 GPU 會把目前不走這條 branch 的 lanes 暫時 mask off。
概念上：
Warp 32 Threads
│
├── 16 Threads 執行 if
└── 16 Threads 暫時 inactive
這種情況稱為：
Warp Divergence
如果同一個 warp 裡大家走不同 branch，GPU 的硬體利用率通常會下降。
因此：
同一個 Warp 的 Threads 越常走相同的 control flow，通常越有效率。
另外，一個 Thread Block 的 thread 數量通常建議是 32 的倍數。
例如：
128
256
512
因為如果 block 有 100 threads：
Warp 0 → 32
Warp 1 → 32
Warp 2 → 32
Warp 3 → 只剩 4 Threads
最後一個 warp 還是存在，只是很多 lanes 沒有工作。
兩者名字很像：
差異可以先簡單理解成：
CUDA programmer 平常主要用 Thread 的角度寫程式，而不是手動操作一個固定寬度的 SIMD vector。
新版 CUDA Programming Guide 還介紹了另一種模式：
Tile Programming
傳統 SIMT：
Programmer
↓
自己寫每個 Thread 要做什麼
↓
Thread 操作資料
Tile Programming：
Programmer
↓
描述整個 Block 要對一塊 Tile 做什麼
↓
Compiler
↓
自動分配給 Block 裡的 Threads
例如 programmer 可以描述：
比較重要的是：
Block 是執行工作的單位，Tile 是資料的單位，兩者不要混在一起。
Tile Programming 不會取代 SIMT。
對 CUDA 初學者來說，目前還是先把 SIMT / Thread / Block / Warp 搞懂比較重要。
GPU performance 不只看算力，memory 怎麼用通常一樣重要。
先把 CUDA 裡最常見的 memory hierarchy 畫出來：
Register
↓
L1 Cache / Shared Memory
↓
L2 Cache
↓
GPU DRAM / Global Memory
通常越上面：
CPU 與 GPU 通常都有自己直接連接的 DRAM。
CPU
↓
System DRAM
（System / Host Memory）
GPU
↓
GPU DRAM
（Global Memory）
從 CUDA device code 的角度：
GPU 直接連接的 DRAM 稱為 Global Memory
因為 GPU 裡所有 SM 都可以存取它。
而 CPU 直接連接的 DRAM，通常叫：
如果是 multi-GPU：
GPU 0 → 自己的 GPU Memory
GPU 1 → 自己的 GPU Memory
GPU 2 → 自己的 GPU Memory
每張 GPU 都可能有自己的 memory。
除了 GPU DRAM 之外，SM 上還有速度很快的 on-chip memory。
概念上：
Thread 0 → Registers
Thread 1 → Registers
Thread 2 → Registers
例如：
__shared__ float data[256];
概念：
Block
├── Thread 0 ─┐
├── Thread 1 ─┼→ Shared Memory
├── Thread 2 ─┤
└── ... ─┘
每個 SM 都有自己的 Unified Data Cache。
它提供：
所需的實體資源。
概念：
SM
└── Unified Data Cache
├── L1 Cache
└── Shared Memory
在部分 GPU architecture 中，可以調整 L1 與 Shared Memory 的配置比例。
另外：
可以先記成：
Thread
↓
Register
↓
L1 / Shared Memory
↓
L2
↓
GPU Global Memory
SM 的資源不是無限的。
假設一個 kernel：
那同一個 SM 能同時放的 Blocks 就可能變少。
簡單想：
SM 的資源
├── Registers：有限
└── Shared Memory：有限
一個 Block 用越多：
每個 SM 可以同時 Resident 的 Blocks
↓
可能變少
這也是之後學：
Occupancy
時非常重要的基礎。
傳統做法可能需要 programmer 明確管理：
CPU Memory
│
│ cudaMemcpy
▼
GPU Memory
CUDA 也提供：
Unified Memory
讓同一個 allocation 可以被 CPU 或 GPU 存取。
概念上：
Unified Memory
↙ ↘
CPU GPU
CUDA runtime / hardware 會幫忙處理資料需要放在哪裡、何時 migration。
但：
Unified Memory 不代表資料搬移完全沒有成本。
要有效能，還是要盡量減少不必要的 CPU ↔ GPU migration。
新版 CUDA 還有 Locality Domain 的概念。
可以簡單理解成：
在大型 GPU 中，把部分 Memory 與部分 SM 視為一個比較靠近彼此的區域。
如果資料和執行它的 SM 都放在同一個 locality domain，有些 workload 可以有更好的效能。
不過它屬於：
Performance optimization，而不是 correctness requirement。
初學時先知道有這個概念即可。
把今天的內容全部串起來：
CPU
│
│ Kernel Launch
▼
Grid
│
├── Thread Block
│ │
│ ├── Warp
│ │ └── 32 Threads
│ │
│ └── Warp
│
└── Thread Block
│
▼
GPU SM
Memory：
Thread
↓
Register
↓
L1 Cache / Shared Memory
↓
L2 Cache
↓
GPU Global Memory
我們用 Grid / Block / Thread 描述「工作怎麼切」，GPU 再把 Block 排到 SM 上，並以 Warp 為重要的執行群組；而效能高不高，很大一部分取決於 Threads 怎麼走 control flow，以及資料放在哪一層 Memory。
下一步再往下學 memory coalescing、shared memory optimization、occupancy 時，就可以直接把概念接回今天這張圖。
iThome鐵人賽
看影片追技術