From dba989733f4d6afb458d9b33f278125a9c94f93d Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Mon, 17 Feb 2025 06:18:46 +0000 Subject: [PATCH 1/2] Update bandwidth and latency calculations - Refined bandwidth and latency calculations for improved accuracy [ROCm/rocshmem commit: c155636da4a979f71fd5a0f9544a6456121f94c5] --- .../functional_tests/alltoall_tester.cpp | 17 +- .../functional_tests/amo_bitwise_tester.cpp | 22 ++- .../functional_tests/amo_extended_tester.cpp | 26 +-- .../functional_tests/amo_standard_tester.cpp | 26 +-- .../functional_tests/barrier_all_tester.cpp | 11 +- .../tests/functional_tests/empty_tester.cpp | 10 +- .../functional_tests/extended_primitives.cpp | 19 +- .../functional_tests/fcollect_tester.cpp | 17 +- .../functional_tests/ping_all_tester.cpp | 12 +- .../functional_tests/ping_all_tester.hpp | 3 +- .../functional_tests/ping_pong_tester.cpp | 12 +- .../functional_tests/ping_pong_tester.hpp | 3 +- .../functional_tests/primitive_mr_tester.cpp | 15 +- .../functional_tests/primitive_tester.cpp | 14 +- .../functional_tests/random_access_tester.cpp | 23 ++- .../functional_tests/random_access_tester.hpp | 5 +- .../signaling_operations_tester.cpp | 22 +-- .../tests/functional_tests/swarm_tester.cpp | 20 ++- .../tests/functional_tests/swarm_tester.hpp | 3 +- .../tests/functional_tests/sync_tester.cpp | 12 +- .../team_broadcast_tester.cpp | 15 +- .../team_ctx_primitive_tester.cpp | 22 ++- .../team_reduction_tester.cpp | 17 +- .../tests/functional_tests/tester.cpp | 167 +++++++----------- .../tests/functional_tests/tester.hpp | 15 +- .../wave_level_primitives.cpp | 25 +-- 26 files changed, 283 insertions(+), 270 deletions(-) diff --git a/projects/rocshmem/tests/functional_tests/alltoall_tester.cpp b/projects/rocshmem/tests/functional_tests/alltoall_tester.cpp index ffd543e28c..e4ff9b7e14 100644 --- a/projects/rocshmem/tests/functional_tests/alltoall_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/alltoall_tester.cpp @@ -58,10 +58,12 @@ rocshmem_team_t team_alltoall_world_dup; * DEVICE TEST KERNEL *****************************************************************************/ template -__global__ void AlltoallTest(int loop, int skip, uint64_t *timer, - T1 *source_buf, T1 *dest_buf, int size, - ShmemContextType ctx_type, rocshmem_team_t team) { +__global__ void AlltoallTest(int loop, int skip, long long int *start_time, + long long int *end_time, T1 *source_buf, + T1 *dest_buf, int size, ShmemContextType ctx_type, + rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -70,10 +72,9 @@ __global__ void AlltoallTest(int loop, int skip, uint64_t *timer, __syncthreads(); - uint64_t start; for (int i = 0; i < loop + skip; i++) { if (i == skip && hipThreadIdx_x == 0) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } wg_alltoall(ctx, team, dest_buf, // T* dest @@ -84,7 +85,7 @@ __global__ void AlltoallTest(int loop, int skip, uint64_t *timer, __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -126,8 +127,8 @@ void AlltoallTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(AlltoallTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, source_buf, dest_buf, size, - _shmem_context, team_alltoall_world_dup); + stream, loop, args.skip, start_time, end_time, source_buf, + dest_buf, size, _shmem_context, team_alltoall_world_dup); num_msgs = loop + args.skip; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp index 426fa2247a..971ad2a484 100644 --- a/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp @@ -29,7 +29,8 @@ using namespace rocshmem; /* Declare the global kernel template with a generic implementation */ template -__global__ void AMOBitwiseTest(int loop, int skip, uint64_t *timer, char *r_buf, +__global__ void AMOBitwiseTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *r_buf, T *s_buf, T *ret_val, TestType type, ShmemContextType ctx_type) { return; @@ -64,8 +65,8 @@ void AMOBitwiseTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(AMOBitwiseTest, gridsize, blocksize, shared_bytes, stream, - loop, args.skip, timer, _r_buf, _s_buf, _ret_val, _type, - _shmem_context); + loop, args.skip, start_time, end_time, _r_buf, _s_buf, + _ret_val, _type, _shmem_context); _gridSize = gridsize; num_msgs = (loop + args.skip) * gridsize.x; @@ -123,17 +124,20 @@ void AMOBitwiseTester::verifyResults(uint64_t size) { #define AMO_BITWISE_DEF_GEN(T, TNAME) \ template <> \ __global__ void AMOBitwiseTest( \ - int loop, int skip, uint64_t *timer, char *r_buf, T *s_buf, T *ret_val, \ + int loop, int skip, long long int *start_time, \ + long long int *end_time, char *r_buf, T *s_buf, T *ret_val, \ TestType type, ShmemContextType ctx_type) { \ __shared__ rocshmem_ctx_t ctx; \ + int wg_id = get_flat_grid_id(); \ rocshmem_wg_init(); \ rocshmem_wg_ctx_create(ctx_type, &ctx); \ if (hipThreadIdx_x == 0) { \ - uint64_t start; \ T ret = 0; \ T cond = 0; \ for (int i = 0; i < loop + skip; i++) { \ - if (i == skip) start = rocshmem_timer(); \ + if (i == skip) { \ + start_time[wg_id] = wall_clock64(); \ + } \ switch (type) { \ case AMO_FetchAndTestType: \ ret = rocshmem_ctx_##TNAME##_atomic_fetch_and(ctx, (T *)r_buf, \ @@ -161,9 +165,9 @@ void AMOBitwiseTester::verifyResults(uint64_t size) { } \ } \ rocshmem_ctx_quiet(ctx); \ - timer[hipBlockIdx_x] = rocshmem_timer() - start; \ - ret_val[hipBlockIdx_x] = ret; \ - rocshmem_ctx_getmem(ctx, &s_buf[hipBlockIdx_x], r_buf, sizeof(T), 1); \ + end_time[wg_id] = wall_clock64(); \ + ret_val[wg_id] = ret; \ + rocshmem_ctx_getmem(ctx, &s_buf[wg_id], r_buf, sizeof(T), 1); \ } \ rocshmem_wg_ctx_destroy(&ctx); \ rocshmem_wg_finalize(); \ diff --git a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp index 44b3f0865c..83da81d6c9 100644 --- a/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp @@ -29,9 +29,10 @@ using namespace rocshmem; /* Declare the global kernel template with a generic implementation */ template -__global__ void AMOExtendedTest(int loop, int skip, uint64_t *timer, - char *r_buf, T *s_buf, T *ret_val, - TestType type, ShmemContextType ctx_type) { +__global__ void AMOExtendedTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *r_buf, + T *s_buf, T *ret_val, TestType type, + ShmemContextType ctx_type) { return; } @@ -64,8 +65,8 @@ void AMOExtendedTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(AMOExtendedTest, gridsize, blocksize, shared_bytes, stream, - loop, args.skip, timer, _r_buf, _s_buf, _ret_val, _type, - _shmem_context); + loop, args.skip, start_time, end_time, _r_buf, _s_buf, + _ret_val, _type, _shmem_context); _gridSize = gridsize; num_msgs = (loop + args.skip) * gridsize.x; @@ -111,17 +112,20 @@ void AMOExtendedTester::verifyResults(uint64_t size) { #define AMO_EXTENDED_DEF_GEN(T, TNAME) \ template <> \ __global__ void AMOExtendedTest( \ - int loop, int skip, uint64_t *timer, char *r_buf, T *s_buf, T *ret_val, \ + int loop, int skip, long long int *start_time, \ + long long int *end_time, char *r_buf, T *s_buf, T *ret_val, \ TestType type, ShmemContextType ctx_type) { \ __shared__ rocshmem_ctx_t ctx; \ + int wg_id = get_flat_grid_id(); \ rocshmem_wg_init(); \ rocshmem_wg_ctx_create(ctx_type, &ctx); \ if (hipThreadIdx_x == 0) { \ - uint64_t start; \ T ret = 0; \ T cond = 0; \ for (int i = 0; i < loop + skip; i++) { \ - if (i == skip) start = rocshmem_timer(); \ + if (i == skip) { \ + start_time[wg_id] = wall_clock64(); \ + } \ switch (type) { \ case AMO_FetchTestType: \ ret = rocshmem_ctx_##TNAME##_atomic_fetch(ctx, (T *)r_buf, 1); \ @@ -138,9 +142,9 @@ void AMOExtendedTester::verifyResults(uint64_t size) { } \ } \ rocshmem_ctx_quiet(ctx); \ - timer[hipBlockIdx_x] = rocshmem_timer() - start; \ - ret_val[hipBlockIdx_x] = ret; \ - rocshmem_ctx_getmem(ctx, &s_buf[hipBlockIdx_x], r_buf, sizeof(T), 1); \ + end_time[wg_id] = wall_clock64(); \ + ret_val[wg_id] = ret; \ + rocshmem_ctx_getmem(ctx, &s_buf[wg_id], r_buf, sizeof(T), 1); \ } \ rocshmem_wg_ctx_destroy(&ctx); \ rocshmem_wg_finalize(); \ diff --git a/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp b/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp index 93009b08de..3eabdf5fe6 100644 --- a/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp @@ -29,9 +29,10 @@ using namespace rocshmem; /* Declare the global kernel template with a generic implementation */ template -__global__ void AMOStandardTest(int loop, int skip, uint64_t *timer, - char *r_buf, T *s_buf, T *ret_val, - TestType type, ShmemContextType ctx_type) { +__global__ void AMOStandardTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *r_buf, + T *s_buf, T *ret_val, TestType type, + ShmemContextType ctx_type) { return; } @@ -64,8 +65,8 @@ void AMOStandardTester::launchKernel(dim3 gridsize, dim3 blocksize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(AMOStandardTest, gridsize, blocksize, shared_bytes, stream, - loop, args.skip, timer, _r_buf, _s_buf, _ret_val, _type, - _shmem_context); + loop, args.skip, start_time, end_time, _r_buf, _s_buf, + _ret_val, _type, _shmem_context); _gridSize = gridsize; num_msgs = (loop + args.skip) * gridsize.x; @@ -119,17 +120,20 @@ void AMOStandardTester::verifyResults(uint64_t size) { #define AMO_STANDARD_DEF_GEN(T, TNAME) \ template <> \ __global__ void AMOStandardTest( \ - int loop, int skip, uint64_t *timer, char *r_buf, T *s_buf, T *ret_val, \ + int loop, int skip, long long int *start_time, \ + long long int *end_time, char *r_buf, T *s_buf, T *ret_val, \ TestType type, ShmemContextType ctx_type) { \ __shared__ rocshmem_ctx_t ctx; \ + int wg_id = get_flat_grid_id(); \ rocshmem_wg_init(); \ rocshmem_wg_ctx_create(ctx_type, &ctx); \ if (hipThreadIdx_x == 0) { \ - uint64_t start; \ T ret = 0; \ T cond = 0; \ for (int i = 0; i < loop + skip; i++) { \ - if (i == skip) start = rocshmem_timer(); \ + if (i == skip) { \ + start_time[wg_id] = wall_clock64(); \ + } \ switch (type) { \ case AMO_FAddTestType: \ ret = rocshmem_ctx_##TNAME##_atomic_fetch_add(ctx, (T *)r_buf, 2, \ @@ -155,9 +159,9 @@ void AMOStandardTester::verifyResults(uint64_t size) { } \ } \ rocshmem_ctx_quiet(ctx); \ - timer[hipBlockIdx_x] = rocshmem_timer() - start; \ - ret_val[hipBlockIdx_x] = ret; \ - rocshmem_ctx_getmem(ctx, &s_buf[hipBlockIdx_x], r_buf, sizeof(T), 1); \ + end_time[wg_id] = wall_clock64(); \ + ret_val[wg_id] = ret; \ + rocshmem_ctx_getmem(ctx, &s_buf[wg_id], r_buf, sizeof(T), 1); \ } \ rocshmem_wg_ctx_destroy(&ctx); \ rocshmem_wg_finalize(); \ diff --git a/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp b/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp index 54faaa9ffd..01bb2bd2a4 100644 --- a/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp @@ -29,16 +29,17 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void BarrierAllTest(int loop, int skip, uint64_t *timer) { +__global__ void BarrierAllTest(int loop, int skip, long long int *start_time, + long long int *end_time) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ROCSHMEM_CTX_WG_PRIVATE, &ctx); - uint64_t start; for (int i = 0; i < loop + skip; i++) { if (hipThreadIdx_x == 0 && i == skip) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } __syncthreads(); @@ -48,7 +49,7 @@ __global__ void BarrierAllTest(int loop, int skip, uint64_t *timer) { __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -67,7 +68,7 @@ void BarrierAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(BarrierAllTest, gridSize, blockSize, shared_bytes, stream, - loop, args.skip, timer); + loop, args.skip, start_time, end_time); num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/empty_tester.cpp b/projects/rocshmem/tests/functional_tests/empty_tester.cpp index 5bc9facd77..119e75952d 100644 --- a/projects/rocshmem/tests/functional_tests/empty_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/empty_tester.cpp @@ -29,8 +29,9 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void EmptyTest(int loop, int skip, uint64_t *timer, int size, - TestType type, ShmemContextType ctx_type) { +__global__ void EmptyTest(int loop, int skip, long long int *start_time, + long long int *end_time, int size, TestType type, + ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -52,8 +53,9 @@ void EmptyTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, uint64_t size) { size_t shared_bytes = 0; - hipLaunchKernelGGL(EmptyTest, gridSize, blockSize, shared_bytes, stream, loop, - args.skip, timer, size, _type, _shmem_context); + hipLaunchKernelGGL(EmptyTest, gridSize, blockSize, shared_bytes, stream, + loop, args.skip, start_time, end_time, size, _type, + _shmem_context); } void EmptyTester::verifyResults(uint64_t size) {} diff --git a/projects/rocshmem/tests/functional_tests/extended_primitives.cpp b/projects/rocshmem/tests/functional_tests/extended_primitives.cpp index d3a4a7ba4f..26a3ec50a1 100644 --- a/projects/rocshmem/tests/functional_tests/extended_primitives.cpp +++ b/projects/rocshmem/tests/functional_tests/extended_primitives.cpp @@ -31,11 +31,13 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer, - char *s_buf, char *r_buf, int size, - TestType type, +__global__ void ExtendedPrimitiveTest(int loop, int skip, + long long int *start_time, + long long int *end_time, char *s_buf, + char *r_buf, int size, TestType type, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -44,13 +46,14 @@ __global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer, * If the number of work groups is greater than 1, this kernel performs a * tiled functional test */ - uint64_t start; uint64_t idx = size * get_flat_grid_id(); s_buf += idx; r_buf += idx; for (int i = 0; i < loop + skip; i++) { - if (i == skip) start = rocshmem_timer(); + if (i == skip) { + start_time[wg_id] = wall_clock64(); + } switch (type) { case WGGetTestType: @@ -73,7 +76,7 @@ __global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer, rocshmem_ctx_quiet(ctx); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -105,8 +108,8 @@ void ExtendedPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, size_t shared_bytes = 0; hipLaunchKernelGGL(ExtendedPrimitiveTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, (char*)s_buf, - (char*)r_buf, size, _type, _shmem_context); + stream, loop, args.skip, start_time, end_time, + (char*)s_buf, (char*)r_buf, size, _type, _shmem_context); num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop * gridSize.x; diff --git a/projects/rocshmem/tests/functional_tests/fcollect_tester.cpp b/projects/rocshmem/tests/functional_tests/fcollect_tester.cpp index 3eaa8fbe40..3f0efa3445 100644 --- a/projects/rocshmem/tests/functional_tests/fcollect_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/fcollect_tester.cpp @@ -58,10 +58,12 @@ FCOLLECT_DEF_GEN(unsigned long long, ulonglong) * DEVICE TEST KERNEL *****************************************************************************/ template -__global__ void FcollectTest(int loop, int skip, uint64_t *timer, - T1 *source_buf, T1 *dest_buf, int size, - ShmemContextType ctx_type, rocshmem_team_t team) { +__global__ void FcollectTest(int loop, int skip, long long int *start_time, + long long int *end_time, T1 *source_buf, + T1 *dest_buf, int size, ShmemContextType ctx_type, + rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -69,10 +71,9 @@ __global__ void FcollectTest(int loop, int skip, uint64_t *timer, int n_pes = rocshmem_ctx_n_pes(ctx); __syncthreads(); - uint64_t start; for (int i = 0; i < loop + skip; i++) { if (i == skip && hipThreadIdx_x == 0) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } wg_fcollect(ctx, team, dest_buf, // T* dest @@ -83,7 +84,7 @@ __global__ void FcollectTest(int loop, int skip, uint64_t *timer, __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -125,8 +126,8 @@ void FcollectTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(FcollectTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, source_buf, dest_buf, size, - _shmem_context, team_fcollect_world_dup); + stream, loop, args.skip, start_time, end_time, source_buf, + dest_buf, size, _shmem_context, team_fcollect_world_dup); num_msgs = loop + args.skip; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp b/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp index 858882c7c7..01111779b9 100644 --- a/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/ping_all_tester.cpp @@ -29,9 +29,11 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf, +__global__ void PingAllTest(int loop, int skip, long long int *start_time, + long long int *end_time, int *r_buf, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -44,12 +46,11 @@ __global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf, } if (hipThreadIdx_x == 0) { - uint64_t start; auto blk_pe_off {hipBlockIdx_x * num_pe}; for (int i = 0; i < loop + skip; i++) { if (i == skip) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } for (int j{0}; j < num_pe; j++) { @@ -57,7 +58,7 @@ __global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf, } rocshmem_int_wait_until_all(&r_buf[blk_pe_off], num_pe, status, ROCSHMEM_CMP_EQ, 1); } - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); rocshmem_wg_finalize(); @@ -83,7 +84,8 @@ void PingAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(PingAllTest, gridSize, blockSize, shared_bytes, stream, - loop, args.skip, timer, r_buf, _shmem_context); + loop, args.skip, start_time, end_time, r_buf, + _shmem_context); num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/ping_all_tester.hpp b/projects/rocshmem/tests/functional_tests/ping_all_tester.hpp index 33f5f1945d..ce4dcd900b 100644 --- a/projects/rocshmem/tests/functional_tests/ping_all_tester.hpp +++ b/projects/rocshmem/tests/functional_tests/ping_all_tester.hpp @@ -28,7 +28,8 @@ /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf); +__global__ void PingAllTest(int loop, int skip, long long int *start_time, + long long int *end_time, int *r_buf); /****************************************************************************** * HOST TESTER CLASS diff --git a/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp b/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp index f0e6ebb1ea..05e32dfa2e 100644 --- a/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp @@ -29,9 +29,11 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf, +__global__ void PingPongTest(int loop, int skip, long long int *start_time, + long long int *end_time, int *r_buf, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -39,11 +41,10 @@ __global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf, int pe = rocshmem_ctx_my_pe(ctx); if (hipThreadIdx_x == 0) { - uint64_t start; for (int i = 0; i < loop + skip; i++) { if (i == skip) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } if (pe == 0) { @@ -56,7 +57,7 @@ __global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf, rocshmem_ctx_int_p(ctx, &r_buf[hipBlockIdx_x], i + 1, 0); } } - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); rocshmem_wg_finalize(); @@ -80,7 +81,8 @@ void PingPongTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(PingPongTest, gridSize, blockSize, shared_bytes, stream, - loop, args.skip, timer, r_buf, _shmem_context); + loop, args.skip, start_time, end_time, r_buf, + _shmem_context); num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/ping_pong_tester.hpp b/projects/rocshmem/tests/functional_tests/ping_pong_tester.hpp index 0669e78477..2b784773a3 100644 --- a/projects/rocshmem/tests/functional_tests/ping_pong_tester.hpp +++ b/projects/rocshmem/tests/functional_tests/ping_pong_tester.hpp @@ -28,7 +28,8 @@ /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf); +__global__ void PingPongTest(int loop, int skip, long long int *start_time, + long long int *end_time, int *r_buf); /****************************************************************************** * HOST TESTER CLASS diff --git a/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp b/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp index bec7a1c9dc..d40656a76b 100644 --- a/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp @@ -29,17 +29,18 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void PrimitiveMRTest(int loop, uint64_t *timer, char *s_buf, +__global__ void PrimitiveMRTest(int loop, long long int *start_time, + long long int *end_time, char *s_buf, char *r_buf, int size, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); if (hipThreadIdx_x == 0) { - uint64_t start; - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); for (int win_i = 0; win_i < 64 * loop; win_i++) { for (int i = 0; i < 64; i++) { @@ -48,7 +49,7 @@ __global__ void PrimitiveMRTest(int loop, uint64_t *timer, char *s_buf, rocshmem_ctx_quiet(ctx); } - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } __syncthreads(); @@ -81,11 +82,13 @@ void PrimitiveMRTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, /* Warmup */ hipLaunchKernelGGL(PrimitiveMRTest, gridSize, blockSize, shared_bytes, stream, - loop, timer, s_buf, r_buf, size, _shmem_context); + loop, start_time, end_time, s_buf, r_buf, size, + _shmem_context); /* Benchmark */ hipLaunchKernelGGL(PrimitiveMRTest, gridSize, blockSize, shared_bytes, stream, - loop, timer, s_buf, r_buf, size, _shmem_context); + loop, start_time, end_time, s_buf, r_buf, size, + _shmem_context); CHECK_HIP(hipDeviceSynchronize()); diff --git a/projects/rocshmem/tests/functional_tests/primitive_tester.cpp b/projects/rocshmem/tests/functional_tests/primitive_tester.cpp index 1b66664b32..135733b2af 100644 --- a/projects/rocshmem/tests/functional_tests/primitive_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/primitive_tester.cpp @@ -29,19 +29,19 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void PrimitiveTest(int loop, int skip, uint64_t *timer, char *s_buf, +__global__ void PrimitiveTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *s_buf, char *r_buf, int size, TestType type, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); - uint64_t start; - for (int i = 0; i < loop + skip; i++) { if (i == skip) { __syncthreads(); - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } switch (type) { @@ -79,7 +79,7 @@ __global__ void PrimitiveTest(int loop, int skip, uint64_t *timer, char *s_buf, __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -109,8 +109,8 @@ void PrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(PrimitiveTest, gridSize, blockSize, shared_bytes, stream, - loop, args.skip, timer, s_buf, r_buf, size, _type, - _shmem_context); + loop, args.skip, start_time, end_time, s_buf, r_buf, + size, _type, _shmem_context); num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp index 3f09a593ce..233cf11c73 100644 --- a/projects/rocshmem/tests/functional_tests/random_access_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/random_access_tester.cpp @@ -49,13 +49,14 @@ __device__ bool thread_passing(int num_bins, uint32_t *bin_threads, return pass; } -__global__ void RandomAccessTest(int loop, int skip, uint64_t *timer, - int *s_buf, int *r_buf, int size, OpType type, +__global__ void RandomAccessTest(int loop, int skip, long long int *start_time, + long long int *end_time, int *s_buf, + int *r_buf, int size, OpType type, int coal_coef, int num_bins, int num_waves, uint32_t *threads_bins, uint32_t *off_bins, uint32_t *PE_bins, ShmemContextType ctx_type) { - uint64_t start; __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -69,7 +70,9 @@ __global__ void RandomAccessTest(int loop, int skip, uint64_t *timer, r_buf = r_buf + offset; for (int i = 0; i < loop + skip; i++) { - if (i == skip) start = rocshmem_timer(); + if (i == skip) { + start_time[wg_id] = wall_clock64(); + } switch (type) { case GetType: rocshmem_ctx_getmem(ctx, r_buf, s_buf, size, PE); @@ -84,8 +87,9 @@ __global__ void RandomAccessTest(int loop, int skip, uint64_t *timer, rocshmem_ctx_quiet(ctx); - atomicAdd((unsigned long long *)&timer[hipBlockIdx_x], - rocshmem_timer() - start); + // atomicAdd((unsigned long long *)&timer[hipBlockIdx_x], + // rocshmem_timer() - start); + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); rocshmem_wg_finalize(); @@ -178,9 +182,10 @@ void RandomAccessTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, if (args.myid == 0) { hipLaunchKernelGGL(RandomAccessTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, s_buf, r_buf, size, - (OpType)args.op_type, _coal_coef, _num_bins, _num_waves, - _threads_bins, _off_bins, _PE_bins, _shmem_context); + stream, loop, args.skip, start_time, end_time, s_buf, + r_buf, size, (OpType)args.op_type, _coal_coef, + _num_bins, _num_waves, _threads_bins, _off_bins, + _PE_bins, _shmem_context); } num_msgs = (loop + args.skip) * _num_waves * _thread_access; num_timed_msgs = loop * _num_waves * _thread_access; diff --git a/projects/rocshmem/tests/functional_tests/random_access_tester.hpp b/projects/rocshmem/tests/functional_tests/random_access_tester.hpp index 0c2df6ed8e..6cbcfed602 100644 --- a/projects/rocshmem/tests/functional_tests/random_access_tester.hpp +++ b/projects/rocshmem/tests/functional_tests/random_access_tester.hpp @@ -28,8 +28,9 @@ /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void RandomAccessTest(int loop, int skip, uint64_t *timer, - int *s_buf, int *r_buf, int size, OpType type, +__global__ void RandomAccessTest(int loop, int skip, long long int *start_time, + long long int *end_time, int *s_buf, + int *r_buf, int size, OpType type, int coal_coef, int num_bins, int num_waves, uint32_t *threads_bins, uint32_t *off_bins, uint32_t *PE_bins); diff --git a/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp index 45ddbd90ab..cd08cc201a 100644 --- a/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp @@ -29,20 +29,21 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void PutmemSignalTest(int loop, int skip, uint64_t *timer, char *s_buf, +__global__ void PutmemSignalTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *s_buf, char *r_buf, int size, uint64_t *sig_addr, TestType type, ShmemContextType ctx_type, int sig_op) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); - uint64_t start; uint64_t signal = 1; for (int i = 0; i < loop + skip; i++) { if (i == skip) { __syncthreads(); - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } switch (type) { @@ -74,23 +75,24 @@ __global__ void PutmemSignalTest(int loop, int skip, uint64_t *timer, char *s_bu __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); rocshmem_wg_finalize(); } -__global__ void SignalFetchTest(int loop, int skip, uint64_t *timer, uint64_t *sig_addr, +__global__ void SignalFetchTest(int loop, int skip, long long int *start_time, + long long int *end_time, uint64_t *sig_addr, uint64_t *fetched_value, TestType type) { rocshmem_wg_init(); - uint64_t start; + int wg_id = get_flat_grid_id(); for (int i = 0; i < loop + skip; i++) { if (i == skip) { __syncthreads(); - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } switch (type) { @@ -111,7 +113,7 @@ __global__ void SignalFetchTest(int loop, int skip, uint64_t *timer, uint64_t *s __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_finalize(); @@ -155,10 +157,10 @@ void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int (_type == WAVESignalFetchTestType) || (_type == WGSignalFetchTestType)) { hipLaunchKernelGGL(SignalFetchTest, gridSize, blockSize, shared_bytes, stream, - loop, args.skip, timer, sig_addr, fetched_value, _type); + loop, args.skip, start_time, end_time, sig_addr, fetched_value, _type); } else { hipLaunchKernelGGL(PutmemSignalTest, gridSize, blockSize, shared_bytes, stream, - loop, args.skip, timer, s_buf, r_buf, size, sig_addr, + loop, args.skip, start_time, end_time, s_buf, r_buf, size, sig_addr, _type, _shmem_context, sig_op); } diff --git a/projects/rocshmem/tests/functional_tests/swarm_tester.cpp b/projects/rocshmem/tests/functional_tests/swarm_tester.cpp index 5e827186cc..831febbfa0 100644 --- a/projects/rocshmem/tests/functional_tests/swarm_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/swarm_tester.cpp @@ -29,9 +29,11 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void GetSwarmTest(int loop, int skip, uint64_t *timer, char *s_buf, +__global__ void GetSwarmTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *s_buf, char *r_buf, int size, ShmemContextType ctx_type) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); int provided; rocshmem_wg_init_thread(ROCSHMEM_THREAD_MULTIPLE, &provided); @@ -42,18 +44,20 @@ __global__ void GetSwarmTest(int loop, int skip, uint64_t *timer, char *s_buf, __syncthreads(); int index = hipThreadIdx_x * size; - uint64_t start = 0; for (int i = 0; i < loop + skip; i++) { - if (i == skip) start = rocshmem_timer(); - + if (i == skip) { + start_time[wg_id] = wall_clock64(); + } rocshmem_ctx_getmem(ctx, &r_buf[index], &s_buf[index], size, 1); __syncthreads(); } - atomicAdd((unsigned long long *)&timer[hipBlockIdx_x], - rocshmem_timer() - start); + // atomicAdd((unsigned long long *)&timer[hipBlockIdx_x], + // rocshmem_timer() - start); + + end_time[wg_id] = wall_clock64(); rocshmem_wg_ctx_destroy(&ctx); rocshmem_wg_finalize(); @@ -71,8 +75,8 @@ void GetSwarmTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, size_t shared_bytes = 0; hipLaunchKernelGGL(GetSwarmTest, gridSize, blockSize, shared_bytes, stream, - loop, args.skip, timer, s_buf, r_buf, size, - _shmem_context); + loop, args.skip, start_time, end_time, s_buf, r_buf, size, + _shmem_context); num_msgs = (loop + args.skip) * gridSize.x * blockSize.x; num_timed_msgs = loop * gridSize.x * blockSize.x; diff --git a/projects/rocshmem/tests/functional_tests/swarm_tester.hpp b/projects/rocshmem/tests/functional_tests/swarm_tester.hpp index 18b6749073..e9c1bf75f7 100644 --- a/projects/rocshmem/tests/functional_tests/swarm_tester.hpp +++ b/projects/rocshmem/tests/functional_tests/swarm_tester.hpp @@ -28,7 +28,8 @@ /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void GetSwarmTest(int loop, int skip, uint64_t *timer, char *s_buf, +__global__ void GetSwarmTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *s_buf, char *r_buf, int size); /****************************************************************************** diff --git a/projects/rocshmem/tests/functional_tests/sync_tester.cpp b/projects/rocshmem/tests/functional_tests/sync_tester.cpp index d481862eaf..a74ae78d8e 100644 --- a/projects/rocshmem/tests/functional_tests/sync_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/sync_tester.cpp @@ -30,16 +30,18 @@ rocshmem_team_t team_sync_world_dup; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void SyncTest(int loop, int skip, uint64_t *timer, TestType type, +__global__ void SyncTest(int loop, int skip, long long int *start_time, + long long int *end_time, TestType type, ShmemContextType ctx_type, rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); + rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); - uint64_t start; for (int i = 0; i < loop + skip; i++) { if (hipThreadIdx_x == 0 && i == skip) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } __syncthreads(); @@ -57,7 +59,7 @@ __global__ void SyncTest(int loop, int skip, uint64_t *timer, TestType type, __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -84,7 +86,7 @@ void SyncTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, &team_sync_world_dup); hipLaunchKernelGGL(SyncTest, gridSize, blockSize, shared_bytes, stream, loop, - args.skip, timer, _type, _shmem_context, + args.skip, start_time, end_time, _type, _shmem_context, team_sync_world_dup); num_msgs = (loop + args.skip) * gridSize.x; diff --git a/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp b/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp index d9e1db44c6..68c981e797 100644 --- a/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp @@ -61,11 +61,13 @@ rocshmem_team_t team_bcast_world_dup; * DEVICE TEST KERNEL *****************************************************************************/ template -__global__ void TeamBroadcastTest(int loop, int skip, uint64_t *timer, - T1 *source_buf, T1 *dest_buf, int size, +__global__ void TeamBroadcastTest(int loop, int skip, long long int *start_time, + long long int *end_time, T1 *source_buf, + T1 *dest_buf, int size, ShmemContextType ctx_type, rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -74,10 +76,9 @@ __global__ void TeamBroadcastTest(int loop, int skip, uint64_t *timer, __syncthreads(); - uint64_t start; for (int i = 0; i < loop; i++) { if (i == skip && hipThreadIdx_x == 0) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } wg_team_broadcast(ctx, team, @@ -91,7 +92,7 @@ __global__ void TeamBroadcastTest(int loop, int skip, uint64_t *timer, __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -131,8 +132,8 @@ void TeamBroadcastTester::launchKernel(dim3 gridSize, dim3 blockSize, size_t shared_bytes = 0; hipLaunchKernelGGL(TeamBroadcastTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, source_buf, dest_buf, size, - _shmem_context, team_bcast_world_dup); + stream, loop, args.skip, start_time, end_time, source_buf, + dest_buf, size, _shmem_context, team_bcast_world_dup); num_msgs = loop + args.skip; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp b/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp index 7f04eb7c69..45c6549e9c 100644 --- a/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp @@ -31,20 +31,23 @@ rocshmem_team_t team_primitive_world_dup; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void TeamCtxPrimitiveTest(int loop, int skip, uint64_t *timer, - char *s_buf, char *r_buf, int size, - TestType type, ShmemContextType ctx_type, +__global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_time, + long long int *end_time, char *s_buf, + char *r_buf, int size, TestType type, + ShmemContextType ctx_type, rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); + rocshmem_wg_init(); rocshmem_wg_team_create_ctx(team, ctx_type, &ctx); if (hipThreadIdx_x == 0) { - uint64_t start; for (int i = 0; i < loop + skip; i++) { - if (i == skip) start = rocshmem_timer(); - + if (i == skip) { + start_time[wg_id] = wall_clock64(); + } switch (type) { case TeamCtxGetTestType: rocshmem_ctx_getmem(ctx, r_buf, s_buf, size, 1); @@ -65,7 +68,7 @@ __global__ void TeamCtxPrimitiveTest(int loop, int skip, uint64_t *timer, rocshmem_ctx_quiet(ctx); - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -104,8 +107,9 @@ void TeamCtxPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, size_t shared_bytes = 0; hipLaunchKernelGGL(TeamCtxPrimitiveTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, s_buf, r_buf, size, _type, - _shmem_context, team_primitive_world_dup); + stream, loop, args.skip, start_time, end_time, s_buf, + r_buf, size, _type, _shmem_context, + team_primitive_world_dup); num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop * gridSize.x; diff --git a/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp b/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp index 683c329ac4..17c14ae91d 100644 --- a/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp @@ -73,11 +73,13 @@ rocshmem_team_t team_reduce_world_dup; * DEVICE TEST KERNEL *****************************************************************************/ template -__global__ void TeamReductionTest(int loop, int skip, uint64_t *timer, - T1 *s_buf, T1 *r_buf, int size, TestType type, +__global__ void TeamReductionTest(int loop, int skip, long long int *start_time, + long long int *end_time, T1 *s_buf, T1 *r_buf, + int size, TestType type, ShmemContextType ctx_type, rocshmem_team_t team) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -86,10 +88,9 @@ __global__ void TeamReductionTest(int loop, int skip, uint64_t *timer, __syncthreads(); - uint64_t start; for (int i = 0; i < loop + skip; i++) { if (i == skip && hipThreadIdx_x == 0) { - start = rocshmem_timer(); + start_time[wg_id] = wall_clock64(); } wg_team_reduce(ctx, team, r_buf, s_buf, size); rocshmem_ctx_wg_barrier_all(ctx); @@ -98,7 +99,7 @@ __global__ void TeamReductionTest(int loop, int skip, uint64_t *timer, __syncthreads(); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[wg_id] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -138,9 +139,9 @@ void TeamReductionTester::launchKernel(dim3 gridSize, dim3 blockSize, size_t shared_bytes = 0; hipLaunchKernelGGL(HIP_KERNEL_NAME(TeamReductionTest), gridSize, - blockSize, shared_bytes, stream, loop, args.skip, timer, - s_buf, r_buf, size, _type, _shmem_context, - team_reduce_world_dup); + blockSize, shared_bytes, stream, loop, args.skip, + start_time, end_time, s_buf, r_buf, size, _type, + _shmem_context, team_reduce_world_dup); num_msgs = loop + args.skip; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/functional_tests/tester.cpp b/projects/rocshmem/tests/functional_tests/tester.cpp index c7f06be3e3..4a7ae9327f 100644 --- a/projects/rocshmem/tests/functional_tests/tester.cpp +++ b/projects/rocshmem/tests/functional_tests/tester.cpp @@ -58,14 +58,22 @@ Tester::Tester(TesterArguments args) : args(args) { _shmem_context = args.shmem_context; CHECK_HIP(hipGetDevice(&device_id)); CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id)); - num_warps = (args.wg_size - 1) / deviceProps.warpSize + 1; + wf_size = deviceProps.warpSize; + num_warps = (args.wg_size - 1) / wf_size + 1; CHECK_HIP(hipStreamCreate(&stream)); CHECK_HIP(hipEventCreate(&start_event)); CHECK_HIP(hipEventCreate(&stop_event)); - CHECK_HIP(hipMalloc((void**)&timer, sizeof(uint64_t) * args.num_wgs)); + CHECK_HIP(hipDeviceGetAttribute(&wall_clk_rate, + hipDeviceAttributeWallClockRate, device_id)); + num_timers = args.num_wgs; + CHECK_HIP(hipMalloc((void**)&timer, sizeof(long long int) * num_timers)); + CHECK_HIP(hipMalloc((void**)&start_time, sizeof(long long int) * num_timers)); + CHECK_HIP(hipMalloc((void**)&end_time, sizeof(long long int) * num_timers)); } Tester::~Tester() { + CHECK_HIP(hipFree(end_time)); + CHECK_HIP(hipFree(start_time)); CHECK_HIP(hipFree(timer)); CHECK_HIP(hipEventDestroy(stop_event)); CHECK_HIP(hipEventDestroy(start_event)); @@ -75,11 +83,6 @@ Tester::~Tester() { std::vector Tester::create(TesterArguments args) { int rank = args.myid; std::vector testers; - hipDeviceProp_t deviceProps; - int device_id, numWarps; - CHECK_HIP(hipGetDevice(&device_id)); - CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id)); - numWarps = (args.wg_size - 1) / deviceProps.warpSize + 1; if (rank == 0) std::cout << "### Creating Test: "; @@ -352,35 +355,23 @@ std::vector Tester::create(TesterArguments args) { testers.push_back(new ShmemPtrTester(args)); return testers; case WGGetTestType: - if (rank == 0) { - if (args.num_wgs > 1) - std::cout << "Tiled Blocking WG level Gets ###" << std::endl; - else std::cout << "Blocking WG level Gets ###" << std::endl; - } + if (rank == 0) + std::cout << "Blocking WG level Gets ###" << std::endl; testers.push_back(new ExtendedPrimitiveTester(args)); return testers; case WGGetNBITestType: - if (rank == 0) { - if (args.num_wgs > 1) - std::cout << "Tiled Non-Blocking WG level Gets ###" << std::endl; - else std::cout << "Non-Blocking WG level Gets ###" << std::endl; - } + if (rank == 0) + std::cout << "Non-Blocking WG level Gets ###" << std::endl; testers.push_back(new ExtendedPrimitiveTester(args)); return testers; case WGPutTestType: - if (rank == 0) { - if (args.num_wgs > 1) - std::cout << "Tiled Blocking WG level Puts ###" << std::endl; - else std::cout << "Blocking WG level Puts ###" << std::endl; - } + if (rank == 0) + std::cout << "Blocking WG level Puts ###" << std::endl; testers.push_back(new ExtendedPrimitiveTester(args)); return testers; case WGPutNBITestType: - if (rank == 0) { - if (args.num_wgs > 1) - std::cout << "Tiled Non-Blocking WG level Puts ###" << std::endl; - else std::cout << "Non-Blocking WG level Puts ###" << std::endl; - } + if (rank == 0) + std::cout << "Non-Blocking WG level Puts ###" << std::endl; testers.push_back(new ExtendedPrimitiveTester(args)); return testers; case PutNBIMRTestType: @@ -389,35 +380,23 @@ std::vector Tester::create(TesterArguments args) { testers.push_back(new PrimitiveMRTester(args)); return testers; case WAVEGetTestType: - if (rank == 0) { - if (args.num_wgs > 1 || numWarps > 1) - std::cout << "Tiled Blocking WAVE level Gets ###" << std::endl; - else std::cout << "Blocking WAVE level Gets ###" << std::endl; - } + if (rank == 0) + std::cout << "Blocking WAVE level Gets ###" << std::endl; testers.push_back(new WaveLevelPrimitiveTester(args)); return testers; case WAVEGetNBITestType: - if (rank == 0) { - if (args.num_wgs > 1 || numWarps > 1) - std::cout << "Tiled Non-Blocking WAVE level Gets ###" << std::endl; - else std::cout << "Non-Blocking WAVE level Gets ###" << std::endl; - } + if (rank == 0) + std::cout << "Non-Blocking WAVE level Gets ###" << std::endl; testers.push_back(new WaveLevelPrimitiveTester(args)); return testers; case WAVEPutTestType: - if (rank == 0) { - if (args.num_wgs > 1 || numWarps > 1) - std::cout << "Tiled Blocking WAVE level Puts ###" << std::endl; - else std::cout << "Blocking WAVE level Puts ###" << std::endl; - } + if (rank == 0) + std::cout << "Blocking WAVE level Puts ###" << std::endl; testers.push_back(new WaveLevelPrimitiveTester(args)); return testers; case WAVEPutNBITestType: - if (rank == 0) { - if (args.num_wgs > 1 || numWarps > 1) - std::cout << "Tiled Non-Blocking WAVE level Puts ###" << std::endl; - else std::cout << "Non-Blocking WAVE level Puts ###" << std::endl; - } + if (rank == 0) + std::cout << "Non-Blocking WAVE level Puts ###" << std::endl; testers.push_back(new WaveLevelPrimitiveTester(args)); return testers; case PutSignalTestType: @@ -501,11 +480,6 @@ void Tester::execute() { * rocshmem pes. */ if (peLaunchesKernel()) { - /** - * TODO: - * Verify that this timer type is actually uint64_t on the - * device side. - */ memset(timer, 0, sizeof(uint64_t) * args.num_wgs); const dim3 blockSize(args.wg_size, 1, 1); @@ -521,9 +495,6 @@ void Tester::execute() { if (err != hipSuccess) { printf("error = %d \n", err); } - - // rocshmem_dump_stats(); - // rocshmem_reset_stats(); } barrier(); @@ -533,32 +504,10 @@ void Tester::execute() { // data validation verifyResults(size); - /** - * Adjust size for *_wg and *_wave functions - */ - uint64_t size_ = size; - TestType type = (TestType)args.algorithm; - switch (type) { - case WAVEGetTestType: - case WAVEGetNBITestType: - case WAVEPutTestType: - case WAVEPutNBITestType: - size_ *= (args.num_wgs * num_warps); - break; - case WGGetTestType: - case WGGetNBITestType: - case WGPutTestType: - case WGPutNBITestType: - size_ *= args.num_wgs; - break; - default: - break; - } - barrier(); if (_type != TeamCtxInfraTestType) { - print(size_); + print(size); } } } @@ -589,30 +538,39 @@ void Tester::print(uint64_t size) { return; } - uint64_t timer_avg = timerAvgInMicroseconds(); - double latency_avg = static_cast(timer_avg) / num_timed_msgs; + /** + * Calculate total amount of data transfered + */ + uint64_t total_size = size * num_timed_msgs; + double timer_avg = timerAvgInMicroseconds(); + double latency_avg = timer_avg / num_timed_msgs; double avg_msg_rate = num_timed_msgs / (timer_avg / 1e6); float total_kern_time_ms; CHECK_HIP(hipEventElapsedTime(&total_kern_time_ms, start_event, stop_event)); float total_kern_time_s = total_kern_time_ms / 1000; + + double time_us = gpuCyclesToMicroseconds(max_end_time - min_start_time); + double time_s = time_us / 1e6; double bandwidth_avg_gbs = - num_msgs * size * bw_factor / total_kern_time_s / pow(2, 30); + static_cast(total_size * bw_factor) / time_s / pow(2, 30); int field_width = 20; int float_precision = 2; if (_print_header) { - printf("%-*s%*s%*s%*s", - 10, "# Size (B)", + printf("%-*s%-*s%*s%*s%*s", + 15, "# Size (B)", + 15, "# of timed Msgs", field_width, "Latency (us)", field_width, "Bandwidth (GB/s)", field_width + 1, "Msg Rate (Msg/s)\n"); _print_header = 0; } - printf("%-*lu%*.*f%*.*f%*.*f\n", - 10, size, + printf("%-*lu%-*d%*.*f%*.*f%*.*f\n", + 15, size, + 15, num_timed_msgs, field_width, float_precision, latency_avg, field_width, float_precision, bandwidth_avg_gbs, field_width, float_precision, avg_msg_rate); @@ -634,33 +592,26 @@ void Tester::barrier() { flush_hdp(); } -uint64_t Tester::gpuCyclesToMicroseconds(uint64_t cycles) { - /** - * The dGPU asm core timer runs at 27MHz. This is different from the - * core clock returned by HIP. For an APU, this is different and might - * need adjusting. - */ - uint64_t gpu_frequency_MHz = 27; - - /** - * hipDeviceGetAttribute(&gpu_frequency_khz, - * hipDeviceAttributeClockRate, - * 0); - */ - - return cycles / gpu_frequency_MHz; +double Tester::gpuCyclesToMicroseconds(long long int cycles) { + return static_cast(cycles) / + (static_cast(wall_clk_rate) * 1e-3); } -uint64_t Tester::timerAvgInMicroseconds() { - uint64_t sum = 0; +double Tester::timerAvgInMicroseconds() { + double sum = 0; + min_start_time = LLONG_MAX; + max_end_time = 0; - /** - * TODO: (bpotter/avinash) Modify the calcuation for the Tiled version of - * puts and gets at wavefront level - */ - for (uint64_t i = 0; i < args.num_wgs; i++) { + for (uint32_t i = 0; i < num_timers; i++) { + timer[i] = end_time[i] - start_time[i]; sum += gpuCyclesToMicroseconds(timer[i]); + min_start_time = (start_time[i] < min_start_time) + ? start_time[i] + : min_start_time; + max_end_time = (end_time[i] > max_end_time) + ? end_time[i] + : max_end_time; } - return sum / args.num_wgs; + return sum / num_timers; } diff --git a/projects/rocshmem/tests/functional_tests/tester.hpp b/projects/rocshmem/tests/functional_tests/tester.hpp index f506441ae3..cb5c1ab333 100644 --- a/projects/rocshmem/tests/functional_tests/tester.hpp +++ b/projects/rocshmem/tests/functional_tests/tester.hpp @@ -25,8 +25,10 @@ #include #include +#include #include "tester_arguments.hpp" +#include "../src/util.hpp" /****************************************************************************** * TESTER CLASS TYPES @@ -126,6 +128,8 @@ class Tester { int num_warps = 0; int bw_factor = 1; int device_id = 0; + int wall_clk_rate = 0; //in kilohertz + int wf_size = 0; TesterArguments args; @@ -135,7 +139,12 @@ class Tester { hipStream_t stream; hipDeviceProp_t deviceProps; - uint64_t *timer = nullptr; + long long int *timer = nullptr; + long long int *start_time = nullptr; + long long int *end_time = nullptr; + long long int min_start_time = 0; + long long int max_end_time = 0; + uint32_t num_timers = 0; private: bool _print_header = 1; @@ -143,9 +152,9 @@ class Tester { void barrier(); - uint64_t gpuCyclesToMicroseconds(uint64_t cycles); + double gpuCyclesToMicroseconds(long long int cycles); - uint64_t timerAvgInMicroseconds(); + double timerAvgInMicroseconds(); bool peLaunchesKernel(); diff --git a/projects/rocshmem/tests/functional_tests/wave_level_primitives.cpp b/projects/rocshmem/tests/functional_tests/wave_level_primitives.cpp index c434fae387..d935834d24 100644 --- a/projects/rocshmem/tests/functional_tests/wave_level_primitives.cpp +++ b/projects/rocshmem/tests/functional_tests/wave_level_primitives.cpp @@ -31,11 +31,14 @@ using namespace rocshmem; /****************************************************************************** * DEVICE TEST KERNEL *****************************************************************************/ -__global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer, - char *s_buf, char *r_buf, int size, - TestType type, ShmemContextType ctx_type, - int wf_size) { +__global__ void WaveLevelPrimitiveTest(int loop, int skip, + long long int *start_time, + long long int *end_time, char *s_buf, + char *r_buf, int size, TestType type, + ShmemContextType ctx_type, int wf_size) { __shared__ rocshmem_ctx_t ctx; + int wg_id = get_flat_grid_id(); + rocshmem_wg_init(); rocshmem_wg_ctx_create(ctx_type, &ctx); @@ -44,7 +47,6 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer, * If the number of wavefronts is greater than 1, this kernel performs a * tiled functional test */ - uint64_t start; int wf_id = get_flat_block_id() / wf_size; int wg_offset = size * get_flat_grid_id() * (get_flat_block_size() / wf_size); int idx = wf_id * size + wg_offset; @@ -52,8 +54,9 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer, r_buf += idx; for (int i = 0; i < loop + skip; i++) { - if (i == skip) start = rocshmem_timer(); - + if (i == skip) { + start_time[wg_id] = wall_clock64(); + } switch (type) { case WAVEGetTestType: rocshmem_ctx_getmem_wave(ctx, r_buf, s_buf, size, 1); @@ -75,7 +78,7 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer, rocshmem_ctx_quiet(ctx); if (hipThreadIdx_x == 0) { - timer[hipBlockIdx_x] = rocshmem_timer() - start; + end_time[hipBlockIdx_x] = wall_clock64(); } rocshmem_wg_ctx_destroy(&ctx); @@ -109,9 +112,9 @@ void WaveLevelPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, size_t shared_bytes = 0; hipLaunchKernelGGL(WaveLevelPrimitiveTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, (char*)s_buf, - (char*)r_buf, size, _type, _shmem_context, - deviceProps.warpSize); + stream, loop, args.skip, start_time, end_time, + (char*)s_buf, (char*)r_buf, size, _type, _shmem_context, + wf_size); num_msgs = (loop + args.skip) * gridSize.x * num_warps; num_timed_msgs = loop * gridSize.x * num_warps; From 6c70aee32e6f4f51a81cae541add112528dc6c81 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Mon, 17 Feb 2025 07:16:49 +0000 Subject: [PATCH 2/2] Remove `rocshmem_timer` function [ROCm/rocshmem commit: 21dbd5cc5e25d933bc1877efa945c99a8c2d489f] --- projects/rocshmem/include/rocshmem/rocshmem.hpp | 11 ----------- projects/rocshmem/src/rocshmem_gpu.cpp | 6 ------ projects/rocshmem/src/stats.hpp | 4 ++-- 3 files changed, 2 insertions(+), 19 deletions(-) diff --git a/projects/rocshmem/include/rocshmem/rocshmem.hpp b/projects/rocshmem/include/rocshmem/rocshmem.hpp index 20fa559e72..629c5a377c 100644 --- a/projects/rocshmem/include/rocshmem/rocshmem.hpp +++ b/projects/rocshmem/include/rocshmem/rocshmem.hpp @@ -512,17 +512,6 @@ __device__ ATTR_NO_INLINE void rocshmem_wg_team_sync(rocshmem_team_t team); */ __device__ ATTR_NO_INLINE void *rocshmem_ptr(const void *dest, int pe); -/** - * @brief Query the current time. Similar to gettimeofday() on the CPU. To use - * this function, rocSHMEM must be configured with profiling support - * (--enable-profile). - * - * Can be called per thread with no performance penalty. - * - * @return Time in micro-seconds. - */ -__device__ uint64_t rocshmem_timer(); - /** * @brief Make all uncacheable GPU data visible to other agents in the sytem. * diff --git a/projects/rocshmem/src/rocshmem_gpu.cpp b/projects/rocshmem/src/rocshmem_gpu.cpp index 78350deb92..5517da4b47 100644 --- a/projects/rocshmem/src/rocshmem_gpu.cpp +++ b/projects/rocshmem/src/rocshmem_gpu.cpp @@ -638,12 +638,6 @@ __device__ int rocshmem_my_pe() { return get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->my_pe; } -__device__ uint64_t rocshmem_timer() { - GPU_DPRINTF("Function: rocshmem_timer\n"); - - return __read_clock(); -} - template __device__ T rocshmem_atomic_fetch_add(rocshmem_ctx_t ctx, T *dest, T val, int pe) { diff --git a/projects/rocshmem/src/stats.hpp b/projects/rocshmem/src/stats.hpp index adc2bb51d5..a4fe1937ef 100644 --- a/projects/rocshmem/src/stats.hpp +++ b/projects/rocshmem/src/stats.hpp @@ -141,10 +141,10 @@ class Stats { StatType stats[I] = {0}; public: - __device__ uint64_t startTimer() const { return rocshmem_timer(); } + __device__ uint64_t startTimer() const { return wall_clock64(); } __device__ void endTimer(uint64_t start, int index) { - incStat(index, rocshmem_timer() - start); + incStat(index, wall_clock64() - start); } __device__ void incStat(int index, int value = 1) {