updated atomic_fetch() parameters

[ROCm/rocshmem commit: 8b3854b252]

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

projects/rocshmem/src/roc_shmem.cpp

@@ -392,8 +392,8 @@ __host__ T roc_shmem_atomic_fetch_inc(T *dest, int pe) {
 }
 
 template <typename T>
-__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 <typename T>
@@ -592,10 +592,10 @@ __host__ T roc_shmem_atomic_fetch_inc(roc_shmem_ctx_t ctx, T *dest, int pe) {
 }
 
 template <typename T>
-__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<T>(dest, 0, pe);
+  return get_internal_ctx(ctx)->amo_fetch_add<T>(source, 0, pe);
 }
 
 template <typename T>
@@ -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<T>(ctx, dest, pe);                           \
+                                                  T *source, int pe) {         \
+    return roc_shmem_atomic_fetch<T>(ctx, source, pe);                         \
   }                                                                            \
-  __host__ T roc_shmem_##TNAME##_atomic_fetch(T *dest, int pe) {               \
-    return roc_shmem_atomic_fetch<T>(dest, pe);                                \
+  __host__ T roc_shmem_##TNAME##_atomic_fetch(T *source, int pe) {             \
+    return roc_shmem_atomic_fetch<T>(source, pe);                              \
   }                                                                            \
   __host__ void roc_shmem_ctx_##TNAME##_atomic_set(roc_shmem_ctx_t ctx,        \
                                                    T *dest, T value, int pe) { \

projects/rocshmem/src/roc_shmem_gpu.cpp

@@ -181,8 +181,8 @@ __device__ T roc_shmem_atomic_fetch_inc(T *dest, int pe) {
 }
 
 template <typename T>
-__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 <typename T>
@@ -650,10 +650,10 @@ __device__ T roc_shmem_atomic_fetch_inc(roc_shmem_ctx_t ctx, T *dest, int pe) {
 }
 
 template <typename T>
-__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<T>(dest, 0, pe);
+  return get_internal_ctx(ctx)->amo_fetch_add<T>(source, 0, pe);
 }
 
 template <typename T>
@@ -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<T>(ctx, dest, pe);                           \
+                                                    T *source, int pe) {       \
+    return roc_shmem_atomic_fetch<T>(ctx, source, pe);                         \
   }                                                                            \
-  __device__ T roc_shmem_##TNAME##_atomic_fetch(T *dest, int pe) {             \
-    return roc_shmem_atomic_fetch<T>(dest, pe);                                \
+  __device__ T roc_shmem_##TNAME##_atomic_fetch(T *source, int pe) {           \
+    return roc_shmem_atomic_fetch<T>(source, pe);                              \
   }                                                                            \
   __device__ void roc_shmem_ctx_##TNAME##_atomic_set(                          \
       roc_shmem_ctx_t ctx, T *dest, T value, int pe) {                         \

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 <typename T>
-__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 <typename T>
-__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.

projects/rocshmem/src/templates_host.hpp

@@ -102,10 +102,10 @@ template <typename T>
 __host__ T roc_shmem_atomic_fetch_inc(T *dest, int pe);
 
 template <typename T>
-__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 <typename T>
-__host__ T roc_shmem_atomic_fetch(T *dest, int pe);
+__host__ T roc_shmem_atomic_fetch(T *source, int pe);
 
 template <typename T>
 __host__ void roc_shmem_atomic_add(roc_shmem_ctx_t ctx, T *dest, T val, int pe);