From f7184aef3fb8b91d69d88b0cac62c0129670e960 Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Tue, 16 May 2023 10:34:47 -0700 Subject: [PATCH] Revert "Ensure memory copy integrity during transport setup (#731)" (#741) * Revert "Ensure memory copy integrity during transport setup (#731)" This reverts commit 0c0c35927a9573bb94021d664181d094e2d5de5d. Add stream synchronization in ncclStrongStreamRelease. * Use event record and wait [ROCm/rccl commit: 4ca7742c61829f96fb19eec702f93a4493a823a1] --- projects/rccl/src/misc/strongstream.cc | 2 ++ projects/rccl/src/transport.cc | 17 ++--------------- 2 files changed, 4 insertions(+), 15 deletions(-) diff --git a/projects/rccl/src/misc/strongstream.cc b/projects/rccl/src/misc/strongstream.cc index faeec4bca3..27186cc8ac 100644 --- a/projects/rccl/src/misc/strongstream.cc +++ b/projects/rccl/src/misc/strongstream.cc @@ -237,6 +237,8 @@ ncclResult_t ncclStrongStreamRelease(struct ncclCudaGraph graph, struct ncclStro } } #endif + CUDACHECK(cudaEventRecord(ss->scratchEvent, ss->cudaStream)); + CUDACHECK(cudaStreamWaitEvent(ss->cudaStream, ss->scratchEvent, 0)); return ncclSuccess; } diff --git a/projects/rccl/src/transport.cc b/projects/rccl/src/transport.cc index 7a2502bf3e..f1c30fa01b 100644 --- a/projects/rccl/src/transport.cc +++ b/projects/rccl/src/transport.cc @@ -10,7 +10,6 @@ #include "bootstrap.h" #define ENABLE_TIMER 0 #include "timer.h" -#include struct ncclTransport* ncclTransports[NTRANSPORTS] = { &p2pTransport, @@ -161,13 +160,7 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* NCCLCHECKGOTO(conn->transportComm->connect(comm, sendData[i] + sendDataOffset++, 1, comm->rank, conn), ret, fail); if (ret == ncclSuccess) { conn->connected = 1; - do { - struct ncclConnInfo connInfo; - CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->hostStream.cudaStream), ret, fail); - CUDACHECKGOTO(cudaMemcpyAsync(&connInfo, &comm->channels[c].devPeers[sendPeer].send[connIndex], sizeof(struct ncclConnInfo), cudaMemcpyDeviceToHost, comm->hostStream.cudaStream), ret, fail); - CUDACHECKGOTO(hipStreamSynchronize(comm->hostStream.cudaStream), ret, fail); - if (memcmp(&connInfo, &conn->conn, sizeof(struct ncclConnInfo)) == 0) break; - } while (true); + CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->hostStream.cudaStream), ret, fail); } else if (ret == ncclInProgress) { allChannelsConnected = false; } @@ -184,13 +177,7 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* NCCLCHECKGOTO(conn->transportComm->connect(comm, recvData[i] + recvDataOffset++, 1, comm->rank, conn), ret, fail); if (ret == ncclSuccess) { conn->connected = 1; - do { - struct ncclConnInfo connInfo; - CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeers[recvPeer].recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->hostStream.cudaStream), ret, fail); - CUDACHECKGOTO(cudaMemcpyAsync(&connInfo, &comm->channels[c].devPeers[recvPeer].recv[connIndex], sizeof(struct ncclConnInfo), cudaMemcpyDeviceToHost, comm->hostStream.cudaStream), ret, fail); - CUDACHECKGOTO(hipStreamSynchronize(comm->hostStream.cudaStream), ret, fail); - if (memcmp(&connInfo, &conn->conn, sizeof(struct ncclConnInfo)) == 0) break; - } while (true); + CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeers[recvPeer].recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->hostStream.cudaStream), ret, fail); } else if (ret == ncclInProgress) { allChannelsConnected = false; }