自從CUDA Graphs在CUDA 10中引入以來，CUDA Graph已經用于各種應用中。圖形將一組CUDA內核和其他CUDA操作組合在一起，并使用指定的依賴關系樹執行它們。它通過結合與CUDA內核啟動和CUDA API調用相關的驅動程序活動來加快工作流。它還通過硬件加速強制實施依賴關系，而不是在可能的情況下僅依賴CUDA流和事件。

構造CUDA圖有兩種主要方法：顯式API調用和流捕獲。

使用顯式API調用構造CUDA圖

通過這種構建CUDA圖的方法，由CUDA內核和CUDA內存操作形成的圖節點通過調用cudaGraphAdd*節點API添加到圖中，其中*被替換為節點類型。節點之間的依賴關系是用API顯式設置的。

使用顯式API構建CUDA圖的好處是，cudaGraphAdd*Node API返回節點句柄（cudaGraph Node_t），可以用作未來節點更新的引用。例如，可以使用cudaGraphExecKernelNodeSetParams以最低成本更新實例化圖中內核節點的內核啟動配置和內核函數參數。

缺點是，在使用CUDA圖加速現有代碼的場景中，使用顯式API調用構造CUDA圖通常需要大量代碼更改，尤其是有關代碼的控制流和函數調用結構的更改。

使用流捕獲構建CUDA圖

通過這種構建CUDA圖的方法，cudaStreamBeginCapture和cudaStream EndCapture被放置在代碼塊的前后。代碼塊啟動的所有設備活動都會被記錄、捕獲并分組到CUDA圖中。節點之間的依賴關系是從流捕獲區域內的CUDA流或事件API調用中推斷出來的。

使用流捕獲構建CUDA圖的好處是，對于現有代碼，需要的代碼更改更少。原始代碼結構可以基本保持不變，圖形構造是以自動方式執行的。

這種構建CUDA圖的方法也有缺點。在流捕獲區域內，所有內核啟動配置和內核函數參數以及CUDA API調用參數都按值記錄。每當任何配置和參數發生更改時，捕獲的然后實例化的圖形就會過期。

在《在動態環境中使用CUDA圖》一文中提供了兩種解決方案：

重新捕獲工作流。當重新捕獲的圖與實例化的圖具有相同的節點拓撲時，不需要重新實例化，并且可以使用cudaGraphExecUpdate執行整個圖更新。

以配置和參數集作為鍵緩存CUDA圖。每組配置和參數都與緩存中的不同CUDA圖相關聯。在運行工作流時，配置和參數集首先被抽象為一個鍵。然后在緩存中找到相應的圖（如果它已經存在）并啟動。

然而，在某些工作流中，兩種解決方案都不能很好地工作。重新捕獲然后更新方法在紙面上很有效，但在某些情況下，重新捕獲和更新本身的成本很高。也有一些情況下，無法將每組參數與CUDA圖相關聯。例如，具有浮點數字參數的情況很難緩存，因為可能存在大量的浮點數字。

用顯式API構造的CUDA圖很容易更新，但這種方法可能過于繁瑣，靈活性較差。CUDA圖可以通過流捕獲靈活地構造，但生成的圖很難更新，而且更新成本很高。

綜合方法

在本文中，我提供了一種使用顯式API和流捕獲方法構建CUDA圖的方法，從而實現兩者的優點，避免兩者的缺點。

例如，在順序啟動三個內核的工作流中，前兩個內核具有靜態啟動配置和參數，而最后一個內核具有動態啟動配置和屬性。

使用流捕獲來記錄前兩個內核的啟動，并調用顯式API將最后一個內核節點添加到捕獲圖中。然后，顯式API返回的節點句柄用于在每次啟動圖之前用動態配置和參數更新實例化圖。

下面的代碼示例說明了這個想法：

cudaStream_t stream;
std::vector _node_list;
cudaGraphExec_t _graph_exec;

if (not using_graph) {
  
  first_static_kernel<<<1, 1, 0, stream>>>(static_parameters);
  second_static_kernel<<<1, 1, 0, stream>>>(static_parameters);
  dynamic_kernel<<<1, 1, 0, stream>>>(dynamic_parameters);

} else {

  if (capturing_graph) {

    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

    first_static_kernel<<<1, 1, 0, stream>>>(static_parameters);
    second_static_kernel<<<1, 1, 0, stream>>>(static_parameters);

    // Get the current stream capturing graph
    cudaGraph_t _capturing_graph;
    cudaStreamCaptureStatus _capture_status;
    const cudaGraphNode_t *_deps;
    size_t _dep_count;
    cudaStreamGetCaptureInfo_v2(stream, &_capture_status, nullptr &_capturing_graph, &_deps, &_dep_count);

    // Manually add a new kernel node
    cudaGraphNode_t new_node;
    cudakernelNodeParams _dynamic_params_cuda;
    cudaGraphAddKernelNode(&new_node, _capturing_graph, _deps, _dep_count, &_dynamic_params_cuda);
    // ... and store the new node for future references
    _node_list.push_back(new_node);

    // Update the stream dependencies
    cudaStreamUpdateCaptureDependencies(stream, &new_node, 1, 1); 

    // End the capture and instantiate the graph
    cudaGraph_t _captured_graph;
    cudaStreamEndCapture(stream, &_captured_graph);

    cudaGraphInstantiate(&_graph_exec, _captured_graph, nullptr, nullptr, 0);

  } else if (updating_graph) {
    cudakernelNodeParams _dynamic_params_updated_cuda;
  
    cudaGraphExecKernelNodeSetParams(_graph_exec, _node_list[0], &_dynamic_params_updated_cuda);

  }
}

在此示例中，cudaStreamGetCaptureInfo_v2提取當前正在記錄并捕獲到的CUDA圖形。在調用cudaStreamUpdateCaptureDependencies更新當前捕獲流的依賴項樹之前，會將一個內核節點添加到此圖中，并返回和存儲節點句柄（new_node）。最后一步是必要的，以確保隨后捕獲的任何其他活動都在這些手動添加的節點上正確設置了它們的依賴項。

使用這種方法，即使參數是動態的，也可以通過輕量級的cudaGraphExecKernelNodeSetParams調用直接重用相同的實例化圖（cudaGraph Exec_t對象）。本文中的第一張圖片顯示了這種用法。

此外，捕獲和更新代碼路徑可以組合成一段代碼，與啟動最后兩個內核的原始代碼相鄰。這會造成最少的代碼更改，并且不會破壞原始的控制流和函數調用結構。

新方法在帶有動態參數的蜂樹/cuda圖獨立代碼示例中詳細顯示。cudaStreamGetCaptureInfo_v2和cudaStream UpdateCaptureDependencies是CUDA 11.3中引入的新CUDA運行時API。

績效結果

使用帶有動態參數的蜂巢樹/cuda圖獨立代碼示例，我用三種不同的方法測量了運行受內核啟動開銷約束的相同動態工作流的性能：

在沒有CUDA圖形加速的情況下運行

使用重新捕獲然后更新方法運行CUDA圖

使用本文介紹的組合方法運行CUDA圖

表1顯示了結果。本文中提到的方法的提速很大程度上取決于底層工作流。

結論

在本文中，我介紹了一種結合顯式API和流捕獲方法構建CUDA圖的方法。它提供了一種以最低成本為具有動態參數的工作流重用實例化圖的方法。

關于作者

Tu Jiqun在加入NVIDIA擔任高級計算機開發技術工程師之前，曾獲得哥倫比亞大學晶格QCD物理學博士學位。在NVIDIA，他致力于在最新的NVIDIAGPU上使用最新的硬件和軟件功能，以加速廣泛的HPC應用程序。

審核編輯：郭婷