使用 PTX 與硬體模型 (Using PTX and Hardware Model)

重點總覽

Using PTX（手寫 PTX）

PTX（Parallel Thread Execution） 是 CUDA 用來抽象各代硬體 ISA 的虛擬機指令集架構（virtual machine ISA）。直接以 PTX 撰寫程式是高度進階的最佳化手法，對大多數開發者並非必要，應視為最後手段（tool of last resort）。

• 適用時機：應用程式中極度效能敏感的部分，連 0.x% 的效能提升都有顯著效益。
• 完整指令列於 PTX ISA document。
• 兩種在 C++ 中使用 PTX 的方式：

inline PTX 透過 asm volatile 嵌入，語法上很像 CPU 的 inline assembly：

__device__ int add_via_ptx(int a, int b) {
    int r;
    // 將兩個 .s32 相加，輸出綁到 %0，輸入綁到 %1、%2
    asm volatile("add.s32 %0, %1, %2;" : "=r"(r) : "r"(a), "r"(b));
    return r;
}

重點：volatile 阻止編譯器將該指令最佳化掉或重排；"=r"/"r" 為輸出/輸入 register constraint。

Hardware Implementation（硬體實作概觀）

一個 streaming multiprocessor（SM） 設計上可同時並行執行數百個 threads。為管理如此大量 threads，它採用獨特的並行計算模型 SIMT（Single-Instruction, Multiple-Thread）。

• 指令是 pipeline 化的：同時利用單一 thread 內的 instruction-level parallelism (ILP)，以及透過 hardware multithreading 的 thread-level parallelism。
• 與 CPU core 不同：SM 按順序（in order）發射指令，不做 branch prediction，也不做 speculative execution。
• SIMT Execution Model 與 Hardware Multithreading 描述所有裝置共通的架構特性；各 compute capability 的細節在 Compute Capabilities。
• NVIDIA GPU 架構使用 little-endian 表示法。

SIMT Execution Model

每個 SM 以 32 個並行 threads 為一組來建立、管理、排程與執行，這一組稱為 warp。

• 同一 warp 的各 thread 起始於同一程式位址，但各自擁有獨立的 instruction address counter 與 register state，因此可自由 branch 與獨立執行。
• 名詞溯源：warp 來自織布（weaving）；half-warp 是 warp 的前半或後半（16 threads），quarter-warp 是四等分之一（8 threads）。
• warp 一次執行一條共同指令，當 32 threads 都走同一執行路徑時達到全效率。
• 若 warp 內 threads 因 data-dependent conditional branch 而 diverge，warp 會逐條執行每個被取用的 branch 路徑，並停用不在該路徑上的 threads。
• Branch divergence 只發生在 warp 內；不同 warp 互相獨立，無論是否走相同路徑。

warp（32 threads）遇到 if(cond) 分歧：
  cond=true 路徑 :  T0 T1 .. T15  ----執行---->  (T16..T31 停用)
  cond=false 路徑:                              T16..T31 ----執行---->  (T0..T15 停用)
  reconverge     :  ============= 全 32 threads 再次一起 =============
  → 兩條路徑「序列化」，divergence 越多效能越差

Independent Thread Scheduling（CC 7.0+）

• CC < 7.0：因共用 PC + active mask，divergent region 中的 threads 無法互相 signal 或交換資料；需要 lock/mutex 細粒度共享資料的演算法可能 deadlock（視競爭 threads 來自哪個 warp 而定）。
• CC 7.0+：independent thread scheduling 允許 thread 間（不分 warp）完全並行。GPU 維護 per-thread 執行狀態（含 program counter 與 call stack），可讓一個 thread 等待另一 thread 產生的資料。
• schedule optimizer 決定如何把同 warp 的 active threads 分組進 SIMT units，保留 SIMT 高吞吐但更有彈性。

共用 PC (CC<7.0):     [PC]  → 全 warp 一個指令指標 + active mask
每 thread PC (CC≥7.0): [PC0][PC1][PC2]...[PC31] + 各自 call stack
                        ↓ schedule optimizer 重新分組成 SIMT units
                        → 可在 sub-warp 粒度 diverge / reconverge

Hardware Multithreading

當 SM 拿到一或多個 thread block，會把它們切分成 warps，每個 warp 由一個 warp scheduler 排程執行。

• 切分方式固定：每個 warp 含連續遞增 thread ID，第一個 warp 含 thread 0。
• 每個 block 的 warp 數：

warps_per_block = ceil(T / Wsize)
  T     = 每 block 的 threads 數
  Wsize = warp size = 32
  ceil(x, y) = x 向上取整到 y 的最近倍數

• 每個 warp 的 execution context（program counters、registers 等） 在 warp 整個生命週期中常駐 on-chip，因此 warp 之間切換零成本（zero-overhead），可隱藏延遲。
• 每個 instruction issue cycle，warp scheduler 挑選一個有 ready threads 的 warp，把指令發射給該 warp 的 active threads。

issue cycle:  cycle0   cycle1   cycle2   cycle3 ...
scheduler →   warpA    warpC    warpA    warpB   （挑 ready 者，無切換開銷）
              ↑ 某 warp 在等記憶體時，立刻切到另一 ready warp 隱藏延遲

• 每個 SM 有一組 32-bit registers（在 warps 間分配）與 shared memory（在 thread blocks 間分配）。
• 一個 kernel 可同時常駐的 block/warp 數，取決於 kernel 用掉的 register 與 shared memory，以及 SM 上可用量；另有每 SM 最大常駐 block 數與 warp 數上限。

Asynchronous Execution Features

近代 NVIDIA GPU 加入非同步執行能力，讓資料搬移、計算、同步在 GPU 內更能重疊。這些能力讓從 GPU code 發起的某些操作，可相對於同一 thread block 內的其他 GPU code 非同步執行。

• async operation 由一個 CUDA thread 發起，但「彷彿由另一個 thread 執行」，那個 thread 稱為 async thread。
• well-formed 程式中，一或多個 CUDA threads 會與該 async operation 同步；發起的 thread 不必屬於同步它的 threads。async thread 永遠關聯到發起它的 CUDA thread。
• async operation 用一個 synchronization object 來 signal 完成，可以是 barrier 或 pipeline（詳見 Advanced Synchronization Primitives 與 Asynchronous Data Copies）。

Async Thread and Async Proxy

非同步操作存取記憶體的方式可能與一般操作不同。為區分，CUDA 引入三個概念：async thread、generic proxy、async proxy。一般 load/store 走 generic proxy。

• Async thread on generic proxy：發起時關聯一個有別於發起者的 async thread。對同位址而言，先前的 normal load/store 保證排在 async op 之前；但後續的 normal load/store 不保證維持順序，在 async thread 完成前可能 race condition。
• Async thread on async proxy：對同位址而言，先前與後續的 normal load/store 都不保證維持順序。需要 proxy fence 跨 proxy 同步，以確保正確的 memory ordering（TMA 一節示範用 proxy fence 確保非同步複製正確）。

generic proxy:  [前 normal store] --ordered--> [async op (LDGSTS)] ...[後 normal load 不保證順序]
async proxy:    [前 normal store]  ?無序?  [async op (TMA)]  ?無序?  [後 normal load]
                                    └────────── 需 proxy fence 才能跨 proxy 排序 ──────────┘

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/01-Advanced-Launch-and-Clusters
• 03-Advanced-CUDA/05-Thread-Scopes-and-Scoped-Atomics
• 03-Advanced-CUDA/06-Asynchronous-Barriers-and-Pipelines
• 03-Advanced-CUDA/07-Async-Data-Copies-and-L1-Config
• 03-Advanced-CUDA/08-CUDA-Driver-API
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 01-Introduction-to-CUDA/05-CUDA-Platform
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps