From 801d2c5012792d6f0886bb4349ec88fdcc12314d Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Tue, 16 Sep 2025 11:54:53 -0400 Subject: [PATCH] Enable GDA+IPC (#249) * Enable GDA+IPC Fix ROCSHMEM_DISABLE_IPC for both RO and GDA * add more functionality to bootstrap class we need a few more functions in the boostrap class to be able to fully handle the rocshmem requirements: - add a function to return the list of local ranks - provide a groupAllgather operation which takes a vector of ranks participating - provide a groupAlltoall operation which takes a vector of ranks participating Also, update the functionality of the gda-Alltoall and gda-Allreduce operations to take advantage of these functions. * ipc_policy adapted to use bootstrap groupallgather * bugfix: there was a mistake in computing sendto in groupallgather * bugfix: shm_size and shm_rank were set in a local variable rather than the class member * mpi-bootstrap: remove an unecessary allgather --------- Co-authored-by: Edgar Gabriel --- src/CMakeLists.txt | 6 +- src/backend_bc.cpp | 18 ++-- src/backend_type.hpp | 74 ++++++++-------- src/bootstrap/bootstrap.cpp | 91 +++++++++++++++++++ src/bootstrap/bootstrap.hpp | 6 ++ src/context_incl.hpp | 8 +- src/context_tmpl_device.hpp | 6 +- src/context_tmpl_host.hpp | 6 +- src/gda/backend_gda.cpp | 63 +++++--------- src/gda/backend_gda.hpp | 8 ++ src/gda/context_gda_device.cpp | 111 +++++++++++++++++++++--- src/gda/context_gda_host.cpp | 28 +++++- src/gda/context_gda_tmpl_device.hpp | 13 +++ src/ipc/context_ipc_host.cpp | 5 +- src/ipc_policy.cpp | 31 ++++--- src/reverse_offload/context_ro_host.cpp | 11 ++- src/rocshmem.cpp | 24 ++--- src/rocshmem_gpu.cpp | 6 +- src/util.cpp | 11 ++- src/util.hpp | 4 +- 20 files changed, 378 insertions(+), 152 deletions(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index c8542da45f..cf39bb3a78 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -63,12 +63,12 @@ target_compile_options(${PROJECT_NAME} PUBLIC ${ROCSHMEM_COMPILE_FLAGS}) ############################################################################### # ROCSHMEM TARGET FOR BACKENDS ############################################################################### -if (USE_RO) +if (USE_GDA) +add_subdirectory(gda) +elseif (USE_RO) add_subdirectory(reverse_offload) elseif (USE_IPC) add_subdirectory(ipc) -elseif (USE_GDA) -add_subdirectory(gda) endif() add_subdirectory(host) add_subdirectory(containers) diff --git a/src/backend_bc.cpp b/src/backend_bc.cpp index c974c11428..beefbeebd8 100644 --- a/src/backend_bc.cpp +++ b/src/backend_bc.cpp @@ -27,12 +27,12 @@ #include "backend_type.hpp" #include "context_incl.hpp" -#if defined(USE_RO) +#if defined(USE_GDA) +#include "gda/backend_gda.hpp" +#elif defined(USE_RO) #include "reverse_offload/backend_ro.hpp" #elif defined(USE_IPC) #include "ipc/backend_ipc.hpp" -#elif defined(USE_GDA) -#include "gda/backend_gda.hpp" #endif #include @@ -249,22 +249,22 @@ void Backend::reset_stats() { } __device__ bool Backend::create_ctx(int64_t option, rocshmem_ctx_t* ctx) { -#if defined(USE_RO) +#if defined(USE_GDA) + return static_cast(this)->create_ctx(option, ctx); +#elif defined(USE_RO) return static_cast(this)->create_ctx(option, ctx); #elif defined(USE_IPC) return static_cast(this)->create_ctx(option, ctx); -#elif defined(USE_GDA) - return static_cast(this)->create_ctx(option, ctx); #endif } __device__ void Backend::destroy_ctx(rocshmem_ctx_t* ctx) { -#if defined(USE_RO) +#if defined(USE_GDA) + static_cast(this)->destroy_ctx(ctx); +#elif defined(USE_RO) static_cast(this)->destroy_ctx(ctx); #elif defined(USE_IPC) static_cast(this)->destroy_ctx(ctx); -#elif defined(USE_GDA) - static_cast(this)->destroy_ctx(ctx); #endif } diff --git a/src/backend_type.hpp b/src/backend_type.hpp index ef9c7d3874..203aaeecc3 100644 --- a/src/backend_type.hpp +++ b/src/backend_type.hpp @@ -46,7 +46,7 @@ namespace rocshmem { * @note Derived classes which use Backend as a base class must add * themselves to this enum class to support static polymorphism. */ -enum class BackendType { RO_BACKEND, IPC_BACKEND, GDA_BACKEND }; +enum class BackendType { GDA_BACKEND, RO_BACKEND, IPC_BACKEND }; /** * @brief Helper macro for some dispatch calls @@ -56,52 +56,52 @@ enum class BackendType { RO_BACKEND, IPC_BACKEND, GDA_BACKEND }; /** * @brief Device static dispatch method call. */ -#if defined(USE_RO) +#if defined(USE_GDA) +#define DISPATCH(Func) \ + static_cast(this)->Func; +#elif defined(USE_RO) #define DISPATCH(Func) \ static_cast(this)->Func; #elif defined(USE_IPC) #define DISPATCH(Func) \ static_cast(this)->Func; -#elif defined(USE_GDA) -#define DISPATCH(Func) \ - static_cast(this)->Func; #endif /** * @brief Device static dispatch method call with a return value. */ -#if defined(USE_RO) -#define DISPATCH_RET(Func) \ - auto ret_val = static_cast(this)->Func; \ +#if defined(USE_GDA) +#define DISPATCH_RET(Func) \ + auto ret_val = static_cast(this)->Func; \ + return ret_val; +#elif defined(USE_RO) +#define DISPATCH_RET(Func) \ + auto ret_val = static_cast(this)->Func; \ return ret_val; #elif defined(USE_IPC) #define DISPATCH_RET(Func) \ auto ret_val = static_cast(this)->Func; \ return ret_val; -#elif defined(USE_GDA) -#define DISPATCH_RET(Func) \ - auto ret_val = static_cast(this)->Func; \ - return ret_val; #endif /** * @brief Device static dispatch method call with a return type of pointer. */ -#if defined(USE_RO) -#define DISPATCH_RET_PTR(Func) \ - void *ret_val{nullptr}; \ - ret_val = static_cast(this)->Func; \ +#if defined(USE_GDA) +#define DISPATCH_RET_PTR(Func) \ + void *ret_val{nullptr}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; +#elif defined(USE_RO) +#define DISPATCH_RET_PTR(Func) \ + void *ret_val{nullptr}; \ + ret_val = static_cast(this)->Func; \ return ret_val; #elif defined(USE_IPC) #define DISPATCH_RET_PTR(Func) \ void *ret_val{nullptr}; \ ret_val = static_cast(this)->Func; \ return ret_val; -#elif defined(USE_GDA) -#define DISPATCH_RET_PTR(Func) \ - void *ret_val{nullptr}; \ - ret_val = static_cast(this)->Func; \ - return ret_val; #endif /** @@ -111,12 +111,12 @@ enum class BackendType { RO_BACKEND, IPC_BACKEND, GDA_BACKEND }; * MPI_THREAD_MULTIPLE (for RMA and AMO operations) and the ordering and * threading semantics of collectives in OpenSHMEM match those of MPI. */ -#if defined(USE_RO) +#if defined(USE_GDA) +#define HOST_DISPATCH(Func) static_cast(this)->Func; +#elif defined(USE_RO) #define HOST_DISPATCH(Func) static_cast(this)->Func; #elif defined(USE_IPC) #define HOST_DISPATCH(Func) static_cast(this)->Func; -#elif defined(USE_GDA) -#define HOST_DISPATCH(Func) static_cast(this)->Func; #endif /** @@ -126,24 +126,29 @@ enum class BackendType { RO_BACKEND, IPC_BACKEND, GDA_BACKEND }; * MPI_THREAD_MULTIPLE (for RMA and AMO operations) and the ordering and * threading semantics of collectives in OpenSHMEM match those of MPI. */ -#if defined(USE_RO) -#define HOST_DISPATCH_RET(Func) \ - auto ret_val = static_cast(this)->Func; \ +#if defined(USE_GDA) +#define HOST_DISPATCH_RET(Func) \ + auto ret_val = static_cast(this)->Func; \ + return ret_val; +#elif defined(USE_RO) +#define HOST_DISPATCH_RET(Func) \ + auto ret_val = static_cast(this)->Func; \ return ret_val; #elif defined(USE_IPC) #define HOST_DISPATCH_RET(Func) \ auto ret_val = static_cast(this)->Func; \ return ret_val; -#elif defined(USE_GDA) -#define HOST_DISPATCH_RET(Func) \ - auto ret_val = static_cast(this)->Func; \ - return ret_val; #endif /** * @brief Host static dispatch method call with a return type of pointer. */ -#if defined(USE_RO) +#if defined(USE_GDA) +#define HOST_DISPATCH_RET_PTR(Func) \ + void *ret_val{nullptr}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; +#elif defined(USE_RO) #define HOST_DISPATCH_RET_PTR(Func) \ void *ret_val{nullptr}; \ ret_val = static_cast(this)->Func; \ @@ -153,11 +158,6 @@ enum class BackendType { RO_BACKEND, IPC_BACKEND, GDA_BACKEND }; void *ret_val{nullptr}; \ ret_val = static_cast(this)->Func; \ return ret_val; -#elif defined(USE_GDA) -#define HOST_DISPATCH_RET_PTR(Func) \ - void *ret_val{nullptr}; \ - ret_val = static_cast(this)->Func; \ - return ret_val; #endif } // namespace rocshmem diff --git a/src/bootstrap/bootstrap.cpp b/src/bootstrap/bootstrap.cpp index ff107695a4..f20d444413 100644 --- a/src/bootstrap/bootstrap.cpp +++ b/src/bootstrap/bootstrap.cpp @@ -74,6 +74,89 @@ struct ExtInfo { } } + void Bootstrap::groupAllGather(void* allData, int size, const std::vector& ranks) { + char* data = static_cast(allData); + int rank = this->getRank(); + int nRanks = ranks.size(); + int rank_pos = -1; + + // Confirm that rank is in the vectors of ranks + for (int i = 0; i < ranks.size(); i++) { + if (rank == ranks[i]) { + rank_pos = i; + break; + } + } + + if (rank_pos == -1) { + printf("Bootstrap::groupAllGather: called with process that is not in list of ranks. Aborting\n"); + abort(); + } + + DPRINTF("groupAllGather: rank %d nranks %d size %d\n", rank, nRanks, size); + + int sendto = (rank_pos + 1 + nRanks) % nRanks; + int recvfrom = (rank_pos - 1 + nRanks) % nRanks; + for (int i = 0; i < nRanks - 1; i++) { + size_t rSlice = (rank_pos - i - 1 + nRanks) % nRanks; + size_t sSlice = (rank_pos - i + nRanks) % nRanks; + + char *tmpsend = data + sSlice * size; + char *tmprecv = data + rSlice * size; + this->send(tmpsend, size, ranks[sendto], i); + this->recv(tmprecv, size, ranks[recvfrom], i); + } + + DPRINTF("groupAllGather: rank %d nranks %d size %d - DONE\n", rank, nRanks, size); + } + + void Bootstrap::groupAlltoall(void* allData, int size, const std::vector& ranks) { + char* data = static_cast(allData); + int num_pes = ranks.size(); + int rank = this->getRank(); + int rank_pos = -1; + + // Confirm that rank is in the vectors of ranks + for (int i = 0; i < ranks.size(); i++) { + if (rank == ranks[i]) { + rank_pos = i; + break; + } + } + + if (rank_pos == -1) { + printf("Bootstrap::groupAlltoall: called with process that is not in list of ranks. Aborting\n"); + abort(); + } + + DPRINTF("groupAlltoall: rank %d nranks %d size %d\n", rank, num_pes, size); + + // Since this is an in-place algorithm, allocate temporary receive buffer + char *recv_buf = new char[size * num_pes]; + std::memset(recv_buf, 0, num_pes * size); + + // Perform pairwise exchange - local copy is ommitted + for (int step = 1; step < num_pes; step++) { + int sendto = (rank_pos + step) % num_pes; + int recvfrom = (rank_pos + num_pes - step) % num_pes; + + char *tmpsend = (char*)data + (ptrdiff_t)sendto * size; + char *tmprecv = (char*)recv_buf + (ptrdiff_t)recvfrom * size; + + this->send(tmpsend, size, ranks[sendto], step /* used as tag */); + this->recv(tmprecv, size, ranks[recvfrom], step); + } + + //Since this is an in_place all-to-all, copy data back into the user buffer + for (int step = 0; step < num_pes; step++) { + if (step == rank_pos) continue; + std::memcpy(&data[step*size], &recv_buf[step*size], size); + } + + DPRINTF("groupAlltoall: rank %d nranks %d size %d DONE \n", rank, num_pes, size); + delete[] recv_buf; + } + void Bootstrap::send(const std::vector& data, int peer, int tag) { size_t size = data.size(); send((void*)&size, sizeof(size_t), peer, tag); @@ -107,6 +190,7 @@ class TcpBootstrap::Impl { int getRank(); int getNranks(); int getNranksPerNode(); + std::vector getLocalRanks(); void allGather(void* allData, int size); void send(void* data, int size, int peer, int tag); void recv(void* data, int size, int peer, int tag); @@ -131,6 +215,7 @@ class TcpBootstrap::Impl { SocketAddress netIfAddr_; std::unordered_map, std::shared_ptr, PairHash> peerSendSockets_; std::unordered_map, std::shared_ptr, PairHash> peerRecvSockets_; + std::vector localRanks_; void netSend(Socket* sock, const void* data, int size); void netRecv(Socket* sock, void* data, int size); @@ -181,6 +266,8 @@ int TcpBootstrap::Impl::getRank() { return rank_; } int TcpBootstrap::Impl::getNranks() { return nRanks_; } +std::vector TcpBootstrap::Impl::getLocalRanks() { return localRanks_; } + void TcpBootstrap::Impl::initialize(const rocshmem_uniqueid_t& uniqueId, int64_t timeoutSec) { if (!netInitialized) { netInit("", "", netIfAddr_); @@ -467,12 +554,14 @@ int TcpBootstrap::Impl::getNranksPerNode() { if (useIpv4) { if (peerCommAddresses_[i].sin.sin_addr.s_addr == peerCommAddresses_[rank_].sin.sin_addr.s_addr) { + localRanks_.push_back(i); nRanksPerNode++; } } else { if (std::memcmp(&(peerCommAddresses_[i].sin6.sin6_addr), &(peerCommAddresses_[rank_].sin6.sin6_addr), sizeof(in6_addr)) == 0) { + localRanks_.push_back(i); nRanksPerNode++; } } @@ -586,6 +675,8 @@ void TcpBootstrap::Impl::close() { int TcpBootstrap::getNranksPerNode() { return pimpl_->getNranksPerNode(); } + std::vector TcpBootstrap::getLocalRanks() { return pimpl_->getLocalRanks(); } + void TcpBootstrap::send(void* data, int size, int peer, int tag) { pimpl_->send(data, size, peer, tag); } diff --git a/src/bootstrap/bootstrap.hpp b/src/bootstrap/bootstrap.hpp index 16e4e9db77..953c5a71c0 100644 --- a/src/bootstrap/bootstrap.hpp +++ b/src/bootstrap/bootstrap.hpp @@ -49,12 +49,15 @@ class Bootstrap { virtual int getRank() = 0; virtual int getNranks() = 0; virtual int getNranksPerNode() = 0; + virtual std::vector getLocalRanks() = 0; virtual void send(void* data, int size, int peer, int tag) = 0; virtual void recv(void* data, int size, int peer, int tag) = 0; virtual void allGather(void* allData, int size) = 0; virtual void barrier() = 0; void groupBarrier(const std::vector& ranks); + void groupAllGather(void* allData, int size, const std::vector& ranks); + void groupAlltoall(void* allData, int size, const std::vector& ranks); void send(const std::vector& data, int peer, int tag); void recv(std::vector& data, int peer, int tag); }; @@ -119,6 +122,9 @@ class TcpBootstrap : public Bootstrap { /// @param tag The tag to receive the data with. void recv(void* data, int size, int peer, int tag) override; + /// Provide list of ranks that are local to the calling process + std::vector getLocalRanks() override; + /// Gather data from all processes. /// /// When called by rank `r`, this sends data from `allData[r * size]` to `allData[(r + 1) * size - 1]` to all other diff --git a/src/context_incl.hpp b/src/context_incl.hpp index 5f8106ebed..efc0fa65f2 100644 --- a/src/context_incl.hpp +++ b/src/context_incl.hpp @@ -28,15 +28,15 @@ #include "context.hpp" #include "context_tmpl_device.hpp" #include "context_tmpl_host.hpp" -#if defined(USE_RO) +#if defined(USE_GDA) +#include "gda/context_gda_device.hpp" +#include "gda/context_gda_host.hpp" +#elif defined(USE_RO) #include "reverse_offload/context_ro_device.hpp" #include "reverse_offload/context_ro_host.hpp" #elif defined(USE_IPC) #include "ipc/context_ipc_device.hpp" #include "ipc/context_ipc_host.hpp" -#elif defined(USE_GDA) -#include "gda/context_gda_device.hpp" -#include "gda/context_gda_host.hpp" #else #error "Select one backend among USE_RO, USE_IPC, USE_GDA" #endif diff --git a/src/context_tmpl_device.hpp b/src/context_tmpl_device.hpp index 3c0b5802ac..ecaaf6c909 100644 --- a/src/context_tmpl_device.hpp +++ b/src/context_tmpl_device.hpp @@ -27,12 +27,12 @@ #include "rocshmem/rocshmem_config.h" // NOLINT(build/include_subdir) #include "backend_type.hpp" -#if defined(USE_RO) +#if defined(USE_GDA) +#include "gda/context_gda_device.hpp" +#elif defined(USE_RO) #include "reverse_offload/context_ro_device.hpp" #elif defined(USE_IPC) #include "ipc/context_ipc_device.hpp" -#elif defined(USE_GDA) -#include "gda/context_gda_device.hpp" #endif namespace rocshmem { diff --git a/src/context_tmpl_host.hpp b/src/context_tmpl_host.hpp index 53236c540c..36ba8a552b 100644 --- a/src/context_tmpl_host.hpp +++ b/src/context_tmpl_host.hpp @@ -27,12 +27,12 @@ #include "rocshmem/rocshmem_config.h" // NOLINT(build/include_subdir) #include "backend_type.hpp" -#if defined(USE_RO) +#if defined(USE_GDA) +#include "gda/context_gda_host.hpp" +#elif defined(USE_RO) #include "reverse_offload/context_ro_host.hpp" #elif defined(USE_IPC) #include "ipc/context_ipc_host.hpp" -#elif defined(USE_GDA) -#include "gda/context_gda_host.hpp" #endif namespace rocshmem { diff --git a/src/gda/backend_gda.cpp b/src/gda/backend_gda.cpp index 6d2aee55a5..a467925b19 100644 --- a/src/gda/backend_gda.cpp +++ b/src/gda/backend_gda.cpp @@ -115,6 +115,8 @@ void GDABackend::init() { setup_team_world(); rte_barrier(); + setup_ipc(); + setup_ibv(); setup_heap_memory_rkey(); setup_gpu_qps(); @@ -133,6 +135,8 @@ GDABackend::~GDABackend() { cleanup_wrk_sync_buffer(); + cleanup_ipc(); + cleanup_gpu_qps(); cleanup_heap_memory_rkey(); cleanup_ibv(); @@ -166,6 +170,18 @@ void GDABackend::read_env() { } } +void GDABackend::setup_ipc() { + const auto &heap_bases{heap.get_heap_bases()}; + + if (MPI_COMM_NULL != backend_comm) + ipcImpl.ipcHostInit(my_pe, heap_bases, backend_comm); + else + ipcImpl.ipcHostInit(my_pe, heap_bases, backend_bootstr); +} + +void GDABackend::cleanup_ipc() { + ipcImpl.ipcHostStop(); +} void GDABackend::setup_host_ctx() { default_host_ctx = std::make_unique(this, 0); @@ -258,41 +274,12 @@ void GDABackend::team_destroy(rocshmem_team_t team) { void GDABackend::Alltoall_char_inplace (char *inoutbuf, size_t num_bytes, rocshmem_team_t team) { // Implement an Alltoall outside of MPI assuming in_place communication GDATeam *team_obj = reinterpret_cast(team); - int num_pes = team_obj->num_pes; - int my_pe = team_obj->my_pe; - int *pes_in_world = new int[num_pes]; + std::vector pes_in_world; - int my_pe_in_world = team_obj->my_pe_in_world; for (int i = 0; i < num_pes; i++) { - pes_in_world[i] = team_obj->get_pe_in_world(i); + pes_in_world.push_back(team_obj->get_pe_in_world(i)); } - - // Since this is an in-place algorithm, allocate the temporary receive buffer first - char *recv_buf = new char[num_bytes * num_pes]; - std::memset(recv_buf, 0, num_pes * num_bytes); - - // Perform pairwise exchange - local copy is ommitted - for (int step = 1; step < num_pes; step++) { - int sendto_team = (my_pe + step) % num_pes; - int recvfrom_team = (my_pe + num_pes - step) % num_pes; - - char *tmpsend = (char*)inoutbuf + (ptrdiff_t)sendto_team * num_bytes; - char *tmprecv = (char*)recv_buf + (ptrdiff_t)recvfrom_team * num_bytes; - - // similarly to the allGather in the bootstrap code, we do send first - // followed by the receive. - // There is a chance for deadlock in my opinion for large messages. - backend_bootstr->send(tmpsend, num_bytes, pes_in_world[sendto_team], step /* used as tag */); - backend_bootstr->recv(tmprecv, num_bytes, pes_in_world[recvfrom_team], step); - } - //Since this is an in_place all-to-all, copy data back into the user buffer - for (int step = 0; step < num_pes; step++) { - if (step == my_pe) continue; - std::memcpy(&inoutbuf[step*num_bytes], &recv_buf[step*num_bytes], num_bytes); - } - - delete[] recv_buf; - delete[] pes_in_world; + backend_bootstr->groupAlltoall(inoutbuf, num_bytes, pes_in_world); } //TODO: factorize somewhere else, maybe backend_bc? @@ -305,18 +292,16 @@ void GDABackend::Allreduce_char_BAND (char* inbuf, char *outbuf, size_t num_byte GDATeam *team_obj = reinterpret_cast(team); int num_pes = team_obj->num_pes; - int my_pe = team_obj->my_pe; + std::vector pes_in_world; char *tmp_buffer = new char[num_pes * num_bytes]; std::memset(tmp_buffer, 0, num_pes * num_bytes); - std::memcpy (&tmp_buffer[my_pe * num_bytes], inbuf, num_bytes); + std::memcpy(&tmp_buffer[my_pe * num_bytes], inbuf, num_bytes); - if (num_pes == backend_bootstr->getNranks() ) { - backend_bootstr->allGather(tmp_buffer, num_bytes); - } else { - printf("GDABackend::create_new_team: non-mpi version only supports parent_teams that contain all processes. Aborting.\n"); - abort(); + for (int i = 0; i < num_pes; i++) { + pes_in_world.push_back(team_obj->get_pe_in_world(i)); } + backend_bootstr->groupAllGather(tmp_buffer, num_bytes, pes_in_world); for (int i = 0; i < num_bytes; i++) { outbuf[i] = tmp_buffer[i]; diff --git a/src/gda/backend_gda.hpp b/src/gda/backend_gda.hpp index f94f7bbe3b..56313a0ac8 100644 --- a/src/gda/backend_gda.hpp +++ b/src/gda/backend_gda.hpp @@ -301,6 +301,14 @@ class GDABackend : public Backend { void setup_host_ctx(); void setup_default_ctx(); + + /** + * @brief Allocation and initialization of resources required to + * support IPC handover. + */ + void setup_ipc(); + void cleanup_ipc(); + /** * @brief Allocate and initialize barrier operation addresses on * symmetric heap. diff --git a/src/gda/context_gda_device.cpp b/src/gda/context_gda_device.cpp index 5ba384b4fe..6bb8dc4032 100644 --- a/src/gda/context_gda_device.cpp +++ b/src/gda/context_gda_device.cpp @@ -49,6 +49,12 @@ __host__ GDAContext::GDAContext(Backend *b, unsigned int ctx_id) CHECK_HIP(hipMemcpy(&qps[i], &backend->gpu_qps[offset], sizeof(QueuePair), hipMemcpyDefault)); qps[i].base_heap = base_heap; } + + ipcImpl_.ipc_bases = backend->ipcImpl.ipc_bases; + ipcImpl_.shm_size = backend->ipcImpl.shm_size; + ipcImpl_.shm_rank = backend->ipcImpl.shm_rank; + ipcImpl_.pes_with_ipc_avail = backend->ipcImpl.pes_with_ipc_avail; + ctx_id_ = ctx_id; } @@ -63,7 +69,13 @@ __device__ void GDAContext::ctx_destroy(){ } __device__ void GDAContext::putmem(void *dest, const void *source, size_t nelems, - int pe) { + int pe) { + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); + return; + } uint64_t L_offset = reinterpret_cast(dest) - base_heap[my_pe]; bool need_turn {true}; uint64_t turns = __ballot(need_turn); @@ -80,8 +92,14 @@ __device__ void GDAContext::putmem(void *dest, const void *source, size_t nelems } __device__ void GDAContext::getmem(void *dest, const void *source, size_t nelems, - int pe) { + int pe) { const char *src_typed = reinterpret_cast(source); + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + return; + } uint64_t L_offset = const_cast(src_typed) - base_heap[my_pe]; bool need_turn {true}; uint64_t turns = __ballot(need_turn); @@ -98,7 +116,13 @@ __device__ void GDAContext::getmem(void *dest, const void *source, size_t nelems } __device__ void GDAContext::putmem_nbi(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); + return; + } uint64_t L_offset = reinterpret_cast(dest) - base_heap[my_pe]; bool need_turn {true}; uint64_t turns = __ballot(need_turn); @@ -114,8 +138,14 @@ __device__ void GDAContext::putmem_nbi(void *dest, const void *source, } __device__ void GDAContext::getmem_nbi(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { const char *src_typed = reinterpret_cast(source); + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + return; + } uint64_t L_offset = const_cast(src_typed) - base_heap[my_pe]; bool need_turn {true}; uint64_t turns = __ballot(need_turn); @@ -148,11 +178,24 @@ __device__ void GDAContext::quiet() { } __device__ void *GDAContext::shmem_ptr(const void *dest, int pe) { - return nullptr; + void *ret = nullptr; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + void *dst = const_cast(dest); + uint64_t L_offset = reinterpret_cast(dst) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ret = ipcImpl_.ipc_bases[local_pe] + L_offset; + } + return ret; } __device__ void GDAContext::putmem_wg(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); + return; + } uint64_t L_offset = reinterpret_cast(dest) - base_heap[my_pe]; if (is_thread_zero_in_block()) { qps[pe].put_nbi(base_heap[pe] + L_offset, source, nelems, pe); @@ -161,8 +204,14 @@ __device__ void GDAContext::putmem_wg(void *dest, const void *source, } __device__ void GDAContext::getmem_wg(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { const char *src_typed = reinterpret_cast(source); + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + return; + } uint64_t L_offset = const_cast(src_typed) - base_heap[my_pe]; if (is_thread_zero_in_block()) { qps[pe].get_nbi(dest, base_heap[pe] + L_offset, nelems, pe); @@ -171,7 +220,13 @@ __device__ void GDAContext::getmem_wg(void *dest, const void *source, } __device__ void GDAContext::putmem_nbi_wg(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); + return; + } uint64_t L_offset = reinterpret_cast(dest) - base_heap[my_pe]; if (is_thread_zero_in_block()) { qps[pe].put_nbi(base_heap[pe] + L_offset, source, nelems, pe); @@ -179,8 +234,14 @@ __device__ void GDAContext::putmem_nbi_wg(void *dest, const void *source, } __device__ void GDAContext::getmem_nbi_wg(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { const char *src_typed = reinterpret_cast(source); + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + return; + } uint64_t L_offset = const_cast(src_typed) - base_heap[my_pe]; if (is_thread_zero_in_block()) { qps[pe].get_nbi(dest, base_heap[pe] + L_offset, nelems, pe); @@ -188,7 +249,13 @@ __device__ void GDAContext::getmem_nbi_wg(void *dest, const void *source, } __device__ void GDAContext::putmem_wave(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); + return; + } uint64_t L_offset = reinterpret_cast(dest) - base_heap[my_pe]; if (is_thread_zero_in_wave()) { qps[pe].put_nbi(base_heap[pe] + L_offset, source, nelems, pe); @@ -197,8 +264,14 @@ __device__ void GDAContext::putmem_wave(void *dest, const void *source, } __device__ void GDAContext::getmem_wave(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { const char *src_typed = reinterpret_cast(source); + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + return; + } uint64_t L_offset = const_cast(src_typed) - base_heap[my_pe]; if (is_thread_zero_in_wave()) { qps[pe].get_nbi(dest, base_heap[pe] + L_offset, nelems, pe); @@ -207,7 +280,13 @@ __device__ void GDAContext::getmem_wave(void *dest, const void *source, } __device__ void GDAContext::putmem_nbi_wave(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); + return; + } uint64_t L_offset = reinterpret_cast(dest) - base_heap[my_pe]; if (is_thread_zero_in_wave()) { qps[pe].put_nbi(base_heap[pe] + L_offset, source, nelems, pe); @@ -215,8 +294,14 @@ __device__ void GDAContext::putmem_nbi_wave(void *dest, const void *source, } __device__ void GDAContext::getmem_nbi_wave(void *dest, const void *source, - size_t nelems, int pe) { + size_t nelems, int pe) { const char *src_typed = reinterpret_cast(source); + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + return; + } uint64_t L_offset = const_cast(src_typed) - base_heap[my_pe]; if (is_thread_zero_in_wave()) { qps[pe].get_nbi(dest, base_heap[pe] + L_offset, nelems, pe); diff --git a/src/gda/context_gda_host.cpp b/src/gda/context_gda_host.cpp index 5190e8637f..ddc2536d11 100644 --- a/src/gda/context_gda_host.cpp +++ b/src/gda/context_gda_host.cpp @@ -42,9 +42,29 @@ __host__ GDAHostContext::GDAHostContext(Backend *backend, host_interface = b->host_interface; context_window_info = host_interface->acquire_window_context(); + + int *pes_with_ipc_avail = new int[backend->ipcImpl.shm_size]; + char** ipc_bases = new char*[b->ipcImpl.shm_size]; + if (backend->ipcImpl.pes_with_ipc_avail != nullptr) { + CHECK_HIP(hipMemcpy(pes_with_ipc_avail, + backend->ipcImpl.pes_with_ipc_avail, + backend->ipcImpl.shm_size * sizeof(int), + hipMemcpyDeviceToHost)); + CHECK_HIP(hipMemcpy(ipc_bases, + backend->ipcImpl.ipc_bases, + backend->ipcImpl.shm_size * sizeof(char *), + hipMemcpyDeviceToHost)); + } + ipcImpl_.pes_with_ipc_avail = pes_with_ipc_avail; + ipcImpl_.ipc_bases = ipc_bases; + ipcImpl_.shm_size = backend->ipcImpl.shm_size; + ipcImpl_.shm_rank = backend->ipcImpl.shm_rank; } __host__ GDAHostContext::~GDAHostContext() { + delete[] ipcImpl_.pes_with_ipc_avail; + delete[] ipcImpl_.ipc_bases; + host_interface->release_window_context(context_window_info); } @@ -78,8 +98,12 @@ __host__ void GDAHostContext::quiet() { __host__ void *GDAHostContext::shmem_ptr(const void *dest, int pe) { void *ret = nullptr; - //not implemented, returning nullptr is spec-valid - //TODO: copy ipc handover from RO when IPC+GDA is implemented + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + void *dst = const_cast(dest); + uint64_t L_offset = reinterpret_cast(dst) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ret = ipcImpl_.ipc_bases[local_pe] + L_offset; + } return ret; } diff --git a/src/gda/context_gda_tmpl_device.hpp b/src/gda/context_gda_tmpl_device.hpp index e28c5e0aa8..12d6d2274d 100644 --- a/src/gda/context_gda_tmpl_device.hpp +++ b/src/gda/context_gda_tmpl_device.hpp @@ -42,6 +42,12 @@ namespace rocshmem { *****************************************************************************/ template __device__ void GDAContext::p(T *dest, T value, int pe) { + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + long L_offset{reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]}; + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, reinterpret_cast(&value), sizeof(T)); + return; + } putmem_nbi(dest, &value, sizeof(T), pe); } @@ -58,6 +64,13 @@ __device__ void GDAContext::put_nbi(T *dest, const T *source, size_t nelems, int template __device__ T GDAContext::g(const T *source, int pe) { T ret; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { + const char *src_typed{reinterpret_cast(source)}; + long L_offset{const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]}; + ipcImpl_.ipcCopy(&ret, ipcImpl_.ipc_bases[local_pe] + L_offset, sizeof(T)); + return ret; + } printf("rocshmem::gda:g not implemented\n"); abort(); //TODO the following is incorrect because ret is not ibv registered memory diff --git a/src/ipc/context_ipc_host.cpp b/src/ipc/context_ipc_host.cpp index e6e630355d..b549f8c426 100644 --- a/src/ipc/context_ipc_host.cpp +++ b/src/ipc/context_ipc_host.cpp @@ -90,9 +90,8 @@ __host__ void IPCHostContext::quiet() { __host__ void *IPCHostContext::shmem_ptr(const void *dest, int pe) { void *ret = nullptr; void *dst = const_cast(dest); - uint64_t L_offset = - reinterpret_cast(dst) - ipcImpl_.ipc_bases[my_pe]; - ret = ipcImpl_.ipc_bases[pe] + L_offset; + uint64_t L_offset = reinterpret_cast(dst) - ipcImpl_.ipc_bases[my_pe]; + ret = ipcImpl_.ipc_bases[pe] + L_offset; return ret; } diff --git a/src/ipc_policy.cpp b/src/ipc_policy.cpp index 7a6d0639e6..cfc169b162 100644 --- a/src/ipc_policy.cpp +++ b/src/ipc_policy.cpp @@ -109,25 +109,29 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, */ free(vec_ipc_handle); - if (0 == rocshmem_env_.get_ro_disable_ipc()) { + if (0 == rocshmem_env_.get_disable_ipc()) { int thread_comm_rank {-1}; CHECK_HIP(hipMalloc(reinterpret_cast(&pes_with_ipc_avail), shm_size * sizeof(int))); - MPI_Comm_rank(thread_comm, &thread_comm_rank); - MPI_Allgather(&thread_comm_rank, 1, MPI_INT, pes_with_ipc_avail, 1, MPI_INT, shmcomm); + MPI_Group thread_grp, shm_grp; + MPI_Comm_group(thread_comm, &thread_grp); + MPI_Comm_group(shmcomm, &shm_grp); + int *seqranks = new int[shm_size]; + for(int i = 0; i < shm_size; i++) + seqranks[i] = i; + MPI_Group_translate_ranks(shm_grp, shm_size, seqranks, thread_grp, pes_with_ipc_avail); + delete [] seqranks; + MPI_Group_free(&shm_grp); + MPI_Group_free(&thread_grp); } } __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, TcpBootstrap *bootstr) { - /* - * The non-MPI based version only works for ipc conduit for now, - * i.e. total number of ranks and number of local ranks have to match. - */ shm_size = bootstr->getNranksPerNode(); - assert (shm_size == bootstr->getNranks()); - shm_rank = my_pe; + auto shm_ranks = bootstr->getLocalRanks(); + shm_rank = std::find(shm_ranks.begin(), shm_ranks.end(), my_pe) - shm_ranks.begin(); /* * Allocate a host-side c-array to hold the IPC handles. @@ -148,7 +152,7 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, * Do an all-to-all exchange with each local processing element to * share the symmetric heap IPC handles. */ - bootstr->allGather(vec_ipc_handle, sizeof(hipIpcMemHandle_t)); + bootstr->groupAllGather(vec_ipc_handle, sizeof(hipIpcMemHandle_t), shm_ranks); /* * Allocate device-side array to hold the IPC symmetric heap base @@ -182,6 +186,13 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, * addresses. */ free(vec_ipc_handle); + + if (0 == rocshmem_env_.get_disable_ipc()) { + int thread_comm_rank {-1}; + + CHECK_HIP(hipMalloc(reinterpret_cast(&pes_with_ipc_avail), shm_size * sizeof(int))); + std::copy(shm_ranks.begin(), shm_ranks.end(), pes_with_ipc_avail); + } } __host__ void IpcOnImpl::ipcHostStop() { diff --git a/src/reverse_offload/context_ro_host.cpp b/src/reverse_offload/context_ro_host.cpp index 8692b8dfe5..27671ac76d 100644 --- a/src/reverse_offload/context_ro_host.cpp +++ b/src/reverse_offload/context_ro_host.cpp @@ -44,17 +44,16 @@ __host__ ROHostContext::ROHostContext(Backend *backend, long options) int *pes_with_ipc_avail = new int[backend->ipcImpl.shm_size]; char** ipc_bases = new char*[b->ipcImpl.shm_size]; - - CHECK_HIP(hipMemcpy(pes_with_ipc_avail, + if (backend->ipcImpl.pes_with_ipc_avail != nullptr) { + CHECK_HIP(hipMemcpy(pes_with_ipc_avail, backend->ipcImpl.pes_with_ipc_avail, backend->ipcImpl.shm_size * sizeof(int), hipMemcpyDeviceToHost)); - - CHECK_HIP(hipMemcpy(ipc_bases, + CHECK_HIP(hipMemcpy(ipc_bases, backend->ipcImpl.ipc_bases, backend->ipcImpl.shm_size * sizeof(char *), hipMemcpyDeviceToHost)); - + } ipcImpl_.pes_with_ipc_avail = pes_with_ipc_avail; ipcImpl_.ipc_bases = ipc_bases; ipcImpl_.shm_size = backend->ipcImpl.shm_size; @@ -62,9 +61,9 @@ __host__ ROHostContext::ROHostContext(Backend *backend, long options) } __host__ ROHostContext::~ROHostContext() { - // host_interface->release_window_context(context_window_info); delete[] ipcImpl_.pes_with_ipc_avail; delete[] ipcImpl_.ipc_bases; + // host_interface->release_window_context(context_window_info); } __host__ void ROHostContext::putmem_nbi(void *dest, const void *source, diff --git a/src/rocshmem.cpp b/src/rocshmem.cpp index d1a8e38b2d..eee3bd3c82 100644 --- a/src/rocshmem.cpp +++ b/src/rocshmem.cpp @@ -35,15 +35,15 @@ #include "backend_bc.hpp" #include "context_incl.hpp" -#if defined(USE_RO) +#if defined(USE_GDA) +#include "gda/backend_gda.hpp" +#include "gda/context_gda_tmpl_host.hpp" +#elif defined(USE_RO) #include "reverse_offload/backend_ro.hpp" #include "reverse_offload/context_ro_tmpl_host.hpp" #elif defined(USE_IPC) #include "ipc/backend_ipc.hpp" #include "ipc/context_ipc_tmpl_host.hpp" -#elif defined(USE_GDA) -#include "gda/backend_gda.hpp" -#include "gda/context_gda_tmpl_host.hpp" #else #error "Select one backend among USE_RO, USE_IPC, USE_GDA" #endif @@ -94,15 +94,15 @@ rocshmem_ctx_t ROCSHMEM_HOST_CTX_DEFAULT; mpi_instance = new MPIInstance(comm); -#if defined(USE_RO) +#if defined(USE_GDA) + CHECK_HIP(hipHostMalloc(&backend, sizeof(GDABackend))); + backend = new (backend) GDABackend(comm); +#elif defined(USE_RO) CHECK_HIP(hipHostMalloc(&backend, sizeof(ROBackend))); backend = new (backend) ROBackend(comm); #elif defined(USE_IPC) CHECK_HIP(hipHostMalloc(&backend, sizeof(IPCBackend))); backend = new (backend) IPCBackend(comm); -#elif defined(USE_GDA) - CHECK_HIP(hipHostMalloc(&backend, sizeof(GDABackend))); - backend = new (backend) GDABackend(comm); #endif if (!backend) { @@ -174,15 +174,15 @@ rocshmem_ctx_t ROCSHMEM_HOST_CTX_DEFAULT; rocm_init(); -#if defined(USE_RO) +#if defined(USE_GDA) + CHECK_HIP(hipHostMalloc(&backend, sizeof(GDABackend))); + backend = new (backend) GDABackend(bootstrap); +#elif defined(USE_RO) printf("RO Backend requires MPI library to be initialized, even when using uniqueId initializations!\n"); abort(); #elif defined(USE_IPC) CHECK_HIP(hipHostMalloc(&backend, sizeof(IPCBackend))); backend = new (backend) IPCBackend(bootstrap); -#elif defined(USE_GDA) - CHECK_HIP(hipHostMalloc(&backend, sizeof(GDABackend))); - backend = new (backend) GDABackend(bootstrap); #endif if (!backend) { diff --git a/src/rocshmem_gpu.cpp b/src/rocshmem_gpu.cpp index 8abe25e7ec..48623ae372 100644 --- a/src/rocshmem_gpu.cpp +++ b/src/rocshmem_gpu.cpp @@ -51,15 +51,15 @@ #include "templates.hpp" #include "util.hpp" -#if defined(USE_RO) +#if defined(USE_GDA) +#include "gda/context_gda_tmpl_device.hpp" +#elif defined(USE_RO) #include "reverse_offload/context_ro_tmpl_device.hpp" #elif defined(USE_IPC) # if defined(ENABLE_IPC_BITCODE) # include "ipc/backend_ipc.hpp" # endif #include "ipc/context_ipc_tmpl_device.hpp" -#elif defined(USE_GDA) -#include "gda/context_gda_tmpl_device.hpp" #else #error "Select one backend among USE_RO, USE_IPC, USE_GDA" #endif diff --git a/src/util.cpp b/src/util.cpp index 255986eb7b..1a4c89ab33 100644 --- a/src/util.cpp +++ b/src/util.cpp @@ -127,9 +127,14 @@ void rocm_memory_lock_to_fine_grain(void* ptr, size_t size, void** gpu_ptr, rocshmem_env_config::rocshmem_env_config() { char* env_value = NULL; + env_value = getenv("ROCSHMEM_DISABLE_IPC"); + if (NULL != env_value) { + disable_ipc = atoi(env_value); + } + // For backward compatibility, synonymous with ROCSHMEM_DISABLE_IPC env_value = getenv("ROCSHMEM_RO_DISABLE_IPC"); if (NULL != env_value) { - ro_disable_ipc = atoi(env_value); + disable_ipc = atoi(env_value); } env_value = getenv("ROCSHMEM_RO_PROGRESS_DELAY"); @@ -163,8 +168,8 @@ rocshmem_env_config::rocshmem_env_config() { } } -int rocshmem_env_config::get_ro_disable_ipc() { - return ro_disable_ipc; +int rocshmem_env_config::get_disable_ipc() { + return disable_ipc; } int rocshmem_env_config::get_ro_progress_delay() { diff --git a/src/util.hpp b/src/util.hpp index cdd70932fa..24ba037e70 100644 --- a/src/util.hpp +++ b/src/util.hpp @@ -361,7 +361,7 @@ class rocshmem_env_config { public: rocshmem_env_config(); - int get_ro_disable_ipc(); + int get_disable_ipc(); int get_ro_progress_delay(); int get_uniqueid_with_mpi(); int get_bootstrap_timeout(); @@ -370,7 +370,7 @@ public: std::string get_bootstrap_socket_ifname(); private: - int ro_disable_ipc = 0; + int disable_ipc = 0; int ro_progress_delay = 3; int bootstrap_timeout = 5; int uniqueid_with_mpi = 0;