diff --git a/projects/rocshmem/src/gda/bnxt/provider_gda_bnxt.hpp b/projects/rocshmem/src/gda/bnxt/provider_gda_bnxt.hpp index ebc079635e..87d26c33f9 100644 --- a/projects/rocshmem/src/gda/bnxt/provider_gda_bnxt.hpp +++ b/projects/rocshmem/src/gda/bnxt/provider_gda_bnxt.hpp @@ -39,10 +39,6 @@ struct bnxt_device_wq { uint32_t tail; uint32_t flags; uint32_t id; - - uint32_t lock; - - uint32_t db_cnt {0}; } __attribute__((packed)); struct bnxt_device_cq : public bnxt_device_wq { @@ -56,6 +52,7 @@ struct bnxt_device_sq : public bnxt_device_wq { uint32_t msn_tbl_sz; uint32_t psn_sz_log2; uint64_t mtu; + uint32_t lock; } __attribute__((packed)); struct bnxt_host_cq { diff --git a/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp b/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp index 72244d95ab..26ad09cfab 100644 --- a/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp +++ b/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp @@ -177,8 +177,6 @@ __device__ void QueuePair::poll_cq_until(uint32_t requested_available_slots) { sq_depth = sq.depth; - aquire_lock(&cq.lock); - do { cqe = (struct bnxt_re_req_cqe *) cq.buf; @@ -197,8 +195,6 @@ __device__ void QueuePair::poll_cq_until(uint32_t requested_available_slots) { consumed_slots = (sq_tail - sq_head + sq_depth) % sq_depth; available_slots = sq_depth - consumed_slots; } while (available_slots < requested_available_slots); - - release_lock(&cq.lock); } __device__ void QueuePair::bnxt_quiet() {