AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Remove the device-side active-set-based reduction API interface from rocSHMEM

Browse Source 
[ROCm/rocshmem commit: c22048112e]
This commit is contained in:
avinashkethineedi
2024-10-28 21:35:14 +00:00
parent 4e71fa5d22
commit 9f9ecd9a38
2 changed files with 0 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocshmem/include/roc_shmem/roc_shmem.hpp
-4
View File
@@ -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); \
projects/rocshmem/src/roc_shmem_gpu.cpp
-20
View File
@@ -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 <typename T, ROC_SHMEM_OP Op>
__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<T, Op>(dest, source, nreduce, PE_start,
logPE_stride, PE_size, pWrk, pSync);
}
template <typename T, ROC_SHMEM_OP Op>
__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<T, Op>( \
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<T, Op>( \
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<T, Op>(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) { \