fix: avoid indefinite hang in gin gdaki init

[dromniscience]

Description

Fix indefinite hang in GIN GDAKI init when ring-neighbor NICs are unreachable

Problem: ncclGinIbAllGather spins forever if the RDMA fabric between neighboring
ranks is disconnected (e.g., NICs on disjoint rails). The receiver's CTS write never
lands, and the loop has no timeout, error CQE, or abort flag, so jobs hang silently in
ncclDevCommCreate until killed externally.

Solution: This PR adds a per-ring-step deadline checked in both the post and completion
loops, gated by NCCL_GIN_IB_ALLGATHER_TIMEOUT_SEC (default 180s, <=0 disables).
On expiry, it issues a WARN with ring-neighbor rank info and returns
ncclRemoteError through existing cleanup paths. This works for all IB collectives during
devComm setup since they call this allgather underneath (allToAll, rkey exchanges, P2P barrier).

Related Issues

None

Changes & Impact

No change of any public or internal NCCL APIs

Changes:

• Adds NCCL_PARAM(GinIbAllGatherTimeoutSec, "GIN_IB_ALLGATHER_TIMEOUT_SEC", 180)
• Adds ncclGinIbAllGatherCheckDeadline() helper
• Inserts deadline checks into both while loops of ncclGinIbAllGather()
• No API changes, no breaking changes
• Backward compatible: set NCCL_GIN_IB_ALLGATHER_TIMEOUT_SEC=0 to disable this check

Performance Impact

• clockNano() is called once per ncclNetIb.isend/irecv in ncclGinIbAllGather() - negligible
• No measurable overhead in the normal (connected) path
• Tested: verified fail-fast behavior on cross-rail disconnected NICs

[kgioioso]

Hi @dromniscience Thanks for the PR.

I want to make sure we fully understand the problem before implementing a solution. I'm a bit stuck in a few places:

1. Why were the QPs able to connect if they are on disconnected rails? In most cases, if NICs are not connected, we get a "Failed to modify QP" error. But here, it seems like we are past the QP connection phase?
2. Why is there no error CQE? If an RDMA transfer fails, I would expect an error CQE on at least the initiator QP. Is it that there is an error CQE on the initiator but not the target? Or is there really no error CQE?
3. Why are we attempting to create cross-rail connections? Are you setting NCCL_CROSS_NIC=0?

[dromniscience]

Thanks for the timely reply! The linked issue mentions the exact hanging issue I experienced.

The issue surfaces in a disconnected rail network topology. Without NCCL_CROSS_NIC=0, the whole group hangs silently in this while loop containing ncclNetIb.isend and irecv. Such silent hang is barely acceptable.

Regarding your questions:

1. Yes, we are past the QP establishment phase. For IB, I think ibv_modify_qp INIT→RTR→RTS does not involve RDMA communication with remote peers. I do not observe a "Failed to modify QP" error.
2. Due to receiver-driven net_ib comm, only the receiver will get a retry-exceeded CQE, and the sender never posts a WR before seeing a CTS. In the disconnected rail setup, all ranks may get stuck in checking srequest == NULLand they never get the chance to enter the next while loop that checks CQE status.
3. Yes, NCCL_CROSS_NIC=0 avoids the hang. However, with cross-NIC allowed, GIN connects the full world ring, and the hang appears. This PR addresses the failure mode when it isn't set. This PR avoids this silent init hang with explicit WARN message.

[kgioioso]

I agree with your analysis and I agree that we should fix the silent hang.

I'm not super convinced about the timeout, for a few reasons:

• Timeouts are notoriously difficult to tune and it would be unmanageable to add a timeout to every part of the code that may hang (how does every timeout environment variable interact?)
• This timeout diverges from the typical NCCL design pattern. Normally, if one side hangs due to an error on the other side, users should (1) detect the error (2) share the error to all ranks and (3) call abort on all ranks
• The error CQE reports information that is useful for debugging.

I think we can fix this with the error-then-abort pattern. What do you think? This would require 2 changes (1) we need to poll for CQEs even if the sender does not yet have a request and (2) we need to abort if abort is called

This is much larger scope than adding a timeout. If you want to give it a shot, it's appreciated, but otherwise we can add it to our queue.

[dromniscience]

Understood. I agree that the error-then-abort pattern is a better fix, and it's a broader change that may affect other parts of the code. I may use this temporary solution myself, and am grateful to be notified when the problem is fixed in a future release.