From 13922baf76c26c7808d5962717b2a57daef5cdc9 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Thu, 27 Jul 2023 09:04:35 -0700
Subject: [PATCH] ll_latency_test: fix time calculation (#825)

* ll_latency_test: fix time calculation

* Measure time after barrier

* Read time stamp only from thread 0

[ROCm/rccl commit: c424979c143031b16c62ad26ce2b75fc6661422c]
---
 .../tools/p2p-latency-test/ll_latency_test.cpp   | 16 ++++++++++------
 1 file changed, 10 insertions(+), 6 deletions(-)

diff --git a/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp b/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp
index 32565c63b6..0f2d5aa987 100644
--- a/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp
+++ b/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp
@@ -84,14 +84,16 @@ __global__ void PingKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint
     storeLL(remote_flag+tid, i, i);
     while (readLL(local_flag+tid, i) != i);
   }
-  uint64_t start_time = wall_clock64();
+  uint64_t start_time, end_time;
+  if (tid == 0) start_time = wall_clock64();
   #pragma unroll
   for (uint32_t i = NUM_LOOPS_WARMUP; i <= NUM_LOOPS_WARMUP + NUM_LOOPS_RUN; i++) {
     storeLL(remote_flag+tid, i, i);
     while (readLL(local_flag+tid, i) != i);
   }
-  uint64_t end_time = wall_clock64();
-  *time_delta = end_time - start_time;
+  __syncthreads();
+  if (tid == 0) end_time = wall_clock64();
+  if (tid == 0) *time_delta = end_time - start_time;
 }
 
 __global__ void PongKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint64_t* time_delta) {
@@ -101,14 +103,16 @@ __global__ void PongKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint
     while (readLL(local_flag+tid, i) != i);
     storeLL(remote_flag+tid, i, i);
   }
-  uint64_t start_time = wall_clock64();
+  uint64_t start_time, end_time;
+  if (tid == 0) start_time = wall_clock64();
   #pragma unroll
   for (uint32_t i = NUM_LOOPS_WARMUP; i <= NUM_LOOPS_WARMUP + NUM_LOOPS_RUN; i++) {
     while (readLL(local_flag+tid, i) != i);
     storeLL(remote_flag+tid, i, i);
   }
-  uint64_t end_time = wall_clock64();
-  *time_delta = end_time - start_time;
+  __syncthreads();
+  if (tid == 0) end_time = wall_clock64();
+  if (tid == 0) *time_delta = end_time - start_time;
 }
 
 int main(int argc, char** argv) {