From c43dc136f3be83c5158e42ea4f419d4425e671aa Mon Sep 17 00:00:00 2001
From: Omri Mor <Omri.Mor@amd.com>
Date: Mon, 22 Dec 2025 12:05:00 -0800
Subject: [PATCH] [Bugfix] GDA/bnxt: release SQ lock before return (#372)

* bnxt_post_wqe_amo_single with fetching = true would return
  before releasing the send queue lock, resulting in a deadlock.
* Release the send queue lock before returning from the function.

[ROCm/rocshmem commit: 016e08120ae4a01eab0bc5d379fd42308765ec51]
---
 projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp b/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp
index 43f0dae8df..ba63b29c3a 100644
--- a/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp
+++ b/projects/rocshmem/src/gda/bnxt/queue_pair_bnxt.cpp
@@ -434,13 +434,13 @@ __device__ uint64_t QueuePair::bnxt_post_wqe_amo_single(uintptr_t raddr, uint8_t
 
   bnxt_ring_doorbell(sq.tail);
 
+  release_lock(&sq.lock);
+
   if (fetching) {
     quiet();
     return fetching_atomic[atomic_idx];
   }
 
-  release_lock(&sq.lock);
-
   return 0;
 }