From de750ddacc7ac6e8ce23a417373691d92fbe9202 Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Thu, 25 Jul 2024 14:48:51 -0700 Subject: [PATCH 1/2] 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 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_ From b1bc4a497a304c1f91f654dd3a2a19c0e0a6c2ca Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 30 Jul 2024 15:22:03 -0700 Subject: [PATCH 2/2] 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 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