Merge pull request #18 from avinashkethineedi/ipc_backend

IPC backend

[ROCm/rocshmem commit: 73365cfefd]

projects/rocshmem/CMakeLists.txt

@@ -54,7 +54,8 @@ string(REGEX REPLACE "([0-9]+)\.([0-9]+)\.([0-9]+)(.*)" "\\1.\\2.\\3" ROCSHMEM_V
 ###############################################################################
 option(DEBUG "Enable debug trace" OFF)
 option(PROFILE "Enable statistics and timing support" OFF)
-option(USE_GPU_IB "Enable GPU_IB conduit. If off, RO_NET will be used" ON)
+option(USE_GPU_IB "Enable GPU_IB conduit." ON)
+option(USE_RO "Enable RO conduit." ON)
 option(USE_DC "Enable IB dynamically connected transport (DC)" OFF)
 option(USE_IPC "Enable IPC support (using HIP)" OFF)
 option(USE_THREADS "Enable workgroup threads to share network queues" OFF)

projects/rocshmem/cmake/config.h.in

@@ -1,6 +1,7 @@
 #cmakedefine DEBUG
 #cmakedefine PROFILE
 #cmakedefine USE_GPU_IB
+#cmakedefine USE_RO
 #cmakedefine USE_DC
 #cmakedefine USE_IPC
 #cmakedefine USE_THREADS

projects/rocshmem/scripts/build_configs/ipc_single

@@ -0,0 +1,30 @@
+#!/bin/bash
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+
+if [ -z $1 ]
+then
+  install_path=~/rocshmem
+else
+  install_path=$1
+fi
+
+src_path=$(dirname "$(realpath $0)")/../../
+
+cmake \
+    -DCMAKE_BUILD_TYPE=Release \
+    -DCMAKE_INSTALL_PREFIX=$install_path \
+    -DCMAKE_VERBOSE_MAKEFILE=OFF \
+    -DDEBUG=OFF \
+    -DPROFILE=OFF \
+    -DUSE_GPU_IB=OFF \
+    -DUSE_RO=OFF \
+    -DUSE_DC=OFF \
+    -DUSE_IPC=ON \
+    -DUSE_COHERENT_HEAP=ON \
+    -DUSE_THREADS=OFF \
+    -DUSE_WF_COAL=OFF \
+    -DUSE_SINGLE_NODE=ON \
+    -DUSE_HOST_SIDE_HDP_FLUSH=OFF\
+    $src_path
+cmake --build . --parallel 8
+cmake --install .

projects/rocshmem/src/CMakeLists.txt

@@ -37,6 +37,7 @@ target_sources(
     team_tracker.cpp
     util.cpp
     wf_coal_policy.cpp
+    ipc_policy.cpp
 )
 
 target_compile_options(
@@ -59,11 +60,12 @@ target_compile_options(
 ###############################################################################
 IF (USE_GPU_IB)
 add_subdirectory(gpu_ib)
-ELSE()
+ELSEIF(USE_RO)
 add_subdirectory(reverse_offload)
+ELSE()
+add_subdirectory(ipc)
 ENDIF()
 add_subdirectory(containers)
 add_subdirectory(host)
 add_subdirectory(memory)
 add_subdirectory(sync)
-add_subdirectory(ipc)

projects/rocshmem/src/backend_bc.cpp

@@ -25,10 +25,12 @@
 #include "backend_type.hpp"
 #include "context_incl.hpp"
 
-#ifndef USE_GPU_IB
+#ifdef USE_GPU_IB
+#include "gpu_ib/backend_ib.hpp"
+#elif defined(USE_RO)
 #include "reverse_offload/backend_ro.hpp"
 #else
-#include "gpu_ib/backend_ib.hpp"
+#include "ipc/backend_ipc.hpp"
 #endif
 
 namespace rocshmem {
@@ -201,18 +203,22 @@ void Backend::reset_stats() {
 }
 
 __device__ bool Backend::create_ctx(int64_t option, roc_shmem_ctx_t* ctx) {
-#ifndef USE_GPU_IB
+#ifdef USE_GPU_IB
+  return static_cast<GPUIBBackend*>(this)->create_ctx(option, ctx);
+#elif defined(USE_RO)
   return static_cast<ROBackend*>(this)->create_ctx(option, ctx);
 #else
-  return static_cast<GPUIBBackend*>(this)->create_ctx(option, ctx);
+  return static_cast<IPCBackend*>(this)->create_ctx(option, ctx);
 #endif
 }
 
 __device__ void Backend::destroy_ctx(roc_shmem_ctx_t* ctx) {
-#ifndef USE_GPU_IB
+#ifdef USE_GPU_IB
+  static_cast<GPUIBBackend*>(this)->destroy_ctx(ctx);
+#elif defined(USE_RO)
   static_cast<ROBackend*>(this)->destroy_ctx(ctx);
 #else
-  static_cast<GPUIBBackend*>(this)->destroy_ctx(ctx);
+  static_cast<IPCBackend*>(this)->destroy_ctx(ctx);
 #endif
 }
 

projects/rocshmem/src/backend_bc.hpp

@@ -38,7 +38,7 @@
 #include "config.h"  // NOLINT(build/include_subdir)
 #include "roc_shmem/roc_shmem.hpp"
 #include "backend_type.hpp"
-#include "ipc/ipc_policy.hpp"
+#include "ipc_policy.hpp"
 #include "memory/symmetric_heap.hpp"
 #include "stats.hpp"
 #include "team_tracker.hpp"

projects/rocshmem/src/backend_type.hpp

@@ -44,7 +44,7 @@ namespace rocshmem {
  * @note Derived classes which use Backend as a base class must add
  * themselves to this enum class to support static polymorphism.
  */
-enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
+enum class BackendType { RO_BACKEND, GPU_IB_BACKEND, IPC_BACKEND };
 
 /**
  * @brief Helper macro for some dispatch calls
@@ -57,9 +57,12 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
 #ifdef USE_GPU_IB
 #define DISPATCH(Func)                     \
   static_cast<GPUIBContext *>(this)->Func;
-#else
+#elif defined(USE_RO)
 #define DISPATCH(Func)                     \
   static_cast<ROContext *>(this)->Func;
+#else
+#define DISPATCH(Func)                     \
+  static_cast<IPCContext *>(this)->Func;
 #endif
 
 /**
@@ -70,11 +73,16 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
   auto ret_val{0};                                   \
   ret_val = static_cast<GPUIBContext *>(this)->Func; \
   return ret_val;
-#else
+#elif defined(USE_RO)
 #define DISPATCH_RET(Func)                        \
   auto ret_val{0};                                \
   ret_val = static_cast<ROContext *>(this)->Func; \
   return ret_val;
+#else
+#define DISPATCH_RET(Func)                        \
+  auto ret_val{0};                                \
+  ret_val = static_cast<IPCContext *>(this)->Func; \
+  return ret_val;
 #endif
 /**
  * @brief Device static dispatch method call with a return type of pointer.
@@ -84,11 +92,16 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
   void *ret_val{nullptr};                            \
   ret_val = static_cast<GPUIBContext *>(this)->Func; \
   return ret_val;
-#else
+#elif defined(USE_RO)
 #define DISPATCH_RET_PTR(Func)                    \
   void *ret_val{nullptr};                         \
   ret_val = static_cast<ROContext *>(this)->Func; \
   return ret_val;
+#else
+#define DISPATCH_RET_PTR(Func)                    \
+  void *ret_val{nullptr};                         \
+  ret_val = static_cast<IPCContext *>(this)->Func; \
+  return ret_val;
 #endif
 
 /**
@@ -100,8 +113,10 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
  */
 #ifdef USE_GPU_IB
 #define HOST_DISPATCH(Func) static_cast<GPUIBHostContext *>(this)->Func;
-#else
+#elif defined(USE_RO)
 #define HOST_DISPATCH(Func) static_cast<ROHostContext *>(this)->Func;
+#else
+#define HOST_DISPATCH(Func) static_cast<IPCHostContext *>(this)->Func;
 #endif
 /**
  * @brief Host static dispatch method call with return value.
@@ -116,11 +131,16 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
   auto ret_val{0};                                       \
   ret_val = static_cast<GPUIBHostContext *>(this)->Func; \
   return ret_val;
-#else
+#elif defined(USE_RO)
 #define HOST_DISPATCH_RET(Func)                       \
   auto ret_val{0};                                    \
   ret_val = static_cast<ROHostContext *>(this)->Func; \
   return ret_val;
+#else
+#define HOST_DISPATCH_RET(Func)                       \
+  auto ret_val{0};                                    \
+  ret_val = static_cast<IPCHostContext *>(this)->Func; \
+  return ret_val;
 #endif
 
 }  // namespace rocshmem

projects/rocshmem/src/context.hpp

@@ -28,7 +28,7 @@
 #include "backend_type.hpp"
 #include "fence_policy.hpp"
 #include "host/host.hpp"
-#include "ipc/ipc_policy.hpp"
+#include "ipc_policy.hpp"
 #include "stats.hpp"
 #include "sync/spin_ebo_block_mutex.hpp"
 #include "wf_coal_policy.hpp"

projects/rocshmem/src/context_incl.hpp

@@ -29,9 +29,12 @@
 #ifdef USE_GPU_IB
 #include "gpu_ib/context_ib_device.hpp"
 #include "gpu_ib/context_ib_host.hpp"
-#else
+#elif defined (USE_RO)
 #include "reverse_offload/context_ro_device.hpp"
 #include "reverse_offload/context_ro_host.hpp"
+#else
+#include "ipc/context_ipc_device.hpp"
+#include "ipc/context_ipc_host.hpp"
 #endif
 
 #endif  // LIBRARY_SRC_CONTEXT_INCL_HPP_

projects/rocshmem/src/context_tmpl_device.hpp

@@ -27,8 +27,10 @@
 #include "backend_type.hpp"
 #ifdef USE_GPU_IB
 #include "gpu_ib/context_ib_device.hpp"
-#else
+#elif defined(USE_RO)
 #include "reverse_offload/context_ro_device.hpp"
+#else
+#include "ipc/context_ipc_device.hpp"
 #endif
 
 namespace rocshmem {

projects/rocshmem/src/context_tmpl_host.hpp

@@ -27,8 +27,10 @@
 #include "backend_type.hpp"
 #ifdef USE_GPU_IB
 #include "gpu_ib/context_ib_host.hpp"
-#else
+#elif defined(USE_RO)
 #include "reverse_offload/context_ro_host.hpp"
+#else
+#include "ipc/context_ipc_host.hpp"
 #endif
 namespace rocshmem {
 

projects/rocshmem/src/ipc/CMakeLists.txt

@@ -26,6 +26,7 @@
 target_sources(
   ${PROJECT_NAME}
   PRIVATE
-    ipc_policy.cpp
-    context_ipc.cpp
+    context_ipc_device.cpp
+    context_ipc_host.cpp
+    backend_ipc.cpp
 )

projects/rocshmem/src/ipc/backend_ipc.cpp

@@ -0,0 +1,183 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#include "backend_ipc.hpp"
+
+namespace rocshmem {
+
+#define NET_CHECK(cmd)                                       \
+  {                                                          \
+    if (cmd != MPI_SUCCESS) {                                \
+      fprintf(stderr, "Unrecoverable error: MPI Failure\n"); \
+      abort() ;                                              \
+    }                                                        \
+  }
+
+extern roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT;
+
+IPCBackend::IPCBackend(MPI_Comm comm)
+    :  Backend() {
+  type = BackendType::IPC_BACKEND;
+    
+  if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) {
+    std::stringstream sstream(maximum_num_contexts_str);
+    sstream >> maximum_num_contexts_;
+  }
+
+  init_mpi_once(comm);
+    
+  initIPC();
+    
+  auto *bp{ipc_backend_proxy.get()};
+    
+  bp->heap_ptr = &heap;
+
+  /* Initialize the host interface */
+  host_interface =
+      new HostInterface(hdp_proxy_.get(), thread_comm, &heap);
+  //free host interface
+
+  default_host_ctx = std::make_unique<IPCHostContext>(this, 0);
+
+  ROC_SHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx.get();
+
+  init_g_ret(&heap, thread_comm, MAX_NUM_BLOCKS, &bp->g_ret);
+
+  allocate_atomic_region(&bp->atomic_ret, MAX_NUM_BLOCKS);
+
+  default_context_proxy_ = IPCDefaultContextProxyT(this);
+
+  setup_ctxs();
+  
+}
+
+IPCBackend::~IPCBackend() {
+  ipc_net_free_runtime();
+  CHECK_HIP(hipFree(ctx_array));
+}
+
+void IPCBackend::setup_ctxs() {
+  CHECK_HIP(hipMalloc(&ctx_array, sizeof(IPCContext) * maximum_num_contexts_));
+  for (int i = 0; i < maximum_num_contexts_; i++) {
+    new (&ctx_array[i]) IPCContext(this);
+    ctx_free_list.get()->push_back(ctx_array + i);
+  }
+}
+
+__device__ bool IPCBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) {
+  IPCContext *ctx_{nullptr};
+
+  auto pop_result = ctx_free_list.get()->pop_front();
+  if (!pop_result.success) {
+    return false;
+  }
+  ctx_ = pop_result.value;
+
+  ctx->ctx_opaque = ctx_;
+  return true;
+}
+
+__device__ void IPCBackend::destroy_ctx(roc_shmem_ctx_t *ctx) {
+  ctx_free_list.get()->push_back(static_cast<IPCContext *>(ctx->ctx_opaque));
+}
+
+void IPCBackend::init_mpi_once(MPI_Comm comm) {
+  int init_done{};
+  NET_CHECK(MPI_Initialized(&init_done));
+
+  int provided{};
+  if (!init_done) {
+    NET_CHECK(MPI_Init_thread(0, 0, MPI_THREAD_MULTIPLE, &provided));
+    if (provided != MPI_THREAD_MULTIPLE) {
+      std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n";
+    }
+  }
+  if (comm == MPI_COMM_NULL) comm = MPI_COMM_WORLD;
+
+  NET_CHECK(MPI_Comm_dup(comm, &thread_comm));
+  NET_CHECK(MPI_Comm_size(thread_comm, &num_pes));
+  NET_CHECK(MPI_Comm_rank(thread_comm, &my_pe));
+}
+
+void IPCBackend::team_destroy(roc_shmem_team_t team) {
+  assert(false);
+}
+
+void IPCBackend::create_new_team(Team *parent_team,
+                                TeamInfo *team_info_wrt_parent,
+                                TeamInfo *team_info_wrt_world, int num_pes,
+                                int my_pe_in_new_team, MPI_Comm team_comm,
+                                roc_shmem_team_t *new_team) {
+  assert(false);
+}
+
+void IPCBackend::ctx_create(int64_t options, void **ctx) {
+  IPCHostContext *new_ctx{nullptr};
+  new_ctx = new IPCHostContext(this, options);
+  *ctx = new_ctx;
+}
+
+IPCHostContext *get_internal_ipc_net_ctx(Context *ctx) {
+  return reinterpret_cast<IPCHostContext *>(ctx);
+}
+
+void IPCBackend::ctx_destroy(Context *ctx) {
+  IPCHostContext *ro_net_host_ctx{get_internal_ipc_net_ctx(ctx)};
+  delete ro_net_host_ctx;
+}
+
+void IPCBackend::reset_backend_stats() {
+  assert(false);
+}
+
+void IPCBackend::dump_backend_stats() {
+  assert(false);
+}
+
+void IPCBackend::initIPC() {
+  const auto &heap_bases{heap.get_heap_bases()};
+
+  ipcImpl.ipcHostInit(my_pe, heap_bases,
+                      thread_comm);
+}
+
+void IPCBackend::ipc_net_free_runtime() {
+  /*
+   * Validate that a handle was passed that is not a nullptr.
+   */
+  auto *bp{ipc_backend_proxy.get()};
+  assert(bp);
+
+  /*
+   * Free the atomic_ret array.
+   */
+  CHECK_HIP(hipFree(bp->atomic_ret->atomic_base_ptr));
+  
+  // TODO(Avinash) Free g_ret
+}
+
+void IPCBackend::global_exit(int status) {
+  assert(false);
+}
+
+
+}  // namespace rocshmem

projects/rocshmem/src/ipc/backend_ipc.hpp

@@ -0,0 +1,190 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef LIBRARY_SRC_IPC_BACKEND_HPP_
+#define LIBRARY_SRC_IPC_BACKEND_HPP_
+
+#include "../backend_bc.hpp"
+#include "../containers/free_list_impl.hpp"
+#include "../hdp_proxy.hpp"
+#include "../memory/hip_allocator.hpp"
+#include "ipc_backend_proxy.hpp"
+#include "../context_incl.hpp"
+#include "ipc_context_proxy.hpp"
+#include "../ipc_policy.hpp"
+
+namespace rocshmem {
+
+class IPCBackend : public Backend {
+  const unsigned MAX_NUM_BLOCKS{65536};
+
+ public:
+  /**
+   * @copydoc Backend::Backend(unsigned)
+   */
+  explicit IPCBackend(MPI_Comm comm);
+
+  /**
+   * @copydoc Backend::~Backend()
+   */
+  virtual ~IPCBackend();
+
+  __device__ bool create_ctx(int64_t options, roc_shmem_ctx_t *ctx);
+
+  /**
+   * @brief Destroy a `roc_shmem_ctx_t` context and returns it back to the
+   * context free list.
+   */
+  __device__ void destroy_ctx(roc_shmem_ctx_t *ctx);
+
+  /**
+   * @copydoc Backend::ctx_create
+   */
+  void ctx_create(int64_t options, void **ctx) override;
+
+  /**
+   * @copydoc Backend::ctx_destroy
+   */
+  void ctx_destroy(Context *ctx) override;
+
+   /**
+   * @brief initialize MPI.
+   *
+   * IPC relies on MPI just to exchange the IPC_handle information.
+   *
+   * todo: remove the dependency on MPI and make it generic to PMI-X or just
+   * to OpenSHMEM to have support for both CPU and GPU
+   */
+  void init_mpi_once(MPI_Comm comm);
+
+  /**
+   * @brief Helper to initialize IPC interface.
+   */
+  void initIPC();
+
+  /**
+   * @brief Allocation and initialization of backend contexts.
+   */
+  void setup_ctxs();
+
+  /**
+   * @brief Free all resources associated with the backend.
+   *
+   * The memory allocated to the handle param is deallocated during this
+   * method. The handle should be treated as a nullptr after the call.
+   *
+   * The destructor treats this method as a helper function to destroy
+   * this object.
+   *
+   * @todo The method needs to be broken into smaller pieces and most
+   * of these internal resources need to be moved into subclasses using
+   * RAII.
+   */
+  void ipc_net_free_runtime();
+
+  /**
+   * @brief Abort the application.
+   *
+   * @param[in] status Exit code.
+   *
+   * @return void.
+   *
+   * @note This routine terminates the entire application.
+   */
+  void global_exit(int status) override;
+
+  /**
+   * @copydoc Backend::create_new_team
+   */
+  void create_new_team(Team *parent_team, TeamInfo *team_info_wrt_parent,
+                       TeamInfo *team_info_wrt_world, int num_pes,
+                       int my_pe_in_new_team, MPI_Comm team_comm,
+                       roc_shmem_team_t *new_team) override;
+  
+  /**
+   * @copydoc Backend::team_destroy(roc_shmem_team_t)
+   */
+  void team_destroy(roc_shmem_team_t team) override;
+
+  /**
+   * @brief Handle to device memory fields.
+   */
+  IPCBackendProxyT ipc_backend_proxy{};
+
+  /**
+   * @brief The host-facing interface that will be used
+   * by all contexts of the ROBackend
+   */
+  HostInterface *host_interface{nullptr};
+
+ protected:
+   /**
+   * @copydoc Backend::dump_backend_stats()
+   */
+  void dump_backend_stats() override;
+
+  /**
+   * @copydoc Backend::reset_backend_stats()
+   */
+  void reset_backend_stats() override;
+
+  /**
+   * @brief Allocates uncacheable host memory for the hdp policy.
+   *
+   * @note Internal data ownership is managed by the proxy
+   */
+  HdpProxy<HIPHostAllocator> hdp_proxy_{};
+
+  /**
+   * @brief Holds a copy of the default context for host functions
+   */
+  std::unique_ptr<IPCHostContext> default_host_ctx{nullptr};
+
+ private:
+  /**
+   * @brief Proxy for the default context
+   *
+   * @note Internal data ownership is managed by the proxy
+   */
+  IPCDefaultContextProxyT default_context_proxy_;  // init handled in constructor
+
+  /**
+   * @brief An array of @ref ROContexts that backs the context FreeList.
+   */
+  IPCContext *ctx_array{nullptr};
+
+  /**
+   * @brief A free-list containing contexts.
+   */
+  FreeListProxy<HIPAllocator, IPCContext *> ctx_free_list{};
+
+  /**
+   * @brief Holds maximum number of contexts used in library
+   */
+  size_t maximum_num_contexts_{1024};
+
+
+};
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_BACKEND_HPP_

projects/rocshmem/src/ipc/context_ipc_device.cpp

@@ -20,7 +20,7 @@
  * IN THE SOFTWARE.
  *****************************************************************************/
 
-#include "context_ipc.hpp"
+#include "context_ipc_device.hpp"
 #include "context_ipc_tmpl_device.hpp"
 
 #include <hip/hip_runtime.h>

projects/rocshmem/src/ipc/context_ipc_host.cpp

@@ -0,0 +1,85 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#include "context_ipc_host.hpp"
+
+#include <mpi.h>
+
+#include "config.h"  // NOLINT(build/include_subdir)
+#include "../backend_type.hpp"
+#include "../context_incl.hpp"
+#include "backend_ipc.hpp"
+#include "../host/host.hpp"
+
+namespace rocshmem {
+
+__host__ IPCHostContext::IPCHostContext(Backend *backend,
+                                            [[maybe_unused]] int64_t options)
+    : Context(backend, true) {
+  IPCBackend *b{static_cast<IPCBackend *>(backend)};
+
+  host_interface = b->host_interface;
+
+  context_window_info = host_interface->acquire_window_context();
+}
+
+__host__ IPCHostContext::~IPCHostContext() {
+  host_interface->release_window_context(context_window_info);
+}
+
+__host__ void IPCHostContext::putmem_nbi(void *dest, const void *source,
+                                           size_t nelems, int pe) {
+  host_interface->putmem_nbi(dest, source, nelems, pe, context_window_info);
+}
+
+__host__ void IPCHostContext::getmem_nbi(void *dest, const void *source,
+                                           size_t nelems, int pe) {
+  host_interface->getmem_nbi(dest, source, nelems, pe, context_window_info);
+}
+
+__host__ void IPCHostContext::putmem(void *dest, const void *source,
+                                       size_t nelems, int pe) {
+  host_interface->putmem(dest, source, nelems, pe, context_window_info);
+}
+
+__host__ void IPCHostContext::getmem(void *dest, const void *source,
+                                       size_t nelems, int pe) {
+  host_interface->getmem(dest, source, nelems, pe, context_window_info);
+}
+
+__host__ void IPCHostContext::fence() {
+  host_interface->fence(context_window_info);
+}
+
+__host__ void IPCHostContext::quiet() {
+  host_interface->quiet(context_window_info);
+}
+
+__host__ void IPCHostContext::sync_all() {
+  host_interface->sync_all(context_window_info);
+}
+
+__host__ void IPCHostContext::barrier_all() {
+  host_interface->barrier_all(context_window_info);
+}
+
+}  // namespace rocshmem

projects/rocshmem/src/ipc/context_ipc_host.hpp

@@ -0,0 +1,149 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef LIBRARY_SRC_IPC_CONTEXT_HOST_HPP_
+#define LIBRARY_SRC_IPC_CONTEXT_HOST_HPP_
+
+#include "../context.hpp"
+
+namespace rocshmem {
+
+class IPCHostContext : public Context {
+ public:
+  __host__ IPCHostContext(Backend *b, int64_t options);
+
+  __host__ ~IPCHostContext();
+
+  template <typename T>
+  __host__ void p(T *dest, T value, int pe);
+
+  template <typename T>
+  __host__ T g(const T *source, int pe);
+
+  template <typename T>
+  __host__ void put(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __host__ void get(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __host__ void put_nbi(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __host__ void get_nbi(T *dest, const T *source, size_t nelems, int pe);
+
+  __host__ void putmem(void *dest, const void *source, size_t nelems, int pe);
+
+  __host__ void getmem(void *dest, const void *source, size_t nelems, int pe);
+
+  __host__ void putmem_nbi(void *dest, const void *source, size_t nelems,
+                           int pe);
+
+  __host__ void getmem_nbi(void *dest, const void *source, size_t size, int pe);
+
+  template <typename T>
+  __host__ void amo_add(void *dst, T value, int pe);
+
+  template <typename T>
+  __host__ void amo_cas(void *dst, T value, T cond, int pe);
+
+  template <typename T>
+  __host__ T amo_fetch_add(void *dst, T value, int pe);
+
+  template <typename T>
+  __host__ T amo_fetch_cas(void *dst, T value, T cond, int pe);
+
+  __host__ void fence();
+
+  __host__ void quiet();
+
+  __host__ void barrier_all();
+
+  __host__ void sync_all();
+
+  template <typename T>
+  __host__ void broadcast(T *dest, const T *source, int nelems, int pe_root,
+                          int pe_start, int log_pe_stride, int pe_size,
+                          long *p_sync);
+
+  template <typename T>
+  __host__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
+                          int nelems, int pe_root);
+
+  template <typename T, ROC_SHMEM_OP Op>
+  __host__ void to_all(T *dest, const T *source, int nreduce, int pe_start,
+                       int log_pe_stride, int pe_size, T *p_wrk,
+                       long *p_sync);
+
+  template <typename T, ROC_SHMEM_OP Op>
+  __host__ void to_all(roc_shmem_team_t team, T *dest, const T *source,
+                       int nreduce);
+
+  template <typename T>
+  __host__ void wait_until(T *ptr, roc_shmem_cmps cmp, T val);
+
+  template <typename T>
+  __host__ size_t wait_until_any(T* ptr, size_t nelems,
+                                 const int *status,
+                                 roc_shmem_cmps cmp, T val);
+
+  template <typename T>
+  __host__ void wait_until_all(T* ptr, size_t nelems,
+                               const int *status,
+                               roc_shmem_cmps cmp, T val);
+
+  template <typename T>
+  __host__ size_t wait_until_some(T* ptr, size_t nelems,
+                                size_t* indices,
+                                const int *status,
+                                roc_shmem_cmps cmp, T val);
+
+  template <typename T>
+  __host__ void wait_until_all_vector(T* ptr, size_t nelems,
+                                      const int *status,
+                                      roc_shmem_cmps cmp, T* vals);
+
+  template <typename T>
+  __host__ size_t wait_until_any_vector(T* ptr, size_t nelems,
+                                        const int *status,
+                                        roc_shmem_cmps cmp, T* vals);
+
+  template <typename T>
+  __host__ size_t wait_until_some_vector(T* ptr, size_t nelems,
+                                         size_t* indices,
+                                         const int *status,
+                                         roc_shmem_cmps cmp, T* vals);
+
+  template <typename T>
+  __host__ int test(T *ptr, roc_shmem_cmps cmp, T val);
+
+ public:
+  /* Pointer to the backend's host interface */
+  HostInterface *host_interface{nullptr};
+
+  /* An MPI Window implements a context */
+  WindowInfo *context_window_info{nullptr};
+};
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_CONTEXT_HOST_HPP_

projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp

@@ -25,7 +25,7 @@
 
 #include "config.h"  // NOLINT(build/include_subdir)
 #include "roc_shmem/roc_shmem.hpp"
-#include "context_ipc.hpp"
+#include "context_ipc_device.hpp"
 #include "../util.hpp"
 
 namespace rocshmem {

projects/rocshmem/src/ipc/context_ipc_tmpl_host.hpp

@@ -0,0 +1,173 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_
+#define LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_
+
+#include "config.h"  // NOLINT(build/include_subdir)
+#include "../host/host_templates.hpp"
+
+namespace rocshmem {
+
+template <typename T>
+__host__ void IPCHostContext::p(T *dest, T value, int pe) {
+  host_interface->p<T>(dest, value, pe, context_window_info);
+}
+
+template <typename T>
+__host__ T IPCHostContext::g(const T *source, int pe) {
+  return host_interface->g<T>(source, pe, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::put(T *dest, const T *source, size_t nelems,
+                                    int pe) {
+  host_interface->put<T>(dest, source, nelems, pe, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::get(T *dest, const T *source, size_t nelems,
+                                    int pe) {
+  host_interface->get<T>(dest, source, nelems, pe, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::put_nbi(T *dest, const T *source, size_t nelems,
+                                        int pe) {
+  host_interface->put_nbi<T>(dest, source, nelems, pe, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::get_nbi(T *dest, const T *source, size_t nelems,
+                                        int pe) {
+  host_interface->get_nbi<T>(dest, source, nelems, pe, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::amo_add(void *dst, T value, int pe) {
+  host_interface->amo_add(dst, value, pe, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::amo_cas(void *dst, T value, T cond, int pe) {
+  host_interface->amo_cas(dst, value, cond, pe, context_window_info);
+}
+
+template <typename T>
+__host__ T IPCHostContext::amo_fetch_add(void *dst, T value, int pe) {
+  return host_interface->amo_fetch_add(dst, value, pe, context_window_info);
+}
+
+template <typename T>
+__host__ T IPCHostContext::amo_fetch_cas(void *dst, T value, T cond, int pe) {
+  return host_interface->amo_fetch_cas(dst, value, cond, pe,
+                                       context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::broadcast(
+    T *dest, const T *source, int nelems, int pe_root, int pe_start,
+    int log_pe_stride, int pe_size,
+    long *p_sync) {  // NOLINT(runtime/int)
+  host_interface->broadcast<T>(dest, source, nelems, pe_root, pe_start,
+                               log_pe_stride, pe_size, p_sync);
+}
+
+template <typename T>
+__host__ void IPCHostContext::broadcast(roc_shmem_team_t team, T *dest,
+                                          const T *source, int nelems,
+                                          int pe_root) {
+  host_interface->broadcast<T>(team, dest, source, nelems, pe_root);
+}
+
+template <typename T, ROC_SHMEM_OP Op>
+__host__ void IPCHostContext::to_all(T *dest, const T *source, int nreduce,
+                                       int pe_start, int log_pe_stride,
+                                       int pe_size, T *p_wrk,
+                                       long *p_sync) {  // NOLINT(runtime/int)
+  host_interface->to_all<T, Op>(dest, source, nreduce, pe_start, log_pe_stride,
+                                pe_size, p_wrk, p_sync);
+}
+
+template <typename T, ROC_SHMEM_OP Op>
+__host__ void IPCHostContext::to_all(roc_shmem_team_t team, T *dest,
+                                       const T *source, int nreduce) {
+  host_interface->to_all<T, Op>(team, dest, source, nreduce);
+}
+
+template <typename T>
+__host__ void IPCHostContext::wait_until(T *ptr, roc_shmem_cmps cmp, T val) {
+  host_interface->wait_until<T>(ptr, cmp, val, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::wait_until_all(T *ptr, size_t nelems,
+                                               const int* status,
+                                               roc_shmem_cmps cmp, T val) {
+  host_interface->wait_until_all<T>(ptr, nelems, status, cmp, val, context_window_info);
+}
+
+template <typename T>
+__host__ size_t IPCHostContext::wait_until_any(T *ptr, size_t nelems,
+                                               const int* status,
+                                               roc_shmem_cmps cmp, T val) {
+  return host_interface->wait_until_any<T>(ptr, nelems, status, cmp, val, context_window_info);
+}
+
+template <typename T>
+__host__ size_t IPCHostContext::wait_until_some(T *ptr, size_t nelems,
+                                                 size_t* indices,
+                                                 const int* status,
+                                                 roc_shmem_cmps cmp, T val) {
+  return host_interface->wait_until_some<T>(ptr, nelems, indices, status, cmp, val, context_window_info);
+}
+
+template <typename T>
+__host__ void IPCHostContext::wait_until_all_vector(T *ptr, size_t nelems,
+                                                      const int* status,
+                                                      roc_shmem_cmps cmp, T* vals) {
+  host_interface->wait_until_all_vector<T>(ptr, nelems, status, cmp, vals, context_window_info);
+}
+
+template <typename T>
+__host__ size_t IPCHostContext::wait_until_any_vector(T *ptr, size_t nelems,
+                                                        const int* status,
+                                                        roc_shmem_cmps cmp, T* vals) {
+  return host_interface->wait_until_any_vector<T>(ptr, nelems, status, cmp, vals, context_window_info);
+}
+
+template <typename T>
+__host__ size_t IPCHostContext::wait_until_some_vector(T *ptr, size_t nelems,
+                                                         size_t* indices,
+                                                         const int* status,
+                                                         roc_shmem_cmps cmp, T* vals) {
+  return host_interface->wait_until_some_vector<T>(ptr, nelems, indices, status, cmp, vals, context_window_info);
+}
+
+template <typename T>
+__host__ int IPCHostContext::test(T *ptr, roc_shmem_cmps cmp, T val) {
+  return host_interface->test<T>(ptr, cmp, val, context_window_info);
+}
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_

projects/rocshmem/src/ipc/ipc_backend_proxy.hpp

@@ -0,0 +1,69 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_
+#define LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_
+
+#include "../device_proxy.hpp"
+#include "../atomic_return.hpp"
+
+namespace rocshmem {
+
+struct IPCBackendRegister {
+  char *g_ret{nullptr};
+  atomic_ret_t *atomic_ret{nullptr};
+  SymmetricHeap *heap_ptr{nullptr};
+};
+
+template <typename ALLOCATOR>
+class IPCBackendProxy {
+  using ProxyT = DeviceProxy<ALLOCATOR, IPCBackendRegister>;
+
+ public:
+  /*
+   * Placement new the memory which is allocated by proxy_
+   */
+  IPCBackendProxy() { new (proxy_.get()) IPCBackendRegister(); }
+
+  /*
+   * Since placement new is called in the constructor, then
+   * delete must be called manually.
+   */
+  ~IPCBackendProxy() { proxy_.get()->~IPCBackendRegister(); }
+
+  /*
+   * @brief Provide access to the memory referenced by the proxy
+   */
+  __host__ __device__ IPCBackendRegister *get() { return proxy_.get(); }
+
+ private:
+  /*
+   * @brief Memory managed by the lifetime of this object
+   */
+  ProxyT proxy_{};
+};
+
+using IPCBackendProxyT = IPCBackendProxy<HIPHostAllocator>;
+
+}  // namespace rocshmem
+
+#endif  // #define LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_

projects/rocshmem/src/ipc/ipc_context_proxy.hpp

@@ -0,0 +1,90 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_
+#define LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_
+
+
+#include "../device_proxy.hpp"
+#include "backend_ipc.hpp"
+
+namespace rocshmem {
+
+class IPCBackend;
+
+template <typename ALLOCATOR>
+class IPCDefaultContextProxy {
+  using ProxyT = DeviceProxy<ALLOCATOR, IPCContext>;
+
+ public:
+  IPCDefaultContextProxy() = default;
+
+  /*
+   * Placement new the memory which is allocated by proxy_
+   */
+  explicit IPCDefaultContextProxy(IPCBackend* backend) : constructed_{true} {
+    auto ctx{proxy_.get()};
+    new (ctx) IPCContext(reinterpret_cast<Backend*>(backend));
+    roc_shmem_ctx_t local{ctx, nullptr};
+    set_internal_ctx(&local);
+  }
+
+  /*
+   * Since placement new is called in the constructor, then
+   * delete must be called manually.
+   */
+  ~IPCDefaultContextProxy() {
+    if (constructed_) {
+      proxy_.get()->~IPCContext();
+    }
+  }
+
+  IPCDefaultContextProxy(const IPCDefaultContextProxy& other) = delete;
+
+  IPCDefaultContextProxy& operator=(const IPCDefaultContextProxy& other) = delete;
+
+  IPCDefaultContextProxy(IPCDefaultContextProxy&& other) = default;
+
+  IPCDefaultContextProxy& operator=(IPCDefaultContextProxy&& other) = default;
+
+  /*
+   * @brief Provide access to the memory referenced by the proxy
+   */
+  __host__ __device__ Context* get() { return proxy_.get(); }
+
+ private:
+  /*
+   * @brief Memory managed by the lifetime of this object
+   */
+  ProxyT proxy_{};
+
+  /*
+   * @brief denotes if an objects was constructed in proxy
+   */
+  bool constructed_{false};
+};
+
+using IPCDefaultContextProxyT = IPCDefaultContextProxy<HIPAllocator>;
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_

projects/rocshmem/src/ipc_policy.cpp

@@ -25,9 +25,9 @@
 #include <mpi.h>
 
 #include "config.h"  // NOLINT(build/include_subdir)
-#include "../backend_bc.hpp"
-#include "../context_incl.hpp"
-#include "../util.hpp"
+#include "backend_bc.hpp"
+#include "context_incl.hpp"
+#include "util.hpp"
 
 namespace rocshmem {
 

projects/rocshmem/src/ipc_policy.hpp

@@ -30,8 +30,8 @@
 #include <vector>
 
 #include "config.h"  // NOLINT(build/include_subdir)
-#include "../memory/hip_allocator.hpp"
-#include "../util.hpp"
+#include "memory/hip_allocator.hpp"
+#include "util.hpp"
 
 namespace rocshmem {
 

projects/rocshmem/src/reverse_offload/block_handle.hpp

@@ -24,7 +24,7 @@
 #define LIBRARY_SRC_REVERSE_OFFLOAD_BLOCK_HANDLE_HPP_
 
 #include "../hdp_policy.hpp"
-#include "../ipc/ipc_policy.hpp"
+#include "../ipc_policy.hpp"
 #include "profiler.hpp"
 #include "queue.hpp"
 

projects/rocshmem/src/reverse_offload/queue_proxy.hpp

@@ -28,7 +28,7 @@
 #include "../atomic_return.hpp"
 #include "../device_proxy.hpp"
 #include "../hdp_policy.hpp"
-#include "../ipc/ipc_policy.hpp"
+#include "../ipc_policy.hpp"
 #include "commands_types.hpp"
 #include "profiler.hpp"
 #include "../sync/abql_block_mutex.hpp"

projects/rocshmem/src/roc_shmem.cpp

@@ -39,12 +39,14 @@
 #ifdef USE_GPU_IB
 #include "gpu_ib/backend_ib.hpp"
 #include "gpu_ib/context_ib_tmpl_host.hpp"
-#endif
-#include "mpi_init_singleton.hpp"
-#ifndef USE_GPU_IB
+#elif defined(USE_RO)
 #include "reverse_offload/backend_ro.hpp"
 #include "reverse_offload/context_ro_tmpl_host.hpp"
+#else
+#include "ipc/backend_ipc.hpp"
+#include "ipc/context_ipc_tmpl_host.hpp"
 #endif
+#include "mpi_init_singleton.hpp"
 #include "team.hpp"
 #include "templates_host.hpp"
 #include "util.hpp"
@@ -82,12 +84,15 @@ roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT;
 
   rocm_init();
 
-#ifndef USE_GPU_IB
+#ifdef USE_GPU_IB
+  CHECK_HIP(hipHostMalloc(&backend, sizeof(GPUIBBackend)));
+  backend = new (backend) GPUIBBackend(comm);
+#elif defined(USE_RO)
   CHECK_HIP(hipHostMalloc(&backend, sizeof(ROBackend)));
   backend = new (backend) ROBackend(comm);
 #else
-  CHECK_HIP(hipHostMalloc(&backend, sizeof(GPUIBBackend)));
-  backend = new (backend) GPUIBBackend(comm);
+  CHECK_HIP(hipHostMalloc(&backend, sizeof(IPCBackend)));
+  backend = new (backend) IPCBackend(comm);
 #endif
 
   if (!backend) {

projects/rocshmem/src/roc_shmem_gpu.cpp

@@ -51,8 +51,10 @@
 
 #ifdef USE_GPU_IB
 #include "gpu_ib/context_ib_tmpl_device.hpp"
-#else
+#elif defined(USE_RO)
 #include "reverse_offload/context_ro_tmpl_device.hpp"
+#else
+#include "ipc/context_ipc_tmpl_device.hpp"
 #endif
 
 /******************************************************************************

projects/rocshmem/tests/unit_tests/CMakeLists.txt

@@ -88,7 +88,7 @@ target_sources(
     notifier_gtest.cpp
     #forward_list_gtest.cpp
     free_list_gtest.cpp
-    context_ipc_gtest.cpp
+    #context_ipc_gtest.cpp
     ipc_impl_simple_coarse_gtest.cpp
 )
 

projects/rocshmem/tests/unit_tests/context_ipc_gtest.hpp

@@ -25,8 +25,8 @@
 
 #include "gtest/gtest.h"
 
-#include "../src/ipc/context_ipc.hpp"
-#include "../src/reverse_offload/backend_ro.hpp"
+#include "../src/ipc/context_ipc_device.hpp"
+#include "../src/ipc/backend_ipc.hpp"
 
 namespace rocshmem {
 
@@ -36,7 +36,7 @@ class ContextIpcTestFixture : public ::testing::Test
     /**
      * @brief Context Ipc Test
      */
-   ROBackend be{MPI_COMM_WORLD};
+   IPCBackend be{MPI_COMM_WORLD};
 
    IPCContext ipc_context_ {&be};
 };

projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp

@@ -29,7 +29,7 @@
 
 #include <mpi.h>
 #include "../src/memory/symmetric_heap.hpp"
-#include "../src/ipc/ipc_policy.hpp"
+#include "../src/ipc_policy.hpp"
 
 namespace rocshmem {