From 1ddc82fc28599999e9cf41c49d3f3081729a862a Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Mon, 9 Sep 2024 07:01:20 -0700 Subject: [PATCH] remove pSync from internal_bcast functions remove the pSync arguments from the internal_broadcast functions, they are not used anyway. [ROCm/rocshmem commit: dfcacdc4a389a0bc32a28438a293981736127f7c] --- projects/rocshmem/src/ipc/context_ipc_device.hpp | 6 ++---- .../rocshmem/src/ipc/context_ipc_tmpl_device.hpp | 16 ++++------------ 2 files changed, 6 insertions(+), 16 deletions(-) diff --git a/projects/rocshmem/src/ipc/context_ipc_device.hpp b/projects/rocshmem/src/ipc/context_ipc_device.hpp index 01fb1c1fb8..4cc6d08d56 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_device.hpp @@ -207,13 +207,11 @@ class IPCContext : public Context { template __device__ void internal_put_broadcast(T *dst, const T *src, int nelems, int pe_root, int PE_start, - int logPE_stride, int PE_size, - long *pSync); // NOLINT(runtime/int) + int logPE_stride, int PE_size); // NOLINT(runtime/int) template __device__ void internal_get_broadcast(T *dst, const T *src, int nelems, - int pe_root, - long *pSync); // NOLINT(runtime/int) + int pe_root); // NOLINT(runtime/int) template __device__ void fcollect_linear(roc_shmem_team_t team, T *dest, diff --git a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp index af79cd20b1..81ec0b2b97 100644 --- a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp @@ -154,8 +154,7 @@ __device__ void IPCContext::to_all(T *dest, const T *source, int nreduce, template __device__ void IPCContext::internal_put_broadcast( T *dst, const T *src, int nelems, int pe_root, int pe_start, - int log_pe_stride, int pe_size, - [[maybe_unused]] long *p_sync) { // NOLINT(runtime/int) + int log_pe_stride, int pe_size) { // NOLINT(runtime/int) if (my_pe == pe_root) { int stride = 1 << log_pe_stride; int finish = pe_start + stride * pe_size; @@ -169,8 +168,7 @@ __device__ void IPCContext::internal_put_broadcast( template __device__ void IPCContext::internal_get_broadcast( - T *dst, const T *src, int nelems, int pe_root, - [[maybe_unused]] long *pSync) { // NOLINT(runtime/int) + T *dst, const T *src, int nelems, int pe_root) { // NOLINT(runtime/int) if (my_pe != pe_root) { get_wg(dst, src, nelems, pe_root); } @@ -202,9 +200,9 @@ __device__ void IPCContext::broadcast(T *dst, const T *src, int nelems, long *p_sync) { // NOLINT(runtime/int) if (num_pes < 4) { internal_put_broadcast(dst, src, nelems, pe_root, pe_start, log_pe_stride, - pe_size, p_sync); + pe_size); } else { - internal_get_broadcast(dst, src, nelems, pe_root, p_sync); + internal_get_broadcast(dst, src, nelems, pe_root); } // Synchronize on completion of broadcast @@ -232,9 +230,6 @@ __device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst, long *pSync = team_obj->alltoall_pSync; int my_pe_in_team = team_obj->my_pe; - /* TODO (egabriel): shouldn't the loop handle pe_start, pe_size, and stride? - * otherwise what is the point of the few lines above? - */ // Have each PE put their designated data to the other PEs for (int j = 0; j < pe_size; j++) { int dest_pe = team_obj->get_pe_in_world(j); @@ -268,9 +263,6 @@ __device__ void IPCContext::fcollect_linear(roc_shmem_team_t team, T *dst, long *pSync = team_obj->alltoall_pSync; int my_pe_in_team = team_obj->my_pe; - /* TODO (egabriel): shouldn't the loop handle pe_start, pe_size, and stride? - * otherwise what is the point of the few lines above? - */ // Have each PE put their designated data to the other PEs for (int j = 0; j < pe_size; j++) { int dest_pe = team_obj->get_pe_in_world(j);