Enter your search text in the box above

Select a result to preview

CUDA 功能導覽 (A Tour of CUDA Features)

#cuda #cuda-platform #concept

重點總覽

本筆記是 part 4（Section 4）各進階功能的「導覽地圖」：每個功能只用 1-3 句點出用途與何時使用，細節留給對應的專屬筆記與 Section 4。原文把功能依「想解決的問題類型」粗略分成五大類。

項目 重點
導覽的定位 Section 1-3 打基礎，part 4 假設你已懂這些基礎；本章只介紹各功能動機，不深入細節
分類原則 依「解決的問題類型」粗略分類；有些功能（如 CUDA graphs）可橫跨多類
3.5.1 Improving Kernel Performance async barriers、async data copies + TMA、pipelines、cluster launch control 的 work stealing
3.5.2 Improving Latencies green contexts、stream-ordered allocation、CUDA graphs、programmatic dependent launch、lazy loading
3.5.3 Functionality Features extended GPU memory (EGM)、dynamic parallelism
3.5.4 CUDA Interoperability 與其他 GPU API 互通（Direct3D / Vulkan）、interprocess communication (CUDA IPC)
3.5.5 Fine-Grained Control virtual memory management、driver entry point access、error log management 
A Tour of CUDA Features
├── 3.5.1 Improving Kernel Performance   ── kernel 內部效能
│   ├── Asynchronous Barriers            (§4.9)
│   ├── Async Data Copies + TMA          (§4.11)
│   ├── Pipelines                        (§4.10)
│   └── Work Stealing / Cluster Launch Control (§4.12)
├── 3.5.2 Improving Latencies            ── launch 層級以上的延遲
│   ├── Green Contexts                   (§4.6)
│   ├── Stream-Ordered Allocation        (§4.3)
│   ├── CUDA Graphs                      (§4.2)
│   ├── Programmatic Dependent Launch    (§4.5)
│   └── Lazy Loading                     (§4.7)
├── 3.5.3 Functionality Features         ── 開啟額外能力
│   ├── Extended GPU Memory (EGM)        (§4.17)
│   └── Dynamic Parallelism              (§4.18)
├── 3.5.4 CUDA Interoperability          ── 與外部 API / 程序互通
│   ├── Interop with other APIs          (§4.19)
│   └── Interprocess Communication (IPC) (§4.15)
└── 3.5.5 Fine-Grained Control           ── 細緻底層控制
    ├── Virtual Memory Management        (§4.16)
    ├── Driver Entry Point Access        (§4.20)
    └── Error Log Management             (§4.8)
Tip

不是每個功能都適用於你的 use case。先用這張地圖判斷「我的瓶頸屬於哪一類」，再跳到對應筆記或 Section 4 深入。

3.5.1 Improving Kernel Performance（提升 kernel 效能）

這一類功能全都是為了幫助 kernel 開發者把單一 kernel 的效能榨到最高。

功能 用途 / 何時使用 細節
Asynchronous Barriers 把 barrier 的「arrive」與「wait」分開，等其他 thread 抵達時可先做不依賴 barrier 的工作；可指定不同 thread scope §4.9
Async Data Copies + TMA 在 kernel 內部搬資料於 shared memory 與 GPU DRAM 之間，同時繼續運算；底層用到 async barriers。Tensor Memory Accelerator (TMA) 為其硬體加速 §4.11
Pipelines 階段化（staging）工作、協調多緩衝的 producer–consumer 模式，常用來把運算與 async data copies 重疊 §4.10
Work Stealing / Cluster Launch Control 不均勻負載下維持利用率：已完工的 worker 去「偷」別人的工作 §4.12
Warning

Async data copies 指的是 kernel 內 shared memory ↔ GPU DRAM 的搬移，不要與 CPU↔GPU 之間的 asynchronous memory copy 混淆。

Work stealing with cluster launch control：cluster launch control 是 compute capability 10.0 (Blackwell) 引入的功能，讓 kernel 直接控制 in-flight block 的排程。一個 thread block 可以取消另一個「尚未開始」的 block 或 cluster 的啟動、奪取它的 index，並立刻開始執行偷來的工作。這在不規則資料或執行期變動下能讓 SM 保持忙碌、減少 idle，達到比硬體排程器更細的負載平衡。

SM busy?  ┌── block A 已完工 ──┐
          │  尚有未啟動的 block B (pending)
          └──► A: cancel(B) → 奪取 B.index → 立刻執行 B 的工作
               => SM 不空轉，細粒度 load balancing

3.5.2 Improving Latencies（降低延遲）

共同主題是降低某種延遲，但各功能針對的延遲類型不同。重點放在 kernel launch 層級以上的延遲。

Important

這一類不包含 kernel 內部的 GPU memory access latency（那屬於 3.5.1 的範疇）。

功能 用途 / 何時使用 細節
Green Contexts（execution contexts） 建立只在 GPU 部分 SM 上執行工作的 context；被保留的 SM，其他 context（含 primary context）不會把 block 排上去 → 為高優先/延遲敏感工作保留 SM。CUDA runtime 自 CUDA 13.1 起支援 §4.6
Stream-Ordered Allocation 把配置/釋放 GPU 記憶體排入 stream；cudaMallocAsync / cudaFreeAsync 不像 cudaMalloc/cudaFree 立即執行 §4.3
CUDA Graphs 預先描述一串 CUDA 操作（kernel launch、memcpy）及其相依，供 GPU 高效重複執行；可由 stream capture 或 graphs API 建立 §4.2
Programmatic Dependent Launch 讓依賴前一 kernel 輸出的 dependent kernel 提早啟動 §4.5
Lazy Loading 控制 startup 時 JIT 編譯行為（PTX→cubin），避免一次編譯全部 kernel 拖慢啟動 §4.7

預設情況下，kernel launch 的 thread block 會被分派到 GPU 上任何能滿足資源需求的 SM（受 shared memory、register、cluster 使用、block 內 thread 總數等因素影響）。Green context 進一步限縮可用 SM。

// Stream-ordered allocation：配置/釋放排入 stream，依 stream 順序生效
cudaMallocAsync(&ptr, bytes, stream);
myKernel<<<grid, block, 0, stream>>>(ptr);
cudaFreeAsync(ptr, stream);

重點：*Async 版本不立即執行，而是插入 stream，遵守 stream ordering。

Programmatic dependent launch 的時序：dependent kernel 可先跑 setup 與無關工作，直到需要 primary kernel 的資料才 block；primary kernel 在資料就緒時 signal，釋放 dependent kernel 繼續。

時間 ──────────────────────────────────────►
primary    : [====== compute ======] signal(data ready)
                                     │
dependent  : [setup / 無關工作 ...] block ──┘ [使用資料繼續執行]
             ▲ 提早啟動，overlap 兩個 kernel，縮短關鍵資料路徑延遲
Tip

CUDA graphs 的效益有兩面：減少 CPU 端的 launch 成本，以及啟用「整個 workload 預先已知」才可能的最佳化。適合會重複執行的 workload。

3.5.3 Functionality Features（功能性能力）

共同點：開啟額外的能力 / 功能（而非單純調效能或延遲）。

功能 用途 / 何時使用 細節
Extended GPU Memory (EGM) 在 NVLink-C2C 連接的系統上，讓 GPU 高效存取系統內所有記憶體 §4.17
Dynamic Parallelism 從 GPU 上執行的 kernel 直接發起新的 kernel 啟動（而非只從 CPU 啟動） §4.18

CUDA 最常見的是從 CPU 端啟動 kernel；dynamic parallelism 讓 kernel 在 GPU 上自行產生新的 kernel 呼叫。

3.5.4 CUDA Interoperability（互通）

功能 用途 / 何時使用 細節
Interop with other APIs 與 Direct3D / Vulkan 等圖形 API 共享 GPU buffer：例如用 CUDA 做模擬、再用 3D API 視覺化結果；做法是讓 buffer 對 CUDA 與圖形 API 皆可讀/寫 §4.19
Interprocess Communication (CUDA IPC) 在不同 host process 之間共享 GPU buffer（跨進程、跨機器協作） §4.15
Important

與圖形 API 共享 buffer 的「同一套機制」，也被用來與通訊機制共享 buffer，以支援多節點環境下快速、直接的 GPU-to-GPU 通訊。

多 GPU 協作的兩種型態：

3.5.5 Fine-Grained Control（細緻控制）

針對需要底層細緻控制的進階開發者；多數應用用 CUDA 預設行為即可。

功能 用途 / 何時使用 細節
Virtual Memory Management 透過 CUDA driver API 細緻控制 unified 虛擬位址空間的佈局；主要用於跨 GPU（含跨系統）共享 buffer 時的行為控制 §4.16
Driver Entry Point Access 自 CUDA 11.3 起取得 Driver / Runtime API 的函式指標，包含特定變體，甚至存取比目前 toolkit 更新的 driver 函式 §4.20
Error Log Management 處理與記錄 CUDA API 錯誤；設定環境變數 CUDA_LOG_FILE 即可把錯誤導向 stderr / stdout / 檔案，也可註冊錯誤回呼 (callback) §4.8

如 Section 2.6.1 所述，系統內所有 GPU 加上 CPU 記憶體共享單一統一虛擬位址空間；virtual memory management 即是對這個空間佈局的進階控制。

# Error log management：單一環境變數即可捕捉 CUDA 錯誤
export CUDA_LOG_FILE=stderr      # 或 stdout，或某個檔案路徑

重點：不必改程式碼，設定 CUDA_LOG_FILE 就能落地錯誤紀錄；另可註冊 callback 在 CUDA 出錯時觸發。

考試/測驗重點

情境/關鍵字 答案
Async data copy 指的是 CPU↔GPU 還是 shared↔DRAM？ kernel 內 shared memory ↔ GPU DRAM；勿與 CPU↔GPU async memcpy 混淆
Work stealing 靠哪個功能？哪個架構引入？ cluster launch control；compute capability 10.0 (Blackwell)
一個 block 偷工作的動作 cancel 尚未啟動的 block/cluster → 奪取其 index → 立即執行
Green context 的另一個名字 / 作用 execution context；限制 kernel 只用部分 SM，並讓其他 context（含 primary）不佔用這些 SM
Green context 何時起在 runtime 支援？ CUDA 13.1 及之後
cudaMallocAsync vs cudaMalloc Async 版排入 stream、依序生效；非 Async 版立即執行
建立 CUDA graph 的兩種方式 stream capture（錄製 stream）或 CUDA graphs API
CUDA graphs 的兩大效益 降低 CPU launch 成本 + 啟用「整個 workload 已知」才有的最佳化
Programmatic dependent launch 解決什麼 dependent kernel 提早啟動，與 primary kernel overlap，縮短關鍵資料路徑延遲
Lazy loading 影響哪個階段？ startup 時的 JIT (PTX→cubin)；用環境變數控制，預設 module 用到才編譯
EGM 需要什麼硬體？ NVLink-C2C 連接的系統
Dynamic parallelism 的定義 從 GPU kernel 內發起新的 kernel 啟動
跨 process 共享 GPU buffer 用什麼 CUDA IPC（interprocess communication）
與 Direct3D/Vulkan 共享 buffer 的機制還能做什麼 同一套機制用於 GPU-to-GPU 直接通訊（多節點）
Virtual memory management 透過哪個 API？ CUDA driver API（控制統一虛擬位址空間佈局）
Driver entry point access 自哪版起？做什麼 CUDA 11.3；取得 Driver/Runtime API 函式指標（含變體、較新 driver 函式）
啟用錯誤紀錄的環境變數 CUDA_LOG_FILE（stderr / stdout / 檔案）+ 可註冊 error callback
CUDA graphs 屬於哪一類？ 可橫跨多類；本章歸在 Improving Latencies