From c158d3a9b40b17dcd1e5ebf7e0304f7af3ad7789 Mon Sep 17 00:00:00 2001
From: corey-derochie-amd
 <161367113+corey-derochie-amd@users.noreply.github.com>
Date: Fri, 3 Jan 2025 13:06:47 -0700
Subject: [PATCH] [SWDEV-497665] Blocked `cudaMemcpyAsync` race condition by
 synchronizing (#1447)

* Switched calls to `cudaMemcpyAsync` to be `cudaMemcpy` in `ncclTransportP2pSetup` to avoid race condition with `cudaIpcOpenMemHandle` inside p2p `connect`. See `ncclP2pImportShareableBuffer`.

* Moved synchronize outside of the loop, as it isn't necessary to sync between every iteration of the loop.
---
 src/transport.cc | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/src/transport.cc b/src/transport.cc
index 916fcf53d9..20aa1ad2be 100644
--- a/src/transport.cc
+++ b/src/transport.cc
@@ -263,6 +263,8 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
     }
   }
 
+  CUDACHECKGOTO(cudaStreamSynchronize(comm->sharedRes->hostStream.cudaStream), ret, fail);
+
   if (timeReported) {
     struct timeval now;
     gettimeofday(&now, NULL);