From a4d799805dff1a88acd684f8651e2e925e59a933 Mon Sep 17 00:00:00 2001 From: unionxic <187358711+unionxic@users.noreply.github.com> Date: Thu, 22 Oct 2026 06:07:54 +0900 Subject: [PATCH] ibgda: Trap on error CQEs in release builds ibgda_poll_cq() detects an error CQE and returns -1, but its two callers, ibgda_quiet() and ibgda_wait_for_slot_availability(), only check the result with assert(), which release builds compile out. After a peer dies, nvshmem_quiet() returns once the IB retries run out, later quiets return as if complete, fetching operations return wrong values, and nothing is reported. Trap right after the existing assert in both callers, so release builds fail the kernel at the error CQE as debug builds already do. The host sees the failure as a CUDA error. Debug builds are unchanged. Limitations: after the trap the CUDA context is unusable; neither NIC handler reads the CQ on the host; the GDAKI device path is not changed. Fixes #120 Signed-off-by: unionxic <187358711+unionxic@users.noreply.github.com> --- src/include/non_abi/device/pt-to-pt/ibgda_device.cuh | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/include/non_abi/device/pt-to-pt/ibgda_device.cuh b/src/include/non_abi/device/pt-to-pt/ibgda_device.cuh index 7fcf3b8e..1035002d 100644 --- a/src/include/non_abi/device/pt-to-pt/ibgda_device.cuh +++ b/src/include/non_abi/device/pt-to-pt/ibgda_device.cuh @@ -1731,6 +1731,10 @@ ibgda_quiet(nvshmemi_ibgda_device_qp_t *qp) { } #endif assert(likely(status == 0)); + // CQE errors are fatal (see ibgda_poll_cq); release builds compile the assert out. + if (unlikely(status != 0)) { + __trap(); + } return prod_idx; } @@ -1749,6 +1753,10 @@ __device__ NVSHMEMI_STATIC NVSHMEMI_DEVICE_ALWAYS_INLINE void ibgda_wait_for_slo printf("ibgda_poll_cq failed with error=%d.\n", err); } assert(likely(status == 0)); + // CQE errors are fatal (see ibgda_poll_cq); release builds compile the assert out. + if (unlikely(status != 0)) { + __trap(); + } } IBGDA_MFENCE(); }