GPU 架構、Block 排程與 Synchronization (Transparent Scalability)

重點總覽 (Overview)

GPU 架構 (Architecture of a Modern GPU)

• GPU 設計目標是 maximize throughput（CPU 則是 minimize latency，見 01-Introduction/01-Heterogeneous-Parallel-Computing-and-the-Demand-for-Speed）。
• 組織成一陣列的 streaming multiprocessors (SMs)。
• 每個 SM 內有多個 streaming processors / CUDA cores（簡稱 cores），它們共用 control logic 與 memory resources。
• On-chip memory（暫存器、shared memory 等）在 SM 內；off-chip global memory 是數 GB 的 DRAM。
  • 舊 GPU 用 GDDR；Pascal 之後可能用 HBM / HBM2（DRAM 模組與 GPU 封裝在一起）。本書統稱 DRAM。

                       GPU
 ┌─────────────────────────────────────────────┐
 │   SM 0        SM 1        ...     SM (N-1)    │
 │ ┌────────┐ ┌────────┐           ┌────────┐   │
 │ │■■■■ cores│■■■■ cores│  ......  │■■■■ cores│   │  cores 共用：
 │ │■■■■     │■■■■     │           │■■■■     │   │   - Control logic
 │ │ Control │ Control │           │ Control │   │   - On-chip Memory
 │ │ Memory  │ Memory  │           │ Memory  │   │
 │ └────────┘ └────────┘           └────────┘   │
 └─────────────────────┬───────────────────────┘
                       │ (high-bandwidth bus)
 ┌─────────────────────┴───────────────────────┐
 │        Global Memory (off-chip DRAM)         │
 └─────────────────────────────────────────────┘

Block 排程 (Block Scheduling)

• Kernel 啟動時，CUDA runtime 發射一個 grid 的 threads。
• Thread 以 block-by-block 為單位指派到 SM：同一 block 的所有 thread 同時被指派到同一個 SM。
• 一個 SM 通常同時容納多個 block（例：圖中每 SM 3 個 block），但因為 block 需保留硬體資源，每個 SM 能同時容納的 block 數有上限（決定因素見 04-Compute-Architecture-And-Scheduling/03-Resource-Partitioning-and-Occupancy）。
• 大多數 grid 的 block 數 遠多於 能同時執行的數量 → runtime 維護一個 待執行 block 清單，當已指派的 block 完成後，把新的 block 補進 SM。

待執行清單:  B0 B1 B2 B3 B4 B5 B6 B7 B8 ...
                    │  runtime 依資源逐塊指派
                    ▼
   SM0          SM1          SM2
 ┌──────┐     ┌──────┐     ┌──────┐
 │ B0   │     │ B1   │     │ B2   │   ← 第一波同時執行
 │ B3   │     │ B4   │     │ B5   │
 │ B6   │     │ B7   │     │ B8   │   ← 某個 block 完成後，
 └──────┘     └──────┘     └──────┘     清單中下一個 block 才補上

Barrier Synchronization (__syncthreads())

• __syncthreads()（注意 __ 是兩個底線）：當一個 thread 呼叫它，會被卡在該位置，直到同一 block 內每個 thread 都到達該位置。
• 效果：確保 block 內所有 thread 都完成某個 phase，才有任何 thread 進入下一個 phase。
• 生活比喻：四個朋友開車去 mall，各自逛街（平行），但離開前必須等所有人回到車上（barrier）—「no one is left behind」。

time ──────────────────────────────────────►
T0 ====●·····················|   (早到 → 等待)
T1 =========●···············|
T2 ==●·····················|
T3 ===============●·········|
 ⋮                          |   最後一個 thread 到達 barrier 後，
Tn ====================●    |   全部 thread 才一起往下執行
                       ▲ __syncthreads() barrier
        「●」= 到達 barrier   「·」= 等待時間

正確使用規則 (Usage Rules)

• 若 kernel 中有 __syncthreads()，它必須被 block 內所有 thread 執行。
• 放在 if 內時：要嘛 block 內所有 thread 走含 __syncthreads() 的路徑，要嘛都不走。
• 放在 if-then-else、且兩條路徑各有一個 __syncthreads()：這是兩個不同的 barrier！要嘛全體走 then、要嘛全體走 else，否則違規。

// ❌ INCORRECT use of __syncthreads() (Fig. 4.4)
if (threadIdx.x % 2 == 0) {       // line 04: 偶數 thread 走 then
    // ... do work A ...
    __syncthreads();              // line 06: barrier #1
} else {                          // 奇數 thread 走 else
    // ... do work B ...
    __syncthreads();              // line 10: barrier #2 (與 #1 不同!)
}
// block 內 thread 被分到兩個不同 barrier → 沒有任何 barrier
// 被全體到達 → undefined behavior / DEADLOCK

硬體如何保證 barrier 不死鎖

• Barrier 要能完成，所有參與 thread 都必須能取得資源、最終抵達 barrier；任何永遠到不了的 thread 都會造成 deadlock。
• CUDA runtime 的解法：把整個 block 的資源當成一個單位來指派。
  • block 內所有 thread 不只要在同一個 SM，還必須同時被指派。
  • 一個 block 只有在 runtime 確保了該 block 所有 thread 所需的全部資源後，才會開始執行。
  • 這保證了 block 內 thread 的「時間鄰近性 (time proximity)」，避免 barrier 等待過久或無限等待。

Transparent Scalability

• 上節的設計帶來一個重要取捨：不允許跨 block 做 barrier synchronization。
• 既然沒有 block 需要互等，runtime 就能以任意順序執行 block（序列、平行、或任意分波）。
• 這個彈性讓同一份應用程式 code 可依硬體資源量以不同速度執行 → transparent scalability。

同一個 grid (8 個 block)，跑在兩種裝置上：

  低階裝置 (每波 2 blocks)        高階裝置 (每波 4 blocks)
 time│  B0  B1                  time│  B0  B1  B2  B3
   │   B2  B3                     │   B4  B5  B6  B7
   │   B4  B5                     ▼  (2 波完成)
   ▼   B6  B7
     (4 波完成 → 較慢)              (較快，code 完全相同)

• 效益：可針對不同市場（手機低功耗、桌機高效能…）用同一份 binary，無需改 code。
• 高階 GPU 今天可同時執行數百個 block。

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence
• 04-Compute-Architecture-And-Scheduling/03-Resource-Partitioning-and-Occupancy
• 02-Heterogeneous-Data-Parallel-Computing/03-Kernel-Functions-and-Threading
• 05-Memory-Architecture-And-Data-Locality/02-Tiling-and-Tiled-Matrix-Multiplication
• 10-Reduction/02-Optimizing-Single-Block-Reduction-Kernel