Refactor Barrier_all and Sync_all APIs to use default context (#159)

* Refactor `Barrier_all` and `Sync_all` to use default context

- Removed context-specific implementations of barrier_all and sync_all
- Added barrier_all and sync_all to the default context implementation
- Updated functional tests to use the default context for barrier_all and sync_all

* Update `Barrier_all` and `Sync_all` API usage in documentation

* Update `CHANGELOG`

---------

Co-authored-by: Yiltan <ytemucin@amd.com>

src/ipc/backend_ipc.cpp

@@ -265,10 +265,9 @@ void IPCBackend::init_wrk_sync_buffer() {
   auto max_num_teams{team_tracker.get_max_num_teams()};
 
   /**
-   * size of barrier sync for all the contexts
+   * size of barrier sync
    */
-  Wrk_Sync_buffer_size_ += sizeof(*barrier_sync) * ROCSHMEM_BARRIER_SYNC_SIZE *
-                           (maximum_num_contexts_ + 1);
+  Wrk_Sync_buffer_size_ += sizeof(*barrier_sync) * ROCSHMEM_BARRIER_SYNC_SIZE;
 
   /**
    * Size of sync arrays for the teams
@@ -367,9 +366,7 @@ void IPCBackend::rocshmem_collective_init() {
    * Allocate heap space for barrier_sync
    */
   size_t one_sync_size_bytes {sizeof(*barrier_sync)};
-  size_t total_sync_elems {
-    ROCSHMEM_BARRIER_SYNC_SIZE * (maximum_num_contexts_ + 1)};
-  size_t sync_size_bytes {one_sync_size_bytes * total_sync_elems};
+  size_t sync_size_bytes {one_sync_size_bytes * ROCSHMEM_BARRIER_SYNC_SIZE};
 
   barrier_sync = reinterpret_cast<int64_t*>(temp_Wrk_Sync_buff_ptr_);
   temp_Wrk_Sync_buff_ptr_ += sync_size_bytes;
@@ -377,7 +374,7 @@ void IPCBackend::rocshmem_collective_init() {
   /*
    * Initialize the barrier synchronization array with default values.
    */
-  for (int i = 0; i < total_sync_elems; i++) {
+  for (int i = 0; i < num_pes; i++) {
     barrier_sync[i] = ROCSHMEM_SYNC_VALUE;
   }
 

src/ipc/context_ipc_device.cpp

@@ -44,9 +44,7 @@ __host__ IPCContext::IPCContext(Backend *b, unsigned int ctx_id)
   ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases;
   ipcImpl_.shm_size = b->ipcImpl.shm_size;
 
-  size_t barrier_sync_offset = ctx_id * ROCSHMEM_BARRIER_SYNC_SIZE;
-
-  barrier_sync = backend->barrier_sync + barrier_sync_offset;
+  barrier_sync = backend->barrier_sync;
   fence_pool = backend->fence_pool;
   Wrk_Sync_buffer_bases_ = backend->get_wrk_sync_bases();
   ctx_id_ = ctx_id;

src/rocshmem_gpu.cpp

@@ -594,25 +594,25 @@ __device__ int rocshmem_test(T *ivars, int cmp, T val) {
   return ctx_internal->test(ivars, cmp, val);
 }
 
-__device__ void rocshmem_ctx_barrier_all(rocshmem_ctx_t ctx) {
-  GPU_DPRINTF("Function: rocshmem_ctx_barrier_all (ctx=%zd)\n",
-    ctx.ctx_opaque);
+__device__ void rocshmem_barrier_all() {
+  GPU_DPRINTF("Function: rocshmem_barrier_all (ctx=%zd)\n",
+    get_internal_ctx(ROCSHMEM_CTX_DEFAULT));
 
-  get_internal_ctx(ctx)->barrier_all();
+  get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->barrier_all();
 }
 
-__device__ void rocshmem_ctx_barrier_all_wave(rocshmem_ctx_t ctx) {
-  GPU_DPRINTF("Function: rocshmem_ctx_barrier_all_wave (ctx=%zd)\n",
-    ctx.ctx_opaque);
+__device__ void rocshmem_barrier_all_wave() {
+  GPU_DPRINTF("Function: rocshmem_barrier_all_wave (ctx=%zd)\n",
+    get_internal_ctx(ROCSHMEM_CTX_DEFAULT));
 
-  get_internal_ctx(ctx)->barrier_all_wave();
+  get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->barrier_all_wave();
 }
 
-__device__ void rocshmem_ctx_barrier_all_wg(rocshmem_ctx_t ctx) {
-  GPU_DPRINTF("Function: rocshmem_ctx_barrier_all_wg (ctx=%zd)\n",
-    ctx.ctx_opaque);
+__device__ void rocshmem_barrier_all_wg() {
+  GPU_DPRINTF("Function: rocshmem_barrier_all_wg (ctx=%zd)\n",
+    get_internal_ctx(ROCSHMEM_CTX_DEFAULT));
 
-  get_internal_ctx(ctx)->barrier_all_wg();
+  get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->barrier_all_wg();
 }
 
 __device__ void rocshmem_ctx_barrier(rocshmem_ctx_t ctx, rocshmem_team_t team) {
@@ -636,25 +636,25 @@ __device__ void rocshmem_ctx_barrier_wg(rocshmem_ctx_t ctx, rocshmem_team_t team
   get_internal_ctx(ctx)->barrier_wg(team);
 }
 
-__device__ void rocshmem_ctx_sync_all(rocshmem_ctx_t ctx) {
-  GPU_DPRINTF("Function: rocshmem_ctx_sync_all (ctx=%zd)\n",
-    ctx.ctx_opaque);
+__device__ void rocshmem_sync_all() {
+  GPU_DPRINTF("Function: rocshmem_sync_all (ctx=%zd)\n",
+    get_internal_ctx(ROCSHMEM_CTX_DEFAULT));
 
-  get_internal_ctx(ctx)->sync_all();
+  get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->sync_all();
 }
 
-__device__ void rocshmem_ctx_sync_all_wave(rocshmem_ctx_t ctx) {
-  GPU_DPRINTF("Function: rocshmem_ctx_sync_all_wave (ctx=%zd)\n",
-    ctx.ctx_opaque);
+__device__ void rocshmem_sync_all_wave() {
+  GPU_DPRINTF("Function: rocshmem_sync_all_wave (ctx=%zd)\n",
+    get_internal_ctx(ROCSHMEM_CTX_DEFAULT));
 
-  get_internal_ctx(ctx)->sync_all_wave();
+  get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->sync_all_wave();
 }
 
-__device__ void rocshmem_ctx_sync_all_wg(rocshmem_ctx_t ctx) {
-  GPU_DPRINTF("Function: rocshmem_ctx_sync_all_wg (ctx=%zd)\n",
-    ctx.ctx_opaque);
+__device__ void rocshmem_sync_all_wg() {
+  GPU_DPRINTF("Function: rocshmem_sync_all_wg (ctx=%zd)\n",
+    get_internal_ctx(ROCSHMEM_CTX_DEFAULT));
 
-  get_internal_ctx(ctx)->sync_all_wg();
+  get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->sync_all_wg();
 }
 
 __device__ void rocshmem_ctx_sync(rocshmem_ctx_t ctx,