From 1183006e20620df7ba097f11925e540d3bcafd03 Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Thu, 25 Jul 2024 07:33:41 -0700 Subject: [PATCH 01/20] src/ipc: IPC folder refactor mv ipc_policy.{hpp,cpp} into a separate folder as a start for the standalone IPC conduit. Unit tests and functional tests pass on the developmpent system. --- src/CMakeLists.txt | 2 +- src/backend_bc.hpp | 2 +- src/context.hpp | 2 +- src/ipc/CMakeLists.txt | 30 ++++++++++++++++++++++++++++ src/{ => ipc}/ipc_policy.cpp | 6 +++--- src/{ => ipc}/ipc_policy.hpp | 4 ++-- src/reverse_offload/block_handle.hpp | 2 +- src/reverse_offload/queue_proxy.hpp | 2 +- 8 files changed, 40 insertions(+), 10 deletions(-) create mode 100644 src/ipc/CMakeLists.txt rename src/{ => ipc}/ipc_policy.cpp (98%) rename src/{ => ipc}/ipc_policy.hpp (98%) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 823102d8a1..15162c2384 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -30,7 +30,6 @@ target_sources( backend_bc.cpp context_host.cpp context_device.cpp - ipc_policy.cpp mpi_init_singleton.cpp roc_shmem_gpu.cpp roc_shmem.cpp @@ -67,3 +66,4 @@ add_subdirectory(containers) add_subdirectory(host) add_subdirectory(memory) add_subdirectory(sync) +add_subdirectory(ipc) diff --git a/src/backend_bc.hpp b/src/backend_bc.hpp index 88cc6b3ba3..b968c0f3bb 100644 --- a/src/backend_bc.hpp +++ b/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_policy.hpp" +#include "ipc/ipc_policy.hpp" #include "memory/symmetric_heap.hpp" #include "stats.hpp" #include "team_tracker.hpp" diff --git a/src/context.hpp b/src/context.hpp index 72c8f38352..4ff1251220 100644 --- a/src/context.hpp +++ b/src/context.hpp @@ -28,7 +28,7 @@ #include "backend_type.hpp" #include "fence_policy.hpp" #include "host/host.hpp" -#include "ipc_policy.hpp" +#include "ipc/ipc_policy.hpp" #include "stats.hpp" #include "sync/spin_ebo_block_mutex.hpp" #include "wf_coal_policy.hpp" diff --git a/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt new file mode 100644 index 0000000000..730b733c63 --- /dev/null +++ b/src/ipc/CMakeLists.txt @@ -0,0 +1,30 @@ +############################################################################### +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to +# deal in the Software without restriction, including without limitation the +# rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +# sell copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +# FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS +# IN THE SOFTWARE. +############################################################################### + +############################################################################### +# ADD ROCSHMEM TARGET FOR FILES IN CURRENT DIRECTORY +############################################################################### +target_sources( + ${PROJECT_NAME} + PRIVATE + ipc_policy.cpp +) diff --git a/src/ipc_policy.cpp b/src/ipc/ipc_policy.cpp similarity index 98% rename from src/ipc_policy.cpp rename to src/ipc/ipc_policy.cpp index 4bf483735d..e4555ac8f2 100644 --- a/src/ipc_policy.cpp +++ b/src/ipc/ipc_policy.cpp @@ -25,9 +25,9 @@ #include #include "config.h" // NOLINT(build/include_subdir) -#include "backend_bc.hpp" -#include "context_incl.hpp" -#include "util.hpp" +#include "../backend_bc.hpp" +#include "../context_incl.hpp" +#include "../util.hpp" namespace rocshmem { diff --git a/src/ipc_policy.hpp b/src/ipc/ipc_policy.hpp similarity index 98% rename from src/ipc_policy.hpp rename to src/ipc/ipc_policy.hpp index f6bb49d409..d608dc8d54 100644 --- a/src/ipc_policy.hpp +++ b/src/ipc/ipc_policy.hpp @@ -30,8 +30,8 @@ #include #include "config.h" // NOLINT(build/include_subdir) -#include "memory/hip_allocator.hpp" -#include "util.hpp" +#include "../memory/hip_allocator.hpp" +#include "../util.hpp" namespace rocshmem { diff --git a/src/reverse_offload/block_handle.hpp b/src/reverse_offload/block_handle.hpp index d2235f8639..f2ba1872b3 100644 --- a/src/reverse_offload/block_handle.hpp +++ b/src/reverse_offload/block_handle.hpp @@ -24,7 +24,7 @@ #define LIBRARY_SRC_REVERSE_OFFLOAD_BLOCK_HANDLE_HPP_ #include "../hdp_policy.hpp" -#include "../ipc_policy.hpp" +#include "../ipc/ipc_policy.hpp" #include "profiler.hpp" #include "queue.hpp" diff --git a/src/reverse_offload/queue_proxy.hpp b/src/reverse_offload/queue_proxy.hpp index 9cde0cb0cb..f3ddcec1ce 100644 --- a/src/reverse_offload/queue_proxy.hpp +++ b/src/reverse_offload/queue_proxy.hpp @@ -28,7 +28,7 @@ #include "../atomic_return.hpp" #include "../device_proxy.hpp" #include "../hdp_policy.hpp" -#include "../ipc_policy.hpp" +#include "../ipc/ipc_policy.hpp" #include "commands_types.hpp" #include "profiler.hpp" #include "../sync/abql_block_mutex.hpp" From de750ddacc7ac6e8ce23a417373691d92fbe9202 Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Thu, 25 Jul 2024 14:48:51 -0700 Subject: [PATCH 02/20] src/ipc: add context_ipc barebone code --- src/ipc/CMakeLists.txt | 1 + src/ipc/context_ipc.cpp | 129 +++++++++++++ src/ipc/context_ipc.hpp | 241 ++++++++++++++++++++++++ src/ipc/context_ipc_tmpl_device.hpp | 280 ++++++++++++++++++++++++++++ 4 files changed, 651 insertions(+) create mode 100644 src/ipc/context_ipc.cpp create mode 100644 src/ipc/context_ipc.hpp create mode 100644 src/ipc/context_ipc_tmpl_device.hpp diff --git a/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt index 730b733c63..e34f62906e 100644 --- a/src/ipc/CMakeLists.txt +++ b/src/ipc/CMakeLists.txt @@ -27,4 +27,5 @@ target_sources( ${PROJECT_NAME} PRIVATE ipc_policy.cpp + context_ipc.cpp ) diff --git a/src/ipc/context_ipc.cpp b/src/ipc/context_ipc.cpp new file mode 100644 index 0000000000..a825a135e9 --- /dev/null +++ b/src/ipc/context_ipc.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.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/src/ipc/context_ipc.hpp b/src/ipc/context_ipc.hpp new file mode 100644 index 0000000000..492d63ca04 --- /dev/null +++ b/src/ipc/context_ipc.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/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp new file mode 100644 index 0000000000..c90e0a3d00 --- /dev/null +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -0,0 +1,280 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ + +#include "config.h" // NOLINT(build/include_subdir) +#include "roc_shmem/roc_shmem.hpp" +#include "context_ipc.hpp" +#include "../util.hpp" + +namespace rocshmem { + +/****************************************************************************** + ************************** TEMPLATE SPECIALIZATIONS ************************** + *****************************************************************************/ +template +__device__ void IPCContext::p(T *dest, T value, int pe) { + putmem_nbi(dest, &value, sizeof(T), pe); +} + +template +__device__ void IPCContext::put(T *dest, const T *source, size_t nelems, + int pe) { + putmem(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::put_nbi(T *dest, const T *source, size_t nelems, + int pe) { + putmem_nbi(dest, source, sizeof(T) * nelems, pe); +} + +template +__device__ T IPCContext::g(const T *source, int pe) { + T ret; + return ret; +} + +template +__device__ void IPCContext::get(T *dest, const T *source, size_t nelems, + int pe) { + getmem(dest, source, sizeof(T) * nelems, pe); +} + +template +__device__ void IPCContext::get_nbi(T *dest, const T *source, size_t nelems, + int pe) { + getmem_nbi(dest, source, sizeof(T) * nelems, pe); +} + +// Atomics +template +__device__ void IPCContext::amo_add(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ void IPCContext::amo_set(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ T IPCContext::amo_swap(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ T IPCContext::amo_fetch_and(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ void IPCContext::amo_and(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ T IPCContext::amo_fetch_or(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ void IPCContext::amo_or(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ T IPCContext::amo_fetch_xor(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ void IPCContext::amo_xor(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ void IPCContext::amo_cas(void *dst, T value, T cond, int pe) { + assert(false); +} + +template +__device__ T IPCContext::amo_fetch_add(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ T IPCContext::amo_fetch_cas(void *dst, T value, T cond, int pe) { + assert(false); + return 0; +} + +// Collectives +template +__device__ void IPCContext::to_all(roc_shmem_team_t team, T *dest, + const T *source, int nreduce) { + //to_all(dest, source, nreduce, pe_start, log_pe_stride, pe_size, pWrk, + // p_sync); +} + +template +__device__ void IPCContext::to_all(T *dest, const T *source, int nreduce, + int PE_start, int logPE_stride, + int PE_size, T *pWrk, + long *pSync) { // NOLINT(runtime/int) +} + +template +__device__ void IPCContext::broadcast(roc_shmem_team_t team, T *dst, + const T *src, int nelems, int pe_root) { + //broadcast(dst, src, nelems, pe_root_world, pe_start, log_pe_stride, + // pe_size, p_sync); +} + +template +__device__ void IPCContext::broadcast(T *dst, const T *src, int nelems, + int pe_root, int pe_start, + int log_pe_stride, int pe_size, + long *p_sync) { // NOLINT(runtime/int) +} + +template +__device__ void IPCContext::alltoall(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { +} + +template +__device__ void IPCContext::fcollect(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { +} + +// Block/wave functions +template +__device__ void IPCContext::put_wg(T *dest, const T *source, size_t nelems, + int pe) { + putmem_wg(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::put_nbi_wg(T *dest, const T *source, + size_t nelems, int pe) { + putmem_nbi_wg(dest, source, nelems * sizeof(T), pe); +} + + template +__device__ void IPCContext::put_wave(T *dest, const T *source, size_t nelems, + int pe) { + putmem_wave(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::put_nbi_wave(T *dest, const T *source, + size_t nelems, int pe) { + putmem_nbi_wave(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_wg(T *dest, const T *source, size_t nelems, + int pe) { + getmem_wg(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_nbi_wg(T *dest, const T *source, + size_t nelems, int pe) { + getmem_nbi_wg(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_wave(T *dest, const T *source, size_t nelems, + int pe) { + getmem_wave(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_nbi_wave(T *dest, const T *source, + size_t nelems, int pe) { + getmem_nbi_wave(dest, source, nelems * sizeof(T), pe); +} + + +//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) { + return 0; +} + +template +__device__ size_t wait_until_some(T* ptr, size_t nelems, + size_t* indices, + const int *status, + roc_shmem_cmps cmp, T val){ + return 0; +} + +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){ + return 0; +} + +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) { + return 0; +} + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ From b1bc4a497a304c1f91f654dd3a2a19c0e0a6c2ca Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 30 Jul 2024 15:22:03 -0700 Subject: [PATCH 03/20] unit_tests: add ipc_context tests add the initial outline of an ipc_context unit test. The current test only invokes the ipc_context constructor, more tests will be added later as the class is being populated. Also, at the moment the unit test takes an ROBackend as an argument for the constructor, not sure whether this will be the final solution. --- tests/unit_tests/CMakeLists.txt | 1 + tests/unit_tests/context_ipc_gtest.cpp | 31 +++++++++++++++++ tests/unit_tests/context_ipc_gtest.hpp | 46 ++++++++++++++++++++++++++ 3 files changed, 78 insertions(+) create mode 100644 tests/unit_tests/context_ipc_gtest.cpp create mode 100644 tests/unit_tests/context_ipc_gtest.hpp diff --git a/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index 7130c2c516..1f71119fd6 100644 --- a/tests/unit_tests/CMakeLists.txt +++ b/tests/unit_tests/CMakeLists.txt @@ -87,6 +87,7 @@ target_sources( notifier_gtest.cpp #forward_list_gtest.cpp free_list_gtest.cpp + context_ipc_gtest.cpp ) ############################################################################### diff --git a/tests/unit_tests/context_ipc_gtest.cpp b/tests/unit_tests/context_ipc_gtest.cpp new file mode 100644 index 0000000000..c963a65673 --- /dev/null +++ b/tests/unit_tests/context_ipc_gtest.cpp @@ -0,0 +1,31 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "context_ipc_gtest.hpp" + +using namespace rocshmem; + +TEST_F(ContextIpcTestFixture, constructor) { + /* do nothing for the moment, I *think* the + ** constructor is invoked automatically + */ +} diff --git a/tests/unit_tests/context_ipc_gtest.hpp b/tests/unit_tests/context_ipc_gtest.hpp new file mode 100644 index 0000000000..fc39504958 --- /dev/null +++ b/tests/unit_tests/context_ipc_gtest.hpp @@ -0,0 +1,46 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef ROCSHMEM_CONTEXT_IPC_GTEST_HPP +#define ROCSHMEM_CONTEXT_IPC_GTEST_HPP + +#include "gtest/gtest.h" + +#include "../src/ipc/context_ipc.hpp" +#include "../src/reverse_offload/backend_ro.hpp" + +namespace rocshmem { + +class ContextIpcTestFixture : public ::testing::Test +{ + protected: + /** + * @brief Context Ipc Test + */ + ROBackend be{MPI_COMM_WORLD}; + + IPCContext ipc_context_ {&be}; +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_CONTEXT_IPC_GTEST_HPP From c865fcd4bb3d04b710f1ba17d0d412be0fcac5b6 Mon Sep 17 00:00:00 2001 From: Brandon Potter Date: Tue, 6 Aug 2024 12:08:35 -0700 Subject: [PATCH 04/20] Fix memcpy_wave bug --- src/util.hpp | 27 +++++++++++++++++++++------ 1 file changed, 21 insertions(+), 6 deletions(-) diff --git a/src/util.hpp b/src/util.hpp index 0d4dad806a..67fdac6786 100644 --- a/src/util.hpp +++ b/src/util.hpp @@ -302,19 +302,34 @@ __device__ __forceinline__ void memcpy_wg(void* dst, void* src, size_t size) { } __device__ __forceinline__ void memcpy_wave(void* dst, void* src, size_t size) { - uint8_t* dst_bytes{static_cast(dst)}; - uint8_t* src_bytes{static_cast(src)}; + int wave_tid = get_flat_block_id() % WF_SIZE; + int wave_size{wave_SZ()}; int cpy_size{}; - int thread_id{get_flat_block_id()}; + uint8_t* dst_bytes{nullptr}; + uint8_t* dst_def{nullptr}; + uint8_t* src_bytes{nullptr}; + uint8_t* src_def{nullptr}; + + dst_def = reinterpret_cast(dst); + src_def = reinterpret_cast(src); + dst_bytes = dst_def; + src_bytes = src_def; + for (int j{8}; j > 1; j >>= 1) { cpy_size = size / j; - for (int i{thread_id}; i < cpy_size; i += WF_SIZE) { - store_asm(src_bytes, dst_bytes, j); + for (int i{wave_tid}; i < cpy_size; i += wave_size) { + dst_bytes = dst_def; + src_bytes = src_def; + src_bytes += i * j; dst_bytes += i * j; - size -= cpy_size * j; + + store_asm(src_bytes, dst_bytes, j); } + size -= cpy_size * j; + dst_def += cpy_size * j; + src_def += cpy_size * j; } if (size == 1) { From 58c5a98b5d9fc11263e69547a567c89d51f7badc Mon Sep 17 00:00:00 2001 From: Brandon Potter Date: Tue, 6 Aug 2024 12:11:08 -0700 Subject: [PATCH 05/20] Add ipc unit_tests --- scripts/build_configs/ro_net | 2 +- src/ipc/ipc_policy.cpp | 11 +- src/ipc/ipc_policy.hpp | 6 + tests/unit_tests/CMakeLists.txt | 2 + tests/unit_tests/hipmalloc_gtest.cpp | 43 + tests/unit_tests/hipmalloc_gtest.hpp | 41 + .../ipc_impl_simple_coarse_gtest.cpp | 1029 +++++++++++++++++ .../ipc_impl_simple_coarse_gtest.hpp | 236 ++++ 8 files changed, 1367 insertions(+), 3 deletions(-) create mode 100644 tests/unit_tests/hipmalloc_gtest.cpp create mode 100644 tests/unit_tests/hipmalloc_gtest.hpp create mode 100644 tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp create mode 100644 tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp diff --git a/scripts/build_configs/ro_net b/scripts/build_configs/ro_net index 30c4ad27bc..17809fa0a9 100755 --- a/scripts/build_configs/ro_net +++ b/scripts/build_configs/ro_net @@ -18,7 +18,7 @@ cmake \ -DPROFILE=OFF \ -DUSE_GPU_IB=OFF \ -DUSE_DC=OFF \ - -DUSE_IPC=OFF \ + -DUSE_IPC=ON \ -DUSE_THREADS=ON \ -DUSE_WF_COAL=OFF \ -DUSE_COHERENT_HEAP=ON \ diff --git a/src/ipc/ipc_policy.cpp b/src/ipc/ipc_policy.cpp index e4555ac8f2..356aa81ad8 100644 --- a/src/ipc/ipc_policy.cpp +++ b/src/ipc/ipc_policy.cpp @@ -50,7 +50,6 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, /* * Figure out how this process' rank among local processes. */ - int shm_rank; MPI_Comm_rank(shmcomm, &shm_rank); /* @@ -92,7 +91,6 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, void **ipc_base_uncast = reinterpret_cast(&ipc_base[i]); CHECK_HIP(hipIpcOpenMemHandle(ipc_base_uncast, vec_ipc_handle[i], hipIpcMemLazyEnablePeerAccess)); - // TODO(bpotter): add some error checking here if happens to fail } else { ipc_base[i] = base_heap; } @@ -110,6 +108,15 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, free(vec_ipc_handle); } +__host__ void IpcOnImpl::ipcHostStop() { + for (size_t i = 0; i < shm_size; i++) { + if (i != shm_rank) { + CHECK_HIP(hipIpcCloseMemHandle(ipc_bases[i])); + } + } + CHECK_HIP(hipFree(ipc_bases)); +} + __device__ void IpcOnImpl::ipcCopy(void *dst, void *src, size_t size) { memcpy(dst, src, size); } diff --git a/src/ipc/ipc_policy.hpp b/src/ipc/ipc_policy.hpp index d608dc8d54..9609aa42cb 100644 --- a/src/ipc/ipc_policy.hpp +++ b/src/ipc/ipc_policy.hpp @@ -42,6 +42,8 @@ class IpcOnImpl { using HEAP_BASES_T = std::vector>; public: + int shm_rank{0}; + uint32_t shm_size{0}; char **ipc_bases{nullptr}; @@ -49,6 +51,8 @@ class IpcOnImpl { __host__ void ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, MPI_Comm thread_comm); + __host__ void ipcHostStop(); + __device__ bool isIpcAvailable(int my_pe, int target_pe) { return my_pe / shm_size == target_pe / shm_size; } @@ -115,6 +119,8 @@ class IpcOffImpl { __host__ void ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, MPI_Comm thread_comm) {} + __host__ void ipcHostStop() {} + __device__ bool isIpcAvailable(int my_pe, int target_pe) { return false; } __device__ void ipcGpuInit(Backend *roc_shmem_handle, Context *ctx, diff --git a/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index 1f71119fd6..3f1cf89a17 100644 --- a/tests/unit_tests/CMakeLists.txt +++ b/tests/unit_tests/CMakeLists.txt @@ -71,6 +71,7 @@ target_sources( PRIVATE shmem_gtest.cpp heap_memory_gtest.cpp + hipmalloc_gtest.cpp bin_gtest.cpp binner_gtest.cpp #bitwise_gtest.cpp # Test is disabled becasue of compilation errors @@ -88,6 +89,7 @@ target_sources( #forward_list_gtest.cpp free_list_gtest.cpp context_ipc_gtest.cpp + ipc_impl_simple_coarse_gtest.cpp ) ############################################################################### diff --git a/tests/unit_tests/hipmalloc_gtest.cpp b/tests/unit_tests/hipmalloc_gtest.cpp new file mode 100644 index 0000000000..598952c0e6 --- /dev/null +++ b/tests/unit_tests/hipmalloc_gtest.cpp @@ -0,0 +1,43 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "hipmalloc_gtest.hpp" + +using namespace rocshmem; + +TEST_F(HipMallocTestFixture, normal_1GBx256) { + void* ptr{nullptr}; + size_t gb {1073741824}; + for (int i{0}; i < 256; i++) { + hip_allocator_.allocate(&ptr, gb); + hip_allocator_.deallocate(ptr); + } +} + +TEST_F(HipMallocTestFixture, fine_1GBx256) { + void* ptr{nullptr}; + size_t gb {1073741824}; + for (int i{0}; i < 256; i++) { + hip_allocator_fg_.allocate(&ptr, gb); + hip_allocator_fg_.deallocate(ptr); + } +} diff --git a/tests/unit_tests/hipmalloc_gtest.hpp b/tests/unit_tests/hipmalloc_gtest.hpp new file mode 100644 index 0000000000..d6d341a57b --- /dev/null +++ b/tests/unit_tests/hipmalloc_gtest.hpp @@ -0,0 +1,41 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef ROCSHMEM_HIPMALLOC_GTEST_HPP +#define ROCSHMEM_HIPMALLOC_GTEST_HPP + +#include "gtest/gtest.h" + +#include "../src/memory/symmetric_heap.hpp" +#include "../src/util.hpp" + +namespace rocshmem { + +class HipMallocTestFixture : public ::testing::Test { + public: + HIPAllocator hip_allocator_ {}; + HIPAllocatorFinegrained hip_allocator_fg_ {}; +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_HIPMALLOC_GTEST_HPP diff --git a/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp b/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp new file mode 100644 index 0000000000..dbf72923b5 --- /dev/null +++ b/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp @@ -0,0 +1,1029 @@ +/****************************************************************************** + * 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_impl_simple_coarse_gtest.hpp" + +using namespace rocshmem; + +TEST_F(IPCImplSimpleCoarseTestFixture, ptr_check) { + ASSERT_NE(heap_mem_.get_ptr(), nullptr); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, MPI_num_pes) { + ASSERT_EQ(mpi_.num_pes(), 2); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, IPC_bases) { + for(int i{0}; i < mpi_.num_pes(); i++) { + ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr); + } +} + +TEST_F(IPCImplSimpleCoarseTestFixture, golden_1048576_int) { + iota_golden(1048576); + validate_golden(1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_1024x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + write_wg(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_1024x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + read_wg(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_128x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {128,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_256x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {256,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_512x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {512,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_768x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {768,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_1024x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + write_wg(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_128x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {128,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_256x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {256,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_512x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {512,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_768x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {768,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_1024x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + read_wg(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_64x1x1_1_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 1); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_64x1x1_1_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 1); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_64x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_64x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_3x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {3,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_5x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {5,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_6x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {6,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_7x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {7,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_9x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {9,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_10x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {10,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_11x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {11,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_12x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {12,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_13x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {13,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_14x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {14,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_15x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {15,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_17x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {17,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_18x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {18,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_19x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {19,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_20x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {20,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_21x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {21,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_22x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {22,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_23x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {23,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_24x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {24,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_25x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {25,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_26x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {26,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_27x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {27,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_28x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {28,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_29x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {29,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_30x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {30,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_31x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {31,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_33x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {33,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_34x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {34,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_35x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {35,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_36x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {36,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_37x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {37,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_38x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {38,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_39x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {39,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_40x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {40,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_41x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {41,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_42x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {42,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_43x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {43,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_44x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {44,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_45x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {45,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_46x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {46,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_47x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {47,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_48x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {48,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_49x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {49,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_50x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {50,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_51x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {51,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_52x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {52,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_53x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {53,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_54x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {54,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_55x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {55,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_56x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {56,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_57x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {57,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_58x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {58,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_59x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {59,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_60x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {60,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_61x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {61,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_62x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {62,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_63x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {63,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_3x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {3,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_5x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {5,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_6x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {6,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_7x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {7,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_9x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {9,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_10x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {10,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_11x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {11,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_12x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {12,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_13x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {13,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_14x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {14,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_15x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {15,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_17x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {17,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_18x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {18,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_19x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {19,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_20x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {20,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_21x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {21,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_22x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {22,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_23x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {23,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_24x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {24,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_25x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {25,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_26x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {26,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_27x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {27,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_28x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {28,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_29x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {29,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_30x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {30,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_31x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {31,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_33x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {33,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_34x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {34,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_35x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {35,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_36x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {36,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_37x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {37,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_38x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {38,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_39x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {39,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_40x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {40,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_41x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {41,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_42x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {42,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_43x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {43,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_44x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {44,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_45x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {45,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_46x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {46,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_47x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {47,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_48x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {48,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_49x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {49,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_50x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {50,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_51x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {51,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_52x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {52,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_53x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {53,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_54x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {54,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_55x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {55,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_56x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {56,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_57x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {57,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_58x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {58,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_59x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {59,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_60x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {60,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_61x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {61,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_62x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {62,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_63x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {63,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read(grid, block, 1048576); +} + diff --git a/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp b/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp new file mode 100644 index 0000000000..f31f406416 --- /dev/null +++ b/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp @@ -0,0 +1,236 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP +#define ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP + +#include "gtest/gtest.h" + +#include + +#include +#include "../src/memory/symmetric_heap.hpp" +#include "../src/ipc/ipc_policy.hpp" + +namespace rocshmem { + +__global__ +void +kernel_simple_coarse_copy(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + if (!threadIdx.x) { + ipc_impl->ipcCopy(dest, src, bytes); + ipc_impl->ipcFence(); + } + __syncthreads(); +} + +__global__ +void +kernel_simple_coarse_copy_wg(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + ipc_impl->ipcCopy_wg(dest, src, bytes); + ipc_impl->ipcFence(); + __syncthreads(); +} + +__global__ +void +kernel_simple_coarse_copy_wave(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + ipc_impl->ipcCopy_wave(dest, src, bytes); + ipc_impl->ipcFence(); + __syncthreads(); +} + +class IPCImplSimpleCoarseTestFixture : public ::testing::Test { + + using HEAP_T = HeapMemory; + + using MPI_T = RemoteHeapInfo; + + using FN_T = void (*)(IpcImpl*, int*, int*, size_t); + + public: + IPCImplSimpleCoarseTestFixture() { + ipc_impl_.ipcHostInit(mpi_.my_pe(), mpi_.get_heap_bases() , MPI_COMM_WORLD); + + assert(ipc_impl_dptr_ == nullptr); + hip_allocator_.allocate((void**)&ipc_impl_dptr_, sizeof(IpcImpl)); + + CHECK_HIP(hipMemcpy(ipc_impl_dptr_, &ipc_impl_, + sizeof(IpcImpl), hipMemcpyHostToDevice)); + } + + ~IPCImplSimpleCoarseTestFixture() { + if (ipc_impl_dptr_) { + hip_allocator_.deallocate(ipc_impl_dptr_); + } + + ipc_impl_.ipcHostStop(); + } + + void launch(FN_T f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes) { + f<<>>(ipc_impl_dptr_, src, dest, bytes); + CHECK_HIP(hipStreamSynchronize(nullptr)); + } + + enum TestType { + READ = 0, + WRITE = 1 + }; + + void write(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(WRITE); + copy(WRITE, grid, block); + validate_dest_buffer(WRITE); + } + + void write_wg(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(WRITE); + copy_wg(WRITE, grid, block); + validate_dest_buffer(WRITE); + } + + void write_wave(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(WRITE); + copy_wave(WRITE, grid, block); + validate_dest_buffer(WRITE); + } + + void read(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(READ); + copy(READ, grid, block); + validate_dest_buffer(READ); + } + + void read_wg(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(READ); + copy_wg(READ, grid, block); + validate_dest_buffer(READ); + } + + void read_wave(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(READ); + copy_wave(READ, grid, block); + validate_dest_buffer(READ); + } + + void iota_golden(size_t elems) { + golden_.resize(elems); + std::iota(golden_.begin(), golden_.end(), 0); + } + + void validate_golden(size_t elems) { + ASSERT_EQ(golden_.size(), elems); + for (int i{0}; i < golden_.size(); i++) { + ASSERT_EQ(golden_[i], i); + } + } + + void initialize_src_buffer(TestType test) { + if (!pe_initializes_src_buffer(test)) { + return; + } + size_t bytes = golden_.size() * sizeof(int); + auto dev_src = reinterpret_cast(ipc_impl_.ipc_bases[mpi_.my_pe()]); + CHECK_HIP(hipMemcpy(dev_src, golden_.data(), bytes, hipMemcpyHostToDevice)); + CHECK_HIP(hipStreamSynchronize(nullptr)); + } + + bool pe_initializes_src_buffer(TestType test) { + bool is_write_test = test; + bool is_read_test = !test; + return (is_write_test && mpi_.my_pe() == 0) || + (is_read_test && mpi_.my_pe() == 1); + } + + void execute(TestType test, FN_T fn, const dim3 grid, const dim3 block) { + if (mpi_.my_pe()) { + mpi_.barrier(); + mpi_.barrier(); + return; + } + int *src{nullptr}; + int *dest{nullptr}; + if (test == WRITE) { + src = reinterpret_cast(ipc_impl_.ipc_bases[0]); + dest = reinterpret_cast(ipc_impl_.ipc_bases[1]); + } else { + src = reinterpret_cast(ipc_impl_.ipc_bases[1]); + dest = reinterpret_cast(ipc_impl_.ipc_bases[0]); + } + size_t bytes = golden_.size() * sizeof(int); + mpi_.barrier(); + launch(fn, grid, block, src, dest, bytes); + mpi_.barrier(); + } + + void copy(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_coarse_copy, grid, block); + } + + void copy_wg(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_coarse_copy_wg, grid, block); + } + + void copy_wave(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_coarse_copy_wave, grid, block); + } + + void validate_dest_buffer(TestType test) { + if (!pe_validates_dest_buffer(test)) { + return; + } + + auto dev_dest = reinterpret_cast(ipc_impl_.ipc_bases[mpi_.my_pe()]); + for (int i{0}; i < golden_.size(); i++) { + ASSERT_EQ(golden_[i], dev_dest[i]); + } + } + + bool pe_validates_dest_buffer(TestType test) { + return !pe_initializes_src_buffer(test); + } + + protected: + std::vector golden_; + + std::vector output_; + + HEAP_T heap_mem_ {}; + + MPI_T mpi_ {heap_mem_.get_ptr(), heap_mem_.get_size()}; + + IpcImpl ipc_impl_ {}; + + IpcImpl *ipc_impl_dptr_ {nullptr}; + + HIPAllocator hip_allocator_ {}; +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP From 24375a949e2668171bb82ac220472728d7973c46 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 14 Aug 2024 20:44:35 -0700 Subject: [PATCH 06/20] Code refactor move ipc_policy.hpp and ipc_policy.cpp files to src, since they are used by all the conduits. --- CMakeLists.txt | 3 +- cmake/config.h.in | 1 + src/CMakeLists.txt | 4 + src/backend_bc.hpp | 2 +- src/context.hpp | 2 +- src/ipc/CMakeLists.txt | 2 +- src/ipc_policy.cpp | 132 ++++++++++++++ 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 src/ipc_policy.cpp create mode 100644 src/ipc_policy.hpp diff --git a/CMakeLists.txt b/CMakeLists.txt index ee7f42fb12..d068315c62 100644 --- a/CMakeLists.txt +++ b/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/cmake/config.h.in b/cmake/config.h.in index 6ec07c2853..3b2732137a 100644 --- a/cmake/config.h.in +++ b/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/src/CMakeLists.txt b/src/CMakeLists.txt index 15162c2384..78b90fe035 100644 --- a/src/CMakeLists.txt +++ b/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/src/backend_bc.hpp b/src/backend_bc.hpp index b968c0f3bb..88cc6b3ba3 100644 --- a/src/backend_bc.hpp +++ b/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/src/context.hpp b/src/context.hpp index 4ff1251220..72c8f38352 100644 --- a/src/context.hpp +++ b/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/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt index e34f62906e..1fd601d43d 100644 --- a/src/ipc/CMakeLists.txt +++ b/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/src/ipc_policy.cpp b/src/ipc_policy.cpp new file mode 100644 index 0000000000..2c3a96fa7c --- /dev/null +++ b/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/src/ipc_policy.hpp b/src/ipc_policy.hpp new file mode 100644 index 0000000000..c0190198ca --- /dev/null +++ b/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/src/reverse_offload/block_handle.hpp b/src/reverse_offload/block_handle.hpp index f2ba1872b3..d2235f8639 100644 --- a/src/reverse_offload/block_handle.hpp +++ b/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/src/reverse_offload/queue_proxy.hpp b/src/reverse_offload/queue_proxy.hpp index f3ddcec1ce..9cde0cb0cb 100644 --- a/src/reverse_offload/queue_proxy.hpp +++ b/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/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp b/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp index f31f406416..02dfd8c55a 100644 --- a/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp +++ b/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 49779863c266d068ba74988275ee985620c85391 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 14 Aug 2024 22:59:02 -0700 Subject: [PATCH 07/20] 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 --- scripts/build_configs/ipc_single | 30 ++++ src/CMakeLists.txt | 8 +- src/backend_bc.cpp | 18 ++- src/backend_type.hpp | 32 +++- src/context_incl.hpp | 5 +- src/context_tmpl_device.hpp | 4 +- src/context_tmpl_host.hpp | 4 +- src/ipc/CMakeLists.txt | 5 +- src/ipc/backend_ipc.cpp | 183 +++++++++++++++++++++ src/ipc/backend_ipc.hpp | 190 ++++++++++++++++++++++ src/ipc/context_ipc_device.cpp | 129 +++++++++++++++ src/ipc/context_ipc_device.hpp | 241 ++++++++++++++++++++++++++++ src/ipc/context_ipc_host.cpp | 85 ++++++++++ src/ipc/context_ipc_host.hpp | 149 +++++++++++++++++ src/ipc/context_ipc_tmpl_device.hpp | 2 +- src/ipc/context_ipc_tmpl_host.hpp | 173 ++++++++++++++++++++ src/ipc/ipc_backend_proxy.hpp | 69 ++++++++ src/ipc/ipc_context_proxy.hpp | 90 +++++++++++ src/roc_shmem.cpp | 19 ++- src/roc_shmem_gpu.cpp | 4 +- 20 files changed, 1410 insertions(+), 30 deletions(-) create mode 100755 scripts/build_configs/ipc_single create mode 100644 src/ipc/backend_ipc.cpp create mode 100644 src/ipc/backend_ipc.hpp create mode 100644 src/ipc/context_ipc_device.cpp create mode 100644 src/ipc/context_ipc_device.hpp create mode 100644 src/ipc/context_ipc_host.cpp create mode 100644 src/ipc/context_ipc_host.hpp create mode 100644 src/ipc/context_ipc_tmpl_host.hpp create mode 100644 src/ipc/ipc_backend_proxy.hpp create mode 100644 src/ipc/ipc_context_proxy.hpp diff --git a/scripts/build_configs/ipc_single b/scripts/build_configs/ipc_single new file mode 100755 index 0000000000..2013b2b67c --- /dev/null +++ b/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/src/CMakeLists.txt b/src/CMakeLists.txt index 78b90fe035..5c31872765 100644 --- a/src/CMakeLists.txt +++ b/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/src/backend_bc.cpp b/src/backend_bc.cpp index 1c3a619887..c55b1cd2d1 100644 --- a/src/backend_bc.cpp +++ b/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/src/backend_type.hpp b/src/backend_type.hpp index ad9e992636..d8b2972504 100644 --- a/src/backend_type.hpp +++ b/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/src/context_incl.hpp b/src/context_incl.hpp index ea39a2ead5..f86ca3dfe8 100644 --- a/src/context_incl.hpp +++ b/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/src/context_tmpl_device.hpp b/src/context_tmpl_device.hpp index 213fb90d3b..d6d0df63fa 100644 --- a/src/context_tmpl_device.hpp +++ b/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/src/context_tmpl_host.hpp b/src/context_tmpl_host.hpp index 56b649e19d..bce99c1d6c 100644 --- a/src/context_tmpl_host.hpp +++ b/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/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt index 1fd601d43d..b8d50d12e2 100644 --- a/src/ipc/CMakeLists.txt +++ b/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/src/ipc/backend_ipc.cpp b/src/ipc/backend_ipc.cpp new file mode 100644 index 0000000000..bb1731add3 --- /dev/null +++ b/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/src/ipc/backend_ipc.hpp b/src/ipc/backend_ipc.hpp new file mode 100644 index 0000000000..1a89dd63f1 --- /dev/null +++ b/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/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp new file mode 100644 index 0000000000..97937fcd1e --- /dev/null +++ b/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/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp new file mode 100644 index 0000000000..492d63ca04 --- /dev/null +++ b/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/src/ipc/context_ipc_host.cpp b/src/ipc/context_ipc_host.cpp new file mode 100644 index 0000000000..0d3464f33d --- /dev/null +++ b/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/src/ipc/context_ipc_host.hpp b/src/ipc/context_ipc_host.hpp new file mode 100644 index 0000000000..f9421ca210 --- /dev/null +++ b/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/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index c90e0a3d00..28cd718f0a 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/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/src/ipc/context_ipc_tmpl_host.hpp b/src/ipc/context_ipc_tmpl_host.hpp new file mode 100644 index 0000000000..76484de34d --- /dev/null +++ b/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/src/ipc/ipc_backend_proxy.hpp b/src/ipc/ipc_backend_proxy.hpp new file mode 100644 index 0000000000..0077fb053d --- /dev/null +++ b/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/src/ipc/ipc_context_proxy.hpp b/src/ipc/ipc_context_proxy.hpp new file mode 100644 index 0000000000..867f199094 --- /dev/null +++ b/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/src/roc_shmem.cpp b/src/roc_shmem.cpp index 8c61d6c864..5ca9051f37 100644 --- a/src/roc_shmem.cpp +++ b/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/src/roc_shmem_gpu.cpp b/src/roc_shmem_gpu.cpp index 269a98a28a..086a6deed1 100644 --- a/src/roc_shmem_gpu.cpp +++ b/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 b68867ee17a5edbab35090bd3dbb0b9b2d90e0ce Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 15 Aug 2024 08:52:06 -0700 Subject: [PATCH 08/20] 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} --- src/ipc/context_ipc.cpp | 129 --------------------- src/ipc/context_ipc.hpp | 241 ---------------------------------------- src/ipc/ipc_policy.cpp | 132 ---------------------- src/ipc/ipc_policy.hpp | 172 ---------------------------- 4 files changed, 674 deletions(-) delete mode 100644 src/ipc/context_ipc.cpp delete mode 100644 src/ipc/context_ipc.hpp delete mode 100644 src/ipc/ipc_policy.cpp delete mode 100644 src/ipc/ipc_policy.hpp diff --git a/src/ipc/context_ipc.cpp b/src/ipc/context_ipc.cpp deleted file mode 100644 index a825a135e9..0000000000 --- a/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/src/ipc/context_ipc.hpp b/src/ipc/context_ipc.hpp deleted file mode 100644 index 492d63ca04..0000000000 --- a/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/src/ipc/ipc_policy.cpp b/src/ipc/ipc_policy.cpp deleted file mode 100644 index 356aa81ad8..0000000000 --- a/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/src/ipc/ipc_policy.hpp b/src/ipc/ipc_policy.hpp deleted file mode 100644 index 9609aa42cb..0000000000 --- a/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 c9dbcf80c20986fa60274effc43ce9e7957fe653 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 15 Aug 2024 11:54:56 -0700 Subject: [PATCH 09/20] Update context_ipc_gtest.cpp to use IPCbackend --- src/roc_shmem.cpp | 2 -- tests/unit_tests/CMakeLists.txt | 2 +- tests/unit_tests/context_ipc_gtest.hpp | 6 +++--- 3 files changed, 4 insertions(+), 6 deletions(-) diff --git a/src/roc_shmem.cpp b/src/roc_shmem.cpp index 5ca9051f37..32ef91bbda 100644 --- a/src/roc_shmem.cpp +++ b/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/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index 3f1cf89a17..349bb2c2bf 100644 --- a/tests/unit_tests/CMakeLists.txt +++ b/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/tests/unit_tests/context_ipc_gtest.hpp b/tests/unit_tests/context_ipc_gtest.hpp index fc39504958..24a80fcedb 100644 --- a/tests/unit_tests/context_ipc_gtest.hpp +++ b/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}; }; From c8b0f2378e7eea965e8e7b82ee69420f1f89b961 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 15 Aug 2024 13:17:44 -0700 Subject: [PATCH 10/20] Add gets and puts functionality to IPC context --- src/ipc/context_ipc_device.cpp | 53 ++++++++++++++++++++++++++++++++-- src/ipc/context_ipc_device.hpp | 7 +++++ 2 files changed, 58 insertions(+), 2 deletions(-) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index 97937fcd1e..7b9f3dd469 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -32,11 +32,19 @@ #include "config.h" // NOLINT(build/include_subdir) #include "roc_shmem/roc_shmem.hpp" +#include "backend_ipc.hpp" namespace rocshmem { __host__ IPCContext::IPCContext(Backend *b) : Context(b, false) { + IPCBackend *backend{static_cast(b)}; + ipcImpl = &backend->ipcImpl; + + auto *bp{backend->ipc_backend_proxy.get()}; + + g_ret = bp->g_ret; + atomic_base_ptr = bp->atomic_ret->atomic_base_ptr; } __device__ void IPCContext::threadfence_system() { @@ -50,18 +58,32 @@ __device__ void IPCContext::ctx_destroy(){ __device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems, int pe) { + // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) + int local_pe = pe % ipcImpl->shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; + ipcImpl->ipcCopy(ipcImpl->ipc_bases[local_pe] + L_offset, + const_cast(source), nelems); } __device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, int pe) { + // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) + int local_pe = pe % ipcImpl->shm_size; + const char *src_typed = reinterpret_cast(source); + uint64_t L_offset = + const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; + ipcImpl->ipcCopy(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems); } __device__ void IPCContext::putmem_nbi(void *dest, const void *source, size_t nelems, int pe) { + putmem(dest, source, nelems, pe); } __device__ void IPCContext::getmem_nbi(void *dest, const void *source, size_t nelems, int pe) { + getmem(dest, source, nelems, pe); } __device__ void IPCContext::fence() { @@ -92,38 +114,65 @@ __device__ void IPCContext::sync(roc_shmem_team_t team) { __device__ void IPCContext::putmem_wg(void *dest, const void *source, size_t nelems, int pe) { + // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) + int local_pe = pe % ipcImpl->shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; + ipcImpl->ipcCopy_wg(ipcImpl->ipc_bases[local_pe] + L_offset, + const_cast(source), nelems); __syncthreads(); } __device__ void IPCContext::getmem_wg(void *dest, const void *source, size_t nelems, int pe) { + // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) + int local_pe = pe % ipcImpl->shm_size; + const char *src_typed = reinterpret_cast(source); + uint64_t L_offset = + const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; + ipcImpl->ipcCopy_wg(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems); __syncthreads(); } __device__ void IPCContext::putmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) { - __syncthreads(); + putmem_wg(dest, source, nelems, pe); } __device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) { - __syncthreads(); + getmem_wg(dest, source, nelems, pe); } __device__ void IPCContext::putmem_wave(void *dest, const void *source, size_t nelems, int pe) { + // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) + int local_pe = pe % ipcImpl->shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; + ipcImpl->ipcCopy_wave(ipcImpl->ipc_bases[local_pe] + L_offset, + const_cast(source), nelems); } __device__ void IPCContext::getmem_wave(void *dest, const void *source, size_t nelems, int pe) { + // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) + int local_pe = pe % ipcImpl->shm_size; + const char *src_typed = reinterpret_cast(source); + uint64_t L_offset = + const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; + ipcImpl->ipcCopy_wave(dest, ipcImpl->ipc_bases[local_pe] + L_offset, + nelems); } __device__ void IPCContext::putmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) { + putmem_wave(dest, source, nelems, pe); } __device__ void IPCContext::getmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) { + getmem_wave(dest, source, nelems, pe); } } // namespace rocshmem diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index 492d63ca04..afadb9828c 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -234,6 +234,13 @@ class IPCContext : public Context { private: + //context class has IpcImpl object (ipcImpl_) + IpcImpl *ipcImpl{nullptr}; + + uint64_t* atomic_base_ptr{nullptr}; + + char* g_ret; + }; } // namespace rocshmem From a59bdd4f6b6792a097342e774e929fa26000f483 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 22 Aug 2024 09:15:44 -0700 Subject: [PATCH 11/20] Add IPC teams --- src/ipc/CMakeLists.txt | 1 + src/ipc/ipc_team.cpp | 56 ++++++++++++++++++++++++++++++++++++++++++ src/ipc/ipc_team.hpp | 50 +++++++++++++++++++++++++++++++++++++ src/team.cpp | 4 +++ src/team.hpp | 3 +++ 5 files changed, 114 insertions(+) create mode 100644 src/ipc/ipc_team.cpp create mode 100644 src/ipc/ipc_team.hpp diff --git a/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt index b8d50d12e2..e8d64ff523 100644 --- a/src/ipc/CMakeLists.txt +++ b/src/ipc/CMakeLists.txt @@ -29,4 +29,5 @@ target_sources( context_ipc_device.cpp context_ipc_host.cpp backend_ipc.cpp + ipc_team.cpp ) diff --git a/src/ipc/ipc_team.cpp b/src/ipc/ipc_team.cpp new file mode 100644 index 0000000000..433ca32d4f --- /dev/null +++ b/src/ipc/ipc_team.cpp @@ -0,0 +1,56 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "ipc_team.hpp" + +#include "../backend_type.hpp" +#include "backend_ipc.hpp" + +namespace rocshmem { + +IPCTeam::IPCTeam(Backend *backend, TeamInfo *team_info_parent, + TeamInfo *team_info_world, int num_pes, int my_pe, + MPI_Comm mpi_comm, int pool_index) + : Team(backend, team_info_parent, team_info_world, num_pes, my_pe, + mpi_comm) { + type = BackendType::IPC_BACKEND; + const IPCBackend *b = static_cast(backend); + + pool_index_ = pool_index; + + barrier_pSync = + &(b->barrier_pSync_pool[pool_index * ROC_SHMEM_BARRIER_SYNC_SIZE]); + reduce_pSync = + &(b->reduce_pSync_pool[pool_index * ROC_SHMEM_REDUCE_SYNC_SIZE]); + bcast_pSync = &(b->bcast_pSync_pool[pool_index * ROC_SHMEM_BCAST_SYNC_SIZE]); + alltoall_pSync = + &(b->alltoall_pSync_pool[pool_index * ROC_SHMEM_ALLTOALL_SYNC_SIZE]); + + pWrk = reinterpret_cast(b->pWrk_pool) + + ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index; + pAta = reinterpret_cast(b->pAta_pool) + + ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index; +} + +IPCTeam::~IPCTeam() {} + +} // namespace rocshmem diff --git a/src/ipc/ipc_team.hpp b/src/ipc/ipc_team.hpp new file mode 100644 index 0000000000..e931adcfe4 --- /dev/null +++ b/src/ipc/ipc_team.hpp @@ -0,0 +1,50 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_IPC_TEAM_HPP_ +#define LIBRARY_SRC_IPC_TEAM_HPP_ + +#include "../team.hpp" + +namespace rocshmem { + +class IPCTeam : public Team { + public: + IPCTeam(Backend* handle, TeamInfo* team_info_wrt_parent, + TeamInfo* team_info_wrt_world, int num_pes, int my_pe, + MPI_Comm team_comm, int pool_index); + + virtual ~IPCTeam(); + + long* barrier_pSync{nullptr}; + long* reduce_pSync{nullptr}; + long* bcast_pSync{nullptr}; + long* alltoall_pSync{nullptr}; + void* pWrk{nullptr}; + void* pAta{nullptr}; + + int pool_index_{-1}; +}; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_TEAM_HPP_ diff --git a/src/team.cpp b/src/team.cpp index fefb520328..21ed00037c 100644 --- a/src/team.cpp +++ b/src/team.cpp @@ -44,6 +44,10 @@ ROTeam* get_internal_ro_team(roc_shmem_team_t team) { return reinterpret_cast(team); } +IPCTeam* get_internal_ipc_team(roc_shmem_team_t team) { + return reinterpret_cast(team); +} + __host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe, roc_shmem_team_t dst_team) { if (src_team == ROC_SHMEM_TEAM_INVALID || diff --git a/src/team.hpp b/src/team.hpp index 56475e799b..0dfaac6ae3 100644 --- a/src/team.hpp +++ b/src/team.hpp @@ -34,6 +34,7 @@ class Backend; class Team; class ROTeam; class GPUIBTeam; +class IPCTeam; class TeamInfo { public: @@ -162,6 +163,8 @@ GPUIBTeam* get_internal_gpu_ib_team(roc_shmem_team_t team); ROTeam* get_internal_ro_team(roc_shmem_team_t team); +IPCTeam* get_internal_ipc_team(roc_shmem_team_t team); + __host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe, roc_shmem_team_t dst_team); From a9571ec002e3e4b70179eeae277818a3dc8ab99d Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 22 Aug 2024 09:28:09 -0700 Subject: [PATCH 12/20] Add buffers required for collectives --- src/ipc/backend_ipc.cpp | 249 ++++++++++++++++++++++++++++++++++++---- src/ipc/backend_ipc.hpp | 92 ++++++++++++--- 2 files changed, 306 insertions(+), 35 deletions(-) diff --git a/src/ipc/backend_ipc.cpp b/src/ipc/backend_ipc.cpp index bb1731add3..58caaf53f7 100644 --- a/src/ipc/backend_ipc.cpp +++ b/src/ipc/backend_ipc.cpp @@ -21,6 +21,7 @@ *****************************************************************************/ #include "backend_ipc.hpp" +#include "ipc_team.hpp" namespace rocshmem { @@ -34,6 +35,24 @@ namespace rocshmem { extern roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT; +roc_shmem_team_t get_external_team(GPUIBTeam *team) { + return reinterpret_cast(team); +} + +int get_ls_non_zero_bit(char *bitmask, int mask_length) { + int position = -1; + + for (int bit_i = 0; bit_i < mask_length; bit_i++) { + int byte_i = bit_i / CHAR_BIT; + if (bitmask[byte_i] & (1 << (bit_i % CHAR_BIT))) { + position = bit_i; + break; + } + } + + return position; +} + IPCBackend::IPCBackend(MPI_Comm comm) : Backend() { type = BackendType::IPC_BACKEND; @@ -54,7 +73,6 @@ IPCBackend::IPCBackend(MPI_Comm comm) /* 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); @@ -66,12 +84,42 @@ IPCBackend::IPCBackend(MPI_Comm comm) default_context_proxy_ = IPCDefaultContextProxyT(this); + setup_team_world(); + + roc_shmem_collective_init(); + + teams_init(); + setup_ctxs(); } IPCBackend::~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 + + // delete host_interface; + // host_interface = nullptr; + + /** + * Destroy teams infrastructure + * and team world + */ + teams_destroy(); + auto *team_world{team_tracker.get_team_world()}; + team_world->~Team(); + CHECK_HIP(hipFree(team_world)); + CHECK_HIP(hipFree(ctx_array)); } @@ -100,6 +148,31 @@ __device__ void IPCBackend::destroy_ctx(roc_shmem_ctx_t *ctx) { ctx_free_list.get()->push_back(static_cast(ctx->ctx_opaque)); } +void IPCBackend::setup_team_world() { + TeamInfo *team_info_wrt_parent, *team_info_wrt_world; + + /** + * Allocate device-side memory for team_world and construct a + * IPC team in it. + */ + CHECK_HIP(hipMalloc(&team_info_wrt_parent, sizeof(TeamInfo))); + CHECK_HIP(hipMalloc(&team_info_wrt_world, sizeof(TeamInfo))); + + new (team_info_wrt_parent) TeamInfo(nullptr, 0, 1, num_pes); + new (team_info_wrt_world) TeamInfo(nullptr, 0, 1, num_pes); + + IPCTeam *team_world{nullptr}; + CHECK_HIP(hipMalloc(&team_world, sizeof(IPCTeam))); + new (team_world) IPCTeam(this, team_info_wrt_parent, team_info_wrt_world, + num_pes, my_pe, thread_comm, 0); + team_tracker.set_team_world(team_world); + + /** + * Copy the address to ROC_SHMEM_TEAM_WORLD. + */ + ROC_SHMEM_TEAM_WORLD = reinterpret_cast(team_world); +} + void IPCBackend::init_mpi_once(MPI_Comm comm) { int init_done{}; NET_CHECK(MPI_Initialized(&init_done)); @@ -119,15 +192,53 @@ void IPCBackend::init_mpi_once(MPI_Comm comm) { } void IPCBackend::team_destroy(roc_shmem_team_t team) { - assert(false); + IPCTeam *team_obj = get_internal_ipc_team(team); + + /* Mark the pool as available */ + int bit = team_obj->pool_index_; + int byte_i = bit / CHAR_BIT; + pool_bitmask_[byte_i] |= 1 << (bit % CHAR_BIT); + + team_obj->~IPCTeam(); + CHECK_HIP(hipFree(team_obj)); } -void IPCBackend::create_new_team(Team *parent_team, +void IPCBackend::create_new_team([[maybe_unused]] Team *parent_team, TeamInfo *team_info_wrt_parent, TeamInfo *team_info_wrt_world, int num_pes, int my_pe_in_new_team, MPI_Comm team_comm, roc_shmem_team_t *new_team) { - assert(false); + /** + * Read the bit mask and find out a common index into + * the pool of available work arrays. + */ + NET_CHECK(MPI_Allreduce(pool_bitmask_, reduced_bitmask_, bitmask_size_, + MPI_CHAR, MPI_BAND, team_comm)); + + /* Pick the least significant non-zero bit (logical layout) in the reduced + * bitmask */ + auto max_num_teams{team_tracker.get_max_num_teams()}; + int common_index = get_ls_non_zero_bit(reduced_bitmask_, max_num_teams); + if (common_index < 0) { + /* No team available */ + abort(); + } + + /* Mark the team as taken (by unsetting the bit in the pool bitmask) */ + int byte = common_index / CHAR_BIT; + pool_bitmask_[byte] &= ~(1 << (common_index % CHAR_BIT)); + + /** + * Allocate device-side memory for team_world and + * construct a GPU_IB team in it + */ + GPUIBTeam *new_team_obj; + CHECK_HIP(hipMalloc(&new_team_obj, sizeof(IPCTeam))); + new (new_team_obj) + IPCTeam(this, team_info_wrt_parent, team_info_wrt_world, num_pes, + my_pe_in_new_team, team_comm, common_index); + + *new_team = get_external_team(new_team_obj); } void IPCBackend::ctx_create(int64_t options, void **ctx) { @@ -160,24 +271,122 @@ void IPCBackend::initIPC() { 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); } +void IPCBackend::teams_destroy() { + roc_shmem_free(barrier_pSync_pool); + roc_shmem_free(reduce_pSync_pool); + roc_shmem_free(bcast_pSync_pool); + roc_shmem_free(alltoall_pSync_pool); + roc_shmem_free(pWrk_pool); + roc_shmem_free(pAta_pool); + + free(pool_bitmask_); + free(reduced_bitmask_); +} + +void IPCBackend::roc_shmem_collective_init() { + /* + * Allocate heap space for barrier_sync + */ + size_t one_sync_size_bytes{sizeof(*barrier_sync)}; + size_t sync_size_bytes{one_sync_size_bytes * ROC_SHMEM_BARRIER_SYNC_SIZE}; + heap.malloc(reinterpret_cast(&barrier_sync), sync_size_bytes); + + /* + * Initialize the barrier synchronization array with default values. + */ + for (int i = 0; i < num_pes; i++) { + barrier_sync[i] = ROC_SHMEM_SYNC_VALUE; + } + + /* + * Make sure that all processing elements have done this before + * continuing. + */ + NET_CHECK(MPI_Barrier(thread_comm)); +} + +void IPCBackend::teams_init() { + /** + * Allocate pools for the teams sync and work arrary from the SHEAP. + */ + auto max_num_teams{team_tracker.get_max_num_teams()}; + barrier_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_BARRIER_SYNC_SIZE * max_num_teams)); + reduce_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_REDUCE_SYNC_SIZE * max_num_teams)); + bcast_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_BCAST_SYNC_SIZE * max_num_teams)); + alltoall_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_ALLTOALL_SYNC_SIZE * max_num_teams)); + + /* Accommodating for largest possible data type for pWrk */ + pWrk_pool = roc_shmem_malloc( + sizeof(double) * ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * max_num_teams); + pAta_pool = roc_shmem_malloc(sizeof(double) * ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * + max_num_teams); + + /** + * Initialize the sync arrays in the pool with default values. + */ + long *barrier_pSync, *reduce_pSync, *bcast_pSync, *alltoall_pSync; + for (int team_i = 0; team_i < max_num_teams; team_i++) { + barrier_pSync = reinterpret_cast( + &barrier_pSync_pool[team_i * ROC_SHMEM_BARRIER_SYNC_SIZE]); + reduce_pSync = reinterpret_cast( + &reduce_pSync_pool[team_i * ROC_SHMEM_REDUCE_SYNC_SIZE]); + bcast_pSync = reinterpret_cast( + &bcast_pSync_pool[team_i * ROC_SHMEM_BCAST_SYNC_SIZE]); + alltoall_pSync = reinterpret_cast( + &alltoall_pSync_pool[team_i * ROC_SHMEM_ALLTOALL_SYNC_SIZE]); + + for (int i = 0; i < ROC_SHMEM_BARRIER_SYNC_SIZE; i++) { + barrier_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + for (int i = 0; i < ROC_SHMEM_REDUCE_SYNC_SIZE; i++) { + reduce_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + for (int i = 0; i < ROC_SHMEM_BCAST_SYNC_SIZE; i++) { + bcast_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + for (int i = 0; i < ROC_SHMEM_ALLTOALL_SYNC_SIZE; i++) { + alltoall_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + } + + /** + * Initialize bit mask + * + * Logical: + * MSB..........................................................................LSB + * Physical: MSB...1st least significant 8 bits...LSB MSB...2nd least + * signifant 8 bits...LSB + * + * Description shows only a 2-byte long mask but idea extends to any + * arbitrary size. + */ + bitmask_size_ = (max_num_teams % CHAR_BIT) ? (max_num_teams / CHAR_BIT + 1) + : (max_num_teams / CHAR_BIT); + pool_bitmask_ = reinterpret_cast(malloc(bitmask_size_)); + reduced_bitmask_ = reinterpret_cast(malloc(bitmask_size_)); + + memset(pool_bitmask_, 0, bitmask_size_); + memset(reduced_bitmask_, 0, bitmask_size_); + /* Set all to available except the 0th one (reserved for TEAM_WORLD) */ + for (int bit_i = 1; bit_i < max_num_teams; bit_i++) { + int byte_i = bit_i / CHAR_BIT; + + pool_bitmask_[byte_i] |= 1 << (bit_i % CHAR_BIT); + } + + /** + * Make sure that all processing elements have done this before + * continuing. + */ + NET_CHECK(MPI_Barrier(thread_comm)); +} } // namespace rocshmem \ No newline at end of file diff --git a/src/ipc/backend_ipc.hpp b/src/ipc/backend_ipc.hpp index 1a89dd63f1..cc385f16ee 100644 --- a/src/ipc/backend_ipc.hpp +++ b/src/ipc/backend_ipc.hpp @@ -86,21 +86,6 @@ class IPCBackend : public Backend { */ 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. * @@ -136,6 +121,41 @@ class IPCBackend : public Backend { */ HostInterface *host_interface{nullptr}; + /** + * @brief Scratchpad for the internal barrier algorithms. + */ + int64_t *barrier_sync{nullptr}; + + /** + * @brief Handle for raw memory for barrier sync + */ + long *barrier_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for reduce sync + */ + long *reduce_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for broadcast sync + */ + long *bcast_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for alltoall sync + */ + long *alltoall_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for work + */ + void *pWrk_pool{nullptr}; + + /** + * @brief Handle for raw memory for alltoall + */ + void *pAta_pool{nullptr}; + protected: /** * @copydoc Backend::dump_backend_stats() @@ -159,6 +179,30 @@ class IPCBackend : public Backend { */ std::unique_ptr default_host_ctx{nullptr}; + /** + * @brief Allocate and initialize team world. + */ + void setup_team_world(); + + /** + * @brief Initialize the resources required to support teams + */ + void teams_init(); + + /** + * @brief Destruct the resources required to support teams + */ + void teams_destroy(); + + /** + * @brief Allocate and initialize barrier operation addresses on + * symmetric heap. + * + * When this method completes, the barrier_sync member will be available + * for use. + */ + void roc_shmem_collective_init(); + private: /** * @brief Proxy for the default context @@ -182,6 +226,24 @@ class IPCBackend : public Backend { */ size_t maximum_num_contexts_{1024}; + /** + * @brief The bitmask representing the availability of teams in the pool + */ + char *pool_bitmask_{nullptr}; + + /** + * @brief Bitmask to store the reduced result of bitmasks on pariticipating + * PEs + * + * With no thread-safety for this bitmask, multithreaded creation of teams is + * not supported. + */ + char *reduced_bitmask_{nullptr}; + + /** + * @brief Size of the bitmask + */ + int bitmask_size_{-1}; }; From e2e30b53395cdc605d4d4d99dc983f15cf8fcebf Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 27 Aug 2024 06:25:58 -0700 Subject: [PATCH 13/20] remove device wait_until functions adding the device versions of the wait_until* and test functions in the ipc folder leads to linking errors of the functional tests. Remove them and use for now the upper level versions of the functions, similarly to the RO conduit. Might have to revisit this later again. --- src/ipc/context_ipc_device.hpp | 38 --------------------- src/ipc/context_ipc_tmpl_device.hpp | 52 ----------------------------- 2 files changed, 90 deletions(-) diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index afadb9828c..a8cefb446a 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -194,44 +194,6 @@ class IPCContext : public Context { 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: //context class has IpcImpl object (ipcImpl_) diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index 28cd718f0a..4bf652781e 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -223,58 +223,6 @@ __device__ void IPCContext::get_nbi_wave(T *dest, const T *source, getmem_nbi_wave(dest, source, nelems * sizeof(T), 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) { - return 0; -} - -template -__device__ size_t wait_until_some(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T val){ - return 0; -} - -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){ - return 0; -} - -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) { - return 0; -} - } // namespace rocshmem #endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ From 0de3b5e6fcdde7413053a55d9ac86b2e4782d2c0 Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 13 Aug 2024 14:09:15 -0700 Subject: [PATCH 14/20] first cut on collectives and sync code is based on the GPUIB implementations of the routines, which seem however generic enough to work also for the IPC conduit. Some code is in for broadcast, fcollect, and alltoall. --- src/ipc/CMakeLists.txt | 1 + src/ipc/context_ipc_device.cpp | 12 --- src/ipc/context_ipc_device.hpp | 39 ++++++++- src/ipc/context_ipc_device_coll.cpp | 120 ++++++++++++++++++++++++++++ src/ipc/context_ipc_tmpl_device.hpp | 114 +++++++++++++++++++++++++- 5 files changed, 268 insertions(+), 18 deletions(-) create mode 100644 src/ipc/context_ipc_device_coll.cpp diff --git a/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt index e8d64ff523..07a40538a5 100644 --- a/src/ipc/CMakeLists.txt +++ b/src/ipc/CMakeLists.txt @@ -30,4 +30,5 @@ target_sources( context_ipc_host.cpp backend_ipc.cpp ipc_team.cpp + context_ipc_device_coll.cpp ) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index 7b9f3dd469..48d7dd9b4f 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -100,18 +100,6 @@ __device__ void *IPCContext::shmem_ptr(const void *dest, int pe) { 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) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index a8cefb446a..01fb1c1fb8 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -196,13 +196,44 @@ class IPCContext : public Context { private: - //context class has IpcImpl object (ipcImpl_) - IpcImpl *ipcImpl{nullptr}; + //context class has IpcImpl object (ipcImpl_) + IpcImpl *ipcImpl{nullptr}; - uint64_t* atomic_base_ptr{nullptr}; + uint64_t* atomic_base_ptr{nullptr}; - char* g_ret; + char* g_ret; + //internal functions used by collective operations + template + __device__ void internal_put_broadcast(T *dst, const T *src, int nelems, + int pe_root, int PE_start, + int logPE_stride, int PE_size, + long *pSync); // NOLINT(runtime/int) + + template + __device__ void internal_get_broadcast(T *dst, const T *src, int nelems, + int pe_root, + long *pSync); // NOLINT(runtime/int) + + template + __device__ void fcollect_linear(roc_shmem_team_t team, T *dest, + const T *source, int nelems); + + template + __device__ void alltoall_linear(roc_shmem_team_t team, T *dest, + const T *source, int nelems); + + __device__ void internal_sync(int pe, int PE_start, int stride, int PE_size, + int64_t *pSync); + + __device__ void internal_direct_barrier(int pe, int PE_start, int stride, + int n_pes, int64_t *pSync); + + __device__ void internal_atomic_barrier(int pe, int PE_start, int stride, + int n_pes, int64_t *pSync); + + //Temporary scratchpad memory used by internal barrier algorithms. + int64_t *barrier_sync{nullptr}; }; } // namespace rocshmem diff --git a/src/ipc/context_ipc_device_coll.cpp b/src/ipc/context_ipc_device_coll.cpp new file mode 100644 index 0000000000..804be2efbf --- /dev/null +++ b/src/ipc/context_ipc_device_coll.cpp @@ -0,0 +1,120 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "roc_shmem/roc_shmem.hpp" +#include "../context_incl.hpp" +#include "context_ipc_tmpl_device.hpp" +#include "../util.hpp" +#include "ipc_team.hpp" + +namespace rocshmem { + +__device__ void IPCContext::internal_direct_barrier(int pe, int PE_start, + int stride, int n_pes, + int64_t *pSync) { + int64_t flag_val = 1; + if (pe == PE_start) { + // Go through all PE offsets (except current offset = 0) + // and wait until they all reach + for (size_t i = 1; i < n_pes; i++) { + wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val); + pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + threadfence_system(); + + // Announce to other PEs that all have reached + for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) { + put_nbi(&pSync[0], &flag_val, 1, j); + } + } else { + // Mark current PE offset as reached + size_t pe_offset = (pe - PE_start) / stride; + put_nbi(&pSync[pe_offset], &flag_val, 1, PE_start); + wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val); + pSync[0] = ROC_SHMEM_SYNC_VALUE; + threadfence_system(); + } +} + +__device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start, + int stride, int n_pes, + int64_t *pSync) { + int64_t flag_val = 1; + if (pe == PE_start) { + wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, (int64_t)(n_pes - 1)); + pSync[0] = ROC_SHMEM_SYNC_VALUE; + threadfence_system(); + + for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) { + put_nbi(&pSync[0], &flag_val, 1, j); + } + } else { + amo_add(&pSync[0], flag_val, PE_start); + wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val); + pSync[0] = ROC_SHMEM_SYNC_VALUE; + threadfence_system(); + } +} + +// Uses PE values that are relative to world +__device__ void IPCContext::internal_sync(int pe, int PE_start, int stride, + int PE_size, int64_t *pSync) { + __syncthreads(); + if (is_thread_zero_in_block()) { + if (PE_size < 64) { + internal_direct_barrier(pe, PE_start, stride, PE_size, pSync); + } else { + internal_atomic_barrier(pe, PE_start, stride, PE_size, pSync); + } + } + __threadfence(); + __syncthreads(); +} + +__device__ void IPCContext::sync(roc_shmem_team_t team) { + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2. + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe = team_obj->my_pe_in_world; + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_stride = (1 << log_pe_stride); + int pe_size = team_obj->num_pes; + + internal_sync(pe, pe_start, pe_stride, pe_size, barrier_sync); +} + +__device__ void IPCContext::sync_all() { + internal_sync(my_pe, 0, 1, num_pes, barrier_sync); +} + +__device__ void IPCContext::barrier_all() { + if (is_thread_zero_in_block()) { + quiet(); + } + sync_all(); + __syncthreads(); +} + +} // namespace rocshmem diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index 4bf652781e..af79cd20b1 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -27,6 +27,7 @@ #include "roc_shmem/roc_shmem.hpp" #include "context_ipc_device.hpp" #include "../util.hpp" +#include "ipc_team.hpp" namespace rocshmem { @@ -150,11 +151,48 @@ __device__ void IPCContext::to_all(T *dest, const T *source, int nreduce, long *pSync) { // NOLINT(runtime/int) } +template +__device__ void IPCContext::internal_put_broadcast( + T *dst, const T *src, int nelems, int pe_root, int pe_start, + int log_pe_stride, int pe_size, + [[maybe_unused]] long *p_sync) { // NOLINT(runtime/int) + if (my_pe == pe_root) { + int stride = 1 << log_pe_stride; + int finish = pe_start + stride * pe_size; + for (int i = pe_start; i < finish; i += stride) { + if (i != my_pe) { + put_nbi_wg(dst, src, nelems, i); + } + } + } +} + +template +__device__ void IPCContext::internal_get_broadcast( + T *dst, const T *src, int nelems, int pe_root, + [[maybe_unused]] long *pSync) { // NOLINT(runtime/int) + if (my_pe != pe_root) { + get_wg(dst, src, nelems, pe_root); + } +} + template __device__ void IPCContext::broadcast(roc_shmem_team_t team, T *dst, const T *src, int nelems, int pe_root) { - //broadcast(dst, src, nelems, pe_root_world, pe_start, log_pe_stride, - // pe_size, p_sync); + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2 . + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_size = team_obj->tinfo_wrt_world->size; + long *p_sync = team_obj->bcast_pSync; + + // Passed pe_root is relative to team, convert to world root + int pe_root_world = team_obj->get_pe_in_world(pe_root); + broadcast(dst, src, nelems, pe_root_world, pe_start, log_pe_stride, + pe_size, p_sync); } template @@ -162,16 +200,88 @@ __device__ void IPCContext::broadcast(T *dst, const T *src, int nelems, int pe_root, int pe_start, int log_pe_stride, int pe_size, long *p_sync) { // NOLINT(runtime/int) + if (num_pes < 4) { + internal_put_broadcast(dst, src, nelems, pe_root, pe_start, log_pe_stride, + pe_size, p_sync); + } else { + internal_get_broadcast(dst, src, nelems, pe_root, p_sync); + } + + // Synchronize on completion of broadcast + internal_sync(my_pe, pe_start, (1 << log_pe_stride), pe_size, p_sync); } template __device__ void IPCContext::alltoall(roc_shmem_team_t team, T *dst, const T *src, int nelems) { + alltoall_linear(team, dst, src, nelems); +} + +template +__device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2 + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_size = team_obj->num_pes; + int stride = 1 << log_pe_stride; + long *pSync = team_obj->alltoall_pSync; + int my_pe_in_team = team_obj->my_pe; + + /* TODO (egabriel): shouldn't the loop handle pe_start, pe_size, and stride? + * otherwise what is the point of the few lines above? + */ + // Have each PE put their designated data to the other PEs + for (int j = 0; j < pe_size; j++) { + int dest_pe = team_obj->get_pe_in_world(j); + put_nbi_wg(&dst[my_pe_in_team * nelems], &src[j * nelems], nelems, dest_pe); + } + if (is_thread_zero_in_block()) { + quiet(); + } + // wait until everyone has obtained their designated data + internal_sync(my_pe, pe_start, stride, pe_size, pSync); } template __device__ void IPCContext::fcollect(roc_shmem_team_t team, T *dst, const T *src, int nelems) { + fcollect_linear(team, dst, src, nelems); +} + +template +__device__ void IPCContext::fcollect_linear(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2. + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_size = team_obj->num_pes; + int stride = 1 << log_pe_stride; + long *pSync = team_obj->alltoall_pSync; + int my_pe_in_team = team_obj->my_pe; + + /* TODO (egabriel): shouldn't the loop handle pe_start, pe_size, and stride? + * otherwise what is the point of the few lines above? + */ + // Have each PE put their designated data to the other PEs + for (int j = 0; j < pe_size; j++) { + int dest_pe = team_obj->get_pe_in_world(j); + put_nbi_wg(&dst[my_pe_in_team * nelems], src, nelems, dest_pe); + } + + if (is_thread_zero_in_block()) { + quiet(); + } + // wait until everyone has obtained their designated data + internal_sync(my_pe, pe_start, stride, pe_size, pSync); } // Block/wave functions From 45a8cb33548cb7f0ecd7a07c6c3f52b1d5260e65 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 28 Aug 2024 08:14:38 -0700 Subject: [PATCH 15/20] Update IPC object * Update the IPC object in the context class with the instance created in the IPC backend --- src/ipc/context_ipc_device.cpp | 39 +++++++++++++++++----------------- src/ipc/context_ipc_device.hpp | 3 --- 2 files changed, 20 insertions(+), 22 deletions(-) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index 7b9f3dd469..c1fa885c99 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -39,7 +39,8 @@ namespace rocshmem { __host__ IPCContext::IPCContext(Backend *b) : Context(b, false) { IPCBackend *backend{static_cast(b)}; - ipcImpl = &backend->ipcImpl; + ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases; + ipcImpl_.shm_size = b->ipcImpl.shm_size; auto *bp{backend->ipc_backend_proxy.get()}; @@ -59,21 +60,21 @@ __device__ void IPCContext::ctx_destroy(){ __device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy(ipcImpl->ipc_bases[local_pe] + L_offset, + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems); + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } __device__ void IPCContext::putmem_nbi(void *dest, const void *source, @@ -115,10 +116,10 @@ __device__ void IPCContext::sync(roc_shmem_team_t team) { __device__ void IPCContext::putmem_wg(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wg(ipcImpl->ipc_bases[local_pe] + L_offset, + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); __syncthreads(); } @@ -126,11 +127,11 @@ __device__ void IPCContext::putmem_wg(void *dest, const void *source, __device__ void IPCContext::getmem_wg(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wg(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems); + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); __syncthreads(); } @@ -147,21 +148,21 @@ __device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, __device__ void IPCContext::putmem_wave(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wave(ipcImpl->ipc_bases[local_pe] + L_offset, + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem_wave(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wave(dest, ipcImpl->ipc_bases[local_pe] + L_offset, + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index afadb9828c..377790ad67 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -234,9 +234,6 @@ class IPCContext : public Context { private: - //context class has IpcImpl object (ipcImpl_) - IpcImpl *ipcImpl{nullptr}; - uint64_t* atomic_base_ptr{nullptr}; char* g_ret; From e1e1ac6df6be452d3ca2e6a030ab44b85e437d09 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 28 Aug 2024 08:30:46 -0700 Subject: [PATCH 16/20] Add atomics * Add atomic_add, atomic_set, atomic_cas, atomic_fetch_add and atomic_fetch_cas to IPC backend --- src/ipc/context_ipc_tmpl_device.hpp | 44 +++++++++++++++++++++-------- 1 file changed, 32 insertions(+), 12 deletions(-) diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index 28cd718f0a..5e697cafaf 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -70,13 +70,21 @@ __device__ void IPCContext::get_nbi(T *dest, const T *source, size_t nelems, // Atomics template -__device__ void IPCContext::amo_add(void *dst, T value, int pe) { - assert(false); +__device__ void IPCContext::amo_add(void *dest, T value, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOAdd( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); } template -__device__ void IPCContext::amo_set(void *dst, T value, int pe) { - assert(false); +__device__ void IPCContext::amo_set(void *dest, T value, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOSet( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); } template @@ -119,20 +127,32 @@ __device__ void IPCContext::amo_xor(void *dst, T value, int pe) { } template -__device__ void IPCContext::amo_cas(void *dst, T value, T cond, int pe) { - assert(false); +__device__ void IPCContext::amo_cas(void *dest, T value, T cond, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOCas( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + value); } template -__device__ T IPCContext::amo_fetch_add(void *dst, T value, int pe) { - assert(false); - return 0; +__device__ T IPCContext::amo_fetch_add(void *dest, T value, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + return ipcImpl_.ipcAMOFetchAdd( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); } template -__device__ T IPCContext::amo_fetch_cas(void *dst, T value, T cond, int pe) { - assert(false); - return 0; +__device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + return ipcImpl_.ipcAMOFetchCas( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + value); } // Collectives From aae62954602f134d6ed92dabda96731dc4683374 Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Wed, 28 Aug 2024 09:41:05 -0700 Subject: [PATCH 17/20] ipc/context_ipc_device.cpp: set barrier_sync set the barrier_sync variable on the context during object creation --- src/ipc/context_ipc_device.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index 48d7dd9b4f..ede7633b06 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -43,6 +43,7 @@ __host__ IPCContext::IPCContext(Backend *b) auto *bp{backend->ipc_backend_proxy.get()}; + barrier_sync = backend->barrier_sync; g_ret = bp->g_ret; atomic_base_ptr = bp->atomic_ret->atomic_base_ptr; } From 1ffe62bb2e747f7136a0c1c555fca93c77599be8 Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Wed, 28 Aug 2024 09:42:14 -0700 Subject: [PATCH 18/20] utils.hpp: silence warnings silence some warnings when compiling with debug flags. --- src/util.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/util.hpp b/src/util.hpp index 67fdac6786..75687a0ac8 100644 --- a/src/util.hpp +++ b/src/util.hpp @@ -212,7 +212,7 @@ __device__ void gpu_dprintf(const char* fmt, const Args&... args) { while (atomicCAS(print_lock, 0, 1) == 1) { } - printf("WG (%lu, %lu, %lu) TH (%lu, %lu, %lu) ", hipBlockIdx_x, + printf("WG (%u, %u, %u) TH (%u, %u, %u) ", hipBlockIdx_x, hipBlockIdx_y, hipBlockIdx_z, hipThreadIdx_x, hipThreadIdx_y, hipThreadIdx_z); printf(fmt, args...); From 7bbf34d33446f253917b00a9dd237916c0940957 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 5 Sep 2024 11:52:00 -0700 Subject: [PATCH 19/20] remove local_pe calculation from puts, gets and atomics functions * All the PEs are assumed to be accessible using IPC backend --- src/ipc/context_ipc_device.cpp | 24 ++++++------------------ src/ipc/context_ipc_tmpl_device.hpp | 15 +++++---------- 2 files changed, 11 insertions(+), 28 deletions(-) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index c1fa885c99..4bf7072aa5 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -59,22 +59,18 @@ __device__ void IPCContext::ctx_destroy(){ __device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); } __device__ void IPCContext::putmem_nbi(void *dest, const void *source, @@ -115,23 +111,19 @@ __device__ void IPCContext::sync(roc_shmem_team_t team) { __device__ void IPCContext::putmem_wg(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[pe] + L_offset, const_cast(source), nelems); __syncthreads(); } __device__ void IPCContext::getmem_wg(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); __syncthreads(); } @@ -147,22 +139,18 @@ __device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, __device__ void IPCContext::putmem_wave(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem_wave(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); } diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index 5e697cafaf..94ef855736 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -71,20 +71,18 @@ __device__ void IPCContext::get_nbi(T *dest, const T *source, size_t nelems, // Atomics template __device__ void IPCContext::amo_add(void *dest, T value, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; ipcImpl_.ipcAMOAdd( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); } template __device__ void IPCContext::amo_set(void *dest, T value, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; ipcImpl_.ipcAMOSet( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); } template @@ -128,30 +126,27 @@ __device__ void IPCContext::amo_xor(void *dst, T value, int pe) { template __device__ void IPCContext::amo_cas(void *dest, T value, T cond, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; ipcImpl_.ipcAMOCas( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), cond, value); } template __device__ T IPCContext::amo_fetch_add(void *dest, T value, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; return ipcImpl_.ipcAMOFetchAdd( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); } template __device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; return ipcImpl_.ipcAMOFetchCas( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), cond, value); } From dfcacdc4a389a0bc32a28438a293981736127f7c Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Mon, 9 Sep 2024 07:01:20 -0700 Subject: [PATCH 20/20] remove pSync from internal_bcast functions remove the pSync arguments from the internal_broadcast functions, they are not used anyway. --- src/ipc/context_ipc_device.hpp | 6 ++---- src/ipc/context_ipc_tmpl_device.hpp | 16 ++++------------ 2 files changed, 6 insertions(+), 16 deletions(-) diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index 01fb1c1fb8..4cc6d08d56 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -207,13 +207,11 @@ class IPCContext : public Context { template __device__ void internal_put_broadcast(T *dst, const T *src, int nelems, int pe_root, int PE_start, - int logPE_stride, int PE_size, - long *pSync); // NOLINT(runtime/int) + int logPE_stride, int PE_size); // NOLINT(runtime/int) template __device__ void internal_get_broadcast(T *dst, const T *src, int nelems, - int pe_root, - long *pSync); // NOLINT(runtime/int) + int pe_root); // NOLINT(runtime/int) template __device__ void fcollect_linear(roc_shmem_team_t team, T *dest, diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index af79cd20b1..81ec0b2b97 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -154,8 +154,7 @@ __device__ void IPCContext::to_all(T *dest, const T *source, int nreduce, template __device__ void IPCContext::internal_put_broadcast( T *dst, const T *src, int nelems, int pe_root, int pe_start, - int log_pe_stride, int pe_size, - [[maybe_unused]] long *p_sync) { // NOLINT(runtime/int) + int log_pe_stride, int pe_size) { // NOLINT(runtime/int) if (my_pe == pe_root) { int stride = 1 << log_pe_stride; int finish = pe_start + stride * pe_size; @@ -169,8 +168,7 @@ __device__ void IPCContext::internal_put_broadcast( template __device__ void IPCContext::internal_get_broadcast( - T *dst, const T *src, int nelems, int pe_root, - [[maybe_unused]] long *pSync) { // NOLINT(runtime/int) + T *dst, const T *src, int nelems, int pe_root) { // NOLINT(runtime/int) if (my_pe != pe_root) { get_wg(dst, src, nelems, pe_root); } @@ -202,9 +200,9 @@ __device__ void IPCContext::broadcast(T *dst, const T *src, int nelems, long *p_sync) { // NOLINT(runtime/int) if (num_pes < 4) { internal_put_broadcast(dst, src, nelems, pe_root, pe_start, log_pe_stride, - pe_size, p_sync); + pe_size); } else { - internal_get_broadcast(dst, src, nelems, pe_root, p_sync); + internal_get_broadcast(dst, src, nelems, pe_root); } // Synchronize on completion of broadcast @@ -232,9 +230,6 @@ __device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst, long *pSync = team_obj->alltoall_pSync; int my_pe_in_team = team_obj->my_pe; - /* TODO (egabriel): shouldn't the loop handle pe_start, pe_size, and stride? - * otherwise what is the point of the few lines above? - */ // Have each PE put their designated data to the other PEs for (int j = 0; j < pe_size; j++) { int dest_pe = team_obj->get_pe_in_world(j); @@ -268,9 +263,6 @@ __device__ void IPCContext::fcollect_linear(roc_shmem_team_t team, T *dst, long *pSync = team_obj->alltoall_pSync; int my_pe_in_team = team_obj->my_pe; - /* TODO (egabriel): shouldn't the loop handle pe_start, pe_size, and stride? - * otherwise what is the point of the few lines above? - */ // Have each PE put their designated data to the other PEs for (int j = 0; j < pe_size; j++) { int dest_pe = team_obj->get_pe_in_world(j);