Repository navigation
Conversation
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 NVIDIA#120 Signed-off-by: unionxic <187358711+unionxic@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes #120.
Summary
In release builds the IBGDA device path ignores error CQEs:
nvshmem_quiet()returnsnormally after a peer dies, later quiets return as if complete, and fetching operations
return wrong values with nothing reported. This change makes release builds fail the
kernel at the error CQE, as debug builds already do through their assert.
Root cause
ibgda_poll_cq()detectsMLX5_CQE_REQ_ERRand returns -1, butibgda_quiet()andibgda_wait_for_slot_availability()check the result only withassert(), which-DNDEBUGremoves. The code comment inibgda_poll_cq()says CQE errors are fatal.Fix
Two lines at each of the two call sites, right after the existing
assert:if (unlikely(status != 0)) { __trap(); }. No structure, host-side or environmentvariable change; the polling loop is untouched. Diagnostic detail stays where it is
today:
NVSHMEM_IBGDA_DEBUGbuilds print the poll error before the assert.Limitations
of the kernel, not as an NVSHMEM message.
NVSHMEM_TIMEOUT_DEVICE_POLLINGbuilds a polling timeout also traps at this point, asthe debug assert already does.
Testing
Builds (CUDA 12.8, sm_75 and sm_86): release with our options and with LTO-IR, Debug
(sm_75), and the header-contract target; warning lists identical to the unpatched tree.
Cluster: 2 nodes, one PE each (sm_75 and sm_86), ConnectX-6 RoCE v2,
NVSHMEM_IB_ENABLE_IBGDA=1 NVSHMEM_IB_TIMEOUT=14 NVSHMEM_IB_RETRY_CNT=7 NVSHMEM_IBGDA_NUM_RC_PER_PE=1, the reproducer attached to #120 and a fetch-add variant of it (quiet loop, 40 iterations250 ms apart; fetch-add loop, 160 iterations 15 ms apart), PE 1 SIGKILLed after iteration 3.
5 trials per case on devel (2dca770) before and after.
unspecified launch failure), exit code 4, no later iterations