[SYCLomatic] Migration of cudaGraphCreate, cudaGraphDestroy

clang/lib/DPCT/APINames.inc

@@ -414,10 +414,10 @@ ENTRY(cudaGraphAddNode, cudaGraphAddNode, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphAddNode_v2, cudaGraphAddNode_v2, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphChildGraphNodeGetGraph, cudaGraphChildGraphNodeGetGraph, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphClone, cudaGraphClone, false, NO_FLAG, P4, "comment")
+ENTRY(cudaGraphCreate, cudaGraphCreate, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cudaGraphConditionalHandleCreate, cudaGraphConditionalHandleCreate, false, NO_FLAG, P4, "comment")
-ENTRY(cudaGraphCreate, cudaGraphCreate, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphDebugDotPrint, cudaGraphDebugDotPrint, false, NO_FLAG, P4, "comment")
-ENTRY(cudaGraphDestroy, cudaGraphDestroy, false, NO_FLAG, P4, "comment")
+ENTRY(cudaGraphDestroy, cudaGraphDestroy, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cudaGraphDestroyNode, cudaGraphDestroyNode, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphEventRecordNodeGetEvent, cudaGraphEventRecordNodeGetEvent, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphEventRecordNodeSetEvent, cudaGraphEventRecordNodeSetEvent, false, NO_FLAG, P4, "comment")

clang/lib/DPCT/APINamesGraph.inc

@@ -66,3 +66,26 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cudaGraphExecUpdate"),
                             ARG("--use-experimental-features=graph"))))
+
+ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
+    UseExtGraph,
+    CALL_FACTORY_ENTRY(
+        "cudaGraphCreate",
+        CALL(MapNames::getDpctNamespace() + "experimental::create_graph",
+             ARG(0), CALL(MapNames::getDpctNamespace() + "get_default_context"),
+             CALL(MapNames::getDpctNamespace() + "get_current_device"))),
+    UNSUPPORT_FACTORY_ENTRY("cudaGraphCreate",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaGraphCreate"),
+                            ARG("--use-experimental-features=graph"))))
+
+ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
+    UseExtGraph,
+    CALL_FACTORY_ENTRY("cudaGraphDestroy",
+                       CALL(MapNames::getDpctNamespace() +
+                                "experimental::destroy_graph",
+                            ARG(0))),
+    UNSUPPORT_FACTORY_ENTRY("cudaGraphDestroy",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaGraphDestroy"),
+                            ARG("--use-experimental-features=graph"))))

clang/lib/DPCT/ASTTraversal.cpp

@@ -15062,7 +15062,8 @@ void GraphRule::registerMatcher(MatchFinder &MF) {
   auto functionName = [&]() {
     return hasAnyName("cudaGraphInstantiate", "cudaGraphLaunch",
                       "cudaGraphExecDestroy", "cudaGraphAddEmptyNode",
-                      "cudaGraphAddDependencies", "cudaGraphExecUpdate");
+                      "cudaGraphAddDependencies", "cudaGraphExecUpdate",
+                      "cudaGraphCreate", "cudaGraphDestroy");
   };
   MF.addMatcher(
       callExpr(callee(functionDecl(functionName()))).bind("FunctionCall"),
@@ -15074,6 +15075,11 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) {
   if (!CE) {
     return;
   }
+  std::string FuncName =
+      CE->getDirectCallee()->getNameInfo().getName().getAsString();
+  if (FuncName == "cudaGraphCreate") {
+    report(CE->getBeginLoc(), Diagnostics::GRAPH_DEVICE_MATCH, false);
+  }
   ExprAnalysis EA(CE);
   emplaceTransformation(EA.getReplacement());
   EA.applyAllSubExprRepl();

clang/lib/DPCT/Diagnostics.inc

@@ -286,6 +286,8 @@ DEF_WARNING(NOT_DEVICE_COPYABLE, 1128, MEDIUM_LEVEL, "The type \"%0\" is not dev
 DEF_COMMENT(NOT_DEVICE_COPYABLE, 1128, MEDIUM_LEVEL, "The type \"{0}\" is not device copyable for {1}. It is used in the SYCL kernel, please rewrite the code.")
 DEF_WARNING(NOT_DEVICE_COPYABLE_ADD_SPECIALIZATION, 1129, MEDIUM_LEVEL, "The type \"%0\" is used in the SYCL kernel, but it is not device copyable. The sycl::is_device_copyable specialization has been added for this type. Please review the code.")
 DEF_COMMENT(NOT_DEVICE_COPYABLE_ADD_SPECIALIZATION, 1129, MEDIUM_LEVEL, "The type \"{0}\" is used in the SYCL kernel, but it is not device copyable. The sycl::is_device_copyable specialization has been added for this type. Please review the code.")
+DEF_WARNING(GRAPH_DEVICE_MATCH, 1130, MEDIUM_LEVEL, "Verify that the device used to create the graph matches the device used to launch the graph.")
+DEF_COMMENT(GRAPH_DEVICE_MATCH, 1130, MEDIUM_LEVEL, "Verify that the device used to create the graph matches the device used to launch the graph.")
 DEF_WARNING(NOT_SUPPORT_DYN_PARALLEL, 1130, HIGH_LEVEL, "SYCL 2020 standard does not support dynamic parallelism (launching kernel in device code). Please rewrite the code.")
 DEF_COMMENT(NOT_SUPPORT_DYN_PARALLEL, 1130, HIGH_LEVEL, "SYCL 2020 standard does not support dynamic parallelism (launching kernel in device code). Please rewrite the code.")
 DEF_WARNING(UNSUPPORT_SYCLCOMPAT, 1131, MEDIUM_LEVEL, "The migration of \"%0\" is not supported with SYCLcompat currently, please adjust the code manually.")

clang/runtime/dpct-rt/include/dpct/graph.hpp

@@ -133,5 +133,26 @@ static void add_dependencies(dpct::experimental::command_graph_ptr graph,
   }
 }
 
+/// Creates a new command graph.
+/// \param [out] graph A pointer to a command_graph_ptr pointer where the
+/// command graph will be assigned.
+/// \param [in] context A SYCL context.
+/// \param [in] device A SYCL device.
+static void create_graph(dpct::experimental::command_graph_ptr *graph,
+                         sycl::context context, dpct::device_ext *device) {
+  *graph = new sycl::ext::oneapi::experimental::command_graph<
+      sycl::ext::oneapi::experimental::graph_state::modifiable>(context,
+                                                                *device);
+}
+
+/// Destroys the command graph.
+/// \param [in] graph A pointer to the command graph.
+static void destroy_graph(dpct::experimental::command_graph_ptr graph) {
+  std::vector<sycl::ext::oneapi::experimental::node> nodes = graph->get_nodes();
+  for (auto node : nodes) {
+    delete &node;
+  }
+  delete graph;
+}
 } // namespace experimental
 } // namespace dpct

clang/test/dpct/cudaGraph_test.cu

@@ -24,6 +24,17 @@ int main() {
   // CHECK: dpct::experimental::command_graph_ptr graph5, *graph6, **graph7;
   cudaGraph_t graph5, *graph6, **graph7;
 
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1130:{{[0-9]+}}: Verify that the device used to create the graph matches the device used to launch the graph.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: dpct::experimental::create_graph(&graph, dpct::get_default_context(), dpct::get_current_device());
+  // CHECK-NEXT: /*
+  // CHECK-NEXT: DPCT1130:{{[0-9]+}}: Verify that the device used to create the graph matches the device used to launch the graph.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: dpct::experimental::create_graph(graph2, dpct::get_default_context(), dpct::get_current_device());
+  cudaGraphCreate(&graph, 0);
+  cudaGraphCreate(graph2, 0);
+
   // CHECK: dpct::experimental::command_graph_exec_ptr execGraph;
   // CHECK-NEXT: dpct::experimental::command_graph_exec_ptr *execGraph2;
   // CHECK-NEXT: dpct::experimental::command_graph_exec_ptr **execGraph3;
@@ -116,5 +127,10 @@ int main() {
   cudaGraphExecDestroy(**execGraph3);
   CUDA_CHECK_THROW(cudaGraphExecDestroy(**execGraph3));
 
+  // CHECK: dpct::experimental::destroy_graph(graph);
+  // CHECK-NEXT: dpct::experimental::destroy_graph(*graph2);
+  cudaGraphDestroy(graph);
+  cudaGraphDestroy(*graph2);
+
   return 0;
 }