utils: only reuse PyTorch's NCCL comm when _comm_ptr() is non-null - #727
utils: only reuse PyTorch's NCCL comm when _comm_ptr() is non-null#727dmvevents wants to merge 2 commits into
Conversation
get_nccl_comm_handle's reuse branch trusts backend._comm_ptr() blindly, but PyTorch creates NCCL communicators lazily: _comm_ptr() is a passive accessor that returns 0 when the ProcessGroup+device has no communicator yet. Comms materialize at the first collective on that group+device, or eagerly only when init_process_group received device_id=... . A consumer that constructs ElasticBuffer before any collective ran on the EP group (vLLM's expert-parallel serve path is a live example: TP=1, DP-sync NCCL disabled under async scheduling, barriers avoid the device group) gets NCCLCommHandle(0), and calculate_elastic_buffer_size then reaches deep_ep::nccl::get_physical_domain_size -> ncclTeamWorld(nullptr) -> deterministic segfault on every rank at serve init (reproduced 16/16 ranks, DP16/EP16). Guard the reuse: only take the borrowed-comm path when the pointer is non-null, otherwise fall through to the existing create-own-comm path. The fallback's all_gather_object is itself a collective on the group, so it lazily materializes torch's comm as a side effect and every rank takes the same branch on a cold boot. The DeepEP test suite cannot hit this because tests/utils init_dist passes device_id to init_process_group, which makes torch eager-create comms — _comm_ptr() is always valid there. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
🤖 ds-review-bot Code Reviewv6该修改修复了所有 rank 都处于冷启动状态的空指针崩溃，但本地分支决定可能导致部分 rank 提前返回、其余 rank 卡在集合通信中。 v5该 MR 修复 #726： v4p该 MR 修复 #726：get_nccl_comm_handle 的复用分支此前无条件信任 backend._comm_ptr()，而 PyTorch 对 NCCL 通信子是惰性创建的——在冷启动（EP group 尚未执行任何 collective）时 _comm_ptr() 返回 0，导致下游 NCCLCommHandle(0) 在 C++ 侧（如 ncclTeamWorld 解引用 comm->nRanks）发生空指针崩溃。变更在复用前增加非空判断，冷启动时回退到既有的自建通信子路径，整体实现正确、改动最小且符合既有代码模式。 Files reviewed: 1 |
Per-rank branching on _comm_ptr() could split the group: ranks with a materialized communicator returned early while ranks without one entered the group-wide all_gather_object and blocked forever (a hang, harder to debug than the null-deref crash this PR fixes). Every rank now gathers every rank's nullness first - the collective is reached unconditionally - and reuse happens only when ALL ranks hold a live communicator, so all ranks take the same branch. Also documents that a cold-start consumer keeps its self-built comm for the lifetime of the group entry, and adds a unit test covering the uniform-reuse, uniform-null, and rank-skew cases with a mocked backend. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
2 participants
Fixes #726.
get_nccl_comm_handle's reuse branch trustsbackend._comm_ptr()blindly, but PyTorch creates NCCL communicators lazily —_comm_ptr()is a passive accessor that returns 0 when the ProcessGroup+device has no communicator yet (comms materialize at the first collective on that group+device; eager creation only happens wheninit_process_groupreceiveddevice_id=...).A consumer that constructs
ElasticBufferbefore any collective has run on the EP group getsNCCLCommHandle(0), andcalculate_elastic_buffer_sizereachesget_physical_domain_size→ncclTeamWorld(nullptr)→comm->nRanksnull-deref. vLLM's expert-parallel serve path is a live reproduction: deterministic segfault on all 16 ranks (DP16/EP16) at serve init — full chain in #726.Change: only take the borrowed-comm path when the pointer is non-null; otherwise fall through to the existing create-own-comm path. On a cold group the fallback's
all_gather_objectis itself a collective on that group, so it lazily materializes torch's comm as a side effect and every rank takes the same branch — no rank divergence on the cold-boot case.Why the test suite doesn't catch it:
tests/utilsinit_distpassesdevice_id=...toinit_process_group, so torch eager-creates comms and_comm_ptr()is always valid there. The null case is only reachable from lazy-init consumers.Validated E2E on the vLLM path (Qwen3-30B-A3B-FP8, DP16/EP16, 2 nodes over EFA): with the create-own-comm path taken (behaviorally identical to this guard's fall-through on a cold group), serve comes up and sustains a c=1→64 load sweep with 0 errors.
Update: the guard is now an explicit all-rank-consistent decision — every rank gathers every rank's communicator-nullness before any rank branches, so reuse happens only when all ranks hold a live communicator and no rank can be left behind in a collective. Covered by
tests/utils/test_comm_reuse.py.