diff --git a/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt index 730b733..e34f629 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 0000000..a825a13 --- /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 0000000..492d63c --- /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 0000000..c90e0a3 --- /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_ diff --git a/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index 7130c2c..1f71119 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 0000000..c963a65 --- /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 0000000..fc39504 --- /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