From 9f9ecd9a3887a6d25066160488b27f99edd906cf Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Mon, 28 Oct 2024 21:35:14 +0000 Subject: [PATCH] Remove the device-side active-set-based reduction API interface from rocSHMEM [ROCm/rocshmem commit: c22048112e3800597ac049e681327ee1a7474cb6] --- .../rocshmem/include/roc_shmem/roc_shmem.hpp | 4 ---- projects/rocshmem/src/roc_shmem_gpu.cpp | 20 ------------------- 2 files changed, 24 deletions(-) diff --git a/projects/rocshmem/include/roc_shmem/roc_shmem.hpp b/projects/rocshmem/include/roc_shmem/roc_shmem.hpp index c0484e8bd7..7a0b58ad6a 100644 --- a/projects/rocshmem/include/roc_shmem/roc_shmem.hpp +++ b/projects/rocshmem/include/roc_shmem/roc_shmem.hpp @@ -837,10 +837,6 @@ __device__ ATTR_NO_INLINE void roc_shmem_threadfence_system(); * MACRO DECLARE SHMEM_REDUCTION APIs */ #define REDUCTION_API_GEN(T, TNAME, Op_API) \ - __device__ ATTR_NO_INLINE void roc_shmem_ctx_##TNAME##_##Op_API##_wg_to_all( \ - roc_shmem_ctx_t ctx, T *dest, const T *source, int nreduce, \ - int PE_start, int logPE_stride, int PE_size, T *pWrk, \ - long *pSync); /* NOLINT */ \ __device__ ATTR_NO_INLINE void roc_shmem_ctx_##TNAME##_##Op_API##_wg_to_all( \ roc_shmem_ctx_t ctx, roc_shmem_team_t team, T *dest, const T *source, \ int nreduce); \ diff --git a/projects/rocshmem/src/roc_shmem_gpu.cpp b/projects/rocshmem/src/roc_shmem_gpu.cpp index 582249eec9..87635a4160 100644 --- a/projects/rocshmem/src/roc_shmem_gpu.cpp +++ b/projects/rocshmem/src/roc_shmem_gpu.cpp @@ -429,17 +429,6 @@ __device__ void *roc_shmem_ptr(const void *dest, int pe) { return get_internal_ctx(ROC_SHMEM_CTX_DEFAULT)->shmem_ptr(dest, pe); } -template -__device__ void roc_shmem_wg_to_all(roc_shmem_ctx_t ctx, T *dest, - const T *source, int nreduce, int PE_start, - int logPE_stride, int PE_size, T *pWrk, - long *pSync) { - GPU_DPRINTF("Function: roc_shmem_to_all\n"); - - get_internal_ctx(ctx)->to_all(dest, source, nreduce, PE_start, - logPE_stride, PE_size, pWrk, pSync); -} - template __device__ void roc_shmem_wg_to_all(roc_shmem_ctx_t ctx, roc_shmem_team_t team, T *dest, const T *source, int nreduce) { @@ -886,9 +875,6 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, * Template generator for reductions */ #define REDUCTION_GEN(T, Op) \ - template __device__ void roc_shmem_wg_to_all( \ - roc_shmem_ctx_t ctx, T * dest, const T *source, int nreduce, \ - int PE_start, int logPE_stride, int PE_size, T *pWrk, long *pSync); \ template __device__ void roc_shmem_wg_to_all( \ roc_shmem_ctx_t ctx, roc_shmem_team_t team, T * dest, const T *source, \ int nreduce); @@ -1100,12 +1086,6 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, **/ #define REDUCTION_DEF_GEN(T, TNAME, Op_API, Op) \ - __device__ void roc_shmem_ctx_##TNAME##_##Op_API##_wg_to_all( \ - roc_shmem_ctx_t ctx, T *dest, const T *source, int nreduce, \ - int PE_start, int logPE_stride, int PE_size, T *pWrk, long *pSync) { \ - roc_shmem_wg_to_all(ctx, dest, source, nreduce, PE_start, \ - logPE_stride, PE_size, pWrk, pSync); \ - } \ __device__ void roc_shmem_ctx_##TNAME##_##Op_API##_wg_to_all( \ roc_shmem_ctx_t ctx, roc_shmem_team_t team, T *dest, const T *source, \ int nreduce) { \