CUDA Graphs：結構與擷取 (Graph Structure and Capture)

重點總覽

CUDA Graphs 提供另一種工作提交模型：把一連串以「依賴關係」連接的操作（kernel launch、資料搬移等）定義一次、重複啟動多次。將定義與執行分離可降低 CPU launch 成本，並讓 CUDA 在掌握整體工作流程後做出 stream 逐段提交無法達成的最佳化。

Graph Structure：節點與邊

graph 由 node（操作） 與 edge（依賴） 構成。edge 約束執行順序：一個操作只要它所依賴的 node 都完成，就可在任意時點被排程，排程交由 CUDA 系統決定。

Node Types（節點類型）

一個 graph node 可以是以下其中一種：

Edge Data（邊資料，CUDA 12.3 引入）

edge data 用來修飾一條依賴邊，由三部分組成：

• outgoing port：指定該邊在上游 node 的「何時」被觸發。
• incoming port：指定下游 node 的「哪一部分」依賴此邊。
• type：修飾兩端點之間的關係。

port 值與 node 類型及方向相關，edge type 也可能限定於特定 node 類型。零初始化的 edge data 代表預設行為：

實務重點：

• edge data 在各 graph API 以「與 node 平行的陣列」選擇性傳入；若省略輸入，使用零初始化資料。
• 作為輸出（查詢）參數省略時，若被忽略的 edge data 全為零初始化則允許；否則回傳 cudaErrorLossyQuery（避免遺失資訊）。
• 部分 stream capture API 也支援 edge data：cudaStreamBeginCaptureToGraph()、cudaStreamGetCaptureInfo()、cudaStreamUpdateCaptureDependencies()。此時尚無下游 node，資料附在 dangling edge（half edge，懸空半邊） 上，未來連到被擷取的 node 或在 capture 結束時被丟棄。
• 唯一的非預設依賴類型是 cudaGraphDependencyTypeProgrammatic，用於兩個 kernel node 之間啟用 Programmatic Dependent Launch (PDL)。目前 non-default edge data 的唯一用途就是啟用 PDL。
• 沒有任何 node 類型定義額外的 incoming port；只有 kernel node 定義額外的 outgoing port。

Building and Running Graphs：三階段

  ┌────────────┐    ┌──────────────────┐    ┌───────────────┐
  │ Definition │ →  │  Instantiation   │ →  │   Execution   │
  │  建立節點   │    │ cudaGraph        │    │ cudaGraph     │
  │  與依賴     │    │ Instantiate      │    │ Launch (×N)   │
  │ (cudaGraph_t)   │ → cudaGraphExec_t│    │  送入 stream  │
  └────────────┘    └──────────────────┘    └───────────────┘

• Definition / Creation：建立 graph 中各操作及其依賴的描述（產生 cudaGraph_t 模板）。
• Instantiation：對 graph 模板取快照、驗證、並完成大部分設定與初始化，目的是把 launch 時要做的事降到最低，產出 executable graph。
• Execution：executable graph 可像一般 CUDA work 一樣送入 stream，且可重複啟動任意次而不需重新 instantiate。

Graph Creation 方式一：顯式 Graph API

用 cudaGraphCreate() 建立空 graph，再以 cudaGraphAddNode()（或 cudaGraphAddKernelNode / cudaGraphAddMemcpyNode / cudaGraphAddMemsetNode / cudaGraphAddHostNode 等）逐一加入節點與依賴。

// 建立空 graph
cudaGraphCreate(&graph, 0);
cudaGraphNode_t nodes[4];
cudaGraphNodeParams kParams = { cudaGraphNodeTypeKernel };
// ... 設定 func / gridDim / blockDim ...
cudaGraphAddNode(&nodes[0], graph, NULL,      NULL, 0, &kParams);
cudaGraphAddNode(&nodes[1], graph, &nodes[0], NULL, 1, &kParams);
cudaGraphAddNode(&nodes[2], graph, &nodes[0], NULL, 1, &kParams);
cudaGraphAddNode(&nodes[3], graph, &nodes[1], NULL, 2, &kParams);

cudaGraphAddNode 的依賴參數（第 3、5 個引數）指定該 node 依賴哪些既有 node。上例建出 nodes[0] 為根，分岔到 nodes[1]/nodes[2]，nodes[3] 再依賴 nodes[1]。實際應用還需加入 memcpy/memset 等記憶體操作節點。

Graph Creation 方式二：Stream Capture

stream capture 從既有 stream-based API 建出 graph：用 cudaStreamBeginCapture() 與 cudaStreamEndCapture() 把一段送入 stream 的程式碼包起來。

cudaGraph_t graph;
cudaStreamBeginCapture(stream);
kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);            // 既有函式庫呼叫也能被擷取
kernel_C<<< ..., stream >>>(...);
cudaStreamEndCapture(stream, &graph);

• 進入 capture 模式後，送入 stream 的 work 不會被排入執行，而是逐步附加到一個內部正在建構的 graph（稱為 capture graph），最後由 cudaStreamEndCapture() 回傳並結束 capture。
• 可用於任何 stream，唯獨不能用於 cudaStreamLegacy（NULL stream）；但可用於 cudaStreamPerThread。若程式使用 legacy stream，常可把 stream 0 重定義為 per-thread stream 而不改變功能。
• cudaStreamIsCapturing() 可查詢某 stream 是否正在被擷取。
• cudaStreamBeginCaptureToGraph() 可把 work 擷取到使用者提供的既有 graph，而非內部新建的 graph。

Cross-stream Dependencies and Events（跨 stream 依賴）

stream capture 能處理以 cudaEventRecord() 與 cudaStreamWaitEvent() 表達的跨 stream 依賴，前提是被等待的 event 記錄於同一個 capture graph。

• 在 capture 模式的 stream 中記錄 event 會產生 captured event，代表 capture graph 中的一組 node。
• 當某 stream 等待一個 captured event：若該 stream 尚未在 capture 模式則被帶入 capture 模式，其下一個項目會額外依賴該 captured event 所代表的 node；兩條 stream 自此被擷取到同一個 capture graph。
• origin stream：呼叫 cudaStreamBeginCapture() 的那條 stream。cudaStreamEndCapture() 必須在 origin stream 呼叫，且因 event 依賴而加入同一 capture graph 的其他 stream 都必須 join 回 origin stream，否則整個 capture 失敗。所有 stream 在 cudaStreamEndCapture() 時一起退出 capture 模式。

// stream1 是 origin stream
cudaStreamBeginCapture(stream1);
kernel_A<<< ..., stream1 >>>(...);
cudaEventRecord(event1, stream1);      // fork：記錄 A
cudaStreamWaitEvent(stream2, event1);  // stream2 等待 A
kernel_B<<< ..., stream1 >>>(...);
kernel_C<<< ..., stream2 >>>(...);
cudaEventRecord(event2, stream2);      // join：記錄 C
cudaStreamWaitEvent(stream1, event2);  // stream1 等待 C
kernel_D<<< ..., stream1 >>>(...);
cudaStreamEndCapture(stream1, &graph); // 兩條 stream 退出 capture 模式

上述 fork/join 擷取出的拓撲（菱形依賴）：

            kernel_A          (origin / root)
           ╱        ╲
      kernel_B    kernel_C    (B 在 stream1、C 在 stream2)
           ╲        ╱
            kernel_D          (依賴 B 與 C，回到 origin)

Prohibited and Unhandled Operations（禁止與不支援的操作）

Invalidation（失效）

當 capture 期間嘗試了無效操作，相關的 capture graph 會被 invalidate。capture graph 失效後，任何正被擷取的 stream 或關聯 captured event 都不可再用、會回傳錯誤，直到以 cudaStreamEndCapture() 結束擷取。該呼叫會把相關 stream 移出 capture 模式，但同時回傳錯誤值與 NULL graph。

Capture Introspection（擷取內省）

cudaStreamGetCaptureInfo() 可檢視作用中的 capture：取得 capture 狀態、每個 process 唯一的 capture ID、底層 graph 物件，以及該 stream 中「下一個待擷取 node」的依賴/edge data。這些依賴資訊可用來取得 stream 中最後被擷取的 node 的 handle。

Graph Instantiation（實例化）

graph 建立後（無論來自 graph API 或 stream capture）須先 instantiate 成 executable graph 才能 launch：

cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);

instantiate 會對 graph 模板取快照、驗證、並完成設定，產出可啟動的 cudaGraphExec_t。

Graph Execution（執行）

executable graph 用 cudaGraphLaunch() 送入指定 stream；一次 instantiate、可多次 launch：

cudaGraphLaunch(graphExec, stream);

完整流程（以 stream capture 為例）：擷取 → 結束擷取得 graph → instantiate → launch：

cudaStreamBeginCapture(stream);
kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);
kernel_C<<< ..., stream >>>(...);
cudaStreamEndCapture(stream, &graph);
cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
cudaGraphLaunch(graphExec, stream);

考試/測驗重點

Related Notes

• 04-CUDA-Features/04-CUDA-Graphs-Updating-and-Conditional
• 04-CUDA-Features/05-CUDA-Graphs-Memory-Nodes-and-Device-Launch
• 04-CUDA-Features/08-Programmatic-Dependent-Launch-Deep-Dive
• 02-Programming-GPUs/15-Async-Callbacks-Ordering-Graphs
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps