Merge pull request #27 from ROCm/ipc_bringup

Ipc bringup

Author: BKP

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,14 +30,14 @@ 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
     team.cpp
     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 {