diff --git a/src/envvar.cpp b/src/envvar.cpp index 9b94e31d28..ea20236bd7 100644 --- a/src/envvar.cpp +++ b/src/envvar.cpp @@ -65,6 +65,7 @@ namespace envvar { } // namespace ro namespace gda { + const var provider("PROVIDER", ""); const var alternate_qp_ports("ALTERNATE_QP_PORTS", "", true); const var traffic_class("TRAFFIC_CLASS", "", 0); } // namespace gda diff --git a/src/envvar.hpp b/src/envvar.hpp index 583901ac08..4838f315ca 100644 --- a/src/envvar.hpp +++ b/src/envvar.hpp @@ -448,6 +448,7 @@ namespace envvar { namespace gda { template using var = var; + extern const var provider; extern const var alternate_qp_ports; extern const var traffic_class; } // namespace gda diff --git a/src/gda/backend_gda.cpp b/src/gda/backend_gda.cpp index 232089cb8f..4d3e95cdc3 100644 --- a/src/gda/backend_gda.cpp +++ b/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]; diff --git a/src/gda/backend_gda.hpp b/src/gda/backend_gda.hpp index 71bb4024c6..d2a83cb3b9 100644 --- a/src/gda/backend_gda.hpp +++ b/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; diff --git a/src/gda/queue_pair.cpp b/src/gda/queue_pair.cpp index 36b5c83cdc..e8c670b90b 100644 --- a/src/gda/queue_pair.cpp +++ b/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 diff --git a/src/gda/queue_pair.hpp b/src/gda/queue_pair.hpp index 3bb17588f4..fd74db41a4 100644 --- a/src/gda/queue_pair.hpp +++ b/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};