[Bug]: Ray orchestrator: NcclEP builds its EP communicator from MPI_COMM_WORLD, so each rank gets a 1-rank communicator (nRanks 1 for an EP8 group) #20083

System Info

• CPU architecture: x86_64. One AWS p5en.48xlarge node, in a Kubernetes pod with hostNetwork: true.
• GPU: 8x NVIDIA H200 (141 GB), NVLink. NVIDIA driver 580.178.04; NCCL reports cudaDriverVersion 13040.
• Container: nvcr.io/nvidia/pytorch:26.08-py3 (Ubuntu 24.04), with the unmodified tensorrt_llm-1.3.0rc28-cp312-cp312-manylinux_2_28_x86_64.whl from pypi.nvidia.com (sha256 4b5244b1df862d7dcf94a88e3fbd444dcceda9ee09252a9675bebaee00a8bb75).
• Versions as reported inside the container: tensorrt_llm 1.3.0rc28 (tag v1.3.0rc28, fca831ea); torch 2.14.0a0+4fdf77b940.nv26.08 (CUDA 13.4); ray 2.58.0 and 2.59.0 (one run each); NCCL runtime 2.30.7 (nvidia-nccl-cu13==2.30.7); nccl4py 0.5.0, whose nccl.core.get_version() reports its bindings as NCCL 2.31.2; NCCL-EP library 0.2.0 (nccl.ep.get_lib_version(), the copy embedded in the TensorRT-LLM wheel); mpi4py 4.1.2.
• Model: Qwen/Qwen3-30B-A3B-Instruct-2507 (a local copy), TP8, EP8, PP1, attention DP, orchestrator_type: ray, TRTLLM_FORCE_COMM_METHOD=NCCL_EP.
• Source re-checked on 2026-10-10: nccl_ep_utils.py, deep_ep_utils.py, communication_factory.py, executor/ray/executor.py, executor/ray/gpu_worker.py, mapping.py and _torch/device_mesh.py are byte-identical at v1.3.0rc28, v1.3.0rc29 and main (e1b0ee30). mpi_comm() in _utils.py is unchanged.

Who can help?

@nv-lschneider @shuyixiong

Information

• The official example scripts
• My own modified scripts

Tasks

• An officially supported task in the examples folder (such as GLUE/SQuAD, ...)
• My own task or dataset (give details below)

Reproduction

Under orchestrator_type: ray, NcclEpContext.__init__ builds the NCCL-EP communicator from MPI:

# tensorrt_llm/_torch/moe/fused_moe/nccl_ep_utils.py:196-205 (v1.3.0rc28, v1.3.0rc29, main; call joined onto one line)
self._ep_mpi_comm = mpi_comm().Split(mapping.pp_rank, mapping.moe_ep_rank)
ep_world_rank = self._ep_mpi_comm.Get_rank()
ep_world_size = self._ep_mpi_comm.Get_size()
unique_id = nccl_core.get_unique_id() if ep_world_rank == 0 else None
unique_id = self._ep_mpi_comm.bcast(unique_id, root=0)
self.comm = nccl_core.Communicator.init(nranks=ep_world_size, rank=ep_world_rank, unique_id=unique_id)

A Ray GPU actor is not started by mpirun, so its MPI_COMM_WORLD has one rank, and every rank gets a 1-rank communicator. To see it, on one node with 8 GPUs:

export TRTLLM_FORCE_COMM_METHOD=NCCL_EP NCCL_EP_NUM_QP_PER_RANK=32
export NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT TLLM_LOG_LEVEL=info
export RAY_EXPERIMENTAL_NOSET_CUDA_VISIBLE_DEVICES=1 RAY_DEDUP_LOGS=0
# Our runs also exported ENABLE_CONFIGURABLE_MOE=1 and the AWS EFA / aws-ofi-nccl settings
# (NCCL_GIN_*, FI_*, OFI_NCCL_*, NCCL_NET_PLUGIN); on one node the MoE traffic is NVLink.

printf 'orchestrator_type: ray\nenable_attention_dp: true\n' > ray.yml
ray start --head --port=6379 --disable-usage-stats --num-cpus=16

trtllm-serve /path/to/Qwen3-30B-A3B-Instruct-2507 \
  --served_model_name Qwen/Qwen3-30B-A3B-Instruct-2507 --host 127.0.0.1 --port 8000 \
  --tp_size 8 --ep_size 8 --pp_size 1 \
  --max_batch_size 8 --max_num_tokens 2048 --max_seq_len 8192 \
  --kv_cache_free_gpu_memory_fraction 0.6 --trust_remote_code \
  --extra_llm_api_options ray.yml 2>&1 | tee serve.log

# in a second shell, once "NCCL EP group created" has been logged:
grep -oE 'nRanks [0-9]+ nNodes [0-9]+' serve.log | sort | uniq -c
# then stop the server: under Ray the ranks also accumulate listening sockets during warmup (#20082)

Every rank logs these lines in this order. Host and timestamp prefixes are removed, and the first line repeats once per MoE layer:

[TRT-LLM] [I] [_torch] Selected communication strategy: NcclEP (TRTLLM_FORCE_COMM_METHOD=NCCL_EP)
NCCL INFO [Rank 0] ncclCommInitRankScalable comm 0x5e35f8b0 rank 0 nranks 1 cudaDev 0 nvmlDev 0 busId 59000 commId 0x748e531e8b22ffd3 - Init START
NCCL INFO comm 0x5e35f8b0 rank 0 nRanks 1 nNodes 1 localRanks 1 localRank 0 MNNVL 0
[TRT-LLM] [I] [_torch] NCCL EP group created: ep_size=8, num_experts=128, max_tokens_per_rank=2048, hidden_size=2048, max_top_k=8, layout=RANK_MAJOR

The grep prints 8 nRanks 1 nNodes 1 and nothing else. All eight ranks report rank 0 nranks 1, each with its own commId.

Expected behavior

Under the Ray orchestrator the NCCL-EP communicator spans the EP group, as it does under the MPI orchestrator: 8 ranks for EP8, so group->nRanks == mapping.moe_ep_size.

actual behavior

Each rank creates a 1-rank NCCL communicator and then an NCCL-EP group with ep_size=8 on top of it. The MoE receive buffers are sized from the mapping ([mapping.moe_ep_size, max_tokens_per_rank, hidden]), but the group has one rank.

We saw this twice on one node, with Ray 2.58.0 and 2.59.0. In both runs the server went on into warmup, autotuning and CUDA graph capture with no error. We stopped both runs before health for a separate problem: listening sockets accumulate during warmup under Ray (#20082). So we did not serve requests in this state, and we did not check MoE output correctness.

An earlier two-node run (TP16/EP16) used a non-release combination: the 1.3.0rc9 container with the rc24 versions of nccl_ep_utils.py and communication/nccl_ep.py, and libnccl_ep 0.1.0. With the same 1-rank communicators, the first MoE dispatch, during autotuner warmup, failed the library's check:

nccl_ep.cc:2472: ncclResult_t ncclEpDispatch(...): Assertion `recv_x->sizes[0] == static_cast<unsigned>(group->nRanks)' failed.
*** SIGABRT received ***

The nccl-extensions revision that rc28 builds (af42925a) has the same assert() in ncclEpDispatch, in the NCCL_EP_LAYOUT_RANK_MAJOR case. We did not see it fire on rc28, and we have not checked whether the wheel's libnccl_ep.so is built with NDEBUG.

additional notes

Why the communicator is 1-rank.

• The Ray executor starts every GPU actor with TLLM_DISABLE_MPI=1 and MASTER_ADDR (tensorrt_llm/executor/ray/executor.py:160-161).
• Each actor initialises torch.distributed with the full world size and asserts it (tensorrt_llm/executor/ray/gpu_worker.py:108-113).
• With mpi_disabled(), Mapping becomes a DeviceMeshTopology (tensorrt_llm/mapping.py:559-560, 753). It exposes torch process groups, including moe_ep_group_pg (tensorrt_llm/_torch/device_mesh.py:71-72).
• mpi_comm() returns a thread-local override if one was set, otherwise the module-level pkl5.Intracomm(MPI.COMM_WORLD); it has no torch fallback (tensorrt_llm/_utils.py:453, 468-472).

The comment above the Split (nccl_ep_utils.py:189-195) says the MPI path was chosen to avoid nccl.ep.get_nccl_comm_from_group(), because it "requires torch.distributed.init_process_group() -- the test infrastructure (mpi_pool_executor) and microbenchmarks use MPI4PY only." Under Ray, torch.distributed is initialised by the executor. The helper ships in the wheel as nccl/ep/interop/torch.py.

Same pattern elsewhere. tensorrt_llm/_torch/moe/fused_moe/deep_ep_utils.py:24 and :104 build the DeepEP communicator with the same mpi_comm().Split(...). We have not run DeepEP under Ray. tensorrt_llm/_torch/moe/fused_moe/communication/communication_factory.py has no mpi_disabled() check, so the factory does not steer Ray runs away from these strategies.

A change that gives the expected size. When mpi_disabled(), seed the communicator from the torch EP group; otherwise keep the MPI split. destroy() already handles _ep_mpi_comm is None (nccl_ep_utils.py:395). This is the change we ran, with its comment shortened. It applies to v1.3.0rc28, v1.3.0rc29 and main:

--- a/tensorrt_llm/_torch/moe/fused_moe/nccl_ep_utils.py
+++ b/tensorrt_llm/_torch/moe/fused_moe/nccl_ep_utils.py
@@ -126,7 +126,7 @@
             Tensor,
         )
 
-        from tensorrt_llm._utils import mpi_comm
+        from tensorrt_llm._utils import mpi_comm, mpi_disabled
 
         self.mapping = mapping
         self.ep_size = mapping.moe_ep_size
@@ -186,23 +186,26 @@
             use_internal_fp8_dispatch or external_fp8 or external_nvfp4
         ) and "zero_copy" in getattr(GroupConfig, "__dataclass_fields__", {})
 
-        # MPI sub-communicator scoped to the EP group. Mirrors the
-        # DeepEPLowLatency pattern (see deep_ep_utils.py:104): split
-        # MPI_COMM_WORLD by pp_rank so each pipeline stage gets its own EP
-        # comm, keyed by moe_ep_rank. Avoids the wheel's
-        # nccl.ep.get_nccl_comm_from_group() helper which requires
-        # torch.distributed.init_process_group() -- the test infrastructure
-        # (mpi_pool_executor) and microbenchmarks use MPI4PY only.
-        self._ep_mpi_comm = mpi_comm().Split(mapping.pp_rank, mapping.moe_ep_rank)
-        ep_world_rank = self._ep_mpi_comm.Get_rank()
-        ep_world_size = self._ep_mpi_comm.Get_size()
-        unique_id = nccl_core.get_unique_id() if ep_world_rank == 0 else None
-        unique_id = self._ep_mpi_comm.bcast(unique_id, root=0)
-        self.comm = nccl_core.Communicator.init(
-            nranks=ep_world_size,
-            rank=ep_world_rank,
-            unique_id=unique_id,
-        )
+        # EP communicator bootstrap. Under the MPI orchestrator (and the MPI-based
+        # test infrastructure / microbenchmarks) split MPI_COMM_WORLD by pp_rank,
+        # keyed by moe_ep_rank. Under the Ray orchestrator (TLLM_DISABLE_MPI=1) each
+        # actor is an MPI singleton, so seed the communicator from the torch EP group.
+        if mpi_disabled():
+            from nccl.ep.interop.torch import get_nccl_comm_from_group
+
+            self._ep_mpi_comm = None
+            self.comm = get_nccl_comm_from_group(mapping.moe_ep_group_pg)
+        else:
+            self._ep_mpi_comm = mpi_comm().Split(mapping.pp_rank, mapping.moe_ep_rank)
+            ep_world_rank = self._ep_mpi_comm.Get_rank()
+            ep_world_size = self._ep_mpi_comm.Get_size()
+            unique_id = nccl_core.get_unique_id() if ep_world_rank == 0 else None
+            unique_id = self._ep_mpi_comm.bcast(unique_id, root=0)
+            self.comm = nccl_core.Communicator.init(
+                nranks=ep_world_size,
+                rank=ep_world_rank,
+                unique_id=unique_id,
+            )
 
         dispatch_token_bytes = (
             hidden_size * 2

With this change, the same one-node run logs nRanks 8 nNodes 1 on every rank. A two-node TP16/EP16 run logs nRanks 16 nNodes 2, becomes healthy, and answers requests, with EFA RDMA counters moving during a request. That run also hit the listening-socket problem above, during warmup. We have not run the TensorRT-LLM unit tests with the change. We can open a PR if this direction looks right.

Searches. We found no existing report: searched open and closed issues and PRs separately for NcclEP ray, nccl_ep ray orchestrator, mpi_comm Split ray, TLLM_DISABLE_MPI deep_ep, ncclEpDispatch, get_nccl_comm_from_group, NcclEpContext, moe_ep_group_pg and variants, here and in NVIDIA/nccl-extensions. Closest related work: #18300, which builds MoE MNNVL/all-to-all workspaces without MPI under Ray, and #18231.

Before submitting a new issue...

• Make sure you already searched for relevant issues, and checked the documentation and examples for answers to frequently asked questions.