diff --git a/include/roc_shmem/roc_shmem.hpp b/include/roc_shmem/roc_shmem.hpp index 16bd0a6892..d5e90ab0db 100644 --- a/include/roc_shmem/roc_shmem.hpp +++ b/include/roc_shmem/roc_shmem.hpp @@ -2086,39 +2086,39 @@ TEST_API_GEN(unsigned long long, ulonglong) // NOLINT(runtime/int) *****************************************************************************/ /* - * MACRO DECLARE SHMEMX_PUT APIs + * MACRO DECLARE SHMEM_PUT APIs */ #define PUT_API_EXT_GEN(GRAN, T, TNAME) \ - __device__ ATTR_NO_INLINE void roc_shmemx_ctx_##TNAME##_put_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_ctx_##TNAME##_put_##GRAN( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); \ - __device__ ATTR_NO_INLINE void roc_shmemx_##TNAME##_put_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_##TNAME##_put_##GRAN( \ T *dest, const T *source, size_t nelems, int pe); /* - * MACRO DECLARE SHMEMX_GET APIs + * MACRO DECLARE SHMEM_GET APIs */ #define GET_API_EXT_GEN(GRAN, T, TNAME) \ - __device__ ATTR_NO_INLINE void roc_shmemx_ctx_##TNAME##_get_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_ctx_##TNAME##_get_##GRAN( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); \ - __device__ ATTR_NO_INLINE void roc_shmemx_##TNAME##_get_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_##TNAME##_get_##GRAN( \ T *dest, const T *source, size_t nelems, int pe); /* - * MACRO DECLARE SHMEMX_PUT_NBI APIs + * MACRO DECLARE SHMEM_PUT_NBI APIs */ #define PUT_NBI_API_EXT_GEN(GRAN, T, TNAME) \ - __device__ ATTR_NO_INLINE void roc_shmemx_ctx_##TNAME##_put_nbi_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_ctx_##TNAME##_put_nbi_##GRAN( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); \ - __device__ ATTR_NO_INLINE void roc_shmemx_##TNAME##_put_nbi_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_##TNAME##_put_nbi_##GRAN( \ T *dest, const T *source, size_t nelems, int pe); /* - * MACRO DECLARE SHMEMX_GET_NBI APIs + * MACRO DECLARE SHMEM_GET_NBI APIs */ #define GET_NBI_API_EXT_GEN(GRAN, T, TNAME) \ - __device__ ATTR_NO_INLINE void roc_shmemx_ctx_##TNAME##_get_nbi_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_ctx_##TNAME##_get_nbi_##GRAN( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); \ - __device__ ATTR_NO_INLINE void roc_shmemx_##TNAME##_get_nbi_##GRAN( \ + __device__ ATTR_NO_INLINE void roc_shmem_##TNAME##_get_nbi_##GRAN( \ T *dest, const T *source, size_t nelems, int pe); /** @@ -2140,12 +2140,12 @@ TEST_API_GEN(unsigned long long, ulonglong) // NOLINT(runtime/int) * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_putmem_wave( +__device__ ATTR_NO_INLINE void roc_shmem_ctx_putmem_wave( roc_shmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_putmem_wave(void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_putmem_wave(void *dest, + const void *source, + size_t nelems, int pe); /** * @brief Writes contiguous data of \p nelems bytes from \p source on the @@ -2166,14 +2166,14 @@ __device__ ATTR_NO_INLINE void roc_shmemx_putmem_wave(void *dest, * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_putmem_wg(roc_shmem_ctx_t ctx, - void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_ctx_putmem_wg(roc_shmem_ctx_t ctx, + void *dest, + const void *source, + size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_putmem_wg(void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_putmem_wg(void *dest, + const void *source, + size_t nelems, int pe); /** * @brief Writes contiguous data of \p nelems elements from \p source on the @@ -2265,12 +2265,12 @@ PUT_API_EXT_GEN(wg, unsigned long long, ulonglong) // NOLINT(runtime/int) * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_getmem_wave( +__device__ ATTR_NO_INLINE void roc_shmem_ctx_getmem_wave( roc_shmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_getmem_wave(void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_getmem_wave(void *dest, + const void *source, + size_t nelems, int pe); /** * @brief Reads contiguous data of \p nelems bytes from \p source on \p pe @@ -2290,14 +2290,14 @@ __device__ ATTR_NO_INLINE void roc_shmemx_getmem_wave(void *dest, * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_getmem_wg(roc_shmem_ctx_t ctx, - void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_ctx_getmem_wg(roc_shmem_ctx_t ctx, + void *dest, + const void *source, + size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_getmem_wg(void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_getmem_wg(void *dest, + const void *source, + size_t nelems, int pe); /** * @brief Reads contiguous data of \p nelems elements from \p source on \p pe @@ -2389,13 +2389,13 @@ GET_API_EXT_GEN(wg, unsigned long long, ulonglong) // NOLINT(runtime/int) * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_putmem_nbi_wave( +__device__ ATTR_NO_INLINE void roc_shmem_ctx_putmem_nbi_wave( roc_shmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_putmem_nbi_wave(void *dest, - const void *source, - size_t nelems, - int pe); +__device__ ATTR_NO_INLINE void roc_shmem_putmem_nbi_wave(void *dest, + const void *source, + size_t nelems, + int pe); /** * @brief Writes contiguous data of \p nelems elements from \p source on the @@ -2454,12 +2454,12 @@ PUT_NBI_API_EXT_GEN(wave, unsigned long long, ulonglong) // NOLINT * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_putmem_nbi_wg( +__device__ ATTR_NO_INLINE void roc_shmem_ctx_putmem_nbi_wg( roc_shmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_putmem_nbi_wg(void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_putmem_nbi_wg(void *dest, + const void *source, + size_t nelems, int pe); /** * @brief Writes contiguous data of \p nelems elements from \p source on the @@ -2518,13 +2518,13 @@ PUT_NBI_API_EXT_GEN(wg, unsigned long long, ulonglong) // NOLINT(runtime/int) * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_getmem_nbi_wave( +__device__ ATTR_NO_INLINE void roc_shmem_ctx_getmem_nbi_wave( roc_shmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_getmem_nbi_wave(void *dest, - const void *source, - size_t nelems, - int pe); +__device__ ATTR_NO_INLINE void roc_shmem_getmem_nbi_wave(void *dest, + const void *source, + size_t nelems, + int pe); /** * @brief Reads contiguous data of \p nelems elements from \p source on \p pe @@ -2583,12 +2583,12 @@ GET_NBI_API_EXT_GEN(wave, unsigned long long, ulonglong) // NOLINT * * @return void. */ -__device__ ATTR_NO_INLINE void roc_shmemx_ctx_getmem_nbi_wg( +__device__ ATTR_NO_INLINE void roc_shmem_ctx_getmem_nbi_wg( roc_shmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe); -__device__ ATTR_NO_INLINE void roc_shmemx_getmem_nbi_wg(void *dest, - const void *source, - size_t nelems, int pe); +__device__ ATTR_NO_INLINE void roc_shmem_getmem_nbi_wg(void *dest, + const void *source, + size_t nelems, int pe); /** * @brief Reads contiguous data of \p nelems elements from \p source on \p pe diff --git a/internal/clients/shmem_rccl/primitive_tester.cpp b/internal/clients/shmem_rccl/primitive_tester.cpp index ef265a7fb4..7982bc5356 100644 --- a/internal/clients/shmem_rccl/primitive_tester.cpp +++ b/internal/clients/shmem_rccl/primitive_tester.cpp @@ -47,7 +47,7 @@ PrimitiveTest(int loop, int block_id = hipBlockIdx_x; for(int i =0; i< loop; i++){ - roc_shmemx_ctx_putmem_nbi_wg(ctx, &r_buf[my_pe*size], &s_buf[block_id * size], size, block_id); + roc_shmem_ctx_putmem_nbi_wg(ctx, &r_buf[my_pe*size], &s_buf[block_id * size], size, block_id); if(hipThreadIdx_x==0){ //roc_shmem_ctx_quiet(ctx); //roc_shmem_ctx_threadfence_system(ctx); diff --git a/src/roc_shmem_gpu.cpp b/src/roc_shmem_gpu.cpp index 186823ac61..2deff5d0b7 100644 --- a/src/roc_shmem_gpu.cpp +++ b/src/roc_shmem_gpu.cpp @@ -716,132 +716,132 @@ __device__ void roc_shmem_atomic_xor(roc_shmem_ctx_t ctx, T *dest, T val, /** * SHMEM X RMA API for WG and Wave level */ -__device__ void roc_shmemx_ctx_putmem_wave(roc_shmem_ctx_t ctx, void *dest, - const void *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_putmem_wave\n"); +__device__ void roc_shmem_ctx_putmem_wave(roc_shmem_ctx_t ctx, void *dest, + const void *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_putmem_wave\n"); get_internal_ctx(ctx)->putmem_wave(dest, source, nelems, pe); } -__device__ void roc_shmemx_ctx_putmem_wg(roc_shmem_ctx_t ctx, void *dest, - const void *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_putmem_wg\n"); +__device__ void roc_shmem_ctx_putmem_wg(roc_shmem_ctx_t ctx, void *dest, + const void *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_putmem_wg\n"); get_internal_ctx(ctx)->putmem_wg(dest, source, nelems, pe); } -__device__ void roc_shmemx_ctx_putmem_nbi_wave(roc_shmem_ctx_t ctx, void *dest, - const void *source, - size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_putmem_nbi_wave\n"); +__device__ void roc_shmem_ctx_putmem_nbi_wave(roc_shmem_ctx_t ctx, void *dest, + const void *source, + size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_putmem_nbi_wave\n"); get_internal_ctx(ctx)->putmem_nbi_wave(dest, source, nelems, pe); } -__device__ void roc_shmemx_ctx_putmem_nbi_wg(roc_shmem_ctx_t ctx, void *dest, - const void *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_putmem_nbi_wg\n"); +__device__ void roc_shmem_ctx_putmem_nbi_wg(roc_shmem_ctx_t ctx, void *dest, + const void *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_putmem_nbi_wg\n"); get_internal_ctx(ctx)->putmem_nbi_wg(dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_wave(roc_shmem_ctx_t ctx, T *dest, - const T *source, size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_put_wave\n"); +__device__ void roc_shmem_put_wave(roc_shmem_ctx_t ctx, T *dest, + const T *source, size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_put_wave\n"); get_internal_ctx(ctx)->put_wave(dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, - size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_put_wg\n"); +__device__ void roc_shmem_put_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, + size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_put_wg\n"); get_internal_ctx(ctx)->put_wg(dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_nbi_wave(roc_shmem_ctx_t ctx, T *dest, - const T *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_put_nbi_wave\n"); +__device__ void roc_shmem_put_nbi_wave(roc_shmem_ctx_t ctx, T *dest, + const T *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_put_nbi_wave\n"); get_internal_ctx(ctx)->put_nbi_wave(dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_nbi_wg(roc_shmem_ctx_t ctx, T *dest, - const T *source, size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_put_nbi_wg\n"); +__device__ void roc_shmem_put_nbi_wg(roc_shmem_ctx_t ctx, T *dest, + const T *source, size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_put_nbi_wg\n"); get_internal_ctx(ctx)->put_nbi_wg(dest, source, nelems, pe); } -__device__ void roc_shmemx_ctx_getmem_wg(roc_shmem_ctx_t ctx, void *dest, - const void *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_getmem_wg\n"); +__device__ void roc_shmem_ctx_getmem_wg(roc_shmem_ctx_t ctx, void *dest, + const void *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_getmem_wg\n"); get_internal_ctx(ctx)->getmem_wg(dest, source, nelems, pe); } -__device__ void roc_shmemx_ctx_getmem_wave(roc_shmem_ctx_t ctx, void *dest, - const void *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_getmem_wave\n"); +__device__ void roc_shmem_ctx_getmem_wave(roc_shmem_ctx_t ctx, void *dest, + const void *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_getmem_wave\n"); get_internal_ctx(ctx)->getmem_wave(dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, - size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_get_wg\n"); +__device__ void roc_shmem_get_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, + size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_get_wg\n"); get_internal_ctx(ctx)->get_wg(dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_wave(roc_shmem_ctx_t ctx, T *dest, - const T *source, size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_get_wave\n"); +__device__ void roc_shmem_get_wave(roc_shmem_ctx_t ctx, T *dest, + const T *source, size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_get_wave\n"); get_internal_ctx(ctx)->get_wave(dest, source, nelems, pe); } -__device__ void roc_shmemx_ctx_getmem_nbi_wg(roc_shmem_ctx_t ctx, void *dest, - const void *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_getmem_nbi_wg\n"); +__device__ void roc_shmem_ctx_getmem_nbi_wg(roc_shmem_ctx_t ctx, void *dest, + const void *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_getmem_nbi_wg\n"); get_internal_ctx(ctx)->getmem_nbi_wg(dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_nbi_wg(roc_shmem_ctx_t ctx, T *dest, - const T *source, size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_get_nbi_wg\n"); +__device__ void roc_shmem_get_nbi_wg(roc_shmem_ctx_t ctx, T *dest, + const T *source, size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_get_nbi_wg\n"); get_internal_ctx(ctx)->get_nbi_wg(dest, source, nelems, pe); } -__device__ void roc_shmemx_ctx_getmem_nbi_wave(roc_shmem_ctx_t ctx, void *dest, - const void *source, - size_t nelems, int pe) { - GPU_DPRINTF("Function: roc_shmemx_ctx_getmem_nbi_wave\n"); +__device__ void roc_shmem_ctx_getmem_nbi_wave(roc_shmem_ctx_t ctx, void *dest, + const void *source, + size_t nelems, int pe) { + GPU_DPRINTF("Function: roc_shmem_ctx_getmem_nbi_wave\n"); get_internal_ctx(ctx)->getmem_nbi_wave(dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_nbi_wave(roc_shmem_ctx_t ctx, T *dest, - const T *source, size_t nelems, - int pe) { - GPU_DPRINTF("Function: roc_shmemx_get_nbi_wave\n"); +__device__ void roc_shmem_get_nbi_wave(roc_shmem_ctx_t ctx, T *dest, + const T *source, size_t nelems, + int pe) { + GPU_DPRINTF("Function: roc_shmem_get_nbi_wave\n"); get_internal_ctx(ctx)->get_nbi_wave(dest, source, nelems, pe); } @@ -903,38 +903,38 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, template __device__ void roc_shmem_wg_fcollect( \ roc_shmem_ctx_t ctx, roc_shmem_team_t team, T * dest, const T *source, \ int nelem); \ - template __device__ void roc_shmemx_put_wave( \ + template __device__ void roc_shmem_put_wave( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_put_wg( \ + template __device__ void roc_shmem_put_wg( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_put_wave(T * dest, const T *source, \ - size_t nelems, int pe); \ - template __device__ void roc_shmemx_put_wg(T * dest, const T *source, \ - size_t nelems, int pe); \ - template __device__ void roc_shmemx_put_nbi_wave( \ + template __device__ void roc_shmem_put_wave(T * dest, const T *source, \ + size_t nelems, int pe); \ + template __device__ void roc_shmem_put_wg(T * dest, const T *source, \ + size_t nelems, int pe); \ + template __device__ void roc_shmem_put_nbi_wave( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_put_nbi_wg( \ + template __device__ void roc_shmem_put_nbi_wg( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_put_nbi_wave( \ + template __device__ void roc_shmem_put_nbi_wave( \ T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_put_nbi_wg(T * dest, const T *source, \ - size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_wave( \ + template __device__ void roc_shmem_put_nbi_wg(T * dest, const T *source, \ + size_t nelems, int pe); \ + template __device__ void roc_shmem_get_wave( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_wg( \ + template __device__ void roc_shmem_get_wg( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_wave(T * dest, const T *source, \ - size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_wg(T * dest, const T *source, \ - size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_nbi_wave( \ + template __device__ void roc_shmem_get_wave(T * dest, const T *source, \ + size_t nelems, int pe); \ + template __device__ void roc_shmem_get_wg(T * dest, const T *source, \ + size_t nelems, int pe); \ + template __device__ void roc_shmem_get_nbi_wave( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_nbi_wg( \ + template __device__ void roc_shmem_get_nbi_wg( \ roc_shmem_ctx_t ctx, T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_nbi_wave( \ + template __device__ void roc_shmem_get_nbi_wave( \ T * dest, const T *source, size_t nelems, int pe); \ - template __device__ void roc_shmemx_get_nbi_wg(T * dest, const T *source, \ - size_t nelems, int pe); + template __device__ void roc_shmem_get_nbi_wg(T * dest, const T *source, \ + size_t nelems, int pe); /** * Declare templates for the standard amo types @@ -1142,69 +1142,69 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, __device__ T roc_shmem_##TNAME##_g(const T *source, int pe) { \ return roc_shmem_g(source, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_put_wave( \ + __device__ void roc_shmem_ctx_##TNAME##_put_wave( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_put_wave(ctx, dest, source, nelems, pe); \ + roc_shmem_put_wave(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_put_wg( \ + __device__ void roc_shmem_ctx_##TNAME##_put_wg( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_put_wg(ctx, dest, source, nelems, pe); \ + roc_shmem_put_wg(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_put_wave(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_put_wave(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_put_wave(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_put_wave(dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_put_wg(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_put_wg(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_put_wg(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_put_wg(dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_put_nbi_wave( \ + __device__ void roc_shmem_ctx_##TNAME##_put_nbi_wave( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_put_nbi_wave(ctx, dest, source, nelems, pe); \ + roc_shmem_put_nbi_wave(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_put_nbi_wg( \ + __device__ void roc_shmem_ctx_##TNAME##_put_nbi_wg( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_put_nbi_wg(ctx, dest, source, nelems, pe); \ + roc_shmem_put_nbi_wg(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_put_nbi_wave(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_put_nbi_wave(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_put_nbi_wave(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_put_nbi_wave(dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_put_nbi_wg(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_put_nbi_wg(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_put_nbi_wg(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_put_nbi_wg(dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_get_wave( \ + __device__ void roc_shmem_ctx_##TNAME##_get_wave( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_get_wave(ctx, dest, source, nelems, pe); \ + roc_shmem_get_wave(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_get_wg( \ + __device__ void roc_shmem_ctx_##TNAME##_get_wg( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_get_wg(ctx, dest, source, nelems, pe); \ + roc_shmem_get_wg(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_get_wave(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_get_wave(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_get_wave(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_get_wave(dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_get_wg(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_get_wg(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_get_wg(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_get_wg(dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_get_nbi_wave( \ + __device__ void roc_shmem_ctx_##TNAME##_get_nbi_wave( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_get_nbi_wave(ctx, dest, source, nelems, pe); \ + roc_shmem_get_nbi_wave(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_ctx_##TNAME##_get_nbi_wg( \ + __device__ void roc_shmem_ctx_##TNAME##_get_nbi_wg( \ roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe) { \ - roc_shmemx_get_nbi_wg(ctx, dest, source, nelems, pe); \ + roc_shmem_get_nbi_wg(ctx, dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_get_nbi_wave(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_get_nbi_wave(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_get_nbi_wave(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_get_nbi_wave(dest, source, nelems, pe); \ } \ - __device__ void roc_shmemx_##TNAME##_get_nbi_wg(T *dest, const T *source, \ - size_t nelems, int pe) { \ - roc_shmemx_get_nbi_wg(dest, source, nelems, pe); \ + __device__ void roc_shmem_##TNAME##_get_nbi_wg(T *dest, const T *source, \ + size_t nelems, int pe) { \ + roc_shmem_get_nbi_wg(dest, source, nelems, pe); \ } \ __device__ void roc_shmem_ctx_##TNAME##_wg_broadcast( \ roc_shmem_ctx_t ctx, roc_shmem_team_t team, T *dest, const T *source, \ @@ -1380,8 +1380,8 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, return roc_shmem_wait_until_some_vector(ivars, nelems, indices, \ status, cmp, vals); \ } \ - __device__ int roc_shmem_##TNAME##_test(T *ivars, int cmp, T val) { \ - return roc_shmem_test(ivars, cmp, val); \ + __device__ int roc_shmem_##TNAME##_test(T *ivars, int cmp, T val) { \ + return roc_shmem_test(ivars, cmp, val); \ } /****************************************************************************** diff --git a/src/templates.hpp b/src/templates.hpp index 025cefdcfa..1b482a219b 100644 --- a/src/templates.hpp +++ b/src/templates.hpp @@ -489,11 +489,11 @@ __device__ void roc_shmem_wg_to_all(roc_shmem_ctx_t ctx, T *dest, * */ template -__device__ void roc_shmemx_put_wave(roc_shmem_ctx_t ctx, T *dest, +__device__ void roc_shmem_put_wave(roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); template -__device__ void roc_shmemx_put_wave(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_put_wave(T *dest, const T *source, size_t nelems, int pe); /** @@ -517,11 +517,11 @@ __device__ void roc_shmemx_put_wave(T *dest, const T *source, size_t nelems, * */ template -__device__ void roc_shmemx_put_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, +__device__ void roc_shmem_put_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); template -__device__ void roc_shmemx_put_wg(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_put_wg(T *dest, const T *source, size_t nelems, int pe); /** @@ -544,11 +544,11 @@ __device__ void roc_shmemx_put_wg(T *dest, const T *source, size_t nelems, * */ template -__device__ void roc_shmemx_get_wave(roc_shmem_ctx_t ctx, T *dest, +__device__ void roc_shmem_get_wave(roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); template -__device__ void roc_shmemx_get_wave(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_wave(T *dest, const T *source, size_t nelems, int pe); /** @@ -571,11 +571,11 @@ __device__ void roc_shmemx_get_wave(T *dest, const T *source, size_t nelems, * */ template -__device__ void roc_shmemx_get_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, +__device__ void roc_shmem_get_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); template -__device__ void roc_shmemx_get_wg(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_wg(T *dest, const T *source, size_t nelems, int pe); /** @@ -599,11 +599,11 @@ __device__ void roc_shmemx_get_wg(T *dest, const T *source, size_t nelems, * */ template -__device__ void roc_shmemx_put_nbi_wave(roc_shmem_ctx_t ctx, T *dest, +__device__ void roc_shmem_put_nbi_wave(roc_shmem_ctx_t ctx, T *dest, const T *src, size_t nelems, int pe); template -__device__ void roc_shmemx_put_nbi_wave(T *dest, const T *src, size_t nelems, +__device__ void roc_shmem_put_nbi_wave(T *dest, const T *src, size_t nelems, int pe); /** @@ -627,11 +627,11 @@ __device__ void roc_shmemx_put_nbi_wave(T *dest, const T *src, size_t nelems, * */ template -__device__ void roc_shmemx_put_nbi_wg(roc_shmem_ctx_t ctx, T *dest, +__device__ void roc_shmem_put_nbi_wg(roc_shmem_ctx_t ctx, T *dest, const T *src, size_t nelems, int pe); template -__device__ void roc_shmemx_put_nbi_wg(T *dest, const T *src, size_t nelems, +__device__ void roc_shmem_put_nbi_wg(T *dest, const T *src, size_t nelems, int pe); /** @@ -655,11 +655,11 @@ __device__ void roc_shmemx_put_nbi_wg(T *dest, const T *src, size_t nelems, * */ template -__device__ void roc_shmemx_get_nbi_wave(roc_shmem_ctx_t ctx, T *dest, +__device__ void roc_shmem_get_nbi_wave(roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); template -__device__ void roc_shmemx_get_nbi_wave(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_nbi_wave(T *dest, const T *source, size_t nelems, int pe); /** @@ -683,101 +683,101 @@ __device__ void roc_shmemx_get_nbi_wave(T *dest, const T *source, size_t nelems, * */ template -__device__ void roc_shmemx_get_nbi_wg(roc_shmem_ctx_t ctx, T *dest, +__device__ void roc_shmem_get_nbi_wg(roc_shmem_ctx_t ctx, T *dest, const T *source, size_t nelems, int pe); template -__device__ void roc_shmemx_get_nbi_wg(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_nbi_wg(T *dest, const T *source, size_t nelems, int pe); -__device__ void roc_shmemx_putmem_wave(void *dest, const void *source, +__device__ void roc_shmem_putmem_wave(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_putmem_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_ctx_putmem_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } -__device__ void roc_shmemx_putmem_wg(void *dest, const void *source, +__device__ void roc_shmem_putmem_wg(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_putmem_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_ctx_putmem_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_wave(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_put_wave(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_put_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_put_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_wg(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_put_wg(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_put_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_put_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } -__device__ void roc_shmemx_getmem_wg(void *dest, const void *source, +__device__ void roc_shmem_getmem_wg(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_getmem_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_ctx_getmem_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_wg(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_wg(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_get_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_get_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } -__device__ void roc_shmemx_getmem_wave(void *dest, const void *source, +__device__ void roc_shmem_getmem_wave(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_getmem_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_ctx_getmem_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_wave(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_wave(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_get_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_get_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } -__device__ void roc_shmemx_putmem_nbi_wg(void *dest, const void *source, +__device__ void roc_shmem_putmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_putmem_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_ctx_putmem_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_nbi_wg(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_put_nbi_wg(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_put_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_put_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } -__device__ void roc_shmemx_putmem_nbi_wave(void *dest, const void *source, +__device__ void roc_shmem_putmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_putmem_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, + roc_shmem_ctx_putmem_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_put_nbi_wave(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_put_nbi_wave(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_put_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_put_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } -__device__ void roc_shmemx_getmem_nbi_wg(void *dest, const void *source, +__device__ void roc_shmem_getmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_getmem_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_ctx_getmem_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_nbi_wg(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_nbi_wg(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_get_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_get_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } -__device__ void roc_shmemx_getmem_nbi_wave(void *dest, const void *source, +__device__ void roc_shmem_getmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) { - roc_shmemx_ctx_getmem_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, + roc_shmem_ctx_getmem_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } template -__device__ void roc_shmemx_get_nbi_wave(T *dest, const T *source, size_t nelems, +__device__ void roc_shmem_get_nbi_wave(T *dest, const T *source, size_t nelems, int pe) { - roc_shmemx_get_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); + roc_shmem_get_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe); } } // namespace rocshmem diff --git a/tests/functional_tests/extended_primitives.cpp b/tests/functional_tests/extended_primitives.cpp index d6c6042faa..8b540a565a 100644 --- a/tests/functional_tests/extended_primitives.cpp +++ b/tests/functional_tests/extended_primitives.cpp @@ -54,16 +54,16 @@ __global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer, switch (type) { case WGGetTestType: - roc_shmemx_ctx_getmem_wg(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_getmem_wg(ctx, r_buf, s_buf, size, 1); break; case WGGetNBITestType: - roc_shmemx_ctx_getmem_nbi_wg(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_getmem_nbi_wg(ctx, r_buf, s_buf, size, 1); break; case WGPutTestType: - roc_shmemx_ctx_putmem_wg(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_putmem_wg(ctx, r_buf, s_buf, size, 1); break; case WGPutNBITestType: - roc_shmemx_ctx_putmem_nbi_wg(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_putmem_nbi_wg(ctx, r_buf, s_buf, size, 1); break; default: break; diff --git a/tests/functional_tests/wave_level_primitives.cpp b/tests/functional_tests/wave_level_primitives.cpp index 56b263ebf4..598e5ad2e7 100644 --- a/tests/functional_tests/wave_level_primitives.cpp +++ b/tests/functional_tests/wave_level_primitives.cpp @@ -56,16 +56,16 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer, switch (type) { case WAVEGetTestType: - roc_shmemx_ctx_getmem_wave(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_getmem_wave(ctx, r_buf, s_buf, size, 1); break; case WAVEGetNBITestType: - roc_shmemx_ctx_getmem_nbi_wave(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_getmem_nbi_wave(ctx, r_buf, s_buf, size, 1); break; case WAVEPutTestType: - roc_shmemx_ctx_putmem_wave(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_putmem_wave(ctx, r_buf, s_buf, size, 1); break; case WAVEPutNBITestType: - roc_shmemx_ctx_putmem_nbi_wave(ctx, r_buf, s_buf, size, 1); + roc_shmem_ctx_putmem_nbi_wave(ctx, r_buf, s_buf, size, 1); break; default: break; diff --git a/tests/sos_tests/bcast_flood.cpp b/tests/sos_tests/bcast_flood.cpp index cef952b073..e8ad1bfd3f 100644 --- a/tests/sos_tests/bcast_flood.cpp +++ b/tests/sos_tests/bcast_flood.cpp @@ -45,7 +45,7 @@ using namespace rocshmem; static int atoi_scaled(char *s); static void usage(char *pgm); -static double shmemx_wtime(void) { +static double shmem_wtime(void) { struct timeval tv; gettimeofday(&tv, NULL); return (double)tv.tv_sec + (double)tv.tv_usec / 1000000.0; @@ -154,14 +154,14 @@ int main(int argc, char **argv) { roc_shmem_barrier_all(); for (time_taken = 0.0, ps = i = 0; i < loops; i++) { - start_time = shmemx_wtime(); + start_time = shmem_wtime(); roc_shmem_ctx_int_broadcast(ROC_SHMEM_CTX_DEFAULT, target, source, elements, 0, 0, 0, npes, &pSync[ps]); if (Serialize) roc_shmem_barrier_all(); - time_taken += (shmemx_wtime() - start_time); + time_taken += (shmem_wtime() - start_time); if (ps_cnt > 1) { ps += ROC_SHMEM_BCAST_SYNC_SIZE; diff --git a/tests/sos_tests/bigget.cpp b/tests/sos_tests/bigget.cpp index d0105b6497..d841f637c1 100644 --- a/tests/sos_tests/bigget.cpp +++ b/tests/sos_tests/bigget.cpp @@ -83,7 +83,7 @@ static void usage(char *pgm) { pgm, NUM_ELEMENTS, DFLT_LOOPS); } -static inline double shmemx_wtime(void) { +static inline double shmem_wtime(void) { struct timeval tv; gettimeofday(&tv, 0); return (double)((tv.tv_usec / 1000000.0) + tv.tv_sec); @@ -186,11 +186,11 @@ int main(int argc, char **argv) { roc_shmem_barrier_all(); for (i = 0; i < loops; i++) { - start_time = shmemx_wtime(); + start_time = shmem_wtime(); roc_shmem_int_get(Target, Source, elements, target_pe); - time_taken += shmemx_wtime() - start_time; + time_taken += shmem_wtime() - start_time; if (me == 0) { if (Track && i > 0 && ((i % 200) == 0)) fprintf(stderr, ".%d", i); diff --git a/tests/sos_tests/bigput.cpp b/tests/sos_tests/bigput.cpp index 635a7d7efd..ffdb3ff98c 100644 --- a/tests/sos_tests/bigput.cpp +++ b/tests/sos_tests/bigput.cpp @@ -83,7 +83,7 @@ static void usage(char *pgm) { pgm, NUM_ELEMENTS, DFLT_LOOPS); } -static inline double shmemx_wtime(void) { +static inline double shmem_wtime(void) { struct timeval tv; gettimeofday(&tv, 0); return (double)((tv.tv_usec / 1000000.0) + tv.tv_sec); @@ -214,11 +214,11 @@ int main(int argc, char **argv) { roc_shmem_barrier_all(); for (i = 0; i < loops; i++) { - start_time = shmemx_wtime(); + start_time = shmem_wtime(); roc_shmem_int_put(Target, Source, elements, target_PE); - time_taken += (shmemx_wtime() - start_time); + time_taken += (shmem_wtime() - start_time); if (me == 0) { if (Track && i > 0 && ((i % 200) == 0)) fprintf(stderr, ".%d", i); diff --git a/tests/sos_tests/lfinc.cpp b/tests/sos_tests/lfinc.cpp index 45fb9aa74f..9402520c7e 100644 --- a/tests/sos_tests/lfinc.cpp +++ b/tests/sos_tests/lfinc.cpp @@ -44,7 +44,7 @@ using namespace rocshmem; #define LOOPS 25000 -static double shmemx_wtime(void) { +static double shmem_wtime(void) { struct timeval tv; gettimeofday(&tv, NULL); return (double)tv.tv_sec + (double)tv.tv_usec / 1000000.0; @@ -86,11 +86,11 @@ int main(int argc, char *argv[]) { roc_shmem_barrier_all(); neighbor = (my_pe + 1) % npes; - start_time = shmemx_wtime(); + start_time = shmem_wtime(); for (j = 0, elapsed = 0.0; j < loops; j++) { - start_time = shmemx_wtime(); + start_time = shmem_wtime(); lval = roc_shmem_int64_atomic_fetch_inc((int64_t *)&data[1], neighbor); - elapsed += shmemx_wtime() - start_time; + elapsed += shmem_wtime() - start_time; if (lval != (long)j) { fprintf(stderr, "[%d] Test: FAIL previous val %ld != %d Exit.\n", my_pe, lval, j);