From 45a8cb33548cb7f0ecd7a07c6c3f52b1d5260e65 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 28 Aug 2024 08:14:38 -0700 Subject: [PATCH 1/3] Update IPC object * Update the IPC object in the context class with the instance created in the IPC backend --- src/ipc/context_ipc_device.cpp | 39 +++++++++++++++++----------------- src/ipc/context_ipc_device.hpp | 3 --- 2 files changed, 20 insertions(+), 22 deletions(-) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index 7b9f3dd469..c1fa885c99 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -39,7 +39,8 @@ namespace rocshmem { __host__ IPCContext::IPCContext(Backend *b) : Context(b, false) { IPCBackend *backend{static_cast(b)}; - ipcImpl = &backend->ipcImpl; + ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases; + ipcImpl_.shm_size = b->ipcImpl.shm_size; auto *bp{backend->ipc_backend_proxy.get()}; @@ -59,21 +60,21 @@ __device__ void IPCContext::ctx_destroy(){ __device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy(ipcImpl->ipc_bases[local_pe] + L_offset, + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems); + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } __device__ void IPCContext::putmem_nbi(void *dest, const void *source, @@ -115,10 +116,10 @@ __device__ void IPCContext::sync(roc_shmem_team_t team) { __device__ void IPCContext::putmem_wg(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wg(ipcImpl->ipc_bases[local_pe] + L_offset, + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); __syncthreads(); } @@ -126,11 +127,11 @@ __device__ void IPCContext::putmem_wg(void *dest, const void *source, __device__ void IPCContext::getmem_wg(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wg(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems); + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); __syncthreads(); } @@ -147,21 +148,21 @@ __device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, __device__ void IPCContext::putmem_wave(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wave(ipcImpl->ipc_bases[local_pe] + L_offset, + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem_wave(void *dest, const void *source, size_t nelems, int pe) { // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl->shm_size; + int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl->ipc_bases[my_pe]; - ipcImpl->ipcCopy_wave(dest, ipcImpl->ipc_bases[local_pe] + L_offset, + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index afadb9828c..377790ad67 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -234,9 +234,6 @@ class IPCContext : public Context { private: - //context class has IpcImpl object (ipcImpl_) - IpcImpl *ipcImpl{nullptr}; - uint64_t* atomic_base_ptr{nullptr}; char* g_ret; From e1e1ac6df6be452d3ca2e6a030ab44b85e437d09 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 28 Aug 2024 08:30:46 -0700 Subject: [PATCH 2/3] Add atomics * Add atomic_add, atomic_set, atomic_cas, atomic_fetch_add and atomic_fetch_cas to IPC backend --- src/ipc/context_ipc_tmpl_device.hpp | 44 +++++++++++++++++++++-------- 1 file changed, 32 insertions(+), 12 deletions(-) diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index 28cd718f0a..5e697cafaf 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -70,13 +70,21 @@ __device__ void IPCContext::get_nbi(T *dest, const T *source, size_t nelems, // Atomics template -__device__ void IPCContext::amo_add(void *dst, T value, int pe) { - assert(false); +__device__ void IPCContext::amo_add(void *dest, T value, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOAdd( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); } template -__device__ void IPCContext::amo_set(void *dst, T value, int pe) { - assert(false); +__device__ void IPCContext::amo_set(void *dest, T value, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOSet( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); } template @@ -119,20 +127,32 @@ __device__ void IPCContext::amo_xor(void *dst, T value, int pe) { } template -__device__ void IPCContext::amo_cas(void *dst, T value, T cond, int pe) { - assert(false); +__device__ void IPCContext::amo_cas(void *dest, T value, T cond, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOCas( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + value); } template -__device__ T IPCContext::amo_fetch_add(void *dst, T value, int pe) { - assert(false); - return 0; +__device__ T IPCContext::amo_fetch_add(void *dest, T value, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + return ipcImpl_.ipcAMOFetchAdd( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); } template -__device__ T IPCContext::amo_fetch_cas(void *dst, T value, T cond, int pe) { - assert(false); - return 0; +__device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) { + int local_pe = pe % ipcImpl_.shm_size; + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + return ipcImpl_.ipcAMOFetchCas( + reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + value); } // Collectives From 7bbf34d33446f253917b00a9dd237916c0940957 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 5 Sep 2024 11:52:00 -0700 Subject: [PATCH 3/3] remove local_pe calculation from puts, gets and atomics functions * All the PEs are assumed to be accessible using IPC backend --- src/ipc/context_ipc_device.cpp | 24 ++++++------------------ src/ipc/context_ipc_tmpl_device.hpp | 15 +++++---------- 2 files changed, 11 insertions(+), 28 deletions(-) diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp index c1fa885c99..4bf7072aa5 100644 --- a/src/ipc/context_ipc_device.cpp +++ b/src/ipc/context_ipc_device.cpp @@ -59,22 +59,18 @@ __device__ void IPCContext::ctx_destroy(){ __device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); } __device__ void IPCContext::putmem_nbi(void *dest, const void *source, @@ -115,23 +111,19 @@ __device__ void IPCContext::sync(roc_shmem_team_t team) { __device__ void IPCContext::putmem_wg(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[pe] + L_offset, const_cast(source), nelems); __syncthreads(); } __device__ void IPCContext::getmem_wg(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); + ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); __syncthreads(); } @@ -147,22 +139,18 @@ __device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, __device__ void IPCContext::putmem_wave(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[pe] + L_offset, const_cast(source), nelems); } __device__ void IPCContext::getmem_wave(void *dest, const void *source, size_t nelems, int pe) { - // TODO (Avinash) check if PE is available for IPC using (isIpcAvailable) - int local_pe = pe % ipcImpl_.shm_size; const char *src_typed = reinterpret_cast(source); uint64_t L_offset = const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; - ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, + ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); } diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index 5e697cafaf..94ef855736 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -71,20 +71,18 @@ __device__ void IPCContext::get_nbi(T *dest, const T *source, size_t nelems, // Atomics template __device__ void IPCContext::amo_add(void *dest, T value, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; ipcImpl_.ipcAMOAdd( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); } template __device__ void IPCContext::amo_set(void *dest, T value, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; ipcImpl_.ipcAMOSet( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); } template @@ -128,30 +126,27 @@ __device__ void IPCContext::amo_xor(void *dst, T value, int pe) { template __device__ void IPCContext::amo_cas(void *dest, T value, T cond, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; ipcImpl_.ipcAMOCas( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), cond, value); } template __device__ T IPCContext::amo_fetch_add(void *dest, T value, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; return ipcImpl_.ipcAMOFetchAdd( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), value); + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); } template __device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) { - int local_pe = pe % ipcImpl_.shm_size; uint64_t L_offset = reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; return ipcImpl_.ipcAMOFetchCas( - reinterpret_cast(ipcImpl_.ipc_bases[local_pe] + L_offset), cond, + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), cond, value); }