UCT/DEVICE: support no delay

Author: JeynmannZ

src/uct/ib/mlx5/gdaki/gdaki.cuh

@@ -263,7 +263,7 @@ UCS_F_DEVICE void uct_rc_mlx5_gda_db(uct_rc_gdaki_dev_ep_t *ep,
 {
     cuda::atomic_ref<uint64_t, cuda::thread_scope_device> ref(
             ep->sq_ready_index);
-    const uint64_t wqe_num = __ldg(&ep->sq_wqe_num);
+    const bool no_delay = (flags & UCT_DEVICE_FLAG_NODELAY);
     const uint64_t wqe_next = wqe_base + count;
     const uint64_t wqe_base_orig = wqe_base;
 
@@ -273,9 +273,17 @@ UCS_F_DEVICE void uct_rc_mlx5_gda_db(uct_rc_gdaki_dev_ep_t *ep,
         wqe_base = wqe_base_orig;
     }
 
+    if (no_delay) {
+        const uint64_t ready_index = ep->sq_ready_index;
+        uct_rc_mlx5_gda_ring_db(ep, ready_index);
+        uct_rc_mlx5_gda_update_dbr(ep, ready_index);
+        uct_rc_mlx5_gda_ring_db(ep, ready_index);
+        return;
+    }
+
     if (READ_ONCE(ep->sq_ready_index) == wqe_next) {
         uct_rc_mlx5_gda_lock(&ep->sq_lock);
-        const uint64_t ready_index = READ_ONCE(ep->sq_ready_index);
+        const uint64_t ready_index = ep->sq_ready_index;
         if (ready_index == wqe_next) {
             uct_rc_mlx5_gda_ring_db(ep, ready_index);
             uct_rc_mlx5_gda_update_dbr(ep, ready_index);