From fe0246ff632222598e44332ebd566ab4ba25a040 Mon Sep 17 00:00:00 2001 From: Jessie Yang Date: Mon, 29 Jun 2026 11:13:02 -0700 Subject: [PATCH 1/7] [v2.6.x] prov/efa: Update ibv_qp_attach_comp_cntr_op names https://github.com/linux-rdma/rdma-core/pull/1701 renamed ibv_comp_cntr_attach_op to ibv_qp_attach_comp_cntr_op. Signed-off-by: Jessie Yang (cherry picked from commit c2a24013fc9d6e3f5f234ffa22eb0defe4857de0) --- prov/efa/src/efa_base_ep.c | 18 +++++++++--------- prov/efa/test/efa_unit_test_mocks.c | 6 +++--- prov/efa/test/efa_unit_test_mocks.h | 8 ++++---- 3 files changed, 16 insertions(+), 16 deletions(-) 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_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 From 5b14b8a5d3ae2092ded37c6e98e8fc409d0d6705 Mon Sep 17 00:00:00 2001 From: Jessie Yang Date: Fri, 26 Jun 2026 13:59:42 -0700 Subject: [PATCH 2/7] [v2.6.x] fabtest/efa: add pytest marker for hw_cntr Deselect test at pytest collect time instead of skipping. Signed-off-by: Jessie Yang (cherry picked from commit 2c8b645f2b5d54808e241309ab74a35c19c961c6) --- fabtests/pytest/efa/conftest.py | 17 +++++++++++++++++ fabtests/pytest/efa/test_hw_cntr.py | 15 ++++----------- fabtests/pytest/pytest.ini | 1 + 3 files changed, 22 insertions(+), 11 deletions(-) diff --git a/fabtests/pytest/efa/conftest.py b/fabtests/pytest/efa/conftest.py index 7e9c3683480..315f30c8008 100644 --- a/fabtests/pytest/efa/conftest.py +++ b/fabtests/pytest/efa/conftest.py @@ -1,3 +1,4 @@ +import os import pytest import time from common import test_selected_by_marker @@ -246,6 +247,22 @@ def cuda_validation_fixture(request, cmdline_args): @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 "" + have_hw_cntr = os.path.exists(os.path.join(binpath, "fi_efa_hw_cntr")) + + 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) + 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_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..4e8c76077c7 100644 --- a/fabtests/pytest/pytest.ini +++ b/fabtests/pytest/pytest.ini @@ -16,6 +16,7 @@ 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) junit_suite_name = fabtests junit_logging = all junit_log_passing_tests = true From d50be4e7454b48da5391d7a4e411ac68c169ce4d Mon Sep 17 00:00:00 2001 From: Jessie Yang Date: Mon, 1 Jun 2026 15:14:05 -0700 Subject: [PATCH 3/7] [v2.6.x] fabtests/efa: Test hardware cntr in GDA fabtests Create hardware counters on cuda memory, exported with dmabuf. Test the counters are updated correctly with GDA for send, recv, read, write, and writedata. Use gda marker to deselect tests not built with efagda. Signed-off-by: Jessie Yang (cherry picked from commit 011b15c7f1659126ddee0a9a983c7f9ffced1cd2) --- fabtests/prov/efa/src/efa_gda.c | 102 +++++++++++++-- .../prov/efa/src/efagda/efa_gda_kernels.cu | 117 ++++++++++++++---- .../prov/efa/src/efagda/efa_gda_kernels.h | 4 + fabtests/pytest/efa/conftest.py | 3 + fabtests/pytest/efa/test_gda.py | 79 ++++++++---- fabtests/pytest/pytest.ini | 1 + 6 files changed, 251 insertions(+), 55 deletions(-) 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 315f30c8008..b911580317d 100644 --- a/fabtests/pytest/efa/conftest.py +++ b/fabtests/pytest/efa/conftest.py @@ -249,6 +249,7 @@ 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 "" have_hw_cntr = os.path.exists(os.path.join(binpath, "fi_efa_hw_cntr")) + have_gda = os.path.exists(os.path.join(binpath, "fi_efa_gda")) deselected = [] remaining = [] @@ -256,6 +257,8 @@ def pytest_collection_modifyitems(session, config, 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) 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/pytest.ini b/fabtests/pytest/pytest.ini index 4e8c76077c7..df6ee3deddd 100644 --- a/fabtests/pytest/pytest.ini +++ b/fabtests/pytest/pytest.ini @@ -17,6 +17,7 @@ markers = 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 From 7f3cfc4940553799cd6becc99f70942d7f9509e8 Mon Sep 17 00:00:00 2001 From: Jessie Yang Date: Mon, 20 Jul 2026 13:56:56 -0700 Subject: [PATCH 4/7] [v2.6.x] prov/efa: Fix missing argument of efa_ibv_post_write efa_ibv_post_write takes inline_data_list and use_inline. Signed-off-by: Jessie Yang (cherry picked from commit 5e32da00185dc46b8d0a0256eba4cd69f187b82d) --- prov/efa/test/efa_unit_test_rma.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/prov/efa/test/efa_unit_test_rma.c b/prov/efa/test/efa_unit_test_rma.c index 29c88b74793..f8d66fab4fd 100644 --- a/prov/efa/test/efa_unit_test_rma.c +++ b/prov/efa/test/efa_unit_test_rma.c @@ -776,7 +776,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); From 11676ee8cab3bd6af6366db15fee06db957a62e6 Mon Sep 17 00:00:00 2001 From: Jessie Yang Date: Mon, 20 Jul 2026 14:18:33 -0700 Subject: [PATCH 5/7] =?UTF-8?q?[v2.6.x]=20prov/efa:=20remove=20unused=20va?= =?UTF-8?q?riable=20=E2=80=98enable=5Fhigh=5Fpps=5Forig=E2=80=99?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Jessie Yang (cherry picked from commit 6e4ea7be5a96cc37042172345b23ad2cdb057c73) --- prov/efa/test/efa_unit_test_rma.c | 1 - 1 file changed, 1 deletion(-) diff --git a/prov/efa/test/efa_unit_test_rma.c b/prov/efa/test/efa_unit_test_rma.c index f8d66fab4fd..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); From f24ee307ca5db02a2e6dd492c92b1916617c47e5 Mon Sep 17 00:00:00 2001 From: Jessie Yang Date: Mon, 20 Jul 2026 14:09:14 -0700 Subject: [PATCH 6/7] [v2.6.x] prov/efa: Fix hw counter unit test bugs Mock device comp_count_max_value and err_count_max_value instead of max_comp_cntr for hardware counter support. max_comp_cntr is the max number of counters supported, not the max value of the counter. Add macro and helper function efa_unit_test_set_hw_cntr_max_values. Signed-off-by: Jessie Yang (cherry picked from commit 7f2f9750924f4dff2fa2b420eb7bb3f448fab62b) --- prov/efa/test/efa_unit_test_cntr.c | 22 ++++++---------------- prov/efa/test/efa_unit_test_domain.c | 4 +--- prov/efa/test/efa_unit_test_info.c | 18 +++++++++--------- prov/efa/test/efa_unit_tests.h | 15 +++++++++++++++ 4 files changed, 31 insertions(+), 28 deletions(-) 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_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) { From 3f81a2204d11c747f115e3f39f813f6118a965b6 Mon Sep 17 00:00:00 2001 From: Jessie Yang Date: Mon, 20 Jul 2026 15:10:51 -0700 Subject: [PATCH 7/7] [v2.6.x] fabtests/efa: Deselect hardware counter tests if device does not support it Check both the RDMA core symbol and cntr_cnt in fi_info for device support. Signed-off-by: Jessie Yang (cherry picked from commit 933315e6207f1dfa5e06a16073f406283748ede9) --- fabtests/pytest/efa/conftest.py | 42 ++++++++++++++++++++++++++++++++- 1 file changed, 41 insertions(+), 1 deletion(-) diff --git a/fabtests/pytest/efa/conftest.py b/fabtests/pytest/efa/conftest.py index b911580317d..48ec1c221b4 100644 --- a/fabtests/pytest/efa/conftest.py +++ b/fabtests/pytest/efa/conftest.py @@ -1,7 +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, @@ -244,11 +251,44 @@ 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 = []