Guard GIN error query on ginState->connected in ncclCommGetAsyncError

[goelayu]

Summary:

Summary

D92104317 added ncclCommGetAsyncError() polling in the TorchComms timeout watchdog thread to detect asynchronous NCCL errors. This broke the DeviceApiTest.DeviceWindowCreationFloat test when GIN (GPU-Initiated Networking) is enabled with GDAKI (NCCL_GIN_TYPE=3).

Root Cause

ncclCommGetAsyncError() internally calls ncclGinQueryLastError(), which iterates over all GIN contexts and queries GDAKI QP error states via doca_gpu_verbs_query_last_error(). The problem is that GIN context creation (ncclGinConnectOnce → ncclGinGdakiCreateContext) happens lazily — triggered inside any collective call (ncclAllReduce, etc.) that needs new transport channels. During this initialization, IB QPs are still transitioning through states (INIT → RTR → RTS) and doca_gpu_verbs_query_last_error() reports spurious errors that surface as ncclRemoteError.

Since the watchdog thread runs concurrently and polls ncclCommGetAsyncError() on a timer, it races with GIN context creation happening on the main thread inside NCCL collective calls. The watchdog sees the false-positive ncclRemoteError and calls abort(), crashing all 8 ranks.

Fix

Gate the GIN error query block in ncclCommGetAsyncError() on ginState->connected, which is only set to true at the end of ncclGinConnectOnce() after all GIN contexts are fully created and QPs have completed their IB state transitions. This skips the query entirely during the vulnerable initialization window while still checking for real GIN errors once the subsystem is ready.

Differential Revision: D92623789

[meta-codesync]

@goelayu has exported this pull request. If you are a Meta employee, you can view the originating Diff in D92623789.

[meta-codesync]

This pull request has been merged in 8de4962.