[WideEP] Update DeepEP version in Dockerfile

docker/Dockerfile

@@ -25,6 +25,10 @@
 ARG CUDA_VERSION=13.0.2
 ARG PYTHON_VERSION=3.12
 ARG UBUNTU_VERSION=22.04
+# DeepEPv2 requires NCCL >= 2.30.4 (GIN backend).
+# This version is only used for CUDA 13+ builds; CUDA 12 falls back to
+# the default NCCL version shipped with the base image.
+ARG NCCL_VERSION=2.30.7
 
 # By parameterizing the base images, we allow third-party to use their own
 # base images. One use case is hermetic builds with base images stored in
@@ -466,10 +470,17 @@ WORKDIR /workspace
 # Build DeepEP wheels
 COPY tools/ep_kernels/install_python_libraries.sh /tmp/install_python_libraries.sh
 # Defaults moved here from tools/ep_kernels/install_python_libraries.sh for centralized version management
-ARG DEEPEP_COMMIT_HASH=73b6ea4
+ARG DEEPEP_COMMIT_HASH=d4f41e4e93
 ARG NVSHMEM_VER
+ARG NCCL_VERSION
 RUN --mount=type=cache,target=/opt/uv/cache \
     mkdir -p /tmp/ep_kernels_workspace/dist && \
+    CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
+    if [ "$CUDA_MAJOR" -ge 13 ] && [ -n "$NCCL_VERSION" ]; then \
+        echo "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" \
+            > /tmp/nccl-override.txt && \
+        export UV_OVERRIDE=/tmp/nccl-override.txt; \
+    fi && \
     export TORCH_CUDA_ARCH_LIST='9.0a 10.0a' && \
     /tmp/install_python_libraries.sh \
         --workspace /tmp/ep_kernels_workspace \
@@ -633,6 +644,7 @@ FROM ${FINAL_BASE_IMAGE} AS vllm-base
 
 ARG CUDA_VERSION
 ARG PYTHON_VERSION
+ARG NCCL_VERSION
 ARG DEADSNAKES_MIRROR_URL
 ARG DEADSNAKES_GPGKEY_URL
 ARG GET_PIP_URL
@@ -685,7 +697,6 @@ RUN apt-get update -y \
 # Install CUDA development tools for runtime JIT compilation
 # (FlashInfer, DeepGEMM, EP kernels all require compilation at runtime)
 RUN CUDA_VERSION_DASH=$(echo $CUDA_VERSION | cut -d. -f1,2 | tr '.' '-') && \
-    CUDA_VERSION_SHORT=$(echo $CUDA_VERSION | cut -d. -f1,2) && \
     apt-get update -y && \
     apt-get install -y --no-install-recommends --allow-change-held-packages \
         cuda-nvcc-${CUDA_VERSION_DASH} \
@@ -698,12 +709,6 @@ RUN CUDA_VERSION_DASH=$(echo $CUDA_VERSION | cut -d. -f1,2 | tr '.' '-') && \
         libnuma-dev \
         # numactl CLI for NUMA binding at runtime
         numactl && \
-    # Fixes nccl_allocator requiring nccl.h at runtime
-    # https://github.com/vllm-project/vllm/blob/1336a1ea244fa8bfd7e72751cabbdb5b68a0c11a/vllm/distributed/device_communicators/pynccl_allocator.py#L22
-    # NCCL packages don't use the cuda-MAJOR-MINOR naming convention,
-    # so we pin the version to match our CUDA version
-    NCCL_VER=$(apt-cache madison libnccl-dev | grep "+cuda${CUDA_VERSION_SHORT}" | head -1 | awk -F'|' '{gsub(/^ +| +$/, "", $2); print $2}') && \
-    apt-get install -y --no-install-recommends --allow-change-held-packages libnccl-dev=${NCCL_VER} libnccl2=${NCCL_VER} && \
     rm -rf /var/lib/apt/lists/*
 
 # Install uv for faster pip installs
@@ -723,6 +728,18 @@ RUN mkdir -p "${UV_PYTHON_INSTALL_DIR}" "${UV_CACHE_DIR}" \
     && chgrp -R 0 /opt/uv \
     && chmod -R g+rwX,a+rX /opt/uv
 
+# DeepEPv2 GIN requires NCCL >= 2.30.4 at both compile and runtime. torch pins
+# an older version as a transitive dep; this override forces uv to use our
+# pinned version whenever nvidia-nccl-cu* is resolved. Empty on CUDA 12 (no-op).
+RUN CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
+    if [ "$CUDA_MAJOR" -ge 13 ]; then \
+        echo "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" \
+            > /etc/uv-overrides.txt; \
+    else \
+        touch /etc/uv-overrides.txt; \
+    fi
+ENV UV_OVERRIDE=/etc/uv-overrides.txt
+
 # ----------------------------------------------------------------------
 # Non-root support (opt-in)
 # ----------------------------------------------------------------------

docker/versions.json

@@ -10,6 +10,9 @@
     "UBUNTU_VERSION": {
       "default": "22.04"
     },
+    "NCCL_VERSION": {
+      "default": "2.30.7"
+    },
     "BUILD_BASE_IMAGE": {
       "default": "nvidia/cuda:13.0.2-devel-ubuntu22.04"
     },
@@ -56,7 +59,7 @@
       "default": "cuda"
     },
     "DEEPEP_COMMIT_HASH": {
-      "default": "73b6ea4"
+      "default": "d4f41e4e93"
     },
     "GIT_REPO_CHECK": {
       "default": "0"

docs/serving/expert_parallel_deployment.md

@@ -12,6 +12,11 @@ Before using EP, you need to install the necessary dependencies. We are actively
 2. **Install DeepGEMM library**: Follow the [official instructions](https://github.com/deepseek-ai/DeepGEMM#installation).
 3. **For disaggregated serving**: Install `gdrcopy` by running the [`install_gdrcopy.sh`](../../tools/install_gdrcopy.sh) script (e.g., `install_gdrcopy.sh "${GDRCOPY_OS_VERSION}" "12.8" "x64"`). You can find available OS versions [here](https://developer.download.nvidia.com/compute/redist/gdrcopy/CUDA%2012.8/).
 
+!!! note "NCCL version (CUDA 13+)"
+    The `deepep_v2` backend requires NCCL >= 2.30.4. PyTorch ships an older
+    NCCL, so you must upgrade it before building or running DeepEP. See the
+    [EP kernels guide](../../tools/ep_kernels) for instructions.
+
 ### Backend Selection Guide
 
 vLLM provides multiple communication backends for EP. Use `--all2all-backend` to select one:

tests/distributed/test_mnnvl_alltoall.py

@@ -73,7 +73,10 @@ def _spawn_workers(worker_fn, world_size, *, dp_size=None):
     err_queue.close()
     err_queue.join_thread()
     if errors:
-        pytest.fail("Worker(s) failed:\n" + "\n---\n".join(errors))
+        combined = "\n---\n".join(errors)
+        if "NCCL GIN" in combined or "ginType=" in combined:
+            pytest.skip("NCCL GIN not available on this system")
+        pytest.fail("Worker(s) failed:\n" + combined)
 
 
 def _run_worker(rank, world_size, port, worker_fn, dp_size, dp_port, err_queue):

tests/kernels/moe/parallel_utils.py

@@ -36,6 +36,10 @@
 P = ParamSpec("P")
 
 
+class GINNotAvailableError(RuntimeError):
+    pass
+
+
 @dataclasses.dataclass
 class ProcessGroupInfo:
     world_size: int
@@ -96,19 +100,26 @@ def parallel_launch(
     **kwargs: P.kwargs,
 ) -> None:
     assert not kwargs
-    spawn(
-        _worker_parallel_launch,
-        args=(
-            world_size,
-            world_size,
-            0,
-            f"tcp://{os.getenv('LOCALHOST', 'localhost')}:{get_open_port()}",
-            worker,
+    try:
+        spawn(
+            _worker_parallel_launch,
+            args=(
+                world_size,
+                world_size,
+                0,
+                f"tcp://{os.getenv('LOCALHOST', 'localhost')}:{get_open_port()}",
+                worker,
+            )
+            + args,
+            nprocs=world_size,
+            join=True,
         )
-        + args,
-        nprocs=world_size,
-        join=True,
-    )
+    except Exception as exc:
+        if "GINNotAvailableError" in str(exc):
+            import pytest
+
+            pytest.skip("NCCL GIN not available (no IBGDA-capable hardware)")
+        raise
 
 
 ## DeepEP specific utils
@@ -225,6 +236,14 @@ def make_deepep_v2_a2a(
 ):
     import deep_ep
 
+    from vllm.utils.nccl import query_nccl_gin_type
+
+    # Unfortunately we can't check GIN availability until the process
+    # group has a live ncclComm_t, so raise here to skip in parallel_launch.
+    gin_type = query_nccl_gin_type(pg)
+    if gin_type is None or gin_type == 0:
+        raise GINNotAvailableError(f"NCCL GIN not available (ginType={gin_type})")
+
     buffer = deep_ep.ElasticBuffer(
         group=pg,
         num_max_tokens_per_rank=v2_args.max_tokens_per_rank,

tools/ep_kernels/README.md

@@ -11,6 +11,38 @@ Step 2 is necessary for multi-node deployment.
 
 All scripts accept a positional argument as workspace path for staging the build, defaulting to `$(pwd)/ep_kernels_workspace`.
 
+## NCCL version requirement (CUDA 13+)
+
+DeepEPv2 uses the NCCL GIN (GPU-Initiated Networking) backend, which requires
+NCCL >= 2.30.4 at both compile time and runtime. PyTorch 2.11 pins
+`nvidia-nccl-cu13==2.28.9` as a transitive dependency, so you need to
+override it.
+
+**With uv** (recommended):
+
+```bash
+# Create an override file
+echo "nvidia-nccl-cu13>=2.30.4" > /tmp/nccl-override.txt
+export UV_OVERRIDE=/tmp/nccl-override.txt
+
+# All subsequent uv pip install commands will respect the override
+uv pip install vllm
+```
+
+**With pip**:
+
+```bash
+pip install vllm
+pip install "nvidia-nccl-cu13>=2.30.4" --no-deps
+```
+
+The override / reinstall must happen before building DeepEP (for GIN device
+headers) and must remain in place at runtime. You can verify with:
+
+```bash
+python -c "from vllm.utils.import_utils import has_deep_ep_v2; print(has_deep_ep_v2())"
+```
+
 ## Usage
 
 ```bash

vllm/distributed/device_communicators/all2all.py

@@ -939,6 +939,7 @@ def __init__(self, cpu_group, tcp_store_group=None, device_group=None):
         self._device_group = device_group
         self.handle_cache = Cache()
         self._num_sms: int | None = None
+        self._gin_checked = False
 
     def _make_all2all_kwargs(
         self,
@@ -961,11 +962,26 @@ def _make_all2all_kwargs(
             explicitly_destroy=True,
         )
 
+    def _check_gin_support(self, group):
+        from vllm.utils.nccl import query_nccl_gin_type
+
+        gin_type = query_nccl_gin_type(group)
+        if gin_type is None or gin_type == 0:
+            raise RuntimeError(
+                "DeepEPv2 requires NCCL GIN (GPU-Initiated Networking) but "
+                f"ginType={gin_type}. This usually means IBGDA-capable "
+                "InfiniBand NICs or drivers are not available. "
+                "See tools/ep_kernels/README.md for requirements."
+            )
+
     def get_handle(self, kwargs):
         import deep_ep  # type: ignore[import-not-found]
 
         num_experts = kwargs.pop("num_experts", 256)
         buffer_kwargs = self._make_all2all_kwargs(**kwargs)
+        if not self._gin_checked:
+            self._check_gin_support(buffer_kwargs["group"])
+            self._gin_checked = True
         logger.debug("DeepEP v2 all2all args %s", buffer_kwargs)
         handle: deep_ep.ElasticBuffer = self.handle_cache.get_or_create(
             buffer_kwargs, deep_ep.ElasticBuffer

vllm/distributed/device_communicators/pynccl_allocator.py

@@ -14,7 +14,7 @@
 from vllm.distributed.device_communicators.pynccl import PyNcclCommunicator
 from vllm.logger import init_logger
 from vllm.platforms import current_platform
-from vllm.utils.nccl import find_nccl_include_paths
+from vllm.utils.nccl import find_nccl_include_paths, find_nccl_library_paths
 
 logger = init_logger(__name__)
 
@@ -74,11 +74,15 @@ def compile_nccl_allocator():
         out_dir = tempfile.gettempdir()
         nccl_allocator_libname = "nccl_allocator"
         nccl_include_paths = find_nccl_include_paths()
+        ldflags = ["-l:libnccl.so.2"]
+        nccl_lib_paths = find_nccl_library_paths()
+        if nccl_lib_paths:
+            ldflags = [f"-L{p}" for p in nccl_lib_paths] + ldflags
         load_inline(
             name=nccl_allocator_libname,
             cpp_sources=nccl_allocator_source,
             with_cuda=True,
-            extra_ldflags=["-lnccl"],
+            extra_ldflags=ldflags,
             verbose=envs.VLLM_LOGGING_LEVEL == "DEBUG",
             is_python_module=False,
             build_directory=out_dir,

vllm/distributed/device_communicators/pynccl_wrapper.py

@@ -51,6 +51,35 @@ class ncclUniqueId(ctypes.Structure):
     _fields_ = [("internal", ctypes.c_byte * 128)]
 
 
+# NCCL 2.30+ ncclCommProperties_t — mirrors src/include/nccl_device/core.h.
+# The struct has a versioned header (size/magic/version) followed by
+# user-visible fields.  We only read ginType so padding mismatches for
+# trailing fields are harmless.
+class ncclCommProperties(ctypes.Structure):
+    _fields_ = [
+        ("size", ctypes.c_size_t),
+        ("magic", ctypes.c_uint),
+        ("version", ctypes.c_uint),
+        ("rank", ctypes.c_int),
+        ("nRanks", ctypes.c_int),
+        ("cudaDev", ctypes.c_int),
+        ("nvmlDev", ctypes.c_int),
+        ("deviceApiSupport", ctypes.c_bool),
+        ("multimemSupport", ctypes.c_bool),
+        ("ginType", ctypes.c_int),
+        ("nLsaTeams", ctypes.c_int),
+        ("hostRmaSupport", ctypes.c_bool),
+        ("railedGinType", ctypes.c_int),
+    ]
+
+
+class ncclGinTypeEnum:
+    NONE = 0
+    PROXY = 2
+    GDAKI = 3
+    GPI = 4
+
+
 cudaStream_t = ctypes.c_void_p
 buffer_type = ctypes.c_void_p
 
@@ -317,6 +346,14 @@ class NCCLLibrary:
         # ncclResult_t ncclCommWindowDeregister(
         #   ncclComm_t comm, ncclWindow_t win);
         Function("ncclCommWindowDeregister", ncclResult_t, [ncclComm_t, ncclWindow_t]),
+        # ncclResult_t ncclCommQueryProperties(
+        #   ncclComm_t comm, ncclCommProperties_t* props);
+        # Available since NCCL 2.29; used to detect GIN support.
+        Function(
+            "ncclCommQueryProperties",
+            ncclResult_t,
+            [ncclComm_t, ctypes.POINTER(ncclCommProperties)],
+        ),
     ]
 
     # class attribute to store the mapping from the path to the library
@@ -375,6 +412,9 @@ def __init__(self, so_file: str | None = None):
                             # Having an exception here on ROCm platform is
                             # not allowed during graph capturing
                             continue
+                    elif func.name == "ncclCommQueryProperties":
+                        # NCCL >= 2.29 only; silently skip on older versions
+                        continue
                     raise
             NCCLLibrary.path_to_dict_mapping[so_file] = _funcs
         self._funcs = NCCLLibrary.path_to_dict_mapping[so_file]