diff --git a/fabtests/prov/efa/src/efa_gda.c b/fabtests/prov/efa/src/efa_gda.c index e4254491b66..3f761150137 100644 --- a/fabtests/prov/efa/src/efa_gda.c +++ b/fabtests/prov/efa/src/efa_gda.c @@ -49,6 +49,53 @@ void *send_cq_buffer, *recv_cq_buffer; struct fid_ep *gda_ep; struct efa_cuda_qp *gda_qp; static enum ibv_wr_opcode gda_op = IBV_WR_SEND; +static bool use_hw_cntr; +static volatile uint64_t *hw_send_cntr_ptr; +static volatile uint64_t *hw_recv_cntr_ptr; + +enum { + LONG_OPT_USE_HW_CNTR, +}; + +static int create_hw_cntr(struct fid_cntr **cntr, + volatile uint64_t **cntr_ptr) +{ + int ret, dmabuf_fd; + uint64_t dmabuf_offset; + void *cuda_buf; + struct fi_cntr_attr attr = {0}; + struct fi_efa_comp_cntr_init_attr efa_attr = {0}; + + ret = ft_hmem_alloc(opts.iface, opts.device, &cuda_buf, + sizeof(uint64_t)); + if (ret) + return ret; + + ret = ft_hmem_get_dmabuf_fd(opts.iface, cuda_buf, sizeof(uint64_t), + &dmabuf_fd, &dmabuf_offset); + if (ret) { + ft_hmem_free(opts.iface, cuda_buf); + return ret; + } + + attr.events = FI_CNTR_EVENTS_COMP; + attr.wait_obj = FI_WAIT_UNSPEC; + + efa_attr.flags = FI_EFA_COMP_CNTR_INIT_WITH_COMP_EXTERNAL_MEM; + efa_attr.comp_cntr_ext_mem.type = FI_EFA_MEMORY_LOCATION_DMABUF; + efa_attr.comp_cntr_ext_mem.dmabuf.fd = dmabuf_fd; + efa_attr.comp_cntr_ext_mem.dmabuf.offset = dmabuf_offset; + + ret = efa_gda_ops->cntr_open_ext(domain, &attr, cntr, NULL, &efa_attr); + if (ret) { + FT_WARN("hw cntr open failed (%s)\n", fi_strerror(-ret)); + ft_hmem_free(opts.iface, cuda_buf); + return ret; + } + + *cntr_ptr = (volatile uint64_t *)cuda_buf; + return 0; +} int create_ext_cq(void **cq_buffer, struct fid_cq **cq_ext, struct efa_cuda_cq **gda_cq) @@ -274,7 +321,9 @@ static int run_bw(struct fi_rma_iov *remote_iov) ft_start(); if (is_client) { /* Client: post writes/writedata/read */ - ret = efagda_run_bw(gda_qp, gda_send_cq, gda_op, + ret = efagda_run_bw(gda_qp, gda_send_cq, + use_hw_cntr ? hw_send_cntr_ptr : NULL, + gda_op, (uintptr_t) (gda_op == IBV_WR_RDMA_READ ? rx_buf : tx_buf), opts.transfer_size, @@ -286,6 +335,7 @@ static int run_bw(struct fi_rma_iov *remote_iov) if (gda_op == IBV_WR_RDMA_WRITE_WITH_IMM || gda_op == IBV_WR_SEND) { ret = efagda_run_bw_recv(gda_qp, gda_recv_cq, + use_hw_cntr ? hw_recv_cntr_ptr : NULL, (uintptr_t) rx_buf, opts.transfer_size, lkey, opts.iterations, rx_depth, @@ -362,10 +412,12 @@ static int run() ft_start(); ret = efagda_run_lat_send(gda_qp, gda_send_cq, gda_recv_cq, - ah, remote_qpn, remote_qkey, - (uintptr_t) rx_buf, opts.transfer_size, lkey, - (uintptr_t) tx_buf, opts.transfer_size, lkey, - opts.iterations, rx_depth, is_client, stream); + use_hw_cntr ? hw_send_cntr_ptr : NULL, + use_hw_cntr ? hw_recv_cntr_ptr : NULL, + ah, remote_qpn, remote_qkey, + (uintptr_t) rx_buf, opts.transfer_size, lkey, + (uintptr_t) tx_buf, opts.transfer_size, lkey, + opts.iterations, rx_depth, is_client, stream); ft_stop(); if (ret) { @@ -400,9 +452,17 @@ int main(int argc, char **argv) if (!hints) return EXIT_FAILURE; - while ((op = getopt(argc, argv, - "vh" ADDR_OPTS INFO_OPTS CS_OPTS API_OPTS)) != -1) { + while ((op = getopt_long(argc, argv, + "vh" ADDR_OPTS INFO_OPTS CS_OPTS API_OPTS, + (struct option[]){ + {"use-hw-cntr", no_argument, NULL, + LONG_OPT_USE_HW_CNTR}, + {0, 0, 0, 0} + }, NULL)) != -1) { switch (op) { + case LONG_OPT_USE_HW_CNTR: + use_hw_cntr = true; + break; default: ft_parse_addr_opts(op, optarg, &opts); ft_parseinfo(op, optarg, hints, &opts); @@ -420,6 +480,8 @@ int main(int argc, char **argv) FT_PRINT_OPTS_USAGE("-o ", "op: msg, write, writedata, read\n"); FT_PRINT_OPTS_USAGE("-v", "Enable data verification"); + FT_PRINT_OPTS_USAGE("--use-hw-cntr", + "Use hardware counter instead of send CQ"); return EXIT_FAILURE; } } @@ -453,6 +515,9 @@ int main(int argc, char **argv) hints->domain_attr->progress = FI_PROGRESS_MANUAL; hints->mode |= FI_CONTEXT | FI_CONTEXT2; + if (use_hw_cntr) + ft_fiversion = FI_VERSION(2, 5); + ret = ft_init_fabric(); if (ret) { FT_PRINTERR("ft_init_fabric", -ret); @@ -489,7 +554,28 @@ int main(int argc, char **argv) return ret; } - ret = ft_enable_ep(gda_ep, eq, av, txcq_ext, rxcq_ext, txcntr, rxcntr, + if (use_hw_cntr) { + ret = create_hw_cntr(&txcntr, &hw_send_cntr_ptr); + if (ret) { + FT_PRINTERR("create_hw_cntr send", -ret); + return ret; + } + ret = create_hw_cntr(&rxcntr, &hw_recv_cntr_ptr); + if (ret) { + FT_PRINTERR("create_hw_cntr recv", -ret); + return ret; + } + ret = fi_ep_bind(gda_ep, &rxcntr->fid, + (gda_op == IBV_WR_RDMA_WRITE_WITH_IMM) ? + FI_REMOTE_WRITE : + FI_RECV); + if (ret) { + FT_PRINTERR("fi_ep_bind rxcntr", -ret); + return ret; + } + } + + ret = ft_enable_ep(gda_ep, eq, av, txcq_ext, rxcq_ext, txcntr, NULL, rma_cntr); if (ret) return ret; diff --git a/fabtests/prov/efa/src/efagda/efa_gda_kernels.cu b/fabtests/prov/efa/src/efagda/efa_gda_kernels.cu index 917d9553ffa..e88343ac78d 100644 --- a/fabtests/prov/efa/src/efagda/efa_gda_kernels.cu +++ b/fabtests/prov/efa/src/efagda/efa_gda_kernels.cu @@ -29,11 +29,25 @@ #include "efa_cuda_dp_impl.cuh" #include "efa_gda_kernels.h" #include +#include + +#define CNTR_POLL_MAX_ITER 100000000 + +static __device__ int efa_cuda_cntr_poll(volatile uint64_t *cntr_ptr, uint64_t target) +{ + for (int i = 0; i < CNTR_POLL_MAX_ITER; i++) { + if (*cntr_ptr >= target) + return 0; + } + return -1; +} __global__ void efagda_lat_send_kernel( efa_cuda_qp *qp, efa_cuda_cq *send_cq, efa_cuda_cq *recv_cq, + volatile uint64_t *send_cntr_ptr, + volatile uint64_t *recv_cntr_ptr, uint16_t ah, uint16_t remote_qpn, uint32_t remote_qkey, @@ -54,11 +68,18 @@ __global__ void efagda_lat_send_kernel( __shared__ efa_cuda_qp local_qp; __shared__ efa_cuda_cq local_send_cq; __shared__ efa_cuda_cq local_recv_cq; + uint64_t send_cntr_start_val = 0; + uint64_t recv_cntr_start_val = 0; local_qp = *qp; local_send_cq = *send_cq; local_recv_cq = *recv_cq; + if (send_cntr_ptr) + send_cntr_start_val = *send_cntr_ptr; + if (recv_cntr_ptr) + recv_cntr_start_val = *recv_cntr_ptr; + /* Post initial receives */ for (int i = 0; i < rx_depth; i++) { if (efa_cuda_post_recv_wr(&local_qp, recv_addr, recv_length, @@ -70,9 +91,16 @@ __global__ void efagda_lat_send_kernel( while (scnt < iters || rcnt < iters) { /* Poll for receive completion (except for first client send) */ if (rcnt < iters && !(scnt < 1 && machine_type == 1)) { - do { - cqe = efa_cuda_cq_poll(&local_recv_cq, 0); - } while (!cqe); + if (recv_cntr_ptr) { + if (efa_cuda_cntr_poll(recv_cntr_ptr, recv_cntr_start_val + rcnt + 1)) { + printf("recv hw cntr timeout at rcnt=%d, recv_cntr_ptr=%lu\n", rcnt, *recv_cntr_ptr); + assert(0); + } + } else { + do { + cqe = efa_cuda_cq_poll(&local_recv_cq, 0); + } while (!cqe); + } rcnt++; efa_cuda_cq_pop(&local_recv_cq, 1); @@ -105,13 +133,20 @@ __global__ void efagda_lat_send_kernel( efa_cuda_flush_sq_wrs(&local_qp); /* Wait for send completion */ - do { - cqe = efa_cuda_cq_poll(&local_send_cq, 0); - } while (!cqe); - - uint32_t err = efa_cuda_wc_read_vendor_err(cqe); - if (err) - printf("send comp err %d\n", err); + if (send_cntr_ptr) { + if (efa_cuda_cntr_poll(send_cntr_ptr, send_cntr_start_val + scnt)) { + printf("send hw cntr timeout at scnt=%d, send_cntr_ptr=%lu\n", scnt, *send_cntr_ptr); + assert(0); + } + } else { + do { + cqe = efa_cuda_cq_poll(&local_send_cq, 0); + } while (!cqe); + + uint32_t err = efa_cuda_wc_read_vendor_err(cqe); + if (err) + printf("send comp err %d\n", err); + } efa_cuda_cq_pop(&local_send_cq, 1); } } @@ -124,6 +159,8 @@ __global__ void efagda_lat_send_kernel( int efagda_run_lat_send(struct efa_cuda_qp *qp, struct efa_cuda_cq *send_cq, struct efa_cuda_cq *recv_cq, + volatile uint64_t *send_cntr_ptr, + volatile uint64_t *recv_cntr_ptr, uint16_t ah, uint16_t remote_qpn, uint32_t remote_qkey, uint64_t recv_addr, uint32_t recv_length, @@ -136,7 +173,9 @@ int efagda_run_lat_send(struct efa_cuda_qp *qp, cudaError_t err; efagda_lat_send_kernel<<<1, 1, 0, stream>>>( - qp, send_cq, recv_cq, ah, remote_qpn, remote_qkey, + qp, send_cq, recv_cq, + send_cntr_ptr, recv_cntr_ptr, + ah, remote_qpn, remote_qkey, recv_addr, recv_length, recv_lkey, send_addr, send_length, send_lkey, iters, rx_depth, is_client); @@ -161,6 +200,7 @@ int efagda_run_lat_send(struct efa_cuda_qp *qp, __global__ void efagda_bw_kernel( efa_cuda_qp *qp, efa_cuda_cq *send_cq, + volatile uint64_t *send_cntr_ptr, enum ibv_wr_opcode opcode, uint64_t send_addr, uint32_t send_length, @@ -179,10 +219,15 @@ __global__ void efagda_bw_kernel( struct efa_io_tx_wqe wr_buf; __shared__ efa_cuda_qp local_qp; __shared__ efa_cuda_cq local_send_cq; + uint64_t send_cntr_start_val = 0; local_qp = *qp; local_send_cq = *send_cq; + if (send_cntr_ptr) { + send_cntr_start_val = *send_cntr_ptr; + } + while (scnt < iters || ccnt < iters) { /* Post writes up to tx_depth */ while (scnt < iters && (scnt - ccnt) < tx_depth) { @@ -227,13 +272,22 @@ __global__ void efagda_bw_kernel( /* Poll completions */ while (ccnt < scnt && (scnt == iters || (scnt - ccnt) >= tx_depth)) { - cqe = efa_cuda_cq_poll(&local_send_cq, 0); - if (cqe) { - if (((efa_io_cdesc_common *)cqe)->status != 0) - printf("bw comp err %d\n", - ((efa_io_cdesc_common *)cqe)->status); + if (send_cntr_ptr) { + if (efa_cuda_cntr_poll(send_cntr_ptr, send_cntr_start_val + ccnt + 1)) { + printf("bw send hw cntr timeout at ccnt=%d, send_cntr_ptr=%lu\n", ccnt, *send_cntr_ptr); + assert(0); + } efa_cuda_cq_pop(&local_send_cq, 1); ccnt++; + } else { + cqe = efa_cuda_cq_poll(&local_send_cq, 0); + if (cqe) { + if (((efa_io_cdesc_common *)cqe)->status != 0) + printf("bw comp err %d\n", + ((efa_io_cdesc_common *)cqe)->status); + efa_cuda_cq_pop(&local_send_cq, 1); + ccnt++; + } } } } @@ -245,6 +299,7 @@ __global__ void efagda_bw_kernel( __global__ void efagda_bw_recv_kernel( efa_cuda_qp *qp, efa_cuda_cq *recv_cq, + volatile uint64_t *recv_cntr_ptr, uint64_t recv_addr, uint32_t recv_length, uint32_t recv_lkey, @@ -255,10 +310,14 @@ __global__ void efagda_bw_recv_kernel( void *cqe; __shared__ efa_cuda_qp local_qp; __shared__ efa_cuda_cq local_recv_cq; + uint64_t recv_cntr_start_val = 0; local_qp = *qp; local_recv_cq = *recv_cq; + if (recv_cntr_ptr) + recv_cntr_start_val = *recv_cntr_ptr; + /* Post initial receives */ for (int i = 0; i < rx_depth; i++) { if (efa_cuda_post_recv_wr(&local_qp, recv_addr, recv_length, @@ -268,13 +327,20 @@ __global__ void efagda_bw_recv_kernel( efa_cuda_flush_rq_wrs(&local_qp); while (rcnt < iters) { - do { - cqe = efa_cuda_cq_poll(&local_recv_cq, 0); - } while (!cqe); + if (recv_cntr_ptr) { + if (efa_cuda_cntr_poll(recv_cntr_ptr, recv_cntr_start_val + (rcnt + 1))) { + printf("bw recv hw cntr timeout at rcnt=%d, recv_cntr_ptr=%lu\n", rcnt, *recv_cntr_ptr); + assert(0); + } + } else { + do { + cqe = efa_cuda_cq_poll(&local_recv_cq, 0); + } while (!cqe); - if (((efa_io_cdesc_common *)cqe)->status != 0) - printf("bw recv err %d\n", - ((efa_io_cdesc_common *)cqe)->status); + if (((efa_io_cdesc_common *)cqe)->status != 0) + printf("bw recv err %d\n", + ((efa_io_cdesc_common *)cqe)->status); + } rcnt++; efa_cuda_cq_pop(&local_recv_cq, 1); @@ -294,6 +360,7 @@ __global__ void efagda_bw_recv_kernel( int efagda_run_bw(struct efa_cuda_qp *qp, struct efa_cuda_cq *send_cq, + volatile uint64_t *send_cntr_ptr, enum ibv_wr_opcode opcode, uint64_t send_addr, uint32_t send_length, uint32_t send_lkey, uint16_t ah, uint32_t remote_qpn, uint32_t remote_qkey, @@ -304,7 +371,7 @@ int efagda_run_bw(struct efa_cuda_qp *qp, cudaError_t err; efagda_bw_kernel<<<1, 1, 0, stream>>>( - qp, send_cq, opcode, + qp, send_cq, send_cntr_ptr, opcode, send_addr, send_length, send_lkey, ah, remote_qpn, remote_qkey, remote_addr, remote_rkey, @@ -329,6 +396,7 @@ int efagda_run_bw(struct efa_cuda_qp *qp, int efagda_run_bw_recv(struct efa_cuda_qp *qp, struct efa_cuda_cq *recv_cq, + volatile uint64_t *recv_cntr_ptr, uint64_t recv_addr, uint32_t recv_length, uint32_t recv_lkey, int iters, int rx_depth, @@ -337,7 +405,8 @@ int efagda_run_bw_recv(struct efa_cuda_qp *qp, cudaError_t err; efagda_bw_recv_kernel<<<1, 1, 0, stream>>>( - qp, recv_cq, recv_addr, recv_length, recv_lkey, + qp, recv_cq, recv_cntr_ptr, + recv_addr, recv_length, recv_lkey, iters, rx_depth); err = cudaGetLastError(); diff --git a/fabtests/prov/efa/src/efagda/efa_gda_kernels.h b/fabtests/prov/efa/src/efagda/efa_gda_kernels.h index f0d959072fa..eee63c71ed3 100644 --- a/fabtests/prov/efa/src/efagda/efa_gda_kernels.h +++ b/fabtests/prov/efa/src/efagda/efa_gda_kernels.h @@ -41,6 +41,8 @@ extern "C" { int efagda_run_lat_send(struct efa_cuda_qp *qp, struct efa_cuda_cq *send_cq, struct efa_cuda_cq *recv_cq, + volatile uint64_t *send_cntr_ptr, + volatile uint64_t *recv_cntr_ptr, uint16_t ah, uint16_t remote_qpn, uint32_t remote_qkey, uint64_t recv_addr, uint32_t recv_length, @@ -52,6 +54,7 @@ int efagda_run_lat_send(struct efa_cuda_qp *qp, int efagda_run_bw(struct efa_cuda_qp *qp, struct efa_cuda_cq *send_cq, + volatile uint64_t *send_cntr_ptr, enum ibv_wr_opcode opcode, uint64_t send_addr, uint32_t send_length, uint32_t send_lkey, uint16_t ah, uint32_t remote_qpn, uint32_t remote_qkey, @@ -61,6 +64,7 @@ int efagda_run_bw(struct efa_cuda_qp *qp, int efagda_run_bw_recv(struct efa_cuda_qp *qp, struct efa_cuda_cq *recv_cq, + volatile uint64_t *recv_cntr_ptr, uint64_t recv_addr, uint32_t recv_length, uint32_t recv_lkey, int iters, int rx_depth, diff --git a/fabtests/pytest/efa/conftest.py b/fabtests/pytest/efa/conftest.py index 7e9c3683480..48ec1c221b4 100644 --- a/fabtests/pytest/efa/conftest.py +++ b/fabtests/pytest/efa/conftest.py @@ -1,6 +1,14 @@ +import os import pytest +import subprocess import time -from common import test_selected_by_marker +from retrying import retry +from common import ( + test_selected_by_marker, + has_ssh_connection_err_msg, + is_ssh_connection_error, + SshConnectionError, +) from efa_common import ( has_rdma, support_cq_interrupts, @@ -243,9 +251,61 @@ def cuda_validation_fixture(request, cmdline_args): print("No CUDA memory mark, skipping validation") +@retry(retry_on_exception=is_ssh_connection_error, stop_max_attempt_number=3, wait_fixed=5000) +def device_has_hw_cntr(host_id): + """ + Return True if the EFA device reports a non-zero hardware + counter count, indiciating hardware counter support + """ + command = "ssh {} 'fi_info -p efa -v || /opt/amazon/efa/bin/fi_info -p efa -v' | grep cntr_cnt".format(host_id) + proc = subprocess.run(command, shell=True, + stdout=subprocess.PIPE, stderr=subprocess.PIPE, + encoding="utf-8", timeout=60) + + if has_ssh_connection_err_msg(proc.stdout) or has_ssh_connection_err_msg(proc.stderr): + raise SshConnectionError() + + for line in proc.stdout.strip().split("\n"): + value = line.split(":")[-1].strip() + if value.isdigit() and int(value) > 0: + return True + return False + + @pytest.hookimpl(hookwrapper=True) def pytest_collection_modifyitems(session, config, items): # Called after collection has been performed, may filter or re-order the items in-place + binpath = config.getoption("--binpath", default="") or "" + server_id = config.getoption("--server-id", default="") + client_id = config.getoption("--client-id", default="") + have_hw_cntr = os.path.exists(os.path.join(binpath, "fi_efa_hw_cntr")) + if have_hw_cntr: + for host_id in filter(None, (server_id, client_id)): + try: + if not device_has_hw_cntr(host_id): + have_hw_cntr = False + break + except SshConnectionError: + pytest.fail( + "Could not determine hw_cntr support: ssh to {} failed after " + "retries. Refusing to silently deselect hw_cntr tests.".format(host_id)) + have_gda = os.path.exists(os.path.join(binpath, "fi_efa_gda")) + + deselected = [] + remaining = [] + for item in items: + markers = {m.name for m in item.iter_markers()} + if "hw_cntr" in markers and not have_hw_cntr: + deselected.append(item) + elif "gda" in markers and not have_gda: + deselected.append(item) + else: + remaining.append(item) + + if deselected: + config.hook.pytest_deselected(items=deselected) + items[:] = remaining + # We use this hook to always run the MR exhaustion test at the end mr_exhaustion_tests, other_tests = [], [] for item in items: diff --git a/fabtests/pytest/efa/test_gda.py b/fabtests/pytest/efa/test_gda.py index b15e9c12d3c..7a302cd9efb 100644 --- a/fabtests/pytest/efa/test_gda.py +++ b/fabtests/pytest/efa/test_gda.py @@ -1,29 +1,20 @@ -import os import pytest from common import ClientServerTest, has_cuda from efa.efa_common import DIRECT_SIZES, DIRECT_RMA_SIZES - -def _skip_if_gda_unavailable(cmdline_args): - binpath = cmdline_args.binpath or "" - if not os.path.exists(os.path.join(binpath, "fi_efa_gda")): - pytest.skip("fi_efa_gda is not built") - - if not cmdline_args.do_dmabuf_reg_for_hmem: - pytest.skip("DMABUF is required for GDA tests") - - if not has_cuda(cmdline_args.client_id) or not has_cuda(cmdline_args.server_id): - pytest.skip("Client and server both need a cuda device") +# Apply to all functions in this file +pytestmark = [ + pytest.mark.short, + pytest.mark.functional, + pytest.mark.cuda_memory, + pytest.mark.gda, + pytest.mark.fabric(params=["efa-direct"]), +] @pytest.mark.pr_ci -@pytest.mark.fabric(params=["efa-direct"]) @pytest.mark.message_sizes(default_efa_direct=DIRECT_SIZES, pr_ci_efa_direct=DIRECT_SIZES) -@pytest.mark.short -@pytest.mark.functional -@pytest.mark.cuda_memory def test_gda_send_recv(cmdline_args, message_sizes, fabric): - _skip_if_gda_unavailable(cmdline_args) test = ClientServerTest(cmdline_args, "fi_efa_gda", message_size=message_sizes, iteration_type="short", @@ -40,8 +31,6 @@ def test_gda_send_recv(cmdline_args, message_sizes, fabric): @pytest.mark.functional @pytest.mark.cuda_memory def test_gda_write_bw(cmdline_args, message_sizes, rma_fabric): - _skip_if_gda_unavailable(cmdline_args) - test = ClientServerTest(cmdline_args, "fi_efa_gda -o write", message_size=message_sizes, iteration_type="short", @@ -58,8 +47,6 @@ def test_gda_write_bw(cmdline_args, message_sizes, rma_fabric): @pytest.mark.functional @pytest.mark.cuda_memory def test_gda_writedata_bw(cmdline_args, message_sizes, rma_fabric): - _skip_if_gda_unavailable(cmdline_args) - test = ClientServerTest(cmdline_args, "fi_efa_gda -o writedata", message_size=message_sizes, iteration_type="short", @@ -76,8 +63,6 @@ def test_gda_writedata_bw(cmdline_args, message_sizes, rma_fabric): @pytest.mark.functional @pytest.mark.cuda_memory def test_gda_read_bw(cmdline_args, message_sizes, rma_fabric): - _skip_if_gda_unavailable(cmdline_args) - test = ClientServerTest(cmdline_args, "fi_efa_gda -o read", message_size=message_sizes, iteration_type="short", @@ -85,3 +70,51 @@ def test_gda_read_bw(cmdline_args, message_sizes, rma_fabric): fabric=rma_fabric, datacheck_type="with_datacheck") test.run() + + +@pytest.mark.hw_cntr +@pytest.mark.message_sizes(default_efa_direct=DIRECT_SIZES) +def test_gda_send_recv_hw_cntr(cmdline_args, message_sizes, fabric): + test = ClientServerTest(cmdline_args, "fi_efa_gda --use-hw-cntr", + message_size=message_sizes, + iteration_type="short", + memory_type="cuda_to_cuda", + fabric=fabric, + additional_env="FI_EFA_USE_HW_CNTR=1") + test.run() + + +@pytest.mark.hw_cntr +@pytest.mark.message_sizes(default_efa_direct=DIRECT_SIZES) +def test_gda_write_bw_hw_cntr(cmdline_args, message_sizes, rma_fabric): + test = ClientServerTest(cmdline_args, "fi_efa_gda -o write --use-hw-cntr", + message_size=message_sizes, + iteration_type="short", + memory_type="cuda_to_cuda", + fabric=rma_fabric, + additional_env="FI_EFA_USE_HW_CNTR=1") + test.run() + + +@pytest.mark.hw_cntr +@pytest.mark.message_sizes(default_efa_direct=DIRECT_SIZES) +def test_gda_writedata_bw_hw_cntr(cmdline_args, message_sizes, rma_fabric): + test = ClientServerTest(cmdline_args, "fi_efa_gda -o writedata --use-hw-cntr", + message_size=message_sizes, + iteration_type="short", + memory_type="cuda_to_cuda", + fabric=rma_fabric, + additional_env="FI_EFA_USE_HW_CNTR=1") + test.run() + + +@pytest.mark.hw_cntr +@pytest.mark.message_sizes(default_efa_direct=DIRECT_SIZES) +def test_gda_read_bw_hw_cntr(cmdline_args, message_sizes, rma_fabric): + test = ClientServerTest(cmdline_args, "fi_efa_gda -o read --use-hw-cntr", + message_size=message_sizes, + iteration_type="short", + memory_type="cuda_to_cuda", + fabric=rma_fabric, + additional_env="FI_EFA_USE_HW_CNTR=1") + test.run() diff --git a/fabtests/pytest/efa/test_hw_cntr.py b/fabtests/pytest/efa/test_hw_cntr.py index cc1abdc7b14..2e98547f641 100644 --- a/fabtests/pytest/efa/test_hw_cntr.py +++ b/fabtests/pytest/efa/test_hw_cntr.py @@ -1,15 +1,8 @@ -import os import pytest from common import ClientServerTest from efa.efa_common import DIRECT_SIZES, DIRECT_RMA_SIZES -@pytest.fixture(autouse=True) -def skip_if_not_built(cmdline_args): - binpath = cmdline_args.binpath or "" - if not os.path.exists(os.path.join(binpath, "fi_efa_hw_cntr")): - pytest.skip("fi_efa_hw_cntr requires efadv_create_comp_cntr") - - +@pytest.mark.hw_cntr @pytest.mark.fabric(params=["efa-direct"]) @pytest.mark.message_sizes(default_efa_direct=DIRECT_SIZES) @pytest.mark.short @@ -22,7 +15,7 @@ def test_efa_hw_cntr_pingpong(cmdline_args, message_sizes, fabric): additional_env="FI_EFA_USE_HW_CNTR=1") test.run() - +@pytest.mark.hw_cntr @pytest.mark.fabric(params=["efa-direct"]) @pytest.mark.message_sizes(default_efa_direct=DIRECT_RMA_SIZES) @pytest.mark.short @@ -35,7 +28,7 @@ def test_efa_hw_cntr_rma_write(cmdline_args, message_sizes, rma_fabric): additional_env="FI_EFA_USE_HW_CNTR=1") test.run() - +@pytest.mark.hw_cntr @pytest.mark.fabric(params=["efa-direct"]) @pytest.mark.message_sizes(default_efa_direct=DIRECT_SIZES) @pytest.mark.short @@ -48,7 +41,7 @@ def test_efa_hw_cntr_pingpong_ext_mem(cmdline_args, message_sizes, fabric): additional_env="FI_EFA_USE_HW_CNTR=1") test.run() - +@pytest.mark.hw_cntr @pytest.mark.fabric(params=["efa-direct"]) @pytest.mark.message_sizes(default_efa_direct=DIRECT_RMA_SIZES) @pytest.mark.short diff --git a/fabtests/pytest/pytest.ini b/fabtests/pytest/pytest.ini index 28a2285e334..df6ee3deddd 100644 --- a/fabtests/pytest/pytest.ini +++ b/fabtests/pytest/pytest.ini @@ -16,6 +16,8 @@ markers = pr_ci: test is included in the PR CI suite message_sizes: decorator specifying default and pr_ci message size lists fabric: decorator declaring the set of libfabric fabric values a test runs on + hw_cntr: test requires fi_efa_hw_cntr binary (built with efadv_create_comp_cntr) + gda: test requires fi_efa_gda binary (built with --enable-efagda) junit_suite_name = fabtests junit_logging = all junit_log_passing_tests = true diff --git a/prov/efa/src/efa_base_ep.c b/prov/efa/src/efa_base_ep.c index 2e125233cde..e98efc6c331 100644 --- a/prov/efa/src/efa_base_ep.c +++ b/prov/efa/src/efa_base_ep.c @@ -451,16 +451,16 @@ static int efa_base_ep_create_qp(struct efa_base_ep *base_ep, } /** - * @brief Map ofi_cntr_index to ibv_comp_cntr_attach_op bit + * @brief Map ofi_cntr_index to ibv_qp_attach_comp_cntr_op bit */ #if HAVE_EFADV_CREATE_COMP_CNTR static const uint32_t efa_cntr_index_to_ibv_op[] = { - [CNTR_TX] = IBV_COMP_CNTR_ATTACH_OP_SEND, - [CNTR_RX] = IBV_COMP_CNTR_ATTACH_OP_RECV, - [CNTR_RD] = IBV_COMP_CNTR_ATTACH_OP_RDMA_READ, - [CNTR_WR] = IBV_COMP_CNTR_ATTACH_OP_RDMA_WRITE, - [CNTR_REM_RD] = IBV_COMP_CNTR_ATTACH_OP_REMOTE_RDMA_READ, - [CNTR_REM_WR] = IBV_COMP_CNTR_ATTACH_OP_REMOTE_RDMA_WRITE, + [CNTR_TX] = IBV_QP_ATTACH_COMP_CNTR_OP_SEND, + [CNTR_RX] = IBV_QP_ATTACH_COMP_CNTR_OP_RECV, + [CNTR_RD] = IBV_QP_ATTACH_COMP_CNTR_OP_RDMA_READ, + [CNTR_WR] = IBV_QP_ATTACH_COMP_CNTR_OP_RDMA_WRITE, + [CNTR_REM_RD] = IBV_QP_ATTACH_COMP_CNTR_OP_REMOTE_RDMA_READ, + [CNTR_REM_WR] = IBV_QP_ATTACH_COMP_CNTR_OP_REMOTE_RDMA_WRITE, }; /** @@ -478,7 +478,7 @@ static int efa_base_ep_attach_comp_cntrs(struct efa_base_ep *base_ep, { struct util_cntr *util_cntr; struct efa_cntr *efa_cntr; - struct ibv_comp_cntr_attach_attr attr; + struct ibv_qp_attach_comp_cntr_attr attr; int i, err; for (i = 0; i < CNTR_CNT; i++) { @@ -489,7 +489,7 @@ static int efa_base_ep_attach_comp_cntrs(struct efa_base_ep *base_ep, efa_cntr = container_of(util_cntr, struct efa_cntr, util_cntr); if (!efa_cntr->ibv_comp_cntr) continue; - attr = (struct ibv_comp_cntr_attach_attr){ + attr = (struct ibv_qp_attach_comp_cntr_attr){ .comp_mask = 0, .op_mask = efa_cntr_index_to_ibv_op[i], }; diff --git a/prov/efa/test/efa_unit_test_cntr.c b/prov/efa/test/efa_unit_test_cntr.c index c7fb6936628..f5780a22519 100644 --- a/prov/efa/test/efa_unit_test_cntr.c +++ b/prov/efa/test/efa_unit_test_cntr.c @@ -261,9 +261,7 @@ void test_efa_hw_cntr_open_unsupported_type_bytes(void **state) efa_domain = container_of(resource->domain, struct efa_domain, util_domain.domain_fid); - efa_domain->device->max_comp_cntr = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_cntr_value = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_err_cntr_value = (1ULL << 31) - 1; + efa_unit_test_set_hw_cntr_max_values(efa_domain); g_efa_unit_test_mocks.efadv_create_comp_cntr = efa_mock_efadv_create_comp_cntr_return_null_enotsup; @@ -294,7 +292,7 @@ void test_efa_hw_cntr_open_max_cntr_value_exceeded(void **state) efa_domain = container_of(resource->domain, struct efa_domain, util_domain.domain_fid); - efa_domain->device->max_comp_cntr = (1ULL << 31) - 1; + efa_domain->device->comp_count_max_value = EFA_HW_CNTR_MAX_VALUE; efa_domain->info->domain_attr->max_cntr_value = UINT64_MAX; cntr = calloc(1, sizeof(*cntr)); @@ -322,9 +320,7 @@ void test_efa_hw_cntr_open_ibv_fail(void **state) efa_domain = container_of(resource->domain, struct efa_domain, util_domain.domain_fid); - efa_domain->device->max_comp_cntr = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_cntr_value = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_err_cntr_value = (1ULL << 31) - 1; + efa_unit_test_set_hw_cntr_max_values(efa_domain); g_efa_unit_test_mocks.efadv_create_comp_cntr = efa_mock_efadv_create_comp_cntr_return_null_enotsup; @@ -365,9 +361,7 @@ static struct fid_cntr *test_efa_hw_cntr_open(struct efa_resource *resource, efa_domain = container_of(resource->domain, struct efa_domain, util_domain.domain_fid); - efa_domain->device->max_comp_cntr = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_cntr_value = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_err_cntr_value = (1ULL << 31) - 1; + efa_unit_test_set_hw_cntr_max_values(efa_domain); g_efa_unit_test_mocks.efadv_create_comp_cntr = efa_mock_efadv_create_comp_cntr_return_mock; g_efa_unit_test_mocks.ibv_destroy_comp_cntr = efa_mock_ibv_destroy_comp_cntr_return_mock; @@ -504,9 +498,7 @@ static int test_efa_hw_cntr_open_with_wait_obj(struct efa_resource *resource, efa_domain = container_of(resource->domain, struct efa_domain, util_domain.domain_fid); - efa_domain->device->max_comp_cntr = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_cntr_value = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_err_cntr_value = (1ULL << 31) - 1; + efa_unit_test_set_hw_cntr_max_values(efa_domain); g_efa_unit_test_mocks.efadv_create_comp_cntr = efa_mock_efadv_create_comp_cntr_return_mock; g_efa_unit_test_mocks.ibv_destroy_comp_cntr = efa_mock_ibv_destroy_comp_cntr_return_mock; @@ -658,9 +650,7 @@ void test_efa_cntr_open_uses_hw_cntr(void **state) efa_domain = container_of(resource->domain, struct efa_domain, util_domain.domain_fid); - efa_domain->device->max_comp_cntr = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_cntr_value = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_err_cntr_value = (1ULL << 31) - 1; + efa_unit_test_set_hw_cntr_max_values(efa_domain); g_efa_unit_test_mocks.efadv_create_comp_cntr = efa_mock_efadv_create_comp_cntr_return_mock; g_efa_unit_test_mocks.ibv_destroy_comp_cntr = efa_mock_ibv_destroy_comp_cntr_return_mock; diff --git a/prov/efa/test/efa_unit_test_domain.c b/prov/efa/test/efa_unit_test_domain.c index 4783fbfdd0d..116cdd3d568 100644 --- a/prov/efa/test/efa_unit_test_domain.c +++ b/prov/efa/test/efa_unit_test_domain.c @@ -847,9 +847,7 @@ void test_efa_domain_open_ops_cntr_open_ext(void **state) struct efa_domain *efa_domain; efa_domain = container_of(resource->domain, struct efa_domain, util_domain.domain_fid); - efa_domain->device->max_comp_cntr = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_cntr_value = (1ULL << 31) - 1; - efa_domain->info->domain_attr->max_err_cntr_value = (1ULL << 31) - 1; + efa_unit_test_set_hw_cntr_max_values(efa_domain); g_efa_unit_test_mocks.efadv_create_comp_cntr = efa_mock_efadv_create_comp_cntr_return_mock; g_efa_unit_test_mocks.ibv_destroy_comp_cntr = efa_mock_ibv_destroy_comp_cntr_return_mock; } diff --git a/prov/efa/test/efa_unit_test_info.c b/prov/efa/test/efa_unit_test_info.c index 84dcff54f0d..219ba40cda1 100644 --- a/prov/efa/test/efa_unit_test_info.c +++ b/prov/efa/test/efa_unit_test_info.c @@ -1217,7 +1217,7 @@ void test_info_max_cntr_value_api_lt_2_5(void **state) struct fi_info *hints, *info; int err; - mock_cntr_max_values((1ULL << 31) - 1); + mock_cntr_max_values(EFA_HW_CNTR_MAX_VALUE); hints = efa_unit_test_alloc_hints(FI_EP_RDM, EFA_DIRECT_FABRIC_NAME); assert_non_null(hints); @@ -1241,7 +1241,7 @@ void test_info_max_cntr_value_api_ge_2_5_within_hw_range(void **state) struct fi_info *hints, *info; int err; - mock_cntr_max_values((1ULL << 31) - 1); + mock_cntr_max_values(EFA_HW_CNTR_MAX_VALUE); hints = efa_unit_test_alloc_hints(FI_EP_RDM, EFA_DIRECT_FABRIC_NAME); assert_non_null(hints); @@ -1249,8 +1249,8 @@ void test_info_max_cntr_value_api_ge_2_5_within_hw_range(void **state) err = fi_getinfo(FI_VERSION(2, 5), NULL, NULL, 0, hints, &info); assert_int_equal(err, 0); - assert_int_equal(info->domain_attr->max_cntr_value, (1ULL << 31) - 1); - assert_int_equal(info->domain_attr->max_err_cntr_value, (1ULL << 31) - 1); + assert_int_equal(info->domain_attr->max_cntr_value, EFA_HW_CNTR_MAX_VALUE); + assert_int_equal(info->domain_attr->max_err_cntr_value, EFA_HW_CNTR_MAX_VALUE); fi_freeinfo(info); fi_freeinfo(hints); @@ -1266,7 +1266,7 @@ void test_info_max_cntr_value_api_ge_2_5_hint_within_hw_range(void **state) struct fi_info *hints, *info; int err; - mock_cntr_max_values((1ULL << 31) - 1); + mock_cntr_max_values(EFA_HW_CNTR_MAX_VALUE); hints = efa_unit_test_alloc_hints(FI_EP_RDM, EFA_DIRECT_FABRIC_NAME); assert_non_null(hints); @@ -1275,8 +1275,8 @@ void test_info_max_cntr_value_api_ge_2_5_hint_within_hw_range(void **state) err = fi_getinfo(FI_VERSION(2, 5), NULL, NULL, 0, hints, &info); assert_int_equal(err, 0); - assert_int_equal(info->domain_attr->max_cntr_value, (1ULL << 31) - 1); - assert_int_equal(info->domain_attr->max_err_cntr_value, (1ULL << 31) - 1); + assert_int_equal(info->domain_attr->max_cntr_value, EFA_HW_CNTR_MAX_VALUE); + assert_int_equal(info->domain_attr->max_err_cntr_value, EFA_HW_CNTR_MAX_VALUE); fi_freeinfo(info); fi_freeinfo(hints); @@ -1292,7 +1292,7 @@ void test_info_max_cntr_value_api_ge_2_5_above_hw_range(void **state) struct fi_info *hints, *info; int err; - mock_cntr_max_values((1ULL << 31) - 1); + mock_cntr_max_values(EFA_HW_CNTR_MAX_VALUE); hints = efa_unit_test_alloc_hints(FI_EP_RDM, EFA_DIRECT_FABRIC_NAME); assert_non_null(hints); @@ -1316,7 +1316,7 @@ void test_info_rdm_max_cntr_value_api_ge_2_5_within_hw_range(void **state) struct fi_info *hints, *info; int err; - mock_cntr_max_values((1ULL << 31) - 1); + mock_cntr_max_values(EFA_HW_CNTR_MAX_VALUE); hints = efa_unit_test_alloc_hints(FI_EP_RDM, EFA_FABRIC_NAME); assert_non_null(hints); diff --git a/prov/efa/test/efa_unit_test_mocks.c b/prov/efa/test/efa_unit_test_mocks.c index 0d7a94b5e39..4f7fa53664f 100644 --- a/prov/efa/test/efa_unit_test_mocks.c +++ b/prov/efa/test/efa_unit_test_mocks.c @@ -883,21 +883,21 @@ int efa_mock_ibv_destroy_comp_cntr_return_mock(struct ibv_comp_cntr *cntr) int __wrap_ibv_qp_attach_comp_cntr(struct ibv_qp *qp, struct ibv_comp_cntr *comp_cntr, - struct ibv_comp_cntr_attach_attr *attr) + struct ibv_qp_attach_comp_cntr_attr *attr) { return g_efa_unit_test_mocks.ibv_qp_attach_comp_cntr(qp, comp_cntr, attr); } int efa_mock_ibv_qp_attach_comp_cntr_return_mock(struct ibv_qp *qp, struct ibv_comp_cntr *comp_cntr, - struct ibv_comp_cntr_attach_attr *attr) + struct ibv_qp_attach_comp_cntr_attr *attr) { return 0; } int efa_mock_ibv_qp_attach_comp_cntr_return_enotsup(struct ibv_qp *qp, struct ibv_comp_cntr *comp_cntr, - struct ibv_comp_cntr_attach_attr *attr) + struct ibv_qp_attach_comp_cntr_attr *attr) { return ENOTSUP; } diff --git a/prov/efa/test/efa_unit_test_mocks.h b/prov/efa/test/efa_unit_test_mocks.h index 41583d5661d..958ffe2b246 100644 --- a/prov/efa/test/efa_unit_test_mocks.h +++ b/prov/efa/test/efa_unit_test_mocks.h @@ -274,7 +274,7 @@ struct efa_unit_test_mocks int (*ibv_read_err_comp_cntr)(struct ibv_comp_cntr *comp_cntr, uint64_t *value); int (*ibv_qp_attach_comp_cntr)(struct ibv_qp *qp, struct ibv_comp_cntr *comp_cntr, - struct ibv_comp_cntr_attach_attr *attr); + struct ibv_qp_attach_comp_cntr_attr *attr); #endif }; @@ -373,13 +373,13 @@ int efa_mock_ibv_read_comp_cntr_return_mock(struct ibv_comp_cntr *comp_cntr, uin int efa_mock_ibv_read_err_comp_cntr_return_mock(struct ibv_comp_cntr *comp_cntr, uint64_t *value); int __real_ibv_qp_attach_comp_cntr(struct ibv_qp *qp, struct ibv_comp_cntr *comp_cntr, - struct ibv_comp_cntr_attach_attr *attr); + struct ibv_qp_attach_comp_cntr_attr *attr); int efa_mock_ibv_qp_attach_comp_cntr_return_mock(struct ibv_qp *qp, struct ibv_comp_cntr *comp_cntr, - struct ibv_comp_cntr_attach_attr *attr); + struct ibv_qp_attach_comp_cntr_attr *attr); int efa_mock_ibv_qp_attach_comp_cntr_return_enotsup(struct ibv_qp *qp, struct ibv_comp_cntr *comp_cntr, - struct ibv_comp_cntr_attach_attr *attr); + struct ibv_qp_attach_comp_cntr_attr *attr); #endif #endif diff --git a/prov/efa/test/efa_unit_test_rma.c b/prov/efa/test/efa_unit_test_rma.c index 29c88b74793..25dcc1dcdfa 100644 --- a/prov/efa/test/efa_unit_test_rma.c +++ b/prov/efa/test/efa_unit_test_rma.c @@ -746,7 +746,6 @@ static void test_efa_ibv_post_write_processing_hints_impl(struct efa_resource *r struct ibv_qp_ex *ibv_qpx; struct efadv_qp *efadv_qp; struct efa_ah fake_ah = {0}; - bool enable_high_pps_orig; efa_unit_test_resource_construct(resource, FI_EP_RDM, EFA_DIRECT_FABRIC_NAME); @@ -776,7 +775,7 @@ static void test_efa_ibv_post_write_processing_hints_impl(struct efa_resource *r sge.length = local_buff.size; sge.lkey = ((struct efa_mr *)fi_mr_desc(local_buff.mr))->ibv_mr->lkey; - efa_ibv_post_write(qp, &sge, 1, 123456, 0x87654321, 0, 0, + efa_ibv_post_write(qp, &sge, 1, NULL, false, 123456, 0x87654321, 0, 0, flags, &fake_ah, 0, 0); efa_unit_test_buff_destruct(&local_buff); diff --git a/prov/efa/test/efa_unit_tests.h b/prov/efa/test/efa_unit_tests.h index 6bf8aa36c7c..533c236200a 100644 --- a/prov/efa/test/efa_unit_tests.h +++ b/prov/efa/test/efa_unit_tests.h @@ -18,6 +18,12 @@ #include "efa.h" #include "efa_unit_test_mocks.h" +/** + * Maximum value of hardware completion counters on the EFA device. + * Used to mock hardware counter support in unit tests. + */ +#define EFA_HW_CNTR_MAX_VALUE (((uint64_t)1 << 31) - 1) + extern int g_ibv_ah_limit; extern int g_ibv_ah_cnt; extern int g_self_ah_cnt; @@ -614,6 +620,15 @@ void test_efa_ibv_post_write_processing_hints_with_high_pps(void **state); void test_efa_ibv_post_write_processing_hints_without_high_pps(void **state); /* end efa_unit_test_rdm_rma.c */ +static inline +void efa_unit_test_set_hw_cntr_max_values(struct efa_domain *efa_domain) +{ + efa_domain->device->comp_count_max_value = EFA_HW_CNTR_MAX_VALUE; + efa_domain->device->err_count_max_value = EFA_HW_CNTR_MAX_VALUE; + efa_domain->info->domain_attr->max_cntr_value = EFA_HW_CNTR_MAX_VALUE; + efa_domain->info->domain_attr->max_err_cntr_value = EFA_HW_CNTR_MAX_VALUE; +} + static inline int efa_unit_test_get_dlist_length(struct dlist_entry *head) {