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

Remove HDP workaround for ROCm 4.2 HIP (#334)

Browse Source 
[ROCm/rccl commit: 0fbb9510a5]
This commit is contained in:
Wenkai Du
2021-03-23 20:11:37 -07:00
committed by GitHub
parent 287ed0f18a
commit 07548845a6
2 changed files with 2 additions and 40 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/transport/coll_net.cc
+1 -20
View File
@@ -83,26 +83,7 @@ ncclResult_t collNetSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph
if (resources->useGdr) {
NCCLCHECK(ncclCudaCalloc((char**)(&resources->devRecvMem), recvSize, 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;
CUDACHECK(hipDeviceGetAttribute((int*)&resources->curr_hdp_reg, hipDeviceAttributeHdpMemFlushCntl, myInfo->cudaDev));
send->conn.curr_hdp_reg = resources->curr_hdp_reg;
}
NCCLCHECK(ncclCudaHostCalloc((char**)&resources->recvMem, recvSize));
projects/rccl/src/transport/net.cc
+1 -20
View File
@@ -112,26 +112,7 @@ ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st
}
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;
CUDACHECK(hipDeviceGetAttribute((int*)&resources->curr_hdp_reg, hipDeviceAttributeHdpMemFlushCntl, myInfo->cudaDev));
send->conn.curr_hdp_reg = resources->curr_hdp_reg;
}
INFO(NCCL_INIT|NCCL_NET,"Channel %02d : %d[%lx] -> %d[%lx] [send] via NET/%s/%d%s%s", channelId, myInfo->rank, myInfo->busId, peerInfo->rank, peerInfo->busId, ncclNetName(), resources->netDev,