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 {