From 1335d98f1f817996f82fbb5aaf0e1bef87cd48fc Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Wed, 28 Aug 2024 08:14:38 -0700 Subject: [PATCH] Update IPC object * Update the IPC object in the context class with the instance created in the IPC backend [ROCm/rocshmem commit: 45a8cb33548cb7f0ecd7a07c6c3f52b1d5260e65] --- .../rocshmem/src/ipc/context_ipc_device.cpp | 39 ++++++++++--------- .../rocshmem/src/ipc/context_ipc_device.hpp | 3 -- 2 files changed, 20 insertions(+), 22 deletions(-) diff --git a/projects/rocshmem/src/ipc/context_ipc_device.cpp b/projects/rocshmem/src/ipc/context_ipc_device.cpp index 7b9f3dd469..c1fa885c99 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.cpp +++ b/projects/rocshmem/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/projects/rocshmem/src/ipc/context_ipc_device.hpp b/projects/rocshmem/src/ipc/context_ipc_device.hpp index afadb9828c..377790ad67 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.hpp +++ b/projects/rocshmem/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;