From e3b0353fa9a4036eb001702fe417b07b03ce87ca Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Thu, 17 Apr 2025 10:11:32 -0500 Subject: [PATCH] use correct id when accessing ipc-bases (#88) we need to use the position of that processes in the local ipc-bases array, not the global rank. --- src/ipc_policy.hpp | 11 ++- src/reverse_offload/context_ro_device.cpp | 80 ++++++++++--------- .../context_ro_tmpl_device.hpp | 14 ++-- 3 files changed, 57 insertions(+), 48 deletions(-) diff --git a/src/ipc_policy.hpp b/src/ipc_policy.hpp index ed13aa1631..c1560ab8ba 100644 --- a/src/ipc_policy.hpp +++ b/src/ipc_policy.hpp @@ -57,11 +57,14 @@ class IpcOnImpl { __host__ void ipcHostStop(); - __device__ bool isIpcAvailable(int my_pe, int target_pe) { + __device__ bool isIpcAvailable(int my_pe, int target_pe, int *local_target_pe) { if (nullptr == pes_with_ipc_avail) { return false; } for (int i=0; i>; public: + int shm_rank{0}; + uint32_t shm_size{0}; char **ipc_bases{nullptr}; @@ -134,7 +139,7 @@ class IpcOffImpl { __host__ void ipcHostStop() {} - __device__ bool isIpcAvailable(int my_pe, int target_pe) { return false; } + __device__ bool isIpcAvailable(int my_pe, int target_pe, int *local_target_pe) { return false; } __device__ void ipcGpuInit(Backend *rocshmem_handle, Context *ctx, int thread_id) {} diff --git a/src/reverse_offload/context_ro_device.cpp b/src/reverse_offload/context_ro_device.cpp index 6d652795be..b38545a368 100644 --- a/src/reverse_offload/context_ro_device.cpp +++ b/src/reverse_offload/context_ro_device.cpp @@ -60,15 +60,16 @@ __host__ ROContext::ROContext(Backend *b, ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases; ipcImpl_.shm_size = b->ipcImpl.shm_size; + ipcImpl_.shm_rank = b->ipcImpl.shm_rank; ipcImpl_.pes_with_ipc_avail = b->ipcImpl.pes_with_ipc_avail; } __device__ void ROContext::putmem(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } else { @@ -85,11 +86,11 @@ __device__ void ROContext::putmem(void *dest, const void *source, size_t nelems, __device__ void ROContext::getmem(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } else { bool must_send_message = wf_coal_.coalesce(pe, source, dest, &nelems); @@ -105,10 +106,10 @@ __device__ void ROContext::getmem(void *dest, const void *source, size_t nelems, __device__ void ROContext::putmem_nbi(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } else { @@ -124,11 +125,11 @@ __device__ void ROContext::putmem_nbi(void *dest, const void *source, __device__ void ROContext::getmem_nbi(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } else { bool must_send_message = wf_coal_.coalesce(pe, source, dest, &nelems); @@ -162,11 +163,12 @@ __device__ void ROContext::quiet() { __device__ void *ROContext::shmem_ptr(const void *dest, int pe) { void *ret = nullptr; - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { + 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[my_pe]; - ret = ipcImpl_.ipc_bases[pe] + L_offset; + reinterpret_cast(dst) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; + ret = ipcImpl_.ipc_bases[local_pe] + L_offset; } return ret; } @@ -291,10 +293,10 @@ __device__ void ROContext::ctx_destroy() { __device__ void ROContext::putmem_wg(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } else { @@ -310,11 +312,11 @@ __device__ void ROContext::putmem_wg(void *dest, const void *source, __device__ void ROContext::getmem_wg(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } else { if (is_thread_zero_in_block()) { @@ -329,10 +331,10 @@ __device__ void ROContext::getmem_wg(void *dest, const void *source, __device__ void ROContext::putmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } else { @@ -347,11 +349,11 @@ __device__ void ROContext::putmem_nbi_wg(void *dest, const void *source, __device__ void ROContext::getmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } else { if (is_thread_zero_in_block()) { @@ -365,10 +367,10 @@ __device__ void ROContext::getmem_nbi_wg(void *dest, const void *source, __device__ void ROContext::putmem_wave(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } else { @@ -383,11 +385,11 @@ __device__ void ROContext::putmem_wave(void *dest, const void *source, __device__ void ROContext::getmem_wave(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } else { @@ -402,10 +404,10 @@ __device__ void ROContext::getmem_wave(void *dest, const void *source, __device__ void ROContext::putmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { uint64_t L_offset = - reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + reinterpret_cast(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset, const_cast(source), nelems); } else { @@ -419,11 +421,11 @@ __device__ void ROContext::putmem_nbi_wave(void *dest, const void *source, __device__ void ROContext::getmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - int local_pe = pe % ipcImpl_.shm_size; + int local_pe{-1}; + if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) { const char *src_typed = reinterpret_cast(source); uint64_t L_offset = - const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]; ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems); } else { diff --git a/src/reverse_offload/context_ro_tmpl_device.hpp b/src/reverse_offload/context_ro_tmpl_device.hpp index d6ce0f2828..aadb3dafdc 100644 --- a/src/reverse_offload/context_ro_tmpl_device.hpp +++ b/src/reverse_offload/context_ro_tmpl_device.hpp @@ -145,9 +145,10 @@ __device__ void ROContext::put_nbi(T *dest, const T *source, size_t nelems, template __device__ void ROContext::p(T *dest, T value, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { - long L_offset{reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]}; - ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[pe] + L_offset, + 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)); } else { build_queue_element(RO_NET_P, dest, &value, sizeof(T), pe, 0, 0, 0, nullptr, @@ -158,11 +159,12 @@ __device__ void ROContext::p(T *dest, T value, int pe) { template __device__ T ROContext::g(const T *source, int pe) { - if (ipcImpl_.isIpcAvailable(my_pe, pe)) { + 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[my_pe]}; + long L_offset{const_cast(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]}; T dest; - ipcImpl_.ipcCopy(&dest, ipcImpl_.ipc_bases[pe] + L_offset, sizeof(T)); + ipcImpl_.ipcCopy(&dest, ipcImpl_.ipc_bases[local_pe] + L_offset, sizeof(T)); return dest; } else { auto dest{get_g_ret_buf()};