Merge pull request #27 from ROCm/ipc_bringup

Ipc bringup

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)

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

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 .

scripts/build_configs/ro_net

@@ -18,7 +18,7 @@ cmake \
     -DPROFILE=OFF \
     -DUSE_GPU_IB=OFF \
     -DUSE_DC=OFF \
-    -DUSE_IPC=OFF \
+    -DUSE_IPC=ON \
     -DUSE_THREADS=ON \
     -DUSE_WF_COAL=OFF \
     -DUSE_COHERENT_HEAP=ON \

src/CMakeLists.txt

@@ -30,7 +30,6 @@ target_sources(
     backend_bc.cpp
     context_host.cpp
     context_device.cpp
-    ipc_policy.cpp
     mpi_init_singleton.cpp
     roc_shmem_gpu.cpp
     roc_shmem.cpp
@@ -38,6 +37,7 @@ target_sources(
     team_tracker.cpp
     util.cpp
     wf_coal_policy.cpp
+    ipc_policy.cpp
 )
 
 target_compile_options(
@@ -60,8 +60,10 @@ 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)

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
 }
 

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
 
 /**
@@ -69,10 +72,15 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
 #define DISPATCH_RET(Func)                           \
   auto ret_val = static_cast<GPUIBContext *>(this)->Func; \
   return ret_val;
-#else
+#elif defined(USE_RO)
 #define DISPATCH_RET(Func)                        \
   auto 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.
@@ -82,11 +90,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
 
 /**
@@ -98,8 +111,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.
@@ -113,10 +128,15 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND };
 #define HOST_DISPATCH_RET(Func)                          \
   auto ret_val = static_cast<GPUIBHostContext *>(this)->Func; \
   return ret_val;
-#else
+#elif defined(USE_RO)
 #define HOST_DISPATCH_RET(Func)                       \
   auto 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

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_

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 {

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 {
 

src/ipc/CMakeLists.txt

@@ -0,0 +1,34 @@
+###############################################################################
+# 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.
+###############################################################################
+
+###############################################################################
+# ADD ROCSHMEM TARGET FOR FILES IN CURRENT DIRECTORY
+###############################################################################
+target_sources(
+  ${PROJECT_NAME}
+  PRIVATE
+    context_ipc_device.cpp
+    context_ipc_host.cpp
+    backend_ipc.cpp
+    ipc_team.cpp
+    context_ipc_device_coll.cpp
+)

src/ipc/backend_ipc.cpp

@@ -0,0 +1,392 @@
+/******************************************************************************
+ * 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"
+#include "ipc_team.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;
+
+roc_shmem_team_t get_external_team(GPUIBTeam *team) {
+  return reinterpret_cast<roc_shmem_team_t>(team);
+}
+
+int get_ls_non_zero_bit(char *bitmask, int mask_length) {
+  int position = -1;
+
+  for (int bit_i = 0; bit_i < mask_length; bit_i++) {
+    int byte_i = bit_i / CHAR_BIT;
+    if (bitmask[byte_i] & (1 << (bit_i % CHAR_BIT))) {
+      position = bit_i;
+      break;
+    }
+  }
+
+  return position;
+}
+
+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);
+
+  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_team_world();
+
+  roc_shmem_collective_init();
+
+  teams_init();
+
+  setup_ctxs();
+  
+}
+
+IPCBackend::~IPCBackend() {
+  /*
+   * 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
+
+  // delete host_interface;
+  // host_interface = nullptr;
+
+  /**
+   * Destroy teams infrastructure
+   * and team world
+   */
+  teams_destroy();
+  auto *team_world{team_tracker.get_team_world()};
+  team_world->~Team();
+  CHECK_HIP(hipFree(team_world));
+
+  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::setup_team_world() {
+  TeamInfo *team_info_wrt_parent, *team_info_wrt_world;
+
+  /**
+   * Allocate device-side memory for team_world and construct a
+   * IPC team in it.
+   */
+  CHECK_HIP(hipMalloc(&team_info_wrt_parent, sizeof(TeamInfo)));
+  CHECK_HIP(hipMalloc(&team_info_wrt_world, sizeof(TeamInfo)));
+
+  new (team_info_wrt_parent) TeamInfo(nullptr, 0, 1, num_pes);
+  new (team_info_wrt_world) TeamInfo(nullptr, 0, 1, num_pes);
+
+  IPCTeam *team_world{nullptr};
+  CHECK_HIP(hipMalloc(&team_world, sizeof(IPCTeam)));
+  new (team_world) IPCTeam(this, team_info_wrt_parent, team_info_wrt_world,
+                             num_pes, my_pe, thread_comm, 0);
+  team_tracker.set_team_world(team_world);
+
+  /**
+   * Copy the address to ROC_SHMEM_TEAM_WORLD.
+   */
+  ROC_SHMEM_TEAM_WORLD = reinterpret_cast<roc_shmem_team_t>(team_world);
+}
+
+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) {
+  IPCTeam *team_obj = get_internal_ipc_team(team);
+
+  /* Mark the pool as available */
+  int bit = team_obj->pool_index_;
+  int byte_i = bit / CHAR_BIT;
+  pool_bitmask_[byte_i] |= 1 << (bit % CHAR_BIT);
+
+  team_obj->~IPCTeam();
+  CHECK_HIP(hipFree(team_obj));
+}
+
+void IPCBackend::create_new_team([[maybe_unused]] 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) {
+  /**
+   * Read the bit mask and find out a common index into
+   * the pool of available work arrays.
+   */
+  NET_CHECK(MPI_Allreduce(pool_bitmask_, reduced_bitmask_, bitmask_size_,
+                          MPI_CHAR, MPI_BAND, team_comm));
+
+  /* Pick the least significant non-zero bit (logical layout) in the reduced
+   * bitmask */
+  auto max_num_teams{team_tracker.get_max_num_teams()};
+  int common_index = get_ls_non_zero_bit(reduced_bitmask_, max_num_teams);
+  if (common_index < 0) {
+    /* No team available */
+    abort();
+  }
+
+  /* Mark the team as taken (by unsetting the bit in the pool bitmask) */
+  int byte = common_index / CHAR_BIT;
+  pool_bitmask_[byte] &= ~(1 << (common_index % CHAR_BIT));
+
+  /**
+   * Allocate device-side memory for team_world and
+   * construct a GPU_IB team in it
+   */
+  GPUIBTeam *new_team_obj;
+  CHECK_HIP(hipMalloc(&new_team_obj, sizeof(IPCTeam)));
+  new (new_team_obj)
+      IPCTeam(this, team_info_wrt_parent, team_info_wrt_world, num_pes,
+                my_pe_in_new_team, team_comm, common_index);
+
+  *new_team = get_external_team(new_team_obj);
+}
+
+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::global_exit(int status) {
+  assert(false);
+}
+
+void IPCBackend::teams_destroy() {
+  roc_shmem_free(barrier_pSync_pool);
+  roc_shmem_free(reduce_pSync_pool);
+  roc_shmem_free(bcast_pSync_pool);
+  roc_shmem_free(alltoall_pSync_pool);
+  roc_shmem_free(pWrk_pool);
+  roc_shmem_free(pAta_pool);
+
+  free(pool_bitmask_);
+  free(reduced_bitmask_);
+}
+
+void IPCBackend::roc_shmem_collective_init() {
+  /*
+   * Allocate heap space for barrier_sync
+   */
+  size_t one_sync_size_bytes{sizeof(*barrier_sync)};
+  size_t sync_size_bytes{one_sync_size_bytes * ROC_SHMEM_BARRIER_SYNC_SIZE};
+  heap.malloc(reinterpret_cast<void **>(&barrier_sync), sync_size_bytes);
+
+  /*
+   * Initialize the barrier synchronization array with default values.
+   */
+  for (int i = 0; i < num_pes; i++) {
+    barrier_sync[i] = ROC_SHMEM_SYNC_VALUE;
+  }
+
+  /*
+   * Make sure that all processing elements have done this before
+   * continuing.
+   */
+  NET_CHECK(MPI_Barrier(thread_comm));
+}
+
+void IPCBackend::teams_init() {
+  /**
+   * Allocate pools for the teams sync and work arrary from the SHEAP.
+   */
+  auto max_num_teams{team_tracker.get_max_num_teams()};
+  barrier_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
+      sizeof(long) * ROC_SHMEM_BARRIER_SYNC_SIZE * max_num_teams));
+  reduce_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
+      sizeof(long) * ROC_SHMEM_REDUCE_SYNC_SIZE * max_num_teams));
+  bcast_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
+      sizeof(long) * ROC_SHMEM_BCAST_SYNC_SIZE * max_num_teams));
+  alltoall_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
+      sizeof(long) * ROC_SHMEM_ALLTOALL_SYNC_SIZE * max_num_teams));
+
+  /* Accommodating for largest possible data type for pWrk */
+  pWrk_pool = roc_shmem_malloc(
+      sizeof(double) * ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * max_num_teams);
+  pAta_pool = roc_shmem_malloc(sizeof(double) * ROC_SHMEM_ATA_MAX_WRKDATA_SIZE *
+                               max_num_teams);
+
+  /**
+   * Initialize the sync arrays in the pool with default values.
+   */
+  long *barrier_pSync, *reduce_pSync, *bcast_pSync, *alltoall_pSync;
+  for (int team_i = 0; team_i < max_num_teams; team_i++) {
+    barrier_pSync = reinterpret_cast<long *>(
+        &barrier_pSync_pool[team_i * ROC_SHMEM_BARRIER_SYNC_SIZE]);
+    reduce_pSync = reinterpret_cast<long *>(
+        &reduce_pSync_pool[team_i * ROC_SHMEM_REDUCE_SYNC_SIZE]);
+    bcast_pSync = reinterpret_cast<long *>(
+        &bcast_pSync_pool[team_i * ROC_SHMEM_BCAST_SYNC_SIZE]);
+    alltoall_pSync = reinterpret_cast<long *>(
+        &alltoall_pSync_pool[team_i * ROC_SHMEM_ALLTOALL_SYNC_SIZE]);
+
+    for (int i = 0; i < ROC_SHMEM_BARRIER_SYNC_SIZE; i++) {
+      barrier_pSync[i] = ROC_SHMEM_SYNC_VALUE;
+    }
+    for (int i = 0; i < ROC_SHMEM_REDUCE_SYNC_SIZE; i++) {
+      reduce_pSync[i] = ROC_SHMEM_SYNC_VALUE;
+    }
+    for (int i = 0; i < ROC_SHMEM_BCAST_SYNC_SIZE; i++) {
+      bcast_pSync[i] = ROC_SHMEM_SYNC_VALUE;
+    }
+    for (int i = 0; i < ROC_SHMEM_ALLTOALL_SYNC_SIZE; i++) {
+      alltoall_pSync[i] = ROC_SHMEM_SYNC_VALUE;
+    }
+  }
+
+  /**
+   * Initialize bit mask
+   *
+   * Logical:
+   * MSB..........................................................................LSB
+   * Physical: MSB...1st least significant 8 bits...LSB  MSB...2nd least
+   * signifant 8 bits...LSB
+   *
+   * Description shows only a 2-byte long mask but idea extends to any
+   * arbitrary size.
+   */
+  bitmask_size_ = (max_num_teams % CHAR_BIT) ? (max_num_teams / CHAR_BIT + 1)
+                                             : (max_num_teams / CHAR_BIT);
+  pool_bitmask_ = reinterpret_cast<char *>(malloc(bitmask_size_));
+  reduced_bitmask_ = reinterpret_cast<char *>(malloc(bitmask_size_));
+
+  memset(pool_bitmask_, 0, bitmask_size_);
+  memset(reduced_bitmask_, 0, bitmask_size_);
+  /* Set all to available except the 0th one (reserved for TEAM_WORLD) */
+  for (int bit_i = 1; bit_i < max_num_teams; bit_i++) {
+    int byte_i = bit_i / CHAR_BIT;
+
+    pool_bitmask_[byte_i] |= 1 << (bit_i % CHAR_BIT);
+  }
+
+  /**
+   * Make sure that all processing elements have done this before
+   * continuing.
+   */
+  NET_CHECK(MPI_Barrier(thread_comm));
+}
+
+}  // namespace rocshmem

src/ipc/backend_ipc.hpp

@@ -0,0 +1,252 @@
+/******************************************************************************
+ * 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 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};
+
+  /**
+   * @brief Scratchpad for the internal barrier algorithms.
+   */
+  int64_t *barrier_sync{nullptr};
+
+  /**
+   * @brief Handle for raw memory for barrier sync
+   */
+  long *barrier_pSync_pool{nullptr};
+
+  /**
+   * @brief Handle for raw memory for reduce sync
+   */
+  long *reduce_pSync_pool{nullptr};
+
+  /**
+   * @brief Handle for raw memory for broadcast sync
+   */
+  long *bcast_pSync_pool{nullptr};
+
+  /**
+   * @brief Handle for raw memory for alltoall sync
+   */
+  long *alltoall_pSync_pool{nullptr};
+
+  /**
+   * @brief Handle for raw memory for work
+   */
+  void *pWrk_pool{nullptr};
+
+  /**
+   * @brief Handle for raw memory for alltoall
+   */
+  void *pAta_pool{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};
+
+  /**
+   * @brief Allocate and initialize team world.
+   */
+  void setup_team_world();
+
+  /**
+   * @brief Initialize the resources required to support teams
+   */
+  void teams_init();
+
+  /**
+   * @brief Destruct the resources required to support teams
+   */
+  void teams_destroy();
+
+  /**
+   * @brief Allocate and initialize barrier operation addresses on
+   * symmetric heap.
+   *
+   * When this method completes, the barrier_sync member will be available
+   * for use.
+   */
+  void roc_shmem_collective_init();
+
+ 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};
+
+  /**
+   * @brief The bitmask representing the availability of teams in the pool
+   */
+  char *pool_bitmask_{nullptr};
+
+  /**
+   * @brief Bitmask to store the reduced result of bitmasks on pariticipating
+   * PEs
+   *
+   * With no thread-safety for this bitmask, multithreaded creation of teams is
+   * not supported.
+   */
+  char *reduced_bitmask_{nullptr};
+
+  /**
+   * @brief Size of the bitmask
+   */
+  int bitmask_size_{-1};
+
+};
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_BACKEND_HPP_

src/ipc/context_ipc_device.cpp

@@ -0,0 +1,156 @@
+/******************************************************************************
+ * 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_device.hpp"
+#include "context_ipc_tmpl_device.hpp"
+
+#include <hip/hip_runtime.h>
+#include <hip/amd_detail/amd_device_functions.h>
+#include <unistd.h>
+
+#include <cstdio>
+#include <cstdlib>
+
+#include "config.h"  // NOLINT(build/include_subdir)
+#include "roc_shmem/roc_shmem.hpp"
+#include "backend_ipc.hpp"
+
+namespace rocshmem {
+
+__host__ IPCContext::IPCContext(Backend *b)
+    : Context(b, false) {
+  IPCBackend *backend{static_cast<IPCBackend *>(b)};
+  ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases;
+  ipcImpl_.shm_size = b->ipcImpl.shm_size;
+
+  auto *bp{backend->ipc_backend_proxy.get()};
+
+  barrier_sync = backend->barrier_sync;
+  g_ret = bp->g_ret;
+  atomic_base_ptr = bp->atomic_ret->atomic_base_ptr;
+}
+
+__device__ void IPCContext::threadfence_system() {
+}
+
+__device__ void IPCContext::ctx_create() {
+}
+
+__device__ void IPCContext::ctx_destroy(){
+}
+
+__device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems,
+                                  int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[pe] + L_offset,
+                   const_cast<void *>(source), nelems);
+}
+
+__device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems,
+                                  int pe) {
+  const char *src_typed = reinterpret_cast<const char *>(source);
+  uint64_t L_offset =
+      const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems);
+}
+
+__device__ void IPCContext::putmem_nbi(void *dest, const void *source,
+                                      size_t nelems, int pe) {
+  putmem(dest, source, nelems, pe);
+}
+
+__device__ void IPCContext::getmem_nbi(void *dest, const void *source,
+                                      size_t nelems, int pe) {
+  getmem(dest, source, nelems, pe);
+}
+
+__device__ void IPCContext::fence() {
+}
+
+__device__ void IPCContext::fence(int pe) {
+}
+
+__device__ void IPCContext::quiet() {
+}
+
+__device__ void *IPCContext::shmem_ptr(const void *dest, int pe) {
+  void *ret = nullptr;
+  return ret;
+}
+
+__device__ void IPCContext::putmem_wg(void *dest, const void *source,
+                                     size_t nelems, int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[pe] + L_offset,
+                      const_cast<void *>(source), nelems);
+  __syncthreads();
+}
+
+__device__ void IPCContext::getmem_wg(void *dest, const void *source,
+                                     size_t nelems, int pe) {
+  const char *src_typed = reinterpret_cast<const char *>(source);
+  uint64_t L_offset =
+      const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems);
+  __syncthreads();
+}
+
+__device__ void IPCContext::putmem_nbi_wg(void *dest, const void *source,
+                                         size_t nelems, int pe) {
+  putmem_wg(dest, source, nelems, pe);
+}
+
+__device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source,
+                                         size_t nelems, int pe) {
+  getmem_wg(dest, source, nelems, pe);
+}
+
+__device__ void IPCContext::putmem_wave(void *dest, const void *source,
+                                       size_t nelems, int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[pe] + L_offset,
+                        const_cast<void *>(source), nelems);
+}
+
+__device__ void IPCContext::getmem_wave(void *dest, const void *source,
+                                       size_t nelems, int pe) {
+  const char *src_typed = reinterpret_cast<const char *>(source);
+  uint64_t L_offset =
+      const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[pe] + L_offset,
+                        nelems);
+}
+
+__device__ void IPCContext::putmem_nbi_wave(void *dest, const void *source,
+                                           size_t nelems, int pe) {
+  putmem_wave(dest, source, nelems, pe);
+}
+
+__device__ void IPCContext::getmem_nbi_wave(void *dest, const void *source,
+                                           size_t nelems, int pe) {
+  getmem_wave(dest, source, nelems, pe);
+}
+
+}  // namespace rocshmem

src/ipc/context_ipc_device.hpp

@@ -0,0 +1,239 @@
+/******************************************************************************
+ * 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_DEVICE_HPP_
+#define LIBRARY_SRC_IPC_CONTEXT_DEVICE_HPP_
+
+#include "../context.hpp"
+
+namespace rocshmem {
+
+class IPCContext : public Context {
+ public:
+  __host__ IPCContext(Backend *b);
+
+  __device__ IPCContext(Backend *b);
+
+  __device__ void threadfence_system();
+
+  __device__ void ctx_create();
+
+  __device__ void ctx_destroy();
+
+  __device__ void putmem(void *dest, const void *source, size_t nelems, int pe);
+
+  __device__ void getmem(void *dest, const void *source, size_t nelems, int pe);
+
+  __device__ void putmem_nbi(void *dest, const void *source, size_t nelems,
+                             int pe);
+
+  __device__ void getmem_nbi(void *dest, const void *source, size_t size,
+                             int pe);
+
+  __device__ void fence();
+
+  __device__ void fence(int pe);
+
+  __device__ void quiet();
+
+  __device__ void *shmem_ptr(const void *dest, int pe);
+
+  __device__ void barrier_all();
+
+  __device__ void sync_all();
+
+  __device__ void sync(roc_shmem_team_t team);
+
+  template <typename T>
+  __device__ void p(T *dest, T value, int pe);
+
+  template <typename T>
+  __device__ void put(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void put_nbi(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ T g(const T *source, int pe);
+
+  template <typename T>
+  __device__ void get(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void get_nbi(T *dest, const T *source, size_t nelems, int pe);
+
+  // Atomic operations
+  template <typename T>
+  __device__ void amo_add(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ void amo_set(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ T amo_swap(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ T amo_fetch_and(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ void amo_and(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ T amo_fetch_or(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ void amo_or(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ T amo_fetch_xor(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ void amo_xor(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ void amo_cas(void *dst, T value, T cond, int pe);
+
+  template <typename T>
+  __device__ T amo_fetch_add(void *dst, T value, int pe);
+
+  template <typename T>
+  __device__ T amo_fetch_cas(void *dst, T value, T cond, int pe);
+
+  // Collectives
+  template <typename T, ROC_SHMEM_OP Op>
+  __device__ void to_all(T *dest, const T *source, int nreduce, int PE_start,
+                         int logPE_stride, int PE_size, T *pWrk,
+                         long *pSync);  // NOLINT(runtime/int)
+
+  template <typename T, ROC_SHMEM_OP Op>
+  __device__ void to_all(roc_shmem_team_t team, T *dest, const T *source,
+                         int nreduce);
+
+  template <typename T>
+  __device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
+                            int nelems, int pe_root);
+
+  template <typename T>
+  __device__ 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);  // NOLINT(runtime/int)
+  template <typename T>
+  __device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source,
+                           int nelems);
+  template <typename T>
+  __device__ void fcollect(roc_shmem_team_t team, T *dest, const T *source,
+                           int nelems);
+
+
+  // Block/wave functions
+  __device__ void putmem_wg(void *dest, const void *source, size_t nelems,
+                            int pe);
+
+  __device__ void getmem_wg(void *dest, const void *source, size_t nelems,
+                            int pe);
+
+  __device__ void putmem_nbi_wg(void *dest, const void *source, size_t nelems,
+                                int pe);
+
+  __device__ void getmem_nbi_wg(void *dest, const void *source, size_t size,
+                                int pe);
+
+  __device__ void putmem_wave(void *dest, const void *source, size_t nelems,
+                              int pe);
+
+  __device__ void getmem_wave(void *dest, const void *source, size_t nelems,
+                              int pe);
+
+  __device__ void putmem_nbi_wave(void *dest, const void *source, size_t nelems,
+                                  int pe);
+
+  __device__ void getmem_nbi_wave(void *dest, const void *source, size_t size,
+                                  int pe);
+
+  template <typename T>
+  __device__ void put_wg(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void put_nbi_wg(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void put_wave(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void put_nbi_wave(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void get_wg(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void get_nbi_wg(T *dest, const T *source, size_t nelems, int pe);
+
+
+  template <typename T>
+  __device__ void get_wave(T *dest, const T *source, size_t nelems, int pe);
+
+  template <typename T>
+  __device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe);
+
+ private:
+
+  //context class has IpcImpl object (ipcImpl_)
+  IpcImpl *ipcImpl{nullptr};
+
+  uint64_t* atomic_base_ptr{nullptr};
+
+  char* g_ret;
+
+  //internal functions used by collective operations
+   template <typename T>
+  __device__ void internal_put_broadcast(T *dst, const T *src, int nelems,
+                                         int pe_root, int PE_start,
+                                         int logPE_stride, int PE_size);  // NOLINT(runtime/int)
+
+  template <typename T>
+  __device__ void internal_get_broadcast(T *dst, const T *src, int nelems,
+                                         int pe_root);  // NOLINT(runtime/int)
+
+  template <typename T>
+  __device__ void fcollect_linear(roc_shmem_team_t team, T *dest,
+                                  const T *source, int nelems);
+
+  template <typename T>
+  __device__ void alltoall_linear(roc_shmem_team_t team, T *dest,
+                                  const T *source, int nelems);
+
+  __device__ void internal_sync(int pe, int PE_start, int stride, int PE_size,
+                                int64_t *pSync);
+
+  __device__ void internal_direct_barrier(int pe, int PE_start, int stride,
+                                          int n_pes, int64_t *pSync);
+
+  __device__ void internal_atomic_barrier(int pe, int PE_start, int stride,
+                                          int n_pes, int64_t *pSync);
+
+  //Temporary scratchpad memory used by internal barrier algorithms.
+  int64_t *barrier_sync{nullptr};
+};
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_GPU_IB_CONTEXT_IB_DEVICE_HPP_

src/ipc/context_ipc_device_coll.cpp

@@ -0,0 +1,120 @@
+/******************************************************************************
+ * 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 "roc_shmem/roc_shmem.hpp"
+#include "../context_incl.hpp"
+#include "context_ipc_tmpl_device.hpp"
+#include "../util.hpp"
+#include "ipc_team.hpp"
+
+namespace rocshmem {
+
+__device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
+                                                      int stride, int n_pes,
+                                                      int64_t *pSync) {
+  int64_t flag_val = 1;
+  if (pe == PE_start) {
+    // Go through all PE offsets (except current offset = 0)
+    // and wait until they all reach
+    for (size_t i = 1; i < n_pes; i++) {
+      wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val);
+      pSync[i] = ROC_SHMEM_SYNC_VALUE;
+    }
+    threadfence_system();
+
+    // Announce to other PEs that all have reached
+    for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
+      put_nbi(&pSync[0], &flag_val, 1, j);
+    }
+  } else {
+    // Mark current PE offset as reached
+    size_t pe_offset = (pe - PE_start) / stride;
+    put_nbi(&pSync[pe_offset], &flag_val, 1, PE_start);
+    wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
+    pSync[0] = ROC_SHMEM_SYNC_VALUE;
+    threadfence_system();
+  }
+}
+
+__device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start,
+                                                      int stride, int n_pes,
+                                                      int64_t *pSync) {
+  int64_t flag_val = 1;
+  if (pe == PE_start) {
+    wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, (int64_t)(n_pes - 1));
+    pSync[0] = ROC_SHMEM_SYNC_VALUE;
+    threadfence_system();
+
+    for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
+      put_nbi(&pSync[0], &flag_val, 1, j);
+    }
+  } else {
+    amo_add<int64_t>(&pSync[0], flag_val, PE_start);
+    wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
+    pSync[0] = ROC_SHMEM_SYNC_VALUE;
+    threadfence_system();
+  }
+}
+
+// Uses PE values that are relative to world
+__device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
+                                          int PE_size, int64_t *pSync) {
+  __syncthreads();
+  if (is_thread_zero_in_block()) {
+    if (PE_size < 64) {
+      internal_direct_barrier(pe, PE_start, stride, PE_size, pSync);
+    } else {
+      internal_atomic_barrier(pe, PE_start, stride, PE_size, pSync);
+    }
+  }
+  __threadfence();
+  __syncthreads();
+}
+
+__device__ void IPCContext::sync(roc_shmem_team_t team) {
+  IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
+
+  /**
+   * Ensure that the stride is a multiple of 2.
+   */
+  int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
+  int pe = team_obj->my_pe_in_world;
+  int pe_start = team_obj->tinfo_wrt_world->pe_start;
+  int pe_stride = (1 << log_pe_stride);
+  int pe_size = team_obj->num_pes;
+
+  internal_sync(pe, pe_start, pe_stride, pe_size, barrier_sync);
+}
+
+__device__ void IPCContext::sync_all() {
+  internal_sync(my_pe, 0, 1, num_pes, barrier_sync);
+}
+
+__device__ void IPCContext::barrier_all() {
+  if (is_thread_zero_in_block()) {
+    quiet();
+  }
+  sync_all();
+  __syncthreads();
+}
+
+}  // namespace 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

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_

src/ipc/context_ipc_tmpl_device.hpp

@@ -0,0 +1,345 @@
+/******************************************************************************
+ * 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_DEVICE_HPP_
+#define LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_
+
+#include "config.h"  // NOLINT(build/include_subdir)
+#include "roc_shmem/roc_shmem.hpp"
+#include "context_ipc_device.hpp"
+#include "../util.hpp"
+#include "ipc_team.hpp"
+
+namespace rocshmem {
+
+/******************************************************************************
+ ************************** TEMPLATE SPECIALIZATIONS **************************
+ *****************************************************************************/
+template <typename T>
+__device__ void IPCContext::p(T *dest, T value, int pe) {
+  putmem_nbi(dest, &value, sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::put(T *dest, const T *source, size_t nelems,
+				int pe) {
+  putmem(dest, source, nelems * sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::put_nbi(T *dest, const T *source, size_t nelems,
+				    int pe) {
+  putmem_nbi(dest, source, sizeof(T) * nelems, pe);
+}
+
+template <typename T>
+__device__ T IPCContext::g(const T *source, int pe) {
+  T ret;
+  return ret;
+}
+
+template <typename T>
+__device__ void IPCContext::get(T *dest, const T *source, size_t nelems,
+				int pe) {
+  getmem(dest, source, sizeof(T) * nelems, pe);
+}
+
+template <typename T>
+__device__ void IPCContext::get_nbi(T *dest, const T *source, size_t nelems,
+				    int pe) {
+  getmem_nbi(dest, source, sizeof(T) * nelems, pe);
+}
+
+// Atomics
+template <typename T>
+__device__ void IPCContext::amo_add(void *dest, T value, int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcAMOAdd(
+      reinterpret_cast<T *>(ipcImpl_.ipc_bases[pe] + L_offset), value);
+}
+
+template <typename T>
+__device__ void IPCContext::amo_set(void *dest, T value, int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcAMOSet(
+      reinterpret_cast<T *>(ipcImpl_.ipc_bases[pe] + L_offset), value);
+}
+
+template <typename T>
+__device__ T IPCContext::amo_swap(void *dst, T value, int pe) {
+  assert(false);
+  return 0;
+}
+
+template <typename T>
+__device__ T IPCContext::amo_fetch_and(void *dst, T value, int pe) {
+  assert(false);
+  return 0;
+}
+
+template <typename T>
+__device__ void IPCContext::amo_and(void *dst, T value, int pe) {
+  assert(false);
+}
+
+template <typename T>
+__device__ T IPCContext::amo_fetch_or(void *dst, T value, int pe) {
+  assert(false);
+  return 0;
+}
+
+template <typename T>
+__device__ void IPCContext::amo_or(void *dst, T value, int pe) {
+  assert(false);
+}
+
+template <typename T>
+__device__ T IPCContext::amo_fetch_xor(void *dst, T value, int pe) {
+  assert(false);
+  return 0;
+}
+
+template <typename T>
+__device__ void IPCContext::amo_xor(void *dst, T value, int pe) {
+  assert(false);
+}
+
+template <typename T>
+__device__ void IPCContext::amo_cas(void *dest, T value, T cond, int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  ipcImpl_.ipcAMOCas(
+      reinterpret_cast<T *>(ipcImpl_.ipc_bases[pe] + L_offset), cond,
+      value);
+}
+
+template <typename T>
+__device__ T IPCContext::amo_fetch_add(void *dest, T value, int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  return ipcImpl_.ipcAMOFetchAdd(
+      reinterpret_cast<T *>(ipcImpl_.ipc_bases[pe] + L_offset), value);
+}
+
+template <typename T>
+__device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) {
+  uint64_t L_offset =
+      reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
+  return ipcImpl_.ipcAMOFetchCas(
+      reinterpret_cast<T *>(ipcImpl_.ipc_bases[pe] + L_offset), cond,
+      value);
+}
+  
+// Collectives
+template <typename T, ROC_SHMEM_OP Op>
+__device__ void IPCContext::to_all(roc_shmem_team_t team, T *dest,
+				   const T *source, int nreduce) {
+  //to_all<T, Op>(dest, source, nreduce, pe_start, log_pe_stride, pe_size, pWrk,
+  //              p_sync);
+}
+
+template <typename T, ROC_SHMEM_OP Op>
+__device__ void IPCContext::to_all(T *dest, const T *source, int nreduce,
+				   int PE_start, int logPE_stride,
+				   int PE_size, T *pWrk,
+				   long *pSync) {  // NOLINT(runtime/int)
+}
+
+template <typename T>
+__device__ void IPCContext::internal_put_broadcast(
+    T *dst, const T *src, int nelems, int pe_root, int pe_start,
+    int log_pe_stride, int pe_size) {  // NOLINT(runtime/int)
+  if (my_pe == pe_root) {
+    int stride = 1 << log_pe_stride;
+    int finish = pe_start + stride * pe_size;
+    for (int i = pe_start; i < finish; i += stride) {
+      if (i != my_pe) {
+        put_nbi_wg(dst, src, nelems, i);
+      }
+    }
+  }
+}
+
+template <typename T>
+__device__ void IPCContext::internal_get_broadcast(
+  T *dst, const T *src, int nelems, int pe_root) {  // NOLINT(runtime/int)
+  if (my_pe != pe_root) {
+    get_wg(dst, src, nelems, pe_root);
+  }
+}
+
+template <typename T>
+__device__ void IPCContext::broadcast(roc_shmem_team_t team, T *dst,
+				      const T *src, int nelems, int pe_root) {
+  IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
+
+  /**
+   * Ensure that the stride is a multiple of 2 .
+   */
+  int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
+  int pe_start = team_obj->tinfo_wrt_world->pe_start;
+  int pe_size = team_obj->tinfo_wrt_world->size;
+  long *p_sync = team_obj->bcast_pSync;
+
+  // Passed pe_root is relative to team, convert to world root
+  int pe_root_world = team_obj->get_pe_in_world(pe_root);
+  broadcast<T>(dst, src, nelems, pe_root_world, pe_start, log_pe_stride,
+               pe_size, p_sync);
+}
+
+template <typename T>
+__device__ void IPCContext::broadcast(T *dst, const T *src, int nelems,
+				      int pe_root, int pe_start,
+				      int log_pe_stride, int pe_size,
+				      long *p_sync) {  // NOLINT(runtime/int)
+  if (num_pes < 4) {
+    internal_put_broadcast(dst, src, nelems, pe_root, pe_start, log_pe_stride,
+                           pe_size);
+  } else {
+    internal_get_broadcast(dst, src, nelems, pe_root);
+  }
+
+  // Synchronize on completion of broadcast
+  internal_sync(my_pe, pe_start, (1 << log_pe_stride), pe_size, p_sync);
+}
+
+template <typename T>
+__device__ void IPCContext::alltoall(roc_shmem_team_t team, T *dst,
+				     const T *src, int nelems) {
+  alltoall_linear(team, dst, src, nelems);
+}
+
+template <typename T>
+__device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst,
+                                            const T *src, int nelems) {
+  IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
+
+  /**
+   * Ensure that the stride is a multiple of 2
+   */
+  int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
+  int pe_start = team_obj->tinfo_wrt_world->pe_start;
+  int pe_size = team_obj->num_pes;
+  int stride = 1 << log_pe_stride;
+  long *pSync = team_obj->alltoall_pSync;
+  int my_pe_in_team = team_obj->my_pe;
+
+  // Have each PE put their designated data to the other PEs
+  for (int j = 0; j < pe_size; j++) {
+    int dest_pe = team_obj->get_pe_in_world(j);
+    put_nbi_wg(&dst[my_pe_in_team * nelems], &src[j * nelems], nelems, dest_pe);
+  }
+  if (is_thread_zero_in_block()) {
+    quiet();
+  }
+  // wait until everyone has obtained their designated data
+  internal_sync(my_pe, pe_start, stride, pe_size, pSync);
+}
+
+template <typename T>
+__device__ void IPCContext::fcollect(roc_shmem_team_t team, T *dst,
+				     const T *src, int nelems) {
+  fcollect_linear(team, dst, src, nelems);
+}
+
+template <typename T>
+__device__ void IPCContext::fcollect_linear(roc_shmem_team_t team, T *dst,
+                                            const T *src, int nelems) {
+  IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
+
+  /**
+   * Ensure that the stride is a multiple of 2.
+   */
+  int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
+  int pe_start = team_obj->tinfo_wrt_world->pe_start;
+  int pe_size = team_obj->num_pes;
+  int stride = 1 << log_pe_stride;
+  long *pSync = team_obj->alltoall_pSync;
+  int my_pe_in_team = team_obj->my_pe;
+
+  // Have each PE put their designated data to the other PEs
+  for (int j = 0; j < pe_size; j++) {
+    int dest_pe = team_obj->get_pe_in_world(j);
+    put_nbi_wg(&dst[my_pe_in_team * nelems], src, nelems, dest_pe);
+  }
+
+  if (is_thread_zero_in_block()) {
+    quiet();
+  }
+  // wait until everyone has obtained their designated data
+  internal_sync(my_pe, pe_start, stride, pe_size, pSync);
+}
+
+// Block/wave functions
+template <typename T>
+__device__ void IPCContext::put_wg(T *dest, const T *source, size_t nelems,
+				   int pe) {
+  putmem_wg(dest, source, nelems * sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::put_nbi_wg(T *dest, const T *source,
+				       size_t nelems, int pe) {
+  putmem_nbi_wg(dest, source, nelems * sizeof(T), pe);
+}
+
+  template <typename T>
+__device__ void IPCContext::put_wave(T *dest, const T *source, size_t nelems,
+				     int pe) {
+  putmem_wave(dest, source, nelems * sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::put_nbi_wave(T *dest, const T *source,
+					 size_t nelems, int pe) {
+  putmem_nbi_wave(dest, source, nelems * sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::get_wg(T *dest, const T *source, size_t nelems,
+				   int pe) {
+  getmem_wg(dest, source, nelems * sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::get_nbi_wg(T *dest, const T *source,
+				       size_t nelems, int pe) {
+  getmem_nbi_wg(dest, source, nelems * sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::get_wave(T *dest, const T *source, size_t nelems,
+				     int pe) {
+  getmem_wave(dest, source, nelems * sizeof(T), pe);
+}
+
+template <typename T>
+__device__ void IPCContext::get_nbi_wave(T *dest, const T *source,
+					 size_t nelems, int pe) {
+  getmem_nbi_wave(dest, source, nelems * sizeof(T), pe);
+}
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_

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_

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_

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_

src/ipc/ipc_team.cpp

@@ -0,0 +1,56 @@
+/******************************************************************************
+ * 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 "ipc_team.hpp"
+
+#include "../backend_type.hpp"
+#include "backend_ipc.hpp"
+
+namespace rocshmem {
+
+IPCTeam::IPCTeam(Backend *backend, TeamInfo *team_info_parent,
+                     TeamInfo *team_info_world, int num_pes, int my_pe,
+                     MPI_Comm mpi_comm, int pool_index)
+    : Team(backend, team_info_parent, team_info_world, num_pes, my_pe,
+           mpi_comm) {
+  type = BackendType::IPC_BACKEND;
+  const IPCBackend *b = static_cast<const IPCBackend *>(backend);
+
+  pool_index_ = pool_index;
+
+  barrier_pSync =
+      &(b->barrier_pSync_pool[pool_index * ROC_SHMEM_BARRIER_SYNC_SIZE]);
+  reduce_pSync =
+      &(b->reduce_pSync_pool[pool_index * ROC_SHMEM_REDUCE_SYNC_SIZE]);
+  bcast_pSync = &(b->bcast_pSync_pool[pool_index * ROC_SHMEM_BCAST_SYNC_SIZE]);
+  alltoall_pSync =
+      &(b->alltoall_pSync_pool[pool_index * ROC_SHMEM_ALLTOALL_SYNC_SIZE]);
+
+  pWrk = reinterpret_cast<char *>(b->pWrk_pool) +
+         ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index;
+  pAta = reinterpret_cast<char *>(b->pAta_pool) +
+         ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index;
+}
+
+IPCTeam::~IPCTeam() {}
+
+}  // namespace rocshmem

src/ipc/ipc_team.hpp

@@ -0,0 +1,50 @@
+/******************************************************************************
+ * 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_TEAM_HPP_
+#define LIBRARY_SRC_IPC_TEAM_HPP_
+
+#include "../team.hpp"
+
+namespace rocshmem {
+
+class IPCTeam : public Team {
+ public:
+  IPCTeam(Backend* handle, TeamInfo* team_info_wrt_parent,
+            TeamInfo* team_info_wrt_world, int num_pes, int my_pe,
+            MPI_Comm team_comm, int pool_index);
+
+  virtual ~IPCTeam();
+
+  long* barrier_pSync{nullptr};
+  long* reduce_pSync{nullptr};
+  long* bcast_pSync{nullptr};
+  long* alltoall_pSync{nullptr};
+  void* pWrk{nullptr};
+  void* pAta{nullptr};
+
+  int pool_index_{-1};
+};
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_TEAM_HPP_

src/ipc_policy.cpp

@@ -50,7 +50,6 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases,
   /*
    * Figure out how this process' rank among local processes.
    */
-  int shm_rank;
   MPI_Comm_rank(shmcomm, &shm_rank);
 
   /*
@@ -92,7 +91,6 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases,
       void **ipc_base_uncast = reinterpret_cast<void **>(&ipc_base[i]);
       CHECK_HIP(hipIpcOpenMemHandle(ipc_base_uncast, vec_ipc_handle[i],
                                     hipIpcMemLazyEnablePeerAccess));
-      // TODO(bpotter): add some error checking here if happens to fail
     } else {
       ipc_base[i] = base_heap;
     }
@@ -110,6 +108,15 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases,
   free(vec_ipc_handle);
 }
 
+__host__ void IpcOnImpl::ipcHostStop() {
+  for (size_t i = 0; i < shm_size; i++) {
+    if (i != shm_rank) {
+      CHECK_HIP(hipIpcCloseMemHandle(ipc_bases[i]));
+    }
+  }
+  CHECK_HIP(hipFree(ipc_bases));
+}
+
 __device__ void IpcOnImpl::ipcCopy(void *dst, void *src, size_t size) {
   memcpy(dst, src, size);
 }

src/ipc_policy.hpp

@@ -42,6 +42,8 @@ class IpcOnImpl {
   using HEAP_BASES_T = std::vector<char *, StdAllocatorHIP<char *>>;
 
  public:
+  int shm_rank{0};
+
   uint32_t shm_size{0};
 
   char **ipc_bases{nullptr};
@@ -49,6 +51,8 @@ class IpcOnImpl {
   __host__ void ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases,
                             MPI_Comm thread_comm);
 
+  __host__ void ipcHostStop();
+
   __device__ bool isIpcAvailable(int my_pe, int target_pe) {
     return my_pe / shm_size == target_pe / shm_size;
   }
@@ -115,6 +119,8 @@ class IpcOffImpl {
   __host__ void ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases,
                             MPI_Comm thread_comm) {}
 
+  __host__ void ipcHostStop() {}
+
   __device__ bool isIpcAvailable(int my_pe, int target_pe) { return false; }
 
   __device__ void ipcGpuInit(Backend *roc_shmem_handle, Context *ctx,

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) {

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
 
 /******************************************************************************

src/team.cpp

@@ -44,6 +44,10 @@ ROTeam* get_internal_ro_team(roc_shmem_team_t team) {
   return reinterpret_cast<ROTeam*>(team);
 }
 
+IPCTeam* get_internal_ipc_team(roc_shmem_team_t team) {
+  return reinterpret_cast<IPCTeam*>(team);
+}
+
 __host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe,
                                           roc_shmem_team_t dst_team) {
   if (src_team == ROC_SHMEM_TEAM_INVALID ||

src/team.hpp

@@ -34,6 +34,7 @@ class Backend;
 class Team;
 class ROTeam;
 class GPUIBTeam;
+class IPCTeam;
 
 class TeamInfo {
  public:
@@ -162,6 +163,8 @@ GPUIBTeam* get_internal_gpu_ib_team(roc_shmem_team_t team);
 
 ROTeam* get_internal_ro_team(roc_shmem_team_t team);
 
+IPCTeam* get_internal_ipc_team(roc_shmem_team_t team);
+
 __host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe,
                                           roc_shmem_team_t dst_team);
 

src/util.hpp

@@ -149,7 +149,7 @@ __device__ void gpu_dprintf(const char* fmt, const Args&... args) {
       while (atomicCAS(print_lock, 0, 1) == 1) {
       }
 
-      printf("WG (%lu, %lu, %lu) TH (%lu, %lu, %lu) ", hipBlockIdx_x,
+      printf("WG (%u, %u, %u) TH (%u, %u, %u) ", hipBlockIdx_x,
              hipBlockIdx_y, hipBlockIdx_z, hipThreadIdx_x, hipThreadIdx_y,
              hipThreadIdx_z);
       printf(fmt, args...);
@@ -216,19 +216,34 @@ __device__ __forceinline__ void memcpy_wg(void* dst, void* src, size_t size) {
 }
 
 __device__ __forceinline__ void memcpy_wave(void* dst, void* src, size_t size) {
-  uint8_t* dst_bytes{static_cast<uint8_t*>(dst)};
-  uint8_t* src_bytes{static_cast<uint8_t*>(src)};
+  int wave_tid = get_flat_block_id() % WF_SIZE;
+  int wave_size{wave_SZ()};
 
   int cpy_size{};
-  int thread_id{get_flat_block_id()};
+  uint8_t* dst_bytes{nullptr};
+  uint8_t* dst_def{nullptr};
+  uint8_t* src_bytes{nullptr};
+  uint8_t* src_def{nullptr};
+
+  dst_def = reinterpret_cast<uint8_t*>(dst);
+  src_def = reinterpret_cast<uint8_t*>(src);
+  dst_bytes = dst_def;
+  src_bytes = src_def;
+
   for (int j{8}; j > 1; j >>= 1) {
     cpy_size = size / j;
-    for (int i{thread_id}; i < cpy_size; i += WF_SIZE) {
-      store_asm(src_bytes, dst_bytes, j);
+    for (int i{wave_tid}; i < cpy_size; i += wave_size) {
+      dst_bytes = dst_def;
+      src_bytes = src_def;
+
       src_bytes += i * j;
       dst_bytes += i * j;
-      size -= cpy_size * j;
+
+      store_asm(src_bytes, dst_bytes, j);
     }
+    size -= cpy_size * j;
+    dst_def += cpy_size * j;
+    src_def += cpy_size * j;
   }
 
   if (size == 1) {

tests/unit_tests/CMakeLists.txt

@@ -71,6 +71,7 @@ target_sources(
   PRIVATE
     shmem_gtest.cpp
     heap_memory_gtest.cpp
+    hipmalloc_gtest.cpp
     bin_gtest.cpp
     binner_gtest.cpp
     #bitwise_gtest.cpp # Test is disabled becasue of compilation errors
@@ -87,6 +88,8 @@ target_sources(
     notifier_gtest.cpp
     #forward_list_gtest.cpp
     free_list_gtest.cpp
+    #context_ipc_gtest.cpp
+    ipc_impl_simple_coarse_gtest.cpp
 )
 
 ###############################################################################

tests/unit_tests/context_ipc_gtest.cpp

@@ -0,0 +1,31 @@
+/******************************************************************************
+ * 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_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(ContextIpcTestFixture, constructor) {
+  /* do nothing for the moment, I *think* the 
+  ** constructor is invoked automatically 
+  */
+}

tests/unit_tests/context_ipc_gtest.hpp

@@ -0,0 +1,46 @@
+/******************************************************************************
+ * 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 ROCSHMEM_CONTEXT_IPC_GTEST_HPP
+#define ROCSHMEM_CONTEXT_IPC_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/ipc/context_ipc_device.hpp"
+#include "../src/ipc/backend_ipc.hpp"
+
+namespace rocshmem {
+
+class ContextIpcTestFixture : public ::testing::Test
+{
+  protected:
+    /**
+     * @brief Context Ipc Test
+     */
+   IPCBackend be{MPI_COMM_WORLD};
+
+   IPCContext ipc_context_ {&be};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_CONTEXT_IPC_GTEST_HPP

tests/unit_tests/hipmalloc_gtest.cpp

@@ -0,0 +1,43 @@
+/******************************************************************************
+ * 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 "hipmalloc_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(HipMallocTestFixture, normal_1GBx256) {
+    void* ptr{nullptr};
+    size_t gb {1073741824};
+    for (int i{0}; i < 256; i++) {
+        hip_allocator_.allocate(&ptr, gb);
+        hip_allocator_.deallocate(ptr);
+    }
+}
+
+TEST_F(HipMallocTestFixture, fine_1GBx256) {
+    void* ptr{nullptr};
+    size_t gb {1073741824};
+    for (int i{0}; i < 256; i++) {
+        hip_allocator_fg_.allocate(&ptr, gb);
+        hip_allocator_fg_.deallocate(ptr);
+    }
+}

tests/unit_tests/hipmalloc_gtest.hpp

@@ -0,0 +1,41 @@
+/******************************************************************************
+ * 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 ROCSHMEM_HIPMALLOC_GTEST_HPP
+#define ROCSHMEM_HIPMALLOC_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/symmetric_heap.hpp"
+#include "../src/util.hpp"
+
+namespace rocshmem {
+
+class HipMallocTestFixture : public ::testing::Test {
+  public:
+    HIPAllocator hip_allocator_ {};
+    HIPAllocatorFinegrained hip_allocator_fg_ {};
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_HIPMALLOC_GTEST_HPP

tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp

@@ -0,0 +1,236 @@
+/******************************************************************************
+ * 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 ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP
+#define ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include <numeric>
+
+#include <mpi.h>
+#include "../src/memory/symmetric_heap.hpp"
+#include "../src/ipc_policy.hpp"
+
+namespace rocshmem {
+
+__global__
+void
+kernel_simple_coarse_copy(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) {
+    if (!threadIdx.x) {
+      ipc_impl->ipcCopy(dest, src, bytes);
+      ipc_impl->ipcFence();
+    }
+    __syncthreads();
+}
+
+__global__
+void
+kernel_simple_coarse_copy_wg(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) {
+    ipc_impl->ipcCopy_wg(dest, src, bytes);
+    ipc_impl->ipcFence();
+    __syncthreads();
+}
+
+__global__
+void
+kernel_simple_coarse_copy_wave(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) {
+    ipc_impl->ipcCopy_wave(dest, src, bytes);
+    ipc_impl->ipcFence();
+    __syncthreads();
+}
+
+class IPCImplSimpleCoarseTestFixture : public ::testing::Test {
+
+    using HEAP_T = HeapMemory<HIPAllocator>;
+
+    using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
+
+    using FN_T = void (*)(IpcImpl*, int*, int*, size_t);
+
+  public:
+    IPCImplSimpleCoarseTestFixture() {
+        ipc_impl_.ipcHostInit(mpi_.my_pe(), mpi_.get_heap_bases() , MPI_COMM_WORLD);
+
+        assert(ipc_impl_dptr_ == nullptr);
+        hip_allocator_.allocate((void**)&ipc_impl_dptr_, sizeof(IpcImpl));
+
+        CHECK_HIP(hipMemcpy(ipc_impl_dptr_, &ipc_impl_,
+                            sizeof(IpcImpl), hipMemcpyHostToDevice));
+    }
+
+    ~IPCImplSimpleCoarseTestFixture() {
+        if (ipc_impl_dptr_) {
+            hip_allocator_.deallocate(ipc_impl_dptr_);
+        }
+
+        ipc_impl_.ipcHostStop();
+    }
+
+    void launch(FN_T f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes) {
+        f<<<grid, block>>>(ipc_impl_dptr_, src, dest, bytes);
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    enum TestType {
+        READ = 0,
+        WRITE = 1
+    };
+
+    void write(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(WRITE);
+        copy(WRITE, grid, block);
+        validate_dest_buffer(WRITE);
+    }
+
+    void write_wg(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(WRITE);
+        copy_wg(WRITE, grid, block);
+        validate_dest_buffer(WRITE);
+    }
+
+    void write_wave(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(WRITE);
+        copy_wave(WRITE, grid, block);
+        validate_dest_buffer(WRITE);
+    }
+
+    void read(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(READ);
+        copy(READ, grid, block);
+        validate_dest_buffer(READ);
+    }
+
+    void read_wg(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(READ);
+        copy_wg(READ, grid, block);
+        validate_dest_buffer(READ);
+    }
+
+    void read_wave(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(READ);
+        copy_wave(READ, grid, block);
+        validate_dest_buffer(READ);
+    }
+
+    void iota_golden(size_t elems) {
+        golden_.resize(elems);
+        std::iota(golden_.begin(), golden_.end(), 0);
+    }
+
+    void validate_golden(size_t elems) {
+        ASSERT_EQ(golden_.size(), elems);
+        for (int i{0}; i < golden_.size(); i++) {
+            ASSERT_EQ(golden_[i], i);
+        }
+    }
+
+    void initialize_src_buffer(TestType test) {
+        if (!pe_initializes_src_buffer(test)) {
+            return;
+        }
+        size_t bytes = golden_.size() * sizeof(int);
+        auto dev_src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_.my_pe()]);
+        CHECK_HIP(hipMemcpy(dev_src, golden_.data(), bytes, hipMemcpyHostToDevice));
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    bool pe_initializes_src_buffer(TestType test) {
+        bool is_write_test = test;
+        bool is_read_test = !test;
+        return (is_write_test && mpi_.my_pe() == 0) ||
+               (is_read_test && mpi_.my_pe() == 1);
+    }
+
+    void execute(TestType test, FN_T fn, const dim3 grid, const dim3 block) {
+        if (mpi_.my_pe()) {
+            mpi_.barrier();
+            mpi_.barrier();
+            return;
+        }
+        int *src{nullptr};
+        int *dest{nullptr};
+        if (test == WRITE) {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+        } else {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+        }
+        size_t bytes = golden_.size() * sizeof(int);
+        mpi_.barrier();
+        launch(fn, grid, block, src, dest, bytes);
+        mpi_.barrier();
+    }
+
+    void copy(TestType test, dim3 grid, dim3 block) {
+        execute(test, kernel_simple_coarse_copy, grid, block);
+    }
+
+    void copy_wg(TestType test, dim3 grid, dim3 block) {
+        execute(test, kernel_simple_coarse_copy_wg, grid, block);
+    }
+
+    void copy_wave(TestType test, dim3 grid, dim3 block) {
+        execute(test, kernel_simple_coarse_copy_wave, grid, block);
+    }
+
+    void validate_dest_buffer(TestType test) {
+        if (!pe_validates_dest_buffer(test)) {
+            return;
+        }
+
+        auto dev_dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_.my_pe()]);
+        for (int i{0}; i < golden_.size(); i++) {
+            ASSERT_EQ(golden_[i], dev_dest[i]);
+        }
+    }
+
+    bool pe_validates_dest_buffer(TestType test) {
+        return !pe_initializes_src_buffer(test);
+    }
+
+  protected:
+    std::vector<int> golden_;
+
+    std::vector<int> output_;
+
+    HEAP_T heap_mem_ {};
+
+    MPI_T mpi_ {heap_mem_.get_ptr(), heap_mem_.get_size()};
+
+    IpcImpl ipc_impl_ {};
+
+    IpcImpl *ipc_impl_dptr_ {nullptr};
+
+    HIPAllocator hip_allocator_ {};
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP