Enable GPU direct RDMA read from GPU

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);
   }
 

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

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 {

src/transport/net.cc

@@ -10,6 +10,7 @@
 #include "graph.h"
 #include <sys/time.h>
 #include "collectives.h"
+#include <hsa/hsa_ext_amd.h>
 
 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;