AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Unify common BNXT and MLX5 initialization code (#233)

Browse Source 
Co-authored-by: Aurelien Bouteiller <aurelien.bouteiller@amd.com>

[ROCm/rocshmem commit: cb39f7a313]
This commit is contained in:
Yiltan
2025-09-10 09:13:36 -04:00
committed by GitHub
parent e30b76e81b
commit b79b9f4b60
10 changed files with 743 additions and 737 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocshmem/CMakeLists.txt
+2 -1
View File
@@ -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
projects/rocshmem/cmake/rocshmem_config.h.in
+1
View File
@@ -44,3 +44,4 @@
#cmakedefine USE_HDP_FLUSH_HOST_SIDE
#cmakedefine GDA_IONIC
#cmakedefine GDA_BNXT
#cmakedefine GDA_MLX5
projects/rocshmem/scripts/build_configs/gda_mlx5
+1 -1
View File
@@ -27,4 +27,4 @@ set -e
script_path=$(dirname "$(realpath $0)")
exec $script_path/gda $*
source $script_path/gda -DGDA_MLX5=ON $*
projects/rocshmem/src/gda/CMakeLists.txt
+4 -3
View File
@@ -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()
projects/rocshmem/src/gda/backend_gda.cpp
+341 -416
View File
@@ -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<char*>(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<char*>(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<uint64_t*>(dvctx.db_page);
uint64_t *gpu_db_page_u64 = reinterpret_cast<uint64_t*>(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 <typename StateType>
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<InitQPState>(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<RtrState>(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<RtsState>(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<uint64_t*>(dvctx.db_page);
uint64_t *gpu_db_page_u64 = reinterpret_cast<uint64_t*>(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<ionic_v1_wqe *>(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, &current_gid);
err = ibv_query_gid(context, port, i, &current_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
projects/rocshmem/src/gda/backend_gda.hpp
+78 -120
View File
@@ -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<ibv_qp*> qps;
std::vector<ibv_cq*> cqs;
std::vector<dest_info_t> 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<struct bnxt_host_qp> bnxt_qps;
std::vector<struct bnxt_host_cq> 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 <typename T>
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_t> 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<ibv_qp*> qps;
std::vector<ibv_cq*> 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<struct bnxt_host_qp> bnxt_qps;
std::vector<struct bnxt_host_cq> bnxt_cqs;
struct bnxt_re_dv_db_region_attr db_region_attr;
#endif
};
} // namespace rocshmem
projects/rocshmem/src/gda/bnxt/backend_gda_bnxt.cpp
+8 -154
View File
@@ -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");
}
}
projects/rocshmem/src/gda/debug_gda.hpp
+240
View File
@@ -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_ */
projects/rocshmem/src/gda/queue_pair.cpp
+55 -39
View File
@@ -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<uint32_t*>(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 *******************************
projects/rocshmem/src/gda/queue_pair.hpp
+13 -3
View File
@@ -44,8 +44,10 @@ extern "C" {
}
#elif defined(GDA_BNXT)
#include "bnxt/provider_gda_bnxt.hpp"
#else
#elif defined(GDA_MLX5)
#include <infiniband/mlx5dv.h>
#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