Add 'projects/rocshmem/' from commit '0496586829058af5cfd7f23acda2a6d0040da584'

git-subtree-dir: projects/rocshmem
git-subtree-mainline: 5fd976da70
git-subtree-split: 0496586829

projects/rocshmem/tests/unit_tests/CMakeLists.txt

@@ -0,0 +1,131 @@
+###############################################################################
+# Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+#
+# SPDX-License-Identifier: MIT
+#
+# 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.
+###############################################################################
+
+cmake_minimum_required(VERSION 3.16.3 FATAL_ERROR)
+
+###############################################################################
+# PROJECT
+###############################################################################
+include(${CMAKE_SOURCE_DIR}/cmake/setup_project.cmake)
+project(rocshmem_unit_tests VERSION 1.0.0 LANGUAGES CXX)
+
+###############################################################################
+# SOURCES
+###############################################################################
+add_executable(${PROJECT_NAME} "")
+
+target_include_directories(
+  ${PROJECT_NAME}
+  PRIVATE
+    ${CMAKE_CURRENT_SOURCE_DIR}
+)
+
+target_sources(
+  ${PROJECT_NAME}
+  PRIVATE
+    shmem_gtest.cpp
+    heap_memory_gtest.cpp
+    hipmalloc_gtest.cpp
+    bin_gtest.cpp
+    binner_gtest.cpp
+    #bitwise_gtest.cpp # Test is disabled because of compilation errors
+    address_record_gtest.cpp
+    index_strategy_gtest.cpp
+    single_heap_gtest.cpp
+    symmetric_heap_gtest.cpp
+    pow2_bins_gtest.cpp
+    dlmalloc_gtest.cpp
+    remote_heap_info_gtest.cpp
+    mpi_instance_gtest.cpp
+    abql_block_mutex_gtest.cpp
+    notifier_gtest.cpp
+    free_list_gtest.cpp
+    wavefront_size_gtest.cpp
+    atomic_wf_queue_gtest.cpp
+    envvar_gtest.cpp
+)
+
+if (USE_IPC)
+  target_sources(
+    ${PROJECT_NAME}
+    PRIVATE
+      ipc_impl_simple_coarse_gtest.cpp
+      ipc_impl_simple_fine_gtest.cpp
+      ipc_impl_tiled_fine_gtest.cpp
+  )
+endif()
+
+###############################################################################
+# ROCSHMEM DEPENDENCY
+###############################################################################
+find_package(hip REQUIRED PATHS /opt/rocm)
+find_package(MPI REQUIRED)
+
+if (BUILD_TESTS_ONLY)
+  find_package(rocshmem REQUIRED PATHS /opt/rocm)
+
+  target_include_directories(
+    ${PROJECT_NAME}
+    PRIVATE
+      $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/..>
+      $<BUILD_INTERFACE:${CMAKE_BINARY_DIR}>
+  )
+endif()
+
+target_link_libraries(
+  ${PROJECT_NAME}
+  PRIVATE
+    MPI::MPI_CXX
+    roc::rocshmem
+)
+
+###############################################################################
+# GTEST DEPENDENCY
+###############################################################################
+include(FetchContent)
+
+FetchContent_Declare(
+    googletest
+    PREFIX extern/googletest
+    GIT_REPOSITORY https://github.com/google/googletest.git
+    GIT_TAG release-1.12.0
+)
+
+set(gtest_force_shared_crt ON CACHE BOOL "" FORCE)
+set(BUILD_GMOCK OFF CACHE BOOL "" FORCE)
+set(BUILD_GTEST ON CACHE BOOL "" FORCE)
+set(INSTALL_GTEST OFF CACHE BOOL "" FORCE)
+
+FetchContent_MakeAvailable(googletest)
+
+configure_file(../../scripts/unit_tests/driver.sh rocshmem_unit_driver.sh COPYONLY)
+rocm_install(PROGRAMS ${CMAKE_CURRENT_BINARY_DIR}/rocshmem_unit_driver.sh COMPONENT tests DESTINATION ${CMAKE_INSTALL_DATADIR}/rocshmem)
+rocm_install(TARGETS rocshmem_unit_tests COMPONENT tests)
+
+target_link_libraries(
+  ${PROJECT_NAME}
+  PRIVATE
+    gtest
+    gtest_main
+)

projects/rocshmem/tests/unit_tests/abql_block_mutex_gtest.cpp

@@ -0,0 +1,75 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "abql_block_mutex_gtest.hpp"
+
+using namespace rocshmem;
+
+/*****************************************************************************
+ ******************************* Fixture Tests *******************************
+ *****************************************************************************/
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_1) {
+  run_all_threads_once(1, 1);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_2) {
+  run_all_threads_once(1, 2);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_8) {
+  run_all_threads_once(1, 8);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_64) {
+  run_all_threads_once(1, 64);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_128) {
+  run_all_threads_once(1, 128);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_256) {
+  run_all_threads_once(1, 256);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_1024) {
+  run_all_threads_once(1, 1024);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_2048) {
+  run_all_threads_once(1, 2048);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_4096) {
+  run_all_threads_once(1, 4096);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_8192) {
+  run_all_threads_once(1, 8192);
+}
+
+TEST_F(ABQLBlockMutexTestFixture, run_all_threads_once_1_65536) {
+  run_all_threads_once(1, 65536);
+}

projects/rocshmem/tests/unit_tests/abql_block_mutex_gtest.hpp

@@ -0,0 +1,141 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_ABQL_BLOCK_MUTEX_GTEST_HPP
+#define ROCSHMEM_ABQL_BLOCK_MUTEX_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/sync/abql_block_mutex.hpp"
+#include "../src/util.hpp"
+
+#include <hip/hip_runtime.h>
+#include <cassert>
+
+namespace rocshmem {
+
+inline __device__
+void
+increment_counter(ABQLBlockMutex *mutex,
+                  size_t *counter) {
+    auto ticket {mutex->lock()};
+    (*counter)++;
+    __threadfence();
+    mutex->unlock(ticket);
+}
+
+__global__
+void
+all_threads_once(ABQLBlockMutex *mutex,
+                 size_t *counter) {
+    increment_counter(mutex, counter);
+}
+
+__global__
+void
+block_leader_once(ABQLBlockMutex *mutex,
+                  size_t *counter) {
+    if (is_thread_zero_in_block()) {
+        increment_counter(mutex, counter);
+    }
+}
+
+__global__
+void
+warp_leader_once(ABQLBlockMutex *mutex,
+                 size_t *counter) {
+    if (is_thread_zero_in_wave()) {
+        increment_counter(mutex, counter);
+    }
+}
+
+class ABQLBlockMutexTestFixture : public ::testing::Test {
+  public:
+    ABQLBlockMutexTestFixture() {
+        assert(mutex_ == nullptr);
+        hip_allocator_.allocate((void**)&mutex_, sizeof(ABQLBlockMutex));
+
+        assert(mutex_);
+        new (mutex_) ABQLBlockMutex();
+
+        assert(counter_ == nullptr);
+        hip_allocator_.allocate((void**)&counter_, sizeof(int));
+
+        assert(counter_);
+        *counter_ = 0;
+    }
+
+    ~ABQLBlockMutexTestFixture() {
+        if (mutex_) {
+            hip_allocator_.deallocate(mutex_);
+        }
+
+        if (counter_) {
+            hip_allocator_.deallocate(counter_);
+        }
+    }
+
+    void
+    run_all_threads_once(uint32_t x_block_dim,
+                         uint32_t x_grid_dim) {
+        const dim3 hip_blocksize(x_block_dim, 1, 1);
+        const dim3 hip_gridsize(x_grid_dim, 1, 1);
+
+        hipLaunchKernelGGL(all_threads_once,
+                           hip_gridsize,
+                           hip_blocksize,
+                           0,
+                           nullptr,
+                           mutex_,
+                           counter_);
+
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+
+        size_t number_threads {x_block_dim * x_grid_dim};
+
+        ASSERT_EQ(*counter_, number_threads);
+    }
+
+  protected:
+    /**
+     * @brief An allocator to create objects in device memory.
+     */
+    HIPAllocator hip_allocator_ {};
+
+    /**
+     * @brief A mutex to prevent data races.
+     */
+    ABQLBlockMutex *mutex_ {nullptr};
+
+    /**
+     * @brief A monotonically increasing counter to track accesses.
+     */
+    size_t *counter_ {nullptr};
+};
+
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_ABQL_BLOCK_MUTEX_GTEST_HPP

projects/rocshmem/tests/unit_tests/address_record_gtest.cpp

@@ -0,0 +1,90 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "address_record_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(AddressRecordTestFixture, split_entry) {
+  auto [ar1, ar2] = split_.split();
+
+  char* a1{reinterpret_cast<char*>(0x200)};
+  size_t size1{0x40};
+  ASSERT_EQ(ar1.get_address(), a1);
+  ASSERT_EQ(ar1.get_size(), size1);
+
+  char* a2{reinterpret_cast<char*>(0x240)};
+  size_t size2{0x40};
+  ASSERT_EQ(ar2.get_address(), a2);
+  ASSERT_EQ(ar2.get_size(), size2);
+}
+
+#ifdef NDEBUG
+TEST_F(AddressRecordTestFixture, DISABLED_split_bad_address)
+#else
+TEST_F(AddressRecordTestFixture, split_bad_address)
+#endif
+{
+  ASSERT_DEATH({ bad_addr_.split(); }, "");
+}
+
+#ifdef NDEBUG
+TEST_F(AddressRecordTestFixture, DISABLED_split_bad_size)
+#else
+TEST_F(AddressRecordTestFixture, split_bad_size)
+#endif
+{
+  ASSERT_DEATH({ bad_size_.split(); }, "");
+}
+
+TEST_F(AddressRecordTestFixture, combine_1_into_combine_2) {
+  AddressRecord ar{combine_1_.combine(combine_2_)};
+  ASSERT_EQ(ar.get_address(), combine_1_.get_address());
+  ASSERT_EQ(ar.get_size(), combine_1_.get_size() << 1);
+}
+
+#ifdef NDEBUG
+TEST_F(AddressRecordTestFixture, DISABLED_combine_nullptr_record)
+#else
+TEST_F(AddressRecordTestFixture, combine_nullptr_record)
+#endif
+{
+  ASSERT_DEATH({ AddressRecord ar = bad_addr_.combine(combine_2_); }, "");
+}
+
+#ifdef NDEBUG
+TEST_F(AddressRecordTestFixture, DISABLED_combine_differerent_sizes)
+#else
+TEST_F(AddressRecordTestFixture, combine_different_sizes)
+#endif
+{
+  char* a1{reinterpret_cast<char*>(0x120)};
+  char* a2{reinterpret_cast<char*>(0x140)};
+  size_t size1{0x20};
+  size_t size2{0x40};
+
+  AddressRecord ar1{a1, size1};
+  AddressRecord ar2{a2, size2};
+  ASSERT_DEATH({ AddressRecord ar = ar1.combine(ar2); }, "");
+}

projects/rocshmem/tests/unit_tests/address_record_gtest.hpp

@@ -0,0 +1,46 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_ADDRESS_RECORD_GTEST_HPP
+#define ROCSHMEM_ADDRESS_RECORD_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/address_record.hpp"
+
+namespace rocshmem {
+
+class AddressRecordTestFixture : public ::testing::Test
+{
+  protected:
+    AddressRecord split_ {reinterpret_cast<char*>(0x200), 0x80};
+    AddressRecord combine_1_ {reinterpret_cast<char*>(0x120), 0x20};
+    AddressRecord combine_2_ {reinterpret_cast<char*>(0x140), 0x20};
+    AddressRecord bad_addr_ {nullptr, 0x20};
+    AddressRecord bad_size_ {reinterpret_cast<char*>(0x800), 0x0};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_ADDRESS_RECORD_GTEST_HPP

projects/rocshmem/tests/unit_tests/atomic_wf_queue_gtest.cpp

@@ -0,0 +1,201 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "atomic_wf_queue_gtest.hpp"
+
+using namespace rocshmem;
+
+/*****************************************************************************
+ ******************************* Fixture Tests *******************************
+ *****************************************************************************/
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_4) {
+  get_thread_lane_ids(4);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_10) {
+  get_thread_lane_ids(10);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_16) {
+  get_thread_lane_ids(16);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_17) {
+  get_thread_lane_ids(17);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_64) {
+  get_thread_lane_ids(64);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_97) {
+  get_thread_lane_ids(97);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_256) {
+  get_thread_lane_ids(256);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_183) {
+  get_thread_lane_ids(183);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_1024) {
+  get_thread_lane_ids(1024);
+}
+
+TEST_F(AtomicWFQueueTestFixture, init_2_64) {
+  init_queue(2, 64);
+}
+
+TEST_F(AtomicWFQueueTestFixture, init_64_256) {
+  init_queue(64, 256);
+}
+
+TEST_F(AtomicWFQueueTestFixture, init_1024_96) {
+  init_queue(1024, 96);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_eq_qsize_2_32) {
+  int num_blocks {2};
+  int block_size {32};
+  int wf_size {this->wf_size};
+  int queue_size {num_blocks * ((block_size - 1) / wf_size + 1)};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_eq_qsize_32_192) {
+  int num_blocks {32};
+  int block_size {192};
+  int wf_size {this->wf_size};
+  int queue_size {num_blocks * ((block_size - 1) / wf_size + 1)};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_eq_qsize_64_256) {
+  int num_blocks {64};
+  int block_size {256};
+  int wf_size {this->wf_size};
+  int queue_size {num_blocks * ((block_size - 1) / wf_size + 1)};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_16_32_4) {
+  int num_blocks {16};
+  int block_size {32};
+  int queue_size {4};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_16_96_8) {
+  int num_blocks {16};
+  int block_size {96};
+  int queue_size {8};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_16_320_8) {
+  int num_blocks {16};
+  int block_size {320};
+  int queue_size {8};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_32_192_16) {
+  int num_blocks {32};
+  int block_size {192};
+  int queue_size {16};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_64_256_64) {
+  int num_blocks {64};
+  int block_size {256};
+  int queue_size {64};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_64_576_64) {
+  int num_blocks {64};
+  int block_size {576};
+  int queue_size {64};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_4_8_16) {
+  int num_blocks {4};
+  int block_size {8};
+  int queue_size {16};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_4_96_16) {
+  int num_blocks {4};
+  int block_size {96};
+  int queue_size {16};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_32_192_128) {
+  int num_blocks {32};
+  int block_size {192};
+  int queue_size {128};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_32_320_256) {
+  int num_blocks {32};
+  int block_size {320};
+  int queue_size {256};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_16_128) {
+  int num_blocks {64};
+  int block_size {16};
+  int queue_size {128};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_96_256) {
+  int num_blocks {64};
+  int block_size {96};
+  int queue_size {256};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_320_512) {
+  int num_blocks {64};
+  int block_size {320};
+  int queue_size {512};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_576_1024) {
+  int num_blocks {64};
+  int block_size {576};
+  int queue_size {1024};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}

projects/rocshmem/tests/unit_tests/atomic_wf_queue_gtest.hpp

@@ -0,0 +1,172 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_ATOMIC_WF_QUEUE_GTEST_HPP
+#define ROCSHMEM_ATOMIC_WF_QUEUE_GTEST_HPP
+
+#include "../src/containers/atomic_wf_queue_impl.hpp"
+#include "gtest/gtest.h"
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/util.hpp"
+#include <iostream>
+
+namespace rocshmem {
+
+template <typename AWFQueue>
+__global__ void wf_lane_ids(AWFQueue* awf_queue, unsigned int* device_array,
+                           int wf_size) {
+
+  int t_id {get_flat_id()};
+  int lane_id {t_id % wf_size};
+  device_array[t_id] = awf_queue->active_logical_lane_id();
+}
+
+template <typename AWFQueue>
+__global__ void concurrent_enqueue_dequeue(
+    AWFQueue* awf_queue,
+    unsigned int* device_array) {
+
+  int t_id {get_flat_id()};
+  int val {awf_queue->dequeue()};
+  device_array[t_id] = val;
+  awf_queue->enqueue(val);
+}
+
+class AtomicWFQueueTestFixture : public ::testing::Test {
+ public:
+  AtomicWFQueueTestFixture() {
+    awf_queue = awf_queue_proxy.get();
+
+    int device_id {};
+    hipDeviceProp_t device_props;
+    CHECK_HIP(hipGetDevice(&device_id));
+    CHECK_HIP(hipGetDeviceProperties(&device_props, device_id));
+    
+    wf_size = device_props.warpSize;
+  }
+
+  ~AtomicWFQueueTestFixture() {}
+
+  void get_thread_lane_ids(unsigned int num_threads = 4) {
+    
+    unsigned int *device_array {nullptr};
+
+    hip_allocator_.allocate(reinterpret_cast<void**>(&device_array),
+                            sizeof(unsigned int) * num_threads);
+  
+
+    hipLaunchKernelGGL(wf_lane_ids, 1, num_threads, 0, nullptr,
+                       awf_queue, device_array, wf_size);
+
+    CHECK_HIP(hipDeviceSynchronize());
+
+    hip_allocator_.deallocate(device_array);
+
+    for (unsigned int i{0}; i < num_threads; i++) {
+      EXPECT_EQ(device_array[i], i % wf_size);
+    }
+  }
+
+  void init_queue(int num_blocks, int block_size) {
+    int num_elems = num_blocks * ((block_size - 1) / wf_size + 1);
+
+    awf_queue->allocate_queue(num_elems);
+
+    for (int i{0}; i < num_elems; i++) {
+      awf_queue->push(i);
+    }
+
+    EXPECT_EQ(awf_queue->get_curr_size(), num_elems);
+    EXPECT_EQ(awf_queue->get_queue_size(), num_elems);
+    EXPECT_EQ(awf_queue->get_tail(), awf_queue->get_head());
+
+    awf_queue->deallocate_queue();
+  }
+
+  void verify(unsigned int *arr, int num_blocks, int block_size) {
+    unsigned int expected_val {};
+    unsigned int lane_id {};
+    unsigned int idx {};
+    for (unsigned int i{0}; i < num_blocks; i++) {
+      for (unsigned int j{0}; j < block_size; j++) {
+        idx = i * block_size + j;
+        lane_id = j % wf_size;
+        if (!lane_id) {
+          expected_val = arr[idx];
+        }
+        EXPECT_EQ(arr[idx], expected_val + lane_id);
+      }
+    }
+  }
+
+  void dequeue_enqueue(int num_blocks, int block_size, int queue_size) {
+
+    int num_threads {num_blocks * block_size};
+    unsigned int *device_array {nullptr};
+
+    hip_allocator_.allocate(reinterpret_cast<void**>(&device_array),
+                            sizeof(unsigned int) * num_threads);
+
+    awf_queue->allocate_queue(queue_size);
+
+    for (int i{0}; i < queue_size; i++) {
+      awf_queue->push(i);
+    }
+
+    EXPECT_EQ(awf_queue->get_curr_size(), queue_size);
+    EXPECT_EQ(awf_queue->get_queue_size(), queue_size);
+    EXPECT_EQ(awf_queue->get_tail(), awf_queue->get_head());
+
+    hipLaunchKernelGGL(concurrent_enqueue_dequeue, num_blocks, block_size, 0,
+                       nullptr, awf_queue, device_array);
+    CHECK_HIP(hipDeviceSynchronize());
+
+    EXPECT_EQ(awf_queue->get_curr_size(), queue_size);
+    EXPECT_EQ(awf_queue->get_tail(), awf_queue->get_head());
+
+    verify(device_array, num_blocks, block_size);
+
+    hip_allocator_.deallocate(device_array);
+    awf_queue->deallocate_queue();
+  }
+
+  protected:
+
+    AtomicWFQueueProxy<HIPAllocator, int> awf_queue_proxy{};
+    AtomicWFQueue<int, HIPAllocator>* awf_queue{};
+
+    /**
+     * @brief An allocator to create objects in device memory.
+     */
+    HIPAllocator hip_allocator_ {};
+
+    /**
+     * @brief Wavefront size.
+     */
+    int wf_size {};
+};
+
+}  // namespace rocshmem
+
+#endif  // ROCSHMEM_ATOMIC_WF_QUEUE_GTEST_HPP

projects/rocshmem/tests/unit_tests/bin_gtest.cpp

@@ -0,0 +1,54 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "bin_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(BinTestFixture, is_empty_check) { ASSERT_TRUE(bin_.empty()); }
+
+TEST_F(BinTestFixture, is_not_empty_check) {
+  bin_.put(nullptr);
+
+  ASSERT_FALSE(bin_.empty());
+}
+
+TEST_F(BinTestFixture, size_check) {
+  ASSERT_EQ(bin_.size(), 0);
+  bin_.put(nullptr);
+  ASSERT_EQ(bin_.size(), 1);
+  bin_.put(nullptr);
+  ASSERT_EQ(bin_.size(), 2);
+}
+
+TEST_F(BinTestFixture, retrieval_check) {
+  char* p_xa{reinterpret_cast<char*>(0xa)};
+  char* p_xb{reinterpret_cast<char*>(0xb)};
+  bin_.put(p_xa);
+  bin_.put(p_xb);
+  auto g_xb = bin_.get();
+  auto g_xa = bin_.get();
+  ASSERT_EQ(p_xa, g_xa);
+  ASSERT_EQ(p_xb, g_xb);
+}

projects/rocshmem/tests/unit_tests/bin_gtest.hpp

@@ -0,0 +1,45 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_BIN_GTEST_HPP
+#define ROCSHMEM_BIN_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/bin.hpp"
+
+namespace rocshmem {
+
+class BinTestFixture : public ::testing::Test
+{
+  protected:
+    /**
+     * @brief A bin object containing pointers
+     */
+    Bin<char*> bin_ {};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_BIN_GTEST_HPP

projects/rocshmem/tests/unit_tests/binner_gtest.cpp

@@ -0,0 +1,141 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "binner_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST(Binner, ffs_0) {
+  size_t val{0x0};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, UINT_MAX);
+}
+
+TEST(Binner, ffs_1) {
+  size_t val{0x1};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 0);
+}
+
+TEST(Binner, ffs_2) {
+  size_t val{0x2};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 1);
+}
+
+TEST(Binner, ffs_4) {
+  size_t val{0x4};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 2);
+}
+
+TEST(Binner, ffs_8) {
+  size_t val{0x8};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 3);
+}
+
+TEST(Binner, ffs_100) {
+  size_t val{0x100};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 8);
+}
+
+TEST(Binner, ffs_8000) {
+  size_t val{0x8000};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 15);
+}
+
+TEST(Binner, ffs_ff80) {
+  size_t val{0xFF80};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 15);
+}
+
+TEST(Binner, ffs_4000_0000) {
+  size_t val{0x40000000};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 30);
+}
+
+TEST(Binner, ffs_0100_0000_0000_0000) {
+  size_t val{0x0100000000000000};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 56);
+}
+
+TEST(Binner, ffs_0200_0000_0000_0000) {
+  size_t val{0x0200000000000000};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 57);
+}
+
+TEST(Binner, ffs_4000_0000_0000_0000) {
+  size_t val{0x4000000000000000};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 62);
+}
+
+TEST(Binner, ffs_8000_0000_0000_0000) {
+  size_t val{0x8000000000000000};
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 63);
+}
+
+TEST(Binner, ffs_max) {
+  size_t val = -1;
+  auto bit_pos{find_first_set_one(val)};
+  ASSERT_EQ(bit_pos, 63);
+}
+
+TEST_F(BinnerTestFixture, correct_bin_sizes_one_gig) {
+  auto bins{binner_.get_bins()};
+  ASSERT_EQ(bins->size(), 24);
+
+  std::array<unsigned, 24> a;
+  for (unsigned i = 0; i < 24; i++) {
+    a[i] = i + 7;
+  }
+
+  for (auto e : a) {
+    bool found = bins->count(std::pow(2, e));
+    ASSERT_TRUE(found);
+  }
+}
+
+TEST_F(BinnerTestFixture, bin_get_one_gig) {
+  binner_.assign_heap_to_bins();
+
+  auto bins{binner_.get_bins()};
+  ASSERT_EQ(bins->size(), 24);
+
+  size_t gibibyte = std::pow(2, 30);
+  auto bin{(*bins)[gibibyte]};
+
+  ASSERT_EQ(bin.size(), 1);
+
+  auto address_record = bin.get();
+  ASSERT_EQ(address_record.get_size(), gibibyte);
+}

projects/rocshmem/tests/unit_tests/binner_gtest.hpp

@@ -0,0 +1,78 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_BINNER_GTEST_HPP
+#define ROCSHMEM_BINNER_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include <map>
+
+#include "../src/../src/memory/address_record.hpp"
+#include "../src/memory/bin.hpp"
+#include "../src/memory/binner.hpp"
+#include "../src/memory/heap_memory.hpp"
+#include "../src/memory/hip_allocator.hpp"
+
+namespace rocshmem {
+
+class BinnerTestFixture : public ::testing::Test
+{
+    /**
+     * @brief Helper type for address records
+     */
+    using AR_T = AddressRecord;
+
+    /**
+     * @brief Helper type for size to bin maps
+     */
+    using BINS_T = std::map<size_t, Bin<AR_T>>;
+
+    /**
+     * @brief Helper type for binner
+     */
+    using BINNER_T = Binner<AR_T, BINS_T>;
+
+  protected:
+    /**
+     * @brief a heap memory object
+     */
+    HeapMemory<HIPAllocator> hm_ {};
+
+    /**
+     * @brief a bins object
+     */
+    BINS_T bins_ {};
+
+    /**
+     * @brief a binner object
+     */
+    BINNER_T binner_ {&bins_,
+                      hm_.get_ptr(),
+                      hm_.get_size()};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_BINNER_GTEST_HPP

projects/rocshmem/tests/unit_tests/bitwise_gtest.hpp

@@ -0,0 +1,351 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_BITWISE_GTEST_HPP
+#define ROCSHMEM_BITWISE_GTEST_HPP
+
+#define HIP_ENABLE_PRINTF
+
+#include "gtest/gtest.h"
+#include "wf_size.hpp"
+
+#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 {
+
+/*****************************************************************************
+ ************************ WarpMatrix Type Helpers ****************************
+ *****************************************************************************/
+typedef Matrix<uint64_t> WarpMatrix;
+
+/*****************************************************************************
+ ***************************** Device Methods ********************************
+ *****************************************************************************/
+class BitwiseDeviceMethods
+{
+  public:
+    /*************************************************************************
+     ************************* Block Strategy Methods ************************
+     *************************************************************************/
+    __device__
+    void
+    lowest_active_lane(WarpMatrix *warp_matrix,
+                       size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            auto low_lane = block.lowest_active_lane();
+            size_t warp_index = hipThreadIdx_x / WF_SIZE;
+            size_t block_index = hipBlockIdx_x;
+            auto *elem = warp_matrix->access(warp_index, block_index);
+            *elem = low_lane;
+        }
+    }
+
+    __device__
+    void
+    is_lowest_active_lane(WarpMatrix *warp_matrix,
+                          size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            if (block.is_lowest_active_lane()) {
+                size_t warp_index = hipThreadIdx_x / WF_SIZE;
+                size_t block_index = hipBlockIdx_x;
+                auto *elem = warp_matrix->access(warp_index, block_index);
+                *elem = block.lane_id();
+            }
+        }
+    }
+
+    __device__
+    void
+    active_logical_lane_id_2(WarpMatrix *warp_matrix,
+                             size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            if (block.active_logical_lane_id() == 2) {
+                size_t warp_index = hipThreadIdx_x / WF_SIZE;
+                size_t block_index = hipBlockIdx_x;
+                auto *elem = warp_matrix->access(warp_index, block_index);
+                *elem = block.lane_id();
+            }
+        }
+    }
+
+    __device__
+    void
+    lane_id(WarpMatrix *warp_matrix,
+            size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            auto lane_id = block.lane_id();
+            size_t warp_index = hipThreadIdx_x / WF_SIZE;
+            size_t block_index = hipBlockIdx_x;
+            auto *elem = warp_matrix->access(warp_index, block_index);
+            *elem = lane_id;
+        }
+    }
+
+    __device__
+    void
+    number_active_lanes(WarpMatrix *warp_matrix,
+                        size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            auto number_active_lanes = block.number_active_lanes();
+            size_t warp_index = hipThreadIdx_x / WF_SIZE;
+            size_t block_index = hipBlockIdx_x;
+            auto *elem = warp_matrix->access(warp_index, block_index);
+            *elem = number_active_lanes;
+        }
+    }
+
+    __device__
+    void
+    broadcast_up_value_42(WarpMatrix *warp_matrix,
+                          size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            uint64_t value = 1;
+            if (block.is_lowest_active_lane()) {
+                value = 42;
+            }
+            value = block.broadcast_up(value);
+            size_t warp_index = hipThreadIdx_x / WF_SIZE;
+            size_t block_index = hipBlockIdx_x;
+            auto *elem = warp_matrix->access(warp_index, block_index);
+            *elem = value;
+        }
+    }
+
+    __device__
+    void
+    fetch_incr_lowest_active_lane(WarpMatrix *warp_matrix,
+                                  size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            auto orig = block.fetch_incr(_fetch_value);
+            if (block.is_lowest_active_lane()) {
+                size_t warp_index = hipThreadIdx_x / WF_SIZE;
+                size_t block_index = hipBlockIdx_x;
+                auto *elem = warp_matrix->access(warp_index, block_index);
+                *elem = orig;
+            }
+        }
+    }
+
+    __device__
+    void
+    fetch_incr_active_logical_lane_1(WarpMatrix *warp_matrix,
+                                     size_t lanes_bitfield)
+    {
+        Block block {};
+        if (activate_lane_helper(lanes_bitfield)) {
+            auto orig = block.fetch_incr(_fetch_value);
+            if (block.active_logical_lane_id() == 1) {
+                size_t warp_index = hipThreadIdx_x / WF_SIZE;
+                size_t block_index = hipBlockIdx_x;
+                auto *elem = warp_matrix->access(warp_index, block_index);
+                *elem = orig;
+            }
+        }
+    }
+
+    /*************************************************************************
+     ************************* Helper Methods ********************************
+     *************************************************************************/
+    __device__
+    bool
+    activate_lane_helper(uint64_t lanes_bitfield)
+    {
+        /*
+         * In the following example, assume the following values:
+         *     hipThreadIdx_x := 66
+         *     _warp_size := 64.
+         *
+         *     index (tens):    0 0 0 0 0 0 0 ... 6 6 . .
+         *           (ones):    0 1 2 3 4 5 6 ... 2 3 . .
+         *     lanes_bitfield: [1 0 1 0 1 0 1 ... 1 0 . .]
+         *
+         * Example:
+         * warp_bit_id := hipThreadIdx_x % _warp_size;
+         * warp_bit_id := 66 % 64
+         * warp_bit_id := 2
+         */
+        uint64_t warp_bit_id = hipThreadIdx_x % WF_SIZE;
+
+        /*
+         * Example (continued):
+         * warp_bitmask := 1 << 2
+         * index (tens):  0 0 0 0 0 0 0 ... 6 6 . .
+         *       (ones):  0 1 2 3 4 5 6 ... 2 3 . .
+         * warp_bitmask: [0 0 1 0 0 0 0 ... 0 0 . .]
+         */
+        uint64_t my_warp_bitmask_id = 1UL << warp_bit_id;
+
+        /*
+         * Example (continued):
+         * index (tens):    0 0 0 0 0 0 0 ... 6 6 . .
+         *       (ones):    0 1 2 3 4 5 6 ... 2 3 . .
+         * lanes_bitfield: [1 0 1 0 1 0 1 ... 1 0 . .]
+         * warp_bitmask:   [0 0 1 0 0 0 0 ... 0 0 . .]
+         */
+        bool is_an_active_lane = lanes_bitfield & my_warp_bitmask_id;
+
+        return is_an_active_lane;
+    }
+
+    long long unsigned *_fetch_value = nullptr;
+};
+
+/*****************************************************************************
+ ***************************** Test Fixture **********************************
+ *****************************************************************************/
+class BitwiseTestFixture : public ::testing::Test
+{
+  public:
+    BitwiseTestFixture() = default;
+
+    ~BitwiseTestFixture()
+    {
+        if (_device_methods) {
+            if (_device_methods->_fetch_value) {
+                _hip_allocator.deallocate(_device_methods->_fetch_value);
+            }
+            _hip_allocator.deallocate(_device_methods);
+        }
+        if (_warp_matrix) {
+            _hip_allocator.deallocate(_warp_matrix);
+        }
+    }
+
+    /*************************************************************************
+     **************************** Setup Methods ******************************
+     *************************************************************************/
+    void
+    setup_fixture(dim3 block_dim, dim3 grid_dim)
+    {
+        _hip_block_dim = block_dim;
+        _hip_grid_dim = grid_dim;
+        _wf_size = get_wf_size();
+
+        assert(_device_methods == nullptr);
+        _hip_allocator.allocate(reinterpret_cast<void**>(&_device_methods),
+                                sizeof(BitwiseDeviceMethods));
+
+        assert(_device_methods);
+
+        _hip_allocator.allocate(reinterpret_cast<void**>(&_device_methods->_fetch_value),
+                                sizeof(long long unsigned));
+
+        assert(_device_methods->_fetch_value);
+
+        *_device_methods->_fetch_value = 0;
+
+        assert(_warp_matrix == nullptr);
+        _hip_allocator.allocate(reinterpret_cast<void**>(&_warp_matrix),
+                                sizeof(WarpMatrix));
+
+        size_t warps_per_block = ceil(float(_hip_block_dim.x) / WF_SIZE);
+
+        const ObjectStrategy *default_object_strategy =
+                DefaultObjectStrategy::instance()->get();
+
+        assert(_warp_matrix);
+        new (_warp_matrix) WarpMatrix(warps_per_block,
+                                      _hip_grid_dim.x,
+                                      _hip_allocator,
+                                      *default_object_strategy);
+    }
+
+    void
+    zero_warp_matrix()
+    {
+        for (size_t row = 0; row < _warp_matrix->rows(); row++) {
+            for (size_t col = 0; col < _warp_matrix->columns(); col++) {
+                auto *entry = _warp_matrix->access(row, col);
+                *entry = 0;
+            }
+        }
+    }
+
+    void
+    verify_zeroed_warp_matrix()
+    {
+        for (size_t row = 0; row < _warp_matrix->rows(); row++) {
+            for (size_t col = 0; col < _warp_matrix->columns(); col++) {
+                auto *entry = _warp_matrix->access(row, col);
+                ASSERT_EQ(*entry, 0);
+            }
+        }
+    }
+
+    /*************************************************************************
+     *********************** Kernel Launch Methods ***************************
+     *************************************************************************/
+    void
+    host_run_device_kernel(void(*fn)(BitwiseDeviceMethods*,
+                                     WarpMatrix*,
+                                     size_t),
+                           size_t activate_lanes_bitfield)
+    {
+        hipLaunchKernelGGL(fn,
+                           _hip_grid_dim,
+                           _hip_block_dim,
+                           0,
+                           nullptr,
+                           _device_methods,
+                           _warp_matrix,
+                           activate_lanes_bitfield);
+
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+  protected:
+    /*************************************************************************
+     ********************** Implementation Variables *************************
+     *************************************************************************/
+    dim3 _hip_block_dim {};
+    dim3 _hip_grid_dim {};
+    HIPAllocator _hip_allocator {};
+    WarpMatrix *_warp_matrix = nullptr;
+    BitwiseDeviceMethods *_device_methods = nullptr;
+    int _wf_size;
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_BITWISE_GTEST_HPP

projects/rocshmem/tests/unit_tests/dlmalloc_gtest.cpp

@@ -0,0 +1,130 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "dlmalloc_gtest.hpp"
+#include <cstdint>
+
+using namespace rocshmem;
+
+// a small portion of the heap is not available due to cost of dlmalloc bookkeeping
+#define DLMALLOC_BOOKKEEPING static_cast<size_t>(128 * ALIGNMENT)
+
+TEST_F(DLMallocTestFixture, used_0_bytes) {
+  size_t heap_size{1 << 30};
+  ASSERT_LE(strat_.get_used(), DLMALLOC_BOOKKEEPING);
+  ASSERT_EQ(strat_.get_used() + strat_.get_avail(), heap_size);
+}
+
+TEST_F(DLMallocTestFixture, alloc_0_bytes) {
+  size_t initial_used{strat_.get_used()};
+  char* c_ptr{nullptr};
+  size_t size{0};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_EQ(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used(), initial_used);
+}
+
+TEST_F(DLMallocTestFixture, alloc_1_byte) {
+  char* c_ptr{nullptr};
+  size_t size{1};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_128_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{128};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_256_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{256};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_512_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{512};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_513_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{513};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_4KB) {
+  char* c_ptr{nullptr};
+  size_t size{4096};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_128KB) {
+  char* c_ptr{nullptr};
+  size_t size{131072};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_1GB) {
+  char* c_ptr{nullptr};
+  size_t heap_size{1 << 30};
+  size_t size{heap_size - DLMALLOC_BOOKKEEPING};
+
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used() + strat_.get_avail(), heap_size);
+}
+
+TEST_F(DLMallocTestFixture, alloc_256_bytes_X2_free_256_bytes_X2) {
+  char* c_ptr_1{nullptr};
+  char* c_ptr_2{nullptr};
+  size_t size{256};
+  strat_.alloc(&c_ptr_1, size);
+  ASSERT_NE(c_ptr_1, nullptr);
+  strat_.alloc(&c_ptr_2, size);
+  ASSERT_NE(c_ptr_2, nullptr);
+
+  strat_.free(c_ptr_1);
+  strat_.free(c_ptr_2);
+}
+
+TEST_F(DLMallocTestFixture, alloc_1GB_free_1GB) {
+  char* c_ptr{nullptr};
+  size_t heap_size{1 << 30};
+  size_t size{heap_size - DLMALLOC_BOOKKEEPING};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used() + strat_.get_avail(), heap_size);
+
+  strat_.free(c_ptr);
+}
+

projects/rocshmem/tests/unit_tests/dlmalloc_gtest.hpp

@@ -0,0 +1,62 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_DLMALLOC_GTEST_HPP
+#define ROCSHMEM_DLMALLOC_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/heap_memory.hpp"
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/memory/dlmalloc.hpp"
+
+namespace rocshmem {
+
+class DLMallocTestFixture : public ::testing::Test
+{
+    /**
+     * @brief Helper type for heap memory
+     */
+    using HEAP_T = HeapMemory<HIPAllocator>;
+
+    /**
+     * @brief Helper type for allocation strategy
+     */
+    using STRAT_T = DLAllocatorStrategy<HEAP_T>;
+
+  protected:
+    /**
+     * @brief Heap memory object
+     */
+    HEAP_T heap_mem_ {};
+
+    /**
+     * @brief Allocation strategy object
+     */
+    STRAT_T strat_ {&heap_mem_};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_DLMALLOC_GTEST_HPP

projects/rocshmem/tests/unit_tests/envvar_gtest.cpp

@@ -0,0 +1,283 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 <limits>
+
+#include "envvar_gtest.hpp"
+
+using namespace rocshmem;
+
+using VarTypes = ::testing::Types<bool,
+                                  uint32_t,
+                                  size_t,
+                                  int64_t,
+                                  std::string,
+                                  envvar::types::socket_family,
+                                  envvar::types::debug_level>;
+
+TYPED_TEST_SUITE(EnvVarUnsetTestFixture, VarTypes);
+TYPED_TEST_SUITE(EnvVarSetTestFixture, VarTypes);
+
+TYPED_TEST(EnvVarUnsetTestFixture, name) {
+  EXPECT_EQ(this->var_.get_name(), this->var_full_name_);
+}
+
+TYPED_TEST(EnvVarUnsetTestFixture, doc) {
+  EXPECT_EQ(this->var_.get_doc(), this->var_doc_);
+}
+
+TYPED_TEST(EnvVarUnsetTestFixture, is_default) {
+  EXPECT_TRUE(this->var_.is_default());
+  EXPECT_EQ(this->var_.get_value(), this->var_.get_default());
+}
+
+TYPED_TEST(EnvVarSetTestFixture, is_default) {
+  EXPECT_FALSE(this->var_.is_default());
+  EXPECT_EQ(this->var_.get_value(), this->var_.get_default());
+}
+
+TEST_F(EnvVarTestFixture, string_custom_default) {
+  const std::string default_value_{"This is the default value."};
+  const envvar::var<std::string> var_{this->var_name_, this->var_doc_, default_value_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), default_value_);
+}
+
+TEST_F(EnvVarTestFixture, parse_integer) {
+  this->setenv("1073741824");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), 1L << 30);
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_notaninteger) {
+  this->setenv("Ceci n'est pas un entier.");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_, 1L << 30};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_value(), 1L << 30);
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_large) {
+  this->setenv("9223372036854775807");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), std::numeric_limits<int64_t>::max());
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_too_large) {
+  this->setenv("9223372036854775808");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), var_.get_default());
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_negative) {
+  this->setenv("-1073741824");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), -1L << 30);
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_negative_large) {
+  this->setenv("-9223372036854775808");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), std::numeric_limits<int64_t>::min());
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_negative_too_large) {
+  this->setenv("-9223372036854775809");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), var_.get_default());
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_hex) {
+  this->setenv("0x40000000");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), 1L << 30);
+}
+
+TEST_F(EnvVarTestFixture, parse_integer_hex_only) {
+  this->setenv("0x40000000");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_,
+                                  envvar::parser::parse_hex<int64_t>{}};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), 1L << 30);
+}
+
+// parse_hex<> interprets input digits as hexadecimal, even without 0x prefix
+TEST_F(EnvVarTestFixture, parse_integer_hex_only_noprefix) {
+  this->setenv("40000000");
+  const envvar::var<int64_t> var_{this->var_name_, this->var_doc_,
+                                  envvar::parser::parse_hex<int64_t>{}};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), 1L << 30);
+}
+
+TEST_F(EnvVarTestFixture, parse_unsigned_integer) {
+  this->setenv("1073741824");
+  const envvar::var<uint32_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), 1L << 30);
+}
+
+TEST_F(EnvVarTestFixture, parse_unsigned_integer_negative) {
+  this->setenv("-1");
+  const envvar::var<uint32_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), var_.get_default());
+}
+
+TEST_F(EnvVarTestFixture, parse_unsigned_large) {
+  this->setenv("4294967296");
+  const envvar::var<uint32_t> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), 0);
+  EXPECT_EQ(var_.get_value(), var_.get_default());
+}
+
+TEST_F(EnvVarTestFixture, parse_bool_zero) {
+  this->setenv("0");
+  const envvar::var<bool> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), false);
+  EXPECT_EQ(var_.get_value(), false);
+}
+
+TEST_F(EnvVarTestFixture, parse_bool_one) {
+  this->setenv("1");
+  const envvar::var<bool> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), false);
+  EXPECT_EQ(var_.get_value(), true);
+}
+
+TEST_F(EnvVarTestFixture, parse_bool_negative) {
+  this->setenv("-1");
+  const envvar::var<bool> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), false);
+  EXPECT_EQ(var_.get_value(), false);
+}
+
+TEST_F(EnvVarTestFixture, parse_bool_two) {
+  this->setenv("2");
+  const envvar::var<bool> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), false);
+  EXPECT_EQ(var_.get_value(), false);
+}
+
+TEST_F(EnvVarTestFixture, parse_bool_false) {
+  this->setenv("false");
+  const envvar::var<bool> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), false);
+  EXPECT_EQ(var_.get_value(), false);
+}
+
+TEST_F(EnvVarTestFixture, parse_bool_true) {
+  this->setenv("true");
+  const envvar::var<bool> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), false);
+  EXPECT_EQ(var_.get_value(), true);
+}
+
+TEST_F(EnvVarTestFixture, parse_bool_other) {
+  this->setenv("other");
+  const envvar::var<bool> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), false);
+  EXPECT_EQ(var_.get_value(), false);
+}
+
+TEST_F(EnvVarTestFixture, parse_socket_family_unspec) {
+  this->setenv("UNSPEC");
+  const envvar::var<envvar::types::socket_family> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), envvar::types::socket_family::UNSPEC);
+  EXPECT_EQ(var_.get_value(), envvar::types::socket_family::UNSPEC);
+}
+
+TEST_F(EnvVarTestFixture, parse_socket_family_af_unspec) {
+  this->setenv("AF_UNSPEC");
+  const envvar::var<envvar::types::socket_family> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), envvar::types::socket_family::UNSPEC);
+  EXPECT_EQ(var_.get_value(), envvar::types::socket_family::UNSPEC);
+}
+
+TEST_F(EnvVarTestFixture, parse_socket_family_inet) {
+  this->setenv("INET");
+  const envvar::var<envvar::types::socket_family> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), envvar::types::socket_family::UNSPEC);
+  EXPECT_EQ(var_.get_value(), envvar::types::socket_family::INET);
+}
+
+TEST_F(EnvVarTestFixture, parse_socket_family_af_inet) {
+  this->setenv("AF_INET");
+  const envvar::var<envvar::types::socket_family> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), envvar::types::socket_family::UNSPEC);
+  EXPECT_EQ(var_.get_value(), envvar::types::socket_family::INET);
+}
+
+TEST_F(EnvVarTestFixture, parse_socket_family_inet6) {
+  this->setenv("INET6");
+  const envvar::var<envvar::types::socket_family> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), envvar::types::socket_family::UNSPEC);
+  EXPECT_EQ(var_.get_value(), envvar::types::socket_family::INET6);
+}
+
+TEST_F(EnvVarTestFixture, parse_socket_family_af_inet6) {
+  this->setenv("AF_INET6");
+  const envvar::var<envvar::types::socket_family> var_{this->var_name_, this->var_doc_};
+  EXPECT_FALSE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), envvar::types::socket_family::UNSPEC);
+  EXPECT_EQ(var_.get_value(), envvar::types::socket_family::INET6);
+}
+
+TEST_F(EnvVarTestFixture, parse_socket_family_nonesense) {
+  this->setenv("'Twas brillig, and the slithy toves, did gyre and gimble in the wabe.");
+  const envvar::var<envvar::types::socket_family> var_{this->var_name_, this->var_doc_};
+  EXPECT_TRUE(var_.is_default());
+  EXPECT_EQ(var_.get_default(), envvar::types::socket_family::UNSPEC);
+  EXPECT_EQ(var_.get_value(), var_.get_default());
+}

projects/rocshmem/tests/unit_tests/envvar_gtest.hpp

@@ -0,0 +1,97 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_ENVVAR_GTEST_HPP
+#define ROCSHMEM_ENVVAR_GTEST_HPP
+
+#include <cstdlib>
+#include <sstream>
+#include <string>
+
+#include "gtest/gtest.h"
+#include "../src/envvar.hpp"
+
+namespace rocshmem {
+
+class EnvVarTestFixture : public ::testing::Test
+{
+  public:
+    static constexpr envvar::category::tag category_ = envvar::category::tag::ROCSHMEM;
+    static constexpr const char* var_name_ = "GTEST";
+    static constexpr const char* var_prefix_ = envvar::category::prefix<category_>;
+    static inline const std::string var_full_name_{std::string(var_prefix_).append("_").append(var_name_)};
+    static constexpr const char* var_doc_ = "Test envvar documentation: documents the test envvar.";
+
+    static int setenv(const char* value) {
+      return ::setenv(var_full_name_.c_str(), value, true);
+    }
+
+    static int setenv(const std::string& value) {
+      return setenv(value.c_str());
+    }
+
+    static int unsetenv() {
+      return ::unsetenv(var_full_name_.c_str());
+    }
+
+  protected:
+    static void SetUpTestSuite() {
+      unsetenv();
+    }
+
+    static void TearDownTestSuite() {
+      unsetenv();
+    }
+
+    void SetUp() override {
+      unsetenv();
+    }
+
+    void TearDown() override {
+      unsetenv();
+    }
+};
+
+template <typename T>
+class EnvVarUnsetTestFixture : public EnvVarTestFixture
+{
+  protected:
+    const envvar::var<T> var_{var_name_, var_doc_};
+};
+
+template <typename T>
+class EnvVarSetTestFixture : public EnvVarTestFixture
+{
+  protected:
+    const envvar::var<T> var_{var_name_, var_doc_};
+
+    static void SetUpTestSuite() {
+      std::ostringstream oss{};
+      oss << T{};
+      setenv(oss.str());
+    }
+};
+} // namespace rocshmem
+
+#endif // ROCSHMEM_ENVVAR_GTEST_HPP

projects/rocshmem/tests/unit_tests/free_list_gtest.cpp

@@ -0,0 +1,309 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "free_list_gtest.hpp"
+
+#include "../src/util.hpp"
+
+using namespace rocshmem;
+
+/*****************************************************************************
+ ******************************* Fixture Tests *******************************
+ *****************************************************************************/
+
+namespace rocshmem {
+
+template <typename List, typename Value>
+__global__ void pop_all(List* list, Value* values, const std::size_t count) {
+  const auto stride = blockDim.x * gridDim.x;
+  const auto thread_index = blockIdx.x * blockDim.x + threadIdx.x;
+  // One push per block. block size is always WF_SIZE
+  for (std::size_t i = thread_index; i < count * WF_SIZE; i += stride) {
+    if (is_thread_zero_in_wave()) {
+      auto last = list->pop_front();
+      if (values != nullptr) {
+        values[i / WF_SIZE] = last.value;
+      }
+    }
+  }
+}
+
+template <typename List, typename Value>
+__global__ void push_all(List* list, const Value* values,
+                         const std::size_t count) {
+  const auto stride = blockDim.x * gridDim.x;
+  const auto thread_index = blockIdx.x * blockDim.x + threadIdx.x;
+  // One push per block. block size is always WF_SIZE
+  for (std::size_t i = thread_index; i < count * WF_SIZE; i += stride) {
+    if (is_thread_zero_in_wave()) {
+      list->push_back(values[i / WF_SIZE]);
+    }
+  }
+}
+
+template <typename List>
+__global__ void pop_empty(List* list, bool* empty) {
+  auto pop_result = list->pop_front();
+  *empty = !pop_result.success;
+}
+}  // namespace rocshmem
+
+TYPED_TEST(FreeListTestFixture, pop_empty_device) {
+  using Allocator = typename TestFixture::Allocator;
+  using T = typename TestFixture::T;
+
+  auto& h_input = this->h_input;
+  auto& free_list = this->free_list;
+  auto& hip_allocator_ = this->hip_allocator_;
+
+  bool *is_empty {nullptr};
+  hip_allocator_.allocate(reinterpret_cast<void**>(&is_empty),
+                          sizeof(bool));
+
+  CHECK_HIP(hipMemset(is_empty, 0, sizeof(bool)));
+  FreeListProxy<Allocator, T> empty_list_proxy{};
+  FreeList<T, Allocator>* empty_free_list{empty_list_proxy.get()};
+
+  rocshmem::pop_empty<<<1, 1>>>(empty_free_list, is_empty);
+  CHECK_HIP(hipDeviceSynchronize());
+  EXPECT_TRUE(is_empty[0]);
+
+  hip_allocator_.deallocate(is_empty);
+}
+
+TYPED_TEST(FreeListTestFixture, push_host_pop_device) {
+  using Allocator = typename TestFixture::Allocator;
+  using T = typename TestFixture::T;
+
+  auto& h_input = this->h_input;
+  auto& free_list = this->free_list;
+  auto& hip_allocator_ = this->hip_allocator_;
+
+  T *results {nullptr};
+  bool *is_empty {nullptr};
+  size_t size_bytes = sizeof(T) * h_input.size() + sizeof(bool);
+  hip_allocator_.allocate(reinterpret_cast<void**>(&results),
+                          size_bytes);
+
+  CHECK_HIP(hipMemset(results, 0, size_bytes));
+  is_empty = reinterpret_cast<bool*>(results + h_input.size());
+  const auto block_size = this->wf_size;
+  rocshmem::pop_all<<<1, block_size>>>(free_list, results, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  for (std::size_t i = 0; i < h_input.size(); i++) {
+    EXPECT_EQ(results[i], h_input[i]);
+  }
+
+  rocshmem::pop_empty<<<1, 1>>>(free_list, is_empty);
+  CHECK_HIP(hipDeviceSynchronize());
+
+  EXPECT_TRUE(is_empty[0]);
+  hip_allocator_.deallocate(results);
+}
+
+TYPED_TEST(FreeListTestFixture, push_host_concurrent_pop_device) {
+  using Allocator = typename TestFixture::Allocator;
+  using T = typename TestFixture::T;
+
+  auto& h_input = this->h_input;
+  auto& free_list = this->free_list;
+  auto& hip_allocator_ = this->hip_allocator_;
+
+  T *results {nullptr};
+  bool *is_empty {nullptr};
+  size_t size_bytes = sizeof(T) * h_input.size() + sizeof(bool);
+  hip_allocator_.allocate(reinterpret_cast<void**>(&results),
+                          size_bytes);
+
+  CHECK_HIP(hipMemset(results, 0, size_bytes));
+  is_empty = reinterpret_cast<bool*>(results + h_input.size());
+  const auto num_blocks = h_input.size();
+  const auto block_size = this->wf_size;
+  rocshmem::pop_all<<<num_blocks, block_size>>>(
+            free_list, results, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  std::vector<T> h_results(h_input.size());
+  CHECK_HIP(hipMemcpy(h_results.data(), results, sizeof(T) * h_input.size(),
+                      hipMemcpyDeviceToHost));
+
+  // sort to guarantee that the ordering is correct
+  std::sort(h_input.begin(), h_input.end());
+  std::sort(h_results.begin(), h_results.end());
+
+
+  for (std::size_t i = 0; i < h_results.size(); i++) {
+    EXPECT_EQ(h_results[i], h_input[i]);
+  }
+
+  rocshmem::pop_empty<<<1, 1>>>(free_list, is_empty);
+  CHECK_HIP(hipDeviceSynchronize());
+
+  EXPECT_TRUE(is_empty[0]);
+  hip_allocator_.deallocate(results);
+}
+
+TYPED_TEST(FreeListTestFixture, push_host_pop_push_device) {
+  using Allocator = typename TestFixture::Allocator;
+  using T = typename TestFixture::T;
+  using FreeListType = FreeList<T, Allocator>;
+
+  auto& h_input = this->h_input;
+  auto& free_list = this->free_list;
+  auto& hip_allocator_ = this->hip_allocator_;
+
+  T *results {nullptr};
+  T *d_input {nullptr};
+  bool *is_empty {nullptr};
+  size_t size_bytes = 2 * sizeof(T) * h_input.size() + sizeof(bool);
+  hip_allocator_.allocate(reinterpret_cast<void**>(&results),
+                          size_bytes);
+
+  CHECK_HIP(hipMemset(results, 0, size_bytes));
+  d_input = reinterpret_cast<T*>(results + h_input.size());
+  is_empty = reinterpret_cast<bool*>(d_input + h_input.size());
+  const auto block_size = this->wf_size;
+
+  CHECK_HIP(hipMemcpy(d_input, h_input.data(), sizeof(T) * h_input.size(),
+                      hipMemcpyHostToDevice));
+
+  rocshmem::pop_all<FreeListType, T><<<1, block_size>>>(
+            free_list, nullptr, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  rocshmem::push_all<<<1, block_size>>>(free_list, d_input, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  rocshmem::pop_all<<<1, block_size>>>(free_list, results, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  for (std::size_t i = 0; i < h_input.size(); i++) {
+    EXPECT_EQ(results[i], h_input[i]);
+  }
+
+  hip_allocator_.deallocate(results);
+}
+
+TYPED_TEST(FreeListTestFixture, push_host_pop_concurrent_push_device) {
+  using Allocator = typename TestFixture::Allocator;
+  using T = typename TestFixture::T;
+  using FreeListType = FreeList<T, Allocator>;
+
+  auto& h_input = this->h_input;
+  auto& free_list = this->free_list;
+  auto& hip_allocator_ = this->hip_allocator_;
+
+  T *results {nullptr};
+  T *d_input {nullptr};
+  size_t size_bytes = 2 * sizeof(T) * h_input.size();
+  hip_allocator_.allocate(reinterpret_cast<void**>(&results),
+                          size_bytes);
+
+  CHECK_HIP(hipMemset(results, 0, size_bytes));
+  d_input = reinterpret_cast<T*>(results + h_input.size());
+  const auto block_size = this->wf_size;
+
+  CHECK_HIP(hipMemcpy(d_input, h_input.data(), sizeof(T) * h_input.size(),
+                      hipMemcpyHostToDevice));
+
+  rocshmem::pop_all<FreeListType, T><<<1, block_size>>>(
+            free_list, nullptr,h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  // Concurrently push all values
+  const auto num_blocks = h_input.size();
+  rocshmem::push_all<<<num_blocks, block_size>>>(
+      free_list, d_input, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  rocshmem::pop_all<<<1, block_size>>>(free_list, results, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  std::vector<T> h_results(h_input.size());
+  CHECK_HIP(hipMemcpy(h_results.data(), results, sizeof(T) * h_input.size(),
+                      hipMemcpyDeviceToHost));
+
+  // sort to guarantee that the ordering is correct
+  std::sort(h_input.begin(), h_input.end());
+  std::sort(h_results.begin(), h_results.end());
+
+  for (std::size_t i = 0; i < h_results.size(); i++) {
+    EXPECT_EQ(h_results[i], h_input[i]);
+  }
+
+  hip_allocator_.deallocate(results);
+}
+
+TYPED_TEST(FreeListTestFixture, push_host_concurrent_pop_push_device) {
+  using Allocator = typename TestFixture::Allocator;
+  using T = typename TestFixture::T;
+  using FreeListType = FreeList<T, Allocator>;
+
+  auto& h_input = this->h_input;
+  auto& free_list = this->free_list;
+  auto& hip_allocator_ = this->hip_allocator_;
+
+  T *results {nullptr};
+  T *d_input {nullptr};
+  size_t size_bytes = 2 * sizeof(T) * h_input.size();
+  hip_allocator_.allocate(reinterpret_cast<void**>(&results),
+                          size_bytes);
+
+  CHECK_HIP(hipMemset(results, 0, size_bytes));
+  d_input = reinterpret_cast<T*>(results + h_input.size());
+
+  CHECK_HIP(hipMemcpy(d_input, h_input.data(), sizeof(T) * h_input.size(),
+                      hipMemcpyHostToDevice));
+
+  const auto block_size = this->wf_size;
+  rocshmem::pop_all<FreeListType, T><<<1, block_size>>>(
+            free_list, nullptr, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  // Concurrently push all values
+  const auto num_blocks = h_input.size();
+  rocshmem::push_all<<<num_blocks, block_size>>>(
+      free_list, d_input, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  // Concurrently pop all values
+  rocshmem::pop_all<<<num_blocks, block_size>>>(
+            free_list, results, h_input.size());
+  CHECK_HIP(hipDeviceSynchronize());
+
+  std::vector<T> h_results(h_input.size());
+  CHECK_HIP(hipMemcpy(h_results.data(), results, sizeof(T) * h_input.size(),
+                      hipMemcpyDeviceToHost));
+
+  // sort to guarantee that the ordering is correct
+  std::sort(h_input.begin(), h_input.end());
+  std::sort(h_results.begin(), h_results.end());
+
+  for (std::size_t i = 0; i < h_results.size(); i++) {
+    EXPECT_EQ(h_results[i], h_input[i]);
+  }
+
+  hip_allocator_.deallocate(results);
+}

projects/rocshmem/tests/unit_tests/free_list_gtest.hpp

@@ -0,0 +1,68 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_FREE_LIST_GTEST_HPP
+#define ROCSHMEM_FREE_LIST_GTEST_HPP
+
+#include <numeric>
+#include <vector>
+
+#include "../src/containers/free_list_impl.hpp"
+#include "gtest/gtest.h"
+#include "../src/memory/hip_allocator.hpp"
+#include "wf_size.hpp"
+
+namespace rocshmem {
+
+template <typename ValueType>
+class FreeListTestFixture : public ::testing::Test {
+ public:
+  FreeListTestFixture() : h_input(num_elements) {
+    std::iota(h_input.begin(), h_input.end(), T{1});
+    free_list = list_proxy.get();
+  }
+
+ protected:
+  void SetUp() override {
+    free_list->push_back_range(h_input.begin(), h_input.end());
+    wf_size = get_wf_size();
+  }
+
+  using T = ValueType;
+  using Allocator = HIPAllocator;
+  Allocator hip_allocator_ {};
+  const std::size_t num_elements{32};
+  std::vector<T> h_input{};
+  int wf_size;
+
+  FreeListProxy<Allocator, T> list_proxy{};
+  FreeList<T, Allocator>* free_list{};
+};
+
+using TestTypes = ::testing::Types<std::uint32_t, std::uint64_t>;
+TYPED_TEST_SUITE(FreeListTestFixture, TestTypes);
+
+}  // namespace rocshmem
+
+#endif  // ROCSHMEM_FREE_LIST_GTEST_HPP

projects/rocshmem/tests/unit_tests/heap_memory_gtest.cpp

@@ -0,0 +1,41 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "heap_memory_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST(HeapMemoryTest, size_constructor) {
+  HeapMemory<HIPAllocator> heap_mem{2048};
+  ASSERT_EQ(heap_mem.get_size(), 2048);
+  ASSERT_NE(heap_mem.get_ptr(), nullptr);
+}
+
+TEST_F(HeapMemoryTestFixture, size_check) {
+  ASSERT_EQ(heap_mem_.get_size(), 1 << 30);
+}
+
+TEST_F(HeapMemoryTestFixture, ptr_check) {
+  ASSERT_NE(heap_mem_.get_ptr(), nullptr);
+}

projects/rocshmem/tests/unit_tests/heap_memory_gtest.hpp

@@ -0,0 +1,48 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_HEAP_MEMORY_GTEST_HPP
+#define ROCSHMEM_HEAP_MEMORY_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include <hip/hip_runtime_api.h>
+
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/memory/heap_memory.hpp"
+
+namespace rocshmem {
+
+class HeapMemoryTestFixture : public ::testing::Test
+{
+  protected:
+    /**
+     * @brief a heap memory object
+     */
+    HeapMemory<HIPAllocator> heap_mem_ {};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_HEAP_MEMORY_GTEST_HPP

projects/rocshmem/tests/unit_tests/hipmalloc_gtest.cpp

@@ -0,0 +1,45 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#include "hipmalloc_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(HipMallocTestFixture, normal_1GBx256) {
+    void* ptr{nullptr};
+    size_t gb {1073741824};
+    for (int i{0}; i < 256; i++) {
+        hip_allocator_.allocate(&ptr, gb);
+        hip_allocator_.deallocate(ptr);
+    }
+}
+
+TEST_F(HipMallocTestFixture, fine_1GBx256) {
+    void* ptr{nullptr};
+    size_t gb {1073741824};
+    for (int i{0}; i < 256; i++) {
+        hip_allocator_fg_.allocate(&ptr, gb);
+        hip_allocator_fg_.deallocate(ptr);
+    }
+}

projects/rocshmem/tests/unit_tests/hipmalloc_gtest.hpp

@@ -0,0 +1,43 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef ROCSHMEM_HIPMALLOC_GTEST_HPP
+#define ROCSHMEM_HIPMALLOC_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/symmetric_heap.hpp"
+#include "../src/util.hpp"
+
+namespace rocshmem {
+
+class HipMallocTestFixture : public ::testing::Test {
+  public:
+    HIPAllocator hip_allocator_ {};
+    HIPAllocatorFinegrained hip_allocator_fg_ {};
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_HIPMALLOC_GTEST_HPP

projects/rocshmem/tests/unit_tests/index_strategy_gtest.cpp

@@ -0,0 +1,157 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "index_strategy_gtest.hpp"
+
+#define HIP_ENABLE_PRINTF
+
+using namespace rocshmem;
+
+TEST_F(IndexStrategyTestFixture,
+       run_TCBA_memory_set_test_grid_1_1_1_block_1_1_1) {
+  using IndexStrategy = Thread_Contiguous_Block_Agnostic;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {1, 1, 1});
+}
+
+//=============================================================================
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1_1_1_block_64_1_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 1, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1_1_1_block_64_2_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 2, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1_1_1_block_64_4_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 4, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1_1_1_block_64_16_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 16, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1_1_1_block_16_16_4) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {16, 16, 4});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_4_1_1_block_16_16_4) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({4, 1, 1}, {16, 16, 4});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_4_4_4_block_16_16_4) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({4, 4, 4}, {16, 16, 4});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1024_8_8_block_1024_1_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1024, 8, 8}, {1024, 1, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1024_8_8_block_1_1024_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1024, 8, 8}, {1, 1024, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBD_memory_set_test_grid_1024_8_8_block_1_1_1024) {
+  using IndexStrategy = Thread_Discontiguous_Block_Discontiguous;
+  run_memory_set_test<IndexStrategy>({1024, 8, 8}, {1, 1, 1024});
+}
+
+//=============================================================================
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1_1_1_block_64_1_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 1, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1_1_1_block_64_2_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 2, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1_1_1_block_64_4_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 4, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1_1_1_block_64_16_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {64, 16, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1_1_1_block_16_16_4) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1, 1, 1}, {16, 16, 4});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_4_1_1_block_16_16_4) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({4, 1, 1}, {16, 16, 4});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_4_4_4_block_16_16_4) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({4, 4, 4}, {16, 16, 4});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1024_8_8_block_1024_1_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1024, 8, 8}, {1024, 1, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1024_8_8_block_1_1024_1) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1024, 8, 8}, {1, 1024, 1});
+}
+
+TEST_F(IndexStrategyTestFixture,
+       run_TDBC_memory_set_test_grid_1024_8_8_block_1_1_1024) {
+  using IndexStrategy = Thread_Discontiguous_Block_Contiguous;
+  run_memory_set_test<IndexStrategy>({1024, 8, 8}, {1, 1, 1024});
+}

projects/rocshmem/tests/unit_tests/index_strategy_gtest.hpp

@@ -0,0 +1,105 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_INDEX_STRATEGY_GTEST_HPP
+#define ROCSHMEM_INDEX_STRATEGY_GTEST_HPP
+
+#include "gtest/gtest.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 {
+
+template <typename INDEX_STRAT>
+__global__
+void
+memory_set(int *raw_mem, size_t num_elems)
+{
+    assert(raw_mem);
+
+    INDEX_STRAT idx_strat(num_elems);
+
+    for (size_t i = idx_strat.start();
+         i < idx_strat.end();
+         i = idx_strat.next(i)) {
+        raw_mem[i]++;
+    }
+}
+
+class IndexStrategyTestFixture : public ::testing::Test
+{
+  public:
+    IndexStrategyTestFixture()
+    {
+        assert(_raw_mem == nullptr);
+        size_t raw_mem_size_bytes = sizeof(int) * _mem_elements;
+        _hip_allocator.allocate(reinterpret_cast<void**>(&_raw_mem),
+                                raw_mem_size_bytes);
+
+        assert(_raw_mem);
+        for (size_t i = 0; i < _mem_elements; i++) {
+            _raw_mem[i] = 0;
+        }
+    }
+
+    ~IndexStrategyTestFixture()
+    {
+        if (_raw_mem) {
+            _hip_allocator.deallocate(_raw_mem);
+        }
+    }
+
+    template <typename INDEX_STRAT>
+    void
+    run_memory_set_test(const dim3 grid_dim, const dim3 block_dim)
+    {
+        hipLaunchKernelGGL(memory_set<INDEX_STRAT>,
+                           grid_dim,
+                           block_dim,
+                           0,
+                           nullptr,
+                           _raw_mem,
+                           _mem_elements);
+
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+
+        for(size_t i = 0; i < _mem_elements; i++) {
+            EXPECT_EQ(_raw_mem[i], 1);
+        }
+    }
+
+  protected:
+    int *_raw_mem = nullptr;
+    HIPAllocator _hip_allocator {};
+    static constexpr size_t _mem_elements = 262144;
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_INDEX_STRATEGY_GTEST_HPP

projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp

@@ -0,0 +1,203 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#include "ipc_impl_simple_coarse_gtest.hpp"
+
+using namespace rocshmem;
+
+
+TEST_P(DegenerateSimpleCoarse, ptr_check) {
+    ASSERT_NE(heap_mem_.get_ptr(), nullptr);
+}
+
+TEST_P(DegenerateSimpleCoarse, MPI_num_pes) {
+    ASSERT_EQ(mpi_->num_pes(), 2);
+}
+
+TEST_P(DegenerateSimpleCoarse, IPC_bases) {
+    ASSERT_EQ(mpi_->num_pes(), 2);
+    ASSERT_NE(ipc_impl_.ipc_bases, nullptr);
+    for(int i{0}; i < mpi_->num_pes(); i++) {
+        ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr);
+    }
+}
+
+TEST_P(DegenerateSimpleCoarse, golden_1048576_int) {
+    iota_golden(1048576);
+    validate_golden(1048576);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleCoarseTestFixture,
+    DegenerateSimpleCoarse,
+    ::testing::Values(
+        std::make_tuple(1,  1,   1))
+);
+
+//=============================================================================
+
+TEST_P(ParameterizedBlockSimpleCoarse, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    write(grid, block, size);
+}
+
+TEST_P(ParameterizedBlockSimpleCoarse, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleCoarseTestFixture,
+    ParameterizedBlockSimpleCoarse,
+    ::testing::Values(
+        std::make_tuple(1, 1024, 32),      // 0
+        std::make_tuple(1, 1,    1048576), // 1
+        std::make_tuple(1, 2,    1048576), // 2
+        std::make_tuple(1, 4,    1048576), // 3
+        std::make_tuple(1, 8,    1048576), // 4
+        std::make_tuple(1, 16,   1048576), // 5
+        std::make_tuple(1, 32,   1048576), // 6
+        std::make_tuple(1, 64,   1048576), // 7
+        std::make_tuple(1, 128,  1048576), // 8
+        std::make_tuple(1, 256,  1048576), // 9
+        std::make_tuple(1, 512,  1048576), // 10
+        std::make_tuple(1, 768,  1048576), // 11
+        std::make_tuple(1, 1024, 1048576)) // 12
+);
+
+//=============================================================================
+
+TEST_P(ParameterizedWarpSimpleCoarse, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    write(grid, block, size);
+}
+
+TEST_P(ParameterizedWarpSimpleCoarse, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleCoarseTestFixture,
+    ParameterizedWarpSimpleCoarse,
+    ::testing::Values(
+        std::make_tuple(1, 64,   1),       // 0
+        std::make_tuple(1, 64,   32),      // 1
+        std::make_tuple(1,  1,   1048576), // 2
+        std::make_tuple(1,  2,   1048576), // 3
+        std::make_tuple(1,  3,   1048576), // 4
+        std::make_tuple(1,  4,   1048576), // 5
+        std::make_tuple(1,  5,   1048576), // 6
+        std::make_tuple(1,  6,   1048576), // 7
+        std::make_tuple(1,  7,   1048576), // 8
+        std::make_tuple(1,  8,   1048576), // 9
+        std::make_tuple(1,  9,   1048576), // 10
+        std::make_tuple(1, 10,   1048576), // 11
+        std::make_tuple(1, 11,   1048576), // 12
+        std::make_tuple(1, 12,   1048576), // 13
+        std::make_tuple(1, 13,   1048576), // 14
+        std::make_tuple(1, 14,   1048576), // 15
+        std::make_tuple(1, 15,   1048576), // 16
+        std::make_tuple(1, 16,   1048576), // 17
+        std::make_tuple(1, 17,   1048576), // 18
+        std::make_tuple(1, 18,   1048576), // 19
+        std::make_tuple(1, 19,   1048576), // 20
+        std::make_tuple(1, 20,   1048576), // 21
+        std::make_tuple(1, 21,   1048576), // 22
+        std::make_tuple(1, 22,   1048576), // 23
+        std::make_tuple(1, 23,   1048576), // 24
+        std::make_tuple(1, 24,   1048576), // 25
+        std::make_tuple(1, 25,   1048576), // 26
+        std::make_tuple(1, 26,   1048576), // 27
+        std::make_tuple(1, 27,   1048576), // 28
+        std::make_tuple(1, 28,   1048576), // 29
+        std::make_tuple(1, 28,   1048576), // 30
+        std::make_tuple(1, 29,   1048576), // 31
+        std::make_tuple(1, 30,   1048576), // 32
+        std::make_tuple(1, 31,   1048576), // 33
+        std::make_tuple(1, 32,   1048576), // 34
+        std::make_tuple(1, 33,   1048576), // 35
+        std::make_tuple(1, 34,   1048576), // 36
+        std::make_tuple(1, 35,   1048576), // 37
+        std::make_tuple(1, 36,   1048576), // 38
+        std::make_tuple(1, 37,   1048576), // 39
+        std::make_tuple(1, 38,   1048576), // 40
+        std::make_tuple(1, 39,   1048576), // 41
+        std::make_tuple(1, 40,   1048576), // 42
+        std::make_tuple(1, 41,   1048576), // 43
+        std::make_tuple(1, 42,   1048576), // 44
+        std::make_tuple(1, 43,   1048576), // 45
+        std::make_tuple(1, 44,   1048576), // 46
+        std::make_tuple(1, 45,   1048576), // 47
+        std::make_tuple(1, 46,   1048576), // 48
+        std::make_tuple(1, 47,   1048576), // 49
+        std::make_tuple(1, 48,   1048576), // 50
+        std::make_tuple(1, 49,   1048576), // 51
+        std::make_tuple(1, 50,   1048576), // 52
+        std::make_tuple(1, 51,   1048576), // 53
+        std::make_tuple(1, 52,   1048576), // 54
+        std::make_tuple(1, 53,   1048576), // 55
+        std::make_tuple(1, 54,   1048576), // 56
+        std::make_tuple(1, 55,   1048576), // 57
+        std::make_tuple(1, 56,   1048576), // 58
+        std::make_tuple(1, 57,   1048576), // 59
+        std::make_tuple(1, 58,   1048576), // 60
+        std::make_tuple(1, 59,   1048576), // 61
+        std::make_tuple(1, 60,   1048576), // 62
+        std::make_tuple(1, 61,   1048576), // 63
+        std::make_tuple(1, 62,   1048576), // 64
+        std::make_tuple(1, 63,   1048576), // 65
+        std::make_tuple(1, 64,   1048576)) // 66
+);
+
+//=============================================================================
+
+TEST_P(ParameterizedThreadSimpleCoarse, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    write(grid, block, size);
+}
+
+TEST_P(ParameterizedThreadSimpleCoarse, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleCoarseTestFixture,
+    ParameterizedThreadSimpleCoarse,
+    ::testing::Values(
+        std::make_tuple(1,  1,   1048576))
+);

projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp

@@ -0,0 +1,233 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP
+#define ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include <numeric>
+#include <tuple>
+
+#include <mpi.h>
+#include "../src/memory/symmetric_heap.hpp"
+#include "../src/ipc_policy.hpp"
+
+#include <hip/hip_runtime.h>
+#include <cassert>
+
+
+namespace rocshmem {
+
+__global__
+void
+kernel_simple_coarse_copy(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) {
+    if (!threadIdx.x) {
+      ipc_impl->ipcCopy(dest, src, bytes);
+      ipc_impl->ipcFence();
+    }
+    __syncthreads();
+}
+
+__global__
+void
+kernel_simple_coarse_copy_block(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) {
+    ipc_impl->ipcCopy_wg(dest, src, bytes);
+    ipc_impl->ipcFence();
+    __syncthreads();
+}
+
+__global__
+void
+kernel_simple_coarse_copy_warp(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) {
+    ipc_impl->ipcCopy_wave(dest, src, bytes);
+    ipc_impl->ipcFence();
+    __syncthreads();
+}
+
+class IPCImplSimpleCoarse : public ::testing::TestWithParam<std::tuple<int, int, int>> {
+    using HEAP_T = HeapMemory<HIPAllocator>;
+    using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
+    using FN_T = void (*)(IpcImpl*, int*, int*, size_t);
+
+  public:
+    IPCImplSimpleCoarse() {
+        MPIInstance::mpilib_dl_init();
+        mpi_ = new MPI_T (heap_mem_.get_ptr(), heap_mem_.get_size(), MPI_COMM_WORLD);
+
+        ipc_impl_.ipcHostInit(mpi_->my_pe(), mpi_->get_heap_bases(), MPI_COMM_WORLD);
+        assert(ipc_impl_dptr_ == nullptr);
+        hip_allocator_.allocate((void**)&ipc_impl_dptr_, sizeof(IpcImpl));
+        CHECK_HIP(hipMemcpy(ipc_impl_dptr_, &ipc_impl_,
+                            sizeof(IpcImpl), hipMemcpyHostToDevice));
+    }
+
+    virtual ~IPCImplSimpleCoarse() {
+        if (ipc_impl_dptr_) {
+            hip_allocator_.deallocate(ipc_impl_dptr_);
+        }
+        ipc_impl_.ipcHostStop();
+        MPIInstance::mpilib_dl_close();
+    }
+
+    void launch(FN_T f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes) {
+        f<<<grid, block>>>(ipc_impl_dptr_, src, dest, bytes);
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    enum TestType {
+        READ = 0,
+        WRITE = 1
+    };
+
+    virtual void copy(TestType test, dim3 grid, dim3 block) {
+        FAIL();
+    }
+
+    void write(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(WRITE);
+        copy(WRITE, grid, block);
+        validate_dest_buffer(WRITE);
+    }
+
+    void read(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_src_buffer(READ);
+        copy(READ, grid, block);
+        validate_dest_buffer(READ);
+    }
+
+    void iota_golden(size_t elems) {
+        golden_.resize(elems);
+        std::iota(golden_.begin(), golden_.end(), 0);
+    }
+
+    void validate_golden(size_t elems) {
+        ASSERT_EQ(golden_.size(), elems);
+        for (int i = 0; i < static_cast<int>(golden_.size()); i++) {
+            ASSERT_EQ(golden_[i], i);
+        }
+    }
+
+    void initialize_src_buffer(TestType test) {
+        if (!pe_initializes_src_buffer(test)) {
+            return;
+        }
+        size_t bytes = golden_.size() * sizeof(int);
+        auto dev_src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_->my_pe()]);
+        CHECK_HIP(hipMemcpy(dev_src, golden_.data(), bytes, hipMemcpyHostToDevice));
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    bool pe_initializes_src_buffer(TestType test) {
+        bool is_write_test = test;
+        bool is_read_test = !test;
+        return (is_write_test && mpi_->my_pe() == 0) ||
+               (is_read_test && mpi_->my_pe() == 1);
+    }
+
+    void execute(TestType test, FN_T fn, const dim3 grid, const dim3 block) {
+        if (mpi_->my_pe()) {
+            mpi_->barrier();
+            mpi_->barrier();
+            return;
+        }
+        int *src{nullptr};
+        int *dest{nullptr};
+        if (test == WRITE) {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+        } else {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+        }
+        size_t bytes = golden_.size() * sizeof(int);
+        mpi_->barrier();
+        launch(fn, grid, block, src, dest, bytes);
+        mpi_->barrier();
+    }
+
+    void validate_dest_buffer(TestType test) {
+        if (!pe_validates_dest_buffer(test)) {
+            return;
+        }
+
+        auto dev_dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_->my_pe()]);
+        for (int i = 0; i < static_cast<int>(golden_.size()); i++) {
+            ASSERT_EQ(golden_[i], dev_dest[i]);
+        }
+    }
+
+    bool pe_validates_dest_buffer(TestType test) {
+        return !pe_initializes_src_buffer(test);
+    }
+
+  protected:
+    std::vector<int> golden_;
+
+    HEAP_T heap_mem_ {};
+    MPI_T *mpi_{nullptr};
+
+    IpcImpl ipc_impl_ {};
+    IpcImpl *ipc_impl_dptr_ {nullptr};
+
+    HIPAllocator hip_allocator_ {};
+};
+
+class DegenerateSimpleCoarse : public IPCImplSimpleCoarse {
+  public:
+    ~DegenerateSimpleCoarse() override {};
+};
+
+class ParameterizedBlockSimpleCoarse : public IPCImplSimpleCoarse {
+  public:
+    ~ParameterizedBlockSimpleCoarse() override {};
+
+    void copy(IPCImplSimpleCoarse::TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_simple_coarse_copy_block, grid, block);
+    }
+};
+
+class ParameterizedWarpSimpleCoarse : public IPCImplSimpleCoarse {
+  public:
+    ~ParameterizedWarpSimpleCoarse() override {};
+
+    void copy(IPCImplSimpleCoarse::TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_simple_coarse_copy_warp, grid, block);
+    }
+};
+
+class ParameterizedThreadSimpleCoarse : public IPCImplSimpleCoarse {
+  public:
+    ~ParameterizedThreadSimpleCoarse() override {};
+
+    void copy(IPCImplSimpleCoarse::TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_simple_coarse_copy, grid, block);
+    }
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP

projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp

@@ -0,0 +1,202 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#include "ipc_impl_simple_fine_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_P(DegenerateSimpleFine, ptr_check) {
+    ASSERT_NE(heap_mem_.get_ptr(), nullptr);
+}
+
+TEST_P(DegenerateSimpleFine, MPI_num_pes) {
+    ASSERT_EQ(mpi_->num_pes(), 2);
+}
+
+TEST_P(DegenerateSimpleFine, IPC_bases) {
+    ASSERT_EQ(mpi_->num_pes(), 2);
+    ASSERT_NE(ipc_impl_.ipc_bases, nullptr);
+    for(int i{0}; i < mpi_->num_pes(); i++) {
+        ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr);
+    }
+}
+
+TEST_P(DegenerateSimpleFine, golden_1048576_int) {
+    iota_golden(1048576);
+    validate_golden(1048576);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleFineTestFixture,
+    DegenerateSimpleFine,
+    ::testing::Values(
+        std::make_tuple(1,  1,   1))
+);
+
+//=============================================================================
+
+TEST_P(ParameterizedBlockSimpleFine, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    write(grid, block, size);
+}
+
+TEST_P(ParameterizedBlockSimpleFine, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleFineTestFixture,
+    ParameterizedBlockSimpleFine,
+    ::testing::Values(
+        std::make_tuple(1, 1024, 32),      // 0
+        std::make_tuple(1, 1,    1048576), // 1
+        std::make_tuple(1, 2,    1048576), // 2
+        std::make_tuple(1, 4,    1048576), // 3
+        std::make_tuple(1, 8,    1048576), // 4
+        std::make_tuple(1, 16,   1048576), // 5
+        std::make_tuple(1, 32,   1048576), // 6
+        std::make_tuple(1, 64,   1048576), // 7
+        std::make_tuple(1, 128,  1048576), // 8
+        std::make_tuple(1, 256,  1048576), // 9
+        std::make_tuple(1, 512,  1048576), // 10
+        std::make_tuple(1, 768,  1048576), // 11
+        std::make_tuple(1, 1024, 1048576)) // 12
+);
+
+//=============================================================================
+
+TEST_P(ParameterizedWarpSimpleFine, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    write(grid, block, size);
+}
+
+TEST_P(ParameterizedWarpSimpleFine, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleFineTestFixture,
+    ParameterizedWarpSimpleFine,
+    ::testing::Values(
+        std::make_tuple(1, 64,   1),       // 0
+        std::make_tuple(1, 64,   32),      // 1
+        std::make_tuple(1,  1,   1048576), // 2
+        std::make_tuple(1,  2,   1048576), // 3
+        std::make_tuple(1,  3,   1048576), // 4
+        std::make_tuple(1,  4,   1048576), // 5
+        std::make_tuple(1,  5,   1048576), // 6
+        std::make_tuple(1,  6,   1048576), // 7
+        std::make_tuple(1,  7,   1048576), // 8
+        std::make_tuple(1,  8,   1048576), // 9
+        std::make_tuple(1,  9,   1048576), // 10
+        std::make_tuple(1, 10,   1048576), // 11
+        std::make_tuple(1, 11,   1048576), // 12
+        std::make_tuple(1, 12,   1048576), // 13
+        std::make_tuple(1, 13,   1048576), // 14
+        std::make_tuple(1, 14,   1048576), // 15
+        std::make_tuple(1, 15,   1048576), // 16
+        std::make_tuple(1, 16,   1048576), // 17
+        std::make_tuple(1, 17,   1048576), // 18
+        std::make_tuple(1, 18,   1048576), // 19
+        std::make_tuple(1, 19,   1048576), // 20
+        std::make_tuple(1, 20,   1048576), // 21
+        std::make_tuple(1, 21,   1048576), // 22
+        std::make_tuple(1, 22,   1048576), // 23
+        std::make_tuple(1, 23,   1048576), // 24
+        std::make_tuple(1, 24,   1048576), // 25
+        std::make_tuple(1, 25,   1048576), // 26
+        std::make_tuple(1, 26,   1048576), // 27
+        std::make_tuple(1, 27,   1048576), // 28
+        std::make_tuple(1, 28,   1048576), // 29
+        std::make_tuple(1, 28,   1048576), // 30
+        std::make_tuple(1, 29,   1048576), // 31
+        std::make_tuple(1, 30,   1048576), // 32
+        std::make_tuple(1, 31,   1048576), // 33
+        std::make_tuple(1, 32,   1048576), // 34
+        std::make_tuple(1, 33,   1048576), // 35
+        std::make_tuple(1, 34,   1048576), // 36
+        std::make_tuple(1, 35,   1048576), // 37
+        std::make_tuple(1, 36,   1048576), // 38
+        std::make_tuple(1, 37,   1048576), // 39
+        std::make_tuple(1, 38,   1048576), // 40
+        std::make_tuple(1, 39,   1048576), // 41
+        std::make_tuple(1, 40,   1048576), // 42
+        std::make_tuple(1, 41,   1048576), // 43
+        std::make_tuple(1, 42,   1048576), // 44
+        std::make_tuple(1, 43,   1048576), // 45
+        std::make_tuple(1, 44,   1048576), // 46
+        std::make_tuple(1, 45,   1048576), // 47
+        std::make_tuple(1, 46,   1048576), // 48
+        std::make_tuple(1, 47,   1048576), // 49
+        std::make_tuple(1, 48,   1048576), // 50
+        std::make_tuple(1, 49,   1048576), // 51
+        std::make_tuple(1, 50,   1048576), // 52
+        std::make_tuple(1, 51,   1048576), // 53
+        std::make_tuple(1, 52,   1048576), // 54
+        std::make_tuple(1, 53,   1048576), // 55
+        std::make_tuple(1, 54,   1048576), // 56
+        std::make_tuple(1, 55,   1048576), // 57
+        std::make_tuple(1, 56,   1048576), // 58
+        std::make_tuple(1, 57,   1048576), // 59
+        std::make_tuple(1, 58,   1048576), // 60
+        std::make_tuple(1, 59,   1048576), // 61
+        std::make_tuple(1, 60,   1048576), // 62
+        std::make_tuple(1, 61,   1048576), // 63
+        std::make_tuple(1, 62,   1048576), // 64
+        std::make_tuple(1, 63,   1048576), // 65
+        std::make_tuple(1, 64,   1048576)) // 66
+);
+
+//=============================================================================
+
+TEST_P(ParameterizedThreadSimpleFine, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    write(grid, block, size);
+}
+
+TEST_P(ParameterizedThreadSimpleFine, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplSimpleFineTestFixture,
+    ParameterizedThreadSimpleFine,
+    ::testing::Values(
+        std::make_tuple(1,  1,   1048576))
+);

projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp

@@ -0,0 +1,345 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP
+#define ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include <numeric>
+#include <mpi.h>
+
+#include "../src/atomic.hpp"
+#include "../src/ipc_policy.hpp"
+#include "../src/memory/notifier.hpp"
+#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};
+
+enum TestType {
+    READ = 0,
+    WRITE = 1
+};
+
+__device__
+void
+simple_validator(bool *error, int *golden, int *dest, size_t bytes) {
+    size_t elements {bytes / sizeof(int)};
+    for (size_t i = get_flat_id(); i < elements; i += get_flat_grid_size()) {
+        if (golden[i] != dest[i]) {
+            printf("golden[%zu] %d != dest[%zu] %d\n", i, golden[i], i, dest[i]);
+            *error = true;
+        }
+    }
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_put_with_signal_simple_validator(bool *error, int *golden, int *dest, size_t bytes, NotifierT *notifier) {
+    detail::atomic::rocshmem_memory_orders orders{};
+    if (!get_flat_id()) {
+        while (detail::atomic::load<int, detail::atomic::memory_scope_system>(dest + SIGNAL_OFFSET, orders) == 0) {
+            ;
+        }
+    }
+    notifier->sync();
+    simple_validator(error, golden, dest, bytes);
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_simple_fine_copy(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
+    if (!get_flat_id()) {
+        ipc_impl->ipcCopy(dest, src, bytes);
+        ipc_impl->ipcFence();
+        if (test == WRITE) {
+            ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1);
+        }
+    }
+    if (test == READ) {
+        notifier->sync();
+        simple_validator(error, golden, dest, bytes);
+    }
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_simple_fine_copy_block(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
+    if (!blockIdx.x) {
+        ipc_impl->ipcCopy_wg(dest, src, bytes);
+        ipc_impl->ipcFence();
+        if (test == WRITE) {
+            if (!threadIdx.x) {
+                ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1);
+            }
+        }
+    }
+    if (test == READ) {
+        notifier->sync();
+        simple_validator(error, golden, dest, bytes);
+    }
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_simple_fine_copy_warp(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
+    if (!blockIdx.x && threadIdx.x < 64) {
+        ipc_impl->ipcCopy_wave(dest, src, bytes);
+        ipc_impl->ipcFence();
+        if (test == WRITE) {
+            if (!threadIdx.x) {
+                ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1);
+            }
+        }
+    }
+    __syncthreads();
+    if (test == READ) {
+        notifier->sync();
+        simple_validator(error, golden, dest, bytes);
+    }
+}
+
+class IPCImplSimpleFine : public ::testing::TestWithParam<std::tuple<int, int, int>> {
+    using HEAP_T = HeapMemory<HIPDefaultFinegrainedAllocator>;
+    using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
+    using NotifierT = Notifier<detail::atomic::memory_scope_agent>;
+    using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent>;
+    using FN_T1 = void (*)(IpcImpl*, bool*, int*, int*, int*, size_t, TestType, NotifierT*);
+    using FN_T2 = void (*)(bool*, int*, int*, size_t, NotifierT*);
+
+  public:
+    IPCImplSimpleFine() {
+        MPIInstance::mpilib_dl_init();
+        mpi_ = new MPI_T (heap_mem_.get_ptr(), heap_mem_.get_size(), MPI_COMM_WORLD);
+
+        ipc_impl_.ipcHostInit(mpi_->my_pe(), mpi_->get_heap_bases(), MPI_COMM_WORLD);
+
+        assert(ipc_impl_dptr_ == nullptr);
+        hip_allocator_.allocate((void**)&ipc_impl_dptr_, sizeof(IpcImpl));
+        CHECK_HIP(hipMemcpy(ipc_impl_dptr_, &ipc_impl_, sizeof(IpcImpl), hipMemcpyHostToDevice));
+
+        assert(error_dptr_ == nullptr);
+        hip_allocator_.allocate((void**)&error_dptr_, sizeof(bool));
+        *error_dptr_ = false;
+    }
+
+    ~IPCImplSimpleFine() {
+        if (ipc_impl_dptr_) {
+            hip_allocator_.deallocate(ipc_impl_dptr_);
+        }
+        if (error_dptr_) {
+            hip_allocator_.deallocate(error_dptr_);
+        }
+        if (golden_dptr_) {
+            hip_allocator_.deallocate(golden_dptr_);
+        }
+
+        ipc_impl_.ipcHostStop();
+        MPIInstance::mpilib_dl_close();
+    }
+
+    void launch(FN_T1 f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes, TestType test) {
+        f<<<grid, block>>>(ipc_impl_dptr_, error_dptr_, golden_dptr_, src, dest, bytes, test, notifier_.get());
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    void launch(FN_T2 f, const dim3 grid, const dim3 block, int* dest, size_t bytes) {
+        f<<<grid, block>>>(error_dptr_, golden_dptr_, dest, bytes, notifier_.get());
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    virtual void copy(TestType test, dim3 grid, dim3 block) {
+        FAIL();
+    }
+
+    void write(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_signal(WRITE);
+        initialize_src_buffer(WRITE);
+        copy(WRITE, grid, block);
+        check_device_validation_errors(WRITE);
+    }
+
+    void read(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_signal(READ);
+        initialize_src_buffer(READ);
+        copy(READ, grid, block);
+        check_device_validation_errors(READ);
+    }
+
+    void iota_golden(size_t elems) {
+        golden_.resize(elems);
+        std::iota(golden_.begin(), golden_.end(), 0);
+
+        assert(golden_dptr_ == nullptr);
+        size_t golden_dptr_bytes {golden_.size() * sizeof(int)};
+        hip_allocator_.allocate((void**)&golden_dptr_, golden_dptr_bytes);
+        CHECK_HIP(hipMemcpy(golden_dptr_, golden_.data(), golden_dptr_bytes, hipMemcpyHostToDevice));
+    }
+
+    void validate_golden(size_t elems) {
+        ASSERT_EQ(golden_.size(), elems);
+        for (int i = 0; i < static_cast<int>(golden_.size()); i++) {
+            ASSERT_EQ(golden_[i], i);
+        }
+    }
+
+    void initialize_signal(TestType test) {
+        bool is_write_test = test;
+        if (is_write_test && mpi_->my_pe() == 0) {
+            int *dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+            *(dest + SIGNAL_OFFSET) = 0;
+        }
+    }
+
+    void initialize_src_buffer(TestType test) {
+        if (!pe_initializes_src_buffer(test)) {
+            return;
+        }
+        size_t bytes = golden_.size() * sizeof(int);
+        auto dev_src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_->my_pe()]);
+        CHECK_HIP(hipMemcpy(dev_src, golden_.data(), bytes, hipMemcpyHostToDevice));
+    }
+
+    bool pe_initializes_src_buffer(TestType test) {
+        bool is_write_test = test;
+        bool is_read_test = !test;
+        return (is_write_test && mpi_->my_pe() == 0) ||
+               (is_read_test && mpi_->my_pe() == 1);
+    }
+
+    void execute(TestType test, FN_T1 fn, const dim3 grid, const dim3 block) {
+        size_t bytes = golden_.size() * sizeof(int);
+        if (mpi_->my_pe()) {
+            mpi_->barrier();
+            if (test == WRITE) {
+                int *dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+                FN_T2 val_fn = kernel_put_with_signal_simple_validator;
+                launch(val_fn, grid, block, dest, bytes);
+            }
+            mpi_->barrier();
+            return;
+        }
+        int *src{nullptr};
+        int *dest{nullptr};
+        if (test == WRITE) {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+        } else {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+        }
+        mpi_->barrier();
+        launch(fn, grid, block, src, dest, bytes, test);
+        mpi_->barrier();
+    }
+
+    void check_device_validation_errors(TestType test) {
+        if (!pe_validates_dest_buffer(test)) {
+            return;
+        }
+        ASSERT_EQ(*error_dptr_, false);
+    }
+
+    void validate_dest_buffer(TestType test) {
+        if (!pe_validates_dest_buffer(test)) {
+            return;
+        }
+
+        auto dev_dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_->my_pe()]);
+        for (int i = 0; i < static_cast<int>(golden_.size()); i++) {
+            ASSERT_EQ(golden_[i], dev_dest[i]);
+        }
+    }
+
+    bool pe_validates_dest_buffer(TestType test) {
+        return !pe_initializes_src_buffer(test);
+    }
+
+  protected:
+    HIPDefaultFinegrainedAllocator hip_allocator_ {};
+
+    NotifierProxyT notifier_ {};
+
+    HEAP_T heap_mem_ {};
+
+    MPI_T *mpi_{nullptr};
+
+    std::vector<int> golden_;
+
+    int *golden_dptr_ {nullptr};
+
+    IpcImpl ipc_impl_ {};
+
+    IpcImpl *ipc_impl_dptr_ {nullptr};
+
+    bool *error_dptr_ {nullptr};
+};
+
+class DegenerateSimpleFine : public IPCImplSimpleFine {
+  public:
+    ~DegenerateSimpleFine() override {};
+};
+
+class ParameterizedBlockSimpleFine : public IPCImplSimpleFine {
+  public:
+    ~ParameterizedBlockSimpleFine() override {};
+
+    void copy(TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_simple_fine_copy_block, grid, block);
+    }
+};
+
+class ParameterizedWarpSimpleFine : public IPCImplSimpleFine {
+  public:
+    ~ParameterizedWarpSimpleFine() override {};
+
+    void copy(TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_simple_fine_copy_warp, grid, block);
+    }
+};
+
+class ParameterizedThreadSimpleFine : public IPCImplSimpleFine {
+  public:
+    ~ParameterizedThreadSimpleFine() override {};
+
+    void copy(TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_simple_fine_copy, grid, block);
+    }
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP

projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.cpp

@@ -0,0 +1,316 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#include "ipc_impl_tiled_fine_gtest.hpp"
+
+using namespace rocshmem;
+
+//=============================================================================
+
+TEST_F(DegenerateTiledFine, ptr_check) {
+    ASSERT_NE(heap_mem_.get_ptr(), nullptr);
+}
+
+TEST_F(DegenerateTiledFine, MPI_num_pes) {
+    ASSERT_EQ(mpi_->num_pes(), 2);
+}
+
+TEST_F(DegenerateTiledFine, IPC_bases) {
+    ASSERT_EQ(mpi_->num_pes(), 2);
+    ASSERT_NE(ipc_impl_.ipc_bases, nullptr);
+    for(int i{0}; i < mpi_->num_pes(); i++) {
+        ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr);
+    }
+}
+
+TEST_F(DegenerateTiledFine, golden_1048576_int) {
+    iota_golden(1048576);
+    validate_golden(1048576);
+}
+
+//=============================================================================
+
+int block_signals_calculation(int grid_dim_x, int block_dim_x, size_t size) {
+    size_t bytes = size * sizeof(int);
+    int total_num_threads {grid_dim_x * block_dim_x};
+    int one_grid_iteration_data_size {total_num_threads * THREAD_TRANSFER_GRANULARITY};
+    int num_grid_iterations = ((bytes - THREAD_TRANSFER_GRANULARITY) + one_grid_iteration_data_size) / one_grid_iteration_data_size;
+    int partial_grid_last_iteration_data_size = bytes % one_grid_iteration_data_size;
+
+    int bytes_per_block {block_dim_x * THREAD_TRANSFER_GRANULARITY};
+
+    int num_signals_for_one_full_iteration = grid_dim_x;
+    int num_signals_for_partial_last_iteration = ((partial_grid_last_iteration_data_size - THREAD_TRANSFER_GRANULARITY) + bytes_per_block) / bytes_per_block;
+
+    int num_signals = 0;
+    if (partial_grid_last_iteration_data_size) {
+        num_signals = num_signals_for_one_full_iteration * (num_grid_iterations - 1);
+	num_signals += num_signals_for_partial_last_iteration;
+    } else {
+        num_signals = num_signals_for_one_full_iteration * num_grid_iterations;
+    }
+
+    return num_signals;
+}
+
+TEST_P(ParameterizedBlockTiledFine, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    auto number_of_signals_required = block_signals_calculation(std::get<0>(GetParam()), std::get<1>(GetParam()), size);
+    write(grid, block, size, number_of_signals_required);
+}
+
+TEST_P(ParameterizedBlockTiledFine, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplTiledFineTestFixture,
+    ParameterizedBlockTiledFine,
+    ::testing::Values(
+        std::make_tuple(1,  1024, 32),       // 0
+        std::make_tuple(1,  1024, 2048),     // 1
+        std::make_tuple(1,  1024, 65536),    // 2
+        std::make_tuple(1,  1,    1048576),  // 3
+        std::make_tuple(1,  2,    1048576),  // 4
+        std::make_tuple(1,  4,    1048576),  // 5
+        std::make_tuple(1,  8,    1048576),  // 6
+        std::make_tuple(1,  16,   1048576),  // 7
+        std::make_tuple(1,  32,   1048576),  // 8
+        std::make_tuple(1,  64,   1048576),  // 9
+        std::make_tuple(1,  128,  1048576),  // 10
+        std::make_tuple(1,  256,  1048576),  // 11
+        std::make_tuple(1,  512,  1048576),  // 12
+        std::make_tuple(1,  768,  1048576),  // 13
+        std::make_tuple(1,  1024, 1048576),  // 14
+        std::make_tuple(2,  1,    1048576),  // 15
+        std::make_tuple(2,  2,    1048576),  // 16
+        std::make_tuple(2,  4,    1048576),  // 17
+        std::make_tuple(2,  8,    1048576),  // 18
+        std::make_tuple(2,  16,   1048576),  // 19
+        std::make_tuple(2,  32,   1048576),  // 20
+        std::make_tuple(2,  64,   1048576),  // 21
+        std::make_tuple(2,  128,  1048576),  // 22
+        std::make_tuple(2,  256,  1048576),  // 23
+        std::make_tuple(2,  512,  1048576),  // 24
+        std::make_tuple(2,  768,  1048576),  // 25
+        std::make_tuple(2,  1024, 1048576),  // 26
+        std::make_tuple(4,  1024, 1048576),  // 27
+        std::make_tuple(8,  1024, 1048576),  // 28
+        std::make_tuple(16, 1024, 1048576),  // 29
+        std::make_tuple(32, 1024, 1048576),  // 30
+        std::make_tuple(38, 1024, 1048576),  // 31
+        std::make_tuple(38, 1024, 2097152),  // 32
+        std::make_tuple(38, 1024, 4194304),  // 33
+        std::make_tuple(38, 1024, 8388608),  // 34
+        std::make_tuple(38, 1024, 16777216), // 35
+        std::make_tuple(38, 1024, 33554432)) // 36
+);
+
+//=============================================================================
+
+int warp_signals_calculation(int grid_dim_x, int block_dim_x, size_t size) {
+    size_t bytes = size * sizeof(int);
+    int total_num_threads {grid_dim_x * block_dim_x};
+    int one_grid_iteration_data_size {total_num_threads * THREAD_TRANSFER_GRANULARITY};
+    int num_grid_iterations = ((bytes - THREAD_TRANSFER_GRANULARITY) + one_grid_iteration_data_size) / one_grid_iteration_data_size;
+    int partial_grid_last_iteration_data_size = bytes % one_grid_iteration_data_size;
+
+    int warps_per_block {warpsPerBlock(block_dim_x)};
+    int total_num_warps_in_grid {grid_dim_x * warps_per_block};
+    int bytes_per_warp {min(block_dim_x, WARP_SIZE) * THREAD_TRANSFER_GRANULARITY};
+
+    int num_signals_for_one_full_iteration = total_num_warps_in_grid;
+    int num_signals_for_partial_last_iteration = ((partial_grid_last_iteration_data_size - THREAD_TRANSFER_GRANULARITY) + bytes_per_warp) / bytes_per_warp;
+
+    int num_signals = 0;
+    if (partial_grid_last_iteration_data_size) {
+        num_signals = num_signals_for_one_full_iteration * (num_grid_iterations - 1);
+	num_signals += num_signals_for_partial_last_iteration;
+    } else {
+        num_signals = num_signals_for_one_full_iteration * num_grid_iterations;
+    }
+
+    return num_signals;
+}
+
+TEST_P(ParameterizedWarpTiledFine, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    auto number_of_signals_required = warp_signals_calculation(std::get<0>(GetParam()), std::get<1>(GetParam()), size);
+    write(grid, block, size, number_of_signals_required);
+}
+
+TEST_P(ParameterizedWarpTiledFine, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplTiledFineTestFixture,
+    ParameterizedWarpTiledFine,
+    ::testing::Values(
+        std::make_tuple(1, 64,           8), // 0
+        std::make_tuple(1, 64,          32), // 1
+        std::make_tuple(1,  1,     1048576), // 2
+        std::make_tuple(1,  2,     1048576), // 3
+        std::make_tuple(1,  3,     1048576), // 4
+        std::make_tuple(1,  4,     1048576), // 5
+        std::make_tuple(1,  5,     1048576), // 6
+        std::make_tuple(1,  6,     1048576), // 7
+        std::make_tuple(1,  7,     1048576), // 8
+        std::make_tuple(1,  8,     1048576), // 9
+        std::make_tuple(1,  9,     1048576), // 10
+        std::make_tuple(1,  10,    1048576), // 11
+        std::make_tuple(1,  11,    1048576), // 12
+        std::make_tuple(1,  12,    1048576), // 13
+        std::make_tuple(1,  13,    1048576), // 14
+        std::make_tuple(1,  14,    1048576), // 15
+        std::make_tuple(1,  15,    1048576), // 16
+        std::make_tuple(1,  16,    1048576), // 17
+        std::make_tuple(1,  17,    1048576), // 18
+        std::make_tuple(1,  18,    1048576), // 19
+        std::make_tuple(1,  19,    1048576), // 20
+        std::make_tuple(1,  20,    1048576), // 21
+        std::make_tuple(1,  21,    1048576), // 22
+        std::make_tuple(1,  22,    1048576), // 23
+        std::make_tuple(1,  23,    1048576), // 24
+        std::make_tuple(1,  24,    1048576), // 25
+        std::make_tuple(1,  25,    1048576), // 26
+        std::make_tuple(1,  26,    1048576), // 27
+        std::make_tuple(1,  27,    1048576), // 28
+        std::make_tuple(1,  28,    1048576), // 29
+        std::make_tuple(1,  28,    1048576), // 30
+        std::make_tuple(1,  29,    1048576), // 31
+        std::make_tuple(1,  30,    1048576), // 32
+        std::make_tuple(1,  31,    1048576), // 33
+        std::make_tuple(1,  32,    1048576), // 34
+        std::make_tuple(1,  33,    1048576), // 35
+        std::make_tuple(1,  34,    1048576), // 36
+        std::make_tuple(1,  35,    1048576), // 37
+        std::make_tuple(1,  36,    1048576), // 38
+        std::make_tuple(1,  37,    1048576), // 39
+        std::make_tuple(1,  38,    1048576), // 40
+        std::make_tuple(1,  39,    1048576), // 41
+        std::make_tuple(1,  40,    1048576), // 42
+        std::make_tuple(1,  41,    1048576), // 43
+        std::make_tuple(1,  42,    1048576), // 44
+        std::make_tuple(1,  43,    1048576), // 45
+        std::make_tuple(1,  44,    1048576), // 46
+        std::make_tuple(1,  45,    1048576), // 47
+        std::make_tuple(1,  46,    1048576), // 48
+        std::make_tuple(1,  47,    1048576), // 49
+        std::make_tuple(1,  48,    1048576), // 50
+        std::make_tuple(1,  49,    1048576), // 51
+        std::make_tuple(1,  50,    1048576), // 52
+        std::make_tuple(1,  51,    1048576), // 53
+        std::make_tuple(1,  52,    1048576), // 54
+        std::make_tuple(1,  53,    1048576), // 55
+        std::make_tuple(1,  54,    1048576), // 56
+        std::make_tuple(1,  55,    1048576), // 57
+        std::make_tuple(1,  56,    1048576), // 58
+        std::make_tuple(1,  57,    1048576), // 59
+        std::make_tuple(1,  58,    1048576), // 60
+        std::make_tuple(1,  59,    1048576), // 61
+        std::make_tuple(1,  60,    1048576), // 62
+        std::make_tuple(1,  61,    1048576), // 63
+        std::make_tuple(1,  62,    1048576), // 64
+        std::make_tuple(1,  63,    1048576), // 65
+        std::make_tuple(1,  64,    1048576), // 66
+        std::make_tuple(1,  256,   1048576), // 67
+        std::make_tuple(1,  512,   1048576), // 68
+        std::make_tuple(1,  768,   1048576), // 69
+        std::make_tuple(1,  1024,  1048576), // 70
+        std::make_tuple(2,  32,    1048576), // 71
+        std::make_tuple(2,  64,    1048576), // 72
+        std::make_tuple(2,  128,   1048576), // 73
+        std::make_tuple(2,  256,   1048576), // 74
+        std::make_tuple(2,  512,   1048576), // 75
+        std::make_tuple(2,  1024,  1048576), // 76
+        std::make_tuple(4,  32,    1048576), // 77
+        std::make_tuple(4,  64,    1048576), // 78
+        std::make_tuple(4,  128,   1048576), // 79
+        std::make_tuple(4,  256,   1048576), // 80
+        std::make_tuple(4,  512,   1048576), // 81
+        std::make_tuple(4,  1024,  1048576), // 82
+        std::make_tuple(8,  32,    1048576), // 83
+        std::make_tuple(8,  64,    1048576), // 84
+        std::make_tuple(8,  128,   1048576), // 85
+        std::make_tuple(8,  256,   1048576), // 86
+        std::make_tuple(8,  512,   1048576), // 87
+        std::make_tuple(8,  1024,  1048576), // 88
+        std::make_tuple(16, 32,    1048576), // 89
+        std::make_tuple(16, 64,    1048576), // 90
+        std::make_tuple(16, 128,   1048576), // 91
+        std::make_tuple(16, 256,   1048576), // 92
+        std::make_tuple(16, 512,   1048576), // 93
+        std::make_tuple(16, 1024,  1048576), // 94
+        std::make_tuple(32, 32,    1048576), // 95
+        std::make_tuple(32, 64,    1048576), // 96
+        std::make_tuple(32, 128,   1048576), // 97
+        std::make_tuple(32, 256,   1048576), // 98
+        std::make_tuple(32, 512,   1048576), // 99
+        std::make_tuple(32, 1024,  1048576), // 100
+        std::make_tuple(38, 32,    1048576), // 101
+        std::make_tuple(38, 64,    1048576), // 102
+        std::make_tuple(38, 128,   1048576), // 103
+        std::make_tuple(38, 256,   1048576), // 104
+        std::make_tuple(38, 512,   1048576), // 105
+        std::make_tuple(38, 1024,  1048576), // 106
+        std::make_tuple(38, 1024,  2097152), // 107
+        std::make_tuple(38, 1024,  4194304), // 108
+        std::make_tuple(38, 1024,  8388608), // 109
+        std::make_tuple(38, 1024, 16777216), // 110
+        std::make_tuple(38, 1024, 33554432)) // 111
+);
+
+//=============================================================================
+
+TEST_P(ParameterizedThreadTiledFine, write) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    write(grid, block, size, 1);
+}
+
+TEST_P(ParameterizedThreadTiledFine, read) {
+    dim3 grid = dim3(std::get<0>(GetParam()), 1, 1);
+    dim3 block = dim3(std::get<1>(GetParam()), 1, 1);
+    size_t size = std::get<2>(GetParam());
+    read(grid, block, size);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    IPCImplTiledFineTestFixture,
+    ParameterizedThreadTiledFine,
+    ::testing::Values(
+        std::make_tuple(1,  1,   1048576))
+);

projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp

@@ -0,0 +1,364 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP
+#define ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include <numeric>
+#include <mpi.h>
+
+#include "../src/atomic.hpp"
+#include "../src/ipc_policy.hpp"
+#include "../src/memory/notifier.hpp"
+#include "../src/memory/symmetric_heap.hpp"
+#include "../src/util.hpp"
+
+#include <hip/hip_runtime.h>
+#include <cassert>
+
+namespace rocshmem {
+
+const int WARP_SIZE = 64;
+
+const int THREAD_TRANSFER_GRANULARITY = 8;  // DWORDX2
+
+int warpsPerBlock(size_t block_size) {
+    return ((block_size - 1) + WARP_SIZE) / WARP_SIZE;
+}
+
+// set signal pointer to ipc_impl.ipc_bases location at unused offset
+// be careful with test size to not overrun this location
+const uint32_t SIGNAL_OFFSET {67108864};
+
+enum TestType {
+    READ = 0,
+    WRITE = 1
+};
+
+__device__
+void
+tiled_validator(bool *error, int *golden, int *dest, size_t bytes) {
+    size_t elements {bytes / sizeof(int)};
+    for (size_t i = get_flat_id(); i < elements; i += get_flat_grid_size()) {
+        if (golden[i] != dest[i]) {
+            printf("golden[%zu] %d != dest[%zu] %d\n", i, golden[i], i, dest[i]);
+            *error = true;
+        }
+    }
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_put_with_signal_tiled_validator(bool *error, int *golden, int *dest, size_t bytes, NotifierT *notifier) {
+    detail::atomic::rocshmem_memory_orders orders{};
+    if (!get_flat_id()) {
+        while (detail::atomic::load<int, detail::atomic::memory_scope_system>(dest + SIGNAL_OFFSET, orders) != 0) {
+            ;
+        }
+    }
+    notifier->sync();
+    tiled_validator(error, golden, dest, bytes);
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_tiled_fine_copy(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
+    if (!get_flat_id()) {
+        ipc_impl->ipcCopy(dest, src, bytes);
+        ipc_impl->ipcFence();
+        if (test == WRITE) {
+            ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, -1);
+        }
+    }
+    if (test == READ) {
+        notifier->sync();
+        tiled_validator(error, golden, dest, bytes);
+    }
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_tiled_fine_copy_block(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
+    size_t block_bytes = blockDim.x * THREAD_TRANSFER_GRANULARITY;
+    size_t block_byte_offset = blockIdx.x * block_bytes;
+    for (size_t i = block_byte_offset; i < bytes; i += get_flat_grid_size() * THREAD_TRANSFER_GRANULARITY) {
+	      int chunk = min(block_bytes, bytes - i);
+        ipc_impl->ipcCopy_wg((char*)dest + i, (char*)src + i, chunk);
+        ipc_impl->ipcFence();
+        __syncthreads();
+        if (test == WRITE) {
+            if (!threadIdx.x) {
+                ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, -1);
+            }
+        }
+    }
+    if (test == READ) {
+        notifier->sync();
+        tiled_validator(error, golden, dest, bytes);
+    }
+}
+
+template <typename NotifierT>
+__global__
+void
+kernel_tiled_fine_copy_warp(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
+    size_t warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / WARP_SIZE;
+    size_t warp_bytes = WARP_SIZE * THREAD_TRANSFER_GRANULARITY;
+    size_t warp_byte_offset = warp_id * warp_bytes;
+    for (size_t i = warp_byte_offset; i < bytes; i += get_flat_grid_size() * THREAD_TRANSFER_GRANULARITY) {
+        int chunk = min(warp_bytes, bytes - i);
+        ipc_impl->ipcCopy_wave(((char*)dest) + i, ((char*)src) + i, chunk);
+        ipc_impl->ipcFence();
+        if (test == WRITE) {
+            if (!(threadIdx.x % WARP_SIZE)) {
+                ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, -1);
+            }
+        }
+    }
+    __syncthreads();
+    if (test == READ) {
+        notifier->sync();
+        tiled_validator(error, golden, dest, bytes);
+    }
+}
+
+class IPCImplTiledFine : public ::testing::TestWithParam<std::tuple<int, int, int>> {
+    using HEAP_T = HeapMemory<HIPDefaultFinegrainedAllocator>;
+    using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
+    using NotifierT = Notifier<detail::atomic::memory_scope_agent>;
+    using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent>;
+    using FN_T1 = void (*)(IpcImpl*, bool*, int*, int*, int*, size_t, TestType, NotifierT*);
+    using FN_T2 = void (*)(bool*, int*, int*, size_t, NotifierT*);
+
+  public:
+    IPCImplTiledFine() {
+        MPIInstance::mpilib_dl_init();
+        mpi_ = new MPI_T (heap_mem_.get_ptr(), heap_mem_.get_size(), MPI_COMM_WORLD);
+
+        ipc_impl_.ipcHostInit(mpi_->my_pe(), mpi_->get_heap_bases(), MPI_COMM_WORLD);
+
+        assert(ipc_impl_dptr_ == nullptr);
+        hip_allocator_.allocate((void**)&ipc_impl_dptr_, sizeof(IpcImpl));
+        CHECK_HIP(hipMemcpy(ipc_impl_dptr_, &ipc_impl_, sizeof(IpcImpl), hipMemcpyHostToDevice));
+
+        assert(error_dptr_ == nullptr);
+        hip_allocator_.allocate((void**)&error_dptr_, sizeof(bool));
+        *error_dptr_ = false;
+    }
+
+    ~IPCImplTiledFine() {
+        if (ipc_impl_dptr_) {
+            hip_allocator_.deallocate(ipc_impl_dptr_);
+        }
+        if (error_dptr_) {
+            hip_allocator_.deallocate(error_dptr_);
+        }
+        if (golden_dptr_) {
+            hip_allocator_.deallocate(golden_dptr_);
+        }
+
+        ipc_impl_.ipcHostStop();
+        MPIInstance::mpilib_dl_close();
+    }
+
+    void launch(FN_T1 f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes, TestType test) {
+        f<<<grid, block>>>(ipc_impl_dptr_, error_dptr_, golden_dptr_, src, dest, bytes, test, notifier_.get());
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    void launch(FN_T2 f, const dim3 grid, const dim3 block, int* dest, size_t bytes) {
+        f<<<grid, block>>>(error_dptr_, golden_dptr_, dest, bytes, notifier_.get());
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+    }
+
+    virtual void copy(TestType test, dim3 grid, dim3 block) {
+        FAIL();
+    }
+
+    void write(const dim3 grid, const dim3 block, size_t elems, int signal_value) {
+        iota_golden(elems);
+        initialize_signal(WRITE, signal_value);
+        initialize_src_buffer(WRITE);
+        copy(WRITE, grid, block);
+        check_device_validation_errors(WRITE);
+    }
+
+    void read(const dim3 grid, const dim3 block, size_t elems) {
+        iota_golden(elems);
+        initialize_signal(READ);
+        initialize_src_buffer(READ);
+        copy(READ, grid, block);
+        check_device_validation_errors(READ);
+    }
+
+    void iota_golden(size_t elems) {
+        golden_.resize(elems);
+        std::iota(golden_.begin(), golden_.end(), 0);
+
+        assert(golden_dptr_ == nullptr);
+        size_t golden_dptr_bytes {golden_.size() * sizeof(int)};
+        hip_allocator_.allocate((void**)&golden_dptr_, golden_dptr_bytes);
+        CHECK_HIP(hipMemcpy(golden_dptr_, golden_.data(), golden_dptr_bytes, hipMemcpyHostToDevice));
+    }
+
+    void validate_golden(size_t elems) {
+        ASSERT_EQ(golden_.size(), elems);
+        for (int i = 0; i < static_cast<int>(golden_.size()); i++) {
+            ASSERT_EQ(golden_[i], i);
+        }
+    }
+
+    void initialize_signal(TestType test, int signal_value = 0) {
+        bool is_write_test = test;
+        if (is_write_test && mpi_->my_pe() == 0) {
+            int *dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+            *(dest + SIGNAL_OFFSET) = signal_value;
+        }
+    }
+
+    void initialize_src_buffer(TestType test) {
+        if (!pe_initializes_src_buffer(test)) {
+            return;
+        }
+        size_t bytes = golden_.size() * sizeof(int);
+        auto dev_src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_->my_pe()]);
+        CHECK_HIP(hipMemcpy(dev_src, golden_.data(), bytes, hipMemcpyHostToDevice));
+    }
+
+    bool pe_initializes_src_buffer(TestType test) {
+        bool is_write_test = test;
+        bool is_read_test = !test;
+        return (is_write_test && mpi_->my_pe() == 0) ||
+               (is_read_test && mpi_->my_pe() == 1);
+    }
+
+    void execute(TestType test, FN_T1 fn, const dim3 grid, const dim3 block) {
+        size_t bytes = golden_.size() * sizeof(int);
+        if (mpi_->my_pe()) {
+            mpi_->barrier();
+            if (test == WRITE) {
+                int *dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+                FN_T2 val_fn = kernel_put_with_signal_tiled_validator;
+                launch(val_fn, grid, block, dest, bytes);
+                ASSERT_EQ(*(dest + SIGNAL_OFFSET), 0);
+            }
+            mpi_->barrier();
+            return;
+        }
+        int *src{nullptr};
+        int *dest{nullptr};
+        if (test == WRITE) {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+        } else {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
+        }
+        mpi_->barrier();
+        launch(fn, grid, block, src, dest, bytes, test);
+        mpi_->barrier();
+    }
+
+    void check_device_validation_errors(TestType test) {
+        if (!pe_validates_dest_buffer(test)) {
+            return;
+        }
+        ASSERT_EQ(*error_dptr_, false);
+    }
+
+    void validate_dest_buffer(TestType test) {
+        if (!pe_validates_dest_buffer(test)) {
+            return;
+        }
+
+        auto dev_dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_->my_pe()]);
+        for (int i = 0; i < static_cast<int>(golden_.size()); i++) {
+            ASSERT_EQ(golden_[i], dev_dest[i]);
+        }
+    }
+
+    bool pe_validates_dest_buffer(TestType test) {
+        return !pe_initializes_src_buffer(test);
+    }
+
+  protected:
+    HIPDefaultFinegrainedAllocator hip_allocator_ {};
+
+    NotifierProxyT notifier_ {};
+
+    HEAP_T heap_mem_ {};
+
+    MPI_T *mpi_ {nullptr};
+
+    std::vector<int> golden_;
+
+    int *golden_dptr_ {nullptr};
+
+    IpcImpl ipc_impl_ {};
+
+    IpcImpl *ipc_impl_dptr_ {nullptr};
+
+    bool *error_dptr_ {nullptr};
+};
+
+class DegenerateTiledFine : public IPCImplTiledFine {
+  public:
+    ~DegenerateTiledFine() override {};
+};
+
+class ParameterizedBlockTiledFine : public IPCImplTiledFine {
+  public:
+    ~ParameterizedBlockTiledFine() override {};
+
+    void copy(TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_tiled_fine_copy_block, grid, block);
+    }
+};
+
+class ParameterizedWarpTiledFine : public IPCImplTiledFine {
+  public:
+    ~ParameterizedWarpTiledFine() override {};
+
+    void copy(TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_tiled_fine_copy_warp, grid, block);
+    }
+};
+
+class ParameterizedThreadTiledFine : public IPCImplTiledFine {
+  public:
+    ~ParameterizedThreadTiledFine() override {};
+
+    void copy(TestType test, dim3 grid, dim3 block) override {
+        execute(test, kernel_tiled_fine_copy, grid, block);
+    }
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP

projects/rocshmem/tests/unit_tests/mpi_instance_gtest.cpp

@@ -0,0 +1,37 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "mpi_instance_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(MPIInstanceTestFixture, library_initialize_destroy) {}
+
+TEST_F(MPIInstanceTestFixture, rank) {
+  ASSERT_NO_FATAL_FAILURE(s_ptr_->get_rank());
+}
+
+TEST_F(MPIInstanceTestFixture, nprocs) {
+  ASSERT_EQ(s_ptr_->get_nprocs(), 4);
+}

projects/rocshmem/tests/unit_tests/mpi_instance_gtest.hpp

@@ -0,0 +1,54 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_MPI_INSTANCE_GTEST_HPP
+#define ROCSHMEM_MPI_INSTANCE_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/mpi_instance.hpp"
+
+namespace rocshmem {
+
+class MPIInstanceTestFixture : public ::testing::Test
+{
+  public:
+    MPIInstanceTestFixture() {
+        s_ptr_ = new MPIInstance(MPI_COMM_WORLD);
+    }
+
+    ~MPIInstanceTestFixture() {
+        delete s_ptr_;
+    }
+
+  protected:
+    /**
+     * @brief A MPI instance object used to initialize MPI
+     */
+    MPIInstance* s_ptr_ {nullptr};
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_MPI_INSTANCE_GTEST_HPP

projects/rocshmem/tests/unit_tests/notifier_gtest.cpp

@@ -0,0 +1,127 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "notifier_gtest.hpp"
+
+using namespace rocshmem;
+
+/*****************************************************************************
+ ******************************* Fixture Tests *******************************
+ *****************************************************************************/
+
+TEST_F(NotifierBlockTestFixture, run_all_threads_once_1_1) {
+  run_all_threads_once(1, 1);
+}
+
+TEST_F(NotifierBlockTestFixture, run_all_threads_once_2_1) {
+  run_all_threads_once(2, 1);
+}
+
+TEST_F(NotifierBlockTestFixture, run_all_threads_once_64_1) {
+  run_all_threads_once(64, 1);
+}
+
+TEST_F(NotifierBlockTestFixture, run_all_threads_once_128_1) {
+  run_all_threads_once(128, 1);
+}
+
+TEST_F(NotifierBlockTestFixture, run_all_threads_once_256_1) {
+  run_all_threads_once(256, 1);
+}
+
+TEST_F(NotifierBlockTestFixture, run_all_threads_once_512_1) {
+  run_all_threads_once(512, 1);
+}
+
+TEST_F(NotifierBlockTestFixture, run_all_threads_once_1024_1) {
+  run_all_threads_once(1024, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_1) {
+  run_all_threads_once(1, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_2_1) {
+  run_all_threads_once(2, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_64_1) {
+  run_all_threads_once(64, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_128_1) {
+  run_all_threads_once(128, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_256_1) {
+  run_all_threads_once(256, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_512_1) {
+  run_all_threads_once(512, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_1) {
+  run_all_threads_once(1024, 1);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_2) {
+  run_all_threads_once(1, 2);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_2) {
+  run_all_threads_once(1024, 2);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_4) {
+  run_all_threads_once(1, 4);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_4) {
+  run_all_threads_once(1024, 4);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_8) {
+  run_all_threads_once(1, 8);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_8) {
+  run_all_threads_once(1024, 8);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_32) {
+  run_all_threads_once(1, 32);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_32) {
+  run_all_threads_once(1024, 32);
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_38) {
+  run_all_threads_once(1, 38);  // MI300 CPX
+}
+
+TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_38) {
+  run_all_threads_once(1024, 38);  // MI300 CPX
+}

projects/rocshmem/tests/unit_tests/notifier_gtest.hpp

@@ -0,0 +1,171 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_NOTIFIER_GTEST_HPP
+#define ROCSHMEM_NOTIFIER_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/memory/notifier.hpp"
+#include "../src/util.hpp"
+
+#include <hip/hip_runtime.h>
+#include <cassert>
+
+namespace rocshmem {
+
+/**
+ * @brief The bit pattern written to memory by each thread.
+ */
+static const uint8_t THREAD_VALUE {0xF9};
+
+/**
+ * @brief The bit pattern written to memory by each thread.
+ */
+static const uint64_t NOTIFIER_OFFSET {0x100B00};
+
+inline __device__
+void
+write_to_memory(uint8_t* raw_memory) {
+    auto thread_idx {get_flat_id()};
+    raw_memory[thread_idx] = THREAD_VALUE;
+    __threadfence();
+}
+
+template <typename NotifierT>
+__global__
+void
+all_threads_once(uint8_t* raw_memory,
+                 NotifierT * notifier) {
+    if (!get_flat_id()) {
+      notifier->store(NOTIFIER_OFFSET);
+      notifier->fence();
+    }
+    notifier->sync();
+    uint64_t offset_u64 {notifier->load()};
+    uint64_t raw_memory_u64 {reinterpret_cast<uint64_t>(raw_memory)};
+    uint64_t address_u64 {raw_memory_u64 + offset_u64};
+    uint8_t* address {reinterpret_cast<uint8_t*>(address_u64)};
+    write_to_memory(address);
+}
+
+class NotifierBase : public ::testing::Test {
+  public:
+    NotifierBase() {
+        assert(raw_memory_ == nullptr);
+        hip_allocator_.allocate((void**)&raw_memory_, GIBIBYTE_);
+        assert(raw_memory_);
+    }
+
+    ~NotifierBase() {
+        if (raw_memory_) {
+            hip_allocator_.deallocate(raw_memory_);
+        }
+    }
+
+    void
+    verify(size_t number_threads) {
+        uint8_t* offset_addr {compute_offset_addr()};
+        for (size_t i {0}; i < number_threads; i++) {
+            ASSERT_EQ(offset_addr[i], THREAD_VALUE);
+        }
+    }
+
+  protected:
+    /**
+     * @brief Helper function to reconstruct device calculation.
+     */
+    uint8_t*
+    compute_offset_addr() {
+        uint64_t raw_memory_u64 {reinterpret_cast<uint64_t>(raw_memory_)};
+        uint64_t address_u64 {raw_memory_u64 + NOTIFIER_OFFSET};
+        uint8_t* address {reinterpret_cast<uint8_t*>(address_u64)};
+        return address;
+    }
+
+    /**
+     * @brief An allocator to create objects in device memory.
+     */
+    HIPAllocator hip_allocator_ {};
+
+    /**
+     * @brief The size of the raw memory block below.
+     */
+    static const size_t GIBIBYTE_ {1 << 30};
+
+    /**
+     * @brief A block of memory used to hold individual writes from threads.
+     */
+    uint8_t *raw_memory_ {nullptr};
+
+};
+
+class NotifierBlockTestFixture : public NotifierBase {
+    using NotifierT = Notifier<detail::atomic::memory_scope_workgroup>;
+    using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_workgroup>;
+
+  public:
+    void
+    run_all_threads_once(uint32_t x_block_dim,
+                         uint32_t x_grid_dim) {
+        new (notifier_.get()) NotifierT();
+        const dim3 block(x_block_dim, 1, 1);
+        const dim3 grid(x_grid_dim, 1, 1);
+        all_threads_once<NotifierT><<<grid, block>>>(raw_memory_, notifier_.get());
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+        verify(x_block_dim * x_grid_dim);
+    }
+
+    /**
+     * @brief Used to broadcast base offset for writing.
+     */
+    NotifierProxyT notifier_ {};
+};
+
+class NotifierAgentTestFixture : public NotifierBase {
+    using NotifierT = Notifier<detail::atomic::memory_scope_agent>;
+    using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent>;
+
+  public:
+    void
+    run_all_threads_once(uint32_t x_block_dim,
+                         uint32_t x_grid_dim) {
+        new (notifier_.get()) NotifierT();
+        const dim3 block(x_block_dim, 1, 1);
+        const dim3 grid(x_grid_dim, 1, 1);
+        all_threads_once<NotifierT><<<grid, block>>>(raw_memory_, notifier_.get());
+        CHECK_HIP(hipStreamSynchronize(nullptr));
+        verify(x_block_dim * x_grid_dim);
+    }
+
+    /**
+     * @brief Used to broadcast base offset for writing.
+     */
+    NotifierProxyT notifier_ {};
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_NOTIFIER_GTEST_HPP

projects/rocshmem/tests/unit_tests/pow2_bins_gtest.cpp

@@ -0,0 +1,207 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "pow2_bins_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(Pow2BinsTestFixture, used_0_bytes) {
+  ASSERT_EQ(strat_.get_used(), 0);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_0_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{0};
+  size_t expected_used{0};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_EQ(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used(), expected_used);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_1_byte) {
+  char* c_ptr{nullptr};
+  size_t size{1};
+  size_t align_size{ALIGNMENT * (1 + (size - 1) / ALIGNMENT)};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  size_t min_size{256};
+  ASSERT_LE(align_size, 256); // test fixture won't work for larger values
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[min_size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), align_size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_128_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{128};
+  size_t align_size{ALIGNMENT * (1 + (size - 1) / ALIGNMENT)};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  size_t min_size{256};
+  ASSERT_LE(align_size, 256);
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[min_size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), align_size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_256_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{256};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_512_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{512};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_513_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{513};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  size_t min_size{1024};
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[min_size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), min_size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_4095_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{4095};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  size_t min_size{4096};
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[min_size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), min_size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_4KB) {
+  char* c_ptr{nullptr};
+  size_t size{4096};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_4097_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{4097};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  size_t min_size{8192};
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[min_size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), min_size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_128KB) {
+  char* c_ptr{nullptr};
+  size_t size{131072};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_1GB) {
+  char* c_ptr{nullptr};
+  size_t size{1 << 30};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[size]};
+  ASSERT_EQ(bin.size(), 0);
+  ASSERT_EQ(strat_.get_used(), size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_256_bytes_X2_free_256_bytes_X2) {
+  char* c_ptr_1{nullptr};
+  char* c_ptr_2{nullptr};
+  size_t size{256};
+  strat_.alloc(&c_ptr_1, size);
+  ASSERT_NE(c_ptr_1, nullptr);
+  strat_.alloc(&c_ptr_2, size);
+  ASSERT_NE(c_ptr_2, nullptr);
+
+  auto bins{strat_.get_bins()};
+  auto& bin{(*bins)[size]};
+  ASSERT_EQ(bin.size(), 0);
+  ASSERT_EQ(strat_.get_used(), 2 * size);
+
+  strat_.free(c_ptr_1);
+  strat_.free(c_ptr_2);
+
+  ASSERT_EQ(bin.size(), 2);
+  ASSERT_EQ(strat_.get_used(), 0);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_1GB_free_1GB) {
+  char* c_ptr{nullptr};
+  size_t size{1 << 30};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  auto bins{strat_.get_bins()};
+  auto& bin{(*bins)[size]};
+  ASSERT_EQ(bin.size(), 0);
+  ASSERT_EQ(strat_.get_used(), size);
+
+  strat_.free(c_ptr);
+
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), 0);
+}

projects/rocshmem/tests/unit_tests/pow2_bins_gtest.hpp

@@ -0,0 +1,68 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_POW2_BINS_GTEST_HPP
+#define ROCSHMEM_POW2_BINS_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/address_record.hpp"
+#include "../src/memory/heap_memory.hpp"
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/memory/pow2_bins.hpp"
+
+namespace rocshmem {
+
+class Pow2BinsTestFixture : public ::testing::Test
+{
+    /**
+     * @brief Helper type for heap memory
+     */
+    using HEAP_T = HeapMemory<HIPAllocator>;
+
+    /**
+     * @brief Helper type for address records
+     */
+    using AR_T = AddressRecord;
+
+    /**
+     * @brief Helper type for allocation strategy
+     */
+    using STRAT_T = Pow2Bins<AR_T, HEAP_T>;
+
+  protected:
+    /**
+     * @brief Heap memory object
+     */
+    HEAP_T heap_mem_ {};
+
+    /**
+     * @brief Allocation strategy object
+     */
+    STRAT_T strat_ {&heap_mem_};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_POW2_BINS_GTEST_HPP

projects/rocshmem/tests/unit_tests/remote_heap_info_gtest.cpp

@@ -0,0 +1,39 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "remote_heap_info_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(RemoteHeapInfoTestFixture, MPI_num_pes) { ASSERT_EQ(mpi_.num_pes(), 4); }
+
+TEST_F(RemoteHeapInfoTestFixture, MPI_barrier) {
+  ASSERT_NO_FATAL_FAILURE(mpi_.barrier());
+}
+
+TEST_F(RemoteHeapInfoTestFixture, MPI_bases) {
+  for (auto base : mpi_.get_heap_bases()) {
+    ASSERT_NE(base, nullptr);
+  }
+}

projects/rocshmem/tests/unit_tests/remote_heap_info_gtest.hpp

@@ -0,0 +1,66 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_REMOTE_HEAP_INFO_GTEST_HPP
+#define ROCSHMEM_REMOTE_HEAP_INFO_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include <mpi.h>
+
+#include "../src/memory/heap_memory.hpp"
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/memory/remote_heap_info.hpp"
+
+namespace rocshmem {
+
+class RemoteHeapInfoTestFixture : public ::testing::Test
+{
+    /**
+     * @brief Helper type for heap memory
+     */
+    using HEAP_T = HeapMemory<HIPAllocator>;
+
+    /**
+     * @brief Helper type for RemoteHeapInfo with MPI
+     */
+    using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
+
+  protected:
+    /**
+     * @brief Heap memory object
+     */
+    HEAP_T heap_mem_ {};
+
+    /**
+     * @brief Remote heap info with MPI Communicator
+     */
+    MPI_T mpi_ {heap_mem_.get_ptr(),
+                heap_mem_.get_size(),
+                MPI_COMM_WORLD};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_REMOTE_HEAP_INFO_GTEST_HPP

projects/rocshmem/tests/unit_tests/shmem_gtest.cpp

@@ -0,0 +1,49 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "gtest/gtest.h"
+#include <mpi.h>
+
+int main(int argc, char **argv) {
+  ::testing::InitGoogleTest(&argc, argv);
+  
+  int initialized;
+  MPI_Initialized(&initialized);
+  if (!initialized) {
+    int provided;
+    MPI_Init_thread(nullptr, nullptr, MPI_THREAD_MULTIPLE, &provided);
+    if (provided != MPI_THREAD_MULTIPLE) {
+      std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n";
+    }
+  }
+
+  int ret_val = RUN_ALL_TESTS();
+  
+  int finalized{0};
+  MPI_Finalized(&finalized);
+  if (!finalized) {
+    MPI_Finalize();
+  }
+  return ret_val;
+}

projects/rocshmem/tests/unit_tests/single_heap_gtest.cpp

@@ -0,0 +1,186 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "single_heap_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(SingleHeapTestFixture, unallocated_size_check) {
+  ASSERT_EQ(single_heap_.get_size(), 1 << 30);
+}
+
+TEST_F(SingleHeapTestFixture, free_null) {
+  void* ptr{nullptr};
+  single_heap_.free(ptr);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_0) {
+  // some allocators (e.g. dlmalloc) use memory for internal bookkeeping
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{0};
+  void* ptr{nullptr};
+
+  single_heap_.malloc(&ptr, request_size);
+  ASSERT_EQ(ptr, nullptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+
+  single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_1) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{1};
+  void* ptr{nullptr};
+
+  single_heap_.malloc(&ptr, request_size);
+  ASSERT_NE(ptr, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_256) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{256};
+  void* ptr{nullptr};
+
+  single_heap_.malloc(&ptr, request_size);
+  ASSERT_NE(ptr, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_1024) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{1024};
+  void* ptr{nullptr};
+
+  single_heap_.malloc(&ptr, request_size);
+  ASSERT_NE(ptr, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_1MB) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{1 << 20};
+  void* ptr{nullptr};
+
+  single_heap_.malloc(&ptr, request_size);
+  ASSERT_NE(ptr, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_4097) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{4097};
+  void* ptr{nullptr};
+
+  single_heap_.malloc(&ptr, request_size);
+  ASSERT_NE(ptr, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_X2_8191) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{8191};
+  void* ptr_1{nullptr};
+  void* ptr_2{nullptr};
+
+  single_heap_.malloc(&ptr_1, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_1) & (ALIGNMENT-1), 0);
+
+  single_heap_.malloc(&ptr_2, request_size);
+  ASSERT_NE(ptr_2, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_2) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_1);
+  single_heap_.free(ptr_2);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_X2_free_alloc_free_X2_1MB) {
+  size_t initial_used{single_heap_.get_used()};
+  void* ptr_1{nullptr};
+  void* ptr_2{nullptr};
+  void* ptr_3{nullptr};
+  size_t request_size{1 << 20};
+
+  single_heap_.malloc(&ptr_1, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_1) & (ALIGNMENT-1), 0);
+
+  single_heap_.malloc(&ptr_2, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_2) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_1);
+
+  single_heap_.malloc(&ptr_3, request_size);
+  ASSERT_NE(ptr_3, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_3) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_3);
+  single_heap_.free(ptr_2);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_X2_free_alloc_free_X2_63) {
+  size_t initial_used{single_heap_.get_used()};
+  void* ptr_1{nullptr};
+  void* ptr_2{nullptr};
+  void* ptr_3{nullptr};
+  size_t request_size{63};
+
+  single_heap_.malloc(&ptr_1, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_1) & (ALIGNMENT-1), 0);
+
+  single_heap_.malloc(&ptr_2, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_2) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_1);
+
+  single_heap_.malloc(&ptr_3, request_size);
+  ASSERT_NE(ptr_3, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_3) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_3);
+  single_heap_.free(ptr_2);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}

projects/rocshmem/tests/unit_tests/single_heap_gtest.hpp

@@ -0,0 +1,45 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_SINGLE_HEAP_GTEST_HPP
+#define ROCSHMEM_SINGLE_HEAP_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/single_heap.hpp"
+
+namespace rocshmem {
+
+class SingleHeapTestFixture : public ::testing::Test
+{
+  protected:
+    /**
+     * @brief Single heap object
+     */
+    SingleHeap single_heap_ {};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_SINGLE_HEAP_GTEST_HPP

projects/rocshmem/tests/unit_tests/symmetric_heap_gtest.cpp

@@ -0,0 +1,57 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "symmetric_heap_gtest.hpp"
+
+using namespace rocshmem;
+
+TEST_F(SymmetricHeapTestFixture, malloc_free) {
+  void *ptr{nullptr};
+  size_t request_bytes{48};
+
+  symmetric_heap_->malloc(&ptr, request_bytes);
+  ASSERT_NE(ptr, nullptr);
+  ASSERT_NO_FATAL_FAILURE(symmetric_heap_->free(ptr));
+}
+
+TEST_F(SymmetricHeapTestFixture, window_info) {
+  auto win_info_ptr{symmetric_heap_->get_window_info()};
+
+  WindowInfoMPI* window_info_mpi = dynamic_cast<WindowInfoMPI*>(win_info_ptr);
+  if (window_info_mpi) {
+    void *window_base_addr{nullptr};
+    int flag{0};
+    MPI_Win_get_attr(window_info_mpi->get_win(), MPI_WIN_BASE, &window_base_addr,
+                     &flag);
+    ASSERT_NE(0, flag);
+    ASSERT_NE(nullptr, window_base_addr);
+  }
+}
+
+TEST_F(SymmetricHeapTestFixture, heap_bases) {
+  auto heap_bases{symmetric_heap_->get_heap_bases()};
+  for (const auto &base : heap_bases) {
+    ASSERT_NE(nullptr, base);
+  }
+}

projects/rocshmem/tests/unit_tests/symmetric_heap_gtest.hpp

@@ -0,0 +1,56 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_SYMMETRIC_HEAP_GTEST_HPP
+#define ROCSHMEM_SYMMETRIC_HEAP_GTEST_HPP
+
+#include <mpi.h>
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/symmetric_heap.hpp"
+
+namespace rocshmem {
+
+class SymmetricHeapTestFixture : public ::testing::Test
+{
+  protected:
+    /**
+     * @brief Symmetric heap object
+     */
+    SymmetricHeap *symmetric_heap_;
+
+    void SetUp() override {
+      MPIInstance::mpilib_dl_init();
+      symmetric_heap_ = new SymmetricHeap(MPI_COMM_WORLD);
+    }
+
+    void TearDown() override {
+      MPIInstance::mpilib_dl_close();
+    }
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_SYMMETRIC_HEAP_GTEST_HPP

projects/rocshmem/tests/unit_tests/wavefront_size_gtest.cpp

@@ -0,0 +1,61 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "wavefront_size_gtest.hpp"
+
+#include "../src/util.hpp"
+
+using namespace rocshmem;
+
+__global__ void check_wf_size(int wf_size_prop, int *ret) {
+  if (wf_size_prop == WF_SIZE) {
+    *ret = 0;
+  } else {
+    *ret = 1;
+  }
+}
+
+
+TEST_F(WavefrontSizeTestFixture, constant_matches_runtime) {
+  int device_count = 0;
+  hipDeviceProp_t prop;
+  int *ret;
+
+  CHECK_HIP(hipGetDeviceCount(&device_count));
+  ASSERT_GT(device_count, 0);
+  CHECK_HIP(hipHostMalloc(&ret, sizeof(int), 0));
+
+  for (int i = 0; i < device_count; i++) {
+    *ret = -1;
+    CHECK_HIP(hipSetDevice(i));
+    CHECK_HIP(hipGetDeviceProperties(&prop, i));
+
+    check_wf_size<<<1, 1>>>(prop.warpSize, ret);
+    CHECK_HIP(hipDeviceSynchronize());
+
+    ASSERT_EQ(*ret, 0);
+  }
+  CHECK_HIP(hipHostFree(ret));
+}
+

projects/rocshmem/tests/unit_tests/wavefront_size_gtest.hpp

@@ -0,0 +1,45 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_WAVEFRONT_SIZE_GTEST_HPP
+#define ROCSHMEM_WAVEFRONT_SIZE_GTEST_HPP
+
+#include "gtest/gtest.h"
+#include "wf_size.hpp"
+
+namespace rocshmem {
+
+class WavefrontSizeTestFixture : public ::testing::Test {
+public:
+  void SetUp() override {
+      wf_size = get_wf_size();
+  }
+
+protected:
+  int wf_size;
+};
+
+} // namespace rocshmem
+
+#endif  // ROCSHMEM_WAVEFRONT_SIZE_GTEST_HPP

projects/rocshmem/tests/unit_tests/wf_size.hpp

@@ -0,0 +1,46 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_WF_SIZE_HPP
+#define ROCSHMEM_WF_SIZE_HPP
+
+#include <hip/hip_runtime.h>
+#include "mpi.h"
+
+#define CHECK_HIP_MPI(cond) {                                      \
+  if(cond != hipSuccess){                                          \
+     fprintf(stderr,"HIP error: %d line: %d\n", cond,  __LINE__);  \
+     MPI_Abort(MPI_COMM_WORLD, 1);  	     	  	 	   \
+  }                                                                \
+}
+
+static int get_wf_size() {
+  int deviceId;
+  hipDeviceProp_t prop;
+  CHECK_HIP_MPI(hipGetDevice(&deviceId));
+  CHECK_HIP_MPI(hipGetDeviceProperties(&prop, deviceId));
+  return prop.warpSize;
+}
+
+#endif