From fe767d9abf82e4ac957e752872845e6b42fd0626 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Wed, 30 Oct 2024 20:10:21 +0000 Subject: [PATCH] remove cooperative groups --- CMakeLists.txt | 1 - cmake/config.h.in | 1 - scripts/build_configs/ipc_single_cg | 31 -------------------------- src/ipc/backend_ipc.hpp | 8 ------- src/ipc/context_ipc_device.cpp | 4 ---- src/ipc/context_ipc_device.hpp | 11 --------- src/ipc/context_ipc_device_coll.cpp | 18 +++------------ tests/functional_tests/sync_tester.cpp | 12 ---------- 8 files changed, 3 insertions(+), 83 deletions(-) delete mode 100755 scripts/build_configs/ipc_single_cg diff --git a/CMakeLists.txt b/CMakeLists.txt index 5cf74fb6cb..f9c3d83d10 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/cmake/config.h.in b/cmake/config.h.in index 21619a6c3e..8067651a01 100644 --- a/cmake/config.h.in +++ b/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 diff --git a/scripts/build_configs/ipc_single_cg b/scripts/build_configs/ipc_single_cg deleted file mode 100755 index c0904bad93..0000000000 --- a/scripts/build_configs/ipc_single_cg +++ /dev/null @@ -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 . diff --git a/src/ipc/backend_ipc.hpp b/src/ipc/backend_ipc.hpp index 2f84b9c534..f029a32fbb 100644 --- a/src/ipc/backend_ipc.hpp +++ b/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 notifier_{}; -#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */ - protected: /** * @copydoc Backend::dump_backend_stats() diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index ee80694578..83f3f609e6 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/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() { diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index dcb9571833..fe8f9ef323 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -27,13 +27,6 @@ #include "../atomic.hpp" #include "../team.hpp" -#ifdef USE_COOPERATIVE_GROUPS -#include -namespace cg = cooperative_groups; -#else -#include "../memory/notifier.hpp" -#endif /* USE_COOPERATIVE_GROUPS */ - namespace rocshmem { class IPCContext : public Context { @@ -272,10 +265,6 @@ class IPCContext : public Context { //Temporary scratchpad memory used by internal barrier algorithms. int64_t *barrier_sync{nullptr}; -#ifndef USE_COOPERATIVE_GROUPS - Notifier *notifier_{nullptr}; -#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */ - //Struct defining memory ordering for atomic operations. detail::atomic::rocshmem_memory_orders orders_{}; diff --git a/src/ipc/context_ipc_device_coll.cpp b/src/ipc/context_ipc_device_coll.cpp index aa7e7be9ae..e0e06ecf7a 100644 --- a/src/ipc/context_ipc_device_coll.cpp +++ b/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) { diff --git a/tests/functional_tests/sync_tester.cpp b/tests/functional_tests/sync_tester.cpp index 1afecaa45c..54d4a99f99 100644 --- a/tests/functional_tests/sync_tester.cpp +++ b/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;