From 95425625261e37b7d5b3480f64170114460a2f27 Mon Sep 17 00:00:00 2001
From: "hongwei.liu" <hongwei.liu@mthreads.com>
Date: Fri, 24 Jul 2026 00:15:02 +0800
Subject: [PATCH] Fix ping-pong flag reset ordering

Complete each local flag reset before entering a global barrier. This
prevents a slower PE from overwriting the peer's first signal after the
timed kernel launches begin.

Signed-off-by: hongwei.liu <hongwei.liu@mthreads.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
---
 perftest/device/pt-to-pt/shmem_p_ping_pong_latency.cu | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/perftest/device/pt-to-pt/shmem_p_ping_pong_latency.cu b/perftest/device/pt-to-pt/shmem_p_ping_pong_latency.cu
index a63d5d09..c1f9f424 100644
--- a/perftest/device/pt-to-pt/shmem_p_ping_pong_latency.cu
+++ b/perftest/device/pt-to-pt/shmem_p_ping_pong_latency.cu
@@ -168,6 +168,8 @@ int main(int argc, char *argv[]) {
             } else {
                 CUDA_CHECK(cudaMemset(flag_d, 0, sizeof(uint64_t)));
             }
+            CUDA_CHECK(cudaDeviceSynchronize());
+            nvshmem_barrier_all();
             cudaEventRecord(start);
             test_ping_pong(args_2, test_cubin, 0);
             if (status != NVSHMEMX_SUCCESS) {