[SYCLomatic] Migration of cudaGraphNodeType, cudaGraphNodeGetType (#2760)

daiyaan-ahmed6 · web-flow · commit 3fdc31ab87ff · 2025-04-07T10:26:13.000+08:00
Signed-off-by: Daiyaan Ahmed &lt;daiyaan.ahmed@intel.com&gt;
diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp
@@ -628,6 +628,11 @@ void MapNames::setExplicitNamespaceMap(
            DpctGlobalInfo::useExtGraph()
                ? getClNamespace() + "ext::oneapi::experimental::queue_state"
                : "cudaStreamCaptureStatus")},
+      {"cudaGraphNodeType",
+       std::make_shared<TypeNameRule>(
+           DpctGlobalInfo::useExtGraph()
+               ? getClNamespace() + "ext::oneapi::experimental::node_type"
+               : "cudaGraphNodeType")},
       {"CUmem_advise", std::make_shared<TypeNameRule>("int")},
       {"CUmemorytype",
        std::make_shared<TypeNameRule>(getClNamespace() + "usm::alloc")},
@@ -1092,6 +1097,43 @@ void MapNames::setExplicitNamespaceMap(
                : "cudaStreamCaptureStatusActive")},
       {"cudaStreamCaptureStatusInvalidated",
        std::make_shared<EnumNameRule>("cudaStreamCaptureStatusInvalidated")},
+      // enum cudaGraphNodeType
+      {"cudaGraphNodeTypeKernel",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtGraph()
+               ? getClNamespace() +
+                     "ext::oneapi::experimental::node_type::kernel"
+               : "cudaGraphNodeTypeKernel")},
+      {"cudaGraphNodeTypeMemcpy",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtGraph()
+               ? getClNamespace() +
+                     "ext::oneapi::experimental::node_type::memcpy"
+               : "cudaGraphNodeTypeMemcpy")},
+      {"cudaGraphNodeTypeMemset",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtGraph()
+               ? getClNamespace() +
+                     "ext::oneapi::experimental::node_type::memset"
+               : "cudaGraphNodeTypeMemset")},
+      {"cudaGraphNodeTypeHost",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtGraph()
+               ? getClNamespace() +
+                     "ext::oneapi::experimental::node_type::host_task"
+               : "cudaGraphNodeTypeHost")},
+      {"cudaGraphNodeTypeGraph",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtGraph()
+               ? getClNamespace() +
+                     "ext::oneapi::experimental::node_type::subgraph"
+               : "cudaGraphNodeTypeGraph")},
+      {"cudaGraphNodeTypeEmpty",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtGraph()
+               ? getClNamespace() +
+                     "ext::oneapi::experimental::node_type::empty"
+               : "cudaGraphNodeTypeEmpty")},
       // enum CUmem_advise_enum
       {"CU_MEM_ADVISE_SET_READ_MOSTLY", std::make_shared<EnumNameRule>("0")},
       {"CU_MEM_ADVISE_UNSET_READ_MOSTLY", std::make_shared<EnumNameRule>("0")},
diff --git a/clang/lib/DPCT/RulesLang/APINamesGraph.inc b/clang/lib/DPCT/RulesLang/APINamesGraph.inc
@@ -66,3 +66,13 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cudaGraphExecUpdate"),
                             ARG("--use-experimental-features=graph"))))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtGraph,
+    ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY("cudaGraphNodeGetType", DEREF(1),
+                                            MEMBER_CALL(ARG(0), true,
+                                                        "get_type"))),
+    UNSUPPORT_FACTORY_ENTRY("cudaGraphNodeGetType",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaGraphNodeGetType"),
+                            ARG("--use-experimental-features=graph")))
diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp
@@ -347,7 +347,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "cublasLtMatrixTransformDesc_t", "cudaGraphicsMapFlags",
               "cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
               "cudaExternalSemaphoreHandleType", "CUstreamCallback",
-              "cudaHostFn_t", "__nv_half2", "__nv_half"))))))
+              "cudaHostFn_t", "__nv_half2", "__nv_half",
+              "cudaGraphNodeType"))))))
           .bind("cudaTypeDef"),
       this);
 
@@ -921,6 +922,13 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
       }
     }
 
+    if (CanonicalTypeStr == "cudaGraphNodeType") {
+      if (!DpctGlobalInfo::useExtGraph()) {
+        report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
+               "cudaGraphNodeType", "--use-experimental-features=graph");
+      }
+    }
+
     if (CanonicalTypeStr == "cudaGraphExecUpdateResult") {
       report(TL->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
              CanonicalTypeStr);
@@ -1909,17 +1917,17 @@ void DeviceInfoVarRule::runRule(const MatchFinder::MatchResult &Result) {
 // Rule for Enums constants.
 void EnumConstantRule::registerMatcher(MatchFinder &MF) {
   MF.addMatcher(
-      declRefExpr(
-          to(enumConstantDecl(anyOf(
-              hasType(enumDecl(hasAnyName(
-                  "cudaComputeMode", "cudaMemcpyKind", "cudaMemoryAdvise",
-                  "cudaStreamCaptureStatus", "cudaDeviceAttr",
-                  "libraryPropertyType_t", "cudaDataType_t",
-                  "CUmem_advise_enum", "cufftType_t",
-                  "cufftType", "cudaMemoryType", "CUctx_flags_enum",
-                  "CUpointer_attribute_enum", "CUmemorytype_enum",
-                  "cudaGraphicsMapFlags", "cudaGraphicsRegisterFlags"))),
-              matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*")))))
+      declRefExpr(to(enumConstantDecl(anyOf(
+                      hasType(enumDecl(hasAnyName(
+                          "cudaComputeMode", "cudaMemcpyKind",
+                          "cudaMemoryAdvise", "cudaStreamCaptureStatus",
+                          "cudaDeviceAttr", "libraryPropertyType_t",
+                          "cudaDataType_t", "CUmem_advise_enum", "cufftType_t",
+                          "cufftType", "cudaMemoryType", "CUctx_flags_enum",
+                          "CUpointer_attribute_enum", "CUmemorytype_enum",
+                          "cudaGraphicsMapFlags", "cudaGraphicsRegisterFlags",
+                          "cudaGraphNodeType"))),
+                      matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*")))))
           .bind("EnumConstant"),
       this);
 }
@@ -1991,7 +1999,14 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) {
       EnumName == "cudaExternalSemaphoreHandleTypeKeyedMutex" ||
       EnumName == "cudaExternalSemaphoreHandleTypeKeyedMutexKmt" ||
       EnumName == "cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd" ||
-      EnumName == "cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32") {
+      EnumName == "cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32" ||
+      EnumName == "cudaGraphNodeTypeWaitEvent" ||
+      EnumName == "cudaGraphNodeTypeEventRecord" ||
+      EnumName == "cudaGraphNodeTypeExtSemaphoreSignal" ||
+      EnumName == "cudaGraphNodeTypeExtSemaphoreWait" ||
+      EnumName == "cudaGraphNodeTypeMemAlloc" ||
+      EnumName == "cudaGraphNodeTypeMemFree" ||
+      EnumName == "cudaGraphNodeTypeConditional") {
     report(E->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumName);
     return;
   } else if (EnumName == "cudaComputeModeDefault" ||
@@ -2024,6 +2039,16 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) {
     report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
            EnumName, "--use-experimental-features=bindless_images");
     return;
+  } else if (!DpctGlobalInfo::useExtGraph() &&
+             (EnumName == "cudaGraphNodeTypeKernel" ||
+              EnumName == "cudaGraphNodeTypeMemcpy" ||
+              EnumName == "cudaGraphNodeTypeMemset" ||
+              EnumName == "cudaGraphNodeTypeHost" ||
+              EnumName == "cudaGraphNodeTypeGraph" ||
+              EnumName == "cudaGraphNodeTypeEmpty")) {
+    report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
+           EnumName, "--use-experimental-features=graph");
+    return;
   } else if (auto ET = dyn_cast<EnumType>(E->getType())) {
     if (auto ETD = ET->getDecl()) {
       auto EnumTypeName = ETD->getName().str();
diff --git a/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp b/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp
@@ -32,7 +32,8 @@ void GraphRule::registerMatcher(MatchFinder &MF) {
   auto functionName = [&]() {
     return hasAnyName("cudaGraphInstantiate", "cudaGraphLaunch",
                       "cudaGraphExecDestroy", "cudaGraphAddEmptyNode",
-                      "cudaGraphAddDependencies", "cudaGraphExecUpdate");
+                      "cudaGraphAddDependencies", "cudaGraphExecUpdate",
+                      "cudaGraphNodeGetType");
   };
   MF.addMatcher(
       callExpr(callee(functionDecl(functionName()))).bind("FunctionCall"),
diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc
@@ -479,7 +479,7 @@ ENTRY(cudaGraphNodeGetDependencies_v2, cudaGraphNodeGetDependencies_v2, false, N
 ENTRY(cudaGraphNodeGetDependentNodes, cudaGraphNodeGetDependentNodes, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphNodeGetDependentNodes_v2, cudaGraphNodeGetDependentNodes_v2, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphNodeGetEnabled, cudaGraphNodeGetEnabled, false, NO_FLAG, P4, "comment")
-ENTRY(cudaGraphNodeGetType, cudaGraphNodeGetType, false, NO_FLAG, P4, "comment")
+ENTRY(cudaGraphNodeGetType, cudaGraphNodeGetType, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cudaGraphNodeSetEnabled, cudaGraphNodeSetEnabled, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphNodeSetParams, cudaGraphNodeSetParams, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphReleaseUserObject, cudaGraphReleaseUserObject, false, NO_FLAG, P4, "comment")
diff --git a/clang/test/dpct/cudaGraphNodeType_test.cu b/clang/test/dpct/cudaGraphNodeType_test.cu
@@ -0,0 +1,82 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3, cuda-11.4, cuda-11.5, cuda-11.6, cuda-11.7, cuda-11.8, cuda-12.0, cuda-12.1, cuda-12.2, cuda-12.3
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3, v11.4, v11.5, v11.6, v11.7, v11.8, v12.0, v12.1, v12.2, v12.3
+// RUN: dpct --use-experimental-features=graph --format-range=none -out-root %T/cudaGraphNodeType_test %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only --std=c++14
+// RUN: FileCheck --input-file %T/cudaGraphNodeType_test/cudaGraphNodeType_test.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -DNO_BUILD_TEST -fsycl %T/cudaGraphNodeType_test/cudaGraphNodeType_test.dp.cpp -o %T/cudaGraphNodeType_test/cudaGraphNodeType_test.dp.o %}
+
+#include <cuda.h>
+#define CUDA_CHECK_THROW(x)  \
+  do {                       \
+    cudaError_t _result = x; \
+  } while (0)
+
+int main() {
+  // CHECK: sycl::ext::oneapi::experimental::node_type nodeType;
+  cudaGraphNodeType nodeType;
+
+  // CHECK: nodeType = sycl::ext::oneapi::experimental::node_type::kernel;
+  nodeType = cudaGraphNodeTypeKernel;
+
+  // CHECK: nodeType = sycl::ext::oneapi::experimental::node_type::memcpy;
+  nodeType = cudaGraphNodeTypeMemcpy;
+
+  // CHECK: nodeType = sycl::ext::oneapi::experimental::node_type::memset;
+  nodeType = cudaGraphNodeTypeMemset;
+
+  // CHECK: nodeType = sycl::ext::oneapi::experimental::node_type::host_task;
+  nodeType = cudaGraphNodeTypeHost;
+
+  // CHECK: nodeType = sycl::ext::oneapi::experimental::node_type::subgraph;
+  nodeType = cudaGraphNodeTypeGraph;
+
+  // CHECK: nodeType = sycl::ext::oneapi::experimental::node_type::empty;
+  nodeType = cudaGraphNodeTypeEmpty;
+
+#ifndef NO_BUILD_TEST
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphNodeTypeWaitEvent is not supported.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: nodeType = cudaGraphNodeTypeWaitEvent;
+  nodeType = cudaGraphNodeTypeWaitEvent;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphNodeTypeEventRecord is not supported.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: nodeType = cudaGraphNodeTypeEventRecord;
+  nodeType = cudaGraphNodeTypeEventRecord;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphNodeTypeExtSemaphoreSignal is not supported.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: nodeType = cudaGraphNodeTypeExtSemaphoreSignal;
+  nodeType = cudaGraphNodeTypeExtSemaphoreSignal;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphNodeTypeExtSemaphoreWait is not supported.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: nodeType = cudaGraphNodeTypeExtSemaphoreWait;
+  nodeType = cudaGraphNodeTypeExtSemaphoreWait;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphNodeTypeMemAlloc is not supported.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: nodeType = cudaGraphNodeTypeMemAlloc;
+  nodeType = cudaGraphNodeTypeMemAlloc;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphNodeTypeMemFree is not supported.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: nodeType = cudaGraphNodeTypeMemFree;
+  nodeType = cudaGraphNodeTypeMemFree;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphNodeTypeConditional is not supported.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: nodeType = cudaGraphNodeTypeConditional;
+  nodeType = cudaGraphNodeTypeConditional;
+
+#endif
+
+  return 0;
+}
diff --git a/clang/test/dpct/cudaGraph_test.cu b/clang/test/dpct/cudaGraph_test.cu
@@ -107,6 +107,13 @@ int main() {
   CUDA_CHECK_THROW(cudaGraphExecUpdate(execGraph, graph, nullptr, nullptr));
 #endif
 
+  // CHECK: sycl::ext::oneapi::experimental::node_type nodeType;
+  // CHECK-NEXT: nodeType = node->get_type();
+  // CHECK-NEXT: CUDA_CHECK_THROW(DPCT_CHECK_ERROR(nodeType = node->get_type()));
+  cudaGraphNodeType nodeType;
+  cudaGraphNodeGetType(node, &nodeType);
+  CUDA_CHECK_THROW(cudaGraphNodeGetType(node, &nodeType));
+
   // CHECK: delete (execGraph);
   // CHECK-NEXT: delete (*execGraph2);
   // CHECK-NEXT:  delete (**execGraph3);
diff --git a/clang/test/dpct/cudaGraph_test_default_option.cu b/clang/test/dpct/cudaGraph_test_default_option.cu
@@ -97,6 +97,21 @@ int main() {
   // CHECK-NEXT: */
   cudaGraphExecDestroy(execGraph);
 
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphNodeType is not supported, please try to remigrate with option: --use-experimental-features=graph.
+  // CHECK-NEXT: */
+  cudaGraphNodeType nodeType;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphNodeGetType is not supported, please try to remigrate with option: --use-experimental-features=graph.
+  // CHECK-NEXT: */
+  cudaGraphNodeGetType(node, &nodeType);
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphNodeTypeKernel is not supported, please try to remigrate with option: --use-experimental-features=graph.
+  // CHECK-NEXT: */
+  nodeType = cudaGraphNodeTypeKernel;
+
   return 0;
 }
 
