Add ROCSHMEM_GDA_PROVIDER envvar (#280)

* Add ROCHSMEM_GDA_PROVIDER env control

* Single name for a single concept: vendor->provider

src/envvar.cpp

@@ -65,6 +65,7 @@ namespace envvar {
   }  // namespace ro
 
   namespace gda {
+    const var<std::string> provider("PROVIDER", "");
     const var<bool> alternate_qp_ports("ALTERNATE_QP_PORTS", "", true);
     const var<uint8_t> traffic_class("TRAFFIC_CLASS", "", 0);
   }  // namespace gda

src/envvar.hpp

@@ -448,6 +448,7 @@ namespace envvar {
 
   namespace gda {
     template <typename T> using var = var<T, category::tag::GDA>;
+    extern const var<std::string> provider;
     extern const var<bool> alternate_qp_ports;
     extern const var<uint8_t> traffic_class;
   }  // namespace gda

src/gda/backend_gda.cpp

@@ -75,7 +75,7 @@ void GDABackend::init() {
 
   type = BackendType::GDA_BACKEND;
 
-  read_env();
+  select_nic();
 
   //TODO setup_host_interface();
   /* Initialize the host interface */
@@ -125,7 +125,7 @@ GDABackend::~GDABackend() {
   close_dv_libs();
 }
 
-void GDABackend::read_env() {
+void GDABackend::select_nic() {
   if (!envvar::requested_dev.is_default()) {
     requested_dev = envvar::requested_dev.get_value().c_str();
   } else {
@@ -528,37 +528,63 @@ void GDABackend::rte_barrier() {
   }
 }
 
+GDAProvider GDABackend::requested_provider() {
+  /* Check whether the user explicitely requests a particular provider type */
+  std::string envstr = envvar::gda::provider;
+  std::transform(envstr.begin(), envstr.end(), envstr.begin(), ::tolower);
+  if (!envstr.empty()) {
+    printf("Found environment variable ROCSHMEM_GDA_PROVIDER, value is %s\n", envstr.c_str());
+    if (envstr.find("bnxt") != std::string::npos) {
+      return GDAProvider::BNXT;
+    }
+    if (envstr.find("ionic") != std::string::npos) {
+      return GDAProvider::IONIC;
+    }
+    if (envstr.find("mlx5") != std::string::npos) {
+      return GDAProvider::MLX5;
+    }
+  }
+  return GDAProvider::UNSET;
+}
+
 /* Currently we only check whether we can dlopen a Direct Verbs library.
  * We might need to extend this logic to check whether we have interfaces that
  * can use those DV libraries
  */
 int GDABackend::backend_can_run() {
   void *handle{nullptr};
+  GDAProvider requested = requested_provider();
 
   /* Try opening bnxt DV libraries */
 #if defined(GDA_BNXT)
-  handle = bnxt_dv_dlopen();
-  if (handle) {
-    dlclose(handle);
-    return ROCSHMEM_SUCCESS;
+  if (requested == GDAProvider::UNSET || requested == GDAProvider::BNXT) {
+    handle = bnxt_dv_dlopen();
+    if (handle) {
+      dlclose(handle);
+      return ROCSHMEM_SUCCESS;
+    }
   }
 #endif //defined(GDA_BNXT)
 
   /* Try opening ionic DV libraries */
 #if defined(GDA_IONIC)
-  handle = ionic_dv_dlopen();
-  if (handle) {
-    dlclose(handle);
-    return ROCSHMEM_SUCCESS;
+  if (requested == GDAProvider::UNSET || requested == GDAProvider::IONIC) {
+    handle = ionic_dv_dlopen();
+    if (handle) {
+      dlclose(handle);
+      return ROCSHMEM_SUCCESS;
+    }
   }
 #endif //defined(GDA_IONIC)
 
   /* Try opening mlx5 DV libraries */
 #if defined(GDA_MLX5)
-  handle = mlx5_dv_dlopen();
-  if (handle) {
-    dlclose(handle);
-    return ROCSHMEM_SUCCESS;
+  if (requested == GDAProvider::UNSET || requested == GDAProvider::MLX5) {
+    handle = mlx5_dv_dlopen();
+    if (handle) {
+      dlclose(handle);
+      return ROCSHMEM_SUCCESS;
+    }
   }
 #endif //defined(GDA_MLX5)
 
@@ -584,7 +610,7 @@ void GDABackend::setup_ibv() {
 void GDABackend::cleanup_ibv() {
   int err;
 
-  if (gda_vendor == GDAVendor::BNXT) {
+  if (gda_provider == GDAProvider::BNXT) {
     CHECK_HIP(hipHostUnregister(db_region_attr.dbr));
 
     for (int i = 0; i < qps.size(); i++) {
@@ -617,7 +643,7 @@ void GDABackend::cleanup_ibv() {
       CHECK_ZERO(err, "ibv_destroy_cqs");
     }
 
-    if (gda_vendor == GDAVendor::IONIC) {
+    if (gda_provider == GDAProvider::IONIC) {
       err = ibv_dealloc_pd(pd_uxdma[0]);
       CHECK_ZERO(err, "ibv_dealloc_pd (uxdma[0])");
 
@@ -636,19 +662,21 @@ void GDABackend::cleanup_ibv() {
   CHECK_ZERO(err, "ibv_close_device");
 }
 
+
 void GDABackend::open_dv_libs() {
   int ret;
+  GDAProvider requested = requested_provider();
 
-  //TODO: environment variable selection/deselection
   //this hardcoded init order will always prefer BNXT>IONIC>MLX5
-  //if all three drivers are installed
+  //if all three drivers are installed and enabled
 
 #if defined(GDA_BNXT)
-  if (gda_vendor == GDAVendor::NONE) {
+  if (gda_provider == GDAProvider::UNSET
+  && (requested == GDAProvider::UNSET || requested == GDAProvider::BNXT)) {
     ret = bnxt_dv_dl_init();
 
     if (ret == ROCSHMEM_SUCCESS) {
-      gda_vendor = GDAVendor::BNXT;
+      gda_provider = GDAProvider::BNXT;
     } else {
       DPRINTF("Initializing rocSHMEM BNXT GDA support failed\n");
     }
@@ -656,11 +684,12 @@ void GDABackend::open_dv_libs() {
 #endif // defined(GDA_BNXT)
 
 #if defined(GDA_IONIC)
-  if (gda_vendor == GDAVendor::NONE) {
+  if (gda_provider == GDAProvider::UNSET
+  && (requested == GDAProvider::UNSET || requested == GDAProvider::IONIC)) {
     ret = ionic_dv_dl_init();
 
     if (ret == ROCSHMEM_SUCCESS) {
-      gda_vendor = GDAVendor::IONIC;
+      gda_provider = GDAProvider::IONIC;
     } else {
       DPRINTF("Initializing rocSHMEM IONIC GDA support failed\n");
     }
@@ -668,18 +697,19 @@ void GDABackend::open_dv_libs() {
 #endif // defined(GDA_IONIC)
 
 #if defined(GDA_MLX5)
-  if (gda_vendor == GDAVendor::NONE) {
+  if (gda_provider == GDAProvider::UNSET
+  && (requested == GDAProvider::UNSET || requested == GDAProvider::MLX5)) {
     ret = mlx5_dv_dl_init();
 
     if (ret == ROCSHMEM_SUCCESS) {
-      gda_vendor = GDAVendor::MLX5;
+      gda_provider = GDAProvider::MLX5;
     } else {
       DPRINTF("Initializing rocSHMEM MLX5 GDA support failed\n");
     }
   }
 #endif // defined(GDA_MLX5)
 
-  if (gda_vendor == GDAVendor::NONE) {
+  if (gda_provider == GDAProvider::UNSET) {
     printf("Initializing rocSHMEM with IONIC, BNXT, or MLX5 GDA support failed: no DV library found\n");
     abort();
   }
@@ -695,7 +725,7 @@ void GDABackend::close_dv_libs() {
   if (mlx5dv_handle_ != nullptr)
     dlclose(mlx5dv_handle_);
 
-  gda_vendor = GDAVendor::NONE;
+  gda_provider = GDAProvider::UNSET;
 }
 
 void GDABackend::exchange_qp_dest_info() {
@@ -766,7 +796,7 @@ void GDABackend::setup_gpu_qps() {
   CHECK_NNULL(host_qps, "malloc (host_qps)");
 
   for (size_t i = 0; i < qp_objs_count; i++) {
-    new (&host_qps[i]) QueuePair(pd_orig, gda_vendor);
+    new (&host_qps[i]) QueuePair(pd_orig, gda_provider);
     CHECK_HIP(hipMemcpy(&gpu_qps[i], &host_qps[i], sizeof(QueuePair), hipMemcpyDefault));
 
     initialize_gpu_qp(&gpu_qps[i], i);
@@ -821,7 +851,7 @@ void GDABackend::open_ib_device() {
   CHECK_NNULL(pd_orig, "ib allocate pd");
   dump_ibv_pd(pd_orig);
 
-  if (gda_vendor == GDAVendor::IONIC || gda_vendor == GDAVendor::MLX5) {
+  if (gda_provider == GDAProvider::IONIC || gda_provider == GDAProvider::MLX5) {
     create_parent_domain();
   }
 
@@ -856,7 +886,7 @@ void GDABackend::modify_qps_reset_to_init() {
             | IBV_QP_ACCESS_FLAGS;
 
   for (int i =0; i < qps.size() ; i++) {
-    if (gda_vendor == GDAVendor::BNXT) {
+    if (gda_provider == GDAProvider::BNXT) {
       err = bnxt_re_dv.modify_qp(qps[i], &attr, attr_mask, 0, 0);
     } else {
       err = ibv_modify_qp(qps[i], &attr, attr_mask);
@@ -876,7 +906,7 @@ void GDABackend::modify_qps_init_to_rtr() {
   attr.min_rnr_timer          = 12;
   attr.ah_attr.port_num       = port;
 
-  if (gda_vendor == GDAVendor::IONIC) {
+  if (gda_provider == GDAProvider::IONIC) {
     attr.max_dest_rd_atomic = 15;
   } else {
     attr.max_dest_rd_atomic = 1;
@@ -908,7 +938,7 @@ void GDABackend::modify_qps_init_to_rtr() {
       attr.ah_attr.dlid = dest_info[i].lid;
     }
 
-    if (gda_vendor == GDAVendor::BNXT) {
+    if (gda_provider == GDAProvider::BNXT) {
       err = bnxt_re_dv.modify_qp(qps[i], &attr, attr_mask, 0, 0);
     } else {
       err = ibv_modify_qp(qps[i], &attr, attr_mask);
@@ -928,7 +958,7 @@ void GDABackend::modify_qps_rtr_to_rts() {
   attr.retry_cnt     = 7;
   attr.rnr_retry     = 7;
 
-  if (gda_vendor == GDAVendor::IONIC) {
+  if (gda_provider == GDAProvider::IONIC) {
     attr.max_rd_atomic = 15;
   } else {
     attr.max_rd_atomic = 1;
@@ -944,7 +974,7 @@ void GDABackend::modify_qps_rtr_to_rts() {
   for (int i = 0; i < qps.size(); i++) {
     attr.sq_psn = dest_info[i].psn;
 
-    if (gda_vendor == GDAVendor::BNXT) {
+    if (gda_provider == GDAProvider::BNXT) {
       err = bnxt_re_dv.modify_qp(qps[i], &attr, attr_mask, 0, 0);
     } else {
       err = ibv_modify_qp(qps[i], &attr, attr_mask);
@@ -957,7 +987,7 @@ void GDABackend::create_queues() {
   int ncqes;
   size_t resize_length;
 
-  if (gda_vendor == GDAVendor::IONIC) {
+  if (gda_provider == GDAProvider::IONIC) {
     ncqes = envvar::sq_size << 1;
   } else {
     ncqes = envvar::sq_size;
@@ -972,7 +1002,7 @@ void GDABackend::create_queues() {
   bnxt_cqs.resize(resize_length);
   bnxt_qps.resize(resize_length);
 
-  if (gda_vendor == GDAVendor::BNXT) {
+  if (gda_provider == GDAProvider::BNXT) {
     bnxt_create_cqs(ncqes);
     bnxt_create_qps(envvar::sq_size);
   } else {
@@ -1060,7 +1090,7 @@ void GDABackend::create_parent_domain() {
   pattr.free       = GDABackend::pd_release;
   pattr.pd_context = nullptr;
 
-  if (gda_vendor == GDAVendor::IONIC) {
+  if (gda_provider == GDAProvider::IONIC) {
     pattr.alloc      = GDABackend::pd_alloc_device_uncached;
   } else {
     pattr.alloc      = GDABackend::pd_alloc_host;
@@ -1070,7 +1100,7 @@ void GDABackend::create_parent_domain() {
   CHECK_NNULL(pd_parent, "ibv_alloc_parent_domain");
   dump_ibv_pd(pd_parent);
 
-  if (gda_vendor == GDAVendor::IONIC) {
+  if (gda_provider == GDAProvider::IONIC) {
     ionic_setup_parent_domain(&pattr);
   }
 }
@@ -1089,7 +1119,7 @@ void GDABackend::create_cqs(int cqe) {
   cq_attr.parent_domain = pd_parent;
 
   for (int i = 0; i < qps.size(); i++) {
-    if (gda_vendor == GDAVendor::IONIC) {
+    if (gda_provider == GDAProvider::IONIC) {
       cq_attr.parent_domain = pd_uxdma[i & 1];
     }
 
@@ -1102,18 +1132,18 @@ void GDABackend::create_cqs(int cqe) {
 }
 
 void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) {
-  switch (gda_vendor) {
-  case GDAVendor::IONIC:
+  switch (gda_provider) {
+  case GDAProvider::IONIC:
     ionic_initialize_gpu_qp(gpu_qp, conn_num);
     break;
-  case GDAVendor::BNXT:
+  case GDAProvider::BNXT:
     bnxt_initialize_gpu_qp(gpu_qp, conn_num);
     break;
-  case GDAVendor::MLX5:
+  case GDAProvider::MLX5:
     mlx5_initialize_gpu_qp(gpu_qp, conn_num);
     break;
   default:
-    assert(false /* GDAVendor initialize_gpu_qp */);
+    assert(false /* GDAProvider initialize_gpu_qp */);
   }
 }
 
@@ -1129,12 +1159,12 @@ void GDABackend::create_qps(int sq_length) {
   attr.comp_mask           = IBV_QP_INIT_ATTR_PD;
   attr.pd                  = pd_parent;
 
-  if (gda_vendor == GDAVendor::IONIC) {
+  if (gda_provider == GDAProvider::IONIC) {
     attr.cap.max_recv_sge    = 1; // TODO allow zero sges in the driver
   }
 
   for (int i = 0; i < qps.size(); i++) {
-    if (gda_vendor == GDAVendor::IONIC) {
+    if (gda_provider == GDAProvider::IONIC) {
       attr.pd      = pd_uxdma[i & 1];
     }
     attr.send_cq = cqs[i];

src/gda/backend_gda.hpp

@@ -47,8 +47,8 @@ class GDAHostContext;
 class QueuePair;
 class HostInterface;
 
-enum GDAVendor {
-  NONE,
+enum GDAProvider {
+  UNSET,
   IONIC,
   BNXT,
   MLX5
@@ -66,7 +66,7 @@ class GDABackend : public Backend {
   const char *requested_dev = nullptr;
   struct ibv_context *context = nullptr;;
   struct ibv_pd *pd_orig = nullptr;
-  enum GDAVendor gda_vendor = GDAVendor::NONE;
+  enum GDAProvider gda_provider = GDAProvider::UNSET;
 
   struct ibv_port_attr portinfo;
   union ibv_gid gid;
@@ -102,9 +102,14 @@ class GDABackend : public Backend {
   /* GDA_IONIC END */
 
  /**
-   * @brief Common code invoked from the different constructors
+   * @brief Choose nic device according to locality/user preferences
    */
-  void read_env();
+  void select_nic();
+
+  /**
+   * @brief return user-preferred GDA provider (or NONE if not specified)
+   */
+  static GDAProvider requested_provider();
 
  public:
   friend GDAContext;

src/gda/queue_pair.cpp

@@ -31,7 +31,7 @@
 
 namespace rocshmem {
 
-QueuePair::QueuePair(struct ibv_pd* pd, int gda_vendor) {
+QueuePair::QueuePair(struct ibv_pd* pd, int gda_provider) {
   int access = IBV_ACCESS_LOCAL_WRITE
              | IBV_ACCESS_REMOTE_WRITE
              | IBV_ACCESS_REMOTE_READ
@@ -51,7 +51,7 @@ QueuePair::QueuePair(struct ibv_pd* pd, int gda_vendor) {
   mr_fetching_atomic = ibv_reg_mr(pd, fetching_atomic, 8 * FETCHING_ATOMIC_CNT, access);
   CHECK_NNULL(mr_fetching_atomic, "ibv_reg_mr");
 
-  if (gda_vendor == GDAVendor::MLX5) {
+  if (gda_provider == GDAProvider::MLX5) {
     nonfetching_atomic_lkey = htobe32(mr_nonfetching_atomic->lkey);
     fetching_atomic_lkey = htobe32(mr_fetching_atomic->lkey);
   } else {
@@ -64,9 +64,9 @@ QueuePair::QueuePair(struct ibv_pd* pd, int gda_vendor) {
   }
 
   /* Set Correct opcodes for each NIC */
-  switch (gda_vendor) {
+  switch (gda_provider) {
 #if defined(GDA_IONIC)
-  case GDAVendor::IONIC:
+  case GDAProvider::IONIC:
     gda_op_rdma_write = IONIC_V2_OP_RDMA_WRITE;
     gda_op_rdma_read  = IONIC_V2_OP_RDMA_READ;
     gda_op_atomic_fa  = IONIC_V2_OP_ATOMIC_FA;
@@ -74,7 +74,7 @@ QueuePair::QueuePair(struct ibv_pd* pd, int gda_vendor) {
     break;
 #endif //defined(GDA_IONIC)
 #if defined(GDA_BNXT)
-  case GDAVendor::BNXT:
+  case GDAProvider::BNXT:
     gda_op_rdma_write = BNXT_RE_WR_OPCD_RDMA_WRITE;
     gda_op_rdma_read  = BNXT_RE_WR_OPCD_RDMA_READ;
     gda_op_atomic_fa  = BNXT_RE_WR_OPCD_ATOMIC_FA;
@@ -82,7 +82,7 @@ QueuePair::QueuePair(struct ibv_pd* pd, int gda_vendor) {
     break;
 #endif //defined(GDA_BNXT)
 #if defined(GDA_MLX5)
-  case GDAVendor::MLX5:
+  case GDAProvider::MLX5:
     gda_op_rdma_write = MLX5_OPCODE_RDMA_WRITE;
     gda_op_rdma_read  = MLX5_OPCODE_RDMA_READ;
     gda_op_atomic_fa  = MLX5_OPCODE_ATOMIC_FA;
@@ -92,7 +92,7 @@ QueuePair::QueuePair(struct ibv_pd* pd, int gda_vendor) {
   default:
     assert(false /* invalid nic provider */);
   }
-  gda_vendor_ = gda_vendor;
+  gda_provider_ = gda_provider;
 }
 
 QueuePair::~QueuePair() {
@@ -115,19 +115,19 @@ QueuePair::~QueuePair() {
  ************************ PROVIDER-SPECIFIC HELPERS ***************************
  *****************************************************************************/
 __device__ void QueuePair::post_wqe_rma(int pe, int32_t size, uintptr_t *laddr, uintptr_t *raddr, uint8_t opcode) {
-  switch (gda_vendor_) {
+  switch (gda_provider_) {
 #if defined(GDA_MLX5)
-  case GDAVendor::MLX5:
+  case GDAProvider::MLX5:
     mlx5_post_wqe_rma(pe, size, laddr, raddr, opcode);
     return;
 #endif
 #if defined(GDA_BNXT)
-  case GDAVendor::BNXT:
+  case GDAProvider::BNXT:
     bnxt_post_wqe_rma(pe, size, laddr, raddr, opcode);
     return;
 #endif
 #if defined(GDA_IONIC)
-  case GDAVendor::IONIC:
+  case GDAProvider::IONIC:
     ionic_post_wqe_rma(pe, size, laddr, raddr, opcode);
     return;
 #endif
@@ -138,17 +138,17 @@ __device__ void QueuePair::post_wqe_rma(int pe, int32_t size, uintptr_t *laddr,
 
 __device__ uint64_t QueuePair::post_wqe_amo(int pe, int32_t size, uintptr_t *raddr, uint8_t opcode,
                                             int64_t atomic_data, int64_t atomic_cmp, bool fetching) {
-  switch (gda_vendor_) {
+  switch (gda_provider_) {
 #if defined(GDA_MLX5)
-  case GDAVendor::MLX5:
+  case GDAProvider::MLX5:
     return mlx5_post_wqe_amo(pe, size, raddr, opcode, atomic_data, atomic_cmp, fetching);
 #endif
 #if defined(GDA_BNXT)
-  case GDAVendor::BNXT:
+  case GDAProvider::BNXT:
     return bnxt_post_wqe_amo(pe, size, raddr, opcode, atomic_data, atomic_cmp, fetching);
 #endif
 #if defined(GDA_IONIC)
-  case GDAVendor::IONIC:
+  case GDAProvider::IONIC:
     return ionic_post_wqe_amo(pe, size, raddr, opcode, atomic_data, atomic_cmp, fetching);
 #endif
   default:
@@ -158,19 +158,19 @@ __device__ uint64_t QueuePair::post_wqe_amo(int pe, int32_t size, uintptr_t *rad
 }
 
 __device__ void QueuePair::quiet() {
-  switch (gda_vendor_) {
+  switch (gda_provider_) {
 #if defined(GDA_MLX5)
-  case GDAVendor::MLX5:
+  case GDAProvider::MLX5:
     mlx5_quiet();
     return;
 #endif
 #if defined(GDA_BNXT)
-  case GDAVendor::BNXT:
+  case GDAProvider::BNXT:
     bnxt_quiet();
     return;
 #endif
 #if defined(GDA_IONIC)
-  case GDAVendor::IONIC:
+  case GDAProvider::IONIC:
     ionic_quiet();
     return;
 #endif

src/gda/queue_pair.hpp

@@ -57,7 +57,7 @@ class QueuePair {
   /**
    * @brief Constructor.
    */
-  explicit QueuePair(struct ibv_pd* pd, int gda_vendor);
+  explicit QueuePair(struct ibv_pd* pd, int gda_provider);
 
   /**
    * @brief Destructor.
@@ -192,18 +192,18 @@ class QueuePair {
   __device__ void ionic_ring_doorbell(uint32_t pos);
 #endif
 
-  int gda_vendor_{0};
+  int gda_provider_{0};
 
-  /* GDAVendor::BNXT START */
+  /* GDAProvider::BNXT START */
   uint64_t *dbr;
   struct bnxt_device_cq cq;
   struct bnxt_device_sq sq;
 
   __device__ int poll_cq();
 
-  /* GDAVendor::BNXT END */
+  /* GDAProvider::BNXT END */
 
-  /* GDAVendor::MLX5 START */
+  /* GDAProvider::MLX5 START */
 
   db_reg_t db{};
 
@@ -264,9 +264,9 @@ class QueuePair {
   static constexpr size_t OUTSTANDING_TABLE_SIZE = 65536;
   uint64_t outstanding_wqes[OUTSTANDING_TABLE_SIZE]{0};
 
-  /* GDAVendor::MLX5 END */
+  /* GDAProvider::MLX5 END */
 
-  /* GDAVendor::IONIC START */
+  /* GDAProvider::IONIC START */
 
   uint64_t *cq_dbreg{nullptr};
   uint64_t cq_dbval{0};
@@ -316,7 +316,7 @@ class QueuePair {
    */
   __device__ __attribute__((noinline)) void ionic_quiet_internal(uint64_t active_lane_mask, uint32_t cons);
 
-  /* GDAVendor::IONIC END */
+  /* GDAProvider::IONIC END */
 
   uint32_t inline_threshold{0};