Determine fine grained memory availability at RCCL bootstrapping (#471)

src/include/alloc.h

@@ -89,16 +89,6 @@ static ncclResult_t ncclCudaMemcpy(T* dst, T* src, size_t nelem) {
   return ncclSuccess;
 }
 
-static bool hasFineGrainVramPcie() {
-  int *ptr;
-  if (hipExtMallocWithFlags((void**)&ptr, sizeof(int), hipDeviceMallocFinegrained) == hipSuccess) {
-    CUDACHECK(hipFree(ptr));
-    return true;
-  }
-  else
-    return false;
-}
-
 // Allocate memory to be potentially ibv_reg_mr'd. This needs to be
 // allocated on separate pages as those pages will be marked DONTFORK
 // and if they are shared, that could cause a crash in a child process

src/include/net.h

@@ -43,7 +43,6 @@ static ncclResult_t ncclGpuGdrSupport(int* gdrSupport) {
     NCCLCHECK(ncclNet->getProperties(dev, &props));
     if ((props.ptrSupport & NCCL_PTR_CUDA) == 0) continue;
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
-    if (!hasFineGrainVramPcie()) continue;
     *gdrSupport = 1;
     break;
 #endif

src/include/transport.h

@@ -29,6 +29,7 @@ struct ncclPeerInfo {
   int rank;
   int cudaDev;
   int gdrSupport;
+  bool hasFineGrain;
   uint64_t hostHash;
   uint64_t pidHash;
   dev_t shmDev;