[GH-ISSUE #726] get_nccl_comm_handle reuses a null _comm_ptr() when the group's NCCL comm is lazily uninitialized — ncclTeamWorld null-deref at ElasticBuffer construction (vLLM serve, 16/16 ranks) #413

New issue

Open

opened 2026-10-26 10:00:18 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

2026-10-26 10:00:18 +00:00

Owner

Originally created by @dmvevents on GitHub (Aug 14, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/726

Summary

get_nccl_comm_handle (deep_ep/utils/comm.py) reuses PyTorch's NCCL communicator when EP_REUSE_NCCL_COMM=1 (the default since #637):

backend = group._get_backend(torch.device('cuda'))
if not force_new_comm and hasattr(backend, '_comm_ptr') and int(os.getenv('EP_REUSE_NCCL_COMM', '1')):
    _storage[group] = NCCLCommHandle(backend._comm_ptr(), False)
    return _storage[group]

backend._comm_ptr() is a passive accessor: PyTorch creates NCCL communicators lazily, and ProcessGroupNCCL::getCommPtr() returns 0 when the group+device has no communicator yet ("ncclComm is a nullptr if the communicator does not exist", ProcessGroupNCCL.cpp, torch 2.11). Comms materialize at the first collective on that group+device — or eagerly only when init_process_group(device_id=...) was used.

So any consumer that constructs an ElasticBuffer before any collective has run on the EP group gets NCCLCommHandle(0), and the constructor immediately calls calculate_elastic_buffer_size(handle.get(), ...) → deep_ep::nccl::get_physical_domain_size (csrc/kernels/backend/nccl.cu) → ncclTeamWorld(comm) which dereferences comm->nRanks (NCCL src/nccl_device/core.cc) → segfault.

Live reproduction (not hypothetical)

vLLM expert-parallel serving (Qwen3-30B-A3B-FP8, DP16/EP16, 2× p5en over EFA) hits this deterministically on all 16 ranks at serve init:

ncclTeamWorld (nccl_device/core.cc:16)
  <- deep_ep::nccl::get_physical_domain_size
  <- ElasticBuffer::calculate_buffer_size (csrc/elastic/buffer.hpp)
  <- pybind

vLLM's default path is exactly the cold-group shape: init_process_group without device_id (all groups lazy), TP=1 (no TP collectives), DP-sync NCCL disabled under async scheduling, gloo-only barriers — so the EP device_group has had zero collectives when ElasticBuffer is constructed from process_weights_after_loading. _comm_ptr() returns 0 and every rank crashes identically.

Setting EP_REUSE_NCCL_COMM=0 (create-own-comm path) fixes the serve on the same substrate — confirmed E2E with sustained load (c=1→64 sweep, 0 errors).

Why DeepEP's own tests don't catch it

tests/utils init_dist passes device_id=torch.device(f'cuda:{local_rank}') to init_process_group, which makes torch eager-create communicators — _comm_ptr() is always valid in the test suite. The reuse branch's null case is only reachable from lazy-init consumers.

Proposed fix

Only take the reuse branch when the pointer is non-null; otherwise fall through to the existing create-own-comm path (whose all_gather_object is itself a collective, so it warms torch's comm as a side effect and all ranks take the same branch on a cold boot):

comm_ptr = backend._comm_ptr()
if comm_ptr != 0:
    _storage[group] = NCCLCommHandle(comm_ptr, False)
    return _storage[group]
# fall through: create our own comm

PR incoming with this guard.

Originally created by @dmvevents on GitHub (Aug 14, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/726 ## Summary `get_nccl_comm_handle` (`deep_ep/utils/comm.py`) reuses PyTorch's NCCL communicator when `EP_REUSE_NCCL_COMM=1` (the default since #637): ```python backend = group._get_backend(torch.device('cuda')) if not force_new_comm and hasattr(backend, '_comm_ptr') and int(os.getenv('EP_REUSE_NCCL_COMM', '1')): _storage[group] = NCCLCommHandle(backend._comm_ptr(), False) return _storage[group] ``` `backend._comm_ptr()` is a **passive accessor**: PyTorch creates NCCL communicators **lazily**, and `ProcessGroupNCCL::getCommPtr()` returns **0** when the group+device has no communicator yet ("ncclComm is a nullptr if the communicator does not exist", `ProcessGroupNCCL.cpp`, torch 2.11). Comms materialize at the first collective on that group+device — or eagerly only when `init_process_group(device_id=...)` was used. So any consumer that constructs an `ElasticBuffer` **before any collective has run on the EP group** gets `NCCLCommHandle(0)`, and the constructor immediately calls `calculate_elastic_buffer_size(handle.get(), ...)` → `deep_ep::nccl::get_physical_domain_size` (`csrc/kernels/backend/nccl.cu`) → `ncclTeamWorld(comm)` which dereferences `comm->nRanks` (NCCL `src/nccl_device/core.cc`) → segfault. ## Live reproduction (not hypothetical) vLLM expert-parallel serving (Qwen3-30B-A3B-FP8, DP16/EP16, 2× p5en over EFA) hits this deterministically on **all 16 ranks** at serve init: ``` ncclTeamWorld (nccl_device/core.cc:16) <- deep_ep::nccl::get_physical_domain_size <- ElasticBuffer::calculate_buffer_size (csrc/elastic/buffer.hpp) <- pybind ``` vLLM's default path is exactly the cold-group shape: `init_process_group` without `device_id` (all groups lazy), TP=1 (no TP collectives), DP-sync NCCL disabled under async scheduling, gloo-only barriers — so the EP `device_group` has had zero collectives when `ElasticBuffer` is constructed from `process_weights_after_loading`. `_comm_ptr()` returns 0 and every rank crashes identically. Setting `EP_REUSE_NCCL_COMM=0` (create-own-comm path) fixes the serve on the same substrate — confirmed E2E with sustained load (c=1→64 sweep, 0 errors). ## Why DeepEP's own tests don't catch it `tests/utils` `init_dist` passes `device_id=torch.device(f'cuda:{local_rank}')` to `init_process_group`, which makes torch **eager-create** communicators — `_comm_ptr()` is always valid in the test suite. The reuse branch's null case is only reachable from lazy-init consumers. ## Proposed fix Only take the reuse branch when the pointer is non-null; otherwise fall through to the existing create-own-comm path (whose `all_gather_object` is itself a collective, so it warms torch's comm as a side effect and all ranks take the same branch on a cold boot): ```python comm_ptr = backend._comm_ptr() if comm_ptr != 0: _storage[group] = NCCLCommHandle(comm_ptr, False) return _storage[group] # fall through: create our own comm ``` PR incoming with this guard.

bsakaguchi commented

2026-10-27 21:02:37 +00:00

Author

Owner

@dmvevents commented on GitHub (Oct 27, 2026):

Status on main @ f6aa38f: this still applies after the V2.5 restructure (#763). get_nccl_comm_handle moved from deep_ep/utils/comm.py to deep_ep/comm/handle.py with the reuse branch unchanged, so a null backend._comm_ptr() is still cached and returned:

github.com/deepseek-ai/DeepEP@f6aa38fe06/deep_ep/comm/handle.py (L63-L66)

The crash path is the same under the new layout. EPBuffer.__init__ takes the handle (deep_ep/buffers/ep.py:299), and the buffer setup calls comm::get_physical_domain_size (csrc/buffers/ep.hpp:263), which evaluates ncclTeamWorld(comm).nRanks without a null check (csrc/kernels/comm/context.cpp:52-56). The get_physical_domain_size and get_logical_domain_size helpers in handle.py pass the handle through the same way.

#727 targets the removed deep_ep/utils/comm.py, which is why it now conflicts. #767 applies the same group-uniform guard to deep_ep/comm/handle.py.
Either one removes the null dereference. One caveat that applies to both: the guard adds an all_gather_object on the group inside the EP_REUSE_NCCL_COMM block, so every rank in the group has to reach it. On a two-node EP16 build of ours that carried #727's change, buffer construction once hung with one node's ranks waiting inside get_nccl_comm_handle and the other node's ranks past it; EP_REUSE_NCCL_COMM=0, which skips the block, avoided it, and we have not root-caused it (it may have been our build rather than the guard). A two-node run on the chosen PR's head before merging would settle that.

Edited 2026-10-27: added the two-node caveat and removed the closing line about #727.

@dmvevents commented on GitHub (Oct 27, 2026): Status on `main` @ f6aa38f: this still applies after the V2.5 restructure (#763). `get_nccl_comm_handle` moved from `deep_ep/utils/comm.py` to `deep_ep/comm/handle.py` with the reuse branch unchanged, so a null `backend._comm_ptr()` is still cached and returned: https://github.com/deepseek-ai/DeepEP/blob/f6aa38fe06cd4a0489e7730365d4e0114970a2f3/deep_ep/comm/handle.py#L63-L66 The crash path is the same under the new layout. `EPBuffer.__init__` takes the handle (`deep_ep/buffers/ep.py:299`), and the buffer setup calls `comm::get_physical_domain_size` (`csrc/buffers/ep.hpp:263`), which evaluates `ncclTeamWorld(comm).nRanks` without a null check (`csrc/kernels/comm/context.cpp:52-56`). The `get_physical_domain_size` and `get_logical_domain_size` helpers in `handle.py` pass the handle through the same way. #727 targets the removed `deep_ep/utils/comm.py`, which is why it now conflicts. #767 applies the same group-uniform guard to `deep_ep/comm/handle.py`. Either one removes the null dereference. One caveat that applies to both: the guard adds an `all_gather_object` on the group inside the `EP_REUSE_NCCL_COMM` block, so every rank in the group has to reach it. On a two-node EP16 build of ours that carried #727's change, buffer construction once hung with one node's ranks waiting inside `get_nccl_comm_handle` and the other node's ranks past it; `EP_REUSE_NCCL_COMM=0`, which skips the block, avoided it, and we have not root-caused it (it may have been our build rather than the guard). A two-node run on the chosen PR's head before merging would settle that. _Edited 2026-10-27: added the two-node caveat and removed the closing line about #727._

bsakaguchi referenced this issue

2026-10-27 21:04:00 +00:00

[PR #413] [MERGED] Fix hidden_size % 128 != 0 in intranode kernels #556