AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Revert "Ensure memory copy integrity during transport setup (#731)" (#741)

Browse Source 
* Revert "Ensure memory copy integrity during transport setup (#731)"

This reverts commit 0c0c35927a.

Add stream synchronization in ncclStrongStreamRelease.

* Use event record and wait

[ROCm/rccl commit: 4ca7742c61]
This commit is contained in:
Wenkai Du
2023-05-16 10:34:47 -07:00
committed by GitHub
parent a26ca6d766
commit f7184aef3f
2 changed files with 4 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/misc/strongstream.cc
+2
View File
@@ -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;
}
projects/rccl/src/transport.cc
+2 -15
View File
@@ -10,7 +10,6 @@
#include "bootstrap.h"
#define ENABLE_TIMER 0
#include "timer.h"
#include <cstring>
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;
}