Swdev/536571 with additional issues found for other various missing includes (#158)

* Revert "SWDEV-536571 - Include assert header. (#157)"

This reverts commit 87d2efa430.

* Fix use of assert/abort and required includes

* Disable IPC AMO testers for non-implemented functions

[ROCm/rocshmem commit: 551603829c]

projects/rocshmem/src/backend_bc.cpp

@@ -33,6 +33,8 @@
 #include "ipc/backend_ipc.hpp"
 #endif
 
+#include <cassert>
+
 namespace rocshmem {
 
 #define NET_CHECK(cmd)                                       \

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 <cassert>
-
 #include "array.hpp"
 #include "../constants.hpp"
 
+#include <hip/hip_runtime.h>
+#include <cassert>
+
 namespace rocshmem {
 
 extern __constant__ int* GLOBAL_DEVICE_PRINT_LOCK;

projects/rocshmem/src/containers/atomic_wf_queue_impl.hpp

@@ -25,6 +25,9 @@
 #include "atomic_wf_queue.hpp"
 #include <iostream>
 
+#include <hip/hip_runtime.h>
+#include <cassert>
+
 namespace rocshmem {
 
 /*****************************************************************************

projects/rocshmem/src/containers/index_strategy.hpp

@@ -133,6 +133,7 @@
 
 #include <hip/hip_runtime.h>
 #include <cassert>
+
 #include <algorithm>
 
 namespace rocshmem {

projects/rocshmem/src/containers/matrix.hpp

@@ -25,12 +25,12 @@
 #ifndef LIBRARY_SRC_CONTAINERS_MATRIX_HPP_
 #define LIBRARY_SRC_CONTAINERS_MATRIX_HPP_
 
-#include <hip/hip_runtime.h>
+#include "index_strategy.hpp"
+#include "memory_allocator.hpp"
 
 #include <vector>
 
-#include "index_strategy.hpp"
-#include "memory_allocator.hpp"
+#include <hip/hip_runtime.h>
 
 namespace rocshmem {
 

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 <vector>
-
 #include "matrix.hpp"
 
+#include <vector>
+#include <hip/hip_runtime.h>
+#include <cassert>
+
 namespace rocshmem {
 
 template <typename TYPE>

projects/rocshmem/src/containers/share_strategy.cpp

@@ -25,6 +25,8 @@
 #include "share_strategy.hpp"
 #include "../constants.hpp"
 
+#include <hip/hip_runtime.h>
+
 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;
   }
 }

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;
     }
   }

projects/rocshmem/src/device_proxy.hpp

@@ -26,9 +26,10 @@
 #define LIBRARY_SRC_DEVICE_PROXY_HPP_
 
 #include <hip/hip_runtime.h>
-
 #include <memory>
 #include <utility>
+#include <cstring>
+#include <cassert>
 
 namespace rocshmem {
 

projects/rocshmem/src/host/host.cpp

@@ -31,6 +31,8 @@
 #include "../memory/window_info.hpp"
 #include "../util.hpp"
 
+#include <cassert>
+
 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;
 }
 

projects/rocshmem/src/host/host_helpers.hpp

@@ -28,6 +28,8 @@
 #include "host.hpp"
 #include "../memory/window_info.hpp"
 
+#include <cassert>
+
 namespace rocshmem {
 
 __host__ inline MPI_Aint HostInterface::compute_offset(

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 <utility>
-
 #include "rocshmem_config.h"  // NOLINT(build/include_subdir)
 #include "host_helpers.hpp"
 #include "../memory/window_info.hpp"
 #include "../team.hpp"
 
+#include <utility>
+#include <cassert>
+
 namespace rocshmem {
 
 template <typename T>

projects/rocshmem/src/ipc/backend_ipc.cpp

@@ -25,6 +25,10 @@
 #include "backend_ipc.hpp"
 #include "ipc_team.hpp"
 
+#include <hip/hip_runtime.h>
+#include <cstdlib>
+#include <cassert>
+
 namespace rocshmem {
 
 #define NET_CHECK(cmd)                                       \

projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp

@@ -32,6 +32,8 @@
 #include "ipc_team.hpp"
 #include "../rocshmem_calc.hpp"
 
+#include <hip/hip_runtime.h>
+
 namespace rocshmem {
 
 /******************************************************************************
@@ -44,13 +46,13 @@ __device__ void IPCContext::p(T *dest, T value, int pe) {
 
 template <typename T>
 __device__ void IPCContext::put(T *dest, const T *source, size_t nelems,
-				int pe) {
+                                int pe) {
   putmem(dest, source, nelems * sizeof(T), pe);
 }
 
 template <typename T>
 __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 <typename T>
 __device__ void IPCContext::get(T *dest, const T *source, size_t nelems,
-				int pe) {
+                                int pe) {
   getmem(dest, source, sizeof(T) * nelems, pe);
 }
 
 template <typename T>
 __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 <typename T>
 __device__ T IPCContext::amo_swap(void *dst, T value, int pe) {
-  assert(false);
+  printf("IPC amo_swap not implemented\n");
+  abort();
   return 0;
 }
 
 template <typename T>
 __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 <typename T>
 __device__ void IPCContext::amo_and(void *dst, T value, int pe) {
-  assert(false);
+  printf("IPC amo_and not implemented\n");
+  abort();
 }
 
 template <typename T>
 __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 <typename T>
 __device__ void IPCContext::amo_or(void *dst, T value, int pe) {
-  assert(false);
+  printf("IPC amo_or not implemented\n");
+  abort();
 }
 
 template <typename T>
 __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 <typename T>
 __device__ void IPCContext::amo_xor(void *dst, T value, int pe) {
-  assert(false);
+  printf("IPC amo_xor not implemented\n");
+  abort();
 }
 
 template <typename T>
@@ -443,7 +452,7 @@ __device__ void IPCContext::internal_get_broadcast(
 
 template <typename T>
 __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<IPCTeam *>(team);
 
   int stride = team_obj->tinfo_wrt_world->stride;
@@ -459,9 +468,9 @@ __device__ void IPCContext::broadcast(rocshmem_team_t team, T *dst,
 
 template <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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 <typename T>
 __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);
 }
 

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 <hip/hip_runtime_api.h>
 
 #include <cstdlib>
 #include <limits>
 
-#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)

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;
     }
 

projects/rocshmem/src/memory/remote_heap_info.hpp

@@ -28,7 +28,6 @@
 #include <hip/hip_runtime_api.h>
 #include <mpi.h>
 
-#include <cassert>
 #include <vector>
 
 #include "hip_allocator.hpp"

projects/rocshmem/src/reverse_offload/backend_ro.cpp

@@ -28,6 +28,7 @@
 #include <smmintrin.h>
 #include <unistd.h>
 
+#include <cassert>
 #include <cstdio>
 #include <cstdlib>
 #include <memory>

projects/rocshmem/src/reverse_offload/mpi_transport.cpp

@@ -28,6 +28,7 @@
 #include <utility>
 #include <vector>
 #include <unistd.h>
+#include <cassert>
 
 #include "../host/host.hpp"
 #include "backend_ro.hpp"

projects/rocshmem/src/reverse_offload/mpi_transport.hpp

@@ -29,6 +29,7 @@
 #include <mutex>  // NOLINT
 #include <queue>
 #include <vector>
+#include <thread>
 
 #include "queue.hpp"
 #include "transport.hpp"

projects/rocshmem/src/rocshmem.cpp

@@ -33,11 +33,6 @@
 
 #include "rocshmem/rocshmem.hpp"
 
-#include <cstdlib>
-#include <cstring>
-#include <functional>
-#include <random>
-
 #include "backend_bc.hpp"
 #include "context_incl.hpp"
 #ifdef USE_RO
@@ -53,6 +48,11 @@
 #include "util.hpp"
 #include "bootstrap/bootstrap.hpp"
 
+#include <cstdlib>
+#include <cstring>
+#include <functional>
+#include <random>
+#include <cassert>
 #include <unistd.h>
 
 namespace rocshmem {

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 <hip/hip_runtime.h>
-
 #include "../device_proxy.hpp"
 
+#include <hip/hip_runtime.h>
+
 namespace rocshmem {
 
 class ABQLBlockMutex {

projects/rocshmem/src/team_tracker.cpp

@@ -23,8 +23,9 @@
  *****************************************************************************/
 
 #include "team_tracker.hpp"
-#include <cassert>
+
 #include <cstdlib>
+#include <cassert>
 
 namespace rocshmem {
 

projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp

@@ -75,9 +75,19 @@ void AMOBitwiseTester<T>::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 <typename T>
 void AMOBitwiseTester<T>::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<T>::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) {                                                \

projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp

@@ -75,6 +75,12 @@ void AMOExtendedTester<T>::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 <typename T>
 void AMOExtendedTester<T>::verifyResults(uint64_t size) {
   T ret;
@@ -89,6 +95,10 @@ void AMOExtendedTester<T>::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<T>::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;                                                            \

projects/rocshmem/tests/functional_tests/random_access_tester.cpp

@@ -25,6 +25,7 @@
 #include "random_access_tester.hpp"
 
 #include <rocshmem/rocshmem.hpp>
+#include <cassert>
 
 using namespace rocshmem;
 

projects/rocshmem/tests/functional_tests/team_ctx_infra_tester.cpp

@@ -24,10 +24,11 @@
 
 #include "team_ctx_infra_tester.hpp"
 
-#include <stdlib.h>
-
 #include <rocshmem/rocshmem.hpp>
 
+#include <cstdlib>
+#include <cassert>
+
 using namespace rocshmem;
 
 /* this constant should equal ROCSHMEM_MAX_NUM_TEAMS-1 */

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 <hip/hip_runtime.h>
+#include <cassert>
+
 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};
 

projects/rocshmem/tests/unit_tests/bitwise_gtest.hpp

@@ -29,13 +29,14 @@
 
 #include "gtest/gtest.h"
 
-#include <hip/hip_runtime.h>
-
 #include "../src/memory/hip_allocator.hpp"
 #include "containers/matrix.hpp"
 #include "containers/share_strategy.hpp"
 #include "containers/strategies.hpp"
 
+#include <hip/hip_runtime.h>
+#include <cassert>
+
 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:

projects/rocshmem/tests/unit_tests/index_strategy_gtest.hpp

@@ -27,10 +27,12 @@
 
 #include "gtest/gtest.h"
 
-#include <hip/hip_runtime_api.h>
-
 #include "../src/containers/index_strategy.hpp"
 #include "../src/memory/hip_allocator.hpp"
+#include "../src/util.hpp"
+
+#include <hip/hip_runtime_api.h>
+#include <cassert>
 
 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);

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 <hip/hip_runtime.h>
+#include <cassert>
+
+
 namespace rocshmem {
 
 __global__

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 <hip/hip_runtime.h>
+#include <cassert>
+
 namespace rocshmem {
 
 const uint32_t SIGNAL_OFFSET {67108864};

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 <hip/hip_runtime.h>
+#include <cassert>
+
 namespace rocshmem {
 
 const int WARP_SIZE = 64;

projects/rocshmem/tests/unit_tests/notifier_gtest.hpp

@@ -31,6 +31,9 @@
 #include "../src/memory/notifier.hpp"
 #include "../src/util.hpp"
 
+#include <hip/hip_runtime.h>
+#include <cassert>
+
 namespace rocshmem {
 
 /**