Fix DMABUF support (#1218)

* Fix DMABUF support

* Reduce log output by moving dmabuf allocation details to TRACE

* Enable peer memory GDR support if ib_umem_get_peer is in kernel

src/init.cc

@@ -722,7 +722,11 @@ static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, u
 #endif
     CUDACHECK(hipFree(ptr));
     info->hasFineGrain = true;
-    NCCLCHECK(ncclGpuGdrSupport(comm, &info->gdrSupport));
+    // GPU supports GDR if DMABUF is supported
+    if (dmaBufSupported(comm) == ncclSuccess)
+      info->gdrSupport = 1;
+    else
+      NCCLCHECK(ncclGpuGdrSupport(comm, &info->gdrSupport));
   }
   else {
     info->hasFineGrain = false;

src/transport/net.cc

@@ -837,7 +837,7 @@ static ncclResult_t sendProxyConnect(struct ncclProxyConnection* connection, str
         CUCHECK(hsa_amd_portable_export_dmabuf((const void*)resources->buffers[p], resources->buffSizes[p], &dmabuf_fd, &offset));
         NCCLCHECK(proxyState->ncclNet->regMrDmaBuf(resources->netSendComm, resources->buffers[p], resources->buffSizes[p], type, offset, dmabuf_fd, &resources->mhandles[p]));
         (void)close(dmabuf_fd);
-        INFO(NCCL_INIT|NCCL_NET, "hsa_amd_portable_export_dmabuf buffer %p size %d handle %x offset %ld",
+        TRACE(NCCL_INIT|NCCL_NET, "hsa_amd_portable_export_dmabuf buffer %p size %d handle %x offset %ld",
           (const void*)resources->buffers[p], resources->buffSizes[p], dmabuf_fd, offset);
       } else // FALL-THROUGH to nv_peermem GDR path
 #endif
@@ -1003,7 +1003,7 @@ static ncclResult_t recvProxyConnect(struct ncclProxyConnection* connection, str
         CUCHECK(hsa_amd_portable_export_dmabuf((const void*)resources->buffers[p], resources->buffSizes[p], &dmabuf_fd, &offset));
         NCCLCHECK(proxyState->ncclNet->regMrDmaBuf(resources->netRecvComm, resources->buffers[p], resources->buffSizes[p], type, offset, dmabuf_fd, &resources->mhandles[p]));
         (void)close(dmabuf_fd);
-        INFO(NCCL_INIT|NCCL_NET, "hsa_amd_portable_export_dmabuf buffer %p size %d handle %x offset %ld",
+        TRACE(NCCL_INIT|NCCL_NET, "hsa_amd_portable_export_dmabuf buffer %p size %d handle %x offset %ld",
           (const void*)resources->buffers[p], resources->buffSizes[p], dmabuf_fd, offset);
       } else // FALL-THROUGH to nv_peermem GDR path
 #endif

src/transport/net_ib.cc

@@ -23,6 +23,7 @@
 #include <unistd.h>
 #define ENABLE_TIMER 0
 #include "timer.h"
+#include <sys/utsname.h>
 
 #include "ibvwrap.h"
 #include "graph/xml.h"
@@ -361,6 +362,31 @@ ncclResult_t ncclIbGdrSupport() {
       NCCLCHECK(ncclTopoGetStrFromSys("/proc/sys/kernel", "numa_balancing", strValue));
       if (strcmp(strValue, "1") == 0 && roMode == 0)
         moduleLoaded = 0;
+    } else {
+      char kernel_header_file[256];
+      struct utsname utsname;
+      moduleLoaded = 0;
+      char buf[256];
+      FILE *fp = NULL;
+      //check for kernel name exists
+      if (uname(&utsname) == -1) {
+        INFO(NCCL_NET,"Could not get kernel name");
+      } else {
+        //format and store the kernel conf file location
+        snprintf(kernel_header_file, sizeof(kernel_header_file), "/lib/modules/%s/build/include/rdma/ib_umem.h", utsname.release);
+        fp = fopen(kernel_header_file, "r");
+        if (fp == NULL) {
+          INFO(NCCL_INIT,"Could not open kernel header file %s", kernel_header_file);
+        } else {
+          //look for kernel_opt1 and kernel_opt2 in the conf file and check
+          while (fgets(buf, sizeof(buf), fp) != NULL) {
+            if (strstr(buf, "ib_umem_get_peer") != NULL) {
+              moduleLoaded = 1;
+              INFO(NCCL_INIT,"Found ib_umem_get_peer in %s", kernel_header_file);
+            }
+          }
+        }
+      }
     }
 #else
     // Check for the nv_peer_mem module being loaded