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: c424979c14]

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) {