dev_runtime: restore capture mode on GIN setup failure

[wanglei875]

Description

ncclDevrCommCreateInternal() switches the CUDA stream capture mode before setting up device communicator resources. Most failure paths after that point use GOTO cleanup labels so the original capture mode is restored.

The GIN setup path used NCCLCHECK(ncclGinDevCommSetup(...)), which returns directly on failure and skips the common fail path. This can leave the thread capture mode unrestored when GIN setup fails.

This change routes the failure through the existing fail label by using NCCLCHECKGOTO.

Changes & Impact

• Restore the existing cleanup behavior on ncclGinDevCommSetup() failure.
• No change to the success path.
• No API or performance impact.

Testing

• git diff --check
• Started make -j$(nproc) src.build; src/dev_runtime.cc compiled successfully. Full device kernel build was not completed due to long compile time.

[Copilot]

Pull request overview

Note

Copilot was unable to run its full agentic suite in this review.

Updates error handling in ncclDevrCommCreateInternal to route a setup failure through the common cleanup path.

Changes:

• Switches ncclGinDevCommSetup failure handling from NCCLCHECK to NCCLCHECKGOTO(..., ret, fail).

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.