From d06394fa2cf9307036fa984105d0e3179619158c Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 15 Aug 2024 08:52:06 -0700 Subject: [PATCH] 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_