Fix ElasticBuffer GIN assertion failure with MLNX_OFED 24.10#680
Fix ElasticBuffer GIN assertion failure with MLNX_OFED 24.10#680Functionhx wants to merge 1 commit into
Conversation
| CUDA_RUNTIME_CHECK(cudaStreamDestroy(stream)); | ||
| CUDA_RUNTIME_CHECK(cudaFree(d_gin)); | ||
| } | ||
|
|
There was a problem hiding this comment.
🔴 critical: The 'graceful fallback to non-GIN path' does not exist for multi-node workloads. buffer.hpp:292 unconditionally calls launch_engram_fetch whenever num_scaleout_ranks > 1, and engram_fetch.cuh:35 unconditionally constructs handle::NCCLGin(nccl_dev_comm, nccl_window, qp_idx, ...) and issues gin.get(). With reqs zeroed here, ncclDevCommCreate succeeds with zero GIN contexts, and the first internode dispatch accesses nonexistent contexts -> illegal memory access / UB inside a JIT'd kernel. This replaces the original clear init-time assertion with a later, opaque device crash: strictly worse debuggability, and for the MLNX_OFED 24.10 users this targets, RDMA still cannot work. The comment on lines 115-117 admits this ('kernel code ... will target nonexistent contexts ... Callers must ensure they do not issue GIN operations') but no such caller-side guard exists. The real fix is either (a) keep a hard, actionable error when GIN is genuinely required (multi-node scaleout, telling users to upgrade OFED), or (b) wire a true runtime non-GIN kernel path. Neither is done.
🤖 v3
| reqs.ginConnectionType = allow_hybrid_mode ? NCCL_GIN_CONNECTION_RAIL: NCCL_GIN_CONNECTION_FULL; | ||
| const int local_gin_type = allow_hybrid_mode ? props.railedGinType : props.ginType; | ||
| int gin_available = (local_gin_type != NCCL_GIN_TYPE_NONE) ? 1 : 0; | ||
|
|
There was a problem hiding this comment.
🔴 critical: Deadlock: this ncclAllReduce consensus is only reached inside if (num_ranks > 1 and get_env("EP_DISABLE_GIN", 0) == 0). In a cluster where some ranks set EP_DISABLE_GIN=1 and others do not (partial rollout / misconfig), the EP_DISABLE_GIN=1 ranks skip the entire block and never call ncclAllReduce, while the other ranks block on it forever. The collective was added to prevent mismatch hangs but introduces a new one. Any collective on an init path must be entered by all ranks unconditionally.
🤖 v3
| // all-reduce to reach consensus. | ||
| { | ||
| int* d_gin; | ||
| cudaStream_t stream; |
There was a problem hiding this comment.
🟡 warning: d_gin (cudaMalloc) and stream (cudaStreamCreate) are not RAII-guarded. If ncclAllReduce fails via NCCL_CHECK (which throws), both the device allocation and the stream leak. Since this is an init path that may be retried, wrap these in a scope guard or free them before the throw path. Also, a dedicated stream + malloc/free for a single-int reduce is heavyweight but acceptable at init.
🤖 v3
| CUDA_RUNTIME_CHECK(cudaFree(d_gin)); | ||
| } | ||
|
|
||
| if (gin_available == 0) { |
There was a problem hiding this comment.
🔵 suggestion: Uses get_env("EP_BUFFER_DEBUG") while the surrounding file also uses get_env("EP_BUFFER_DEBUG", 0). The file is already inconsistent so this isn't newly introduced, but consider unifying. Printing both ginType and railedGinType in the debug message is a genuine improvement.
🤖 v3
🤖 ds-review-bot Code Reviewv4v3This PR replaces a hard Files reviewed: 1 |
The ElasticBuffer constructor (NCCLSymmetricMemoryContext) had a hard assertion that ncclCommQueryProperties returns a non-NONE ginType. With MLNX_OFED 24.10 (libmlx5 < MLX5_1.25), the mlx5dv_reg_dmabuf_mr and mlx5dv_get_data_direct_sysfs_path symbols are unavailable, causing NCCL to report ginType=NCCL_GIN_TYPE_NONE even when RDMA hardware is present. Changes: - Replace EP_HOST_ASSERT with a graceful fallback: when GIN is unavailable, print an informational message (behind EP_BUFFER_DEBUG) and skip GIN configuration. The existing zero-initialized reqs provide a valid non-GIN path. When GIN is available, configure as before. - Add an ncclAllReduce consensus step (ncclMin over a gin_available flag) so all ranks agree on whether to use GIN. Without consensus, heterogeneous environments (some ranks with GIN, some without) would cause ncclDevCommCreate to hang instead of failing cleanly. - Document the coupling between this GIN-configuration path and kernel code that accesses ncclGin(dev_comm, qp_idx) with qp_idx > 0. - Print both ginType and railedGinType in the debug message so the output is correct regardless of allow_hybrid_mode. Test Plan: - Code review: the non-GIN fallback mirrors the existing path taken when EP_DISABLE_GIN=1 or num_ranks <= 1. The all-reduce consensus uses the existing ncclComm_t which already supports collectives. - Verified no cuda_runtime.h include needed (transitively available via compiled.cuh). Signed-off-by: Yuchen Fan <functionhx@gmail.com>
Functionhx
force-pushed
the
fix/mlnx-ofed-24.10-compat
branch
from
64407ae to
f36b0ec
Compare
2 participants
What Problem This Solves
Fixes #628:
EP_HOST_ASSERT(GIN != NONE)crashes on MLNX_OFED 24.10 becauselibmlx5lacks MLX5_1.25 symbols for GIN support.Fix
Instead of asserting, perform a consensus check across all ranks via
ncclAllReduce(ncclMin):EP_DISABLE_GIN=1)EP_BUFFER_DEBUGguard (consistent with file convention)ginTypeandrailedGinTypefor accurate diagnosticsThis prevents both the crash AND the collective mismatch hang that would occur if ranks independently decided GIN state.
Evidence
Single file changed:
csrc/kernels/backend/nccl.cu(+22/-5 lines).🤖 Generated with Claude Code