diff --git a/projects/rocshmem/src/backend_bc.cpp b/projects/rocshmem/src/backend_bc.cpp index cb9a484c47..d2a451dfe6 100644 --- a/projects/rocshmem/src/backend_bc.cpp +++ b/projects/rocshmem/src/backend_bc.cpp @@ -33,6 +33,8 @@ #include "ipc/backend_ipc.hpp" #endif +#include + namespace rocshmem { #define NET_CHECK(cmd) \ diff --git a/projects/rocshmem/src/containers/array_impl.hpp b/projects/rocshmem/src/containers/array_impl.hpp index 1e606c61ca..6697ace54a 100644 --- a/projects/rocshmem/src/containers/array_impl.hpp +++ b/projects/rocshmem/src/containers/array_impl.hpp @@ -25,11 +25,12 @@ #ifndef LIBRARY_SRC_CONTAINERS_ARRAY_IMPL_HPP_ #define LIBRARY_SRC_CONTAINERS_ARRAY_IMPL_HPP_ -#include - #include "array.hpp" #include "../constants.hpp" +#include +#include + namespace rocshmem { extern __constant__ int* GLOBAL_DEVICE_PRINT_LOCK; diff --git a/projects/rocshmem/src/containers/atomic_wf_queue_impl.hpp b/projects/rocshmem/src/containers/atomic_wf_queue_impl.hpp index eb53db9e29..95aeb127e2 100644 --- a/projects/rocshmem/src/containers/atomic_wf_queue_impl.hpp +++ b/projects/rocshmem/src/containers/atomic_wf_queue_impl.hpp @@ -25,6 +25,9 @@ #include "atomic_wf_queue.hpp" #include +#include +#include + namespace rocshmem { /***************************************************************************** diff --git a/projects/rocshmem/src/containers/index_strategy.hpp b/projects/rocshmem/src/containers/index_strategy.hpp index 682c7821d6..dfbb92e247 100644 --- a/projects/rocshmem/src/containers/index_strategy.hpp +++ b/projects/rocshmem/src/containers/index_strategy.hpp @@ -133,6 +133,7 @@ #include #include + #include namespace rocshmem { diff --git a/projects/rocshmem/src/containers/matrix.hpp b/projects/rocshmem/src/containers/matrix.hpp index 944a172d85..cd359536f3 100644 --- a/projects/rocshmem/src/containers/matrix.hpp +++ b/projects/rocshmem/src/containers/matrix.hpp @@ -25,12 +25,12 @@ #ifndef LIBRARY_SRC_CONTAINERS_MATRIX_HPP_ #define LIBRARY_SRC_CONTAINERS_MATRIX_HPP_ -#include +#include "index_strategy.hpp" +#include "memory_allocator.hpp" #include -#include "index_strategy.hpp" -#include "memory_allocator.hpp" +#include namespace rocshmem { diff --git a/projects/rocshmem/src/containers/matrix_impl.hpp b/projects/rocshmem/src/containers/matrix_impl.hpp index 745f6d98b6..fd3203159e 100644 --- a/projects/rocshmem/src/containers/matrix_impl.hpp +++ b/projects/rocshmem/src/containers/matrix_impl.hpp @@ -25,10 +25,12 @@ #ifndef LIBRARY_SRC_CONTAINERS_MATRIX_IMPL_HPP_ #define LIBRARY_SRC_CONTAINERS_MATRIX_IMPL_HPP_ -#include - #include "matrix.hpp" +#include +#include +#include + namespace rocshmem { template diff --git a/projects/rocshmem/src/containers/share_strategy.cpp b/projects/rocshmem/src/containers/share_strategy.cpp index 00cf4b7f6e..1e6ee557b7 100644 --- a/projects/rocshmem/src/containers/share_strategy.cpp +++ b/projects/rocshmem/src/containers/share_strategy.cpp @@ -25,6 +25,8 @@ #include "share_strategy.hpp" #include "../constants.hpp" +#include + namespace rocshmem { __device__ uint64_t Block::lane_id() { @@ -217,10 +219,10 @@ __device__ void ShareStrategy::syncthreads() const { __syncthreads(); return; case ShareStrategyEnum::DEVICE: - assert(false); + abort(); return; case ShareStrategyEnum::UNUSED: - assert(false); + abort(); return; } } diff --git a/projects/rocshmem/src/containers/share_strategy.hpp b/projects/rocshmem/src/containers/share_strategy.hpp index 7cdb75f72b..9d07adb4dd 100644 --- a/projects/rocshmem/src/containers/share_strategy.hpp +++ b/projects/rocshmem/src/containers/share_strategy.hpp @@ -179,10 +179,10 @@ class ShareStrategy { } return value; case ShareStrategyEnum::DEVICE: - assert(false); + abort(); return 0; case ShareStrategyEnum::UNUSED: - assert(false); + abort(); return 0; } } @@ -199,10 +199,10 @@ class ShareStrategy { case ShareStrategyEnum::BLOCK: return true; case ShareStrategyEnum::DEVICE: - assert(false); + abort(); return false; case ShareStrategyEnum::UNUSED: - assert(false); + abort(); return false; } } diff --git a/projects/rocshmem/src/device_proxy.hpp b/projects/rocshmem/src/device_proxy.hpp index 554f09a781..333328ac2c 100644 --- a/projects/rocshmem/src/device_proxy.hpp +++ b/projects/rocshmem/src/device_proxy.hpp @@ -26,9 +26,10 @@ #define LIBRARY_SRC_DEVICE_PROXY_HPP_ #include - #include #include +#include +#include namespace rocshmem { diff --git a/projects/rocshmem/src/host/host.cpp b/projects/rocshmem/src/host/host.cpp index edcd8953c1..41ddb957e0 100644 --- a/projects/rocshmem/src/host/host.cpp +++ b/projects/rocshmem/src/host/host.cpp @@ -31,6 +31,8 @@ #include "../memory/window_info.hpp" #include "../util.hpp" +#include + namespace rocshmem { __host__ HostContextWindowInfo::HostContextWindowInfo(MPI_Comm comm_world, @@ -45,6 +47,8 @@ __host__ HostContextWindowInfo::~HostContextWindowInfo() { WindowInfo* HostInterface::acquire_window_context() { auto index{find_avail_pool_entry()}; + /* Entry should have been available; consider this as an error. */ + assert(index >= 0); HostContextWindowInfo* acquired_win_info = host_window_context_pool_[index]; @@ -55,6 +59,8 @@ WindowInfo* HostInterface::acquire_window_context() { __host__ void HostInterface::release_window_context(WindowInfo* window_info) { auto index{find_win_info_in_pool(window_info)}; + /* Entry should have been present; consider this as an error. */ + assert(index >= 0); host_window_context_pool_[index]->mark_avail(); } @@ -65,8 +71,6 @@ int HostInterface::find_avail_pool_entry() { return i; } } - /* Entry should have been available; consider this as an error. */ - assert(false); return -1; } @@ -79,8 +83,6 @@ int HostInterface::find_win_info_in_pool(WindowInfo* window_info) { return i; } } - /* Entry should have been present; consider this as an error. */ - assert(false); return -1; } diff --git a/projects/rocshmem/src/host/host_helpers.hpp b/projects/rocshmem/src/host/host_helpers.hpp index 2a20d3f363..4da5e7e43c 100644 --- a/projects/rocshmem/src/host/host_helpers.hpp +++ b/projects/rocshmem/src/host/host_helpers.hpp @@ -28,6 +28,8 @@ #include "host.hpp" #include "../memory/window_info.hpp" +#include + namespace rocshmem { __host__ inline MPI_Aint HostInterface::compute_offset( diff --git a/projects/rocshmem/src/host/host_templates.hpp b/projects/rocshmem/src/host/host_templates.hpp index bb08cd7838..56b38dbe71 100644 --- a/projects/rocshmem/src/host/host_templates.hpp +++ b/projects/rocshmem/src/host/host_templates.hpp @@ -25,13 +25,14 @@ #ifndef LIBRARY_SRC_HOST_HOST_TEMPLATES_HPP_ #define LIBRARY_SRC_HOST_HOST_TEMPLATES_HPP_ -#include - #include "rocshmem_config.h" // NOLINT(build/include_subdir) #include "host_helpers.hpp" #include "../memory/window_info.hpp" #include "../team.hpp" +#include +#include + namespace rocshmem { template diff --git a/projects/rocshmem/src/ipc/backend_ipc.cpp b/projects/rocshmem/src/ipc/backend_ipc.cpp index ac1378fc6d..6f3f3673f1 100644 --- a/projects/rocshmem/src/ipc/backend_ipc.cpp +++ b/projects/rocshmem/src/ipc/backend_ipc.cpp @@ -25,6 +25,10 @@ #include "backend_ipc.hpp" #include "ipc_team.hpp" +#include +#include +#include + namespace rocshmem { #define NET_CHECK(cmd) \ diff --git a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp index 563bab9d7e..fa500b63f6 100644 --- a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp @@ -32,6 +32,8 @@ #include "ipc_team.hpp" #include "../rocshmem_calc.hpp" +#include + namespace rocshmem { /****************************************************************************** @@ -44,13 +46,13 @@ __device__ void IPCContext::p(T *dest, T value, int pe) { template __device__ void IPCContext::put(T *dest, const T *source, size_t nelems, - int pe) { + 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) { + int pe) { putmem_nbi(dest, source, sizeof(T) * nelems, pe); } @@ -63,13 +65,13 @@ __device__ T IPCContext::g(const T *source, int pe) { template __device__ void IPCContext::get(T *dest, const T *source, size_t nelems, - int pe) { + 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) { + int pe) { getmem_nbi(dest, source, sizeof(T) * nelems, pe); } @@ -92,41 +94,48 @@ __device__ void IPCContext::amo_set(void *dest, T value, int pe) { template __device__ T IPCContext::amo_swap(void *dst, T value, int pe) { - assert(false); + printf("IPC amo_swap not implemented\n"); + abort(); return 0; } template __device__ T IPCContext::amo_fetch_and(void *dst, T value, int pe) { - assert(false); + printf("IPC amo_fetch_and not implemented\n"); + abort(); return 0; } template __device__ void IPCContext::amo_and(void *dst, T value, int pe) { - assert(false); + printf("IPC amo_and not implemented\n"); + abort(); } template __device__ T IPCContext::amo_fetch_or(void *dst, T value, int pe) { - assert(false); + printf("IPC amo_fetch_or not implemented\n"); + abort(); return 0; } template __device__ void IPCContext::amo_or(void *dst, T value, int pe) { - assert(false); + printf("IPC amo_or not implemented\n"); + abort(); } template __device__ T IPCContext::amo_fetch_xor(void *dst, T value, int pe) { - assert(false); + printf("IPC amo_fetch_xor not implemented\n"); + abort(); return 0; } template __device__ void IPCContext::amo_xor(void *dst, T value, int pe) { - assert(false); + printf("IPC amo_xor not implemented\n"); + abort(); } template @@ -443,7 +452,7 @@ __device__ void IPCContext::internal_get_broadcast( template __device__ void IPCContext::broadcast(rocshmem_team_t team, T *dst, - const T *src, int nelems, int pe_root) { + const T *src, int nelems, int pe_root) { IPCTeam *team_obj = reinterpret_cast(team); int stride = team_obj->tinfo_wrt_world->stride; @@ -459,9 +468,9 @@ __device__ void IPCContext::broadcast(rocshmem_team_t team, T *dst, template __device__ void IPCContext::internal_broadcast(T *dst, const T *src, int nelems, - int pe_root, int pe_start, - int stride, int pe_size, - long *p_sync) { // NOLINT(runtime/int) + int pe_root, int pe_start, + int stride, int pe_size, + long *p_sync) { // NOLINT(runtime/int) if (num_pes < 4) { internal_put_broadcast(dst, src, nelems, pe_root, pe_start, stride, pe_size); @@ -475,7 +484,7 @@ __device__ void IPCContext::internal_broadcast(T *dst, const T *src, int nelems, template __device__ void IPCContext::alltoall(rocshmem_team_t team, T *dst, - const T *src, int nelems) { + const T *src, int nelems) { alltoall_linear(team, dst, src, nelems); } @@ -504,7 +513,7 @@ __device__ void IPCContext::alltoall_linear(rocshmem_team_t team, T *dst, template __device__ void IPCContext::fcollect(rocshmem_team_t team, T *dst, - const T *src, int nelems) { + const T *src, int nelems) { fcollect_linear(team, dst, src, nelems); } @@ -535,49 +544,49 @@ __device__ void IPCContext::fcollect_linear(rocshmem_team_t team, T *dst, // Block/wave functions template __device__ void IPCContext::put_wg(T *dest, const T *source, size_t nelems, - int pe) { + 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) { + 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) { + 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) { + 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) { + 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) { + 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) { + 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) { + size_t nelems, int pe) { getmem_nbi_wave(dest, source, nelems * sizeof(T), pe); } diff --git a/projects/rocshmem/src/memory/hip_allocator.hpp b/projects/rocshmem/src/memory/hip_allocator.hpp index 0fe18edf6f..ae67e2ad8f 100644 --- a/projects/rocshmem/src/memory/hip_allocator.hpp +++ b/projects/rocshmem/src/memory/hip_allocator.hpp @@ -31,14 +31,14 @@ * @brief Contains HIP wrapper class for memory allocator */ +#include "rocshmem_config.h" // NOLINT(build/include_subdir) +#include "memory_allocator.hpp" + #include #include #include -#include "rocshmem_config.h" // NOLINT(build/include_subdir) -#include "memory_allocator.hpp" - // `hipDeviceMallocUncached` was introduced at ROCm 5.5 #if (HIP_VERSION_MAJOR > 5) || \ (HIP_VERSION_MAJOR == 5 && HIP_VERSION_MINOR >= 5) diff --git a/projects/rocshmem/src/memory/notifier.hpp b/projects/rocshmem/src/memory/notifier.hpp index a78d304af5..6805b86cbb 100644 --- a/projects/rocshmem/src/memory/notifier.hpp +++ b/projects/rocshmem/src/memory/notifier.hpp @@ -57,7 +57,7 @@ class Notifier { return; } if constexpr (scope == detail::atomic::memory_scope_system) { - assert(false); + static_assert(false); return; } diff --git a/projects/rocshmem/src/memory/remote_heap_info.hpp b/projects/rocshmem/src/memory/remote_heap_info.hpp index e64a48c023..f941e146e8 100644 --- a/projects/rocshmem/src/memory/remote_heap_info.hpp +++ b/projects/rocshmem/src/memory/remote_heap_info.hpp @@ -28,7 +28,6 @@ #include #include -#include #include #include "hip_allocator.hpp" diff --git a/projects/rocshmem/src/reverse_offload/backend_ro.cpp b/projects/rocshmem/src/reverse_offload/backend_ro.cpp index 8195d351a0..c696fc5004 100644 --- a/projects/rocshmem/src/reverse_offload/backend_ro.cpp +++ b/projects/rocshmem/src/reverse_offload/backend_ro.cpp @@ -28,6 +28,7 @@ #include #include +#include #include #include #include diff --git a/projects/rocshmem/src/reverse_offload/mpi_transport.cpp b/projects/rocshmem/src/reverse_offload/mpi_transport.cpp index 53812b565d..1757137905 100644 --- a/projects/rocshmem/src/reverse_offload/mpi_transport.cpp +++ b/projects/rocshmem/src/reverse_offload/mpi_transport.cpp @@ -28,6 +28,7 @@ #include #include #include +#include #include "../host/host.hpp" #include "backend_ro.hpp" diff --git a/projects/rocshmem/src/reverse_offload/mpi_transport.hpp b/projects/rocshmem/src/reverse_offload/mpi_transport.hpp index c9bca9755e..7effbfc071 100644 --- a/projects/rocshmem/src/reverse_offload/mpi_transport.hpp +++ b/projects/rocshmem/src/reverse_offload/mpi_transport.hpp @@ -29,6 +29,7 @@ #include // NOLINT #include #include +#include #include "queue.hpp" #include "transport.hpp" diff --git a/projects/rocshmem/src/rocshmem.cpp b/projects/rocshmem/src/rocshmem.cpp index 94e42aa410..37a8826b1c 100644 --- a/projects/rocshmem/src/rocshmem.cpp +++ b/projects/rocshmem/src/rocshmem.cpp @@ -33,11 +33,6 @@ #include "rocshmem/rocshmem.hpp" -#include -#include -#include -#include - #include "backend_bc.hpp" #include "context_incl.hpp" #ifdef USE_RO @@ -53,6 +48,11 @@ #include "util.hpp" #include "bootstrap/bootstrap.hpp" +#include +#include +#include +#include +#include #include namespace rocshmem { diff --git a/projects/rocshmem/src/sync/abql_block_mutex.hpp b/projects/rocshmem/src/sync/abql_block_mutex.hpp index 6e2d97ed4e..7bf95edc4a 100644 --- a/projects/rocshmem/src/sync/abql_block_mutex.hpp +++ b/projects/rocshmem/src/sync/abql_block_mutex.hpp @@ -25,10 +25,10 @@ #ifndef LIBRARY_SRC_SYNC_ABQL_BLOCK_MUTEX_HPP_ #define LIBRARY_SRC_SYNC_ABQL_BLOCK_MUTEX_HPP_ -#include - #include "../device_proxy.hpp" +#include + namespace rocshmem { class ABQLBlockMutex { diff --git a/projects/rocshmem/src/team_tracker.cpp b/projects/rocshmem/src/team_tracker.cpp index 343b414164..4d6fc1bb53 100644 --- a/projects/rocshmem/src/team_tracker.cpp +++ b/projects/rocshmem/src/team_tracker.cpp @@ -23,8 +23,9 @@ *****************************************************************************/ #include "team_tracker.hpp" -#include + #include +#include namespace rocshmem { diff --git a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp index 692085b4af..318386a756 100644 --- a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp @@ -75,9 +75,19 @@ void AMOBitwiseTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, num_timed_msgs = loop; } +#if defined(USE_IPC) and not defined(USE_RO) +#define DISABLE_IPC_TEST 1 +#else +#define DISABLE_IPC_TEST 0 +#endif + template void AMOBitwiseTester::verifyResults(uint64_t size) { T ret; + if(DISABLE_IPC_TEST) { + printf("AMO binary ops not implemented for IPC: values were not verified\n"); + return; + } if (args.myid == 0) { T expected_val = 0; @@ -131,6 +141,10 @@ void AMOBitwiseTester::verifyResults(uint64_t size) { TestType type, ShmemContextType ctx_type) { \ __shared__ rocshmem_ctx_t ctx; \ int wg_id = get_flat_grid_id(); \ + if(DISABLE_IPC_TEST) { \ + printf("AMO binary ops not implemented for IPC: test was not run\n"); \ + return; \ + } \ rocshmem_wg_init(); \ rocshmem_wg_ctx_create(ctx_type, &ctx); \ if (hipThreadIdx_x == 0) { \ diff --git a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp index 8e0c696175..4fab12fcbe 100644 --- a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp @@ -75,6 +75,12 @@ void AMOExtendedTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, num_timed_msgs = loop; } +#if defined(USE_IPC) and not defined(USE_RO) +#define DISABLE_IPC_TEST 1 +#else +#define DISABLE_IPC_TEST 0 +#endif + template void AMOExtendedTester::verifyResults(uint64_t size) { T ret; @@ -89,6 +95,10 @@ void AMOExtendedTester::verifyResults(uint64_t size) { expected_val = 44; break; case AMO_SwapTestType: + if(DISABLE_IPC_TEST) { + printf("AMO Swap not implemented for IPC: values were not verified\n"); + return; + } expected_val = num_msgs / 2; break; default: @@ -136,6 +146,10 @@ void AMOExtendedTester::verifyResults(uint64_t size) { rocshmem_ctx_##TNAME##_atomic_set(ctx, (T *)r_buf, 44, 1); \ break; \ case AMO_SwapTestType: \ + if(DISABLE_IPC_TEST) { \ + printf("AMO Swap not implemented for IPC: test was not run\n"); \ + break; \ + } \ ret = rocshmem_ctx_##TNAME##_atomic_swap(ctx, (T *)r_buf, \ ret + 1, 1); \ break; \ diff --git a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp index 960046368d..0467c30b95 100644 --- a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp @@ -25,6 +25,7 @@ #include "random_access_tester.hpp" #include +#include using namespace rocshmem; diff --git a/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp b/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp index 73322eb109..d1847607bb 100644 --- a/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp @@ -24,10 +24,11 @@ #include "team_ctx_infra_tester.hpp" -#include - #include +#include +#include + using namespace rocshmem; /* this constant should equal ROCSHMEM_MAX_NUM_TEAMS-1 */ diff --git a/projects/rocshmem/tests/unit_tests/abql_block_mutex_gtest.hpp b/projects/rocshmem/tests/unit_tests/abql_block_mutex_gtest.hpp index b52ed0dcee..7b200b42ed 100644 --- a/projects/rocshmem/tests/unit_tests/abql_block_mutex_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/abql_block_mutex_gtest.hpp @@ -31,6 +31,9 @@ #include "../src/sync/abql_block_mutex.hpp" #include "../src/util.hpp" +#include +#include + namespace rocshmem { inline __device__ @@ -108,11 +111,7 @@ class ABQLBlockMutexTestFixture : public ::testing::Test { mutex_, counter_); - hipError_t return_code = hipStreamSynchronize(nullptr); - if (return_code != hipSuccess) { - printf("Failed in stream synchronize\n"); - assert(return_code == hipSuccess); - } + CHECK_HIP(hipStreamSynchronize(nullptr)); size_t number_threads {x_block_dim * x_grid_dim}; diff --git a/projects/rocshmem/tests/unit_tests/bitwise_gtest.hpp b/projects/rocshmem/tests/unit_tests/bitwise_gtest.hpp index ee480200a1..cb69b579a0 100644 --- a/projects/rocshmem/tests/unit_tests/bitwise_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/bitwise_gtest.hpp @@ -29,13 +29,14 @@ #include "gtest/gtest.h" -#include - #include "../src/memory/hip_allocator.hpp" #include "containers/matrix.hpp" #include "containers/share_strategy.hpp" #include "containers/strategies.hpp" +#include +#include + namespace rocshmem { /***************************************************************************** @@ -328,12 +329,7 @@ class BitwiseTestFixture : public ::testing::Test _warp_matrix, activate_lanes_bitfield); - hipError_t return_code = hipStreamSynchronize(nullptr); - if (return_code != hipSuccess) { - printf("Failed in stream synchronize\n"); - assert(return_code == hipSuccess); - } - + CHECK_HIP(hipStreamSynchronize(nullptr)); } protected: diff --git a/projects/rocshmem/tests/unit_tests/index_strategy_gtest.hpp b/projects/rocshmem/tests/unit_tests/index_strategy_gtest.hpp index 0cd9e7c3f2..6f66ecf6b5 100644 --- a/projects/rocshmem/tests/unit_tests/index_strategy_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/index_strategy_gtest.hpp @@ -27,10 +27,12 @@ #include "gtest/gtest.h" -#include - #include "../src/containers/index_strategy.hpp" #include "../src/memory/hip_allocator.hpp" +#include "../src/util.hpp" + +#include +#include namespace rocshmem { @@ -85,11 +87,7 @@ class IndexStrategyTestFixture : public ::testing::Test _raw_mem, _mem_elements); - hipError_t return_code = hipStreamSynchronize(nullptr); - if (return_code != hipSuccess) { - printf("Failed in stream synchronize\n"); - assert(return_code == hipSuccess); - } + CHECK_HIP(hipStreamSynchronize(nullptr)); for(size_t i = 0; i < _mem_elements; i++) { EXPECT_EQ(_raw_mem[i], 1); diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp index 23b46bd293..d78063234f 100644 --- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp @@ -34,6 +34,10 @@ #include "../src/memory/symmetric_heap.hpp" #include "../src/ipc_policy.hpp" +#include +#include + + namespace rocshmem { __global__ diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp index d5204d5dd9..bb7b4fd31b 100644 --- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp @@ -36,6 +36,9 @@ #include "../src/memory/symmetric_heap.hpp" #include "../src/util.hpp" +#include +#include + namespace rocshmem { const uint32_t SIGNAL_OFFSET {67108864}; diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp index 842886bbc1..f65e5fbde2 100644 --- a/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp @@ -36,6 +36,9 @@ #include "../src/memory/symmetric_heap.hpp" #include "../src/util.hpp" +#include +#include + namespace rocshmem { const int WARP_SIZE = 64; diff --git a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp index 8d74fb08cc..8163ec9703 100644 --- a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp @@ -31,6 +31,9 @@ #include "../src/memory/notifier.hpp" #include "../src/util.hpp" +#include +#include + namespace rocshmem { /**