From 6ac3c4c90ff627dec0d7fb8ac1ae76458f595b66 Mon Sep 17 00:00:00 2001
From: Changho Hwang <changhohwang@microsoft.com>
Date: Thu, 9 Mar 2023 07:24:09 +0000
Subject: [PATCH] Relaxed sync

---
 tests/p2p_test.cu | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/tests/p2p_test.cu b/tests/p2p_test.cu
index eb962cc0..c6528ef7 100644
--- a/tests/p2p_test.cu
+++ b/tests/p2p_test.cu
@@ -67,13 +67,13 @@ __global__ void kernel(int rank, int world_size)
   // Each warp receives data from different ranks
 #if (USE_DMA_FOR_P2P == 1)
 
+  // Wait until the proxy have sent my data and flag
+  while (*trig != 0) {}
+
   // Trigger sending data and flag
   uint64_t dataOffset = rank * sizeof(int);
   uint64_t dataSize = sizeof(int);
-  *trig = TRIGGER_VALUE(mscclppSync | mscclppFlag | mscclppData, dataOffset, dataSize);
-
-  // Wait until the proxy have sent my data and flag
-  while (*trig != 0) {}
+  *trig = TRIGGER_VALUE(mscclppFlag | mscclppData, dataOffset, dataSize);
 
   // Wait for receiving data from remote rank
   while (*proxyFlag == baseFlag) {}