Fixes NCCLComm::abort() to use correct deregister API for window-registered handles by syed-ahmed · Pull Request #181626 · pytorch/pytorch

syed-ahmed · 2026-04-27T18:49:49Z

NCCLComm::abort() unconditionally called ncclCommDeregister() on all entries in registeredSegmentHandles_. However, handles registered via ncclCommWindowRegister() (symmetric registration) require ncclCommWindowDeregister() instead. Using the wrong API causes NCCL to fail with "Deregister: Could not find handle" followed by a SIGABRT crash during ProcessGroupNCCL destruction.

The fix tracks whether each handle was window-registered by changing registeredSegmentHandles_ from map<void*, void*> to map<void*, pair<void*, bool>>, and dispatches to the correct deregister call in abort(), matching the existing logic in deregisterSegment().

Fixes: #181610

Assistant used: Claude Opus 4.7 (1M)

pytorch-bot · 2026-04-27T18:49:53Z

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/181626

📄 Preview Python docs built from this PR
📄 Preview C++ docs built from this PR
❓ Need help or want to give feedback on the CI? Visit the bot commands wiki

Note: Links to docs will display an error until the docs builds have been completed.

❗ 1 Active SEVs

There are 1 currently active SEVs. If your PR is affected, please view them below:

Rolling out OSDC (ARC) runners on pull & trunk workflows in PyTorch main

✅ No Failures

As of commit 2dc9e57 with merge base fc69434 ():
💚 Looks good so far! There are no failures yet. 💚

This comment was automatically generated by Dr. CI and updates every 15 minutes.

linux-foundation-easycla · 2026-04-27T18:49:57Z

The committers listed above are authorized under a signed CLA.

✅ login: syed-ahmed / name: Syed Tousif Ahmed (2dc9e57)

kwen2501

The code change looks better than making the wrong dereg call.

One thing I am unsure is whether we can guarantee successful wind dereg when we are about to abort a communicator (most likely the communicator is already corrupted).

kwen2501 · 2026-04-27T19:19:20Z

  for (auto& it : registeredSegmentHandles_) {
-    void* handle = it.second;
+    void* handle = it.second.first;
+    bool window = it.second.second;


nit: use is_window so that it is less ambiguous?

Thank you! Addressed.

kwen2501 · 2026-04-27T19:20:59Z

+    void* handle = it.second.first;
+    bool window = it.second.second;
+#ifdef NCCL_HAS_COMM_WINDOW_REGISTER
+    if (window) {


nit: use if (window) as outer condition and #ifdef NCCL_HAS_COMM_WINDOW_REGISTER as inner condition, so that we can avoid duplicated code of the non-window call.

Also addressed.

syed-ahmed · 2026-04-27T23:05:21Z

@pytorchbot merge

pytorchmergebot · 2026-04-27T23:07:29Z

Merge started

Your change will be merged once all checks pass (ETA 0-4 Hours).

Learn more about merging in the wiki.

Questions? Feedback? Please reach out to the PyTorch DevX Team

Advanced Debugging

Check the merge workflow status
here

pytorchmergebot · 2026-04-27T23:33:55Z

Merge failed

Reason: 8 mandatory check(s) failed. The first few are:

Dig deeper by viewing the failures on hud

Details for Dev Infra team

Raised by workflow job

Failing merge rule: Core Maintainers

Skylion007 · 2026-04-28T11:30:38Z

-  std::unordered_map<void*, void*> registeredSegmentHandles_;
+  // Stores handlers for tensors registered by NCCL.
+  // Maps ptr -> (handle, is_window_registered).
+  std::unordered_map<void*, std::pair<void*, bool>> registeredSegmentHandles_;


May want update to ska_flat_hashmap at some point. Seems like this could be perf sensitive.

Skylion007 · 2026-04-28T11:31:32Z

@pytorchbot merge

Skylion007 · 2026-04-28T11:32:20Z

@pytorchbot merge -r

pytorchmergebot · 2026-04-28T11:34:28Z

@pytorchbot started a rebase job onto refs/remotes/origin/viable/strict. Check the current status here

…ered handles NCCLComm::abort() unconditionally called ncclCommDeregister() on all entries in registeredSegmentHandles_. However, handles registered via ncclCommWindowRegister() (symmetric registration) require ncclCommWindowDeregister() instead. Using the wrong API causes NCCL to fail with "Deregister: Could not find handle" followed by a SIGABRT crash during ProcessGroupNCCL destruction. The fix tracks whether each handle was window-registered by changing registeredSegmentHandles_ from map<void*, void*> to map<void*, pair<void*, bool>>, and dispatches to the correct deregister call in abort(), matching the existing logic in deregisterSegment(). Fixes: pytorch#181610

pytorchmergebot · 2026-04-28T11:34:31Z

Successfully rebased fix-window-deregister-abort-v2 onto refs/remotes/origin/viable/strict, please pull locally before adding more changes (for example, via git checkout fix-window-deregister-abort-v2 && git pull --rebase)

pytorchmergebot · 2026-04-28T11:35:49Z

Merge started

Your change will be merged once all checks pass (ETA 0-4 Hours).

Learn more about merging in the wiki.

Questions? Feedback? Please reach out to the PyTorch DevX Team

Advanced Debugging

Check the merge workflow status
here

pytorchmergebot · 2026-04-28T14:10:18Z

Merge failed

Reason: Commit 284031d was HEAD when comment 4334824401 was posted but now the latest commit on the PR is 2dc9e57. Please re-issue the merge command to merge the latest commit.

Details for Dev Infra team

Raised by workflow job

Skylion007 · 2026-04-28T14:43:41Z

@pytorchbot merge

pytorchmergebot · 2026-04-28T14:45:57Z

Merge started

Your change will be merged once all checks pass (ETA 0-4 Hours).

Learn more about merging in the wiki.

Questions? Feedback? Please reach out to the PyTorch DevX Team

Advanced Debugging

Check the merge workflow status
here

syed-ahmed · 2026-04-28T15:29:09Z

Thanks @Skylion007 for merging!

…stered handles (pytorch#181626) `NCCLComm::abort()` unconditionally called `ncclCommDeregister()` on all entries in `registeredSegmentHandles_`. However, handles registered via `ncclCommWindowRegister()` (symmetric registration) require `ncclCommWindowDeregister()` instead. Using the wrong API causes NCCL to fail with "Deregister: Could not find handle" followed by a SIGABRT crash during ProcessGroupNCCL destruction. The fix tracks whether each handle was window-registered by changing `registeredSegmentHandles_` from `map<void*, void*>` to `map<void*, pair<void*, bool>>`, and dispatches to the correct deregister call in `abort()`, matching the existing logic in `deregisterSegment()`. Fixes: pytorch#181610 Assistant used: Claude Opus 4.7 (1M) Pull Request resolved: pytorch#181626 Approved by: https://github.com/kwen2501, https://github.com/Skylion007

pytorch-bot Bot added the release notes: distributed (c10d) release notes category label Apr 27, 2026

syed-ahmed force-pushed the fix-window-deregister-abort-v2 branch from 08625fc to 7fa73cf Compare April 27, 2026 18:52

pytorchbot added the open source label Apr 27, 2026

syed-ahmed requested a review from kwen2501 April 27, 2026 19:08

kwen2501 approved these changes Apr 27, 2026

View reviewed changes

syed-ahmed force-pushed the fix-window-deregister-abort-v2 branch from 7fa73cf to 284031d Compare April 27, 2026 19:49

pytorch-bot Bot added the ciflow/trunk Trigger trunk jobs on your pull request label Apr 27, 2026

pytorchmergebot added the merging label Apr 27, 2026

pytorchmergebot removed the merging label Apr 27, 2026

xrennvidia mentioned this pull request Apr 28, 2026

NCCL UB fix: reduce memory cost and correctly deregister NCCL mem pool NVIDIA/Megatron-LM#4492

Merged

5 tasks

Skylion007 reviewed Apr 28, 2026

View reviewed changes

Skylion007 approved these changes Apr 28, 2026

View reviewed changes

pytorchmergebot force-pushed the fix-window-deregister-abort-v2 branch from 284031d to 2dc9e57 Compare April 28, 2026 11:34

pytorchmergebot added the merging label Apr 28, 2026

pytorchmergebot removed the merging label Apr 28, 2026

pytorchmergebot added the merging label Apr 28, 2026

pytorchmergebot closed this in 573bd8c Apr 28, 2026

pytorchmergebot added Merged and removed merging labels Apr 28, 2026

syed-ahmed deleted the fix-window-deregister-abort-v2 branch April 28, 2026 15:28

Uh oh!

Conversation

syed-ahmed commented Apr 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

pytorch-bot Bot commented Apr 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/181626

❗ 1 Active SEVs

✅ No Failures

Uh oh!

linux-foundation-easycla Bot commented Apr 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

kwen2501 left a comment

Choose a reason for hiding this comment

Uh oh!

kwen2501 Apr 27, 2026

Choose a reason for hiding this comment

Uh oh!

syed-ahmed Apr 27, 2026

Choose a reason for hiding this comment

Uh oh!

kwen2501 Apr 27, 2026

Choose a reason for hiding this comment

Uh oh!

syed-ahmed Apr 27, 2026

Choose a reason for hiding this comment

Uh oh!

syed-ahmed commented Apr 27, 2026

Uh oh!

pytorchmergebot commented Apr 27, 2026

Merge started

Uh oh!

pytorchmergebot commented Apr 27, 2026

Merge failed

Uh oh!

Skylion007 Apr 28, 2026

Choose a reason for hiding this comment

Uh oh!

Skylion007 commented Apr 28, 2026

Uh oh!

Skylion007 commented Apr 28, 2026

Uh oh!

pytorchmergebot commented Apr 28, 2026

Uh oh!

pytorchmergebot commented Apr 28, 2026

Uh oh!

pytorchmergebot commented Apr 28, 2026

Merge started

Uh oh!

pytorchmergebot commented Apr 28, 2026

Merge failed

Uh oh!

Skylion007 commented Apr 28, 2026

Uh oh!

pytorchmergebot commented Apr 28, 2026

Merge started

Uh oh!

syed-ahmed commented Apr 28, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

syed-ahmed commented Apr 27, 2026 •

edited

Loading

pytorch-bot Bot commented Apr 27, 2026 •

edited

Loading

linux-foundation-easycla Bot commented Apr 27, 2026 •

edited

Loading