[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.
Этот коммит содержится в:
corey-derochie-amd
2025-01-03 13:06:47 -07:00
коммит произвёл GitHub
родитель a9d6e7661c
Коммит c158d3a9b4
+2
Просмотреть файл
@@ -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);