Added new logic for init attr flags

Author: TejaX-Alaghari

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -1671,9 +1671,11 @@ void MapNames::setExplicitNamespaceMap(
       {"NVSHMEM_TEAM_INVALID",
        std::make_shared<EnumNameRule>("ISHMEM_TEAM_INVALID")},
       {"NVSHMEMX_INIT_WITH_MPI_COMM",
-       std::make_shared<EnumNameRule>("ISHMEMX_RUNTIME_MPI")},
+       std::make_shared<EnumNameRule>(MapNames::getDpctNamespace() +
+                                      "shmemx::RUNTIME_MPI")},
       {"NVSHMEMX_INIT_WITH_SHMEM",
-       std::make_shared<EnumNameRule>("ISHMEMX_RUNTIME_OPENSHMEM")},
+       std::make_shared<EnumNameRule>(MapNames::getDpctNamespace() +
+                                      "shmemx::RUNTIME_OPENSHMEM")},
       {"NVSHMEM_SIGNAL_SET",
        std::make_shared<EnumNameRule>("ISHMEM_SIGNAL_SET")},
       {"NVSHMEM_SIGNAL_ADD",

clang/lib/DPCT/RulesSHMEM/APINamesNvshmem.inc

@@ -203,18 +203,13 @@
 FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                         CALL_FACTORY_ENTRY("nvshmem_init", CALL("ishmem_init")))
 
-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",
-                 makeCombinedArg(
-                     makeCombinedArg(
-                         ARG("static_cast<ishmemx_runtime_type_t>("), ARG(0)),
-                     ARG(")")),
-                 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() +
+                                                      "shmemx::init_attr",
+                                                  ARG(0), ARG(1)))))
 
 FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                         CALL_FACTORY_ENTRY("nvshmem_my_pe",
@@ -513,7 +508,7 @@ HEADER_INSERT_FACTORY(HeaderType::HT_DPCT_SHMEM_Utils,
                           HelperFeatureEnum::device_ext,
                           CALL_FACTORY_ENTRY("nvshmemx_signal_op",
                                              CALL(MapNames::getDpctNamespace() +
-                                                      "shmem::signal_op",
+                                                      "shmemx::signal_op",
                                                   ARG(0), ARG(1), ARG(2),
                                                   ARG(3)))))
 

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

@@ -9,18 +9,27 @@
 #ifndef __DPCT_SHMEM_UTILS_HPP__
 #define __DPCT_SHMEM_UTILS_HPP__
 
-namespace dpct::shmem {
+namespace dpct::shmemx {
+
+enum flags { RUNTIME_MPI = 1 << 1, RUNTIME_OPENSHMEM = 1 << 2 };
 
 /// Initialize Intel SHMEM library based on exisiting communicator
 /// \param [in] runtime_flags specify the underlying communicator like MPI/
 /// OpenSHMEM to initialize iSHMEM for launcher agnostic bootstrapping
 /// \param [in] attr Additional attributes for initializing the iSHMEM library.
-void init_attr(ishmemx_runtime_type_t runtime_flags, ishmemx_attr_t *attr) {
+void init_attr(unsigned runtime_flags, ishmemx_attr_t *attr) {
   if (runtime_flags == 0) {
     // if no runtime flags are present, initialize iSHMEM normally
     ishmem_init();
   } else {
-    attr->runtime = runtime_flags;
+    unsigned ishmem_runtime_flags = 0;
+
+    if (runtime_flags & RUNTIME_MPI)
+      ishmem_runtime_flags |= ISHMEMX_RUNTIME_MPI;
+    if (runtime_flags & RUNTIME_OPENSHMEM)
+      ishmem_runtime_flags |= ISHMEMX_RUNTIME_OPENSHMEM;
+
+    attr->runtime = static_cast<ishmemx_runtime_type_t>(ishmem_runtime_flags);
     ishmemx_init_attr(attr);
   }
 }
@@ -42,6 +51,6 @@ void signal_op(uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) {
   }
 }
 
-} // namespace dpct::shmem
+} // namespace dpct::shmemx
 
 #endif // __DPCT_SHMEM_UTILS_HPP__

clang/test/dpct/nvshmem/library_setup_exit_query.cu

@@ -60,21 +60,21 @@ int main() {
   attr.args = args;
 #endif // NO_BUILD_TEST
 
-  // CHECK: dpct::shmem::init_attr(static_cast<ishmemx_runtime_type_t>(ISHMEMX_RUNTIME_MPI), &attr);
-  // CHECK-NEXT: dpct::shmem::init_attr(static_cast<ishmemx_runtime_type_t>(ISHMEMX_RUNTIME_OPENSHMEM), &attr);
-  // CHECK-NEXT: dpct::shmem::init_attr(static_cast<ishmemx_runtime_type_t>(ISHMEMX_RUNTIME_MPI | ISHMEMX_RUNTIME_OPENSHMEM), &attr);
+  // CHECK: dpct::shmemx::init_attr(dpct::shmemx::RUNTIME_MPI, &attr);
+  // CHECK-NEXT: dpct::shmemx::init_attr(dpct::shmemx::RUNTIME_OPENSHMEM, &attr);
+  // CHECK-NEXT: dpct::shmemx::init_attr(dpct::shmemx::RUNTIME_MPI | dpct::shmemx::RUNTIME_OPENSHMEM, &attr);
   nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM, &attr);
   nvshmemx_init_attr(NVSHMEMX_INIT_WITH_SHMEM, &attr);
   nvshmemx_init_attr(NVSHMEMX_INIT_WITH_MPI_COMM | NVSHMEMX_INIT_WITH_SHMEM, &attr);
 
-  // CHECK: unsigned int rt = ISHMEMX_RUNTIME_MPI;
-  // CHECK-NEXT: rt = ISHMEMX_RUNTIME_OPENSHMEM;
-  // CHECK-NEXT: dpct::shmem::init_attr(static_cast<ishmemx_runtime_type_t>(rt), &attr);
+  // CHECK: unsigned int rt = dpct::shmemx::RUNTIME_MPI;
+  // CHECK-NEXT: rt = dpct::shmemx::RUNTIME_OPENSHMEM;
+  // CHECK-NEXT: dpct::shmemx::init_attr(rt, &attr);
   unsigned int rt = NVSHMEMX_INIT_WITH_MPI_COMM;
   rt = NVSHMEMX_INIT_WITH_SHMEM;
   nvshmemx_init_attr(rt, &attr);
 
-  // CHECK: dpct::shmem::init_attr(static_cast<ishmemx_runtime_type_t>(0), &attr);
+  // CHECK: dpct::shmemx::init_attr(0, &attr);
   nvshmemx_init_attr(0, &attr);
 
   // CHECK: ishmem_init();

clang/test/dpct/nvshmem/sig_ops.cu

@@ -48,9 +48,9 @@ __host__ __device__ void test() {
 
   // nvshmemx_signal_op
   // ishmemx_signal_set/add
-  // CHECK: dpct::shmem::signal_op(&sig_addr, signal, sig_op, pe);
-  // CHECK-NEXT: dpct::shmem::signal_op(&sig_addr, signal, ISHMEM_SIGNAL_SET, pe);
-  // CHECK-NEXT: dpct::shmem::signal_op(&sig_addr, signal, ISHMEM_SIGNAL_ADD, pe);
+  // CHECK: dpct::shmemx::signal_op(&sig_addr, signal, sig_op, pe);
+  // CHECK-NEXT: dpct::shmemx::signal_op(&sig_addr, signal, ISHMEM_SIGNAL_SET, pe);
+  // CHECK-NEXT: dpct::shmemx::signal_op(&sig_addr, signal, ISHMEM_SIGNAL_ADD, pe);
   nvshmemx_signal_op(&sig_addr, signal, sig_op, pe);
   nvshmemx_signal_op(&sig_addr, signal, NVSHMEM_SIGNAL_SET, pe);
   nvshmemx_signal_op(&sig_addr, signal, NVSHMEM_SIGNAL_ADD, pe);