From 325ce3cba74e68345d027bb1798bb819b7a60fc0 Mon Sep 17 00:00:00 2001
From: Brandon Potter <brandon.potter@amd.com>
Date: Thu, 26 Sep 2024 13:40:05 -0500
Subject: [PATCH] Bugfixes for the ipc unit tests

[ROCm/rocshmem commit: f85c46ec0a4c34ac599949c09c65da9fa4d279e3]
---
 projects/rocshmem/src/ipc_policy.hpp              |  2 +-
 projects/rocshmem/src/memory/notifier.hpp         |  8 ++++++++
 .../unit_tests/ipc_impl_simple_fine_gtest.cpp     |  6 ++++++
 .../unit_tests/ipc_impl_simple_fine_gtest.hpp     | 15 +++++++++++++++
 4 files changed, 30 insertions(+), 1 deletion(-)

diff --git a/projects/rocshmem/src/ipc_policy.hpp b/projects/rocshmem/src/ipc_policy.hpp
index 48b7b7fd40..1b84c52613 100644
--- a/projects/rocshmem/src/ipc_policy.hpp
+++ b/projects/rocshmem/src/ipc_policy.hpp
@@ -64,7 +64,7 @@ class IpcOnImpl {
 
   __device__ void ipcCopy_wave(void *dst, void *src, size_t size);
 
-  __device__ void ipcFence() { __threadfence(); }
+  __device__ void ipcFence() { __threadfence_system(); }
 
   template <typename T>
   __device__ T ipcAMOFetchAdd(T *val, T value) {
diff --git a/projects/rocshmem/src/memory/notifier.hpp b/projects/rocshmem/src/memory/notifier.hpp
index bcc030ce0a..b72a256f15 100644
--- a/projects/rocshmem/src/memory/notifier.hpp
+++ b/projects/rocshmem/src/memory/notifier.hpp
@@ -98,6 +98,14 @@ class NotifierProxy {
   using ProxyT = DeviceProxy<ALLOCATOR, Notifier<scope>>;
 
  public:
+  NotifierProxy() {
+    new (proxy_.get()) Notifier<scope>();
+  }
+
+  ~NotifierProxy() {
+    proxy_.get()->~Notifier<scope>();
+  }
+
   __host__ __device__ Notifier<scope>* get() { return proxy_.get(); }
 
  private:
diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp
index 7c752e2f6a..71cb85dc9e 100644
--- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp
+++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp
@@ -46,6 +46,12 @@ TEST_F(IPCImplSimpleFineTestFixture, golden_1048576_int) {
 
 //=============================================================================
 
+TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_1x1x1_32_int) {
+    dim3 grid {1,1,1};
+    dim3 block {1,1,1};
+    write_wg(grid, block, 32);
+}
+
 TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_1024x1x1_32_int) {
     dim3 grid {1,1,1};
     dim3 block {1024,1,1};
diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp
index 8e9c28ac46..ba4c20e288 100644
--- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp
+++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp
@@ -49,6 +49,7 @@ validator(bool *error, int *golden, int *dest, size_t bytes) {
     size_t elements {bytes / sizeof(int)};
     for (int i {get_flat_id()}; i < elements; i += get_flat_grid_size()) {
         if (golden[i] != dest[i]) {
+            printf("golden[%d] %d != dest[%d] %d\n", i, golden[i], i, dest[i]);
             *error = true;
         }
     }
@@ -171,6 +172,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
 
     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);
@@ -178,6 +180,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
 
     void write_wg(const dim3 grid, const dim3 block, size_t elems) {
         iota_golden(elems);
+        initialize_signal(WRITE);
         initialize_src_buffer(WRITE);
         copy_wg(WRITE, grid, block);
         check_device_validation_errors(WRITE);
@@ -185,6 +188,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
 
     void write_wave(const dim3 grid, const dim3 block, size_t elems) {
         iota_golden(elems);
+        initialize_signal(WRITE);
         initialize_src_buffer(WRITE);
         copy_wave(WRITE, grid, block);
         check_device_validation_errors(WRITE);
@@ -192,6 +196,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
 
     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);
@@ -199,6 +204,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
 
     void read_wg(const dim3 grid, const dim3 block, size_t elems) {
         iota_golden(elems);
+        initialize_signal(READ);
         initialize_src_buffer(READ);
         copy_wg(READ, grid, block);
         check_device_validation_errors(READ);
@@ -206,6 +212,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
 
     void read_wave(const dim3 grid, const dim3 block, size_t elems) {
         iota_golden(elems);
+        initialize_signal(READ);
         initialize_src_buffer(READ);
         copy_wave(READ, grid, block);
         check_device_validation_errors(READ);
@@ -228,6 +235,14 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test {
         }
     }
 
+    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;