計算架構與排程 練習題 (Practice - GPU Architecture, Block Scheduling, Synchronization and Transparent Scalability)

GPU 組織成一陣列的 streaming multiprocessors (SMs)；每個 SM 內有多個 streaming processors / CUDA cores，它們共用 control logic 與 on-chip memory resources。
A100：108 SMs × 64 cores/SM = 6912 cores。考題公式：core 數 = SM 數 × 每 SM core 數。

以 block-by-block 為單位指派：同一 block 的所有 thread 同時被指派到同一個 SM。一個 SM 可同時容納多個 block，但有上限。
block 數通常遠多於能同時執行者 → runtime 維護一個待執行 block 清單，當已指派的 block 完成、釋放資源後，才把新的 block 補進 SM。這個「同 block 同時在同一 SM」的保證，正是 __syncthreads() 與 shared memory 能運作的前提。

不合法。if-then-else 兩個分支裡的 __syncthreads() 是兩個不同的 barrier。由於 block 內並非所有 thread 都會到達同一個 barrier（偶數走 then-barrier、奇數走 else-barrier），違反「barrier 必須由 block 內每個 thread 執行」的規則 → undefined behavior 或 deadlock。
規則：放在 if 內時，block 內所有 thread 要嘛都走含 __syncthreads() 的路徑、要嘛都不走。

這叫 transparent scalability。因為 CUDA 不允許跨 block 做 barrier synchronization，沒有任何 block 需要等待別的 block，runtime 就能以任意順序、任意分波執行 block：低階裝置每波執行少量 block，高階裝置每波執行大量 block。
代價：不同 block 的 thread 無法互相同步。波數公式：waves = ⌈grid_blocks / blocks_per_wave⌉。

48 threads → ⌈48/32⌉ = 2 個 warp。warp n 含 thread 32·n 到 32·(n+1)−1：warp 0 = thread 0–31，warp 1 = thread 32–47。
因為 48 不是 32 的倍數，最後一個 warp 以 16 個 inactive thread 補滿 32 格。warp 是 SM 中 thread 排程的基本單位（目前所有 CUDA device warp size = 32）。

先將多維 thread 投影成 row-major 線性順序（z 大者在後、同 z 內 y 大者在後、同 y 內 x 遞增），再每 32 個切一個 warp。
64 threads → 2 個 warp。Warp 0 = T(0,0) .. T(3,7)（y=0,1,2,3 各 8 個 = 32），Warp 1 = T(4,0) .. T(7,7)。記法 T(y,x)，y 為 threadIdx.y、x 為 threadIdx.x。

SM 把 core 分組成 processing block，每組共用一個 instruction fetch/dispatch unit。同一 warp 的 thread 指派到同一 processing block，硬體取一條指令、同時對 warp 內所有 thread 執行（同指令、不同資料，如 add r1,r2,r3 各 core 的 r2/r3 不同）→ SIMD / SIMT (Single-Instruction, Multiple-Thread)。
A100 SM 64 cores → 4 個 processing block × 16 cores。好處：控制硬體成本被眾多執行單元攤提，晶片面積可多用於算術吞吐。

Oversubscription = 2048 / 64 = 32×，目的是 latency tolerance（延遲容忍）：當某 warp 卡在長延遲操作（如 global memory access），硬體可立即換上另一個就緒 warp 執行，把延遲「藏」起來。駐留 warp 越多，越可能隨時找到就緒的 warp。
Zero-overhead scheduling：GPU 把所有 resident warp 的狀態（PC、registers）常駐於硬體暫存器，切換 warp 不需 save/restore、不引入任何 idle cycle（CPU context-switch 則要存/載狀態 → 有開銷）。

不是好主意 / 不安全。warp 內 thread 是否同步是硬體實作行為，程式正確性絕不可依賴這個假設。尤其從 Volta 架構起的 independent thread scheduling，divergent 的多個 pass 可交錯執行，warp 內不再保證 lockstep。
若確實需要 warp 內同步，應明確使用 __syncwarp()（warp 層級），而非省略同步。

先用 cudaGetDeviceCount(&devCount) 取得 device 數量，再逐一用 cudaGetDeviceProperties(&devProp, i) 查詢每個 device 的屬性後挑選合適者（PC 常含 integrated GPU，CUDA 效能差）。
關鍵欄位：multiProcessorCount（SM 數）、warpSize（warp 大小）、maxThreadsPerBlock、maxThreadsDim[0..2]、maxGridSize[0..2]、regsPerBlock、clockRate。裝置資源量由 compute capability 決定（A100 = 8.0）。

(a) ⌈2000/512⌉ = 4 個 block → 4 × 512 = 2048 threads。
(b) 邊界落在 i = 2000，位於最後一個 block（global thread 1536–2047）內。涵蓋 i = 1984–2015 的那個 warp 同時含 active（i<2000）與 inactive（i≥2000）thread → 僅 1 個 warp 分歧。其餘 warp 要嘛全 active 要嘛全 inactive，不分歧。

(a) 填滿 2048 threads 需 2048/32 = 64 個 block，但撞到 32 block-slot 上限 → 只能放 32 × 32 = 1024 threads → occupancy = 1024/2048 = 50%。（要 ≥64 threads/block 才可能 100%。）
(b) 2048/768 = 2.67 → 只能放 2 blocks = 1536 threads，剩 512 slot 閒置（block 上限與 thread 上限都沒到）→ occupancy = 1536/2048 = 75%。

register 限制下可駐 thread 數 = 65,536 / 64 = 1024 threads → occupancy ≤ 1024/2048 = 50%，且與 block size 無關（register 是限制因子）。
要 full occupancy：regs/thread ≤ 65,536 / 2048 = 32 registers/thread。編譯器可能用 register spilling 降低每 thread register 需求以提升 occupancy，但會增加記憶體存取、可能反而變慢。

兩種情況都只有一個 warp（跨越資料尾端者）分歧，但該 warp 佔總 warp 數的比例決定影響大小：

• 長度 100：⌈100/64⌉=2 blocks，約 4 個 warp，1/4 分歧 → 影響顯著（即使該 warp 時間翻倍，整體可能 +25%）。
• 長度 10,000：約 313 個 warp，1/313 分歧 → 影響 < 1%。
  結論：divergence 的相對成本 ≈ divergent_warps / total_warps，隨資料規模增大而被攤薄。優化時應優先關注大規模下仍顯著的 divergence，而非邊界處的少數 warp。

• 原始 31 reg：4 blocks/SM × 512 = 2048 threads，用 2048×31 = 63,488 ≤ 65,536 registers → 100% occupancy。
• 改成 33 reg：4 blocks 需 2048×33 = 67,584 > 65,536 registers，塞不下 → runtime 降為 3 blocks = 1536 threads（用 50,688 registers）→ occupancy = 1536/2048 = 75%。
  成因：資源用量只小幅增加（+2 registers），卻跨過 register 邊界，使每 SM 的 block 數從 4 驟降為 3 → 平行度與 latency hiding 能力陡降，即 performance cliff (Ryoo et al., 2008)。啟示：優化時必須警覺資源邊界，並用 CUDA Occupancy Calculator 評估；同時注意高 occupancy 只是手段，記憶體頻寬等仍可能才是真正瓶頸。