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

Move HDP flush to CPU

Browse Source
This commit is contained in:
Wenkai Du
2021-02-12 18:06:19 +00:00
parent f1a9ce3fa5
commit bf8eb40705
3 changed files with 7 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/collectives/device/primitives.h
-1
View File
@@ -158,7 +158,6 @@ class ncclPrimitives {
inline __device__ void postSend() {
if (conn->next_hdp_reg) STORE(conn->next_hdp_reg, 0x1);
if (conn->curr_hdp_reg) STORE(conn->curr_hdp_reg, 0x1);
STORE(connTailPtr, step += SLICESTEPS);
}
src/include/proxy.h
+1
View File
@@ -36,6 +36,7 @@ struct ncclProxyArgs {
uint64_t transmitted;
uint64_t done;
uint64_t end;
uint64_t hdp_flushed;
void* requests[NCCL_STEPS];
int idle;
src/transport/net.cc
+6 -1
View File
@@ -283,7 +283,7 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
if (args->state == ncclProxyOpReady) {
// Round to next multiple of sliceSteps
resources->step = ROUNDUP(resources->step, args->chunkSteps);
args->posted = args->transmitted = args->done = resources->step;
args->posted = args->transmitted = args->done = args->hdp_flushed = resources->step;
args->end = resources->step + args->nsteps;
args->state = ncclProxyOpProgress;
}
@@ -345,6 +345,11 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
}
}
if (ready) {
// flush HDP if not done
if (resources->curr_hdp_reg && args->hdp_flushed < LOAD(recvTail)) {
args->hdp_flushed = LOAD(recvTail);
STORE(resources->curr_hdp_reg, 1);
}
// Data is ready, try to send.
NCCLCHECK(ncclNetIsend(resources->netSendComm, buff, size, mhandle, args->requests+buffSlot));
if (args->requests[buffSlot] != NULL) {