CUDA 圖 可以顯著提高性能，因為驅動程序能夠使用任務和依賴關系的完整描述來優化執行。特別是在靜態工作流中，圖形可以提供難以置信的好處，其創建的開銷可以在多次連續啟動中分攤，從而提高整體性能。
然而，幾乎所有的問題都涉及某種形式的決策，這可能需要分解圖并將控制權返回給 CPU，以決定下一步要啟動哪些工作。像這樣分解工作會損害 CUDA 進行優化的能力，占用 CPU 資源，并增加每次圖形啟動的開銷。
從 CUDA 12.4 開始，CUDA Graphs 支持 條件節點，這使得圖形的部分能夠有條件地或重復地執行，而不需要將控制返回到 CPU。這釋放了 CPU 資源，使得更多的工作流能夠在單個圖形中表示，從而提高了計算效率。
條件節點
條件節點有兩種風格：
- IF 節點：如果條件值為 true，則每次評估節點時執行一次主體。
- WHILE 節點：只要條件值為 true，當對節點求值時，就會重復執行主體。
條件節點是容器節點，類似于子圖節點，但節點中包含的圖的執行取決于條件變量的值。與節點關聯的條件值由必須在節點之前創建的句柄訪問，該句柄還可以指定在圖形的每個開始處應用的初始化。可以通過調用在 CUDA 內核中設置條件值的 cudaGraphSetConditional 函數。
創建條件節點時，還會創建一個空圖，并將句柄返回給用戶。此圖與節點綁定，并將根據條件值執行。此條件體圖可以使用 圖形 API 或者通過使用捕獲異步 CUDA 調用的 cudaStreamBeginCaptureToGraph 函數。
條件節點也可以嵌套。例如，可以使用包含條件 IF 節點的體圖創建條件 WHILE 節點。
條件節點體圖可以包含以下任意一項：
- 內核節點（CNP，當前不支持協同）
- 空節點
- 子圖節點
- Memset 節點
- Memcopy 節點
- 條件節點
這遞歸地應用于子圖和條件體。所有內核，包括嵌套條件句或任何級別的子圖中的內核，都必須屬于同一個 CUDA 上下文。Memcopies 和 memset 必須作用于可從條件節點的上下文訪問的內存。
完整的樣品代碼可在 CUDA 樣本庫 中找到。下一節將通過一些示例來展示如何處理條件節點的方法。
條件 IF 節點
如果條件為非零，則無論何時評估 IF 節點，都將執行一次 IF 節點的體圖。圖 1 描述了一個圖，其中中間節點 B 是包含四節點圖的 IF 條件節點：
為了顯示如何創建此圖，以下示例使用條件節點 B 上游的內核節點 A，根據該內核所做的工作結果設置條件的值。條件的主體是使用圖 API 填充的。
首先，定義節點 A 內核。這個內核根據用戶執行的一些任意計算的結果來設置條件句柄。
__global__ void setHandle(cudaGraphConditionalHandle handle){ unsigned int value = 0; // We could perform some work here and set value based on the result of that work. if (someCondition) { // Set ‘value’ to non-zero if we want the conditional body to execute value = 1; } cudaGraphSetConditional(handle, value);} |
接下來，定義一個函數來構建圖。此函數分配條件句柄、創建節點并填充條件圖的主體。為了清楚起見，省略了啟動和執行圖形的代碼。
cudaGraph_t createGraph() { cudaGraph_t graph; cudaGraphNode_t node; void *kernelArgs[1]; cudaGraphCreate(&graph, 0); cudaGraphConditionalHandle handle; cudaGraphConditionalHandleCreate(&handle, graph); // Use a kernel upstream of the conditional to set the handle value cudaGraphNodeParams kParams = { cudaGraphNodeTypeKernel }; kParams.kernel.func = (void *)setHandle; kParams.kernel.gridDim.x = kParams.kernel.gridDim.y = kParams.kernel.gridDim.z = 1; kParams.kernel.blockDim.x = kParams.kernel.blockDim.y = kParams.kernel.blockDim.z = 1; kParams.kernel.kernelParams = kernelArgs; kernelArgs[0] = &handle cudaGraphAddNode(&node, graph, NULL, 0, &kParams); cudaGraphNodeParams cParams = { cudaGraphNodeTypeConditional }; cParams.conditional.handle = handle; cParams.conditional.type = cudaGraphCondTypeIf; cParams.conditional.size = 1; cudaGraphAddNode(&node, graph, &node, 1, &cParams); cudaGraph_t bodyGraph = cParams.conditional.phGraph_out[0]; // Populate the body of the conditional node cudaGraphNode_t bodyNodes[4]; cudaGraphNodeParams params[4] = { ... }; // Setup kernel parameters as needed. cudaGraphAddNode(&bodyNodes[0], bodyGraph, NULL, 0, ¶ms[0]); cudaGraphAddNode(&bodyNodes[1], bodyGraph, &bodyNodes[0], 1, ¶ms[1]); cudaGraphAddNode(&bodyNodes[2], bodyGraph, &bodyNodes[0], 1, ¶ms[2]); cudaGraphAddNode(&bodyNodes[3], bodyGraph, &bodyNodes[1], 2, ¶ms[3]); return graph;} |
條件 WHILE 節點
只要條件為非零，WHILE 節點的體圖就會重復執行。將在執行節點時以及每次完成體圖后評估條件。下圖描述了一個三節點圖，其中中間節點 B 是包含三節點圖的 WHILE 條件節點。
為了了解如何創建此圖，以下示例將句柄的默認值設置為非零值，以便在默認情況下執行 WHILE 循環。將默認值設置為非零，并將條件值保留在條件上游的內核中不修改，這有效地產生了一個 do-while 循環，其中條件體總是至少執行一次。創建 WHILE 循環，其中循環體僅在條件為 true 時執行，需要執行一些計算并在節點 a 中適當設置條件句柄。
在上一個示例中，條件主體由圖 API 填充。在本例中，使用流捕獲填充條件的主體。
第一步是定義一個內核，在每次執行條件體時設置條件值。在本例中，句柄是基于遞減計數器的值設置的。
__global__ void loopKernel(cudaGraphConditionalHandle handle){ static int count = 10; cudaGraphSetConditional(handle, --count ? 1 : 0);} |
接下來，定義一個函數來構建圖。此函數分配條件句柄、創建節點并填充條件圖的主體。為了清楚起見，省略了啟動和執行圖形的代碼。
cudaGraph_t createGraph() { cudaGraph_t graph; cudaGraphNode_t nodes[3]; cudaGraphCreate(&graph, 0); // Insert kernel node A cudaGraphNodeParams params = ...; cudaGraphAddNode(&nodes[0], graph, NULL, 0, ¶ms); cudaGraphConditionalHandle handle; cudaGraphConditionalHandleCreate(&handle, graph, 1, cudaGraphCondAssignDefault); // Insert conditional node B cudaGraphNodeParams cParams = { cudaGraphNodeTypeConditional }; cParams.conditional.handle = handle; cParams.conditional.type = cudaGraphCondTypeWhile; cParams.conditional.size = 1; cudaGraphAddNode(&nodes[1], graph, &nodes[0], 1, &cParams); cudaGraph_t bodyGraph = cParams.conditional.phGraph_out[0]; cudaStream_t captureStream; cudaStreamCreate(&captureStream); // Fill out body graph with stream capture. cudaStreamBeginCaptureToGraph(captureStream, bodyGraph, nullptr, nullptr, 0, cudaStreamCaptureModeRelaxed); myKernel1<<<..., captureStream>>>(...); myKernel2<<<..., captureStream>>>(...); loopKernel<<<1, 1, 0, captureStream>>>(handle); cudaStreamEndCapture(captureStream, nullptr); cudaStreamDestroy(captureStream); // Insert kernel node C. params = ...; cudaGraphAddNode(&nodes[2], graph, &nodes[1], 1, ¶ms); return graph;} |
此示例使用 cudaStreamBeginCaptureToGraph，這是 CUDA 12.3 中新增加的一個 API，它使流捕獲能夠將節點插入到現有的圖中。借助這個 API，可以將多個單獨的捕獲組合到一個圖形對象中。此外，這個 API 還允許填充與條件節點一起創建的條件體圖對象。
結論
CUDA Graphs 為靜態工作流提供了難以置信的好處，在靜態工作流中，圖形創建的開銷可以在多次連續啟動中分攤。消除對圖形的分割并將控制權返回給 CPU，以決定優先啟動哪個，這有助于減少 CPU 開銷和延遲。使用具有條件節點的 CUDA 圖形可以有條件地或重復地執行圖形的部分，而無需將控制權返回給 CPU。這釋放了 CPU 資源，并使單個圖形能夠表示更復雜的工作流。
欲了解更多關于條件節點的信息，請參閱 CUDA 編程指南。要探索簡單、完整的示例代碼，請訪問 NVIDIA/cuda-samples 在 GitHub 上。同時，您也可以加入 NVIDIA 開發者 CUDA 論壇 中的討論。
?
