From cb39f7a31350dd0eac67036c12938b63b16e2221 Mon Sep 17 00:00:00 2001 From: Yiltan Date: Wed, 10 Sep 2025 09:13:36 -0400 Subject: [PATCH] Unify common BNXT and MLX5 initialization code (#233) Co-authored-by: Aurelien Bouteiller --- CMakeLists.txt | 3 +- cmake/rocshmem_config.h.in | 1 + scripts/build_configs/gda_mlx5 | 2 +- src/gda/CMakeLists.txt | 7 +- src/gda/backend_gda.cpp | 757 ++++++++++++++---------------- src/gda/backend_gda.hpp | 198 +++----- src/gda/bnxt/backend_gda_bnxt.cpp | 162 +------ src/gda/debug_gda.hpp | 240 ++++++++++ src/gda/queue_pair.cpp | 94 ++-- src/gda/queue_pair.hpp | 16 +- 10 files changed, 743 insertions(+), 737 deletions(-) create mode 100644 src/gda/debug_gda.hpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 47c9d3f620..6b12f59688 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -70,7 +70,8 @@ option(BUILD_LOCAL_GPU_TARGET_ONLY "Build only for GPUs detected on this machine option(BUILD_CODE_COVERAGE "Build with code coverage flags (gcc only)" OFF) option(GDA_IONIC "Build for AMD Pensando IONIC RDMA provider" OFF) -option(GDA_BNXT "Build for Broadcom" OFF) +option(GDA_BNXT "Build for Broadcom RDMA provider" OFF) +option(GDA_MLX5 "Build for Mellanox MLX5 RDMA provider" OFF) ############################################################################### # PROJECT diff --git a/cmake/rocshmem_config.h.in b/cmake/rocshmem_config.h.in index 644a87a69e..7b54099d38 100644 --- a/cmake/rocshmem_config.h.in +++ b/cmake/rocshmem_config.h.in @@ -44,3 +44,4 @@ #cmakedefine USE_HDP_FLUSH_HOST_SIDE #cmakedefine GDA_IONIC #cmakedefine GDA_BNXT +#cmakedefine GDA_MLX5 diff --git a/scripts/build_configs/gda_mlx5 b/scripts/build_configs/gda_mlx5 index 9337aed9d1..c99cf21f50 100755 --- a/scripts/build_configs/gda_mlx5 +++ b/scripts/build_configs/gda_mlx5 @@ -27,4 +27,4 @@ set -e script_path=$(dirname "$(realpath $0)") -exec $script_path/gda $* +source $script_path/gda -DGDA_MLX5=ON $* diff --git a/src/gda/CMakeLists.txt b/src/gda/CMakeLists.txt index 056f5f2423..70480ecc43 100644 --- a/src/gda/CMakeLists.txt +++ b/src/gda/CMakeLists.txt @@ -47,9 +47,10 @@ target_link_libraries( numa ) -if (GDA_IONIC) -elseif (GDA_BNXT) +if (GDA_BNXT) add_subdirectory(bnxt) -else() +endif() + +if (GDA_MLX5) target_sources(${PROJECT_NAME} PRIVATE segment_builder.cpp) endif() diff --git a/src/gda/backend_gda.cpp b/src/gda/backend_gda.cpp index fde76b2ecd..242e8005c3 100644 --- a/src/gda/backend_gda.cpp +++ b/src/gda/backend_gda.cpp @@ -545,69 +545,97 @@ void GDABackend::rte_barrier() { } } -static void dump_ibv_context(struct ibv_context *x); -static void dump_ibv_device(struct ibv_device *x); -static void dump_ibv_pd(struct ibv_pd *x); -static void dump_ibv_port_attr(struct ibv_port_attr *x); -static void dump_ibv_qp(struct ibv_qp *qp, int conn_num); -static void dump_mlx5dv_qp(struct mlx5dv_qp *qp_dv, int conn_num); -static void dump_mlx5dv_cq(struct mlx5dv_cq *cq_dv, int conn_num); - void GDABackend::setup_ibv() { - dest_info.resize(num_pes * (maximum_num_contexts_ + 1)); - int ib_devices{0}; - dev_list = ibv_get_device_list(&ib_devices); - CHECK_NNULL(dev_list, "ibv_get_device"); - struct ibv_device* ib_dev = dev_list[0]; //TODO default to HIP selected device? - if (requested_dev) { - for (int i = 0; i < ib_devices; i++) { - const char* select_dev{ibv_get_device_name(dev_list[i])}; - CHECK_NNULL(select_dev, "ibv_get_device_name"); - if (strstr(select_dev, requested_dev)) { - ib_dev = dev_list[i]; - break; - } - } - } - uint8_t port{1}; - ib_init(ib_dev, port); - create_qps(port, &ib_state->portinfo); + open_ib_device(); - auto npes = num_pes; - auto dinfo = dest_info.data(); - for (int i = 0; i < maximum_num_contexts_ + 1; i++) { - if (backend_comm != MPI_COMM_NULL) { - MPI_Alltoall(MPI_IN_PLACE, sizeof(dest_info_t), MPI_CHAR, dinfo + i * npes, sizeof(dest_info_t), MPI_CHAR, backend_comm); - } else { - Alltoall_char_inplace(reinterpret_cast(dinfo + i * npes), sizeof(dest_info_t), ROCSHMEM_TEAM_WORLD); - } - } + create_queues(); + + exchange_qp_dest_info(); + + modify_qps_reset_to_init(); + modify_qps_init_to_rtr(); + modify_qps_rtr_to_rts(); - for (int i = 0; i < qps.size(); i++) { - change_status_rtr(qps[i], &dest_info[i], port); - } - rte_barrier(); - for (int i = 0; i < qps.size(); i++) { - change_status_rts(qps[i], &dest_info[i]); - dump_ibv_qp(qps[i], i); - } rte_barrier(); } void GDABackend::cleanup_ibv() { - ibv_free_device_list(dev_list); + int err; - delete ib_state; - if (requested_dev != nullptr) - free(requested_dev); +#ifdef GDA_BNXT + CHECK_HIP(hipHostUnregister(db_region_attr.dbr)); + + for (int i = 0; i < qps.size(); i++) { + err = bnxt_re_dv_destroy_qp(qps[i]); + CHECK_ZERO(err, "bnxt_re_dv_destroy_qp"); + + err = bnxt_re_dv_umem_dereg(bnxt_qps[i].attr.rq_umem_handle); + CHECK_ZERO(err, "bnxt_re_dv_umem_dereg (RQ)"); + + err = bnxt_re_dv_umem_dereg(bnxt_qps[i].attr.sq_umem_handle); + CHECK_ZERO(err, "bnxt_re_dv_umem_dereg (SQ)"); + + CHECK_HIP(hipFree(bnxt_qps[i].sq_buf)); + CHECK_HIP(hipFree(bnxt_qps[i].rq_buf)); + + err = bnxt_re_dv_destroy_cq(cqs[i]); + CHECK_ZERO(err, "bnxt_re_dv_destroy_cq"); + + err = bnxt_re_dv_umem_dereg(bnxt_cqs[i].umem_handle); + CHECK_ZERO(err, "bnxt_re_dv_umem_dereg"); + + CHECK_HIP(hipFree(bnxt_cqs[i].buf)); + } +#else + for (int i = 0; i < qps.size(); i++) { + err = ibv_destroy_qp(qps[i]); + CHECK_ZERO(err, "ibv_destroy_qp"); + + err = ibv_destroy_cq(cqs[i]); + CHECK_ZERO(err, "ibv_destroy_cqs"); + } + +#ifdef GDA_IONIC + err = ibv_dealloc_pd(pd_uxdma[0]); + CHECK_ZERO(err, "ibv_dealloc_pd (uxdma[0])"); + + err = ibv_dealloc_pd(pd_uxdma[1]); + CHECK_ZERO(err, "ibv_dealloc_pd (uxdma[1])"); +#endif + + err = ibv_dealloc_pd(pd_parent); + CHECK_ZERO(err, "ibv_dealloc_pd (pd_parent)"); +#endif + + err = ibv_dealloc_pd(pd_orig); + CHECK_ZERO(err, "ibv_dealloc_pd (pd_orig)"); + + err = ibv_close_device(context); + CHECK_ZERO(err, "ibv_close_device"); } +void GDABackend::exchange_qp_dest_info() { + for (int i = 0; i < qps.size(); i++) { + dest_info[i].lid = portinfo.lid; + dest_info[i].qpn = qps[i]->qp_num; + dest_info[i].psn = 0; + dest_info[i].gid = gid; + } + + for (int i = 0; i < maximum_num_contexts_ + 1; i++) { + if (backend_comm != MPI_COMM_NULL) { + MPI_Alltoall(MPI_IN_PLACE, sizeof(dest_info_t), MPI_CHAR, dest_info.data() + i * num_pes, sizeof(dest_info_t), MPI_CHAR, backend_comm); + } else { + Alltoall_char_inplace(reinterpret_cast(dest_info.data() + i * num_pes), sizeof(dest_info_t), ROCSHMEM_TEAM_WORLD); + } + } +} void GDABackend::setup_heap_memory_rkey() { auto *base_heap = heap.get_local_heap_base(); int access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ | IBV_ACCESS_REMOTE_ATOMIC; - heap_mr = ibv_reg_mr(ib_state->pd_orig, base_heap, heap.get_size(), access); + heap_mr = ibv_reg_mr(pd_orig, base_heap, heap.get_size(), access); CHECK_NNULL(heap_mr, "ibv_reg_mr"); const size_t rkeys_size = sizeof(uint32_t) * num_pes; @@ -644,7 +672,7 @@ void GDABackend::cleanup_heap_memory_rkey() { void GDABackend::setup_gpu_qps() { CHECK_HIP(hipMalloc(&gpu_qps, sizeof(QueuePair) * (maximum_num_contexts_ + 1) * num_pes)); for (int i = 0; i < (maximum_num_contexts_ + 1) * num_pes; i++) { - QueuePair qp(ib_state->pd_orig); + QueuePair qp(pd_orig); CHECK_HIP(hipMemcpy(&gpu_qps[i], &qp, sizeof(QueuePair), hipMemcpyDefault)); initialize_gpu_qp(&gpu_qps[i], i); } @@ -657,119 +685,185 @@ void GDABackend::cleanup_gpu_qps() { } //TODO this ifdef sequence should go in a nic-specific file, like it is for bnxt, maybe whats above too? -#ifndef GDA_BNXT -void GDABackend::ib_init(struct ibv_device* ib_dev, uint8_t port) { - ib_state = new ib_state_t; - CHECK_NNULL(ib_state, "ib_state object create"); +void GDABackend::open_ib_device() { + struct ibv_device **device_list = nullptr; + struct ibv_device *device = nullptr; + int num_devices = 0; + int err; - ib_state->context = ibv_open_device(ib_dev); - CHECK_NNULL(ib_state->context, "ib open device"); - dump_ibv_context(ib_state->context); - dump_ibv_device(ib_state->context->device); + device_list = ibv_get_device_list(&num_devices); + CHECK_NNULL(device_list, "ibv_get_device_list"); - ib_state->pd_orig = ibv_alloc_pd(ib_state->context); - CHECK_NNULL(ib_state->pd_orig, "ib allocate pd"); - dump_ibv_pd(ib_state->pd_orig); + device = device_list[0]; //TODO default to HIP selected device? - ibv_parent_domain_init_attr pattr{}; - init_parent_domain_attr(&pattr); - ib_state->pd_parent = ibv_alloc_parent_domain(ib_state->context, &pattr); - CHECK_NNULL(ib_state->pd_parent, "ibv_alloc_parent_domain"); - dump_ibv_pd(ib_state->pd_parent); + if (requested_dev) { + for (int i = 0; i < num_devices; i++) { + const char *select_device = ibv_get_device_name(device_list[i]); + CHECK_NNULL(select_device, "ibv_get_device_name"); -#ifdef GDA_IONIC - ionic_dv_pd_set_sqcmb(ib_state->pd_parent, false, false, false); - ionic_dv_pd_set_rqcmb(ib_state->pd_parent, false, false, false); - - for (int uxdma_i = 0; uxdma_i < 2; ++uxdma_i) { - ib_state->pd_uxdma[uxdma_i] = ibv_alloc_parent_domain(ib_state->context, &pattr); - CHECK_NNULL(ib_state->pd_uxdma[uxdma_i], "ibv_alloc_parent_domain (uxdma)"); - - ionic_dv_pd_set_sqcmb(ib_state->pd_uxdma[uxdma_i], false, false, false); - ionic_dv_pd_set_rqcmb(ib_state->pd_uxdma[uxdma_i], false, false, false); - ionic_dv_pd_set_udma_mask(ib_state->pd_uxdma[uxdma_i], 1u << uxdma_i); + if (strstr(select_device, requested_dev)) { + device = device_list[i]; + break; + } + } } + + context = ibv_open_device(device); + CHECK_NNULL(context, "ib open device"); + dump_ibv_context(context); + dump_ibv_device(context->device); + + pd_orig = ibv_alloc_pd(context); + CHECK_NNULL(pd_orig, "ib allocate pd"); + dump_ibv_pd(pd_orig); + +#ifndef GDA_BNXT + create_parent_domain(); #endif - int err = ibv_query_port(ib_state->context, port, &ib_state->portinfo); + err = ibv_query_port(context, port, &portinfo); CHECK_ZERO(err, "ibv_query_port"); - dump_ibv_port_attr(&ib_state->portinfo); + dump_ibv_port_attr(&portinfo); /* Must init after querying port */ - init_gid_index(port); + select_gid_index(); -#ifdef GDA_IONIC - ionic_dv_ctx dvctx; - ionic_dv_get_ctx(&dvctx, ib_state->context); - - int hip_dev_id = 0; - CHECK_HIP(hipGetDevice(&hip_dev_id)); - - void* gpu_db_page = nullptr; - rocm_memory_lock_to_fine_grain(dvctx.db_page, 0x1000, &gpu_db_page, hip_dev_id); - - uint64_t *db_page_u64 = reinterpret_cast(dvctx.db_page); - uint64_t *gpu_db_page_u64 = reinterpret_cast(gpu_db_page); - - uint64_t *gpu_db_ptr = &gpu_db_page_u64[dvctx.db_ptr - db_page_u64]; - - ib_state->gpu_db_page = gpu_db_page; - ib_state->gpu_db_cq = &gpu_db_ptr[dvctx.cq_qtype]; - ib_state->gpu_db_sq = &gpu_db_ptr[dvctx.sq_qtype]; -#endif + ibv_free_device_list(device_list); } -template -void GDABackend::try_to_modify_qp(ibv_qp* qp, StateType state) { - int err = ibv_modify_qp(qp, &state.exp_qp_attr, state.exp_attr_mask); - CHECK_ZERO(err, "ibv_modify_qp"); -} +void GDABackend::modify_qps_reset_to_init() { + int err; + struct ibv_qp_attr attr; + int attr_mask; -void GDABackend::init_qp_status(ibv_qp* qp, uint8_t port) { - try_to_modify_qp(qp, initqp(port)); -} + memset(&attr, 0, sizeof(struct ibv_qp_attr)); -void GDABackend::change_status_rtr(ibv_qp* qp, dest_info_t* dest, uint8_t port) { - try_to_modify_qp(qp, rtr(dest, port)); -} + attr.qp_state = IBV_QPS_INIT; + attr.pkey_index = 0; + attr.port_num = port; + attr.qp_access_flags = IBV_ACCESS_REMOTE_WRITE + | IBV_ACCESS_LOCAL_WRITE + | IBV_ACCESS_REMOTE_READ + | IBV_ACCESS_REMOTE_ATOMIC; -void GDABackend::change_status_rts(ibv_qp* qp, dest_info_t* dest) { - try_to_modify_qp(qp, rts(dest)); -} + attr_mask = IBV_QP_STATE + | IBV_QP_PKEY_INDEX + | IBV_QP_PORT + | IBV_QP_ACCESS_FLAGS; -void GDABackend::create_qps(uint8_t port, ibv_port_attr* ib_port_att) { - ibv_qp_cap cap{}; - cap.max_send_wr = sq_size; - cap.max_send_sge = 1; - cap.max_inline_data = 0; -#ifdef GDA_IONIC - // TODO allow zero sges in the driver - cap.max_recv_sge = 1; -#endif - QPInitAttr qp_init_attr{qpattr(cap)}; - cqs.resize((maximum_num_contexts_ + 1) * num_pes); - qps.resize((maximum_num_contexts_ + 1) * num_pes); - int max_num_cqe = qp_init_attr.attr.cap.max_send_wr; - for (int i = 0; i < qps.size(); i++) { -#ifdef GDA_IONIC - int uxdma_i = ((i + 1) / 2) & 1; - cqs[i] = create_cq(ib_state->context, ib_state->pd_uxdma[uxdma_i], max_num_cqe << 1); - CHECK_NNULL(cqs[i], "create_cq"); - qps[i] = create_qp(ib_state->pd_uxdma[uxdma_i], ib_state->context, &qp_init_attr.attr, cqs[i]); + for (int i =0; i < qps.size() ; i++) { +#ifdef GDA_BNXT + err = bnxt_re_dv_modify_qp(qps[i], &attr, attr_mask, 0, 0); #else - cqs[i] = create_cq(ib_state->context, ib_state->pd_parent, max_num_cqe); - CHECK_NNULL(cqs[i], "create_cq"); - qps[i] = create_qp(ib_state->pd_parent, ib_state->context, &qp_init_attr.attr, cqs[i]); + err = ibv_modify_qp(qps[i], &attr, attr_mask); #endif - CHECK_NNULL(qps[i], "create_qp"); - init_qp_status(qps[i], port); - dest_info[i].lid = ib_port_att->lid; - dest_info[i].qpn = qps[i]->qp_num; - dest_info[i].psn = 0; - dest_info[i].gid = gid; + CHECK_ZERO(err, "modify_qp (INIT)"); } } +void GDABackend::modify_qps_init_to_rtr() { + struct ibv_qp_attr attr; + int attr_mask; + int err; + + memset(&attr, 0, sizeof(struct ibv_qp_attr)); + attr.qp_state = IBV_QPS_RTR; + attr.path_mtu = portinfo.active_mtu; + attr.max_dest_rd_atomic = GDA_MAX_ATOMIC; + attr.min_rnr_timer = 12; + attr.ah_attr.port_num = port; + + if (portinfo.link_layer == IBV_LINK_LAYER_ETHERNET) { + attr.ah_attr.grh.sgid_index = gid_index; + attr.ah_attr.is_global = 1; + attr.ah_attr.grh.hop_limit = 1; + attr.ah_attr.sl = 1; + } + + attr_mask = IBV_QP_STATE + | IBV_QP_PATH_MTU + | IBV_QP_RQ_PSN + | IBV_QP_DEST_QPN + | IBV_QP_AV + | IBV_QP_MAX_DEST_RD_ATOMIC + | IBV_QP_MIN_RNR_TIMER; + + for (int i = 0; i < qps.size(); i++) { + attr.rq_psn = dest_info[i].psn; + attr.dest_qp_num = dest_info[i].qpn; + + if (portinfo.link_layer == IBV_LINK_LAYER_ETHERNET) { + memcpy(&attr.ah_attr.grh.dgid, &dest_info[i].gid, 16); + } else { + attr.ah_attr.dlid = dest_info[i].lid; + } + +#ifdef GDA_BNXT + err = bnxt_re_dv_modify_qp(qps[i], &attr, attr_mask, 0, 0); +#else + err = ibv_modify_qp(qps[i], &attr, attr_mask); +#endif + CHECK_ZERO(err, "modify_qp (RTR)"); + } +} + +void GDABackend::modify_qps_rtr_to_rts() { + struct ibv_qp_attr attr; + int attr_mask; + int err; + + memset(&attr, 0, sizeof(struct ibv_qp_attr)); + attr.qp_state = IBV_QPS_RTS; + attr.max_rd_atomic = GDA_MAX_ATOMIC; + attr.timeout = 14; + attr.retry_cnt = 7; + attr.rnr_retry = 7; + + attr_mask = IBV_QP_STATE + | IBV_QP_SQ_PSN + | IBV_QP_MAX_QP_RD_ATOMIC + | IBV_QP_TIMEOUT + | IBV_QP_RETRY_CNT + | IBV_QP_RNR_RETRY; + + for (int i = 0; i < qps.size(); i++) { + attr.sq_psn = dest_info[i].psn; + +#ifdef GDA_BNXT + err = bnxt_re_dv_modify_qp(qps[i], &attr, attr_mask, 0, 0); +#else + err = ibv_modify_qp(qps[i], &attr, attr_mask); +#endif + CHECK_ZERO(err, "modify_qp (RTS)"); + } +} + +void GDABackend::create_queues() { + int ncqes; + int resize_length; + +#ifdef GDA_IONIC + ncqes = sq_size << 1; +#else + ncqes = sq_size; +#endif + + resize_length = (maximum_num_contexts_ + 1) * num_pes; + + dest_info.resize(resize_length); + cqs.resize(resize_length); + qps.resize(resize_length); + +#ifdef GDA_BNXT + bnxt_cqs.resize(resize_length); + bnxt_qps.resize(resize_length); +#endif + + create_cqs(ncqes); + create_qps(sq_size); +} + +#ifndef GDA_BNXT void* GDABackend::pd_alloc(struct ibv_pd* pd, void* pd_context, size_t size, size_t alignment, uint64_t resource_type) { void* dev_ptr{nullptr}; //TODO make this configurable, presumably we want it on device for all types? @@ -786,30 +880,60 @@ void GDABackend::pd_release(struct ibv_pd* pd, void* pd_context, void* ptr, uint CHECK_HIP(hipFree(ptr)); } -void GDABackend::init_parent_domain_attr(ibv_parent_domain_init_attr* attr1) { - attr1->pd = ib_state->pd_orig; - attr1->td = nullptr; - attr1->comp_mask = IBV_PARENT_DOMAIN_INIT_ATTR_ALLOCATORS; - attr1->alloc = GDABackend::pd_alloc; - attr1->free = GDABackend::pd_release; - attr1->pd_context = nullptr; +void GDABackend::create_parent_domain() { + struct ibv_parent_domain_init_attr pattr; + + memset(&pattr, 0, sizeof(struct ibv_parent_domain_init_attr)); + pattr.pd = pd_orig, + pattr.td = nullptr, + pattr.comp_mask = IBV_PARENT_DOMAIN_INIT_ATTR_ALLOCATORS, + pattr.alloc = GDABackend::pd_alloc, + pattr.free = GDABackend::pd_release, + pattr.pd_context = nullptr, + + pd_parent = ibv_alloc_parent_domain(context, &pattr); + CHECK_NNULL(pd_parent, "ibv_alloc_parent_domain"); + dump_ibv_pd(pd_parent); + +#ifdef GDA_IONIC + ionic_dv_pd_set_sqcmb(pd_parent, false, false, false); + ionic_dv_pd_set_rqcmb(pd_parent, false, false, false); + + for (int uxdma_i = 0; uxdma_i < 2; ++uxdma_i) { + pd_uxdma[uxdma_i] = ibv_alloc_parent_domain(context, &pattr); + CHECK_NNULL(pd_uxdma[uxdma_i], "ibv_alloc_parent_domain (uxdma)"); + + ionic_dv_pd_set_sqcmb(pd_uxdma[uxdma_i], false, false, false); + ionic_dv_pd_set_rqcmb(pd_uxdma[uxdma_i], false, false, false); + ionic_dv_pd_set_udma_mask(pd_uxdma[uxdma_i], 1u << uxdma_i); + } +#endif } -ibv_cq* GDABackend::create_cq(ibv_context* context, ibv_pd* pd, int cqe) { - ibv_cq_init_attr_ex cq_attr; - memset(&cq_attr, 0, sizeof(ibv_cq_init_attr_ex)); - cq_attr.cqe = cqe; - cq_attr.cq_context = nullptr; - cq_attr.channel = nullptr; - cq_attr.comp_vector = 0; - cq_attr.flags = 0; // see ibv_exp_cq_create_flags - cq_attr.comp_mask = IBV_CQ_INIT_ATTR_MASK_PD; - cq_attr.parent_domain = pd; - ibv_cq_ex* cq_ex = ibv_create_cq_ex(context, &cq_attr); - CHECK_NNULL(cq_ex, "ibv_create_cq_ex"); - ibv_cq *cq = ibv_cq_ex_to_cq(cq_ex); - CHECK_NNULL(cq, "ibv_cq_ex_to_cq"); - return cq; +void GDABackend::create_cqs(int cqe) { + struct ibv_cq_init_attr_ex cq_attr; + struct ibv_cq_ex *cq_ex; + + memset(&cq_attr, 0, sizeof(struct ibv_cq_init_attr_ex)); + cq_attr.cqe = cqe; + cq_attr.cq_context = nullptr; + cq_attr.channel = nullptr; + cq_attr.comp_vector = 0; + cq_attr.flags = 0; + cq_attr.comp_mask = IBV_CQ_INIT_ATTR_MASK_PD; + cq_attr.parent_domain = pd_parent; + + for (int i = 0; i < qps.size(); i++) { +#ifdef GDA_IONIC + cq_attr.parent_domain = pd_uxdma[((i + 1) / 2) & 1]; +#endif + + cq_ex = ibv_create_cq_ex(context, &cq_attr); + CHECK_NNULL(cq_ex, "ibv_create_cq_ex"); + + cqs[i] = ibv_cq_ex_to_cq(cq_ex); + CHECK_NNULL(cqs[i], "ibv_cq_ex_to_cq"); + } } void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) { @@ -817,12 +941,27 @@ void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) { CHECK_HIP(hipGetDevice(&hip_dev_id)); #ifdef GDA_IONIC + ionic_dv_ctx dvctx; + ionic_dv_get_ctx(&dvctx, context); + + void* gpu_db_page = nullptr; + rocm_memory_lock_to_fine_grain(dvctx.db_page, 0x1000, &gpu_db_page, hip_dev_id); + + uint64_t *db_page_u64 = reinterpret_cast(dvctx.db_page); + uint64_t *gpu_db_page_u64 = reinterpret_cast(gpu_db_page); + + uint64_t *gpu_db_ptr = &gpu_db_page_u64[dvctx.db_ptr - db_page_u64]; + + gpu_db_page = gpu_db_page; + gpu_db_cq = &gpu_db_ptr[dvctx.cq_qtype]; + gpu_db_sq = &gpu_db_ptr[dvctx.sq_qtype]; + uint8_t udma_idx = ionic_dv_qp_get_udma_idx(qps[conn_num]); ionic_dv_cq dvcq; ionic_dv_get_cq(&dvcq, cqs[conn_num], udma_idx); - gpu_qp->cq_dbreg = ib_state->gpu_db_cq; + gpu_qp->cq_dbreg = gpu_db_cq; gpu_qp->cq_dbval = dvcq.q.db_val; gpu_qp->cq_mask = dvcq.q.mask; @@ -831,7 +970,7 @@ void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) { ionic_dv_qp dvqp; ionic_dv_get_qp(&dvqp, qps[conn_num]); - gpu_qp->sq_dbreg = ib_state->gpu_db_sq; + gpu_qp->sq_dbreg = gpu_db_sq; gpu_qp->sq_dbval = dvqp.sq.db_val; gpu_qp->sq_mask = dvqp.sq.mask; gpu_qp->sq_buf = reinterpret_cast(dvqp.sq.ptr); @@ -911,61 +1050,35 @@ void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) { #endif // !GDA_IONIC } -ibv_qp* GDABackend::create_qp(ibv_pd* pd, ibv_context* context, ibv_qp_init_attr_ex* qp_attr, ibv_cq* cq) { - ibv_qp* qp{nullptr}; - assert(pd); - assert(context); - assert(qp_attr); - qp_attr->send_cq = cq; - qp_attr->recv_cq = cq; - qp_attr->pd = pd; - qp_attr->comp_mask = IBV_QP_INIT_ATTR_PD; - qp = ibv_create_qp_ex(context, qp_attr); - CHECK_NNULL(qp, "ibv_create_qp_ex"); - return qp; -} +void GDABackend::create_qps(int sq_length) { + struct ibv_qp_init_attr_ex attr; -GDABackend::InitQPState GDABackend::initqp(uint8_t port) { - InitQPState init{}; - init.exp_qp_attr.qp_access_flags = IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_READ | IBV_ACCESS_REMOTE_ATOMIC; - init.exp_qp_attr.port_num = port; - init.exp_attr_mask |= IBV_QP_ACCESS_FLAGS; - return init; -} + memset(&attr, 0, sizeof(struct ibv_qp_init_attr_ex)); + attr.cap.max_send_wr = sq_length; + attr.cap.max_send_sge = 1; + attr.cap.max_inline_data = 0; +#ifdef GDA_IONIC + attr.cap.max_recv_sge = 1; // TODO allow zero sges in the driver +#endif + attr.sq_sig_all = 0; + attr.qp_type = IBV_QPT_RC; + attr.comp_mask = IBV_QP_INIT_ATTR_PD; + attr.pd = pd_parent; -GDABackend::RtrState GDABackend::rtr(dest_info_t* dest, uint8_t port) { - RtrState rtr{}; - rtr.exp_qp_attr.dest_qp_num = dest->qpn; - rtr.exp_qp_attr.rq_psn = dest->psn; - rtr.exp_qp_attr.ah_attr.port_num = port; - rtr.exp_qp_attr.path_mtu = ib_state->portinfo.active_mtu; - if (ib_state->portinfo.link_layer == IBV_LINK_LAYER_INFINIBAND) { - rtr.exp_qp_attr.ah_attr.dlid = dest->lid; - } else { - rtr.exp_qp_attr.ah_attr.is_global = 1; - rtr.exp_qp_attr.ah_attr.grh.dgid = dest->gid; - rtr.exp_qp_attr.ah_attr.grh.sgid_index = gid_index; - rtr.exp_qp_attr.ah_attr.grh.hop_limit = 1; + for (int i = 0; i < qps.size(); i++) { +#ifdef GDA_IONIC + attr.pd = pd_uxdma[((i + 1) / 2) & 1]; +#endif + attr.send_cq = cqs[i]; + attr.recv_cq = cqs[i]; + + qps[i] = ibv_create_qp_ex(context, &attr); + CHECK_NNULL(qps[i], "ibv_create_qp_ex"); } - rtr.exp_attr_mask |= IBV_QP_DEST_QPN | IBV_QP_RQ_PSN | IBV_QP_MAX_DEST_RD_ATOMIC | IBV_QP_MIN_RNR_TIMER; - return rtr; -} - -GDABackend::RtsState GDABackend::rts(dest_info_t* dest) { - RtsState rts{}; - rts.exp_qp_attr.sq_psn = dest->psn; - rts.exp_attr_mask |= IBV_QP_SQ_PSN; - return rts; -} - -GDABackend::QPInitAttr GDABackend::qpattr(ibv_qp_cap cap) { - QPInitAttr qpattr(cap); - qpattr.attr.qp_type = IBV_QPT_RC; - return qpattr; } #endif -void GDABackend::init_gid_index(uint8_t port_num) { +void GDABackend::select_gid_index() { struct ibv_gid_entry *gid_entries; struct ibv_gid_entry *gid_entry; union ibv_gid current_gid; @@ -978,12 +1091,11 @@ void GDABackend::init_gid_index(uint8_t port_num) { int selected_gid_index = -1; ssize_t gid_tbl_entries = 0; - int gid_tbl_len = ib_state->portinfo.gid_tbl_len; - struct ibv_context *ctx = ib_state->context; + int gid_tbl_len = portinfo.gid_tbl_len; gid_entries = (struct ibv_gid_entry*) calloc(gid_tbl_len, sizeof(struct ibv_gid_entry)); - gid_tbl_entries = ibv_query_gid_table(ctx, gid_entries, gid_tbl_len, 0); + gid_tbl_entries = ibv_query_gid_table(context, gid_entries, gid_tbl_len, 0); if (gid_tbl_entries < 0) { fprintf(stderr, "[Warning] ibv_query_gid_table failed. No available GIDs\n"); free(gid_entries); @@ -1000,7 +1112,7 @@ void GDABackend::init_gid_index(uint8_t port_num) { current_gid = gid_entries[i].gid; - err = ibv_query_gid(ctx, port_num, i, ¤t_gid); + err = ibv_query_gid(context, port, i, ¤t_gid); CHECK_ZERO(err, "ibv_query_gid"); /* We don't want local GIDs */ @@ -1028,205 +1140,18 @@ void GDABackend::init_gid_index(uint8_t port_num) { free(gid_entries); } -static void dump_ibv_context(struct ibv_context* x) { - /* - * struct ibv_context { - * struct ibv_device *device; - * struct ibv_context_ops ops; - * int cmd_fd; - * int async_fd; - * int num_comp_vectors; - * pthread_mutex_t mutex; - * void *abi_compat; - * }; - */ - DPRINTF("\n" - "===============================================\n" - " IBV_CONTEXT\n" - "===============================================\n" - " (ibv_device*) device = %p\n" - " (int) cmd_fd = %d\n" - " (int) async_fd = %d\n" - " (int) num_comp_vectors = %d\n" - " (void*) abi_compat = %p\n", - x->device, x->cmd_fd, x->async_fd, x->num_comp_vectors, x->abi_compat); -}; - -static void dump_ibv_device(struct ibv_device* x) { - /* - * struct ibv_device { - * struct _ibv_device_ops _ops; - * enum ibv_node_type node_type; - * enum ibv_transport_type transport_type; - * char name[IBV_SYSFS_NAME_MAX]; - * char dev_name[IBV_SYSFS_NAME_MAX]; - * char dev_path[IBV_SYSFS_PATH_MAX]; - * char ibdev_path[IBV_SYSFS_PATH_MAX]; - * }; - */ - DPRINTF("\n" - "===============================================\n" - " IBV_DEVICE\n" - "===============================================\n" - " (enum ibv_node_type) node_type = %d\n" - " (enum ibv_transport_type) transport_type = %d\n" - " (char[]) name = %s\n" - " (char[]) dev_name = %s\n" - " (char[]) dev_path = %s\n" - " (char[]) ibdev_path = %s\n", - x->node_type, x->transport_type, x->name, x->dev_name, x->dev_path, x->ibdev_path); +int GDABackend::ibv_mtu_to_int(enum ibv_mtu mtu) { + switch (mtu) { + case IBV_MTU_256: return 256; + case IBV_MTU_512: return 512; + case IBV_MTU_1024: return 1024; + case IBV_MTU_2048: return 2048; + case IBV_MTU_4096: return 4096; + default: { + fprintf(stderr, "[ERROR] Invalid ibv_mtu\n"); + return 0; + } + } } -static void dump_ibv_pd(struct ibv_pd* x) { - /* - * struct ibv_pd { - * struct ibv_context *context; - * uint32_t handle; - * }; - */ - DPRINTF("\n" - "===============================================\n" - " IBV_PD\n" - "===============================================\n" - " (ibv_context*) context = %p\n" - " (uint32_t) handle = 0x%x\n", - x->context, x->handle); -} - -static void dump_ibv_port_attr(struct ibv_port_attr* x) { - /* - * struct ibv_port_attr { - * enum ibv_port_state state; - * enum ibv_mtu max_mtu; - * enum ibv_mtu active_mtu; - * int gid_tbl_len; - * uint32_t port_cap_flags; - * uint32_t max_msg_sz; - * uint32_t bad_pkey_cntr; - * uint32_t qkey_viol_cntr; - * uint16_t pkey_tbl_len; - * uint16_t lid; - * uint16_t sm_lid; - * uint8_t lmc; - * uint8_t max_vl_num; - * uint8_t sm_sl; - * uint8_t subnet_timeout; - * uint8_t init_type_reply; - * uint8_t active_width; - * uint8_t active_speed; - * uint8_t phys_state; - * uint8_t link_layer; - * uint8_t flags; - * uint16_t port_cap_flags2; - * }; - */ - DPRINTF("\n" - "===============================================\n" - " IBV_PORT_ATTR\n" - "===============================================\n" - " (enum ibv_port_state) state = %u\n" - " (enum ibv_mtu) max_mtu = %u\n" - " (enum ibv_mtu) active_mtu = %u\n" - " (int) gid_tbl_len = %u\n" - " (uint32_t) port_cap_flags = 0x%x\n" - " (uint32_t) max_msg_sz = %u\n" - " (uint32_t) bad_pkey_cntr = %u\n" - " (uint32_t) qkey_viol_cntr = %u\n" - " (uint16_t) pkey_tbl_len = %u\n" - " (uint16_t) lid = 0x%x\n" - " (uint16_t) sm_lid = 0x%x\n" - " (uint8_t) lmc = 0x%x\n" - " (uint8_t) max_vl_num = 0x%x\n" - " (uint8_t) sm_sl = 0x%x\n" - " (uint8_t) subnet_timeout = 0x%x\n" - " (uint8_t) init_type_reply = 0x%x\n" - " (uint8_t) active_width = 0x%x\n" - " (uint8_t) active_speed = 0x%x\n" - " (uint8_t) phys_state = 0x%x\n" - " (uint8_t) link_layer = 0x%x\n" - " (uint8_t) flags = 0x%x\n" - " (uint16_t) port_cap_flags2 = 0x%x\n", - x->state, x->max_mtu, x->active_mtu, x->gid_tbl_len, x->port_cap_flags, x->max_msg_sz, - x->bad_pkey_cntr, x->qkey_viol_cntr, x->pkey_tbl_len, x->lid, x->sm_lid, x->lmc, x->max_vl_num, - x->sm_sl, x->subnet_timeout, x->init_type_reply, x->active_width, x->active_speed, x->phys_state, - x->link_layer, x->flags, x->port_cap_flags2); -} - -void dump_ibv_qp(struct ibv_qp *qp, int conn_num) { - /* - * struct ibv_qp { - * struct ibv_context *context; - * void *qp_context; - * struct ibv_pd *pd; - * struct ibv_cq *send_cq; - * struct ibv_cq *recv_cq; - * struct ibv_srq *srq; - * uint32_t handle; - * uint32_t qp_num; - * enum ibv_qp_state state; - * enum ibv_qp_type qp_type; - * pthread_mutex_t mutex; - * pthread_cond_t cond; - * uint32_t events_completed; - * }; - */ - DPRINTF("\n"); - DPRINTF("============== QP_DUMP CONNECTION#%d ==========\n", conn_num); - DPRINTF(" (ibv_context*) context = %p\n", qp->context); - DPRINTF(" (void*) qp_context = %p\n", qp->qp_context); - DPRINTF(" (ibv_pd*) pd = %p\n", qp->pd); - DPRINTF(" (ibv_cq*) send_cq = %p\n", qp->send_cq); - DPRINTF(" (ibv_cq*) recv_cq = %p\n", qp->recv_cq); - DPRINTF(" (ibv_srq*) srq = %p\n", qp->srq); - DPRINTF(" (uint32_t) handle = 0x%x\n", qp->handle); - DPRINTF(" (uint32_t) qp_num = 0x%x\n", qp->qp_num); - DPRINTF(" (enum ibv_qp_state) state = %u\n", qp->state); - DPRINTF(" (enum_ibv_qp_type) qp_type = %u\n", qp->qp_type); - DPRINTF(" (uint32_t) events_completed = %u\n", qp->events_completed); - DPRINTF("=========== QP_DUMP_END CONNECTION#%d ========\n", conn_num); -} - -#if !defined(GDA_IONIC) && !defined(GDA_BNXT) -void dump_mlx5dv_qp(struct mlx5dv_qp *qp_dv, int conn_num) { - DPRINTF("\n"); - DPRINTF("===============================================\n"); - DPRINTF(" INITIALIZED MLXDV_QP FOR CONNECTION#%d\n", conn_num); - DPRINTF("===============================================\n"); - DPRINTF("=================== QP_DUMP ===================\n"); - DPRINTF(" (__be32*) dbrec = %p\n", qp_dv->dbrec); - DPRINTF(" (void*) sq.buf = %p\n", qp_dv->sq.buf); - DPRINTF(" (uint32_t) sq.wqe_cnt = %u\n", qp_dv->sq.wqe_cnt); - DPRINTF(" (uint32_t) sq.stride = %u\n", qp_dv->sq.stride); - DPRINTF(" (void*) rq.buf = %p\n", qp_dv->rq.buf); - DPRINTF(" (uint32_t) rq.wqe_cnt = %u\n", qp_dv->rq.wqe_cnt); - DPRINTF(" (uint32_t) rq.stride = %u\n", qp_dv->rq.stride); - DPRINTF(" (void*) bf.reg = %p\n", qp_dv->bf.reg); - DPRINTF(" (uint32_t) bf.size = 0x%x\n", qp_dv->bf.size); - DPRINTF(" (uint64_t) comp_mask = 0x%lx\n", qp_dv->comp_mask); - DPRINTF(" (off_t) uar_mmap_offset = 0x%lx\n", qp_dv->uar_mmap_offset); - DPRINTF(" (uint32_t) tirn = 0x%x\n", qp_dv->tirn); - DPRINTF(" (uint32_t) tisn = 0x%x\n", qp_dv->tisn); - DPRINTF(" (uint32_t) rqn = 0x%x\n", qp_dv->rqn); - DPRINTF(" (uint32_t) sqn = 0x%x\n", qp_dv->sqn); - DPRINTF(" (uint64_t) tir_icm_addr = 0x%lx\n", qp_dv->tir_icm_addr); - DPRINTF("================== QP_DUMP_END ================\n"); -} - -void dump_mlx5dv_cq(struct mlx5dv_cq *cq_dv, int conn_num) { - DPRINTF("\n"); - DPRINTF("===============================================\n"); - DPRINTF(" INITIALIZED MLX5DV_CQ FOR CONNECTION#%d\n", conn_num); - DPRINTF("===============================================\n"); - DPRINTF("=================== CQ_DUMP ===================\n"); - DPRINTF(" (void*) buf = %p\n", cq_dv->buf); - DPRINTF(" (__be32*) dbrec = %p\n", cq_dv->dbrec); - DPRINTF(" (uint32_t) cqe_cnt = %u\n", cq_dv->cqe_cnt); - DPRINTF(" (uint32_t) cqe_size = %u\n", cq_dv->cqe_size); - DPRINTF(" (void*) cq_uar = %p\n", cq_dv->cq_uar); - DPRINTF(" (uint32_t) cqn = 0x%x\n", cq_dv->cqn); - DPRINTF(" (uint64_t) comp_mask = 0x%lx\n", cq_dv->comp_mask); - DPRINTF("================== CQ_DUMP_END ================\n"); -} -#endif // !GDA_IONIC - } // namespace rocshmem diff --git a/src/gda/backend_gda.hpp b/src/gda/backend_gda.hpp index becdcfdb46..db055c1a3c 100644 --- a/src/gda/backend_gda.hpp +++ b/src/gda/backend_gda.hpp @@ -33,6 +33,7 @@ #include "gda_context_proxy.hpp" #include "queue_pair.hpp" #include "bootstrap/bootstrap.hpp" +#include "debug_gda.hpp" namespace rocshmem { @@ -43,25 +44,6 @@ class HostInterface; class GDABackend : public Backend { private: - typedef struct ib_state { - struct ibv_context* context; - struct ibv_pd* pd_orig; -#ifndef GDA_BNXT - struct ibv_pd* pd_parent; -#endif -#ifdef GDA_IONIC - struct ibv_pd* pd_uxdma[2]; -#endif - struct ibv_mr* mr; - struct ibv_port_attr portinfo; - -#ifdef GDA_IONIC - void *gpu_db_page; - uint64_t *gpu_db_cq; - uint64_t *gpu_db_sq; -#endif - } ib_state_t; - typedef struct dest_info { int lid; int qpn; @@ -69,61 +51,44 @@ class GDABackend : public Backend { union ibv_gid gid; } dest_info_t; -#ifndef GDA_BNXT - class State { - public: - ibv_qp_attr exp_qp_attr{}; - uint64_t exp_attr_mask{}; - }; + char *requested_dev = nullptr; + struct ibv_context *context = nullptr;; + struct ibv_pd *pd_orig = nullptr; - class InitQPState : public State { - public: - InitQPState() { - exp_qp_attr.qp_state = IBV_QPS_INIT; - exp_qp_attr.qp_access_flags = IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_READ | IBV_ACCESS_REMOTE_ATOMIC; - exp_attr_mask = IBV_QP_STATE | IBV_QP_PKEY_INDEX | IBV_QP_PORT; - } - }; + struct ibv_port_attr portinfo; + union ibv_gid gid; + int port = 1; + int gid_index; - class RtrState : public State { - public: - RtrState() { - exp_qp_attr.qp_state = IBV_QPS_RTR; - exp_qp_attr.ah_attr.sl = 1; - exp_qp_attr.max_dest_rd_atomic = GDA_MAX_ATOMIC; - exp_qp_attr.min_rnr_timer = 12; - exp_attr_mask = IBV_QP_STATE | IBV_QP_AV | IBV_QP_PATH_MTU; - } - }; + uint32_t *heap_rkey = nullptr; + struct ibv_mr *heap_mr = nullptr; - class RtsState : public State { - public: - RtsState() { - exp_qp_attr.qp_state = IBV_QPS_RTS; - exp_qp_attr.timeout = 14; - exp_qp_attr.retry_cnt = 7; - exp_qp_attr.rnr_retry = 7; - exp_qp_attr.max_rd_atomic = GDA_MAX_ATOMIC; - exp_attr_mask = IBV_QP_STATE | IBV_QP_TIMEOUT | IBV_QP_RETRY_CNT | IBV_QP_RNR_RETRY | IBV_QP_MAX_QP_RD_ATOMIC; - } - }; + uint32_t sq_size = 1024; + QueuePair *gpu_qps = nullptr; + std::vector qps; + std::vector cqs; + std::vector dest_info; - class QPInitAttr { - public: - explicit QPInitAttr(ibv_qp_cap cap) { - attr.cap = cap; - attr.sq_sig_all = 0; - } - ibv_qp_init_attr_ex attr{}; - }; +#ifdef GDA_BNXT + std::vector bnxt_qps; + std::vector bnxt_cqs; + + struct bnxt_re_dv_db_region_attr db_region_attr; +#else + struct ibv_pd *pd_parent = nullptr; +#endif + +#ifdef GDA_IONIC + struct ibv_pd *pd_uxdma[2]; + void *gpu_db_page = nullptr; + uint64_t *gpu_db_cq = nullptr; + uint64_t *gpu_db_sq = nullptr; #endif /** * @brief Common code invoked from the different constructors */ void read_env(); - void setup_ibv(); - void cleanup_ibv(); public: friend GDAContext; @@ -299,62 +264,77 @@ class GDABackend : public Backend { void initialize_gpu_qp(QueuePair* qp, int conn_num); -#ifndef GDA_BNXT - InitQPState initqp(uint8_t port); + /** + * @brief Setup InfiniBand Resources + */ + void setup_ibv(); - RtrState rtr(dest_info_t* dest, uint8_t port); + /** + * @brief Cleanup InfiniBand Resources + */ + void cleanup_ibv(); - RtsState rts(dest_info_t* dest); + /** + * @brief Open InfiniBand Device and create common structures + */ + void open_ib_device(); - QPInitAttr qpattr(ibv_qp_cap cap); + /** + * @brief Selects the best GID index + */ + void select_gid_index(); - void init_qp_status(ibv_qp* qp, uint8_t port); -#endif + /** + * @brief Create all CQs and QPs + */ + void create_queues(); - void change_status_rtr(ibv_qp* qp, dest_info_t* dest, uint8_t port); + /** + * @brief Create all CQs with a of length ncqes + */ + void create_cqs(int ncqes); - void change_status_rts(ibv_qp* qp, dest_info_t* dest); + /** + * @brief Create all QPs with a SQ of length sq_length + */ + void create_qps(int sq_length); - void create_qps(uint8_t port, ibv_port_attr* ib_port_att); + /** + * @brief Exchange QP information for connection + */ + void exchange_qp_dest_info(); -#ifdef GDA_BNXT - void init_qp_status(uint8_t port); + /** + * @brief Modify all QPs from RESET to INIT state + */ + void modify_qps_reset_to_init(); - void create_cqs(int ncqs, int cqe); + /** + * @brief Modify all QPs from INIT to RTR state + */ + void modify_qps_init_to_rtr(); - void create_qps_impl(int nqps); + /** + * @brief Modify all QPs from RTR to RTs state + */ + void modify_qps_rtr_to_rts(); + /** + * @brief Converts an ibv_mtu to an integer + */ int ibv_mtu_to_int(enum ibv_mtu mtu); -#else - template - void try_to_modify_qp(ibv_qp* qp, T state); +#ifndef GDA_BNXT static void* pd_alloc(ibv_pd* pd, void* pd_context, size_t size, size_t alignment, uint64_t resource_type); static void pd_release(ibv_pd* pd, void* pd_context, void* ptr, uint64_t resource_type); - void init_parent_domain_attr(ibv_parent_domain_init_attr* attr); - - ibv_cq* create_cq(ibv_context* context, ibv_pd* pd, int cqe); - - ibv_qp* create_qp(ibv_pd* pd, ibv_context* context, ibv_qp_init_attr_ex* qp_attr, ibv_cq* rcq); + void create_parent_domain(); #endif - void ib_init(ibv_device* ib_dev, uint8_t port); - - void init_gid_index(uint8_t port); - void setup_gpu_qps(); void cleanup_gpu_qps(); - char* requested_dev{nullptr}; - - ibv_device** dev_list{nullptr}; - - ib_state_t* ib_state{nullptr}; - - std::vector dest_info; - private: /** * @brief Common code invoked from the different constructors @@ -456,28 +436,6 @@ class GDABackend : public Backend { * @brief rte barrier for initialization */ void rte_barrier(); - - QueuePair *gpu_qps{nullptr}; - - std::vector qps; - - std::vector cqs; - - uint32_t sq_size{1024}; - - uint32_t *heap_rkey{nullptr}; - - ibv_mr *heap_mr{nullptr}; - - union ibv_gid gid; - int gid_index; - -#ifdef GDA_BNXT - std::vector bnxt_qps; - std::vector bnxt_cqs; - - struct bnxt_re_dv_db_region_attr db_region_attr; -#endif }; } // namespace rocshmem diff --git a/src/gda/bnxt/backend_gda_bnxt.cpp b/src/gda/bnxt/backend_gda_bnxt.cpp index b9ea02dec5..10eeb0dd3b 100644 --- a/src/gda/bnxt/backend_gda_bnxt.cpp +++ b/src/gda/bnxt/backend_gda_bnxt.cpp @@ -28,151 +28,13 @@ namespace rocshmem { -int GDABackend::ibv_mtu_to_int(enum ibv_mtu mtu) { - switch (mtu) { - case IBV_MTU_256: return 256; - case IBV_MTU_512: return 512; - case IBV_MTU_1024: return 1024; - case IBV_MTU_2048: return 2048; - case IBV_MTU_4096: return 4096; - default: { - fprintf(stderr, "[ERROR] Invalid ibv_mtu\n"); - return 0; - } - } -} - -void GDABackend::ib_init(struct ibv_device* ib_dev, uint8_t port) { - int err; - - ib_state = new ib_state_t; - CHECK_NNULL(ib_state, "ib_state object create"); - - ib_state->context = ibv_open_device(ib_dev); - CHECK_NNULL(ib_state->context, "ibv_open_device"); - - ib_state->pd_orig = ibv_alloc_pd(ib_state->context); - CHECK_NNULL(ib_state->pd_orig, "ibv_alloc_pd"); - - err = ibv_query_port(ib_state->context, port, &ib_state->portinfo); - CHECK_ZERO(err, "ibv_query_port"); - - init_gid_index(port); -} - -void GDABackend::init_qp_status(uint8_t port) { - int err; - struct ibv_qp_attr attr; - int attr_mask; - - memset(&attr, 0, sizeof(struct ibv_qp_attr)); - - attr.qp_state = IBV_QPS_INIT; - attr.pkey_index = 0; - attr.port_num = port; - attr.qp_access_flags = IBV_ACCESS_REMOTE_WRITE - | IBV_ACCESS_LOCAL_WRITE - | IBV_ACCESS_REMOTE_READ - | IBV_ACCESS_REMOTE_ATOMIC; - - attr_mask = IBV_QP_STATE - | IBV_QP_PKEY_INDEX - | IBV_QP_PORT - | IBV_QP_ACCESS_FLAGS; - - for (int i =0; i < qps.size() ; i++) { - err = bnxt_re_dv_modify_qp(qps[i], &attr, attr_mask, 0, 0); - CHECK_ZERO(err, "bnxt_re_dv_modify_qp"); - } -} - -void GDABackend::change_status_rtr(ibv_qp *qp, dest_info_t *dest, uint8_t port) { - int err; - struct ibv_qp_attr attr; - int attr_mask; - - memset(&attr, 0, sizeof(struct ibv_qp_attr)); - attr.qp_state = IBV_QPS_RTR; - attr.path_mtu = ib_state->portinfo.active_mtu; - attr.rq_psn = dest->psn; - attr.dest_qp_num = dest->qpn; - - memcpy(&attr.ah_attr.grh.dgid, &dest->gid, 16); - attr.ah_attr.grh.sgid_index = gid_index; - attr.ah_attr.grh.hop_limit = 1; - attr.ah_attr.sl = 1; - attr.ah_attr.is_global = 1; - attr.ah_attr.port_num = port; - - attr.max_dest_rd_atomic = GDA_MAX_ATOMIC; - attr.min_rnr_timer = 12; - - attr_mask = IBV_QP_STATE - | IBV_QP_PATH_MTU - | IBV_QP_RQ_PSN - | IBV_QP_DEST_QPN - | IBV_QP_AV - | IBV_QP_MAX_DEST_RD_ATOMIC - | IBV_QP_MIN_RNR_TIMER; - - err = bnxt_re_dv_modify_qp(qp, &attr, attr_mask, 0, 0); - CHECK_ZERO(err, "bnxt_re_dv_modify_qp"); -} - -void GDABackend::change_status_rts(ibv_qp* qp, dest_info_t* dest) { - int err; - struct ibv_qp_attr attr; - int attr_mask; - - memset(&attr, 0, sizeof(struct ibv_qp_attr)); - attr.qp_state = IBV_QPS_RTS; - attr.sq_psn = dest->psn; - attr.max_rd_atomic = GDA_MAX_ATOMIC; - attr.timeout = 14; - attr.retry_cnt = 7; - attr.rnr_retry = 7; - - attr_mask = IBV_QP_STATE - | IBV_QP_SQ_PSN - | IBV_QP_MAX_QP_RD_ATOMIC - | IBV_QP_TIMEOUT - | IBV_QP_RETRY_CNT - | IBV_QP_RNR_RETRY; - - err = bnxt_re_dv_modify_qp(qp, &attr, attr_mask, 0, 0); - CHECK_ZERO(err, "bnxt_re_dv_modify_qp"); -} - -void GDABackend::create_qps(uint8_t port, ibv_port_attr* ib_port_att) { - int resize_length = (maximum_num_contexts_ + 1) * num_pes; - - cqs.resize(resize_length); - bnxt_cqs.resize(resize_length); - - bnxt_qps.resize(resize_length); - qps.resize(resize_length); - - create_cqs(qps.size(), sq_size); - create_qps_impl(qps.size()); - init_qp_status(port); - - for (int i{0}; i < qps.size(); i++) { - dest_info[i].lid = ib_port_att->lid; - dest_info[i].qpn = qps[i]->qp_num; - dest_info[i].psn = 0; - dest_info[i].gid = gid; - } -} - void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) { struct bnxt_re_dv_obj dv_obj; struct bnxt_re_dv_cq dv_cq; struct bnxt_re_dv_qp dv_qp; - struct ibv_context *context; struct ibv_qp *ib_qp; int err; - context = ib_state->context; ib_qp = qps[conn_num]; /* Export CQ */ @@ -211,7 +73,7 @@ void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) { gpu_qp->sq.msntbl = bnxt_qps[conn_num].msntbl; gpu_qp->sq.msn_tbl_sz = bnxt_qps[conn_num].msn_tbl_sz; gpu_qp->sq.psn_sz_log2 = std::log2(bnxt_qps[conn_num].mem_info.sq_psn_sz); - gpu_qp->sq.mtu = ibv_mtu_to_int(ib_state->portinfo.active_mtu); + gpu_qp->sq.mtu = ibv_mtu_to_int(portinfo.active_mtu); /* Export DB */ err = bnxt_re_dv_get_default_db_region(context, &db_region_attr); @@ -225,15 +87,12 @@ void GDABackend::initialize_gpu_qp(QueuePair* gpu_qp, int conn_num) { gpu_qp->rkey = heap_rkey[conn_num % num_pes]; } -void GDABackend::create_cqs(int ncqs, int cqe) { +void GDABackend::create_cqs(int cqe) { struct bnxt_re_dv_cq_attr cq_attr; struct bnxt_re_dv_cq_init_attr cq_init_attr; struct bnxt_re_dv_umem_reg_attr umem_attr; - struct ibv_context *context; - context = ib_state->context; - - for (int i = 0; i < ncqs; i++) { + for (int i = 0; i < qps.size(); i++) { /* Allocate CQ mem */ memset(&cq_attr, 0, sizeof(struct bnxt_re_dv_cq_attr)); bnxt_cqs[i].handle = bnxt_re_dv_cq_mem_alloc(context, cqe, &cq_attr); @@ -264,9 +123,7 @@ void GDABackend::create_cqs(int ncqs, int cqe) { } } -void GDABackend::create_qps_impl(int nqps) { - struct ibv_pd *pd; - struct ibv_context *context; +void GDABackend::create_qps(int sq_length) { struct ibv_qp_init_attr ib_qp_attr; struct bnxt_re_dv_umem_reg_attr umem_attr; void *sq_ptr; @@ -277,15 +134,12 @@ void GDABackend::create_qps_impl(int nqps) { uint64_t msntbl_offset; int err; - pd = ib_state->pd_orig; - context = ib_state->context; - - for (int i = 0; i < nqps; i++) { + for (int i = 0; i < qps.size(); i++) { /* IB QP Init Attr */ memset(&ib_qp_attr, 0, sizeof(struct ibv_qp_init_attr)); ib_qp_attr.send_cq = cqs[i]; ib_qp_attr.recv_cq = cqs[i]; - ib_qp_attr.cap.max_send_wr = sq_size; + ib_qp_attr.cap.max_send_wr = sq_length; ib_qp_attr.cap.max_recv_wr = 0; ib_qp_attr.cap.max_send_sge = 1; ib_qp_attr.cap.max_recv_sge = 0; @@ -295,7 +149,7 @@ void GDABackend::create_qps_impl(int nqps) { /* Alloc qp_mem_info */ memset(&bnxt_qps[i].mem_info, 0, sizeof(struct bnxt_re_dv_qp_mem_info)); - err = bnxt_re_dv_qp_mem_alloc(pd, &ib_qp_attr, &bnxt_qps[i].mem_info); + err = bnxt_re_dv_qp_mem_alloc(pd_orig, &ib_qp_attr, &bnxt_qps[i].mem_info); CHECK_ZERO(err, "bnxt_re_dv_qp_mem_alloc"); /* Alloc SQ */ @@ -357,7 +211,7 @@ void GDABackend::create_qps_impl(int nqps) { bnxt_qps[i].attr.comp_mask = bnxt_qps[i].mem_info.comp_mask; /* Alloc QP */ - qps[i] = bnxt_re_dv_create_qp(pd, &bnxt_qps[i].attr); + qps[i] = bnxt_re_dv_create_qp(pd_orig, &bnxt_qps[i].attr); CHECK_NNULL(qps[i], "bnxt_re_dv_create_qp"); } } diff --git a/src/gda/debug_gda.hpp b/src/gda/debug_gda.hpp new file mode 100644 index 0000000000..214f6d0d5c --- /dev/null +++ b/src/gda/debug_gda.hpp @@ -0,0 +1,240 @@ +/****************************************************************************** + * Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef LIBRARY_SRC_GDA_DEBUG_GDA_HPP_ +#define LIBRARY_SRC_GDA_DEBUG_GDA_HPP_ + +static void dump_ibv_context(struct ibv_context *x); +static void dump_ibv_device(struct ibv_device *x); +static void dump_ibv_pd(struct ibv_pd *x); +static void dump_ibv_port_attr(struct ibv_port_attr *x); +static void dump_ibv_qp(struct ibv_qp *qp, int conn_num); + +#if defined(GDA_MLX5) +static void dump_mlx5dv_qp(struct mlx5dv_qp *qp_dv, int conn_num); +static void dump_mlx5dv_cq(struct mlx5dv_cq *cq_dv, int conn_num); +#endif + +static void dump_ibv_context(struct ibv_context* x) { + /* + * struct ibv_context { + * struct ibv_device *device; + * struct ibv_context_ops ops; + * int cmd_fd; + * int async_fd; + * int num_comp_vectors; + * pthread_mutex_t mutex; + * void *abi_compat; + * }; + */ + DPRINTF("\n" + "===============================================\n" + " IBV_CONTEXT\n" + "===============================================\n" + " (ibv_device*) device = %p\n" + " (int) cmd_fd = %d\n" + " (int) async_fd = %d\n" + " (int) num_comp_vectors = %d\n" + " (void*) abi_compat = %p\n", + x->device, x->cmd_fd, x->async_fd, x->num_comp_vectors, x->abi_compat); +}; + +static void dump_ibv_device(struct ibv_device* x) { + /* + * struct ibv_device { + * struct _ibv_device_ops _ops; + * enum ibv_node_type node_type; + * enum ibv_transport_type transport_type; + * char name[IBV_SYSFS_NAME_MAX]; + * char dev_name[IBV_SYSFS_NAME_MAX]; + * char dev_path[IBV_SYSFS_PATH_MAX]; + * char ibdev_path[IBV_SYSFS_PATH_MAX]; + * }; + */ + DPRINTF("\n" + "===============================================\n" + " IBV_DEVICE\n" + "===============================================\n" + " (enum ibv_node_type) node_type = %d\n" + " (enum ibv_transport_type) transport_type = %d\n" + " (char[]) name = %s\n" + " (char[]) dev_name = %s\n" + " (char[]) dev_path = %s\n" + " (char[]) ibdev_path = %s\n", + x->node_type, x->transport_type, x->name, x->dev_name, x->dev_path, x->ibdev_path); +} + +static void dump_ibv_pd(struct ibv_pd* x) { + /* + * struct ibv_pd { + * struct ibv_context *context; + * uint32_t handle; + * }; + */ + DPRINTF("\n" + "===============================================\n" + " IBV_PD\n" + "===============================================\n" + " (ibv_context*) context = %p\n" + " (uint32_t) handle = 0x%x\n", + x->context, x->handle); +} + +static void dump_ibv_port_attr(struct ibv_port_attr* x) { + /* + * struct ibv_port_attr { + * enum ibv_port_state state; + * enum ibv_mtu max_mtu; + * enum ibv_mtu active_mtu; + * int gid_tbl_len; + * uint32_t port_cap_flags; + * uint32_t max_msg_sz; + * uint32_t bad_pkey_cntr; + * uint32_t qkey_viol_cntr; + * uint16_t pkey_tbl_len; + * uint16_t lid; + * uint16_t sm_lid; + * uint8_t lmc; + * uint8_t max_vl_num; + * uint8_t sm_sl; + * uint8_t subnet_timeout; + * uint8_t init_type_reply; + * uint8_t active_width; + * uint8_t active_speed; + * uint8_t phys_state; + * uint8_t link_layer; + * uint8_t flags; + * uint16_t port_cap_flags2; + * }; + */ + DPRINTF("\n" + "===============================================\n" + " IBV_PORT_ATTR\n" + "===============================================\n" + " (enum ibv_port_state) state = %u\n" + " (enum ibv_mtu) max_mtu = %u\n" + " (enum ibv_mtu) active_mtu = %u\n" + " (int) gid_tbl_len = %u\n" + " (uint32_t) port_cap_flags = 0x%x\n" + " (uint32_t) max_msg_sz = %u\n" + " (uint32_t) bad_pkey_cntr = %u\n" + " (uint32_t) qkey_viol_cntr = %u\n" + " (uint16_t) pkey_tbl_len = %u\n" + " (uint16_t) lid = 0x%x\n" + " (uint16_t) sm_lid = 0x%x\n" + " (uint8_t) lmc = 0x%x\n" + " (uint8_t) max_vl_num = 0x%x\n" + " (uint8_t) sm_sl = 0x%x\n" + " (uint8_t) subnet_timeout = 0x%x\n" + " (uint8_t) init_type_reply = 0x%x\n" + " (uint8_t) active_width = 0x%x\n" + " (uint8_t) active_speed = 0x%x\n" + " (uint8_t) phys_state = 0x%x\n" + " (uint8_t) link_layer = 0x%x\n" + " (uint8_t) flags = 0x%x\n" + " (uint16_t) port_cap_flags2 = 0x%x\n", + x->state, x->max_mtu, x->active_mtu, x->gid_tbl_len, x->port_cap_flags, x->max_msg_sz, + x->bad_pkey_cntr, x->qkey_viol_cntr, x->pkey_tbl_len, x->lid, x->sm_lid, x->lmc, x->max_vl_num, + x->sm_sl, x->subnet_timeout, x->init_type_reply, x->active_width, x->active_speed, x->phys_state, + x->link_layer, x->flags, x->port_cap_flags2); +} + +void dump_ibv_qp(struct ibv_qp *qp, int conn_num) { + /* + * struct ibv_qp { + * struct ibv_context *context; + * void *qp_context; + * struct ibv_pd *pd; + * struct ibv_cq *send_cq; + * struct ibv_cq *recv_cq; + * struct ibv_srq *srq; + * uint32_t handle; + * uint32_t qp_num; + * enum ibv_qp_state state; + * enum ibv_qp_type qp_type; + * pthread_mutex_t mutex; + * pthread_cond_t cond; + * uint32_t events_completed; + * }; + */ + DPRINTF("\n"); + DPRINTF("============== QP_DUMP CONNECTION#%d ==========\n", conn_num); + DPRINTF(" (ibv_context*) context = %p\n", qp->context); + DPRINTF(" (void*) qp_context = %p\n", qp->qp_context); + DPRINTF(" (ibv_pd*) pd = %p\n", qp->pd); + DPRINTF(" (ibv_cq*) send_cq = %p\n", qp->send_cq); + DPRINTF(" (ibv_cq*) recv_cq = %p\n", qp->recv_cq); + DPRINTF(" (ibv_srq*) srq = %p\n", qp->srq); + DPRINTF(" (uint32_t) handle = 0x%x\n", qp->handle); + DPRINTF(" (uint32_t) qp_num = 0x%x\n", qp->qp_num); + DPRINTF(" (enum ibv_qp_state) state = %u\n", qp->state); + DPRINTF(" (enum_ibv_qp_type) qp_type = %u\n", qp->qp_type); + DPRINTF(" (uint32_t) events_completed = %u\n", qp->events_completed); + DPRINTF("=========== QP_DUMP_END CONNECTION#%d ========\n", conn_num); +} + +#if defined(GDA_MLX5) +void dump_mlx5dv_qp(struct mlx5dv_qp *qp_dv, int conn_num) { + DPRINTF("\n"); + DPRINTF("===============================================\n"); + DPRINTF(" INITIALIZED MLXDV_QP FOR CONNECTION#%d\n", conn_num); + DPRINTF("===============================================\n"); + DPRINTF("=================== QP_DUMP ===================\n"); + DPRINTF(" (__be32*) dbrec = %p\n", qp_dv->dbrec); + DPRINTF(" (void*) sq.buf = %p\n", qp_dv->sq.buf); + DPRINTF(" (uint32_t) sq.wqe_cnt = %u\n", qp_dv->sq.wqe_cnt); + DPRINTF(" (uint32_t) sq.stride = %u\n", qp_dv->sq.stride); + DPRINTF(" (void*) rq.buf = %p\n", qp_dv->rq.buf); + DPRINTF(" (uint32_t) rq.wqe_cnt = %u\n", qp_dv->rq.wqe_cnt); + DPRINTF(" (uint32_t) rq.stride = %u\n", qp_dv->rq.stride); + DPRINTF(" (void*) bf.reg = %p\n", qp_dv->bf.reg); + DPRINTF(" (uint32_t) bf.size = 0x%x\n", qp_dv->bf.size); + DPRINTF(" (uint64_t) comp_mask = 0x%lx\n", qp_dv->comp_mask); + DPRINTF(" (off_t) uar_mmap_offset = 0x%lx\n", qp_dv->uar_mmap_offset); + DPRINTF(" (uint32_t) tirn = 0x%x\n", qp_dv->tirn); + DPRINTF(" (uint32_t) tisn = 0x%x\n", qp_dv->tisn); + DPRINTF(" (uint32_t) rqn = 0x%x\n", qp_dv->rqn); + DPRINTF(" (uint32_t) sqn = 0x%x\n", qp_dv->sqn); + DPRINTF(" (uint64_t) tir_icm_addr = 0x%lx\n", qp_dv->tir_icm_addr); + DPRINTF("================== QP_DUMP_END ================\n"); +} + +void dump_mlx5dv_cq(struct mlx5dv_cq *cq_dv, int conn_num) { + DPRINTF("\n"); + DPRINTF("===============================================\n"); + DPRINTF(" INITIALIZED MLX5DV_CQ FOR CONNECTION#%d\n", conn_num); + DPRINTF("===============================================\n"); + DPRINTF("=================== CQ_DUMP ===================\n"); + DPRINTF(" (void*) buf = %p\n", cq_dv->buf); + DPRINTF(" (__be32*) dbrec = %p\n", cq_dv->dbrec); + DPRINTF(" (uint32_t) cqe_cnt = %u\n", cq_dv->cqe_cnt); + DPRINTF(" (uint32_t) cqe_size = %u\n", cq_dv->cqe_size); + DPRINTF(" (void*) cq_uar = %p\n", cq_dv->cq_uar); + DPRINTF(" (uint32_t) cqn = 0x%x\n", cq_dv->cqn); + DPRINTF(" (uint64_t) comp_mask = 0x%lx\n", cq_dv->comp_mask); + DPRINTF("================== CQ_DUMP_END ================\n"); +} +#endif // GDA_MLX5 + +#endif /* LIBRARY_SRC_GDA_DEBUG_GDA_HPP_ */ diff --git a/src/gda/queue_pair.cpp b/src/gda/queue_pair.cpp index e44292d364..e021db5b28 100644 --- a/src/gda/queue_pair.cpp +++ b/src/gda/queue_pair.cpp @@ -28,7 +28,7 @@ #include "backend_gda.hpp" #include "endian.hpp" -#if !defined(GDA_IONIC) && !defined(GDA_BNXT) +#if defined(GDA_MLX5) #include "segment_builder.hpp" #endif #include "util.hpp" @@ -37,37 +37,53 @@ namespace rocshmem { QueuePair::QueuePair(struct ibv_pd* pd) { + int access = IBV_ACCESS_LOCAL_WRITE + | IBV_ACCESS_REMOTE_WRITE + | IBV_ACCESS_REMOTE_READ + | IBV_ACCESS_REMOTE_ATOMIC; + allocator.allocate((void**)&nonfetching_atomic, 8); - CHECK_HIP(hipMemset(nonfetching_atomic, 0, 8)); - int access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ | IBV_ACCESS_REMOTE_ATOMIC; - - ibv_mr *mr = ibv_reg_mr(pd, nonfetching_atomic, 8, access); - CHECK_NNULL(mr, "ibv_reg_mr"); - -#if defined(GDA_IONIC) || defined(GDA_BNXT) - nonfetching_atomic_lkey = mr->lkey; -#else - nonfetching_atomic_lkey = htobe32(mr->lkey); -#endif - allocator.allocate((void**)&fetching_atomic, 8 * FETCHING_ATOMIC_CNT); - CHECK_HIP(hipMemset(fetching_atomic, 0, 8 * FETCHING_ATOMIC_CNT)); - access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ | IBV_ACCESS_REMOTE_ATOMIC; - mr = ibv_reg_mr(pd, fetching_atomic, 8 * FETCHING_ATOMIC_CNT, access); - CHECK_NNULL(mr, "ibv_reg_mr"); -#if defined(GDA_IONIC) || defined(GDA_BNXT) - fetching_atomic_lkey = mr->lkey; -#else - fetching_atomic_lkey = htobe32(mr->lkey); -#endif - allocator.allocate((void**)&fetching_atomic_freelist, sizeof(FreeListT*)); new (fetching_atomic_freelist) FreeListT(); + + CHECK_HIP(hipMemset(nonfetching_atomic, 0, 8)); + CHECK_HIP(hipMemset(fetching_atomic, 0, 8 * FETCHING_ATOMIC_CNT)); + + mr_nonfetching_atomic = ibv_reg_mr(pd, nonfetching_atomic, 8, access); + CHECK_NNULL(mr_nonfetching_atomic, "ibv_reg_mr"); + + mr_fetching_atomic = ibv_reg_mr(pd, fetching_atomic, 8 * FETCHING_ATOMIC_CNT, access); + CHECK_NNULL(mr_fetching_atomic, "ibv_reg_mr"); + +#if defined(GDA_MLX5) + nonfetching_atomic_lkey = htobe32(mr_nonfetching_atomic->lkey); + fetching_atomic_lkey = htobe32(mr_fetching_atomic->lkey); +#else + nonfetching_atomic_lkey = mr_nonfetching_atomic->lkey; + fetching_atomic_lkey = mr_fetching_atomic->lkey; +#endif + for(int i{0}; i < FETCHING_ATOMIC_CNT; i+=WF_SIZE) { fetching_atomic_freelist->push_back(fetching_atomic + i); } } +QueuePair::~QueuePair() { + int err; + + err = ibv_dereg_mr(mr_nonfetching_atomic); + CHECK_ZERO(err, "ibv_dereg_mr (nonfetching_atomic)"); + + err = ibv_dereg_mr(mr_fetching_atomic); + CHECK_ZERO(err, "ibv_dereg_mr (fetching_atomic)"); + + allocator.deallocate((void*)nonfetching_atomic); + allocator.deallocate((void*)fetching_atomic); + + fetching_atomic_freelist->~FreeListT(); + allocator.deallocate((void*)fetching_atomic_freelist); +} /****************************************************************************** ************************ PROVIDER-SPECIFIC HELPERS *************************** @@ -232,7 +248,6 @@ __device__ void QueuePair::quiet_internal(uint64_t activemask, uint32_t cons) { } #endif // GDA_IONIC -#ifndef GDA_BNXT #ifdef GDA_IONIC __device__ void QueuePair::ring_doorbell(uint32_t pos) { // TODO When threads write at once to the same address, not all writes reach the bus. @@ -244,7 +259,9 @@ __device__ void QueuePair::ring_doorbell(uint32_t pos) { } __threadfence(); } -#else // !GDA_IONIC +#endif + +#if defined(GDA_MLX5) __device__ void QueuePair::ring_doorbell(uint64_t db_val, uint64_t my_sq_counter) { swap_endian_store(const_cast(dbrec), (uint32_t)my_sq_counter); __atomic_signal_fence(__ATOMIC_SEQ_CST); @@ -254,15 +271,15 @@ __device__ void QueuePair::ring_doorbell(uint64_t db_val, uint64_t my_sq_counter db_uint ^= 0x100; __hip_atomic_store(&db.uint, db_uint, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); } -#endif // !GDA_IONIC -#endif // !GDA_BNXT +#endif // GDA_MLX5 -#ifndef GDA_BNXT #ifdef GDA_IONIC __device__ void QueuePair::quiet() { quiet_internal(get_same_qp_lane_mask(), sq_prod); } -#else // !GDA_IONIC +#endif + +#if defined(GDA_MLX5) __device__ void QueuePair::quiet() { constexpr size_t BROADCAST_SIZE = 1024 / WF_SIZE; __shared__ uint64_t wqe_broadcast[BROADCAST_SIZE]; @@ -344,10 +361,8 @@ __device__ void QueuePair::quiet() { } } } -#endif // !GDA_IONIC -#endif // !GDA_BNXT +#endif // GDA_MLX5 -#ifndef GDA_BNXT #ifdef GDA_IONIC __device__ void QueuePair::post_wqe_rma(int pe, int32_t size, uintptr_t *laddr, uintptr_t *raddr, uint8_t opcode) { uint64_t activemask = get_same_qp_lane_mask(); @@ -392,7 +407,9 @@ __device__ void QueuePair::post_wqe_rma(int pe, int32_t size, uintptr_t *laddr, commit_sq(is_last_active_lane(activemask), my_sq_prod, num_wqes, wqe); } -#else // !GDA_IONIC +#endif + +#if defined (GDA_MLX5) __device__ void QueuePair::post_wqe_rma(int pe, int32_t size, uintptr_t *laddr, uintptr_t *raddr, uint8_t opcode) { uint64_t activemask = get_active_lane_mask(); uint8_t num_active_lanes = get_active_lane_count(activemask); @@ -446,10 +463,8 @@ __device__ void QueuePair::post_wqe_rma(int pe, int32_t size, uintptr_t *laddr, __hip_atomic_store(&sq_db_touched, wave_sq_counter + num_wqes, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); } } -#endif // !GDA_IONIC -#endif // !GDA_BNXT +#endif // GDA_MLX5 -#ifndef GDA_BNXT #ifdef GDA_IONIC __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) { @@ -510,7 +525,9 @@ __device__ uint64_t QueuePair::post_wqe_amo(int pe, int32_t size, uintptr_t *rad } return ret; } -#else // !GDA_IONIC || !GDA_BNXT +#endif + +#if defined(GDA_MLX5) __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) { uint64_t activemask = get_active_lane_mask(); @@ -598,8 +615,7 @@ __device__ uint64_t QueuePair::post_wqe_amo(int pe, int32_t size, uintptr_t *rad } return ret; } -#endif // !GDA_IONIC -#endif // !GDA_BNXT +#endif // GDA_MLX5 /****************************************************************************** ****************************** SHMEM INTERFACE ******************************* diff --git a/src/gda/queue_pair.hpp b/src/gda/queue_pair.hpp index eac8107648..24516c1520 100644 --- a/src/gda/queue_pair.hpp +++ b/src/gda/queue_pair.hpp @@ -44,8 +44,10 @@ extern "C" { } #elif defined(GDA_BNXT) #include "bnxt/provider_gda_bnxt.hpp" -#else +#elif defined(GDA_MLX5) #include +#else +#error "Please select an RDMA provider" #endif #include "containers/free_list.hpp" @@ -56,7 +58,7 @@ extern "C" { #define GDA_OP_RDMA_WRITE IONIC_V2_OP_RDMA_WRITE #define GDA_OP_ATOMIC_FA IONIC_V2_OP_ATOMIC_FA #define GDA_OP_ATOMIC_CS IONIC_V2_OP_ATOMIC_CS -#elif !defined(GDA_BNXT) +#elif defined(GDA_MLX5) #define GDA_MAX_ATOMIC 1 #define GDA_OP_RDMA_WRITE MLX5_OPCODE_RDMA_WRITE #define GDA_OP_ATOMIC_FA MLX5_OPCODE_ATOMIC_FA @@ -85,6 +87,11 @@ class QueuePair { */ explicit QueuePair(struct ibv_pd* pd); + /** + * @brief Destructor. + */ + virtual ~QueuePair(); + /** * @brief Create and enqueue a non-blocking put work queue entry (wqe). * @@ -224,7 +231,7 @@ class QueuePair { struct bnxt_device_sq sq; __device__ int poll_cq(); -#else // !GDA_IONIC && !GDA_BNXT +#else // GDA_MLX5 db_reg_t db{}; @@ -303,6 +310,9 @@ class QueuePair { FreeListT* fetching_atomic_freelist{nullptr}; HIPAllocator allocator{}; + + struct ibv_mr *mr_nonfetching_atomic; + struct ibv_mr *mr_fetching_atomic; }; } // namespace rocshmem