fixed notifier bug

src/ipc/backend_ipc.hpp

@@ -32,6 +32,10 @@
 #include "ipc_context_proxy.hpp"
 #include "../ipc_policy.hpp"
 
+#ifndef USE_COOPERATIVE_GROUPS
+#include "../memory/notifier.hpp"
+#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
+
 namespace rocshmem {
 
 class IPCBackend : public Backend {
@@ -104,7 +108,7 @@ class IPCBackend : public Backend {
                        TeamInfo *team_info_wrt_world, int num_pes,
                        int my_pe_in_new_team, MPI_Comm team_comm,
                        roc_shmem_team_t *new_team) override;
-  
+
   /**
    * @copydoc Backend::team_destroy(roc_shmem_team_t)
    */
@@ -161,6 +165,10 @@ class IPCBackend : public Backend {
   */
   int *fence_pool{nullptr};
 
+#ifndef USE_COOPERATIVE_GROUPS
+  NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent> notifier_{};
+#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
+
  protected:
    /**
    * @copydoc Backend::dump_backend_stats()
@@ -259,4 +267,4 @@ class IPCBackend : public Backend {
 
 }  // namespace rocshmem
 
-#endif  // LIBRARY_SRC_IPC_BACKEND_HPP_
+#endif  // LIBRARY_SRC_IPC_BACKEND_HPP_

src/ipc/context_ipc_device.cpp

@@ -50,6 +50,10 @@ __host__ IPCContext::IPCContext(Backend *b)
   fence_pool = backend->fence_pool;
 
   orders_.store = detail::atomic::rocshmem_memory_order::memory_order_seq_cst;
+
+#ifndef USE_COOPERATIVE_GROUPS
+  notifier_ = backend->notifier_.get();
+#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
 }
 
 __device__ void IPCContext::threadfence_system() {

src/ipc/context_ipc_device.hpp

@@ -243,8 +243,7 @@ class IPCContext : public Context {
   int64_t *barrier_sync{nullptr};
 
 #ifndef USE_COOPERATIVE_GROUPS
-  using NOTIFIER_PROXY_T = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent>;
-  NOTIFIER_PROXY_T notifier_{};
+  Notifier<detail::atomic::memory_scope_agent> *notifier_{nullptr};
 #endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
 
   //Struct defining memory ordering for atomic operations.

src/ipc/context_ipc_device_coll.cpp

@@ -89,9 +89,8 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
   grid.sync();
   if (0 == grid.thread_rank())
 #else
-  auto notifier{notifier_.get()};
-  notifier->sync();
-  if (is_thread_zero_in_block() && is_block_zero_in_grid())
+  notifier_->sync();
+  if (0 == get_flat_id())
 #endif /* USE_COOPERATIVE_GROUPS */
   {
     if (PE_size < 64) {
@@ -104,7 +103,7 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
 #ifdef USE_COOPERATIVE_GROUPS
   grid.sync();
 #else
-  notifier->sync();
+  notifier_->sync();
 #endif /* USE_COOPERATIVE_GROUPS */
 }