Migratiion of cudaGraphSetKernelNodeParams

Signed-off-by: Daiyaan Ahmed <[email protected]>

Author: daiyaan-ahmed6

clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc

@@ -507,6 +507,15 @@ TYPE_REWRITE_ENTRY(
         WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
                         STR("--use-experimental-features=graph"))))
 
+TYPE_REWRITE_ENTRY(
+    "cudaKernelNodeParams",
+    TYPE_CONDITIONAL_FACTORY(
+        checkEnableGraphForType(),
+        TYPE_FACTORY(STR(MapNames::getDpctNamespace() +
+                         "experimental::kernel_node_params")),
+        WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
+                        STR("--use-experimental-features=graph"))))
+
 // Graphics Interop Handle
 TYPE_REWRITE_ENTRY(
     "cudaGraphicsResource",

clang/lib/DPCT/RulesLang/APINamesGraph.inc

@@ -104,3 +104,25 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cudaGraphDestroy"),
                             ARG("--use-experimental-features=graph"))))
+
+ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
+    UseExtGraph,
+    CALL_FACTORY_ENTRY("cudaGraphKernelNodeGetParams",
+                       CALL(MapNames::getDpctNamespace() +
+                                "experimental::get_kernel_node_params",
+                            ARG(0), ARG(1))),
+    UNSUPPORT_FACTORY_ENTRY("cudaGraphKernelNodeGetParams",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaGraphKernelNodeGetParams"),
+                            ARG("--use-experimental-features=graph"))))
+
+ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
+    UseExtGraph,
+    CALL_FACTORY_ENTRY("cudaGraphKernelNodeSetParams",
+                       CALL(MapNames::getDpctNamespace() +
+                                "experimental::set_kernel_node_params",
+                            ARG(0), ARG(1))),
+    UNSUPPORT_FACTORY_ENTRY("cudaGraphKernelNodeSetParams",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaGraphKernelNodeSetParams"),
+                            ARG("--use-experimental-features=graph"))))

clang/lib/DPCT/RulesLang/MapNamesLang.cpp

@@ -362,5 +362,13 @@ const std::unordered_map<std::string, HelperFeatureEnum>
         {"sampler", HelperFeatureEnum::device_ext},
 };
 
+// Graph kernel node params mapping
+MapNamesLang::MapTy GraphRule::KernelNodeParamNames{
+    {"gridDim", "grid_dim"},
+    {"blockDim", "block_dim"},
+    {"kernelParams", "kernel_params"},
+    {"sharedMemBytes", "shared_mem_bytes"},
+    {"func", "func"}};
+
 } // namespace dpct
 } // namespace clang

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -353,18 +353,19 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
       this);
 
   MF.addMatcher(
-      typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName(
-                  "cooperative_groups::__v1::coalesced_group",
-                  "cooperative_groups::__v1::grid_group",
-                  "cooperative_groups::__v1::thread_block_tile", "cudaGraph_t",
-                  "cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource",
-                  "cudaGraphicsResource_t", "CUgraphicsResource",
-                  "cudaExternalMemory_t", "cudaExternalMemoryHandleDesc",
-                  "cudaExternalMemoryMipmappedArrayDesc",
-                  "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t",
-                  "cudaExternalSemaphoreHandleDesc",
-                  "cudaExternalSemaphoreSignalParams",
-                  "cudaExternalSemaphoreWaitParams"))))))
+      typeLoc(
+          loc(qualType(hasDeclaration(namedDecl(hasAnyName(
+              "cooperative_groups::__v1::coalesced_group",
+              "cooperative_groups::__v1::grid_group",
+              "cooperative_groups::__v1::thread_block_tile", "cudaGraph_t",
+              "cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource",
+              "cudaGraphicsResource_t", "CUgraphicsResource",
+              "cudaExternalMemory_t", "cudaExternalMemoryHandleDesc",
+              "cudaExternalMemoryMipmappedArrayDesc",
+              "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t",
+              "cudaExternalSemaphoreHandleDesc",
+              "cudaExternalSemaphoreSignalParams",
+              "cudaExternalSemaphoreWaitParams", "cudaKernelNodeParams"))))))
           .bind("cudaTypeDefEA"),
       this);
   MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl(

clang/lib/DPCT/RulesLang/RulesLang.h

@@ -999,6 +999,10 @@ class CompatWithClangRule : public NamedMigrationRule<CompatWithClangRule> {
 };
 
 class GraphRule : public NamedMigrationRule<GraphRule> {
+  static MapNames::MapTy KernelNodeParamNames;
+  const Expr *getAssignedBO(const Expr *E, ASTContext &Context);
+  const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context);
+
 public:
   void registerMatcher(ast_matchers::MatchFinder &MF) override;
   void runRule(const ast_matchers::MatchFinder::MatchResult &Result);

clang/lib/DPCT/RulesLang/RulesLangGraph.cpp

@@ -34,14 +34,48 @@ void GraphRule::registerMatcher(MatchFinder &MF) {
                       "cudaGraphExecDestroy", "cudaGraphAddEmptyNode",
                       "cudaGraphAddDependencies", "cudaGraphExecUpdate",
                       "cudaGraphNodeGetType", "cudaGraphGetNodes",
-                      "cudaGraphGetRootNodes", "cudaGraphDestroy");
+                      "cudaGraphGetRootNodes", "cudaGraphDestroy", "cudaGraphKernelNodeGetParams",
+                      "cudaGraphKernelNodeSetParams");
   };
   MF.addMatcher(
       callExpr(callee(functionDecl(functionName()))).bind("FunctionCall"),
       this);
+
+  auto typeName = [&]() { return hasAnyName("cudaKernelNodeParams"); };
+  MF.addMatcher(
+      memberExpr(hasObjectExpression(hasType(type(hasUnqualifiedDesugaredType(
+                     recordType(hasDeclaration(recordDecl(typeName()))))))))
+          .bind("Type"),
+      this);
 }
 
 void GraphRule::runRule(const MatchFinder::MatchResult &Result) {
+  if (auto ME = getNodeAsType<MemberExpr>(Result, "Type")) {
+    auto BaseTy = DpctGlobalInfo::getUnqualifiedTypeName(
+        ME->getBase()->getType().getDesugaredType(*Result.Context),
+        *Result.Context);
+    auto MemberName = ME->getMemberNameInfo().getAsString();
+    if (BaseTy == "cudaKernelNodeParams") {
+
+      auto FieldName = KernelNodeParamNames[MemberName];
+      if (FieldName.empty()) {
+        report(ME->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
+               DpctGlobalInfo::getOriginalTypeName(ME->getBase()->getType()) +
+                   "::" + ME->getMemberDecl()->getName().str());
+        return;
+      }
+      requestFeature(HelperFeatureEnum::device_ext);
+      if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) {
+        StringRef ReplacedArg = "";
+        emplaceTransformation(
+            ReplaceMemberAssignAsSetMethod(BO, ME, FieldName, ReplacedArg));
+      } else {
+        emplaceTransformation(new RenameFieldInMemberExpr(
+            ME, buildString("get_", FieldName, "()")));
+      }
+    }
+    return;
+  }
   const CallExpr *CE = getNodeAsType<CallExpr>(Result, "FunctionCall");
   if (!CE) {
     return;
@@ -51,5 +85,47 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) {
   EA.applyAllSubExprRepl();
 }
 
+const Expr *GraphRule::getParentAsAssignedBO(const Expr *E,
+                                             ASTContext &Context) {
+  auto Parents = Context.getParents(*E);
+  if (Parents.size() > 0)
+    return getAssignedBO(Parents[0].get<Expr>(), Context);
+  return nullptr;
+}
+
+// Return the binary operator if E is the lhs of an assign expression, otherwise
+// nullptr.
+const Expr *GraphRule::getAssignedBO(const Expr *E, ASTContext &Context) {
+  if (dyn_cast<MemberExpr>(E)) {
+    // Continue finding parents when E is MemberExpr.
+    return getParentAsAssignedBO(E, Context);
+  } else if (auto ICE = dyn_cast<ImplicitCastExpr>(E)) {
+    // Stop finding parents and return nullptr when E is ImplicitCastExpr,
+    // except for ArrayToPointerDecay cast.
+    if (ICE->getCastKind() == CK_ArrayToPointerDecay) {
+      return getParentAsAssignedBO(E, Context);
+    }
+  } else if (auto ASE = dyn_cast<ArraySubscriptExpr>(E)) {
+    // Continue finding parents when E is ArraySubscriptExpr, and remove
+    // subscript operator anyway for texture object's member.
+    emplaceTransformation(new ReplaceToken(
+        Lexer::getLocForEndOfToken(ASE->getLHS()->getEndLoc(), 0,
+                                   Context.getSourceManager(),
+                                   Context.getLangOpts()),
+        ASE->getRBracketLoc(), ""));
+    return getParentAsAssignedBO(E, Context);
+  } else if (auto BO = dyn_cast<BinaryOperator>(E)) {
+    // If E is BinaryOperator, return E only when it is assign expression,
+    // otherwise return nullptr.
+    if (BO->getOpcode() == BO_Assign)
+      return BO;
+  } else if (auto COCE = dyn_cast<CXXOperatorCallExpr>(E)) {
+    if (COCE->getOperator() == OO_Equal) {
+      return COCE;
+    }
+  }
+  return nullptr;
+}
+
 } // namespace dpct
 } // namespace clang

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -456,12 +456,12 @@ ENTRY(cudaGraphInstantiateWithFlags, cudaGraphInstantiateWithFlags, false, NO_FL
 ENTRY(cudaGraphInstantiateWithParams, cudaGraphInstantiateWithParams, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphKernelNodeCopyAttributes, cudaGraphKernelNodeCopyAttributes, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphKernelNodeGetAttribute, cudaGraphKernelNodeGetAttribute, false, NO_FLAG, P4, "comment")
-ENTRY(cudaGraphKernelNodeGetParams, cudaGraphKernelNodeGetParams, false, NO_FLAG, P4, "comment")
+ENTRY(cudaGraphKernelNodeGetParams, cudaGraphKernelNodeGetParams, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cudaGraphKernelNodeSetAttribute, cudaGraphKernelNodeSetAttribute, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphKernelNodeSetEnabled, cudaGraphKernelNodeSetEnabled, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphKernelNodeSetGridDim, cudaGraphKernelNodeSetGridDim, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphKernelNodeSetParam, cudaGraphKernelNodeSetParam, false, NO_FLAG, P4, "comment")
-ENTRY(cudaGraphKernelNodeSetParams, cudaGraphKernelNodeSetParams, false, NO_FLAG, P4, "comment")
+ENTRY(cudaGraphKernelNodeSetParams, cudaGraphKernelNodeSetParams, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cudaGraphKernelNodeUpdatesApply, cudaGraphKernelNodeUpdatesApply, false, NO_FLAG, P4, "comment")
 ENTRY(cudaGraphLaunch, cudaGraphLaunch, true, NO_FLAG, P4, "Successful/DPCT1119")
 ENTRY(cudaGraphMemAllocNodeGetParams, cudaGraphMemAllocNodeGetParams, false, NO_FLAG, P4, "comment")

clang/lib/DPCT/SrcAPI/TypeNames.inc

@@ -41,7 +41,7 @@ ENTRY_TYPE(CUgraphNode, false, NO_FLAG, P4, "comment")
 ENTRY_TYPE(CUgraphicsResource, true, NO_FLAG, P4, "successful")
 
 // CUDA Runtime Library
-ENTRY_TYPE(cudaKernelNodeParams, false, NO_FLAG, P4, "comment")
+ENTRY_TYPE(cudaKernelNodeParams, true, NO_FLAG, P4, "Successful/DPCT1119")
 
 // cuDNN Library
 ENTRY_TYPE(cudnnReduceTensorIndices_t, false, NO_FLAG, P4, "comment")

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

@@ -8,9 +8,12 @@
 
 #pragma once
 
+#include "dpct/util.hpp"
+#include "sycl/handler.hpp"
 #include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/sycl.hpp>
 #include <unordered_map>
+#include <unordered_map>
 
 namespace dpct {
 namespace experimental {
@@ -25,6 +28,28 @@ typedef sycl::ext::oneapi::experimental::command_graph<
 
 typedef sycl::ext::oneapi::experimental::node *node_ptr;
 
+struct kernel_node_params {
+  dpct::dim3 block_dim;
+  dpct::dim3 grid_dim;
+  void *kernel_params;
+  void* func;
+  unsigned int shared_mem_bytes;
+
+public:
+  void set_block_dim(dpct::dim3 block_dim) { block_dim = block_dim; }
+  void set_grid_dim(dpct::dim3 grid_dim) { grid_dim = grid_dim; }
+  void set_kernel_params(void *kernel_params) { kernel_params = kernel_params; }
+  void set_func(void *func) { func = func; }
+  void set_shared_mem_bytes(unsigned int shared_mem_bytes) {
+    shared_mem_bytes = shared_mem_bytes;
+  }
+  dpct::dim3 get_block_dim() { return block_dim; }
+  dpct::dim3 get_grid_dim() { return grid_dim; }
+  void *get_kernel_params() { return kernel_params; }
+  void *get_func() { return func; }
+  unsigned int get_shared_mem_bytes() { return shared_mem_bytes; }
+};
+
 namespace detail {
 class graph_mgr {
 public:
@@ -39,6 +64,10 @@ class graph_mgr {
     return instance;
   }
 
+  std::unordered_map<dpct::experimental::node_ptr,
+                     dpct::experimental::kernel_node_params>
+      kernel_node_params_map;
+
   void begin_recording(sycl::queue *queue_ptr) {
     // Calling begin_recording on an already recording queue is a no-op in SYCL
     if (queue_graph_map.find(queue_ptr) != queue_graph_map.end()) {
@@ -94,6 +123,18 @@ class graph_mgr {
     }
   }
 
+  void kernel_node_set_params(
+      dpct::experimental::node_ptr node,
+      dpct::experimental::kernel_node_params *kernel_node_params) {
+    kernel_node_params_map[node] = kernel_node_params;
+  }
+
+  void get_kernel_node_get_params(
+      dpct::experimental::node_ptr node,
+      dpct::experimental::kernel_node_params *kernel_node_params) {
+    kernel_node_params = kernel_node_params_map[node];
+  }
+
 private:
   std::unordered_map<sycl::queue *, command_graph_ptr> queue_graph_map;
   std::unordered_map<dpct::experimental::command_graph_ptr,
@@ -174,9 +215,9 @@ static void add_dependencies(dpct::experimental::command_graph_ptr graph,
 /// nodes will be assigned.
 /// \param [out] numberOfNodes The number of nodes in the graph.
 static void get_nodes(dpct::experimental::command_graph_ptr graph,
-                      dpct::experimental::node_ptr *nodesArray,
-                      std::size_t *numberOfNodes) {
-  detail::graph_mgr::instance().get_nodes(graph, nodesArray, numberOfNodes);
+  dpct::experimental::node_ptr *nodesArray,
+  std::size_t *numberOfNodes) {
+detail::graph_mgr::instance().get_nodes(graph, nodesArray, numberOfNodes);
 }
 
 /// Gets the root nodes in the command graph.
@@ -185,10 +226,27 @@ static void get_nodes(dpct::experimental::command_graph_ptr graph,
 /// root nodes will be assigned.
 /// \param [out] numberOfNodes The number of root nodes in the graph.
 static void get_root_nodes(dpct::experimental::command_graph_ptr graph,
-                           dpct::experimental::node_ptr *nodesArray,
-                           std::size_t *numberOfNodes) {
-  detail::graph_mgr::instance().get_root_nodes(graph, nodesArray,
-                                               numberOfNodes);
+       dpct::experimental::node_ptr *nodesArray,
+       std::size_t *numberOfNodes) {
+detail::graph_mgr::instance().get_root_nodes(graph, nodesArray,
+                           numberOfNodes);
+}
+
+static void
+kernel_node_set_params(dpct::experimental::node_ptr node,
+                       dpct::experimental::kernel_node_params *params) {
+  detail::graph_mgr::instance().kernel_node_set_params(node, params);
+}
+
+static void
+kernel_node_get_params(dpct::experimental::node_ptr node,
+                       dpct::experimental::kernel_node_params *params) {
+  detail::graph_mgr::instance().kernel_node_set_params(node, params);
+}
+
+
+static void add_kernel_node(dpct::experimental::node_ptr node, dpct::experimental::command_graph_ptr graph, dpct::experimental::node_ptr *dependencies, std::size_t numberOfDependencies, dpct::experimental::kernel_node_params &kernelNodeParams){
+  
 }
 
 } // namespace experimental