Merge branch 'ROCm:develop' into active_set_APIs

CMakeLists.txt

@@ -68,7 +68,6 @@ option(USE_FUNC_CALL "Force compiler to use function calls on library API" OFF)
 option(USE_SHARED_CTX "Request support for shared ctx between WG" OFF)
 option(USE_SINGLE_NODE "Enable single node support only." OFF)
 option(USE_HOST_SIDE_HDP_FLUSH "Use a polling thread to flush the HDP cache on the host." OFF)
-option(USE_COOPERATIVE_GROUPS "Use cooperative groups for internal syncronization" OFF)
 option(BUILD_FUNCTIONAL_TESTS "Build the functional tests" ON)
 option(BUILD_SOS_TESTS "Build the host-facing tests" OFF)
 option(BUILD_UNIT_TESTS "Build the unit tests" ON)

cmake/config.h.in

@@ -14,4 +14,3 @@
 #cmakedefine USE_FUNC_CALL
 #cmakedefine USE_SINGLE_NODE
 #cmakedefine USE_HOST_SIDE_HDP_FLUSH
-#cmakedefine USE_COOPERATIVE_GROUPS

examples/rocshmem_allreduce_test.cc

@@ -60,7 +60,7 @@ static bool check_recvbuf(int *dest, int nelem, int my_pe, int npes)
         if (dest[i] != result) {
             res = false;
 #ifdef VERBOSE
-            printf("recvbuf[%d] = %d expected %d \n", i, recvbuf[i], result);
+            printf("recvbuf[%d] = %d expected %d \n", i, dest[i], result);
 #endif
         }
     }

examples/rocshmem_alltoall_test.cc

@@ -64,7 +64,7 @@ static bool check_recvbuf(int *dest, int nelem, int my_pe, int npes)
             if (dest[idx] != result) {
                 res = false;
 #ifdef VERBOSE
-                printf("recvbuf[%d] = %d expected %d \n", i, recvbuf[i], result);
+                printf("recvbuf[%d] = %d expected %d \n", i, dest[i], result);
 #endif
             }
         }

scripts/build_configs/ipc_single_cg

@@ -1,31 +0,0 @@
-#!/bin/bash
-# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
-
-if [ -z $1 ]
-then
-  install_path=~/rocshmem
-else
-  install_path=$1
-fi
-
-src_path=$(dirname "$(realpath $0)")/../../
-
-cmake \
-    -DCMAKE_BUILD_TYPE=Release \
-    -DCMAKE_INSTALL_PREFIX=$install_path \
-    -DCMAKE_VERBOSE_MAKEFILE=OFF \
-    -DDEBUG=OFF \
-    -DPROFILE=OFF \
-    -DUSE_GPU_IB=OFF \
-    -DUSE_RO=OFF \
-    -DUSE_DC=OFF \
-    -DUSE_IPC=ON \
-    -DUSE_COHERENT_HEAP=ON \
-    -DUSE_THREADS=OFF \
-    -DUSE_WF_COAL=OFF \
-    -DUSE_SINGLE_NODE=ON \
-    -DUSE_HOST_SIDE_HDP_FLUSH=OFF \
-    -DUSE_COOPERATIVE_GROUPS=ON \
-    $src_path
-cmake --build . --parallel 8
-cmake --install .

src/ipc/backend_ipc.hpp

@@ -32,10 +32,6 @@
 #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 {
@@ -172,10 +168,6 @@ 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()

src/ipc/context_ipc_device.cpp

@@ -51,10 +51,6 @@ __host__ IPCContext::IPCContext(Backend *b)
   Wrk_Sync_buffer_bases_ = backend->get_wrk_sync_bases();
 
   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

@@ -27,13 +27,6 @@
 #include "../atomic.hpp"
 #include "../team.hpp"
 
-#ifdef USE_COOPERATIVE_GROUPS
-#include <hip/hip_cooperative_groups.h>
-namespace cg = cooperative_groups;
-#else
-#include "../memory/notifier.hpp"
-#endif /* USE_COOPERATIVE_GROUPS */
-
 namespace rocshmem {
 
 class IPCContext : public Context {
@@ -273,10 +266,6 @@ class IPCContext : public Context {
   //Temporary scratchpad memory used by internal barrier algorithms.
   int64_t *barrier_sync{nullptr};
 
-#ifndef USE_COOPERATIVE_GROUPS
-  Notifier<detail::atomic::memory_scope_agent> *notifier_{nullptr};
-#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
-
   //Struct defining memory ordering for atomic operations.
   detail::atomic::rocshmem_memory_orders orders_{};
 

src/ipc/context_ipc_device_coll.cpp

@@ -87,27 +87,15 @@ __device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start,
 // Uses PE values that are relative to world
 __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
                                           int PE_size, int64_t *pSync) {
-#ifdef USE_COOPERATIVE_GROUPS
-  cg::grid_group grid = cg::this_grid();
-  grid.sync();
-  if (0 == grid.thread_rank())
-#else
-  notifier_->sync();
-  if (0 == get_flat_id())
-#endif /* USE_COOPERATIVE_GROUPS */
-  {
+  __syncthreads();
+  if (is_thread_zero_in_block()) {
     if (PE_size < 64) {
       internal_direct_barrier(pe, PE_start, stride, PE_size, pSync);
     } else {
       internal_atomic_barrier(pe, PE_start, stride, PE_size, pSync);
     }
   }
-  __threadfence();
-#ifdef USE_COOPERATIVE_GROUPS
-  grid.sync();
-#else
-  notifier_->sync();
-#endif /* USE_COOPERATIVE_GROUPS */
+  __syncthreads();
 }
 
 __device__ void IPCContext::sync(roc_shmem_team_t team) {

tests/functional_tests/sync_tester.cpp

@@ -83,21 +83,9 @@ void SyncTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
   roc_shmem_team_split_strided(ROC_SHMEM_TEAM_WORLD, 0, 1, n_pes, nullptr, 0,
                                &team_sync_world_dup);
 
-#ifdef USE_COOPERATIVE_GROUPS
-  void* kernelParams[] = {(void*)&loop,
-                          (void*)&args.skip,
-                          (void*)&timer,
-                          (void*)&_type,
-                          (void*)&_shmem_context,
-                          (void*)&team_sync_world_dup};
-
-  CHECK_HIP(hipLaunchCooperativeKernel(SyncTest, gridSize, blockSize,
-                                       kernelParams, shared_bytes, stream));
-#else
   hipLaunchKernelGGL(SyncTest, gridSize, blockSize, shared_bytes, stream, loop,
                      args.skip, timer, _type, _shmem_context,
                      team_sync_world_dup);
-#endif /* USE_COOPERATIVE_GROUPS */
 
   num_msgs = (loop + args.skip) * gridSize.x;
   num_timed_msgs = loop;