[BugFix] Prevent orphaned process on NCCL destroy

[jeffreywang-anyscale]

Purpose

PyNcclCommunicator.destroy() currently calls ncclCommDestroy (introduced by #37131), which is a collective and deadlocks during uncoordinated shutdown — when one rank's worker has already exited, surviving ranks block forever in destroy(), leaving orphaned GPU processes that a subsequent test/run can't reclaim. This PR switches to ncclCommAbort and runs it in a daemon thread with a join timeout.

Why we need to run ncclCommAbort in a background thread

ncclCommAbort sets the device abort flag so in-flight kernels can exit without peer rendezvous — that's why it solves the original collective-hang. But its async job still funnels through commDestroySync, which spins on:

// nccl/src/init.cc — commDestroySync
while (comm->localPersistentRefs != 0) {
    ncclCommPollCallbacks(comm, /*waitSome=*/true);
}

localPersistentRefs is incremented for every CUDA graph that captured a NCCL op on this comm and is only decremented when the captured graph object is destroyed (CUDA fires the cuUserObject destructor). The abort flag does not drop that ref count. With CUDAGraphMode.FULL_AND_PIECEWISE, vLLM holds dozens of such graphs at shutdown. Those graphs are released later in the same main-thread teardown chain (model runner cleanup, gc, __del__, interpreter shutdown), so a direct main-thread ncclCommAbort self-deadlocks: the thread that needs to free the graphs is blocked waiting for them to be freed.

Prior art

PyTorch's c10d uses the same daemon thread pattern in ProcessGroupNCCL::abort() (code):

void ProcessGroupNCCL::abort() {
  // This will log counter for how long the abort actually takes.
  STATIC_SCOPED_WAIT_COUNTER(pytorch.ProcessGroupNCCL__abort);

  dumpExtraDebuggingInfo();
  // Don't join threads here since the purpose of this method is to abort all
  // communicators and signal the threads to exit. Joining on the threads could
  // potentially block and hence avoid it in this method.
  terminateProcessGroup_.store(true);
  watchdog_->notify();
  // launch abort asynchronously and wait for it to complete or timeout
  LOG(INFO) << logPrefix()
            << "Launching ProcessGroupNCCL abort asynchronously.";
  std::future<bool> fut =
      std::async(std::launch::async, [this]() { return this->abortComms(); });

  ::c10d::C10dLoggingData debugLog;
  waitForFutureOrTimeout(
      fut, options_->timeout, "ProcessGroup abort", debugLog, true);
  LOG(INFO) << logPrefix() << "ProcessGroupNCCL aborts successfully.";

  // We need to wait for abort to finish before we can safely shut down
  // heartbeat monitoring thread.
  heartbeatMonitor_->stop();
}

Related thread: https://github.com/vllm-project/vllm/pull/37131/changes#r2967953725.

Test Plan

This bug presents when instantiating vLLM with the mp backend and TP > 1 inside a Ray actor and relying on Ray for process lifecycle management. After the fix, the failure no longer reproduces.

Test Result

The second test in https://github.com/ray-project/ray/blob/c91063a871f60550a0c226685a6979b8057a9105/release/llm_tests/batch/test_batch_vllm.py#L511-L520 would've failed without the fix.

---

Essential Elements of an Effective PR Description Checklist

• The purpose of the PR, such as "Fix some issue (link existing issues this PR will resolve)".
• The test plan, such as providing test command.
• The test results, such as pasting the results comparison before and after, or e2e results
• (Optional) The necessary documentation update, such as updating supported_models.md and examples for a new model.
• (Optional) Release notes update. If your change is user facing, please update the release notes draft in the Google Doc.

[gemini-code-assist]

Code Review

This pull request modifies the destroy method in pynccl.py to skip the blocking ncclCommDestroy call, aiming to prevent hangs during uncoordinated shutdowns. Feedback indicates that simply nullifying the communicator causes resource leaks in long-running processes and suggests using ncclCommAbort as a safer, non-blocking alternative. Additionally, a race condition was identified where clearing the communicator reference before updating the disabled state could lead to crashes in multi-threaded environments.

[jeffreywang-anyscale]

cc: @itayalroy @tlrmchlsmth

[mergify]

Hi @jeffreywang-anyscale, the pre-commit checks have failed. Please run:

uv pip install pre-commit>=4.5.1
pre-commit install
pre-commit run --all-files

Then, commit the changes and push to your branch.

For future commits, pre-commit will run automatically on changed files before each commit.

Tip

Is mypy failing?

mypy is run differently in CI. If the failure is related to this check, please use the following command to run it locally:

# For mypy (substitute "3.10" with the failing version if needed)
pre-commit run --hook-stage manual mypy-3.10

[mergify]

Hi @jeffreywang-anyscale, the pre-commit checks have failed. Please run:

uv pip install pre-commit>=4.5.1
pre-commit install
pre-commit run --all-files

Then, commit the changes and push to your branch.

For future commits, pre-commit will run automatically on changed files before each commit.

Tip

Is mypy failing?

mypy is run differently in CI. If the failure is related to this check, please use the following command to run it locally:

# For mypy (substitute "3.10" with the failing version if needed)
pre-commit run --hook-stage manual mypy-3.10

[itayalroy]

We introduced the nccl destroy call because these stale nccl comm instances broke repeated scale up/down tests. When scaling down from 4 gpus to 2 for instance, the GPU memory of ranks 2,3 is not freed even after procs 2, 3 are killed since the stale nccl communicators on ranks 0, 1 hold cuda ipc references to it, eventually causing OOM on GPUs 2, 3. Merging this PR will most probably break Elastic EP, I think we need some way to properly clean these nccl comms.

[jeffreywang-anyscale]

When scaling down from 4 gpus to 2 for instance, the GPU memory of ranks 2,3 is not freed even after procs 2, 3 are killed since the stale nccl communicators on ranks 0, 1 hold cuda ipc references to it, eventually causing OOM on GPUs 2, 3. Merging this PR will most probably break Elastic EP, I think we need some way to properly clean these nccl comms.

Thanks @itayalroy, sounds like aborting the nccl collective group is a more promising approach. Summarize the problems here:

• Uncoordinated shutdown: If one of the process shuts down earlier, the remaining processes aren't able to complete nccl_destroy.
• Stale communicators: If we only set the communicator to null, the remaining processes still hold CUDA IPC references to the original communicator.

elastic-ep-scaling-test seems to succeed though: https://buildkite.com/vllm/ci/builds/61651/steps/canvas?jid=019d9702-b357-48d1-931c-db0951680bbe&tab=output#019d9702-b357-48d1-931c-db0951680bbe. Are we missing test coverages?

[itayalroy]

When scaling down from 4 gpus to 2 for instance, the GPU memory of ranks 2,3 is not freed even after procs 2, 3 are killed since the stale nccl communicators on ranks 0, 1 hold cuda ipc references to it, eventually causing OOM on GPUs 2, 3. Merging this PR will most probably break Elastic EP, I think we need some way to properly clean these nccl comms.

Thanks @itayalroy, sounds like aborting the nccl collective group is a more promising approach. Summarize the problems here:

• Uncoordinated shutdown: If one of the process shuts down earlier, the remaining processes aren't able to complete nccl_destroy.
• Stale communicators: If we only set the communicator to null, the remaining processes still hold CUDA IPC references to the original communicator.

elastic-ep-scaling-test seems to succeed though: https://buildkite.com/vllm/ci/builds/61651/steps/canvas?jid=019d9702-b357-48d1-931c-db0951680bbe&tab=output#019d9702-b357-48d1-931c-db0951680bbe. Are we missing test coverages?

elastic-ep-scaling-test runs a single bench -> scale-up -> bench -> scale-down -> bench iteration, which does not leak enough GPU memory to reach OOM. Internally we have similar tests that run multiple iterations, eventually reaching OOM. Such tests are not included in the CI as they take too long to complete (the single iteration test already takes around 5 minutes).

However, I will look into enhancing elastic-ep-scaling-test with some GPU memory verifications before/after scale up/down to catch memory leaks w/o running extremely long tests.

[jeffreywang-anyscale]

@tlrmchlsmth @itayalroy This PR is ready for another pass. I switched to using ncclCommAbort , following PyTorch's approach to tear down a NCCL comm on cleanup.

[jeffreywang-anyscale]

The failed test (entrypoints/openai/chat_completion/test_chat.py::test_chat_completion_n_parameter_non_streaming[HuggingFaceH4/zephyr-7b-beta]) seems to be irrelevant.

[itayalroy]

@rtourgeman has tested repeated scale up/down scenarios with this approach and it works

[jeffreywang-anyscale]

https://buildkite.com/vllm/ci/builds/64131#019de6dd-799e-47e8-ae0c-2a825fdc3436 failure results from the flakiness that #41421 intends to resolved.

[jeffreywang-anyscale]

@tlrmchlsmth @SageMoore this PR should be ready to go. @itayalroy has verified that this won't break Nvidia's internal elastic EP scaling tests.