From 120453c75c890c324f8d90d0678fd9ae7b9dc095 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Tue, 1 Oct 2024 13:27:45 -0700 Subject: [PATCH 1/6] added .gitignore, we do not want to include the build directory in our commits --- .gitignore | 1 + 1 file changed, 1 insertion(+) create mode 100644 .gitignore diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000000..567609b123 --- /dev/null +++ b/.gitignore @@ -0,0 +1 @@ +build/ From e2f6a65284bc335ef9a684e21bbf02828b17433e Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Wed, 2 Oct 2024 09:20:07 -0700 Subject: [PATCH 2/6] fixed barrier issue on MI250X --- src/ipc/context_ipc_device_coll.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/src/ipc/context_ipc_device_coll.cpp b/src/ipc/context_ipc_device_coll.cpp index 804be2efbf..1465b8ffe0 100644 --- a/src/ipc/context_ipc_device_coll.cpp +++ b/src/ipc/context_ipc_device_coll.cpp @@ -35,6 +35,9 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start, if (pe == PE_start) { // Go through all PE offsets (except current offset = 0) // and wait until they all reach +#if defined(__gfx90a__) + __threadfence_system(); +#endif /* __gfx90a__ */ for (size_t i = 1; i < n_pes; i++) { wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val); pSync[i] = ROC_SHMEM_SYNC_VALUE; @@ -49,6 +52,9 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start, // Mark current PE offset as reached size_t pe_offset = (pe - PE_start) / stride; put_nbi(&pSync[pe_offset], &flag_val, 1, PE_start); +#if defined(__gfx90a__) + __threadfence_system(); +#endif /* __gfx90a__ */ wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val); pSync[0] = ROC_SHMEM_SYNC_VALUE; threadfence_system(); From 1baa071edf525572fe4de6c03883ed8fb69047a1 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Wed, 9 Oct 2024 13:10:50 -0700 Subject: [PATCH 3/6] Fix initialization order bug --- src/ipc/backend_ipc.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/src/ipc/backend_ipc.cpp b/src/ipc/backend_ipc.cpp index 70a9295486..7475fb2c8b 100644 --- a/src/ipc/backend_ipc.cpp +++ b/src/ipc/backend_ipc.cpp @@ -56,18 +56,18 @@ int get_ls_non_zero_bit(char *bitmask, int mask_length) { IPCBackend::IPCBackend(MPI_Comm comm) : Backend() { type = BackendType::IPC_BACKEND; - + if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) { std::stringstream sstream(maximum_num_contexts_str); sstream >> maximum_num_contexts_; } init_mpi_once(comm); - + initIPC(); - + auto *bp{ipc_backend_proxy.get()}; - + bp->heap_ptr = &heap; /* Initialize the host interface */ @@ -85,7 +85,6 @@ IPCBackend::IPCBackend(MPI_Comm comm) setup_team_world(); TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world; - default_context_proxy_ = IPCDefaultContextProxyT(this, tinfo); roc_shmem_collective_init(); @@ -93,8 +92,9 @@ IPCBackend::IPCBackend(MPI_Comm comm) teams_init(); + default_context_proxy_ = IPCDefaultContextProxyT(this, tinfo); + setup_ctxs(); - } IPCBackend::~IPCBackend() { @@ -108,7 +108,7 @@ IPCBackend::~IPCBackend() { * Free the atomic_ret array. */ CHECK_HIP(hipFree(bp->atomic_ret->atomic_base_ptr)); - + // TODO(Avinash) Free g_ret // delete host_interface; @@ -402,4 +402,4 @@ void IPCBackend::teams_init() { NET_CHECK(MPI_Barrier(thread_comm)); } -} // namespace rocshmem \ No newline at end of file +} // namespace rocshmem From 63667a3167cd4c1a774331ee8985a708faa77e87 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Tue, 1 Oct 2024 13:26:43 -0700 Subject: [PATCH 4/6] Added Cooperative Groups configure option and header --- CMakeLists.txt | 1 + cmake/config.h.in | 1 + scripts/build_configs/ipc_single_cg | 31 ++++++++++++++++++++++++++ src/ipc/context_ipc_device_coll.cpp | 15 +++++++++++++ tests/functional_tests/sync_tester.cpp | 12 ++++++++++ 5 files changed, 60 insertions(+) create mode 100755 scripts/build_configs/ipc_single_cg diff --git a/CMakeLists.txt b/CMakeLists.txt index f9c3d83d10..5cf74fb6cb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,6 +68,7 @@ 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 8067651a01..21619a6c3e 100644 --- a/cmake/config.h.in +++ b/cmake/config.h.in @@ -14,3 +14,4 @@ #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 new file mode 100755 index 0000000000..c0904bad93 --- /dev/null +++ b/scripts/build_configs/ipc_single_cg @@ -0,0 +1,31 @@ +#!/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/context_ipc_device_coll.cpp b/src/ipc/context_ipc_device_coll.cpp index 1465b8ffe0..5e1f950c3b 100644 --- a/src/ipc/context_ipc_device_coll.cpp +++ b/src/ipc/context_ipc_device_coll.cpp @@ -26,6 +26,11 @@ #include "../util.hpp" #include "ipc_team.hpp" +#ifdef USE_COOPERATIVE_GROUPS +#include +namespace cg = cooperative_groups; +#endif /* USE_COOPERATIVE_GROUPS */ + namespace rocshmem { __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start, @@ -84,8 +89,14 @@ __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 __syncthreads(); if (is_thread_zero_in_block()) { +#endif /* USE_COOPERATIVE_GROUPS */ if (PE_size < 64) { internal_direct_barrier(pe, PE_start, stride, PE_size, pSync); } else { @@ -93,7 +104,11 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride, } } __threadfence(); +#ifdef USE_COOPERATIVE_GROUPS + grid.sync(); +#else __syncthreads(); +#endif /* USE_COOPERATIVE_GROUPS */ } __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 54d4a99f99..1afecaa45c 100644 --- a/tests/functional_tests/sync_tester.cpp +++ b/tests/functional_tests/sync_tester.cpp @@ -83,9 +83,21 @@ 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; From b1134e8633ecb1f9670bad7540b22857f1d87892 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Mon, 7 Oct 2024 06:06:42 -0700 Subject: [PATCH 5/6] added notifier->sync() when we are not using cooperative groups updated scope bug --- src/ipc/context_ipc_device.hpp | 12 ++++++++++++ src/ipc/context_ipc_device_coll.cpp | 15 ++++++--------- 2 files changed, 18 insertions(+), 9 deletions(-) diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index 48fe5acbdf..8046fa2088 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -27,6 +27,13 @@ #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 { @@ -235,6 +242,11 @@ class IPCContext : public Context { //Temporary scratchpad memory used by internal barrier algorithms. int64_t *barrier_sync{nullptr}; +#ifndef USE_COOPERATIVE_GROUPS + using NOTIFIER_PROXY_T = NotifierProxy; + NOTIFIER_PROXY_T notifier_{}; +#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 5e1f950c3b..02bb3ffcec 100644 --- a/src/ipc/context_ipc_device_coll.cpp +++ b/src/ipc/context_ipc_device_coll.cpp @@ -26,11 +26,6 @@ #include "../util.hpp" #include "ipc_team.hpp" -#ifdef USE_COOPERATIVE_GROUPS -#include -namespace cg = cooperative_groups; -#endif /* USE_COOPERATIVE_GROUPS */ - namespace rocshmem { __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start, @@ -92,11 +87,13 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride, #ifdef USE_COOPERATIVE_GROUPS cg::grid_group grid = cg::this_grid(); grid.sync(); - if (0 == grid.thread_rank()) { + if (0 == grid.thread_rank()) #else - __syncthreads(); - if (is_thread_zero_in_block()) { + auto notifier{notifier_.get()}; + notifier->sync(); + if (is_thread_zero_in_block() && is_block_zero_in_grid()) #endif /* USE_COOPERATIVE_GROUPS */ + { if (PE_size < 64) { internal_direct_barrier(pe, PE_start, stride, PE_size, pSync); } else { @@ -107,7 +104,7 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride, #ifdef USE_COOPERATIVE_GROUPS grid.sync(); #else - __syncthreads(); + notifier->sync(); #endif /* USE_COOPERATIVE_GROUPS */ } From 509277c034a5def013b89b35ff4d8ec1638447e1 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Thu, 10 Oct 2024 06:45:43 -0700 Subject: [PATCH 6/6] fixed notifier bug --- src/ipc/backend_ipc.hpp | 12 ++++++++++-- src/ipc/context_ipc_device.cpp | 4 ++++ src/ipc/context_ipc_device.hpp | 3 +-- src/ipc/context_ipc_device_coll.cpp | 7 +++---- 4 files changed, 18 insertions(+), 8 deletions(-) diff --git a/src/ipc/backend_ipc.hpp b/src/ipc/backend_ipc.hpp index 7d00e6cf07..de8cb9c617 100644 --- a/src/ipc/backend_ipc.hpp +++ b/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 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_ \ No newline at end of file +#endif // LIBRARY_SRC_IPC_BACKEND_HPP_ diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index 4d891cbc53..d208a070fb 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/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() { diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index 8046fa2088..065d37cec4 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/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; - NOTIFIER_PROXY_T notifier_{}; + Notifier *notifier_{nullptr}; #endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */ //Struct defining memory ordering for atomic operations. diff --git a/src/ipc/context_ipc_device_coll.cpp b/src/ipc/context_ipc_device_coll.cpp index 02bb3ffcec..f128d76d3f 100644 --- a/src/ipc/context_ipc_device_coll.cpp +++ b/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 */ }