前言
SM (Streaming Multiprocessor) 是 GPU 並行算力的核心單元。本文對比 Fermi (CC 2.x) 與 Kepler (CC 3.x) 架構的 SM 微架構、執行管線、資源配置，並詳細解析 Kepler 引入的 Dynamic Parallelism 與 Hyper-Q 兩大創新特性。
Fermi 架構 (CC 2.0, 2010)
- 512 CUDA Cores：16 SM × 32 cores/SM。
- SM 結構：32 CUDA Cores、16 Load/Store、4 SFU、2 Warp Scheduler、2 Dispatch Unit、64KB 共享記憶體/L1 (可配置 16/48)、32K 暫存器檔。
- GigaThread Engine：全域 Block 調度分發。
- 6×384-bit GDDR5：支援 6GB 顯存、ECC。
- 768KB L2 Cache：統一快取。
- 並行能力：每 SM 48 Warps (1536 線程) 駐留。
Kepler 架構 (CC 3.0/3.5, 2012)
- 15 SMX (K20X)：每 SMX 192 單精度 Cores、64 雙精度、32 SFU、32 LD/ST。
- SMX 結構：4 Warp Scheduler、8 Dispatch Unit → 每週期最多發射 4 Warps (128 指令)。
- 64K 暫存器檔：較 Fermi 翻倍。
- 6×64-bit Memory Controller：頻寬優化。
- 並行能力：每 SM 64 Warps (2048 線程) 駐留。
架構關鍵參數對比表
| 參數 | Fermi (GF100) | Kepler (GK110) |
|---|
| SM 數量 | 16 | 15 |
| Cores/SM | 32 | 192 (SP) |
| 雙精度比率 | 1/2 | 1/3 |
| Warp Scheduler/SM | 2 | 4 |
| Dispatch Unit/SM | 2 | 8 |
| Max Warps/SM | 48 | 64 |
| Max Threads/SM | 1536 | 2048 |
| 暫存器檔/SM | 32K | 64K |
| 共享記憶體/SM | 64KB (可配置) | 64KB (可配置) |
| L2 Cache | 768KB | 1536KB |
Dynamic Parallelism (動態並行)
痛點
Fermi 要求所有 Kernel 從 Host 啟動，遞歸、自適應演算法、動態任務圖需 Host 參與 → 高延遲、複雜同步。
Kepler 解決方案
- Kernel 內可直接啟動子 Kernel (
<<<>>>))。
- 硬體級工作隊列，無需 Host 干預。
- 支援遞歸、巢狀並行、動態負載平衡。
應用場景
- 自適應網格剖分 (AMR)。
- 光線追蹤遞歸反射/折射。
- 動態任務圖 (DAG) 就地展開。
代碼範例
__global__ void childKernel(int *data) { ... }
__global__ void parentKernel(int *data) {
if (condition) childKernel<<<1, 256>>>(data);
// 需配合 cudaDeviceSynchronize() 或事件同步
}
Hyper-Q (超級隊列)
Fermi 瓶頸
單一硬體工作隊列 → 任務序列化 → 長任務阻塞短任務 → GPU 利用率低。
Kepler 創新
- 32 硬體工作隊列 (Hardware Work Queues)。
- CPU 多線程可同時提交 Kernel 至不同隊列。
- 硬體調度器跨隊列公平調度，消除 Head-of-Line Blocking。
效能提升
- MPI 多進程每進程綁定一隊列 → 並行啟動無干擾。
- Hyper-Q + MPS (Multi-Process Service) → 單 GPU 服務多進程零干擾。
SIMT vs SIMD 核心差異
