Added helper function for nvshmemx_init_attr

Author: TejaX-Alaghari

clang/lib/DPCT/RulesSHMEM/APINamesNvshmem.inc

@@ -203,9 +203,13 @@
 FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                         CALL_FACTORY_ENTRY("nvshmem_init", CALL("ishmem_init")))
 
-FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
-                        CALL_FACTORY_ENTRY("nvshmemx_init_attr",
-                                           CALL("ishmemx_init_attr", ARG(1))))
+HEADER_INSERT_FACTORY(HeaderType::HT_DPCT_SHMEM_Utils,
+                      FEATURE_REQUEST_FACTORY(
+                          HelperFeatureEnum::device_ext,
+                          CALL_FACTORY_ENTRY("nvshmemx_init_attr",
+                                             CALL(MapNames::getDpctNamespace() +
+                                                      "shmem::init_attr",
+                                                  ARG(0), ARG(1)))))
 
 FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                         CALL_FACTORY_ENTRY("nvshmem_my_pe",

clang/lib/DPCT/RulesSHMEM/NVSHMEMAPIMigration.cpp

@@ -216,91 +216,6 @@ void clang::dpct::NVSHMEMRule::runRule(
 
     EA.analyze(*TL);
   } else if (const CallExpr *CE = getNodeAsType<CallExpr>(Result, "call")) {
-    auto &SM = DpctGlobalInfo::getSourceManager();
-    auto LO = LangOptions();
-
-    std::string FuncName = "";
-    const FunctionDecl *FD = CE->getDirectCallee();
-    if (FD) {
-      FuncName = FD->getNameInfo().getName().getAsString();
-    }
-
-    if (!FuncName.empty()) {
-      if (FuncName == "nvshmemx_init_attr") {
-        // Get function arguments
-        std::string NvshmemRT = "";
-        std::string NvshmemInitRT = "";
-        std::string AttrArg = "";
-
-        // Get the first argument's data
-        const Expr *Arg0 = CE->getArg(0);
-
-        // Binary op on first argument is not supported
-        if (dyn_cast<BinaryOperator>(Arg0->IgnoreImpCasts())) {
-          report(CE->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
-                 FuncName);
-          return;
-        }
-
-        // Get first argument's init value
-        if (auto DRE = dyn_cast<DeclRefExpr>(Arg0->IgnoreImpCasts())) {
-          if (const VarDecl *VD = dyn_cast<VarDecl>(DRE->getDecl())) {
-            if (VD->hasInit()) {
-              // get the init value from definition
-              if (auto Init =
-                      dyn_cast<DeclRefExpr>(VD->getInit()->IgnoreImplicit())) {
-                NvshmemInitRT = Init->getNameInfo().getName().getAsString();
-              }
-            }
-          }
-        }
-
-        // Get the first argument as string
-        NvshmemRT = Lexer::getSourceText(
-            CharSourceRange::getTokenRange(Arg0->getSourceRange()), SM, LO);
-
-        if (NvshmemRT == "0" || NvshmemInitRT == "0") {
-          emplaceTransformation(new ReplaceStmt(CE, "ishmem_init()"));
-          return;
-        }
-
-        std::string IshmemRT = "";
-        if (NvshmemRT == "NVSHMEMX_INIT_WITH_MPI_COMM") {
-          IshmemRT = "ISHMEMX_RUNTIME_MPI";
-        } else if (NvshmemRT == "NVSHMEMX_INIT_WITH_SHMEM") {
-          IshmemRT = "ISHMEMX_RUNTIME_OPENSHMEM";
-        }
-
-        if (NvshmemInitRT == "NVSHMEMX_INIT_WITH_MPI_COMM" ||
-            NvshmemInitRT == "NVSHMEMX_INIT_WITH_SHMEM") {
-          IshmemRT = "static_cast<ishmemx_runtime_type_t>(" + NvshmemRT + ")";
-        }
-
-        // Get the second argument as string
-        AttrArg = Lexer::getSourceText(
-            CharSourceRange::getTokenRange(CE->getArg(1)->getSourceRange()), SM,
-            LO);
-
-        if (IshmemRT.empty()) {
-          report(CE->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
-                 FuncName);
-          return;
-        } else {
-          auto IndentLoc = CE->getBeginLoc();
-          if (IndentLoc.isMacroID())
-            IndentLoc = SM.getExpansionLoc(IndentLoc);
-
-          std::string SetIshmemRT =
-              "(" + AttrArg + ")->runtime = " + IshmemRT + ";";
-          SetIshmemRT += getNL();
-          SetIshmemRT += getIndent(IndentLoc, SM).str();
-
-          emplaceTransformation(
-              new InsertBeforeStmt(CE, std::move(SetIshmemRT)));
-        }
-      }
-    }
-
     EA.analyze(CE);
   } else if (const DeclRefExpr *DRE =
                  getNodeAsType<DeclRefExpr>(Result, "enumConstant")) {

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

@@ -11,7 +11,20 @@
 
 namespace dpct::shmem {
 
-/// Signal Operations
+/// Setup Operations
+/// Initialize Intel SHMEM runtime based on exisiting communicator in attributes
+/// \param [in] runtime_flags specifies the exisiting communicator
+/// \param [in] attr attributes of existing communicator
+void init_attr(unsigned runtime_flags, ishmemx_attr_t *attr) {
+  if (runtime_flags == 0) {
+    ishmem_init();
+  } else {
+    attr->runtime = static_cast<ishmemx_runtime_type_t>(runtime_flags);
+    ishmemx_init_attr(attr);
+  }
+}
+
+/// Signal Operationss
 /// Update signal address with signal using signal operation
 /// \param [out] sig_addr Address of the signal to be updated
 /// \param [in] signal Value of the signal to update

clang/test/dpct/nvshmem/library_setup_exit_query.cu

@@ -57,30 +57,21 @@ int main() {
   attr.args = args;
 #endif // NO_BUILD_TEST
 
-  // CHECK: (&attr)->runtime = ISHMEMX_RUNTIME_MPI;
-  // CHECK-NEXT: ishmemx_init_attr(&attr);
+  // CHECK: dpct::shmem::init_attr(ISHMEMX_RUNTIME_MPI, &attr);
+  // CHECK-NEXT: dpct::shmem::init_attr(ISHMEMX_RUNTIME_OPENSHMEM, &attr);
+  // CHECK-NEXT: dpct::shmem::init_attr(ISHMEMX_RUNTIME_MPI | ISHMEMX_RUNTIME_OPENSHMEM, &attr);
   nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM, &attr);
-
-#ifndef NO_BUILD_TEST
-  // CHECK: /*
-  // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of nvshmemx_init_attr is not supported.
-  // CHECK-NEXT: */
+  nvshmemx_init_attr(NVSHMEMX_INIT_WITH_SHMEM, &attr);
   nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM | NVSHMEMX_INIT_WITH_SHMEM, &attr);
-#endif // NO_BUILD_TEST
 
   // CHECK: unsigned int rt = ISHMEMX_RUNTIME_MPI;
-  // CHECK-NEXT: (&attr)->runtime = static_cast<ishmemx_runtime_type_t>(rt);
-  // CHECK-NEXT: ishmemx_init_attr(&attr);
+  // CHECK-NEXT: rt = ISHMEMX_RUNTIME_OPENSHMEM;
+  // CHECK-NEXT: dpct::shmem::init_attr(rt, &attr);
   unsigned int rt = NVSHMEMX_INIT_WITH_MPI_COMM;
-  nvshmemx_init_attr(rt, &attr);
-
-  // CHECK: rt = ISHMEMX_RUNTIME_OPENSHMEM;
-  // CHECK-NEXT: (&attr)->runtime = static_cast<ishmemx_runtime_type_t>(rt);
-  // CHECK-NEXT: ishmemx_init_attr(&attr);
   rt = NVSHMEMX_INIT_WITH_SHMEM;
   nvshmemx_init_attr(rt, &attr);
 
-  // CHECK: ishmem_init();
+  // CHECK: dpct::shmem::init_attr(0, &attr);
   nvshmemx_init_attr(0, &attr);
 
   // CHECK: ishmem_init();