(JIT) Detect CCQE support in runtime instead of compiling

examples/local_rdma_ops/atomic_gpu.cpp

@@ -72,18 +72,24 @@ __device__ void SendThreadKernel(RdmaEndpoint& epSend, RdmaMemoryRegion sendMr,
 }
 
 __device__ void RecvThreadKernel(RdmaEndpoint& epRecv, RdmaMemoryRegion mr) {
-  uint32_t postIdx = 0;
   uint32_t* addr = reinterpret_cast<uint32_t*>(mr.addr);
   uint32_t val = core::AtomicLoadSeqCst(addr);
   printf("val = %u\n", val);
-  while (val != 2) {
+
+  // Cross-block, lock-free observation: there is no sync between the send block
+  // (which issues CAS then FETCH_ADD) and this recv block, so by the time we
+  // start polling either both atomics or only CAS may have landed.  Just wait
+  // until the value leaves its initial 0, and then until it reaches the final
+  // expected sum (CAS_swap + FETCH_ADD_value == 2 + 2 == 4).
+  while (val == 0) {
     val = core::AtomicLoadSeqCst(addr);
-    printf("after compare and swap val = %u\n", val);
   }
+  printf("after compare and swap val = %u\n", val);
+
   while (val != 4) {
     val = core::AtomicLoadSeqCst(addr);
-    printf("after fetch add val = %u\n", val);
   }
+  printf("after fetch add val = %u\n", val);
 }
 
 __global__ void SendRecvOnGpu(RdmaEndpoint& epSend, RdmaEndpoint& epRecv, RdmaMemoryRegion mrSend,

examples/local_rdma_ops/send_recv_gpu.cpp

@@ -53,9 +53,17 @@ __device__ void SendThreadKernel(RdmaEndpoint& epSend, RdmaMemoryRegion mr, int
     printf("RingDoorbell is done\n");
     __threadfence_system();
 
-    int snd_opcode =
-        PollCq<P>(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, &epSend.cqHandle.consIdx);
-    printf("send PollCq is done\n");
+    // PSD 4-arg PollCq is non-blocking (returns -1 when the CCQE msg_msn isn't
+    // there yet); BNXT/MLX5 spin internally.  Wrap in a busy-wait loop so this
+    // example works uniformly across all providers.
+    uint32_t snd_wqeIdx = 0;
+    int snd_opcode;
+    do {
+      snd_opcode = PollCq<P>(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum,
+                             &epSend.cqHandle.consIdx, &snd_wqeIdx);
+    } while (snd_opcode < 0);
+    epSend.cqHandle.consIdx += 1;
+    printf("send PollCq is done, wqeIdx %u\n", snd_wqeIdx);
     UpdateCqDbrRecord<P>(epSend.cqHandle, epSend.cqHandle.consIdx);
     printf("send UpdateCqDbrRecord is done\n");
     // printf("snd_opcode %d val %d\n", snd_opcode, reinterpret_cast<char*>(mrSend.addr)[0]);
@@ -84,9 +92,14 @@ __device__ void RecvThreadKernel(RdmaEndpoint& epRecv, RdmaMemoryRegion mr, int
       printf("recv RingDoorbell is done\n");
     }
 
-    int rcv_opcode =
-        PollCq<P>(epRecv.cqHandle.cqAddr, epRecv.cqHandle.cqeNum, &epRecv.cqHandle.consIdx);
-    printf("recv PollCq is done\n");
+    uint32_t rcv_wqeIdx = 0;
+    int rcv_opcode;
+    do {
+      rcv_opcode = PollCq<P>(epRecv.cqHandle.cqAddr, epRecv.cqHandle.cqeNum,
+                             &epRecv.cqHandle.consIdx, &rcv_wqeIdx);
+    } while (rcv_opcode < 0);
+    epRecv.cqHandle.consIdx += 1;
+    printf("recv PollCq is done, wqeIdx %u\n", rcv_wqeIdx);
     UpdateCqDbrRecord<P>(epRecv.cqHandle, epRecv.cqHandle.consIdx);
     printf("recv UpdateCqDbrRecord is done\n");
 

examples/local_rdma_ops/write_inline_gpu.cpp

@@ -53,11 +53,19 @@ __device__ void SendThreadKernel(RdmaEndpoint& epSend, RdmaMemoryRegion mr) {
     RingDoorbell<P>(epSend.wqHandle.dbrAddr, dbr_val);
     __threadfence_system();
 
-    int opcode =
-        PollCq<P>(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, &epSend.cqHandle.consIdx);
+    // PSD 4-arg PollCq is non-blocking (returns -1 when CQE / CCQE msg_msn not
+    // ready yet), while BNXT/MLX5 internally spin.  Wrap in a busy-wait loop so
+    // this example works uniformly across all providers.
+    uint32_t wqeIdx = 0;
+    int opcode;
+    do {
+      opcode = PollCq<P>(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, &epSend.cqHandle.consIdx,
+                         &wqeIdx);
+    } while (opcode < 0);
+    epSend.cqHandle.consIdx += 1;
     __threadfence_system();
     UpdateCqDbrRecord<P>(epSend.cqHandle, epSend.cqHandle.consIdx);
-    // printf("round %d snd_opcode %d\n", i, opcode);
+    // printf("round %d snd_opcode %d wqeIdx %u\n", i, opcode, wqeIdx);
 
     raddr += i;
   }

include/mori/application/transport/rdma/providers/dv_loader.hpp

@@ -181,6 +181,8 @@ struct IonicDvApi {
   using pd_set_sqcmb_t = int (*)(struct ibv_pd*, bool, bool, bool);
   using pd_set_rqcmb_t = int (*)(struct ibv_pd*, bool, bool, bool);
   using pd_set_udma_mask_t = int (*)(struct ibv_pd*, uint32_t);
+  using create_cq_ex_t = struct ibv_cq_ex* (*)(struct ibv_context*, struct ibv_cq_init_attr_ex*,
+                                               struct ionic_cq_init_attr_ex*);
 
   get_ctx_t get_ctx = nullptr;
   qp_get_udma_idx_t qp_get_udma_idx = nullptr;
@@ -189,6 +191,7 @@ struct IonicDvApi {
   pd_set_sqcmb_t pd_set_sqcmb = nullptr;
   pd_set_rqcmb_t pd_set_rqcmb = nullptr;
   pd_set_udma_mask_t pd_set_udma_mask = nullptr;
+  create_cq_ex_t create_cq_ex = nullptr;
 
   void* handle = nullptr;
 
@@ -203,7 +206,9 @@ struct IonicDvApi {
     pd_set_sqcmb = (pd_set_sqcmb_t)DvLoadSymbol(handle, "ionic_dv_pd_set_sqcmb");
     pd_set_rqcmb = (pd_set_rqcmb_t)DvLoadSymbol(handle, "ionic_dv_pd_set_rqcmb");
     pd_set_udma_mask = (pd_set_udma_mask_t)DvLoadSymbol(handle, "ionic_dv_pd_set_udma_mask");
+    create_cq_ex = (create_cq_ex_t)DvLoadSymbol(handle, "ionic_dv_create_cq_ex");
 
+    // create_cq_ex is optional: nullptr means CCQE not supported by this driver version
     return get_ctx && qp_get_udma_idx && get_cq && get_qp && pd_set_sqcmb && pd_set_rqcmb &&
            pd_set_udma_mask;
   }

include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp

@@ -26,7 +26,6 @@ namespace core {
 
 #define QUEUE_SIZE 1
 #define MAX_INLINE_SIZE 32
-// #define IONIC_CCQE           1
-#undef IONIC_CCQE
+
 }  // namespace core
 }  // namespace mori