From 31fe937259fa895c87a7bf07eedcf336e430f8a6 Mon Sep 17 00:00:00 2001 From: Yiltan Hassan Temucin Date: Thu, 10 Oct 2024 10:24:11 -0700 Subject: [PATCH] updated atomic_fetch() parameters [ROCm/rocshmem commit: 8b3854b2525ab3d6f9a225312930479fc51850a6] --- .../rocshmem/include/roc_shmem/roc_shmem.hpp | 8 ++++---- projects/rocshmem/src/roc_shmem.cpp | 16 ++++++++-------- projects/rocshmem/src/roc_shmem_gpu.cpp | 16 ++++++++-------- projects/rocshmem/src/templates.hpp | 6 +++--- projects/rocshmem/src/templates_host.hpp | 4 ++-- 5 files changed, 25 insertions(+), 25 deletions(-) diff --git a/projects/rocshmem/include/roc_shmem/roc_shmem.hpp b/projects/rocshmem/include/roc_shmem/roc_shmem.hpp index c06d18a6e3..c0484e8bd7 100644 --- a/projects/rocshmem/include/roc_shmem/roc_shmem.hpp +++ b/projects/rocshmem/include/roc_shmem/roc_shmem.hpp @@ -1021,12 +1021,12 @@ __device__ ATTR_NO_INLINE void roc_shmem_threadfence_system(); */ #define ATOMIC_FETCH_API_GEN(T, TNAME) \ __device__ ATTR_NO_INLINE T roc_shmem_ctx_##TNAME##_atomic_fetch( \ - roc_shmem_ctx_t ctx, T *dest, int pe); \ - __device__ ATTR_NO_INLINE T roc_shmem_##TNAME##_atomic_fetch(T *dest, \ + roc_shmem_ctx_t ctx, T *source, int pe); \ + __device__ ATTR_NO_INLINE T roc_shmem_##TNAME##_atomic_fetch(T *source, \ int pe); \ __host__ T roc_shmem_ctx_##TNAME##_atomic_fetch(roc_shmem_ctx_t ctx, \ - T *dest, int pe); \ - __host__ T roc_shmem_##TNAME##_atomic_fetch(T *dest, int pe); + T *source, int pe); \ + __host__ T roc_shmem_##TNAME##_atomic_fetch(T *source, int pe); /* * MACRO DECLARE SHMEM_ATOMIC_ADD APIs diff --git a/projects/rocshmem/src/roc_shmem.cpp b/projects/rocshmem/src/roc_shmem.cpp index 23e2b25a7b..cdf76fa9d3 100644 --- a/projects/rocshmem/src/roc_shmem.cpp +++ b/projects/rocshmem/src/roc_shmem.cpp @@ -392,8 +392,8 @@ __host__ T roc_shmem_atomic_fetch_inc(T *dest, int pe) { } template -__host__ T roc_shmem_atomic_fetch(T *dest, int pe) { - return roc_shmem_atomic_fetch(ROC_SHMEM_HOST_CTX_DEFAULT, dest, pe); +__host__ T roc_shmem_atomic_fetch(T *source, int pe) { + return roc_shmem_atomic_fetch(ROC_SHMEM_HOST_CTX_DEFAULT, source, pe); } template @@ -592,10 +592,10 @@ __host__ T roc_shmem_atomic_fetch_inc(roc_shmem_ctx_t ctx, T *dest, int pe) { } template -__host__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *dest, int pe) { +__host__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *source, int pe) { DPRINTF("Host function: roc_shmem_atomic_fetch\n"); - return get_internal_ctx(ctx)->amo_fetch_add(dest, 0, pe); + return get_internal_ctx(ctx)->amo_fetch_add(source, 0, pe); } template @@ -1099,11 +1099,11 @@ __host__ int roc_shmem_test(T *ivars, int cmp, T val) { #define AMO_EXTENDED_DEF_GEN(T, TNAME) \ __host__ T roc_shmem_ctx_##TNAME##_atomic_fetch(roc_shmem_ctx_t ctx, \ - T *dest, int pe) { \ - return roc_shmem_atomic_fetch(ctx, dest, pe); \ + T *source, int pe) { \ + return roc_shmem_atomic_fetch(ctx, source, pe); \ } \ - __host__ T roc_shmem_##TNAME##_atomic_fetch(T *dest, int pe) { \ - return roc_shmem_atomic_fetch(dest, pe); \ + __host__ T roc_shmem_##TNAME##_atomic_fetch(T *source, int pe) { \ + return roc_shmem_atomic_fetch(source, pe); \ } \ __host__ void roc_shmem_ctx_##TNAME##_atomic_set(roc_shmem_ctx_t ctx, \ T *dest, T value, int pe) { \ diff --git a/projects/rocshmem/src/roc_shmem_gpu.cpp b/projects/rocshmem/src/roc_shmem_gpu.cpp index 9a3a79f6e5..4ff19b20a2 100644 --- a/projects/rocshmem/src/roc_shmem_gpu.cpp +++ b/projects/rocshmem/src/roc_shmem_gpu.cpp @@ -181,8 +181,8 @@ __device__ T roc_shmem_atomic_fetch_inc(T *dest, int pe) { } template -__device__ T roc_shmem_atomic_fetch(T *dest, int pe) { - return roc_shmem_atomic_fetch(ROC_SHMEM_CTX_DEFAULT, dest, pe); +__device__ T roc_shmem_atomic_fetch(T *source, int pe) { + return roc_shmem_atomic_fetch(ROC_SHMEM_CTX_DEFAULT, source, pe); } template @@ -650,10 +650,10 @@ __device__ T roc_shmem_atomic_fetch_inc(roc_shmem_ctx_t ctx, T *dest, int pe) { } template -__device__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *dest, int pe) { +__device__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *source, int pe) { GPU_DPRINTF("Function: roc_shmem_atomic_fetch\n"); - return get_internal_ctx(ctx)->amo_fetch_add(dest, 0, pe); + return get_internal_ctx(ctx)->amo_fetch_add(source, 0, pe); } template @@ -1303,11 +1303,11 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, #define AMO_EXTENDED_DEF_GEN(T, TNAME) \ __device__ T roc_shmem_ctx_##TNAME##_atomic_fetch(roc_shmem_ctx_t ctx, \ - T *dest, int pe) { \ - return roc_shmem_atomic_fetch(ctx, dest, pe); \ + T *source, int pe) { \ + return roc_shmem_atomic_fetch(ctx, source, pe); \ } \ - __device__ T roc_shmem_##TNAME##_atomic_fetch(T *dest, int pe) { \ - return roc_shmem_atomic_fetch(dest, pe); \ + __device__ T roc_shmem_##TNAME##_atomic_fetch(T *source, int pe) { \ + return roc_shmem_atomic_fetch(source, pe); \ } \ __device__ void roc_shmem_ctx_##TNAME##_atomic_set( \ roc_shmem_ctx_t ctx, T *dest, T value, int pe) { \ diff --git a/projects/rocshmem/src/templates.hpp b/projects/rocshmem/src/templates.hpp index bba411151b..025cefdcfa 100644 --- a/projects/rocshmem/src/templates.hpp +++ b/projects/rocshmem/src/templates.hpp @@ -287,7 +287,7 @@ __device__ T roc_shmem_atomic_fetch_inc(T *dest, int pe); * granularity. * * @param[in] ctx Context with which to perform this operation. - * @param[in] dest Destination address. Must be an address on the symmetric + * @param[in] source Source address. Must be an address on the symmetric heap. * @param[in] val The value to be atomically added. * @param[in] pe PE of the remote process. @@ -296,10 +296,10 @@ __device__ T roc_shmem_atomic_fetch_inc(T *dest, int pe); * */ template -__device__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *dest, int pe); +__device__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *source, int pe); template -__device__ T roc_shmem_atomic_fetch(T *dest, int pe); +__device__ T roc_shmem_atomic_fetch(T *source, int pe); /** * @brief Atomically add the value \p val to \p dest on \p pe. diff --git a/projects/rocshmem/src/templates_host.hpp b/projects/rocshmem/src/templates_host.hpp index 97c92e87da..dabd6a3748 100644 --- a/projects/rocshmem/src/templates_host.hpp +++ b/projects/rocshmem/src/templates_host.hpp @@ -102,10 +102,10 @@ template __host__ T roc_shmem_atomic_fetch_inc(T *dest, int pe); template -__host__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *dest, int pe); +__host__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *source, int pe); template -__host__ T roc_shmem_atomic_fetch(T *dest, int pe); +__host__ T roc_shmem_atomic_fetch(T *source, int pe); template __host__ void roc_shmem_atomic_add(roc_shmem_ctx_t ctx, T *dest, T val, int pe);