From 6e021f1adcec791c026119450bf597c43eaabfce Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Wed, 21 Sep 2022 14:17:35 +0530
Subject: [PATCH] SWDEV-356299 - Fix hipPeerToPeer_simple Tests (#2953)

- Add hipDeviceSynchronize after hipMemset with device memory, as hipMemset is asynchronous in this case

Change-Id: Ia55d106931984fed450be29c1b745c8069151f76
---
 tests/src/p2p/hipPeerToPeer_simple.cpp | 11 ++++++++++-
 1 file changed, 10 insertions(+), 1 deletion(-)

diff --git a/tests/src/p2p/hipPeerToPeer_simple.cpp b/tests/src/p2p/hipPeerToPeer_simple.cpp
index f2f3592df2..1f86fe4eb5 100644
--- a/tests/src/p2p/hipPeerToPeer_simple.cpp
+++ b/tests/src/p2p/hipPeerToPeer_simple.cpp
@@ -159,12 +159,15 @@ void enablePeerFirst(bool useAsyncCopy) {
     HIPCHECK(hipSetDevice(g_currentDevice));
     HIPCHECK(hipMalloc(&A_d0, Nbytes));
     HIPCHECK(hipMemset(A_d0, memsetval, Nbytes));
+    // hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
+    HIPCHECK(hipDeviceSynchronize());
 
     // allocate and initialize memory on peer device
     HIPCHECK(hipSetDevice(g_peerDevice));
     HIPCHECK(hipMalloc(&A_d1, Nbytes));
     HIPCHECK(hipMemset(A_d1, 0x13, Nbytes));
-
+    // hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
+    HIPCHECK(hipDeviceSynchronize());
 
     // Device0 push to device1, using P2P:
     // NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for
@@ -216,11 +219,15 @@ void allocMemoryFirst(bool useAsyncCopy) {
     HIPCHECK(hipSetDevice(g_currentDevice));
     HIPCHECK(hipMalloc(&A_d0, Nbytes));
     HIPCHECK(hipMemset(A_d0, memsetval, Nbytes));
+    // hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
+    HIPCHECK(hipDeviceSynchronize());
 
     // allocate and initialize memory on peer device
     HIPCHECK(hipSetDevice(g_peerDevice));
     HIPCHECK(hipMalloc(&A_d1, Nbytes));
     HIPCHECK(hipMemset(A_d1, 0x13, Nbytes));
+    // hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
+    HIPCHECK(hipDeviceSynchronize());
 
 
     //---
@@ -305,6 +312,8 @@ void testPeerHostToDevice(bool useAsyncCopy) {
     HIPCHECK(hipSetDevice(g_peerDevice));
     HIPCHECK(hipMalloc(&A_d1, Nbytes));
     HIPCHECK(hipMemset(A_d1, 0x13, Nbytes));
+    // hipDeviceSynchronize as hipMemset is asynchronous when destination memory is device memory
+    HIPCHECK(hipDeviceSynchronize());
 
     bool firstAsyncCopy = useAsyncCopy; /*TODO - should be useAsyncCopy*/