diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h index 5ffb02601c..a5e52aca7a 100644 --- a/src/collectives/device/primitives.h +++ b/src/collectives/device/primitives.h @@ -158,6 +158,7 @@ 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); } diff --git a/src/graph/paths.cc b/src/graph/paths.cc index 567043d4f6..c9924229be 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -316,9 +316,6 @@ ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* system, int64_t busId, int if (read) { // For reads (sends) only enable under certain conditions int gdrReadParam = ncclParamNetGdrRead(); if (gdrReadParam == 0) return ncclSuccess; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) - return ncclSuccess; -#else if (gdrReadParam < 0) { int nvlink = 0; // Since we don't know whether there are other communicators, @@ -333,7 +330,6 @@ ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* system, int64_t busId, int } if (!nvlink) return ncclSuccess; } -#endif } // Check if we are close enough that it makes sense to enable GDR diff --git a/src/include/devcomm.h b/src/include/devcomm.h index 065d7590e1..bbccb3ac84 100644 --- a/src/include/devcomm.h +++ b/src/include/devcomm.h @@ -115,6 +115,7 @@ struct ncclConnInfo { // allows software to explicitly initiate a flush read to HDP memory. See more // descriptions in primitives.h. uint32_t* next_hdp_reg; // Next GPU in ring (for p2p transport use only) + uint32_t* curr_hdp_reg; // Current GPU's HDP register }; struct ncclConnector { diff --git a/src/transport/net.cc b/src/transport/net.cc index e4178366cf..1f04d3ef13 100644 --- a/src/transport/net.cc +++ b/src/transport/net.cc @@ -10,6 +10,7 @@ #include "graph.h" #include #include "collectives.h" +#include struct netConnectInfo { ncclNetHandle_t netHandle; @@ -32,6 +33,7 @@ struct netSendResources { void** mhandlesProto[NCCL_NUM_PROTOCOLS]; uint64_t step; uint64_t llLastCleaning; + uint32_t* curr_hdp_reg; // Curr GPU in ring (for rdma transport use only) }; struct netRecvResources { @@ -99,6 +101,29 @@ ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st if (resources->buffSizes[LOC_HOSTMEM]) { NCCLCHECK(ncclCudaHostCalloc(resources->buffers+LOC_HOSTMEM, resources->buffSizes[LOC_HOSTMEM])); } + if (resources->useGdr) { + //CUDACHECK(hipDeviceGetAttribute((int*)&resources->curr_hdp_reg, hipDeviceAttributeHdpMemFlushCntl, myInfo->cudaDev)); + struct data_struct {hsa_agent_t agent; int counter;} out; + out.counter = 0; + out.agent.handle = myInfo->cudaDev; + hsa_iterate_agents([](hsa_agent_t agent, void* data) { + int devId = ((struct data_struct *)data)->agent.handle; + hsa_device_type_t type; + hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); + if(type != HSA_DEVICE_TYPE_GPU) + return HSA_STATUS_SUCCESS; + if(((struct data_struct *)data)->counter!=devId) { + ((struct data_struct *)data)->counter++; + return HSA_STATUS_SUCCESS; + } + ((struct data_struct *)data)->agent = agent; + return HSA_STATUS_SUCCESS; + }, (void*)&out); + hsa_amd_hdp_flush_t hdpinfo; + hsa_status_t err = hsa_agent_get_info(out.agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_HDP_FLUSH, &hdpinfo); + resources->curr_hdp_reg = hdpinfo.HDP_MEM_FLUSH_CNTL; + send->conn.curr_hdp_reg = resources->curr_hdp_reg; + } int offsets[LOC_COUNT]; offsets[LOC_HOSTMEM] = offsets[LOC_DEVMEM] = 0;