From 71604cb78fe3d071349c1ea4e94d9571095b4cec Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Thu, 25 Jun 2026 16:18:44 -0400
Subject: [PATCH 01/26] [WideEP] Update DeepEP version and pin NCCL >= 2.30.4

DeepEPv2 (ElasticBuffer API) requires NCCL >= 2.30.4 at both compile
and runtime. NCCL backward compatibility is guaranteed by NVIDIA, so
installing a newer NCCL after PyTorch is safe.

- Pin NCCL 2.30.4-1+cuda13.2 in the base build stage so the
  extensions-build stage inherits it for DeepEP compilation
- Pin the same version in the vllm-base runtime stage, replacing the
  dynamic apt-cache madison lookup
- Update DEEPEP_COMMIT_HASH from 73b6ea4 to d4f41e4e93 to pick up
  DeepEPv2

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile    | 18 ++++++++++--------
 docker/versions.json |  5 ++++-
 2 files changed, 14 insertions(+), 9 deletions(-)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index ef166665c5f3..ce91639de92b 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -25,6 +25,9 @@
 ARG CUDA_VERSION=13.0.2
 ARG PYTHON_VERSION=3.12
 ARG UBUNTU_VERSION=22.04
+# DeepEPv2 requires NCCL >= 2.30.4. NVIDIA does not publish a
+# 2.30.4+cuda13.0 package, so use the first available CUDA 13 package.
+ARG NCCL_VERSION=2.30.4-1+cuda13.2
 
 # 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
@@ -102,6 +105,7 @@ FROM ${BUILD_BASE_IMAGE} AS base
 ARG CUDA_VERSION
 ARG PYTHON_VERSION
 ARG BUILD_OS
+ARG NCCL_VERSION
 
 ENV DEBIAN_FRONTEND=noninteractive
 
@@ -154,6 +158,8 @@ RUN if [ "${BUILD_OS}" = "manylinux" ]; then \
             # (pytorch/pytorch#167929) are not compatible with GCC < 11.3.
             gcc-11 \
             g++-11 \
+            libnccl-dev=${NCCL_VERSION} \
+            libnccl2=${NCCL_VERSION} \
         && update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 110 --slave /usr/bin/g++ g++ /usr/bin/g++-11 \
         # Install python dev headers if available (needed for cmake FindPython on Ubuntu 24.04
         # which ships cmake 3.28 and requires Development.SABIModule; silently skipped on
@@ -465,7 +471,7 @@ 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
 RUN --mount=type=cache,target=/opt/uv/cache \
     mkdir -p /tmp/ep_kernels_workspace/dist && \
@@ -632,6 +638,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
@@ -684,7 +691,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} \
@@ -697,12 +703,8 @@ 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} && \
+    apt-get install -y --no-install-recommends --allow-change-held-packages \
+        libnccl-dev=${NCCL_VERSION} libnccl2=${NCCL_VERSION} && \
     rm -rf /var/lib/apt/lists/*
 
 # Install uv for faster pip installs
diff --git a/docker/versions.json b/docker/versions.json
index 3145cfcc53ef..d98602a1e838 100644
--- a/docker/versions.json
+++ b/docker/versions.json
@@ -10,6 +10,9 @@
     "UBUNTU_VERSION": {
       "default": "22.04"
     },
+    "NCCL_VERSION": {
+      "default": "2.30.4-1+cuda13.2"
+    },
     "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"

From 1152444f58a4b904b1ef81e0585ee54e212f6e40 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Thu, 25 Jun 2026 16:40:07 -0400
Subject: [PATCH 02/26] Move NCCL install from base stage to extensions-build
 stage

The CUDA base image holds NCCL packages, so installing them in the
base stage without --allow-change-held-packages fails. Move NCCL
installation to the extensions-build stage where it's actually needed
for DeepEP compilation, with the correct apt flag.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile | 10 +++++++---
 1 file changed, 7 insertions(+), 3 deletions(-)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index ce91639de92b..2039c7392f37 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -105,7 +105,6 @@ FROM ${BUILD_BASE_IMAGE} AS base
 ARG CUDA_VERSION
 ARG PYTHON_VERSION
 ARG BUILD_OS
-ARG NCCL_VERSION
 
 ENV DEBIAN_FRONTEND=noninteractive
 
@@ -158,8 +157,6 @@ RUN if [ "${BUILD_OS}" = "manylinux" ]; then \
             # (pytorch/pytorch#167929) are not compatible with GCC < 11.3.
             gcc-11 \
             g++-11 \
-            libnccl-dev=${NCCL_VERSION} \
-            libnccl2=${NCCL_VERSION} \
         && update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 110 --slave /usr/bin/g++ g++ /usr/bin/g++-11 \
         # Install python dev headers if available (needed for cmake FindPython on Ubuntu 24.04
         # which ships cmake 3.28 and requires Development.SABIModule; silently skipped on
@@ -460,6 +457,7 @@ RUN --mount=type=cache,target=/root/.cache/ccache \
 # This stage is independent and doesn't affect csrc cache
 FROM base AS extensions-build
 ARG CUDA_VERSION
+ARG NCCL_VERSION
 
 # This timeout (in seconds) is necessary when installing some dependencies via uv since it's likely to time out
 ENV UV_HTTP_TIMEOUT=500
@@ -468,6 +466,12 @@ ENV UV_LINK_MODE=copy
 
 WORKDIR /workspace
 
+# Install NCCL >= 2.30.4 headers/libs for DeepEPv2 compilation
+RUN apt-get update -y && \
+    apt-get install -y --no-install-recommends --allow-change-held-packages \
+        libnccl-dev=${NCCL_VERSION} libnccl2=${NCCL_VERSION} && \
+    rm -rf /var/lib/apt/lists/*
+
 # 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

From 2763c72998b77f2673a733bdcc64d44601d08d41 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Thu, 25 Jun 2026 17:16:19 -0400
Subject: [PATCH 03/26] Use pip-installed NCCL for DeepEPv2 GIN device headers

The apt libnccl-dev package only ships host API headers (nccl.h), not
the GIN device headers (ncclGinRequest_t etc.) that DeepEPv2 compiles
against. Install nvidia-nccl-cu13 via pip instead, which includes the
full header set and is auto-detected by DeepEP's build system via
EP_NCCL_ROOT_DIR.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile | 10 ++++------
 1 file changed, 4 insertions(+), 6 deletions(-)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index 2039c7392f37..95be66152880 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -457,7 +457,6 @@ RUN --mount=type=cache,target=/root/.cache/ccache \
 # This stage is independent and doesn't affect csrc cache
 FROM base AS extensions-build
 ARG CUDA_VERSION
-ARG NCCL_VERSION
 
 # This timeout (in seconds) is necessary when installing some dependencies via uv since it's likely to time out
 ENV UV_HTTP_TIMEOUT=500
@@ -466,11 +465,10 @@ ENV UV_LINK_MODE=copy
 
 WORKDIR /workspace
 
-# Install NCCL >= 2.30.4 headers/libs for DeepEPv2 compilation
-RUN apt-get update -y && \
-    apt-get install -y --no-install-recommends --allow-change-held-packages \
-        libnccl-dev=${NCCL_VERSION} libnccl2=${NCCL_VERSION} && \
-    rm -rf /var/lib/apt/lists/*
+# DeepEPv2 compiles against NCCL GIN device headers (ncclGinRequest_t etc.)
+# which are only in the pip package, not the apt libnccl-dev package.
+RUN --mount=type=cache,target=/opt/uv/cache \
+    uv pip install --python /opt/venv/bin/python3 "nvidia-nccl-cu13>=2.30.4" --no-deps
 
 # Build DeepEP wheels
 COPY tools/ep_kernels/install_python_libraries.sh /tmp/install_python_libraries.sh

From 1783ac71eb9b2599b2e2be82e96e6010a6270ac2 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Thu, 25 Jun 2026 17:19:31 -0400
Subject: [PATCH 04/26] Pin NCCL to 2.30.7 and guard against CUDA 12

- Pin NCCL_VERSION to exactly 2.30.7
- Guard pip install (extensions-build) and apt install (vllm-base)
  behind CUDA major >= 13 checks
- CUDA 12 builds fall back to whatever NCCL the base image ships

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile    | 29 ++++++++++++++++++++++++-----
 docker/versions.json |  2 +-
 2 files changed, 25 insertions(+), 6 deletions(-)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index 95be66152880..360dd7daf8a7 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -25,9 +25,10 @@
 ARG CUDA_VERSION=13.0.2
 ARG PYTHON_VERSION=3.12
 ARG UBUNTU_VERSION=22.04
-# DeepEPv2 requires NCCL >= 2.30.4. NVIDIA does not publish a
-# 2.30.4+cuda13.0 package, so use the first available CUDA 13 package.
-ARG NCCL_VERSION=2.30.4-1+cuda13.2
+# 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
@@ -467,8 +468,14 @@ WORKDIR /workspace
 
 # DeepEPv2 compiles against NCCL GIN device headers (ncclGinRequest_t etc.)
 # which are only in the pip package, not the apt libnccl-dev package.
+# Only needed for CUDA 13+; DeepEPv2 GIN is not supported on CUDA 12.
+ARG NCCL_VERSION
 RUN --mount=type=cache,target=/opt/uv/cache \
-    uv pip install --python /opt/venv/bin/python3 "nvidia-nccl-cu13>=2.30.4" --no-deps
+    CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
+    if [ "$CUDA_MAJOR" -ge 13 ]; then \
+        uv pip install --python /opt/venv/bin/python3 \
+            "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" --no-deps; \
+    fi
 
 # Build DeepEP wheels
 COPY tools/ep_kernels/install_python_libraries.sh /tmp/install_python_libraries.sh
@@ -693,6 +700,7 @@ 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_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
     apt-get update -y && \
     apt-get install -y --no-install-recommends --allow-change-held-packages \
         cuda-nvcc-${CUDA_VERSION_DASH} \
@@ -705,8 +713,19 @@ RUN CUDA_VERSION_DASH=$(echo $CUDA_VERSION | cut -d. -f1,2 | tr '.' '-') && \
         libnuma-dev \
         # numactl CLI for NUMA binding at runtime
         numactl && \
+    # Pin NCCL for DeepEPv2 GIN on CUDA 13+; CUDA 12 uses whatever
+    # version the base image ships.
+    if [ "$CUDA_MAJOR" -ge 13 ]; then \
+        NCCL_VER=$(apt-cache madison libnccl-dev \
+            | grep "${NCCL_VERSION}" | head -1 \
+            | awk -F'|' '{gsub(/^ +| +$/, "", $2); print $2}') ; \
+    else \
+        NCCL_VER=$(apt-cache madison libnccl-dev \
+            | grep "+cuda${CUDA_MAJOR}" | head -1 \
+            | awk -F'|' '{gsub(/^ +| +$/, "", $2); print $2}') ; \
+    fi && \
     apt-get install -y --no-install-recommends --allow-change-held-packages \
-        libnccl-dev=${NCCL_VERSION} libnccl2=${NCCL_VERSION} && \
+        libnccl-dev=${NCCL_VER} libnccl2=${NCCL_VER} && \
     rm -rf /var/lib/apt/lists/*
 
 # Install uv for faster pip installs
diff --git a/docker/versions.json b/docker/versions.json
index d98602a1e838..0672c1aa41c7 100644
--- a/docker/versions.json
+++ b/docker/versions.json
@@ -11,7 +11,7 @@
       "default": "22.04"
     },
     "NCCL_VERSION": {
-      "default": "2.30.4-1+cuda13.2"
+      "default": "2.30.7"
     },
     "BUILD_BASE_IMAGE": {
       "default": "nvidia/cuda:13.0.2-devel-ubuntu22.04"

From 72ab9d7eb7d23b3178a5a37dfd16597a5d2fca2f Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Thu, 25 Jun 2026 17:36:14 -0400
Subject: [PATCH 05/26] Use pip for all NCCL installs, drop apt libnccl-dev

The previous approach failed because:
1. install_python_libraries.sh runs `uv pip install torch` which
   re-pins nvidia-nccl-cu13==2.28.9, overwriting our 2.30.7
2. apt libnccl-dev only has host API headers, not GIN device headers

Fix both by:
- Adding NCCL re-install to install_python_libraries.sh after torch,
  gated on NCCL_VERSION env var and CUDA >= 13
- Replacing apt NCCL install in vllm-base with pip install after all
  wheels are installed
- Removing separate pre-build NCCL RUN from extensions-build

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile                            | 35 ++++++--------------
 tools/ep_kernels/install_python_libraries.sh | 12 +++++++
 2 files changed, 22 insertions(+), 25 deletions(-)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index 360dd7daf8a7..f349587350cd 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -466,22 +466,12 @@ ENV UV_LINK_MODE=copy
 
 WORKDIR /workspace
 
-# DeepEPv2 compiles against NCCL GIN device headers (ncclGinRequest_t etc.)
-# which are only in the pip package, not the apt libnccl-dev package.
-# Only needed for CUDA 13+; DeepEPv2 GIN is not supported on CUDA 12.
-ARG NCCL_VERSION
-RUN --mount=type=cache,target=/opt/uv/cache \
-    CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
-    if [ "$CUDA_MAJOR" -ge 13 ]; then \
-        uv pip install --python /opt/venv/bin/python3 \
-            "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" --no-deps; \
-    fi
-
 # 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=d4f41e4e93
 ARG NVSHMEM_VER
+ARG NCCL_VERSION
 RUN --mount=type=cache,target=/opt/uv/cache \
     mkdir -p /tmp/ep_kernels_workspace/dist && \
     export TORCH_CUDA_ARCH_LIST='9.0a 10.0a' && \
@@ -700,7 +690,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_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
     apt-get update -y && \
     apt-get install -y --no-install-recommends --allow-change-held-packages \
         cuda-nvcc-${CUDA_VERSION_DASH} \
@@ -713,19 +702,6 @@ RUN CUDA_VERSION_DASH=$(echo $CUDA_VERSION | cut -d. -f1,2 | tr '.' '-') && \
         libnuma-dev \
         # numactl CLI for NUMA binding at runtime
         numactl && \
-    # Pin NCCL for DeepEPv2 GIN on CUDA 13+; CUDA 12 uses whatever
-    # version the base image ships.
-    if [ "$CUDA_MAJOR" -ge 13 ]; then \
-        NCCL_VER=$(apt-cache madison libnccl-dev \
-            | grep "${NCCL_VERSION}" | head -1 \
-            | awk -F'|' '{gsub(/^ +| +$/, "", $2); print $2}') ; \
-    else \
-        NCCL_VER=$(apt-cache madison libnccl-dev \
-            | grep "+cuda${CUDA_MAJOR}" | head -1 \
-            | awk -F'|' '{gsub(/^ +| +$/, "", $2); print $2}') ; \
-    fi && \
-    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
@@ -908,6 +884,15 @@ RUN --mount=type=bind,from=build,src=/tmp/ep_kernels_workspace/dist,target=/vllm
     uv pip install --system ep_kernels/dist/*.whl --verbose \
         --extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
 
+# DeepEPv2 GIN requires NCCL >= 2.30.4 at runtime. torch pins an older
+# version, so upgrade after all wheel installs are done.
+RUN --mount=type=cache,target=/opt/uv/cache \
+    CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
+    if [ "$CUDA_MAJOR" -ge 13 ]; then \
+        uv pip install --system \
+            "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" --no-deps; \
+    fi
+
 # nvidia-cutlass-dsl[cu13] installs -libs-base and -libs-cu13 wheels that
 # share paths with different content. Force -libs-cu13 last after runtime
 # dependency installs so uv cannot leave base files behind.
diff --git a/tools/ep_kernels/install_python_libraries.sh b/tools/ep_kernels/install_python_libraries.sh
index 94beef897d08..efa212b8f7ca 100755
--- a/tools/ep_kernels/install_python_libraries.sh
+++ b/tools/ep_kernels/install_python_libraries.sh
@@ -82,6 +82,18 @@ else
   uv pip install cmake torch ninja
 fi
 
+# DeepEPv2 compiles against NCCL GIN device headers that are only in the
+# pip package, not the apt libnccl-dev package. torch pins an older NCCL,
+# so re-install the required version after torch.
+if [ -n "${NCCL_VERSION:-}" ] && [ "${CUDA_VERSION_MAJOR}" -ge 13 ]; then
+  echo "Installing nvidia-nccl-cu${CUDA_VERSION_MAJOR}==${NCCL_VERSION} for DeepEPv2 GIN headers"
+  if [ -z "$VIRTUAL_ENV" ]; then
+    uv pip install --system "nvidia-nccl-cu${CUDA_VERSION_MAJOR}==${NCCL_VERSION}" --no-deps
+  else
+    uv pip install "nvidia-nccl-cu${CUDA_VERSION_MAJOR}==${NCCL_VERSION}" --no-deps
+  fi
+fi
+
 # fetch nvshmem
 ARCH=$(uname -m)
 case "${ARCH,,}" in

From 2d16eab90f350b1fccd1fcca4db470574ce696fc Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Fri, 26 Jun 2026 10:33:27 -0400
Subject: [PATCH 06/26] Re-upgrade NCCL in test stage after dev requirements

The test stage installs requirements/dev.txt which includes
requirements/test/cuda.txt pinning nvidia-nccl-cu13==2.28.9,
overwriting the 2.30.7 upgrade from vllm-base. Add another
NCCL upgrade after all test dependencies are installed.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile | 10 ++++++++++
 1 file changed, 10 insertions(+)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index f349587350cd..3a3f4af8f745 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -979,6 +979,16 @@ RUN --mount=type=cache,target=/opt/uv/cache \
 RUN --mount=type=cache,target=/opt/uv/cache \
     uv pip install --system -e tests/vllm_test_utils
 
+# Re-upgrade NCCL after dev requirements (which re-pin torch's older NCCL).
+ARG CUDA_VERSION
+ARG NCCL_VERSION
+RUN --mount=type=cache,target=/opt/uv/cache \
+    CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
+    if [ "$CUDA_MAJOR" -ge 13 ]; then \
+        uv pip install --system \
+            "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" --no-deps; \
+    fi
+
 # enable fast downloads from hf (for testing)
 ENV HF_XET_HIGH_PERFORMANCE 1
 

From 54f604af25c56db337a98f30d7e9bb508c8ba303 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Fri, 26 Jun 2026 10:51:53 -0400
Subject: [PATCH 07/26] Use UV_OVERRIDE to pin NCCL instead of separate
 re-upgrade lines

Replace the three separate `uv pip install nvidia-nccl-cu13==2.30.7
--no-deps` commands (in install script, vllm-base, and test stage) with
a single UV_OVERRIDE env var that forces uv to resolve nvidia-nccl-cu13
to our pinned version whenever it appears as a transitive dependency.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile                            | 37 ++++++++++----------
 tools/ep_kernels/install_python_libraries.sh | 12 -------
 2 files changed, 18 insertions(+), 31 deletions(-)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index 3a3f4af8f745..8ff31ed7a80f 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -474,6 +474,12 @@ 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 \
@@ -721,6 +727,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)
 # ----------------------------------------------------------------------
@@ -884,15 +902,6 @@ RUN --mount=type=bind,from=build,src=/tmp/ep_kernels_workspace/dist,target=/vllm
     uv pip install --system ep_kernels/dist/*.whl --verbose \
         --extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
 
-# DeepEPv2 GIN requires NCCL >= 2.30.4 at runtime. torch pins an older
-# version, so upgrade after all wheel installs are done.
-RUN --mount=type=cache,target=/opt/uv/cache \
-    CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
-    if [ "$CUDA_MAJOR" -ge 13 ]; then \
-        uv pip install --system \
-            "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" --no-deps; \
-    fi
-
 # nvidia-cutlass-dsl[cu13] installs -libs-base and -libs-cu13 wheels that
 # share paths with different content. Force -libs-cu13 last after runtime
 # dependency installs so uv cannot leave base files behind.
@@ -979,16 +988,6 @@ RUN --mount=type=cache,target=/opt/uv/cache \
 RUN --mount=type=cache,target=/opt/uv/cache \
     uv pip install --system -e tests/vllm_test_utils
 
-# Re-upgrade NCCL after dev requirements (which re-pin torch's older NCCL).
-ARG CUDA_VERSION
-ARG NCCL_VERSION
-RUN --mount=type=cache,target=/opt/uv/cache \
-    CUDA_MAJOR=$(echo $CUDA_VERSION | cut -d. -f1) && \
-    if [ "$CUDA_MAJOR" -ge 13 ]; then \
-        uv pip install --system \
-            "nvidia-nccl-cu${CUDA_MAJOR}==${NCCL_VERSION}" --no-deps; \
-    fi
-
 # enable fast downloads from hf (for testing)
 ENV HF_XET_HIGH_PERFORMANCE 1
 
diff --git a/tools/ep_kernels/install_python_libraries.sh b/tools/ep_kernels/install_python_libraries.sh
index efa212b8f7ca..94beef897d08 100755
--- a/tools/ep_kernels/install_python_libraries.sh
+++ b/tools/ep_kernels/install_python_libraries.sh
@@ -82,18 +82,6 @@ else
   uv pip install cmake torch ninja
 fi
 
-# DeepEPv2 compiles against NCCL GIN device headers that are only in the
-# pip package, not the apt libnccl-dev package. torch pins an older NCCL,
-# so re-install the required version after torch.
-if [ -n "${NCCL_VERSION:-}" ] && [ "${CUDA_VERSION_MAJOR}" -ge 13 ]; then
-  echo "Installing nvidia-nccl-cu${CUDA_VERSION_MAJOR}==${NCCL_VERSION} for DeepEPv2 GIN headers"
-  if [ -z "$VIRTUAL_ENV" ]; then
-    uv pip install --system "nvidia-nccl-cu${CUDA_VERSION_MAJOR}==${NCCL_VERSION}" --no-deps
-  else
-    uv pip install "nvidia-nccl-cu${CUDA_VERSION_MAJOR}==${NCCL_VERSION}" --no-deps
-  fi
-fi
-
 # fetch nvshmem
 ARCH=$(uname -m)
 case "${ARCH,,}" in

From 9892bcea115b9d102ae657ec2a2cb1c3d54dc966 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Fri, 26 Jun 2026 11:29:21 -0400
Subject: [PATCH 08/26] Document NCCL override for non-Docker DeepEPv2 installs

Add instructions for pinning NCCL >= 2.30.4 when installing vLLM via
uv/pip rather than the Docker image, covering both UV_OVERRIDE and
pip --no-deps approaches.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docs/serving/expert_parallel_deployment.md |  5 ++++
 tools/ep_kernels/README.md                 | 32 ++++++++++++++++++++++
 2 files changed, 37 insertions(+)

diff --git a/docs/serving/expert_parallel_deployment.md b/docs/serving/expert_parallel_deployment.md
index b7c2ee873750..ee8bf3e165da 100644
--- a/docs/serving/expert_parallel_deployment.md
+++ b/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:
diff --git a/tools/ep_kernels/README.md b/tools/ep_kernels/README.md
index b4eabe18ca1d..04bc3ef27680 100644
--- a/tools/ep_kernels/README.md
+++ b/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

From 11feadcbfcb14a141fe95ed0343484f279d63d42 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Fri, 26 Jun 2026 17:37:17 -0400
Subject: [PATCH 09/26] Fix NCCL allocator link failure with pip-only NCCL

The NCCL symmetric memory allocator JIT-compiles a C++ extension with
`-lnccl`, but the system linker can't find `libnccl.so` when NCCL is
installed only via pip (in site-packages/nvidia/nccl/lib/). Add
`find_nccl_library_paths()` to discover the pip package's lib directory
and pass `-L<path>` to the linker.

This was a pre-existing issue masked by the apt `libnccl-dev` install;
removing apt NCCL in favor of pip-only exposed it.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 .../device_communicators/pynccl_allocator.py   |  8 ++++++--
 vllm/utils/nccl.py                             | 18 ++++++++++++++++++
 2 files changed, 24 insertions(+), 2 deletions(-)

diff --git a/vllm/distributed/device_communicators/pynccl_allocator.py b/vllm/distributed/device_communicators/pynccl_allocator.py
index 27445b81411e..1e35c6a886b5 100644
--- a/vllm/distributed/device_communicators/pynccl_allocator.py
+++ b/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 = ["-lnccl"]
+        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,
diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index 4807bc076f82..093d038aff80 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -62,3 +62,21 @@ def find_nccl_include_paths() -> list[str] | None:
             out.append(p)
             seen.add(p)
     return out or None
+
+
+def find_nccl_library_paths() -> list[str] | None:
+    """Return possible library paths containing `libnccl.so`.
+
+    Looks inside the `nvidia-nccl-cuXX` pip package.
+    """
+    paths: list[str] = []
+    try:
+        spec = importlib.util.find_spec("nvidia.nccl")
+        if spec and (locs := getattr(spec, "submodule_search_locations", None)):
+            for loc in locs:
+                lib_dir = os.path.join(loc, "lib")
+                if os.path.isdir(lib_dir):
+                    paths.append(lib_dir)
+    except Exception as e:
+        logger.debug("Failed to find nccl library path from nvidia.nccl package: %s", e)
+    return paths or None

From 58859ae6b9ac8f78b75535936874b21c3aa70882 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Fri, 26 Jun 2026 23:30:38 -0400
Subject: [PATCH 10/26] Guard against DeepEPv2 SIGSEGV when NCCL GIN is
 unavailable

DeepEPv2's ElasticBuffer constructor calls ncclTeamWorld() to query
cluster topology before checking GIN availability, causing a SIGSEGV
on systems without IBGDA-capable hardware (e.g. B200-k8s CI pods).

Add a pre-flight check using ncclCommQueryProperties (NCCL >= 2.29)
to detect ginType=NONE before creating the ElasticBuffer:
- Serving: raises RuntimeError with actionable message
- Tests: pytest.skip when GIN hardware is unavailable

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 tests/kernels/moe/parallel_utils.py           |  8 +++
 .../device_communicators/all2all.py           | 16 ++++++
 .../device_communicators/pynccl_wrapper.py    | 40 +++++++++++++
 vllm/utils/nccl.py                            | 56 +++++++++++++++++++
 4 files changed, 120 insertions(+)

diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py
index bb2f9efc7c48..ebbbd23ec7e1 100644
--- a/tests/kernels/moe/parallel_utils.py
+++ b/tests/kernels/moe/parallel_utils.py
@@ -225,6 +225,14 @@ def make_deepep_v2_a2a(
 ):
     import deep_ep
 
+    from vllm.utils.nccl import query_nccl_gin_type
+
+    gin_type = query_nccl_gin_type(pg)
+    if gin_type is not None and gin_type == 0:
+        import pytest
+
+        pytest.skip("NCCL GIN not available (no IBGDA-capable hardware)")
+
     buffer = deep_ep.ElasticBuffer(
         group=pg,
         num_max_tokens_per_rank=v2_args.max_tokens_per_rank,
diff --git a/vllm/distributed/device_communicators/all2all.py b/vllm/distributed/device_communicators/all2all.py
index 0066a60dd02b..9582499c0023 100644
--- a/vllm/distributed/device_communicators/all2all.py
+++ b/vllm/distributed/device_communicators/all2all.py
@@ -893,6 +893,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,
@@ -915,11 +916,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 not None and gin_type == 0:
+            raise RuntimeError(
+                "DeepEPv2 requires NCCL GIN (GPU-Initiated Networking) but "
+                "the current environment reports ginType=NONE. 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
diff --git a/vllm/distributed/device_communicators/pynccl_wrapper.py b/vllm/distributed/device_communicators/pynccl_wrapper.py
index 5ca8cc7c77f4..8dd37db4659c 100644
--- a/vllm/distributed/device_communicators/pynccl_wrapper.py
+++ b/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]
diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index 093d038aff80..3d63872d13cf 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -3,6 +3,7 @@
 
 from __future__ import annotations
 
+import ctypes
 import importlib.util
 import os
 
@@ -13,6 +14,8 @@
 
 logger = init_logger(__name__)
 
+_NCCL_GIN_TYPE_NONE = 0
+
 
 def find_nccl_library() -> str:
     """Return NCCL/RCCL shared library name to load.
@@ -80,3 +83,56 @@ def find_nccl_library_paths() -> list[str] | None:
     except Exception as e:
         logger.debug("Failed to find nccl library path from nvidia.nccl package: %s", e)
     return paths or None
+
+
+def query_nccl_gin_type(
+    group: torch.distributed.ProcessGroup,
+) -> int | None:
+    """Query NCCL GIN (GPU-Initiated Networking) support for a process group.
+
+    Returns the ``ncclGinType_t`` value (0 = NONE, 2 = PROXY, 3 = GDAKI,
+    4 = GPI), or ``None`` if the query could not be performed (e.g. NCCL
+    too old or the PyTorch backend does not expose the comm pointer).
+    """
+    from vllm.distributed.device_communicators.pynccl_wrapper import (
+        NCCLLibrary,
+        ncclCommProperties,
+    )
+
+    try:
+        backend = group._get_backend(torch.device("cuda"))
+        if not hasattr(backend, "_comm_ptr"):
+            logger.debug("PyTorch NCCL backend does not expose _comm_ptr")
+            return None
+        comm_ptr = backend._comm_ptr()
+    except Exception:
+        logger.debug(
+            "Failed to extract NCCL comm pointer from process group", exc_info=True
+        )
+        return None
+
+    try:
+        nccl = NCCLLibrary()
+    except Exception:
+        logger.debug("Failed to load NCCL library", exc_info=True)
+        return None
+
+    query_fn = nccl._funcs.get("ncclCommQueryProperties")
+    if query_fn is None:
+        logger.debug("ncclCommQueryProperties not available (NCCL < 2.29)")
+        return None
+
+    props = ncclCommProperties()
+    ctypes.memset(ctypes.addressof(props), 0, ctypes.sizeof(props))
+
+    try:
+        result = query_fn(ctypes.c_void_p(comm_ptr), ctypes.byref(props))
+    except Exception:
+        logger.debug("ncclCommQueryProperties call failed", exc_info=True)
+        return None
+
+    if result != 0:
+        logger.debug("ncclCommQueryProperties returned error %d", result)
+        return None
+
+    return props.ginType

From 7d6c8bc0662947dd7b5d37e39e22f55f8a7721ce Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Mon, 29 Jun 2026 12:20:55 -0400
Subject: [PATCH 11/26] Fix GIN skip crashing in spawned subprocess

pytest.skip() raises a BaseException that bypasses the except-Exception
handler in _worker_parallel_launch, causing torch.multiprocessing.spawn
workers to SIGSEGV instead of cleanly skipping. Raise a regular
GINNotAvailableError inside the subprocess and convert it to
pytest.skip() in parallel_launch on the parent side.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 tests/kernels/moe/parallel_utils.py | 41 ++++++++++++++++++-----------
 1 file changed, 26 insertions(+), 15 deletions(-)

diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py
index ebbbd23ec7e1..8728a3685d42 100644
--- a/tests/kernels/moe/parallel_utils.py
+++ b/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
@@ -229,9 +240,9 @@ def make_deepep_v2_a2a(
 
     gin_type = query_nccl_gin_type(pg)
     if gin_type is not None and gin_type == 0:
-        import pytest
-
-        pytest.skip("NCCL GIN not available (no IBGDA-capable hardware)")
+        raise GINNotAvailableError(
+            "NCCL GIN not available (no IBGDA-capable hardware)"
+        )
 
     buffer = deep_ep.ElasticBuffer(
         group=pg,

From a8fa73f202b07a91a9b01d41a232cea839a21288 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Mon, 29 Jun 2026 12:37:56 -0400
Subject: [PATCH 12/26] Add NCCL GIN diagnostics to B200 fused MOE CI step

Temporary diagnostic to determine why ginType == NONE on B200-k8s CI
pods: checks NIC model, OFED version, nvidia-peermem, kernel version,
RDMA device access, and NCCL env vars.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 .buildkite/test_areas/kernels.yaml | 1 +
 1 file changed, 1 insertion(+)

diff --git a/.buildkite/test_areas/kernels.yaml b/.buildkite/test_areas/kernels.yaml
index 10c132da095a..326bbf1f6fd8 100644
--- a/.buildkite/test_areas/kernels.yaml
+++ b/.buildkite/test_areas/kernels.yaml
@@ -345,5 +345,6 @@ steps:
   - vllm/distributed/device_communicators/
   - vllm/config
   commands:
+    - echo "=== NCCL GIN Diagnostics ===" && (ibstat 2>/dev/null || echo "no ibstat") && (lspci | grep -i mellanox || echo "no mellanox NICs") && (ofed_info -s 2>/dev/null || echo "no OFED info") && (lsmod | grep nvidia_peermem || echo "no nvidia-peermem") && uname -r && (ls -la /dev/infiniband/ 2>/dev/null || echo "no /dev/infiniband") && (env | grep -i NCCL || echo "no NCCL env vars") || true
     - pytest -v -s kernels/moe/test_moe_layer.py
     - pytest -v -s kernels/moe/test_deepep_v2_moe.py

From 4bc4ecf158ac42306df2d9ee3b594994425751e9 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Mon, 29 Jun 2026 14:01:36 -0400
Subject: [PATCH 13/26] Fix ruff formatting for GINNotAvailableError raise

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 tests/kernels/moe/parallel_utils.py | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py
index 8728a3685d42..cc14beb34c07 100644
--- a/tests/kernels/moe/parallel_utils.py
+++ b/tests/kernels/moe/parallel_utils.py
@@ -240,9 +240,7 @@ def make_deepep_v2_a2a(
 
     gin_type = query_nccl_gin_type(pg)
     if gin_type is not None and gin_type == 0:
-        raise GINNotAvailableError(
-            "NCCL GIN not available (no IBGDA-capable hardware)"
-        )
+        raise GINNotAvailableError("NCCL GIN not available (no IBGDA-capable hardware)")
 
     buffer = deep_ep.ElasticBuffer(
         group=pg,

From 40e191a246ab53a6a929d878981d7f1c33cb14e5 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 30 Jun 2026 10:21:09 -0400
Subject: [PATCH 14/26] Fix GIN check: treat query failure as GIN-unavailable

query_nccl_gin_type() returns None when it can't query (e.g. _comm_ptr
missing, ncclCommQueryProperties unavailable). The previous check only
guarded against ginType==0, allowing None to pass through and hit the
ElasticBuffer segfault. Now None is treated as GIN-unavailable too.

Also promote debug logs to warning so failures are visible in CI.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 tests/kernels/moe/parallel_utils.py           |  6 +++--
 .../device_communicators/all2all.py           |  8 +++----
 vllm/utils/nccl.py                            | 22 +++++++++++++------
 3 files changed, 23 insertions(+), 13 deletions(-)

diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py
index cc14beb34c07..ddb31b3c0629 100644
--- a/tests/kernels/moe/parallel_utils.py
+++ b/tests/kernels/moe/parallel_utils.py
@@ -239,8 +239,10 @@ def make_deepep_v2_a2a(
     from vllm.utils.nccl import query_nccl_gin_type
 
     gin_type = query_nccl_gin_type(pg)
-    if gin_type is not None and gin_type == 0:
-        raise GINNotAvailableError("NCCL GIN not available (no IBGDA-capable hardware)")
+    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,
diff --git a/vllm/distributed/device_communicators/all2all.py b/vllm/distributed/device_communicators/all2all.py
index 9582499c0023..d813a4a6529e 100644
--- a/vllm/distributed/device_communicators/all2all.py
+++ b/vllm/distributed/device_communicators/all2all.py
@@ -920,12 +920,12 @@ 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 not None and gin_type == 0:
+        if gin_type is None or gin_type == 0:
             raise RuntimeError(
                 "DeepEPv2 requires NCCL GIN (GPU-Initiated Networking) but "
-                "the current environment reports ginType=NONE. This usually "
-                "means IBGDA-capable InfiniBand NICs or drivers are not "
-                "available. See tools/ep_kernels/README.md for requirements."
+                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):
diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index 3d63872d13cf..c2aec83903a0 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -102,24 +102,27 @@ def query_nccl_gin_type(
     try:
         backend = group._get_backend(torch.device("cuda"))
         if not hasattr(backend, "_comm_ptr"):
-            logger.debug("PyTorch NCCL backend does not expose _comm_ptr")
+            logger.warning("PyTorch NCCL backend does not expose _comm_ptr")
             return None
         comm_ptr = backend._comm_ptr()
     except Exception:
-        logger.debug(
-            "Failed to extract NCCL comm pointer from process group", exc_info=True
+        logger.warning(
+            "Failed to extract NCCL comm pointer from process group",
+            exc_info=True,
         )
         return None
 
     try:
         nccl = NCCLLibrary()
     except Exception:
-        logger.debug("Failed to load NCCL library", exc_info=True)
+        logger.warning("Failed to load NCCL library", exc_info=True)
         return None
 
     query_fn = nccl._funcs.get("ncclCommQueryProperties")
     if query_fn is None:
-        logger.debug("ncclCommQueryProperties not available (NCCL < 2.29)")
+        logger.warning(
+            "ncclCommQueryProperties not available (NCCL < 2.29)"
+        )
         return None
 
     props = ncclCommProperties()
@@ -128,11 +131,16 @@ def query_nccl_gin_type(
     try:
         result = query_fn(ctypes.c_void_p(comm_ptr), ctypes.byref(props))
     except Exception:
-        logger.debug("ncclCommQueryProperties call failed", exc_info=True)
+        logger.warning(
+            "ncclCommQueryProperties call failed", exc_info=True
+        )
         return None
 
     if result != 0:
-        logger.debug("ncclCommQueryProperties returned error %d", result)
+        logger.warning(
+            "ncclCommQueryProperties returned error %d", result
+        )
         return None
 
+    logger.info("NCCL GIN type: %d", props.ginType)
     return props.ginType

From 97dbca7b1c22eeea4b5ffce7e8076a60a70c1b51 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 30 Jun 2026 10:22:09 -0400
Subject: [PATCH 15/26] Apply ruff formatting fixes

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 tests/kernels/moe/parallel_utils.py |  4 +---
 vllm/utils/nccl.py                  | 12 +++---------
 2 files changed, 4 insertions(+), 12 deletions(-)

diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py
index ddb31b3c0629..e75c10068184 100644
--- a/tests/kernels/moe/parallel_utils.py
+++ b/tests/kernels/moe/parallel_utils.py
@@ -240,9 +240,7 @@ def make_deepep_v2_a2a(
 
     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})"
-        )
+        raise GINNotAvailableError(f"NCCL GIN not available (ginType={gin_type})")
 
     buffer = deep_ep.ElasticBuffer(
         group=pg,
diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index c2aec83903a0..805849887680 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -120,9 +120,7 @@ def query_nccl_gin_type(
 
     query_fn = nccl._funcs.get("ncclCommQueryProperties")
     if query_fn is None:
-        logger.warning(
-            "ncclCommQueryProperties not available (NCCL < 2.29)"
-        )
+        logger.warning("ncclCommQueryProperties not available (NCCL < 2.29)")
         return None
 
     props = ncclCommProperties()
@@ -131,15 +129,11 @@ def query_nccl_gin_type(
     try:
         result = query_fn(ctypes.c_void_p(comm_ptr), ctypes.byref(props))
     except Exception:
-        logger.warning(
-            "ncclCommQueryProperties call failed", exc_info=True
-        )
+        logger.warning("ncclCommQueryProperties call failed", exc_info=True)
         return None
 
     if result != 0:
-        logger.warning(
-            "ncclCommQueryProperties returned error %d", result
-        )
+        logger.warning("ncclCommQueryProperties returned error %d", result)
         return None
 
     logger.info("NCCL GIN type: %d", props.ginType)

From 5d469f4248d34c34308eeeb1b353ffa2b84b586c Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 30 Jun 2026 11:57:36 -0400
Subject: [PATCH 16/26] Create libnccl.so linker symlink from pip-installed
 NCCL

The nvidia-nccl-cu* pip package ships libnccl.so.2 (runtime) but not
the libnccl.so dev symlink. The nccl_allocator JIT extension needs
-lnccl at link time. The previous apt-based libnccl-dev install
provided this symlink; now that we use the pip package instead, we
create it explicitly.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index ea8b58644a25..ef9bb3706423 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -903,6 +903,14 @@ RUN --mount=type=bind,from=build,src=/tmp/ep_kernels_workspace/dist,target=/vllm
     uv pip install --system ep_kernels/dist/*.whl --verbose \
         --extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
 
+# The nvidia-nccl-cu* pip package ships libnccl.so.2 but not the libnccl.so
+# linker symlink. The nccl_allocator JIT extension links with -lnccl, so we
+# create the symlink after pip installs have resolved the final NCCL version.
+RUN NCCL_LIB=$(python3 -c "import nvidia.nccl; import os; print(os.path.join(os.path.dirname(nvidia.nccl.__file__), 'lib'))" 2>/dev/null) && \
+    if [ -n "$NCCL_LIB" ] && [ -f "$NCCL_LIB/libnccl.so.2" ] && [ ! -e "$NCCL_LIB/libnccl.so" ]; then \
+        ln -s libnccl.so.2 "$NCCL_LIB/libnccl.so"; \
+    fi
+
 # nvidia-cutlass-dsl[cu13] installs -libs-base and -libs-cu13 wheels that
 # share paths with different content. Force -libs-cu13 last after runtime
 # dependency installs so uv cannot leave base files behind.

From 441a52296390fa08eb3d50b58284cf43c64d99d3 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 30 Jun 2026 12:03:30 -0400
Subject: [PATCH 17/26] Link nccl_allocator against libnccl.so.2 directly

Use -l:libnccl.so.2 instead of -lnccl so the JIT extension links
against the SONAME directly. This removes the need for the libnccl.so
dev symlink that the pip nvidia-nccl-cu* package doesn't ship, and
that libnccl-dev (apt) would otherwise need to provide.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 docker/Dockerfile                                         | 8 --------
 vllm/distributed/device_communicators/pynccl_allocator.py | 2 +-
 2 files changed, 1 insertion(+), 9 deletions(-)

diff --git a/docker/Dockerfile b/docker/Dockerfile
index ef9bb3706423..ea8b58644a25 100644
--- a/docker/Dockerfile
+++ b/docker/Dockerfile
@@ -903,14 +903,6 @@ RUN --mount=type=bind,from=build,src=/tmp/ep_kernels_workspace/dist,target=/vllm
     uv pip install --system ep_kernels/dist/*.whl --verbose \
         --extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
 
-# The nvidia-nccl-cu* pip package ships libnccl.so.2 but not the libnccl.so
-# linker symlink. The nccl_allocator JIT extension links with -lnccl, so we
-# create the symlink after pip installs have resolved the final NCCL version.
-RUN NCCL_LIB=$(python3 -c "import nvidia.nccl; import os; print(os.path.join(os.path.dirname(nvidia.nccl.__file__), 'lib'))" 2>/dev/null) && \
-    if [ -n "$NCCL_LIB" ] && [ -f "$NCCL_LIB/libnccl.so.2" ] && [ ! -e "$NCCL_LIB/libnccl.so" ]; then \
-        ln -s libnccl.so.2 "$NCCL_LIB/libnccl.so"; \
-    fi
-
 # nvidia-cutlass-dsl[cu13] installs -libs-base and -libs-cu13 wheels that
 # share paths with different content. Force -libs-cu13 last after runtime
 # dependency installs so uv cannot leave base files behind.
diff --git a/vllm/distributed/device_communicators/pynccl_allocator.py b/vllm/distributed/device_communicators/pynccl_allocator.py
index 1e35c6a886b5..c4a199647ca0 100644
--- a/vllm/distributed/device_communicators/pynccl_allocator.py
+++ b/vllm/distributed/device_communicators/pynccl_allocator.py
@@ -74,7 +74,7 @@ def compile_nccl_allocator():
         out_dir = tempfile.gettempdir()
         nccl_allocator_libname = "nccl_allocator"
         nccl_include_paths = find_nccl_include_paths()
-        ldflags = ["-lnccl"]
+        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

From 27121dd9ab7e914ec3a5caaf7ac4fa94106c7610 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 30 Jun 2026 17:11:53 -0400
Subject: [PATCH 18/26] Skip test_deepep_v2_manager_lifecycle when GIN is
 unavailable

The lifecycle test now runs because NCCL >= 2.30.4 is present, but on
pods without IBGDA-capable hardware it hits the GIN RuntimeError from
_check_gin_support. Detect this in _spawn_workers and convert to
pytest.skip.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 tests/distributed/test_mnnvl_alltoall.py | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/tests/distributed/test_mnnvl_alltoall.py b/tests/distributed/test_mnnvl_alltoall.py
index 95c905fc0803..ec326c37d53f 100644
--- a/tests/distributed/test_mnnvl_alltoall.py
+++ b/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):

From 20b07f6090a166b94c83e635717bec497db5171f Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 30 Jun 2026 21:25:34 -0400
Subject: [PATCH 19/26] Address review comments

- Remove debug logger.info for GIN type (was temporary instrumentation)
- Remove NCCL GIN diagnostics echo from B200 CI step
- Add comment explaining why GIN check happens post-spawn

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 .buildkite/test_areas/kernels.yaml  | 1 -
 tests/kernels/moe/parallel_utils.py | 2 ++
 vllm/utils/nccl.py                  | 1 -
 3 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/.buildkite/test_areas/kernels.yaml b/.buildkite/test_areas/kernels.yaml
index 326bbf1f6fd8..10c132da095a 100644
--- a/.buildkite/test_areas/kernels.yaml
+++ b/.buildkite/test_areas/kernels.yaml
@@ -345,6 +345,5 @@ steps:
   - vllm/distributed/device_communicators/
   - vllm/config
   commands:
-    - echo "=== NCCL GIN Diagnostics ===" && (ibstat 2>/dev/null || echo "no ibstat") && (lspci | grep -i mellanox || echo "no mellanox NICs") && (ofed_info -s 2>/dev/null || echo "no OFED info") && (lsmod | grep nvidia_peermem || echo "no nvidia-peermem") && uname -r && (ls -la /dev/infiniband/ 2>/dev/null || echo "no /dev/infiniband") && (env | grep -i NCCL || echo "no NCCL env vars") || true
     - pytest -v -s kernels/moe/test_moe_layer.py
     - pytest -v -s kernels/moe/test_deepep_v2_moe.py
diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py
index e75c10068184..735578722a89 100644
--- a/tests/kernels/moe/parallel_utils.py
+++ b/tests/kernels/moe/parallel_utils.py
@@ -238,6 +238,8 @@ def make_deepep_v2_a2a(
 
     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})")
diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index 805849887680..db5ca6c276c0 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -136,5 +136,4 @@ def query_nccl_gin_type(
         logger.warning("ncclCommQueryProperties returned error %d", result)
         return None
 
-    logger.info("NCCL GIN type: %d", props.ginType)
     return props.ginType

From a6ab383771f5158eda473678d8cb4d07cdaa7d36 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Wed, 1 Jul 2026 13:34:00 -0400
Subject: [PATCH 20/26] Initialize ncclCommProperties with required header
 fields

NCCL validates magic, version, and size on the ncclCommProperties
struct (same pattern as NCCL_CONFIG_INITIALIZER). Without these,
ncclCommQueryProperties rejects the struct.

Co-authored-by: Ilya Markov <ilmarkov@users.noreply.github.com>
Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 vllm/utils/nccl.py | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index db5ca6c276c0..3e250d80446f 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -125,6 +125,10 @@ def query_nccl_gin_type(
 
     props = ncclCommProperties()
     ctypes.memset(ctypes.addressof(props), 0, ctypes.sizeof(props))
+    # NCCL validates these fields (mirrors NCCL_COMM_PROPERTIES_INITIALIZER)
+    props.size = ctypes.sizeof(props)
+    props.magic = 0xCAFEBEEF
+    props.version = nccl.ncclGetRawVersion()
 
     try:
         result = query_fn(ctypes.c_void_p(comm_ptr), ctypes.byref(props))

From 0f698a01c38bca4b45461436ceae91a2f3be6cb8 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Thu, 2 Jul 2026 10:32:11 -0400
Subject: [PATCH 21/26] Fix GIN check: defer when NCCL comm not yet initialized

The NCCL communicator is lazily created on the first collective.
When _check_gin_support runs during process_weights_after_loading,
_comm_ptr() returns 0 (NULL) and the query fails, incorrectly
blocking startup.

Fix: return None from query_nccl_gin_type when comm_ptr is 0, and
retry the check on subsequent get_handle calls once the comm exists
after real collectives.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 vllm/distributed/device_communicators/all2all.py | 12 ++++++++----
 vllm/utils/nccl.py                               |  2 ++
 2 files changed, 10 insertions(+), 4 deletions(-)

diff --git a/vllm/distributed/device_communicators/all2all.py b/vllm/distributed/device_communicators/all2all.py
index 4445cdb16628..1bb0af20bfbb 100644
--- a/vllm/distributed/device_communicators/all2all.py
+++ b/vllm/distributed/device_communicators/all2all.py
@@ -962,17 +962,22 @@ def _make_all2all_kwargs(
             explicitly_destroy=True,
         )
 
-    def _check_gin_support(self, group):
+    def _check_gin_support(self, group) -> bool:
+        """Check GIN support. Returns True if confirmed, False if
+        inconclusive (e.g. NCCL comm not yet initialized)."""
         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:
+        if gin_type is None:
+            return False
+        if 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."
             )
+        return True
 
     def get_handle(self, kwargs):
         import deep_ep  # type: ignore[import-not-found]
@@ -980,8 +985,7 @@ def get_handle(self, kwargs):
         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
+            self._gin_checked = self._check_gin_support(buffer_kwargs["group"])
         logger.debug("DeepEP v2 all2all args %s", buffer_kwargs)
         handle: deep_ep.ElasticBuffer = self.handle_cache.get_or_create(
             buffer_kwargs, deep_ep.ElasticBuffer
diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index 3e250d80446f..fb700c28372b 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -105,6 +105,8 @@ def query_nccl_gin_type(
             logger.warning("PyTorch NCCL backend does not expose _comm_ptr")
             return None
         comm_ptr = backend._comm_ptr()
+        if comm_ptr == 0:
+            return None
     except Exception:
         logger.warning(
             "Failed to extract NCCL comm pointer from process group",

From d7012acf3a10e3cd2b00259529e586c60ae43c63 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Thu, 2 Jul 2026 15:44:53 -0400
Subject: [PATCH 22/26] Fix NCCL downgrade from bare pip install in CI

`pip install tensorizer` resolves torch's transitive dep
nvidia-nccl-cu13==2.28.9, downgrading the 2.30.7 installed by the
Docker image. This breaks deep_ep import (undefined symbol:
ncclCommQueryProperties). Switch to `uv pip install` so the
UV_OVERRIDE NCCL pin is respected.

Also: remove unused _NCCL_GIN_TYPE_NONE constant, improve GIN check
comment in test utils.

Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 .buildkite/test_areas/misc.yaml            | 2 +-
 .buildkite/test_areas/model_runner_v2.yaml | 2 +-
 tests/kernels/moe/parallel_utils.py        | 6 ++++--
 vllm/utils/nccl.py                         | 2 --
 4 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/.buildkite/test_areas/misc.yaml b/.buildkite/test_areas/misc.yaml
index 450cfcbc26de..bb0e27a44308 100644
--- a/.buildkite/test_areas/misc.yaml
+++ b/.buildkite/test_areas/misc.yaml
@@ -202,7 +202,7 @@ steps:
   - vllm/multimodal
   - examples/
   commands:
-    - pip install tensorizer # for tensorizer test
+    - uv pip install tensorizer # for tensorizer test
      # for basic
     - python3 basic/offline_inference/chat.py
     - python3 basic/offline_inference/generate.py --model facebook/opt-125m
diff --git a/.buildkite/test_areas/model_runner_v2.yaml b/.buildkite/test_areas/model_runner_v2.yaml
index dbb35df80be9..8356d7ee97e6 100644
--- a/.buildkite/test_areas/model_runner_v2.yaml
+++ b/.buildkite/test_areas/model_runner_v2.yaml
@@ -41,7 +41,7 @@ steps:
   commands:
     - set -x
     - export VLLM_USE_V2_MODEL_RUNNER=1
-    - pip install tensorizer # for tensorizer test
+    - uv pip install tensorizer # for tensorizer test
     - python3 basic/offline_inference/chat.py # for basic
     - python3 basic/offline_inference/generate.py --model facebook/opt-125m
     #- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10  # TODO
diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py
index 735578722a89..cdbbf2048718 100644
--- a/tests/kernels/moe/parallel_utils.py
+++ b/tests/kernels/moe/parallel_utils.py
@@ -238,8 +238,10 @@ def make_deepep_v2_a2a(
 
     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 availability can only be checked after NCCL comm init, which
+    # happens inside the spawned workers. Raise GINNotAvailableError so
+    # parallel_launch() can catch it (via string match across the process
+    # boundary) and convert to pytest.skip.
     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})")
diff --git a/vllm/utils/nccl.py b/vllm/utils/nccl.py
index fb700c28372b..0fc5379c4d2d 100644
--- a/vllm/utils/nccl.py
+++ b/vllm/utils/nccl.py
@@ -14,8 +14,6 @@
 
 logger = init_logger(__name__)
 
-_NCCL_GIN_TYPE_NONE = 0
-
 
 def find_nccl_library() -> str:
     """Return NCCL/RCCL shared library name to load.

From 54bcda09ad30c4dad80a845d626afdb0d760c294 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Mon, 6 Jul 2026 10:55:07 -0400
Subject: [PATCH 23/26] [CI] Use pip --no-deps for tensorizer to fix AMD CI

uv pip install fails on AMD CI where no virtual environment exists.
Switch to pip install --no-deps to avoid both the venv requirement and
the risk of torch pulling in an older NCCL as a transitive dependency.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tyler@vllm.ai>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 .buildkite/test_areas/misc.yaml            | 2 +-
 .buildkite/test_areas/model_runner_v2.yaml | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/.buildkite/test_areas/misc.yaml b/.buildkite/test_areas/misc.yaml
index 9e9f1ed30bbd..8950c0fdd7e7 100644
--- a/.buildkite/test_areas/misc.yaml
+++ b/.buildkite/test_areas/misc.yaml
@@ -202,7 +202,7 @@ steps:
   - vllm/multimodal
   - examples/
   commands:
-    - uv pip install tensorizer # for tensorizer test
+    - pip install --no-deps tensorizer # for tensorizer test
      # for basic
     - python3 basic/offline_inference/chat.py
     - python3 basic/offline_inference/generate.py --model facebook/opt-125m
diff --git a/.buildkite/test_areas/model_runner_v2.yaml b/.buildkite/test_areas/model_runner_v2.yaml
index 8356d7ee97e6..914d4ed609c2 100644
--- a/.buildkite/test_areas/model_runner_v2.yaml
+++ b/.buildkite/test_areas/model_runner_v2.yaml
@@ -41,7 +41,7 @@ steps:
   commands:
     - set -x
     - export VLLM_USE_V2_MODEL_RUNNER=1
-    - uv pip install tensorizer # for tensorizer test
+    - pip install --no-deps tensorizer # for tensorizer test
     - python3 basic/offline_inference/chat.py # for basic
     - python3 basic/offline_inference/generate.py --model facebook/opt-125m
     #- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10  # TODO

From b9e57d5647e2f8a9f471fa62dd1af52112ec6843 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 7 Jul 2026 13:55:45 -0400
Subject: [PATCH 24/26] [Test] Skip DeepEP MoE layer tests without P2P access

DeepEP backends require peer access between the participating GPUs. Reuse the existing gpu_p2p_access_check helper to skip only those backend parametrizations when the visible device set does not satisfy that requirement.

Co-authored-by: OpenAI Codex <codex@openai.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 tests/kernels/moe/test_moe_layer.py | 24 ++++++++++++++++++++++++
 1 file changed, 24 insertions(+)

diff --git a/tests/kernels/moe/test_moe_layer.py b/tests/kernels/moe/test_moe_layer.py
index 552063988fac..cc8d9c36dc0a 100644
--- a/tests/kernels/moe/test_moe_layer.py
+++ b/tests/kernels/moe/test_moe_layer.py
@@ -36,6 +36,9 @@
     get_eplb_group,
     tensor_model_parallel_all_gather,
 )
+from vllm.distributed.device_communicators.all_reduce_utils import (
+    gpu_p2p_access_check,
+)
 from vllm.distributed.eplb.eplb_communicator import create_eplb_communicator
 from vllm.distributed.eplb.rebalance_execute import rearrange_expert_weights_inplace
 from vllm.forward_context import set_forward_context
@@ -106,6 +109,8 @@
 if has_nixl_ep():
     BACKENDS += ["nixl_ep"]
 
+DEEPEP_BACKENDS = {"deepep_high_throughput", "deepep_low_latency"}
+
 QUANT_METHODS = [
     None,
     "fp8",
@@ -416,6 +421,22 @@ def generate_valid_test_configs(
     return configs
 
 
+@functools.cache
+def visible_devices_have_peer_access(world_size: int) -> bool:
+    if not current_platform.is_cuda():
+        return True
+
+    try:
+        return all(
+            gpu_p2p_access_check(src, dst)
+            for src in range(world_size)
+            for dst in range(world_size)
+            if src != dst
+        )
+    except RuntimeError:
+        return False
+
+
 # TODO: break this up into sections
 def is_valid_config(config: MoETestConfig) -> tuple[bool, str | None]:
     # routed_input_transform only makes sense with shared_experts (latent MoE)
@@ -1802,6 +1823,9 @@ def test_moe_layer(
     if enable_eplb and not use_ep:
         pytest.skip("EPLB requires EP.")
 
+    if backend in DEEPEP_BACKENDS and not visible_devices_have_peer_access(world_size):
+        pytest.skip("DeepEP backends require peer access between visible GPUs.")
+
     verbosity = pytestconfig.getoption("verbose")
 
     if os.environ.get("VLLM_LOGGING_LEVEL") is None:

From af75196876c92cc5bb46f2bf7d95f9c5aa736a79 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 7 Jul 2026 16:18:15 -0400
Subject: [PATCH 25/26] [Test] Add async TP device diagnostics

Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 .../passes/distributed/test_async_tp.py       | 113 ++++++++++++++++++
 1 file changed, 113 insertions(+)

diff --git a/tests/compile/passes/distributed/test_async_tp.py b/tests/compile/passes/distributed/test_async_tp.py
index 7d1320bccaaf..d1f0cd5f76c6 100644
--- a/tests/compile/passes/distributed/test_async_tp.py
+++ b/tests/compile/passes/distributed/test_async_tp.py
@@ -1,6 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+import os
 
 import pytest
 import torch
@@ -44,6 +45,78 @@
 ]
 
 
+def _safe_debug_value(label: str, fn):
+    try:
+        return f"{label}={fn()}"
+    except Exception as e:
+        return f"{label}=<error {type(e).__name__}: {e}>"
+
+
+def _torch_nccl_version():
+    return torch.cuda.nccl.version()
+
+
+def _log_async_tp_debug(stage: str, local_rank: int, world_size: int) -> None:
+    env_keys = [
+        "CUDA_VISIBLE_DEVICES",
+        "NVIDIA_VISIBLE_DEVICES",
+        "LOCAL_RANK",
+        "RANK",
+        "WORLD_SIZE",
+        "MASTER_ADDR",
+        "MASTER_PORT",
+        "VLLM_RAY_PER_WORKER_GPUS",
+        "VLLM_RAY_BUNDLE_INDICES",
+    ]
+    env_values = ", ".join(f"{key}={os.environ.get(key)!r}" for key in env_keys)
+
+    debug_values = [
+        f"stage={stage}",
+        f"local_rank={local_rank}",
+        f"world_size={world_size}",
+        f"DEVICE_TYPE={DEVICE_TYPE}",
+        _safe_debug_value("torch.__version__", lambda: torch.__version__),
+        _safe_debug_value("torch.version.cuda", lambda: torch.version.cuda),
+        _safe_debug_value(
+            "torch.accelerator.is_available", torch.accelerator.is_available
+        ),
+        _safe_debug_value(
+            "torch.accelerator.device_count", torch.accelerator.device_count
+        ),
+        _safe_debug_value(
+            "torch.accelerator.current_device_index",
+            torch.accelerator.current_device_index,
+        ),
+        _safe_debug_value(
+            "current_platform.get_device_name",
+            current_platform.get_device_name,
+        ),
+        _safe_debug_value(
+            "torch cuda nccl.version",
+            _torch_nccl_version,
+        ),
+        _safe_debug_value(
+            "torch.accelerator.current_accelerator",
+            torch.accelerator.current_accelerator,
+        ),
+        _safe_debug_value(
+            "current_platform.device_count", current_platform.device_count
+        ),
+        _safe_debug_value(
+            "current_platform.logical_device_id_to_visible_device_id(local_rank)",
+            lambda: current_platform.logical_device_id_to_visible_device_id(local_rank),
+        ),
+        _safe_debug_value(
+            "current_platform.visible_device_id_to_physical_device_id(local_rank)",
+            lambda: current_platform.visible_device_id_to_physical_device_id(
+                local_rank
+            ),
+        ),
+        f"env=({env_values})",
+    ]
+    print("[async_tp_debug] " + ", ".join(debug_values), flush=True)
+
+
 class TestMMRSModel(torch.nn.Module):
     def __init__(self, hidden_size=16, dtype=torch.float16):
         super().__init__()
@@ -333,10 +406,24 @@ def async_tp_pass_on_test_model(
 ):
     set_random_seed(0)
 
+    print(
+        "[async_tp_debug] "
+        f"test_model={test_model_cls.__name__}, batch_size={batch_size}, "
+        f"seq_len={seq_len}, hidden_size={hidden_size}, dtype={dtype}, "
+        f"dynamic={dynamic}",
+        flush=True,
+    )
+    _log_async_tp_debug("before_set_device", local_rank, world_size)
+
     device = torch.device(f"{DEVICE_TYPE}:{local_rank}")
+    print(
+        f"[async_tp_debug] setting accelerator/default device to {device}",
+        flush=True,
+    )
     torch.accelerator.set_device_index(device)
     torch.set_default_device(device)
     torch.set_default_dtype(dtype)
+    _log_async_tp_debug("after_set_device", local_rank, world_size)
 
     update_environment_variables(
         {
@@ -347,9 +434,11 @@ def async_tp_pass_on_test_model(
             "MASTER_PORT": master_port,
         }
     )
+    _log_async_tp_debug("after_update_environment", local_rank, world_size)
 
     # initialize distributed
     init_distributed_environment()
+    _log_async_tp_debug("after_init_distributed_environment", local_rank, world_size)
 
     # configure vllm config for SequenceParallelismPass
     vllm_config = VllmConfig()
@@ -369,6 +458,7 @@ def async_tp_pass_on_test_model(
 
     with set_current_vllm_config(vllm_config):
         initialize_model_parallel(tensor_model_parallel_size=world_size)
+        _log_async_tp_debug("after_initialize_model_parallel", local_rank, world_size)
 
         async_tp_pass = AsyncTPPass(vllm_config)
         backend = TestBackend(async_tp_pass)
@@ -383,18 +473,41 @@ def async_tp_pass_on_test_model(
         )
 
         model = test_model_cls(hidden_size, dtype)  # Pass dtype to model constructor
+        print(
+            "[async_tp_debug] "
+            f"created model={test_model_cls.__name__}, "
+            f"model_ops_before={model.ops_in_model_before()}, "
+            f"model_ops_after={model.ops_in_model_after()}",
+            flush=True,
+        )
 
         hidden_states = torch.randn(
             (batch_size * seq_len, hidden_size), dtype=dtype, requires_grad=False
         )
+        print(
+            "[async_tp_debug] "
+            f"hidden_states.shape={tuple(hidden_states.shape)}, "
+            f"hidden_states.dtype={hidden_states.dtype}, "
+            f"hidden_states.device={hidden_states.device}, "
+            f"hidden_states.is_cuda={hidden_states.is_cuda}",
+            flush=True,
+        )
 
         if dynamic:
             torch._dynamo.mark_dynamic(hidden_states, 0)
+            print("[async_tp_debug] marked hidden_states dim 0 dynamic", flush=True)
 
+        _log_async_tp_debug("before_torch_compile", local_rank, world_size)
         compiled_model = torch.compile(model, backend=backend)
+        _log_async_tp_debug("before_compiled_model_call", local_rank, world_size)
         compiled_model(hidden_states)
+        _log_async_tp_debug("after_compiled_model_call", local_rank, world_size)
 
         assert async_tp_pass.matched_count == 1
+        print(
+            f"[async_tp_debug] matched_count={async_tp_pass.matched_count}",
+            flush=True,
+        )
 
         # In pre-nodes, all gather or reduce scatter should exist,
         # fused_matmul_reduce_scatter or fused_all_gather_matmul should not

From 583cb55742be27dccb5ca0dceb66e593928e2918 Mon Sep 17 00:00:00 2001
From: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Date: Tue, 7 Jul 2026 22:23:28 -0400
Subject: [PATCH 26/26] Remove async TP debug diagnostics
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

The diagnostics served their purpose — the B200 CI issue is resolved and
the test now passes cleanly.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tyler@tylermsmith.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
---
 .../passes/distributed/test_async_tp.py       | 114 ------------------
 1 file changed, 114 deletions(-)

diff --git a/tests/compile/passes/distributed/test_async_tp.py b/tests/compile/passes/distributed/test_async_tp.py
index d1f0cd5f76c6..607e21bc8326 100644
--- a/tests/compile/passes/distributed/test_async_tp.py
+++ b/tests/compile/passes/distributed/test_async_tp.py
@@ -1,8 +1,6 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-import os
-
 import pytest
 import torch
 
@@ -45,78 +43,6 @@
 ]
 
 
-def _safe_debug_value(label: str, fn):
-    try:
-        return f"{label}={fn()}"
-    except Exception as e:
-        return f"{label}=<error {type(e).__name__}: {e}>"
-
-
-def _torch_nccl_version():
-    return torch.cuda.nccl.version()
-
-
-def _log_async_tp_debug(stage: str, local_rank: int, world_size: int) -> None:
-    env_keys = [
-        "CUDA_VISIBLE_DEVICES",
-        "NVIDIA_VISIBLE_DEVICES",
-        "LOCAL_RANK",
-        "RANK",
-        "WORLD_SIZE",
-        "MASTER_ADDR",
-        "MASTER_PORT",
-        "VLLM_RAY_PER_WORKER_GPUS",
-        "VLLM_RAY_BUNDLE_INDICES",
-    ]
-    env_values = ", ".join(f"{key}={os.environ.get(key)!r}" for key in env_keys)
-
-    debug_values = [
-        f"stage={stage}",
-        f"local_rank={local_rank}",
-        f"world_size={world_size}",
-        f"DEVICE_TYPE={DEVICE_TYPE}",
-        _safe_debug_value("torch.__version__", lambda: torch.__version__),
-        _safe_debug_value("torch.version.cuda", lambda: torch.version.cuda),
-        _safe_debug_value(
-            "torch.accelerator.is_available", torch.accelerator.is_available
-        ),
-        _safe_debug_value(
-            "torch.accelerator.device_count", torch.accelerator.device_count
-        ),
-        _safe_debug_value(
-            "torch.accelerator.current_device_index",
-            torch.accelerator.current_device_index,
-        ),
-        _safe_debug_value(
-            "current_platform.get_device_name",
-            current_platform.get_device_name,
-        ),
-        _safe_debug_value(
-            "torch cuda nccl.version",
-            _torch_nccl_version,
-        ),
-        _safe_debug_value(
-            "torch.accelerator.current_accelerator",
-            torch.accelerator.current_accelerator,
-        ),
-        _safe_debug_value(
-            "current_platform.device_count", current_platform.device_count
-        ),
-        _safe_debug_value(
-            "current_platform.logical_device_id_to_visible_device_id(local_rank)",
-            lambda: current_platform.logical_device_id_to_visible_device_id(local_rank),
-        ),
-        _safe_debug_value(
-            "current_platform.visible_device_id_to_physical_device_id(local_rank)",
-            lambda: current_platform.visible_device_id_to_physical_device_id(
-                local_rank
-            ),
-        ),
-        f"env=({env_values})",
-    ]
-    print("[async_tp_debug] " + ", ".join(debug_values), flush=True)
-
-
 class TestMMRSModel(torch.nn.Module):
     def __init__(self, hidden_size=16, dtype=torch.float16):
         super().__init__()
@@ -406,24 +332,10 @@ def async_tp_pass_on_test_model(
 ):
     set_random_seed(0)
 
-    print(
-        "[async_tp_debug] "
-        f"test_model={test_model_cls.__name__}, batch_size={batch_size}, "
-        f"seq_len={seq_len}, hidden_size={hidden_size}, dtype={dtype}, "
-        f"dynamic={dynamic}",
-        flush=True,
-    )
-    _log_async_tp_debug("before_set_device", local_rank, world_size)
-
     device = torch.device(f"{DEVICE_TYPE}:{local_rank}")
-    print(
-        f"[async_tp_debug] setting accelerator/default device to {device}",
-        flush=True,
-    )
     torch.accelerator.set_device_index(device)
     torch.set_default_device(device)
     torch.set_default_dtype(dtype)
-    _log_async_tp_debug("after_set_device", local_rank, world_size)
 
     update_environment_variables(
         {
@@ -434,11 +346,9 @@ def async_tp_pass_on_test_model(
             "MASTER_PORT": master_port,
         }
     )
-    _log_async_tp_debug("after_update_environment", local_rank, world_size)
 
     # initialize distributed
     init_distributed_environment()
-    _log_async_tp_debug("after_init_distributed_environment", local_rank, world_size)
 
     # configure vllm config for SequenceParallelismPass
     vllm_config = VllmConfig()
@@ -458,7 +368,6 @@ def async_tp_pass_on_test_model(
 
     with set_current_vllm_config(vllm_config):
         initialize_model_parallel(tensor_model_parallel_size=world_size)
-        _log_async_tp_debug("after_initialize_model_parallel", local_rank, world_size)
 
         async_tp_pass = AsyncTPPass(vllm_config)
         backend = TestBackend(async_tp_pass)
@@ -473,41 +382,18 @@ def async_tp_pass_on_test_model(
         )
 
         model = test_model_cls(hidden_size, dtype)  # Pass dtype to model constructor
-        print(
-            "[async_tp_debug] "
-            f"created model={test_model_cls.__name__}, "
-            f"model_ops_before={model.ops_in_model_before()}, "
-            f"model_ops_after={model.ops_in_model_after()}",
-            flush=True,
-        )
 
         hidden_states = torch.randn(
             (batch_size * seq_len, hidden_size), dtype=dtype, requires_grad=False
         )
-        print(
-            "[async_tp_debug] "
-            f"hidden_states.shape={tuple(hidden_states.shape)}, "
-            f"hidden_states.dtype={hidden_states.dtype}, "
-            f"hidden_states.device={hidden_states.device}, "
-            f"hidden_states.is_cuda={hidden_states.is_cuda}",
-            flush=True,
-        )
 
         if dynamic:
             torch._dynamo.mark_dynamic(hidden_states, 0)
-            print("[async_tp_debug] marked hidden_states dim 0 dynamic", flush=True)
 
-        _log_async_tp_debug("before_torch_compile", local_rank, world_size)
         compiled_model = torch.compile(model, backend=backend)
-        _log_async_tp_debug("before_compiled_model_call", local_rank, world_size)
         compiled_model(hidden_states)
-        _log_async_tp_debug("after_compiled_model_call", local_rank, world_size)
 
         assert async_tp_pass.matched_count == 1
-        print(
-            f"[async_tp_debug] matched_count={async_tp_pass.matched_count}",
-            flush=True,
-        )
 
         # In pre-nodes, all gather or reduce scatter should exist,
         # fused_matmul_reduce_scatter or fused_all_gather_matmul should not