Update bandwidth and latency calculations

- Refined bandwidth and latency calculations for improved accuracy


[ROCm/rocshmem commit: c155636da4]

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;