From 07548845a6eca9ca2b180e1c9c97fa544d7a1e93 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Tue, 23 Mar 2021 20:11:37 -0700
Subject: [PATCH] Remove HDP workaround for ROCm 4.2 HIP (#334)

[ROCm/rccl commit: 0fbb9510a555af2bf92117d93066266fb09e5969]
---
 projects/rccl/src/transport/coll_net.cc | 21 +--------------------
 projects/rccl/src/transport/net.cc      | 21 +--------------------
 2 files changed, 2 insertions(+), 40 deletions(-)

diff --git a/projects/rccl/src/transport/coll_net.cc b/projects/rccl/src/transport/coll_net.cc
index e23873b65c..9399ba2af1 100644
--- a/projects/rccl/src/transport/coll_net.cc
+++ b/projects/rccl/src/transport/coll_net.cc
@@ -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));
diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc
index 4fe4423d43..ff9480d688 100644
--- a/projects/rccl/src/transport/net.cc
+++ b/projects/rccl/src/transport/net.cc
@@ -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,