From a13e2599d9a20fb2d5d8af7ea1733c7ae22c5043 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 14 Aug 2024 20:44:35 -0700 Subject: [PATCH 1/4] Code refactor move ipc_policy.hpp and ipc_policy.cpp files to src, since they are used by all the conduits. [ROCm/rocshmem commit: 24375a949e2668171bb82ac220472728d7973c46] --- projects/rocshmem/CMakeLists.txt | 3 +- projects/rocshmem/cmake/config.h.in | 1 + projects/rocshmem/src/CMakeLists.txt | 4 + projects/rocshmem/src/backend_bc.hpp | 2 +- projects/rocshmem/src/context.hpp | 2 +- projects/rocshmem/src/ipc/CMakeLists.txt | 2 +- projects/rocshmem/src/ipc_policy.cpp | 132 ++++++++++++++ projects/rocshmem/src/ipc_policy.hpp | 172 ++++++++++++++++++ .../src/reverse_offload/block_handle.hpp | 2 +- .../src/reverse_offload/queue_proxy.hpp | 2 +- .../ipc_impl_simple_coarse_gtest.hpp | 2 +- 11 files changed, 317 insertions(+), 7 deletions(-) create mode 100644 projects/rocshmem/src/ipc_policy.cpp create mode 100644 projects/rocshmem/src/ipc_policy.hpp diff --git a/projects/rocshmem/CMakeLists.txt b/projects/rocshmem/CMakeLists.txt index ee7f42fb12..d068315c62 100644 --- a/projects/rocshmem/CMakeLists.txt +++ b/projects/rocshmem/CMakeLists.txt @@ -54,7 +54,8 @@ string(REGEX REPLACE "([0-9]+)\.([0-9]+)\.([0-9]+)(.*)" "\\1.\\2.\\3" ROCSHMEM_V ############################################################################### option(DEBUG "Enable debug trace" OFF) option(PROFILE "Enable statistics and timing support" OFF) -option(USE_GPU_IB "Enable GPU_IB conduit. If off, RO_NET will be used" ON) +option(USE_GPU_IB "Enable GPU_IB conduit." ON) +option(USE_RO "Enable RO conduit." ON) option(USE_DC "Enable IB dynamically connected transport (DC)" OFF) option(USE_IPC "Enable IPC support (using HIP)" OFF) option(USE_THREADS "Enable workgroup threads to share network queues" OFF) diff --git a/projects/rocshmem/cmake/config.h.in b/projects/rocshmem/cmake/config.h.in index 6ec07c2853..3b2732137a 100644 --- a/projects/rocshmem/cmake/config.h.in +++ b/projects/rocshmem/cmake/config.h.in @@ -1,6 +1,7 @@ #cmakedefine DEBUG #cmakedefine PROFILE #cmakedefine USE_GPU_IB +#cmakedefine USE_RO #cmakedefine USE_DC #cmakedefine USE_IPC #cmakedefine USE_THREADS diff --git a/projects/rocshmem/src/CMakeLists.txt b/projects/rocshmem/src/CMakeLists.txt index 15162c2384..78b90fe035 100644 --- a/projects/rocshmem/src/CMakeLists.txt +++ b/projects/rocshmem/src/CMakeLists.txt @@ -23,6 +23,9 @@ ############################################################################### # ADD ROCSHMEM TARGET FOR FILES IN CURRENT DIRECTORY ############################################################################### +message(STATUS "Compiler flag USE_GPU_IB: ${USE_GPU_IB}") +message(STATUS "Compiler flag USE_RO: ${USE_RO}") +message(STATUS "Compiler flag USE_IPC: ${USE_IPC}") target_sources( ${PROJECT_NAME} PRIVATE @@ -37,6 +40,7 @@ target_sources( team_tracker.cpp util.cpp wf_coal_policy.cpp + ipc_policy.cpp ) target_compile_options( diff --git a/projects/rocshmem/src/backend_bc.hpp b/projects/rocshmem/src/backend_bc.hpp index b968c0f3bb..88cc6b3ba3 100644 --- a/projects/rocshmem/src/backend_bc.hpp +++ b/projects/rocshmem/src/backend_bc.hpp @@ -38,7 +38,7 @@ #include "config.h" // NOLINT(build/include_subdir) #include "roc_shmem/roc_shmem.hpp" #include "backend_type.hpp" -#include "ipc/ipc_policy.hpp" +#include "ipc_policy.hpp" #include "memory/symmetric_heap.hpp" #include "stats.hpp" #include "team_tracker.hpp" diff --git a/projects/rocshmem/src/context.hpp b/projects/rocshmem/src/context.hpp index 4ff1251220..72c8f38352 100644 --- a/projects/rocshmem/src/context.hpp +++ b/projects/rocshmem/src/context.hpp @@ -28,7 +28,7 @@ #include "backend_type.hpp" #include "fence_policy.hpp" #include "host/host.hpp" -#include "ipc/ipc_policy.hpp" +#include "ipc_policy.hpp" #include "stats.hpp" #include "sync/spin_ebo_block_mutex.hpp" #include "wf_coal_policy.hpp" diff --git a/projects/rocshmem/src/ipc/CMakeLists.txt b/projects/rocshmem/src/ipc/CMakeLists.txt index e34f62906e..1fd601d43d 100644 --- a/projects/rocshmem/src/ipc/CMakeLists.txt +++ b/projects/rocshmem/src/ipc/CMakeLists.txt @@ -26,6 +26,6 @@ target_sources( ${PROJECT_NAME} PRIVATE - ipc_policy.cpp + # ipc_policy.cpp context_ipc.cpp ) diff --git a/projects/rocshmem/src/ipc_policy.cpp b/projects/rocshmem/src/ipc_policy.cpp new file mode 100644 index 0000000000..2c3a96fa7c --- /dev/null +++ b/projects/rocshmem/src/ipc_policy.cpp @@ -0,0 +1,132 @@ +/****************************************************************************** + * 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_policy.hpp" + +#include + +#include "config.h" // NOLINT(build/include_subdir) +#include "backend_bc.hpp" +#include "context_incl.hpp" +#include "util.hpp" + +namespace rocshmem { + +__host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, + MPI_Comm thread_comm) { + /* + * Create an MPI communicator that deals only with local processes. + */ + MPI_Comm shmcomm; + MPI_Comm_split_type(thread_comm, MPI_COMM_TYPE_SHARED, 0, MPI_INFO_NULL, + &shmcomm); + + /* + * Figure out how many local process there are. + */ + int Shm_size; + MPI_Comm_size(shmcomm, &Shm_size); + shm_size = Shm_size; + + /* + * Figure out how this process' rank among local processes. + */ + MPI_Comm_rank(shmcomm, &shm_rank); + + /* + * Allocate a host-side c-array to hold the IPC handles. + */ + void *ipc_mem_handle_uncast = malloc(shm_size * sizeof(hipIpcMemHandle_t)); + hipIpcMemHandle_t *vec_ipc_handle = + reinterpret_cast(ipc_mem_handle_uncast); + + /* + * Call into the hip runtime to get an IPC handle for my symmetric + * heap and store that IPC handle into the host-side c-array which was + * just allocated. + */ + char *base_heap = heap_bases[my_pe]; + CHECK_HIP(hipIpcGetMemHandle(&vec_ipc_handle[shm_rank], base_heap)); + + /* + * Do an all-to-all exchange with each local processing element to + * share the symmetric heap IPC handles. + */ + MPI_Allgather(MPI_IN_PLACE, sizeof(hipIpcMemHandle_t), MPI_CHAR, + vec_ipc_handle, sizeof(hipIpcMemHandle_t), MPI_CHAR, shmcomm); + + /* + * Allocate device-side array to hold the IPC symmetric heap base + * addresses. + */ + char **ipc_base; + CHECK_HIP(hipMalloc(reinterpret_cast(&ipc_base), + shm_size * sizeof(char **))); + + /* + * For all local processing elements, initialize the device-side array + * with the IPC symmetric heap base addresses. + */ + for (size_t i = 0; i < shm_size; i++) { + if (i != shm_rank) { + void **ipc_base_uncast = reinterpret_cast(&ipc_base[i]); + CHECK_HIP(hipIpcOpenMemHandle(ipc_base_uncast, vec_ipc_handle[i], + hipIpcMemLazyEnablePeerAccess)); + } else { + ipc_base[i] = base_heap; + } + } + + /* + * Set member variables used by subsequent method calls. + */ + ipc_bases = ipc_base; + + /* + * Free the host-side memory used to exchange the symmetric heap base + * addresses. + */ + 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); +} + +__device__ void IpcOnImpl::ipcCopy_wave(void *dst, void *src, size_t size) { + memcpy_wave(dst, src, size); +} + +__device__ void IpcOnImpl::ipcCopy_wg(void *dst, void *src, size_t size) { + memcpy_wg(dst, src, size); +} + +} // namespace rocshmem diff --git a/projects/rocshmem/src/ipc_policy.hpp b/projects/rocshmem/src/ipc_policy.hpp new file mode 100644 index 0000000000..c0190198ca --- /dev/null +++ b/projects/rocshmem/src/ipc_policy.hpp @@ -0,0 +1,172 @@ +/****************************************************************************** + * 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_POLICY_HPP_ +#define LIBRARY_SRC_IPC_POLICY_HPP_ + +#include +#include + +#include +#include + +#include "config.h" // NOLINT(build/include_subdir) +#include "memory/hip_allocator.hpp" +#include "util.hpp" + +namespace rocshmem { + +class Backend; +class Context; + +class IpcOnImpl { + using HEAP_BASES_T = std::vector>; + + public: + int shm_rank{0}; + + uint32_t shm_size{0}; + + char **ipc_bases{nullptr}; + + __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; + } + __device__ void ipcGpuInit(Backend *gpu_backend, Context *ctx, int thread_id); + + __device__ void ipcCopy(void *dst, void *src, size_t size); + + __device__ void ipcCopy_wg(void *dst, void *src, size_t size); + + __device__ void ipcCopy_wave(void *dst, void *src, size_t size); + + __device__ void ipcFence() { __threadfence(); } + + template + __device__ T ipcAMOFetchAdd(T *val, T value) { + return __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); + } + + template + __device__ T ipcAMOFetchCas(T *val, T cond, T value) { + __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED, + __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); + return cond; + } + + template + __device__ void ipcAMOAdd(T *val, T value) { + __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); + } + + template + __device__ void ipcAMOCas(T *val, T cond, T value) { + __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED, + __ATOMIC_RELAXED, + __HIP_MEMORY_SCOPE_AGENT); + } + + template + __device__ void ipcAMOSet(T *val, T value) { + __hip_atomic_store(val, value, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + } + + __device__ void zero_byte_read(int pe) { + int local_pe = pe % shm_size; + uint32_t *pe_ipc_base = reinterpret_cast(ipc_bases[local_pe]); + volatile uint32_t read_value = __hip_atomic_load( + pe_ipc_base, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SYSTEM); + } +}; + +// clang-format off +NOWARN(-Wunused-parameter, +class IpcOffImpl { + using HEAP_BASES_T = std::vector>; + + public: + uint32_t shm_size{0}; + + char **ipc_bases{nullptr}; + + __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, + int thread_id) {} + + __device__ void ipcCopy(void *dst, void *src, size_t size) {} + + __device__ void ipcCopy_wg(void *dst, void *src, size_t size) {} + + __device__ void ipcCopy_wave(void *dst, void *src, size_t size) {} + + __device__ void ipcFence() {} + + template + __device__ T ipcAMOFetchAdd(T *val, T value) { + return T(); + } + + template + __device__ T ipcAMOFetchCas(T *val, T cond, T value) { + return T(); + } + + template + __device__ void ipcAMOAdd(T *val, T value) {} + + template + __device__ void ipcAMOSet(T *val, T value) {} + + template + __device__ void ipcAMOCas(T *val, T cond, T value) {} + + __device__ void zero_byte_read(int pe) {} +}; +) +// clang-format on + +/* + * Select which one of our IPC policies to use at compile time. + */ +#ifdef USE_IPC +typedef IpcOnImpl IpcImpl; +#else +typedef IpcOffImpl IpcImpl; +#endif + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_POLICY_HPP_ diff --git a/projects/rocshmem/src/reverse_offload/block_handle.hpp b/projects/rocshmem/src/reverse_offload/block_handle.hpp index f2ba1872b3..d2235f8639 100644 --- a/projects/rocshmem/src/reverse_offload/block_handle.hpp +++ b/projects/rocshmem/src/reverse_offload/block_handle.hpp @@ -24,7 +24,7 @@ #define LIBRARY_SRC_REVERSE_OFFLOAD_BLOCK_HANDLE_HPP_ #include "../hdp_policy.hpp" -#include "../ipc/ipc_policy.hpp" +#include "../ipc_policy.hpp" #include "profiler.hpp" #include "queue.hpp" diff --git a/projects/rocshmem/src/reverse_offload/queue_proxy.hpp b/projects/rocshmem/src/reverse_offload/queue_proxy.hpp index f3ddcec1ce..9cde0cb0cb 100644 --- a/projects/rocshmem/src/reverse_offload/queue_proxy.hpp +++ b/projects/rocshmem/src/reverse_offload/queue_proxy.hpp @@ -28,7 +28,7 @@ #include "../atomic_return.hpp" #include "../device_proxy.hpp" #include "../hdp_policy.hpp" -#include "../ipc/ipc_policy.hpp" +#include "../ipc_policy.hpp" #include "commands_types.hpp" #include "profiler.hpp" #include "../sync/abql_block_mutex.hpp" diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp index f31f406416..02dfd8c55a 100644 --- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp @@ -29,7 +29,7 @@ #include #include "../src/memory/symmetric_heap.hpp" -#include "../src/ipc/ipc_policy.hpp" +#include "../src/ipc_policy.hpp" namespace rocshmem { From 60fb77de7399c7a5ce36c4e81d2651c0e1e367d5 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 14 Aug 2024 22:59:02 -0700 Subject: [PATCH 2/4] Add IPC backend * add backend_ipc.{cpp & hpp} * rename context_ipc.{cpp & hpp} to context_ipc_device.{cpp & hpp} * add host interface to IPC backend * add context_ipc_host.{cpp & hpp} to support host interface * add USE_RO compile flag to enable support for single backend interface at a time * add ipc_single script to build rocSHMEM with IPC backend [ROCm/rocshmem commit: 49779863c266d068ba74988275ee985620c85391] --- .../rocshmem/scripts/build_configs/ipc_single | 30 +++ projects/rocshmem/src/CMakeLists.txt | 8 +- projects/rocshmem/src/backend_bc.cpp | 18 +- projects/rocshmem/src/backend_type.hpp | 32 ++- projects/rocshmem/src/context_incl.hpp | 5 +- projects/rocshmem/src/context_tmpl_device.hpp | 4 +- projects/rocshmem/src/context_tmpl_host.hpp | 4 +- projects/rocshmem/src/ipc/CMakeLists.txt | 5 +- projects/rocshmem/src/ipc/backend_ipc.cpp | 183 +++++++++++++ projects/rocshmem/src/ipc/backend_ipc.hpp | 190 ++++++++++++++ .../rocshmem/src/ipc/context_ipc_device.cpp | 129 ++++++++++ .../rocshmem/src/ipc/context_ipc_device.hpp | 241 ++++++++++++++++++ .../rocshmem/src/ipc/context_ipc_host.cpp | 85 ++++++ .../rocshmem/src/ipc/context_ipc_host.hpp | 149 +++++++++++ .../src/ipc/context_ipc_tmpl_device.hpp | 2 +- .../src/ipc/context_ipc_tmpl_host.hpp | 173 +++++++++++++ .../rocshmem/src/ipc/ipc_backend_proxy.hpp | 69 +++++ .../rocshmem/src/ipc/ipc_context_proxy.hpp | 90 +++++++ projects/rocshmem/src/roc_shmem.cpp | 19 +- projects/rocshmem/src/roc_shmem_gpu.cpp | 4 +- 20 files changed, 1410 insertions(+), 30 deletions(-) create mode 100755 projects/rocshmem/scripts/build_configs/ipc_single create mode 100644 projects/rocshmem/src/ipc/backend_ipc.cpp create mode 100644 projects/rocshmem/src/ipc/backend_ipc.hpp create mode 100644 projects/rocshmem/src/ipc/context_ipc_device.cpp create mode 100644 projects/rocshmem/src/ipc/context_ipc_device.hpp create mode 100644 projects/rocshmem/src/ipc/context_ipc_host.cpp create mode 100644 projects/rocshmem/src/ipc/context_ipc_host.hpp create mode 100644 projects/rocshmem/src/ipc/context_ipc_tmpl_host.hpp create mode 100644 projects/rocshmem/src/ipc/ipc_backend_proxy.hpp create mode 100644 projects/rocshmem/src/ipc/ipc_context_proxy.hpp diff --git a/projects/rocshmem/scripts/build_configs/ipc_single b/projects/rocshmem/scripts/build_configs/ipc_single new file mode 100755 index 0000000000..2013b2b67c --- /dev/null +++ b/projects/rocshmem/scripts/build_configs/ipc_single @@ -0,0 +1,30 @@ +#!/bin/bash +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +if [ -z $1 ] +then + install_path=~/rocshmem +else + install_path=$1 +fi + +src_path=$(dirname "$(realpath $0)")/../../ + +cmake \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_PREFIX=$install_path \ + -DCMAKE_VERBOSE_MAKEFILE=OFF \ + -DDEBUG=OFF \ + -DPROFILE=OFF \ + -DUSE_GPU_IB=OFF \ + -DUSE_RO=OFF \ + -DUSE_DC=OFF \ + -DUSE_IPC=ON \ + -DUSE_COHERENT_HEAP=ON \ + -DUSE_THREADS=OFF \ + -DUSE_WF_COAL=OFF \ + -DUSE_SINGLE_NODE=ON \ + -DUSE_HOST_SIDE_HDP_FLUSH=OFF\ + $src_path +cmake --build . --parallel 8 +cmake --install . \ No newline at end of file diff --git a/projects/rocshmem/src/CMakeLists.txt b/projects/rocshmem/src/CMakeLists.txt index 78b90fe035..5c31872765 100644 --- a/projects/rocshmem/src/CMakeLists.txt +++ b/projects/rocshmem/src/CMakeLists.txt @@ -23,9 +23,6 @@ ############################################################################### # ADD ROCSHMEM TARGET FOR FILES IN CURRENT DIRECTORY ############################################################################### -message(STATUS "Compiler flag USE_GPU_IB: ${USE_GPU_IB}") -message(STATUS "Compiler flag USE_RO: ${USE_RO}") -message(STATUS "Compiler flag USE_IPC: ${USE_IPC}") target_sources( ${PROJECT_NAME} PRIVATE @@ -63,11 +60,12 @@ target_compile_options( ############################################################################### IF (USE_GPU_IB) add_subdirectory(gpu_ib) -ELSE() +ELSEIF(USE_RO) add_subdirectory(reverse_offload) +ELSE() +add_subdirectory(ipc) ENDIF() add_subdirectory(containers) add_subdirectory(host) add_subdirectory(memory) add_subdirectory(sync) -add_subdirectory(ipc) diff --git a/projects/rocshmem/src/backend_bc.cpp b/projects/rocshmem/src/backend_bc.cpp index 1c3a619887..c55b1cd2d1 100644 --- a/projects/rocshmem/src/backend_bc.cpp +++ b/projects/rocshmem/src/backend_bc.cpp @@ -25,10 +25,12 @@ #include "backend_type.hpp" #include "context_incl.hpp" -#ifndef USE_GPU_IB +#ifdef USE_GPU_IB +#include "gpu_ib/backend_ib.hpp" +#elif defined(USE_RO) #include "reverse_offload/backend_ro.hpp" #else -#include "gpu_ib/backend_ib.hpp" +#include "ipc/backend_ipc.hpp" #endif namespace rocshmem { @@ -201,18 +203,22 @@ void Backend::reset_stats() { } __device__ bool Backend::create_ctx(int64_t option, roc_shmem_ctx_t* ctx) { -#ifndef USE_GPU_IB +#ifdef USE_GPU_IB + return static_cast(this)->create_ctx(option, ctx); +#elif defined(USE_RO) return static_cast(this)->create_ctx(option, ctx); #else - return static_cast(this)->create_ctx(option, ctx); + return static_cast(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(this)->destroy_ctx(ctx); +#elif defined(USE_RO) static_cast(this)->destroy_ctx(ctx); #else - static_cast(this)->destroy_ctx(ctx); + static_cast(this)->destroy_ctx(ctx); #endif } diff --git a/projects/rocshmem/src/backend_type.hpp b/projects/rocshmem/src/backend_type.hpp index ad9e992636..d8b2972504 100644 --- a/projects/rocshmem/src/backend_type.hpp +++ b/projects/rocshmem/src/backend_type.hpp @@ -44,7 +44,7 @@ namespace rocshmem { * @note Derived classes which use Backend as a base class must add * themselves to this enum class to support static polymorphism. */ -enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; +enum class BackendType { RO_BACKEND, GPU_IB_BACKEND, IPC_BACKEND }; /** * @brief Helper macro for some dispatch calls @@ -57,9 +57,12 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; #ifdef USE_GPU_IB #define DISPATCH(Func) \ static_cast(this)->Func; -#else +#elif defined(USE_RO) #define DISPATCH(Func) \ static_cast(this)->Func; +#else +#define DISPATCH(Func) \ + static_cast(this)->Func; #endif /** @@ -70,11 +73,16 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; auto ret_val{0}; \ ret_val = static_cast(this)->Func; \ return ret_val; -#else +#elif defined(USE_RO) #define DISPATCH_RET(Func) \ auto ret_val{0}; \ ret_val = static_cast(this)->Func; \ return ret_val; +#else +#define DISPATCH_RET(Func) \ + auto ret_val{0}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; #endif /** * @brief Device static dispatch method call with a return type of pointer. @@ -84,11 +92,16 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; void *ret_val{nullptr}; \ ret_val = static_cast(this)->Func; \ return ret_val; -#else +#elif defined(USE_RO) #define DISPATCH_RET_PTR(Func) \ void *ret_val{nullptr}; \ ret_val = static_cast(this)->Func; \ return ret_val; +#else +#define DISPATCH_RET_PTR(Func) \ + void *ret_val{nullptr}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; #endif /** @@ -100,8 +113,10 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; */ #ifdef USE_GPU_IB #define HOST_DISPATCH(Func) static_cast(this)->Func; -#else +#elif defined(USE_RO) #define HOST_DISPATCH(Func) static_cast(this)->Func; +#else +#define HOST_DISPATCH(Func) static_cast(this)->Func; #endif /** * @brief Host static dispatch method call with return value. @@ -116,11 +131,16 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; auto ret_val{0}; \ ret_val = static_cast(this)->Func; \ return ret_val; -#else +#elif defined(USE_RO) #define HOST_DISPATCH_RET(Func) \ auto ret_val{0}; \ ret_val = static_cast(this)->Func; \ return ret_val; +#else +#define HOST_DISPATCH_RET(Func) \ + auto ret_val{0}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; #endif } // namespace rocshmem diff --git a/projects/rocshmem/src/context_incl.hpp b/projects/rocshmem/src/context_incl.hpp index ea39a2ead5..f86ca3dfe8 100644 --- a/projects/rocshmem/src/context_incl.hpp +++ b/projects/rocshmem/src/context_incl.hpp @@ -29,9 +29,12 @@ #ifdef USE_GPU_IB #include "gpu_ib/context_ib_device.hpp" #include "gpu_ib/context_ib_host.hpp" -#else +#elif defined (USE_RO) #include "reverse_offload/context_ro_device.hpp" #include "reverse_offload/context_ro_host.hpp" +#else +#include "ipc/context_ipc_device.hpp" +#include "ipc/context_ipc_host.hpp" #endif #endif // LIBRARY_SRC_CONTEXT_INCL_HPP_ diff --git a/projects/rocshmem/src/context_tmpl_device.hpp b/projects/rocshmem/src/context_tmpl_device.hpp index 213fb90d3b..d6d0df63fa 100644 --- a/projects/rocshmem/src/context_tmpl_device.hpp +++ b/projects/rocshmem/src/context_tmpl_device.hpp @@ -27,8 +27,10 @@ #include "backend_type.hpp" #ifdef USE_GPU_IB #include "gpu_ib/context_ib_device.hpp" -#else +#elif defined(USE_RO) #include "reverse_offload/context_ro_device.hpp" +#else +#include "ipc/context_ipc_device.hpp" #endif namespace rocshmem { diff --git a/projects/rocshmem/src/context_tmpl_host.hpp b/projects/rocshmem/src/context_tmpl_host.hpp index 56b649e19d..bce99c1d6c 100644 --- a/projects/rocshmem/src/context_tmpl_host.hpp +++ b/projects/rocshmem/src/context_tmpl_host.hpp @@ -27,8 +27,10 @@ #include "backend_type.hpp" #ifdef USE_GPU_IB #include "gpu_ib/context_ib_host.hpp" -#else +#elif defined(USE_RO) #include "reverse_offload/context_ro_host.hpp" +#else +#include "ipc/context_ipc_host.hpp" #endif namespace rocshmem { diff --git a/projects/rocshmem/src/ipc/CMakeLists.txt b/projects/rocshmem/src/ipc/CMakeLists.txt index 1fd601d43d..b8d50d12e2 100644 --- a/projects/rocshmem/src/ipc/CMakeLists.txt +++ b/projects/rocshmem/src/ipc/CMakeLists.txt @@ -26,6 +26,7 @@ target_sources( ${PROJECT_NAME} PRIVATE - # ipc_policy.cpp - context_ipc.cpp + context_ipc_device.cpp + context_ipc_host.cpp + backend_ipc.cpp ) diff --git a/projects/rocshmem/src/ipc/backend_ipc.cpp b/projects/rocshmem/src/ipc/backend_ipc.cpp new file mode 100644 index 0000000000..bb1731add3 --- /dev/null +++ b/projects/rocshmem/src/ipc/backend_ipc.cpp @@ -0,0 +1,183 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "backend_ipc.hpp" + +namespace rocshmem { + +#define NET_CHECK(cmd) \ + { \ + if (cmd != MPI_SUCCESS) { \ + fprintf(stderr, "Unrecoverable error: MPI Failure\n"); \ + abort() ; \ + } \ + } + +extern roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT; + +IPCBackend::IPCBackend(MPI_Comm comm) + : Backend() { + type = BackendType::IPC_BACKEND; + + if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) { + std::stringstream sstream(maximum_num_contexts_str); + sstream >> maximum_num_contexts_; + } + + init_mpi_once(comm); + + initIPC(); + + auto *bp{ipc_backend_proxy.get()}; + + bp->heap_ptr = &heap; + + /* Initialize the host interface */ + host_interface = + new HostInterface(hdp_proxy_.get(), thread_comm, &heap); + //free host interface + + default_host_ctx = std::make_unique(this, 0); + + ROC_SHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx.get(); + + init_g_ret(&heap, thread_comm, MAX_NUM_BLOCKS, &bp->g_ret); + + allocate_atomic_region(&bp->atomic_ret, MAX_NUM_BLOCKS); + + default_context_proxy_ = IPCDefaultContextProxyT(this); + + setup_ctxs(); + +} + +IPCBackend::~IPCBackend() { + ipc_net_free_runtime(); + CHECK_HIP(hipFree(ctx_array)); +} + +void IPCBackend::setup_ctxs() { + CHECK_HIP(hipMalloc(&ctx_array, sizeof(IPCContext) * maximum_num_contexts_)); + for (int i = 0; i < maximum_num_contexts_; i++) { + new (&ctx_array[i]) IPCContext(this); + ctx_free_list.get()->push_back(ctx_array + i); + } +} + +__device__ bool IPCBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) { + IPCContext *ctx_{nullptr}; + + auto pop_result = ctx_free_list.get()->pop_front(); + if (!pop_result.success) { + return false; + } + ctx_ = pop_result.value; + + ctx->ctx_opaque = ctx_; + return true; +} + +__device__ void IPCBackend::destroy_ctx(roc_shmem_ctx_t *ctx) { + ctx_free_list.get()->push_back(static_cast(ctx->ctx_opaque)); +} + +void IPCBackend::init_mpi_once(MPI_Comm comm) { + int init_done{}; + NET_CHECK(MPI_Initialized(&init_done)); + + int provided{}; + if (!init_done) { + NET_CHECK(MPI_Init_thread(0, 0, MPI_THREAD_MULTIPLE, &provided)); + if (provided != MPI_THREAD_MULTIPLE) { + std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n"; + } + } + if (comm == MPI_COMM_NULL) comm = MPI_COMM_WORLD; + + NET_CHECK(MPI_Comm_dup(comm, &thread_comm)); + NET_CHECK(MPI_Comm_size(thread_comm, &num_pes)); + NET_CHECK(MPI_Comm_rank(thread_comm, &my_pe)); +} + +void IPCBackend::team_destroy(roc_shmem_team_t team) { + assert(false); +} + +void IPCBackend::create_new_team(Team *parent_team, + TeamInfo *team_info_wrt_parent, + TeamInfo *team_info_wrt_world, int num_pes, + int my_pe_in_new_team, MPI_Comm team_comm, + roc_shmem_team_t *new_team) { + assert(false); +} + +void IPCBackend::ctx_create(int64_t options, void **ctx) { + IPCHostContext *new_ctx{nullptr}; + new_ctx = new IPCHostContext(this, options); + *ctx = new_ctx; +} + +IPCHostContext *get_internal_ipc_net_ctx(Context *ctx) { + return reinterpret_cast(ctx); +} + +void IPCBackend::ctx_destroy(Context *ctx) { + IPCHostContext *ro_net_host_ctx{get_internal_ipc_net_ctx(ctx)}; + delete ro_net_host_ctx; +} + +void IPCBackend::reset_backend_stats() { + assert(false); +} + +void IPCBackend::dump_backend_stats() { + assert(false); +} + +void IPCBackend::initIPC() { + const auto &heap_bases{heap.get_heap_bases()}; + + ipcImpl.ipcHostInit(my_pe, heap_bases, + thread_comm); +} + +void IPCBackend::ipc_net_free_runtime() { + /* + * Validate that a handle was passed that is not a nullptr. + */ + auto *bp{ipc_backend_proxy.get()}; + assert(bp); + + /* + * Free the atomic_ret array. + */ + CHECK_HIP(hipFree(bp->atomic_ret->atomic_base_ptr)); + + // TODO(Avinash) Free g_ret +} + +void IPCBackend::global_exit(int status) { + assert(false); +} + + +} // namespace rocshmem \ No newline at end of file diff --git a/projects/rocshmem/src/ipc/backend_ipc.hpp b/projects/rocshmem/src/ipc/backend_ipc.hpp new file mode 100644 index 0000000000..1a89dd63f1 --- /dev/null +++ b/projects/rocshmem/src/ipc/backend_ipc.hpp @@ -0,0 +1,190 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_IPC_BACKEND_HPP_ +#define LIBRARY_SRC_IPC_BACKEND_HPP_ + +#include "../backend_bc.hpp" +#include "../containers/free_list_impl.hpp" +#include "../hdp_proxy.hpp" +#include "../memory/hip_allocator.hpp" +#include "ipc_backend_proxy.hpp" +#include "../context_incl.hpp" +#include "ipc_context_proxy.hpp" +#include "../ipc_policy.hpp" + +namespace rocshmem { + +class IPCBackend : public Backend { + const unsigned MAX_NUM_BLOCKS{65536}; + + public: + /** + * @copydoc Backend::Backend(unsigned) + */ + explicit IPCBackend(MPI_Comm comm); + + /** + * @copydoc Backend::~Backend() + */ + virtual ~IPCBackend(); + + __device__ bool create_ctx(int64_t options, roc_shmem_ctx_t *ctx); + + /** + * @brief Destroy a `roc_shmem_ctx_t` context and returns it back to the + * context free list. + */ + __device__ void destroy_ctx(roc_shmem_ctx_t *ctx); + + /** + * @copydoc Backend::ctx_create + */ + void ctx_create(int64_t options, void **ctx) override; + + /** + * @copydoc Backend::ctx_destroy + */ + void ctx_destroy(Context *ctx) override; + + /** + * @brief initialize MPI. + * + * IPC relies on MPI just to exchange the IPC_handle information. + * + * todo: remove the dependency on MPI and make it generic to PMI-X or just + * to OpenSHMEM to have support for both CPU and GPU + */ + void init_mpi_once(MPI_Comm comm); + + /** + * @brief Helper to initialize IPC interface. + */ + void initIPC(); + + /** + * @brief Allocation and initialization of backend contexts. + */ + void setup_ctxs(); + + /** + * @brief Free all resources associated with the backend. + * + * The memory allocated to the handle param is deallocated during this + * method. The handle should be treated as a nullptr after the call. + * + * The destructor treats this method as a helper function to destroy + * this object. + * + * @todo The method needs to be broken into smaller pieces and most + * of these internal resources need to be moved into subclasses using + * RAII. + */ + void ipc_net_free_runtime(); + + /** + * @brief Abort the application. + * + * @param[in] status Exit code. + * + * @return void. + * + * @note This routine terminates the entire application. + */ + void global_exit(int status) override; + + /** + * @copydoc Backend::create_new_team + */ + void create_new_team(Team *parent_team, TeamInfo *team_info_wrt_parent, + TeamInfo *team_info_wrt_world, int num_pes, + int my_pe_in_new_team, MPI_Comm team_comm, + roc_shmem_team_t *new_team) override; + + /** + * @copydoc Backend::team_destroy(roc_shmem_team_t) + */ + void team_destroy(roc_shmem_team_t team) override; + + /** + * @brief Handle to device memory fields. + */ + IPCBackendProxyT ipc_backend_proxy{}; + + /** + * @brief The host-facing interface that will be used + * by all contexts of the ROBackend + */ + HostInterface *host_interface{nullptr}; + + protected: + /** + * @copydoc Backend::dump_backend_stats() + */ + void dump_backend_stats() override; + + /** + * @copydoc Backend::reset_backend_stats() + */ + void reset_backend_stats() override; + + /** + * @brief Allocates uncacheable host memory for the hdp policy. + * + * @note Internal data ownership is managed by the proxy + */ + HdpProxy hdp_proxy_{}; + + /** + * @brief Holds a copy of the default context for host functions + */ + std::unique_ptr default_host_ctx{nullptr}; + + private: + /** + * @brief Proxy for the default context + * + * @note Internal data ownership is managed by the proxy + */ + IPCDefaultContextProxyT default_context_proxy_; // init handled in constructor + + /** + * @brief An array of @ref ROContexts that backs the context FreeList. + */ + IPCContext *ctx_array{nullptr}; + + /** + * @brief A free-list containing contexts. + */ + FreeListProxy ctx_free_list{}; + + /** + * @brief Holds maximum number of contexts used in library + */ + size_t maximum_num_contexts_{1024}; + + +}; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_BACKEND_HPP_ \ No newline at end of file diff --git a/projects/rocshmem/src/ipc/context_ipc_device.cpp b/projects/rocshmem/src/ipc/context_ipc_device.cpp new file mode 100644 index 0000000000..97937fcd1e --- /dev/null +++ b/projects/rocshmem/src/ipc/context_ipc_device.cpp @@ -0,0 +1,129 @@ +/****************************************************************************** + * 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 +#include +#include + +#include +#include + +#include "config.h" // NOLINT(build/include_subdir) +#include "roc_shmem/roc_shmem.hpp" + +namespace rocshmem { + +__host__ IPCContext::IPCContext(Backend *b) + : Context(b, false) { +} + +__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) { +} + +__device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, + int pe) { +} + +__device__ void IPCContext::putmem_nbi(void *dest, const void *source, + size_t nelems, int pe) { +} + +__device__ void IPCContext::getmem_nbi(void *dest, const void *source, + size_t nelems, int 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::barrier_all() { + __syncthreads(); +} + +__device__ void IPCContext::sync_all() { + __syncthreads(); +} + +__device__ void IPCContext::sync(roc_shmem_team_t team) { + __syncthreads(); +} + +__device__ void IPCContext::putmem_wg(void *dest, const void *source, + size_t nelems, int pe) { + __syncthreads(); +} + +__device__ void IPCContext::getmem_wg(void *dest, const void *source, + size_t nelems, int pe) { + __syncthreads(); +} + +__device__ void IPCContext::putmem_nbi_wg(void *dest, const void *source, + size_t nelems, int pe) { + __syncthreads(); +} + +__device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, + size_t nelems, int pe) { + __syncthreads(); +} + +__device__ void IPCContext::putmem_wave(void *dest, const void *source, + size_t nelems, int pe) { +} + +__device__ void IPCContext::getmem_wave(void *dest, const void *source, + size_t nelems, int pe) { +} + +__device__ void IPCContext::putmem_nbi_wave(void *dest, const void *source, + size_t nelems, int pe) { +} + +__device__ void IPCContext::getmem_nbi_wave(void *dest, const void *source, + size_t nelems, int pe) { +} + +} // namespace rocshmem diff --git a/projects/rocshmem/src/ipc/context_ipc_device.hpp b/projects/rocshmem/src/ipc/context_ipc_device.hpp new file mode 100644 index 0000000000..492d63ca04 --- /dev/null +++ b/projects/rocshmem/src/ipc/context_ipc_device.hpp @@ -0,0 +1,241 @@ +/****************************************************************************** + * 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 + __device__ void p(T *dest, T value, int pe); + + template + __device__ void put(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_nbi(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ T g(const T *source, int pe); + + template + __device__ void get(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_nbi(T *dest, const T *source, size_t nelems, int pe); + + // Atomic operations + template + __device__ void amo_add(void *dst, T value, int pe); + + template + __device__ void amo_set(void *dst, T value, int pe); + + template + __device__ T amo_swap(void *dst, T value, int pe); + + template + __device__ T amo_fetch_and(void *dst, T value, int pe); + + template + __device__ void amo_and(void *dst, T value, int pe); + + template + __device__ T amo_fetch_or(void *dst, T value, int pe); + + template + __device__ void amo_or(void *dst, T value, int pe); + + template + __device__ T amo_fetch_xor(void *dst, T value, int pe); + + template + __device__ void amo_xor(void *dst, T value, int pe); + + template + __device__ void amo_cas(void *dst, T value, T cond, int pe); + + template + __device__ T amo_fetch_add(void *dst, T value, int pe); + + template + __device__ T amo_fetch_cas(void *dst, T value, T cond, int pe); + + // Collectives + template + __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 + __device__ void to_all(roc_shmem_team_t team, T *dest, const T *source, + int nreduce); + + template + __device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source, + int nelems, int pe_root); + + template + __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 + __device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source, + int nelems); + template + __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 + __device__ void put_wg(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_nbi_wg(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_wave(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_nbi_wave(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_wg(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_nbi_wg(T *dest, const T *source, size_t nelems, int pe); + + + template + __device__ void get_wave(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe); + + // Wait / Test functions + template + __device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val); + + template + __device__ void wait_until_all(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __device__ size_t wait_until_any(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __device__ size_t wait_until_some(T* ptr, size_t nelems, + size_t* indices, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __device__ void wait_until_all_vector(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T* vals); + + template + __device__ size_t wait_until_any_vector(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T* vals); + template + __device__ size_t wait_until_some_vector(T* ptr, size_t nelems, + size_t* indices, + const int *status, + roc_shmem_cmps cmp, T* vals); + + template + __device__ int test(T* ptr, roc_shmem_cmps cmp, T val); + + private: + +}; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_GPU_IB_CONTEXT_IB_DEVICE_HPP_ diff --git a/projects/rocshmem/src/ipc/context_ipc_host.cpp b/projects/rocshmem/src/ipc/context_ipc_host.cpp new file mode 100644 index 0000000000..0d3464f33d --- /dev/null +++ b/projects/rocshmem/src/ipc/context_ipc_host.cpp @@ -0,0 +1,85 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "context_ipc_host.hpp" + +#include + +#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(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 diff --git a/projects/rocshmem/src/ipc/context_ipc_host.hpp b/projects/rocshmem/src/ipc/context_ipc_host.hpp new file mode 100644 index 0000000000..f9421ca210 --- /dev/null +++ b/projects/rocshmem/src/ipc/context_ipc_host.hpp @@ -0,0 +1,149 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_IPC_CONTEXT_HOST_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_HOST_HPP_ + +#include "../context.hpp" + +namespace rocshmem { + +class IPCHostContext : public Context { + public: + __host__ IPCHostContext(Backend *b, int64_t options); + + __host__ ~IPCHostContext(); + + template + __host__ void p(T *dest, T value, int pe); + + template + __host__ T g(const T *source, int pe); + + template + __host__ void put(T *dest, const T *source, size_t nelems, int pe); + + template + __host__ void get(T *dest, const T *source, size_t nelems, int pe); + + template + __host__ void put_nbi(T *dest, const T *source, size_t nelems, int pe); + + template + __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 + __host__ void amo_add(void *dst, T value, int pe); + + template + __host__ void amo_cas(void *dst, T value, T cond, int pe); + + template + __host__ T amo_fetch_add(void *dst, T value, int pe); + + template + __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 + __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 + __host__ void broadcast(roc_shmem_team_t team, T *dest, const T *source, + int nelems, int pe_root); + + template + __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 + __host__ void to_all(roc_shmem_team_t team, T *dest, const T *source, + int nreduce); + + template + __host__ void wait_until(T *ptr, roc_shmem_cmps cmp, T val); + + template + __host__ size_t wait_until_any(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __host__ void wait_until_all(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __host__ size_t wait_until_some(T* ptr, size_t nelems, + size_t* indices, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __host__ void wait_until_all_vector(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T* vals); + + template + __host__ size_t wait_until_any_vector(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T* vals); + + template + __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 + __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_ diff --git a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp index c90e0a3d00..28cd718f0a 100644 --- a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp @@ -25,7 +25,7 @@ #include "config.h" // NOLINT(build/include_subdir) #include "roc_shmem/roc_shmem.hpp" -#include "context_ipc.hpp" +#include "context_ipc_device.hpp" #include "../util.hpp" namespace rocshmem { diff --git a/projects/rocshmem/src/ipc/context_ipc_tmpl_host.hpp b/projects/rocshmem/src/ipc/context_ipc_tmpl_host.hpp new file mode 100644 index 0000000000..76484de34d --- /dev/null +++ b/projects/rocshmem/src/ipc/context_ipc_tmpl_host.hpp @@ -0,0 +1,173 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_ + +#include "config.h" // NOLINT(build/include_subdir) +#include "../host/host_templates.hpp" + +namespace rocshmem { + +template +__host__ void IPCHostContext::p(T *dest, T value, int pe) { + host_interface->p(dest, value, pe, context_window_info); +} + +template +__host__ T IPCHostContext::g(const T *source, int pe) { + return host_interface->g(source, pe, context_window_info); +} + +template +__host__ void IPCHostContext::put(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->put(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::get(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->get(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::put_nbi(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->put_nbi(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::get_nbi(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->get_nbi(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::amo_add(void *dst, T value, int pe) { + host_interface->amo_add(dst, value, pe, context_window_info); +} + +template +__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 +__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 +__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 +__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(dest, source, nelems, pe_root, pe_start, + log_pe_stride, pe_size, p_sync); +} + +template +__host__ void IPCHostContext::broadcast(roc_shmem_team_t team, T *dest, + const T *source, int nelems, + int pe_root) { + host_interface->broadcast(team, dest, source, nelems, pe_root); +} + +template +__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(dest, source, nreduce, pe_start, log_pe_stride, + pe_size, p_wrk, p_sync); +} + +template +__host__ void IPCHostContext::to_all(roc_shmem_team_t team, T *dest, + const T *source, int nreduce) { + host_interface->to_all(team, dest, source, nreduce); +} + +template +__host__ void IPCHostContext::wait_until(T *ptr, roc_shmem_cmps cmp, T val) { + host_interface->wait_until(ptr, cmp, val, context_window_info); +} + +template +__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(ptr, nelems, status, cmp, val, context_window_info); +} + +template +__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(ptr, nelems, status, cmp, val, context_window_info); +} + +template +__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(ptr, nelems, indices, status, cmp, val, context_window_info); +} + +template +__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(ptr, nelems, status, cmp, vals, context_window_info); +} + +template +__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(ptr, nelems, status, cmp, vals, context_window_info); +} + +template +__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(ptr, nelems, indices, status, cmp, vals, context_window_info); +} + +template +__host__ int IPCHostContext::test(T *ptr, roc_shmem_cmps cmp, T val) { + return host_interface->test(ptr, cmp, val, context_window_info); +} + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_ diff --git a/projects/rocshmem/src/ipc/ipc_backend_proxy.hpp b/projects/rocshmem/src/ipc/ipc_backend_proxy.hpp new file mode 100644 index 0000000000..0077fb053d --- /dev/null +++ b/projects/rocshmem/src/ipc/ipc_backend_proxy.hpp @@ -0,0 +1,69 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_ +#define LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_ + +#include "../device_proxy.hpp" +#include "../atomic_return.hpp" + +namespace rocshmem { + +struct IPCBackendRegister { + char *g_ret{nullptr}; + atomic_ret_t *atomic_ret{nullptr}; + SymmetricHeap *heap_ptr{nullptr}; +}; + +template +class IPCBackendProxy { + using ProxyT = DeviceProxy; + + 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; + +} // namespace rocshmem + +#endif // #define LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_ \ No newline at end of file diff --git a/projects/rocshmem/src/ipc/ipc_context_proxy.hpp b/projects/rocshmem/src/ipc/ipc_context_proxy.hpp new file mode 100644 index 0000000000..867f199094 --- /dev/null +++ b/projects/rocshmem/src/ipc/ipc_context_proxy.hpp @@ -0,0 +1,90 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_ + + +#include "../device_proxy.hpp" +#include "backend_ipc.hpp" + +namespace rocshmem { + +class IPCBackend; + +template +class IPCDefaultContextProxy { + using ProxyT = DeviceProxy; + + 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)); + 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; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_ \ No newline at end of file diff --git a/projects/rocshmem/src/roc_shmem.cpp b/projects/rocshmem/src/roc_shmem.cpp index 8c61d6c864..5ca9051f37 100644 --- a/projects/rocshmem/src/roc_shmem.cpp +++ b/projects/rocshmem/src/roc_shmem.cpp @@ -39,12 +39,14 @@ #ifdef USE_GPU_IB #include "gpu_ib/backend_ib.hpp" #include "gpu_ib/context_ib_tmpl_host.hpp" -#endif -#include "mpi_init_singleton.hpp" -#ifndef USE_GPU_IB +#elif defined(USE_RO) #include "reverse_offload/backend_ro.hpp" #include "reverse_offload/context_ro_tmpl_host.hpp" +#else +#include "ipc/backend_ipc.hpp" +#include "ipc/context_ipc_tmpl_host.hpp" #endif +#include "mpi_init_singleton.hpp" #include "team.hpp" #include "templates_host.hpp" #include "util.hpp" @@ -82,12 +84,17 @@ 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) + printf("RO Backend\n"); CHECK_HIP(hipHostMalloc(&backend, sizeof(ROBackend))); backend = new (backend) ROBackend(comm); #else - CHECK_HIP(hipHostMalloc(&backend, sizeof(GPUIBBackend))); - backend = new (backend) GPUIBBackend(comm); + printf("IPC Backend\n"); + CHECK_HIP(hipHostMalloc(&backend, sizeof(IPCBackend))); + backend = new (backend) IPCBackend(comm); #endif if (!backend) { diff --git a/projects/rocshmem/src/roc_shmem_gpu.cpp b/projects/rocshmem/src/roc_shmem_gpu.cpp index 269a98a28a..086a6deed1 100644 --- a/projects/rocshmem/src/roc_shmem_gpu.cpp +++ b/projects/rocshmem/src/roc_shmem_gpu.cpp @@ -51,8 +51,10 @@ #ifdef USE_GPU_IB #include "gpu_ib/context_ib_tmpl_device.hpp" -#else +#elif defined(USE_RO) #include "reverse_offload/context_ro_tmpl_device.hpp" +#else +#include "ipc/context_ipc_tmpl_device.hpp" #endif /****************************************************************************** From d06394fa2cf9307036fa984105d0e3179619158c Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 15 Aug 2024 08:52:06 -0700 Subject: [PATCH 3/4] remove ipc_policy.{hpp & cpp} and context_ipc.{hpp & cpp} * move ipc_policy.{hpp & cpp} to src * rename context_ipc.{hpp & cpp} to context_ipc_device.{hpp & cpp} [ROCm/rocshmem commit: b68867ee17a5edbab35090bd3dbb0b9b2d90e0ce] --- projects/rocshmem/src/ipc/context_ipc.cpp | 129 ------------ projects/rocshmem/src/ipc/context_ipc.hpp | 241 ---------------------- projects/rocshmem/src/ipc/ipc_policy.cpp | 132 ------------ projects/rocshmem/src/ipc/ipc_policy.hpp | 172 --------------- 4 files changed, 674 deletions(-) delete mode 100644 projects/rocshmem/src/ipc/context_ipc.cpp delete mode 100644 projects/rocshmem/src/ipc/context_ipc.hpp delete mode 100644 projects/rocshmem/src/ipc/ipc_policy.cpp delete mode 100644 projects/rocshmem/src/ipc/ipc_policy.hpp diff --git a/projects/rocshmem/src/ipc/context_ipc.cpp b/projects/rocshmem/src/ipc/context_ipc.cpp deleted file mode 100644 index a825a135e9..0000000000 --- a/projects/rocshmem/src/ipc/context_ipc.cpp +++ /dev/null @@ -1,129 +0,0 @@ -/****************************************************************************** - * 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.hpp" -#include "context_ipc_tmpl_device.hpp" - -#include -#include -#include - -#include -#include - -#include "config.h" // NOLINT(build/include_subdir) -#include "roc_shmem/roc_shmem.hpp" - -namespace rocshmem { - -__host__ IPCContext::IPCContext(Backend *b) - : Context(b, false) { -} - -__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) { -} - -__device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, - int pe) { -} - -__device__ void IPCContext::putmem_nbi(void *dest, const void *source, - size_t nelems, int pe) { -} - -__device__ void IPCContext::getmem_nbi(void *dest, const void *source, - size_t nelems, int 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::barrier_all() { - __syncthreads(); -} - -__device__ void IPCContext::sync_all() { - __syncthreads(); -} - -__device__ void IPCContext::sync(roc_shmem_team_t team) { - __syncthreads(); -} - -__device__ void IPCContext::putmem_wg(void *dest, const void *source, - size_t nelems, int pe) { - __syncthreads(); -} - -__device__ void IPCContext::getmem_wg(void *dest, const void *source, - size_t nelems, int pe) { - __syncthreads(); -} - -__device__ void IPCContext::putmem_nbi_wg(void *dest, const void *source, - size_t nelems, int pe) { - __syncthreads(); -} - -__device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, - size_t nelems, int pe) { - __syncthreads(); -} - -__device__ void IPCContext::putmem_wave(void *dest, const void *source, - size_t nelems, int pe) { -} - -__device__ void IPCContext::getmem_wave(void *dest, const void *source, - size_t nelems, int pe) { -} - -__device__ void IPCContext::putmem_nbi_wave(void *dest, const void *source, - size_t nelems, int pe) { -} - -__device__ void IPCContext::getmem_nbi_wave(void *dest, const void *source, - size_t nelems, int pe) { -} - -} // namespace rocshmem diff --git a/projects/rocshmem/src/ipc/context_ipc.hpp b/projects/rocshmem/src/ipc/context_ipc.hpp deleted file mode 100644 index 492d63ca04..0000000000 --- a/projects/rocshmem/src/ipc/context_ipc.hpp +++ /dev/null @@ -1,241 +0,0 @@ -/****************************************************************************** - * 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 - __device__ void p(T *dest, T value, int pe); - - template - __device__ void put(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void put_nbi(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ T g(const T *source, int pe); - - template - __device__ void get(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void get_nbi(T *dest, const T *source, size_t nelems, int pe); - - // Atomic operations - template - __device__ void amo_add(void *dst, T value, int pe); - - template - __device__ void amo_set(void *dst, T value, int pe); - - template - __device__ T amo_swap(void *dst, T value, int pe); - - template - __device__ T amo_fetch_and(void *dst, T value, int pe); - - template - __device__ void amo_and(void *dst, T value, int pe); - - template - __device__ T amo_fetch_or(void *dst, T value, int pe); - - template - __device__ void amo_or(void *dst, T value, int pe); - - template - __device__ T amo_fetch_xor(void *dst, T value, int pe); - - template - __device__ void amo_xor(void *dst, T value, int pe); - - template - __device__ void amo_cas(void *dst, T value, T cond, int pe); - - template - __device__ T amo_fetch_add(void *dst, T value, int pe); - - template - __device__ T amo_fetch_cas(void *dst, T value, T cond, int pe); - - // Collectives - template - __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 - __device__ void to_all(roc_shmem_team_t team, T *dest, const T *source, - int nreduce); - - template - __device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source, - int nelems, int pe_root); - - template - __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 - __device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source, - int nelems); - template - __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 - __device__ void put_wg(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void put_nbi_wg(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void put_wave(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void put_nbi_wave(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void get_wg(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void get_nbi_wg(T *dest, const T *source, size_t nelems, int pe); - - - template - __device__ void get_wave(T *dest, const T *source, size_t nelems, int pe); - - template - __device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe); - - // Wait / Test functions - template - __device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val); - - template - __device__ void wait_until_all(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ size_t wait_until_any(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ size_t wait_until_some(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ void wait_until_all_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals); - - template - __device__ size_t wait_until_any_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals); - template - __device__ size_t wait_until_some_vector(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T* vals); - - template - __device__ int test(T* ptr, roc_shmem_cmps cmp, T val); - - private: - -}; - -} // namespace rocshmem - -#endif // LIBRARY_SRC_GPU_IB_CONTEXT_IB_DEVICE_HPP_ diff --git a/projects/rocshmem/src/ipc/ipc_policy.cpp b/projects/rocshmem/src/ipc/ipc_policy.cpp deleted file mode 100644 index 356aa81ad8..0000000000 --- a/projects/rocshmem/src/ipc/ipc_policy.cpp +++ /dev/null @@ -1,132 +0,0 @@ -/****************************************************************************** - * 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_policy.hpp" - -#include - -#include "config.h" // NOLINT(build/include_subdir) -#include "../backend_bc.hpp" -#include "../context_incl.hpp" -#include "../util.hpp" - -namespace rocshmem { - -__host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, - MPI_Comm thread_comm) { - /* - * Create an MPI communicator that deals only with local processes. - */ - MPI_Comm shmcomm; - MPI_Comm_split_type(thread_comm, MPI_COMM_TYPE_SHARED, 0, MPI_INFO_NULL, - &shmcomm); - - /* - * Figure out how many local process there are. - */ - int Shm_size; - MPI_Comm_size(shmcomm, &Shm_size); - shm_size = Shm_size; - - /* - * Figure out how this process' rank among local processes. - */ - MPI_Comm_rank(shmcomm, &shm_rank); - - /* - * Allocate a host-side c-array to hold the IPC handles. - */ - void *ipc_mem_handle_uncast = malloc(shm_size * sizeof(hipIpcMemHandle_t)); - hipIpcMemHandle_t *vec_ipc_handle = - reinterpret_cast(ipc_mem_handle_uncast); - - /* - * Call into the hip runtime to get an IPC handle for my symmetric - * heap and store that IPC handle into the host-side c-array which was - * just allocated. - */ - char *base_heap = heap_bases[my_pe]; - CHECK_HIP(hipIpcGetMemHandle(&vec_ipc_handle[shm_rank], base_heap)); - - /* - * Do an all-to-all exchange with each local processing element to - * share the symmetric heap IPC handles. - */ - MPI_Allgather(MPI_IN_PLACE, sizeof(hipIpcMemHandle_t), MPI_CHAR, - vec_ipc_handle, sizeof(hipIpcMemHandle_t), MPI_CHAR, shmcomm); - - /* - * Allocate device-side array to hold the IPC symmetric heap base - * addresses. - */ - char **ipc_base; - CHECK_HIP(hipMalloc(reinterpret_cast(&ipc_base), - shm_size * sizeof(char **))); - - /* - * For all local processing elements, initialize the device-side array - * with the IPC symmetric heap base addresses. - */ - for (size_t i = 0; i < shm_size; i++) { - if (i != shm_rank) { - void **ipc_base_uncast = reinterpret_cast(&ipc_base[i]); - CHECK_HIP(hipIpcOpenMemHandle(ipc_base_uncast, vec_ipc_handle[i], - hipIpcMemLazyEnablePeerAccess)); - } else { - ipc_base[i] = base_heap; - } - } - - /* - * Set member variables used by subsequent method calls. - */ - ipc_bases = ipc_base; - - /* - * Free the host-side memory used to exchange the symmetric heap base - * addresses. - */ - 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); -} - -__device__ void IpcOnImpl::ipcCopy_wave(void *dst, void *src, size_t size) { - memcpy_wave(dst, src, size); -} - -__device__ void IpcOnImpl::ipcCopy_wg(void *dst, void *src, size_t size) { - memcpy_wg(dst, src, size); -} - -} // namespace rocshmem diff --git a/projects/rocshmem/src/ipc/ipc_policy.hpp b/projects/rocshmem/src/ipc/ipc_policy.hpp deleted file mode 100644 index 9609aa42cb..0000000000 --- a/projects/rocshmem/src/ipc/ipc_policy.hpp +++ /dev/null @@ -1,172 +0,0 @@ -/****************************************************************************** - * 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_POLICY_HPP_ -#define LIBRARY_SRC_IPC_POLICY_HPP_ - -#include -#include - -#include -#include - -#include "config.h" // NOLINT(build/include_subdir) -#include "../memory/hip_allocator.hpp" -#include "../util.hpp" - -namespace rocshmem { - -class Backend; -class Context; - -class IpcOnImpl { - using HEAP_BASES_T = std::vector>; - - public: - int shm_rank{0}; - - uint32_t shm_size{0}; - - char **ipc_bases{nullptr}; - - __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; - } - __device__ void ipcGpuInit(Backend *gpu_backend, Context *ctx, int thread_id); - - __device__ void ipcCopy(void *dst, void *src, size_t size); - - __device__ void ipcCopy_wg(void *dst, void *src, size_t size); - - __device__ void ipcCopy_wave(void *dst, void *src, size_t size); - - __device__ void ipcFence() { __threadfence(); } - - template - __device__ T ipcAMOFetchAdd(T *val, T value) { - return __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); - } - - template - __device__ T ipcAMOFetchCas(T *val, T cond, T value) { - __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED, - __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); - return cond; - } - - template - __device__ void ipcAMOAdd(T *val, T value) { - __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); - } - - template - __device__ void ipcAMOCas(T *val, T cond, T value) { - __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED, - __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); - } - - template - __device__ void ipcAMOSet(T *val, T value) { - __hip_atomic_store(val, value, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); - } - - __device__ void zero_byte_read(int pe) { - int local_pe = pe % shm_size; - uint32_t *pe_ipc_base = reinterpret_cast(ipc_bases[local_pe]); - volatile uint32_t read_value = __hip_atomic_load( - pe_ipc_base, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SYSTEM); - } -}; - -// clang-format off -NOWARN(-Wunused-parameter, -class IpcOffImpl { - using HEAP_BASES_T = std::vector>; - - public: - uint32_t shm_size{0}; - - char **ipc_bases{nullptr}; - - __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, - int thread_id) {} - - __device__ void ipcCopy(void *dst, void *src, size_t size) {} - - __device__ void ipcCopy_wg(void *dst, void *src, size_t size) {} - - __device__ void ipcCopy_wave(void *dst, void *src, size_t size) {} - - __device__ void ipcFence() {} - - template - __device__ T ipcAMOFetchAdd(T *val, T value) { - return T(); - } - - template - __device__ T ipcAMOFetchCas(T *val, T cond, T value) { - return T(); - } - - template - __device__ void ipcAMOAdd(T *val, T value) {} - - template - __device__ void ipcAMOSet(T *val, T value) {} - - template - __device__ void ipcAMOCas(T *val, T cond, T value) {} - - __device__ void zero_byte_read(int pe) {} -}; -) -// clang-format on - -/* - * Select which one of our IPC policies to use at compile time. - */ -#ifdef USE_IPC -typedef IpcOnImpl IpcImpl; -#else -typedef IpcOffImpl IpcImpl; -#endif - -} // namespace rocshmem - -#endif // LIBRARY_SRC_IPC_POLICY_HPP_ From fe866aaea61bd5f86e8e6c576c71968236f8d7e7 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 15 Aug 2024 11:54:56 -0700 Subject: [PATCH 4/4] Update context_ipc_gtest.cpp to use IPCbackend [ROCm/rocshmem commit: c9dbcf80c20986fa60274effc43ce9e7957fe653] --- projects/rocshmem/src/roc_shmem.cpp | 2 -- projects/rocshmem/tests/unit_tests/CMakeLists.txt | 2 +- projects/rocshmem/tests/unit_tests/context_ipc_gtest.hpp | 6 +++--- 3 files changed, 4 insertions(+), 6 deletions(-) diff --git a/projects/rocshmem/src/roc_shmem.cpp b/projects/rocshmem/src/roc_shmem.cpp index 5ca9051f37..32ef91bbda 100644 --- a/projects/rocshmem/src/roc_shmem.cpp +++ b/projects/rocshmem/src/roc_shmem.cpp @@ -88,11 +88,9 @@ roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT; CHECK_HIP(hipHostMalloc(&backend, sizeof(GPUIBBackend))); backend = new (backend) GPUIBBackend(comm); #elif defined(USE_RO) - printf("RO Backend\n"); CHECK_HIP(hipHostMalloc(&backend, sizeof(ROBackend))); backend = new (backend) ROBackend(comm); #else - printf("IPC Backend\n"); CHECK_HIP(hipHostMalloc(&backend, sizeof(IPCBackend))); backend = new (backend) IPCBackend(comm); #endif diff --git a/projects/rocshmem/tests/unit_tests/CMakeLists.txt b/projects/rocshmem/tests/unit_tests/CMakeLists.txt index 3f1cf89a17..349bb2c2bf 100644 --- a/projects/rocshmem/tests/unit_tests/CMakeLists.txt +++ b/projects/rocshmem/tests/unit_tests/CMakeLists.txt @@ -88,7 +88,7 @@ target_sources( notifier_gtest.cpp #forward_list_gtest.cpp free_list_gtest.cpp - context_ipc_gtest.cpp + #context_ipc_gtest.cpp ipc_impl_simple_coarse_gtest.cpp ) diff --git a/projects/rocshmem/tests/unit_tests/context_ipc_gtest.hpp b/projects/rocshmem/tests/unit_tests/context_ipc_gtest.hpp index fc39504958..24a80fcedb 100644 --- a/projects/rocshmem/tests/unit_tests/context_ipc_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/context_ipc_gtest.hpp @@ -25,8 +25,8 @@ #include "gtest/gtest.h" -#include "../src/ipc/context_ipc.hpp" -#include "../src/reverse_offload/backend_ro.hpp" +#include "../src/ipc/context_ipc_device.hpp" +#include "../src/ipc/backend_ipc.hpp" namespace rocshmem { @@ -36,7 +36,7 @@ class ContextIpcTestFixture : public ::testing::Test /** * @brief Context Ipc Test */ - ROBackend be{MPI_COMM_WORLD}; + IPCBackend be{MPI_COMM_WORLD}; IPCContext ipc_context_ {&be}; };