Add another Rome model (#249)

* Add another Rome model

* Add gfx908 4P3L models and support

* Revert "Use cached value for detecting GDR support only once"

This reverts commit 0108a1219d.

* Skip using ibverb for GPU direct RDMA detection

* Fine tune one Rome model

[ROCm/rccl commit: a51e4071e3]

projects/rccl/src/include/net.h

@@ -36,14 +36,6 @@ static ncclResult_t ncclNetCloseListen(void* listenComm) { NCCLCHECK(ncclNet->cl
 static ncclResult_t ncclGpuGdrSupport(int* gdrSupport) {
   int netDevs;
   NCCLCHECK(ncclNetDevices(&netDevs));
-  pthread_mutex_t ncclParamMutexGpuGdrSupport = PTHREAD_MUTEX_INITIALIZER;
-  static int gdrSupportCached[16] = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
-  int cudaDev;
-  CUDACHECK(hipGetDevice(&cudaDev));
-  if (gdrSupportCached[cudaDev] != -1) {
-    *gdrSupport = gdrSupportCached[cudaDev];
-    return ncclSuccess;
-  }
   *gdrSupport = 0;
   for (int dev=0; dev<netDevs; dev++) {
     // Find a net device which is GDR-capable
@@ -52,6 +44,8 @@ static ncclResult_t ncclGpuGdrSupport(int* gdrSupport) {
     if ((props.ptrSupport & NCCL_PTR_CUDA) == 0) continue;
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
     if (!hasFineGrainVramPcie()) continue;
+    *gdrSupport = 1;
+    break;
 #endif
 
     // Allocate memory on the GPU and try to register it on the NIC.
@@ -59,7 +53,6 @@ static ncclResult_t ncclGpuGdrSupport(int* gdrSupport) {
     ncclNetHandle_t handle;
     void* gpuPtr = NULL;
     void* mHandle = NULL;
-    pthread_mutex_lock(&ncclParamMutexGpuGdrSupport);
     NCCLCHECK(ncclNetListen(dev, &handle, &lComm));
     NCCLCHECK(ncclNetConnect(dev, &handle, &sComm));
     NCCLCHECK(ncclNetAccept(lComm, &rComm));
@@ -76,10 +69,8 @@ static ncclResult_t ncclGpuGdrSupport(int* gdrSupport) {
     NCCLCHECK(ncclNetCloseRecv(rComm));
     NCCLCHECK(ncclNetCloseSend(sComm));
     NCCLCHECK(ncclNetCloseListen(lComm));
-    pthread_mutex_unlock(&ncclParamMutexGpuGdrSupport);
     break;
   }
-  gdrSupportCached[cudaDev] = *gdrSupport;
   return ncclSuccess;
 }