CUDA Graphs：更新與條件節點 (Updating and Conditional Nodes)

#cuda #async-execution #cuda-graphs #concept

當工作流程改變時，graph 會過期而必須修改。拓撲或節點類型等重大改變需要重新 instantiate（拓撲相關的最佳化必須重跑）；但若僅節點參數（kernel 參數、記憶體位址）改變而拓撲不變，CUDA 提供輕量的「Graph Update」機制可就地修改，遠比 re-instantiation 高效。更新會在下一次 launch 時生效，不影響先前已排入或正在執行的 launch。

重點總覽

項目	重點
何時可 update	拓撲、節點類型不變，只改參數時；否則必須 re-instantiate
生效時機	更新於下一次 `cudaGraphLaunch` 生效，不影響進行中的 launch
Whole graph update	`cudaGraphExecUpdate()`：用拓撲相同的新 `cudaGraph_t` 整批更新
Individual node update	用 `cudaGraphExec*NodeSetParams()` 逐節點更新，跳過拓撲檢查
Individual node enable	`cudaGraphNodeSetEnabled()` 啟用／停用 kernel/memset/memcpy 節點
拓撲配對規則	API 呼叫順序、dependency 陣列順序、sink node 順序必須一致
Conditional node	在 graph 內做條件執行與迴圈，評估在 device 上進行
三種類型	IF（含可選 else）、WHILE（迴圈）、SWITCH（n 選 1）
Conditional handle	`cudaGraphConditionalHandle`，device 端用 `cudaGraphSetConditional()` 設值

Whole Graph Update (cudaGraphExecUpdate)

cudaGraphExecUpdate() 讓已 instantiate 的 graph（original graph）以一個拓撲完全相同的 updating graph 的參數來更新。它會把原圖與新圖的節點做配對；配對成功才能套用新參數。

updating graph 的拓撲必須與用來 instantiate cudaGraphExec_t 的原圖一致，且 dependency 指定的順序也要一致。
CUDA 依賴特定 API 的呼叫順序來一致地排序 sink node（無 outgoing edge 的節點）。
失敗時回傳 errorNode（出錯節點）與 updateResult（原因碼）；只有 cudaGraphExecUpdateSuccess 才代表成功。

要使節點配對具決定性，須遵守以下規則：

同一個 capturing stream 上的 API 呼叫順序須相同（含 event wait 等不直接對應 node 建立的呼叫）。
直接操作某節點 incoming edge 的 API（captured stream API、node add API、edge 增刪 API）順序須相同；dependency 以陣列傳入時，陣列內順序也要相同。
Sink node 須一致排序。影響 sink node 排序的操作（作為一個整體集合）順序須相同：產生 sink node 的 node add、使節點變 sink 的 edge removal、移除 sink 的 cudaStreamUpdateCaptureDependencies()、cudaStreamEndCapture()。

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
do_cuda_work(stream);                 // 重新擷取同一份工作流程
cudaStreamEndCapture(stream, &graph);
if (graphExec != NULL) {              // 已 instantiate 過：嘗試就地更新
    cudaGraphExecUpdate(graphExec, graph, &errorNode, &updateResult);
}
if (graphExec == NULL || updateResult != cudaGraphExecUpdateSuccess) {
    if (graphExec != NULL) cudaGraphExecDestroy(graphExec);  // 更新失敗才重建
    cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
}
cudaGraphLaunch(graphExec, stream);

典型流程：每輪重新建一個拓撲相同的 cudaGraph_t，先試 cudaGraphExecUpdate()，成功就直接 launch，失敗才 cudaGraphExecDestroy() + cudaGraphInstantiate() 重建。

何時用 whole graph update

當更新節點數多、或 graph 拓撲對呼叫端未知（如 library 呼叫經 stream capture 產生）時，整圖更新最方便。Conditional handle 的 flags 與 default value 也會在整圖更新時一併更新。

Individual Node Update

可直接更新 instantiate 後的節點參數，省去建立新 cudaGraph_t 與 instantiation 的開銷。當需更新的節點數相對於總節點數很小時，逐節點更新更佳；它會跳過拓撲檢查與未變節點的比對。

API	節點類型
`cudaGraphExecKernelNodeSetParams()`	Kernel node
`cudaGraphExecMemcpyNodeSetParams()`	Memory copy node
`cudaGraphExecMemsetNodeSetParams()`	Memory set node
`cudaGraphExecHostNodeSetParams()`	Host node
`cudaGraphExecChildGraphNodeSetParams()`	Child graph node
`cudaGraphExecEventRecordNodeSetEvent()`	Event record node
`cudaGraphExecEventWaitNodeSetEvent()`	Event wait node
`cudaGraphExecExternalSemaphoresSignalNodeSetParams()`	External semaphore signal node
`cudaGraphExecExternalSemaphoresWaitNodeSetParams()`	External semaphore wait node

也可改 cudaGraph_t 節點（如 cudaGraphKernelNodeSetParams()）後再 cudaGraphExecUpdate()，但用上表的 exec 級 API 更有效率。
持有節點 handle、變更數量少時優先選 individual node update。

Individual Node Enable

cudaGraphNodeSetEnabled() 可啟用／停用 instantiate 後的 kernel、memset、memcpy 節點；用 cudaGraphNodeGetEnabled() 查詢狀態。這讓你建立一個「功能超集」的 graph，再針對每次 launch 客製化。

被停用的節點功能上等同空節點（empty node），直到重新啟用。
啟用／停用不影響節點參數；enable 狀態也不受 individual node update 或 whole graph update 影響。
節點停用期間做的參數更新，會在重新啟用時生效。

Graph Update Limitations

各節點類型的更新限制

Graph update 只能改參數，不能改變「身分」性質的屬性。違反限制會導致 update 失敗（須 fallback 重建）。

節點類型	限制
Kernel	不能改 function 的 owning context；原本不用 dynamic parallelism 的 function 不能改成用 dynamic parallelism 的
Memset / Memcpy	運算元分配／映射的 device 不能變；source/destination 須來自與原本相同的 context；只有 1D memset/memcpy 可改
Memcpy（額外）	不能改 source/destination 的 memory type（`cudaPitchedPtr`、`cudaArray_t` 等）或 transfer 類型（`cudaMemcpyKind`）
External semaphore wait/record	不能改 semaphore 數量
Conditional	handle 建立與指派的順序須兩圖一致；不能改節點參數（conditional 內的 graph 數量、node context 等）；body graph 內節點的參數變更受上述規則約束
Memory node	若該 `cudaGraph_t` 目前已 instantiate 成另一個 `cudaGraphExec_t`，則不能拿它去 update
Host / Event record / Event wait	無更新限制

Conditional Graph Nodes 概觀

Conditional node 讓「條件執行」與「迴圈」完全表達在 graph 內，使 host CPU 釋放出來並行做其他工作。condition 的評估在 device 上進行，發生於 conditional node 的 dependency 都滿足時。

三種類型：IF（condition 非零執行 body 一次，可選第二個 body 於 condition 為零時執行）、WHILE（condition 非零就反覆執行 body 直到為零）、SWITCH（執行第 n 個 zero-indexed body；若 condition 不對應任何 body 則不執行）。
建立 conditional node 時會建立一個空的 body graph 並回傳給使用者填充；可用 graph API 或 cudaStreamBeginCaptureToGraph() 填充。
Conditional node 可以巢狀。

Conditional Handles

condition value 由 cudaGraphConditionalHandle 表示，以 cudaGraphConditionalHandleCreate() 建立。

handle 必須先於節點建立，且只能關聯到單一 conditional node。
handle 無法被 destroy，故無須追蹤其生命週期。
device 端用 cudaGraphSetConditional(handle, value) 設定 condition value。
若建立時指定 cudaGraphCondAssignDefault，每次 graph execution 開始時 condition 會初始化為指定 default；未指定則 condition value 在每次執行開始時為 undefined，不可假設跨執行保留。
handle 的 default value 與 flags 會在 whole graph update 時更新。

設值有兩種途徑

用 upstream kernel 呼叫 cudaGraphSetConditional() 動態設值（如 IF/SWITCH 範例），或在 handle 建立時給 cudaGraphCondAssignDefault 預設值（如 WHILE 範例，避免額外 upstream kernel）。

Conditional Node Body Graph Requirements

body graph 的內容限制

General：所有節點須在單一 device；只能含 kernel、empty、memcpy、memset、child graph、conditional 節點。

Kernel：body 內 kernel 不可使用 CUDA Dynamic Parallelism 或 Device Graph Launch；Cooperative launch 允許，但前提是未使用 MPS。
Memcpy/Memset：只允許涉及 device memory 與／或 pinned device-mapped host memory 的拷貝/設值；不允許涉及 CUDA arrays。
兩個運算元在 instantiation 時都須能被 current device 存取；即使目標是另一 device 的記憶體，拷貝動作仍由 graph 所在的 device 執行。

Conditional IF Nodes

IF node 的 body graph 在 condition 非零時執行一次。size = 1 只有 if body；size = 2 同時有 if 與 else body（else 在 condition 為零時執行一次）。

   A (set handle via upstream kernel)
   |
   B (conditional IF)
   |  condition != 0 ? ─► [ if body graph ]   (phGraph_out[0])
   |  condition == 0 ? ─► [ else body graph ] (phGraph_out[1], 需 size==2)
   |
   C

cudaGraphConditionalHandleCreate(&handle, graph);        // 無 default → undefined
// ... 加入 upstream setHandle kernel（kernelArgs = {&handle, &value}）...
cudaGraphNodeParams cParams = { cudaGraphNodeTypeConditional };
cParams.conditional.handle = handle;
cParams.conditional.type   = cudaGraphCondTypeIf;
cParams.conditional.size   = 1;                          // 只有 "if" body
cudaGraphAddNode(&node, graph, &node, 1, &cParams);      // 依賴 upstream node
cudaGraph_t bodyGraph = cParams.conditional.phGraph_out[0];  // 取出空 body 填充
cudaGraphAddNode(&node, bodyGraph, NULL, 0, &params);

重點：phGraph_out[] 是 cudaGraphAddNode 回填的 body graph 陣列；size = 2 時 phGraph_out[0] 是 if body、phGraph_out[1] 是 else body。

Conditional WHILE Nodes

WHILE node 的 body 只要 condition 非零就持續執行；condition 在節點執行時評估一次，並在每次 body 完成後再評估。

   A
   |
   B (conditional WHILE) ◄──┐
   |  condition != 0 ─► [ body graph ] ─┘  (執行後重新評估 condition)
   |  condition == 0 ─► 離開迴圈
   C

// 用 default value = 1，免去 upstream kernel
cudaGraphConditionalHandleCreate(&handle, graph, 1, cudaGraphCondAssignDefault);
cParams.conditional.type = cudaGraphCondTypeWhile;
cParams.conditional.size = 1;
cudaGraphAddNode(&node, graph, NULL, 0, &cParams);
cudaGraph_t bodyGraph = cParams.conditional.phGraph_out[0];
// body 內的 loopKernel：--(*dPtr)==0 時呼叫 cudaGraphSetConditional(handle, 0) 結束迴圈

重點：cudaGraphSetConditional() 在 device 端被 body kernel 呼叫來改變 condition，迴圈由 graph 自身在 device 上控制，host 不介入。

Conditional SWITCH Nodes

SWITCH node 執行第 n 個 zero-indexed body（當 condition == n）一次；若 condition 不對應任何 body，則不啟動任何 body graph。

   A (set handle = value)
   |
   B (conditional SWITCH, size = 5)
   |  condition==0 ─► phGraph_out[0]
   |  condition==1 ─► phGraph_out[1]
   |     ...
   |  condition==4 ─► phGraph_out[4]
   |  其他          ─► 不執行任何 body
   C

cParams.conditional.type = cudaGraphCondTypeSwitch;
cParams.conditional.size = 5;                            // 5 個 body graph
cudaGraphAddNode(&node, graph, &node, 1, &cParams);
cudaGraph_t *bodyGraphs = cParams.conditional.phGraph_out;  // 取出整個陣列
cudaGraphAddNode(&node, bodyGraphs[0], NULL, 0, &params);
// ...
cudaGraphAddNode(&node, bodyGraphs[4], NULL, 0, &params);

重點：SWITCH 的 size 即 body graph 數量，由 upstream kernel 設定的 condition value 決定走哪一個分支。

考試/測驗重點

主題	常考點
何時必須 re-instantiate	拓撲 / 節點類型改變；只改參數則可用 graph update
更新生效時機	下一次 launch 生效，不影響進行中的 launch
Whole vs individual	whole 用於更新多節點 / 拓撲未知；individual 用於少量、已持有 handle，跳過拓撲檢查
配對三規則	stream API 順序 / incoming-edge API 與陣列順序 / sink node 順序須一致
node enable	僅 kernel / memset / memcpy 可 enable；停用 = 空節點；不受 update 影響
1D 限制	memset / memcpy 節點只有 1D 才能 update
handle 特性	先於節點建立、關聯單一節點、無法 destroy；無 default → undefined
設值 API	device 端 `cudaGraphSetConditional()`；建立時 `cudaGraphCondAssignDefault` 給預設
IF size	size1 只有 if；size2 才有 else（condition==0 執行）
WHILE 評估時機	進入時評估一次 + 每次 body 完成後再評估
SWITCH 越界	condition 不對應任何 body → 不執行任何 body（非 fallthrough）
body graph 禁用	不可用 Dynamic Parallelism / Device Graph Launch；不可涉及 CUDA arrays；MPS 下不可 cooperative launch