From 5d8e90a96616c4fe339ff0b0c2a2d470f6eb24bf Mon Sep 17 00:00:00 2001 From: Tyler Michael Smith Date: Fri, 24 Jul 2026 16:00:02 -0400 Subject: [PATCH] [WideEP] Update NCCL to 2.30.7 to enable DeepEPv2 in the vllm/vllm-openai image (#45321) Signed-off-by: Tyler Michael Smith Signed-off-by: Tyler Michael Smith Signed-off-by: Tyler Michael Smith Co-authored-by: Claude Co-authored-by: Ilya Markov Co-authored-by: OpenAI Codex Co-authored-by: Codex --- .buildkite/test_areas/misc.yaml | 2 +- .buildkite/test_areas/model_runner_v2.yaml | 2 +- docker/Dockerfile | 33 +++- docker/versions.json | 5 +- docs/serving/expert_parallel_deployment.md | 5 + tests/distributed/test_mnnvl_alltoall.py | 12 +- tests/kernels/moe/parallel_utils.py | 48 ++++-- tests/kernels/moe/test_deepep_v2_moe.py | 151 ++++++++++++------ tools/ep_kernels/README.md | 32 ++++ .../device_communicators/all2all.py | 27 ++++ .../device_communicators/pynccl_allocator.py | 8 +- .../device_communicators/pynccl_wrapper.py | 29 ++++ vllm/utils/nccl.py | 62 +++++++ 13 files changed, 338 insertions(+), 78 deletions(-) diff --git a/.buildkite/test_areas/misc.yaml b/.buildkite/test_areas/misc.yaml index 1c2000a9e5..cadbaed5b0 100644 --- a/.buildkite/test_areas/misc.yaml +++ b/.buildkite/test_areas/misc.yaml @@ -212,7 +212,7 @@ steps: - vllm/multimodal - examples/ commands: - - 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 3601aeee11..bc0b56f4ca 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 + - 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 diff --git a/docker/Dockerfile b/docker/Dockerfile index f8c018ef8a..c5d8f006c9 100644 --- a/docker/Dockerfile +++ b/docker/Dockerfile @@ -25,6 +25,10 @@ ARG CUDA_VERSION=13.0.3 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 @@ -477,10 +481,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 \ @@ -644,6 +655,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 @@ -696,7 +708,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} \ @@ -709,12 +720,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 @@ -734,6 +739,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) # ---------------------------------------------------------------------- diff --git a/docker/versions.json b/docker/versions.json index cc145da93b..783d80aace 100644 --- a/docker/versions.json +++ b/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.3-devel-ubuntu22.04" }, @@ -56,7 +59,7 @@ "default": "cuda" }, "DEEPEP_COMMIT_HASH": { - "default": "73b6ea4" + "default": "d4f41e4e93" }, "GIT_REPO_CHECK": { "default": "0" diff --git a/docs/serving/expert_parallel_deployment.md b/docs/serving/expert_parallel_deployment.md index 9a8bb68bc1..b348c5dd96 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/tests/distributed/test_mnnvl_alltoall.py b/tests/distributed/test_mnnvl_alltoall.py index fb2d9bb832..f9b76ea1b0 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: + 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): @@ -916,8 +919,11 @@ def _deepep_v2_lifecycle_worker(rank, world_size): DeepEPV2All2AllManager, ) - cpu_group = get_ep_group().cpu_group - manager = DeepEPV2All2AllManager(cpu_group) + ep_group = get_ep_group() + manager = DeepEPV2All2AllManager( + ep_group.cpu_group, + device_group=ep_group.device_group, + ) assert manager.rank == rank assert manager.world_size == world_size diff --git a/tests/kernels/moe/parallel_utils.py b/tests/kernels/moe/parallel_utils.py index bb2f9efc7c..35a2f198f5 100644 --- a/tests/kernels/moe/parallel_utils.py +++ b/tests/kernels/moe/parallel_utils.py @@ -36,6 +36,10 @@ if has_deep_ep_v2(): P = ParamSpec("P") +class GINNotAvailableError(RuntimeError): + pass + + @dataclasses.dataclass class ProcessGroupInfo: world_size: int @@ -96,19 +100,27 @@ 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: + # pytest.skip cannot propagate directly through torch.multiprocessing. + if "GINNotAvailableError" in str(exc): + import pytest + + pytest.skip("NCCL GIN not available (no IBGDA-capable hardware)") + raise ## DeepEP specific utils @@ -225,6 +237,18 @@ def make_deepep_v2_a2a( ): import deep_ep + from vllm.utils.nccl import query_nccl_gin_type + + # ElasticBuffer can segfault when GIN is unavailable. Initialize the + # lazy communicator and reject unsupported systems before entering DeepEP. + probe = torch.zeros(1, device=pgi.device) + torch.distributed.all_reduce(probe, group=pg) + gin_type = query_nccl_gin_type(pg) + if gin_type is None: + raise RuntimeError("Failed to determine NCCL GIN support") + if gin_type == 0: + raise GINNotAvailableError("NCCL GIN not available") + buffer = deep_ep.ElasticBuffer( group=pg, num_max_tokens_per_rank=v2_args.max_tokens_per_rank, diff --git a/tests/kernels/moe/test_deepep_v2_moe.py b/tests/kernels/moe/test_deepep_v2_moe.py index 93b7c13660..e4cb5bff09 100644 --- a/tests/kernels/moe/test_deepep_v2_moe.py +++ b/tests/kernels/moe/test_deepep_v2_moe.py @@ -14,6 +14,7 @@ from torch.distributed import ProcessGroup from tests.kernels.moe.utils import make_dummy_moe_config, make_test_weights from tests.kernels.utils import torch_experts from vllm.config import VllmConfig, set_current_vllm_config +from vllm.forward_context import set_forward_context from vllm.model_executor.layers.fused_moe import TritonExperts from vllm.model_executor.layers.fused_moe.activation import MoEActivation from vllm.model_executor.layers.fused_moe.config import ( @@ -36,6 +37,14 @@ requires_deep_ep_v2 = pytest.mark.skipif( ) +def assert_fp8_close(actual: torch.Tensor, expected: torch.Tensor) -> None: + close = torch.isclose(actual, expected, atol=2e-1, rtol=2e-1) + close_fraction = close.float().mean().item() + assert close_fraction > 0.99, ( + f"Only {close_fraction:.1%} of FP8 outputs are within tolerance" + ) + + @dataclasses.dataclass class TestConfig: dtype: torch.dtype @@ -50,6 +59,7 @@ class TestConfig: class TestTensors: rank_tokens: torch.Tensor rank_token_scales: torch.Tensor | None + intermediate_scales: torch.Tensor | None topk: torch.Tensor topk_weights: torch.Tensor config: TestConfig @@ -63,6 +73,12 @@ class TestTensors: rank_tokens = ( torch.randn((config.m, config.k), device="cuda", dtype=token_dtype) / 10 ) + if config.dtype == torch.float8_e4m3fn: + rank_token_scales = torch.tensor(1 / 448, device="cuda") + intermediate_scales = torch.tensor(8 / 448, device="cuda") + else: + rank_token_scales = None + intermediate_scales = None topk = torch.stack( [ @@ -73,7 +89,8 @@ class TestTensors: topk_weights = torch.randn(topk.shape, dtype=torch.float32, device="cuda") return TestTensors( rank_tokens=rank_tokens, - rank_token_scales=None, + rank_token_scales=rank_token_scales, + intermediate_scales=intermediate_scales, topk=topk, topk_weights=topk_weights, config=config, @@ -124,7 +141,6 @@ def make_modular_kernel( mk = FusedMoEKernel( prepare_finalize=a2a, fused_experts=fused_experts, - inplace=False, ) return mk @@ -162,6 +178,7 @@ def deepep_v2_moe_impl( w2_scale=w2_scale, per_act_token_quant=per_act_token_quant, a1_scale=test_tensors.rank_token_scales, + a2_scale=test_tensors.intermediate_scales, ) hidden_size = test_tensors.rank_tokens.size(1) @@ -231,6 +248,8 @@ def _deep_ep_v2_moe( test_tensors.topk, w1_scale=w1_scale, w2_scale=w2_scale, + a1_scale=test_tensors.rank_token_scales, + a2_scale=test_tensors.intermediate_scales, quant_dtype=q_dtype, per_act_token_quant=per_act_token_quant, ) @@ -262,12 +281,15 @@ def _deep_ep_v2_moe( per_act_token_quant, ) - torch.testing.assert_close( - torch_combined, - deepep_combined, - atol=6e-2, - rtol=6e-2, - ) + if is_quantized: + assert_fp8_close(torch_combined, deepep_combined) + else: + torch.testing.assert_close( + torch_combined, + deepep_combined, + atol=6e-2, + rtol=6e-2, + ) MNKs = [ @@ -355,17 +377,29 @@ def _deep_ep_v2_moe_cudagraph( num_local_experts = config.num_experts // pgi.world_size hidden_size = config.k - # Create FP8 weights directly, then dequantize for bf16 reference. - w1_fp8 = torch.randn( - (config.num_experts, 2 * config.n, config.k), - device="cuda", - dtype=torch.bfloat16, - ).to(torch.float8_e4m3fn) - w2_fp8 = torch.randn( - (config.num_experts, config.k, config.n), - device="cuda", - dtype=torch.bfloat16, - ).to(torch.float8_e4m3fn) + # All ranks must use the same global weights before taking their EP slice. + w1_bf16 = ( + torch.randn( + (config.num_experts, 2 * config.n, config.k), + device="cuda", + dtype=torch.bfloat16, + ) + / 15 + ) + w2_bf16 = ( + torch.randn( + (config.num_experts, config.k, config.n), + device="cuda", + dtype=torch.bfloat16, + ) + / 15 + ) + torch.distributed.broadcast(w1_bf16, src=0, group=pg) + torch.distributed.broadcast(w2_bf16, src=0, group=pg) + + # Round-trip through FP8 before constructing the reference and kernel weights. + w1_fp8 = w1_bf16.to(torch.float8_e4m3fn) + w2_fp8 = w2_bf16.to(torch.float8_e4m3fn) w1_ref = w1_fp8.to(torch.bfloat16) w2_ref = w2_fp8.to(torch.bfloat16) @@ -385,21 +419,8 @@ def _deep_ep_v2_moe_cudagraph( backend="nccl", ) initialize_model_parallel(tensor_model_parallel_size=1) - # Reference MoE using dequantized bf16 weights - torch_combined = torch_experts( - test_tensors.rank_tokens, - w1_ref, - w2_ref, - test_tensors.topk_weights, - test_tensors.topk, - ) - - # Use the production pipeline: make_fused_moe_layer creates - # a FusedMoE layer, quantizes weights, runs - # process_weights_after_loading (TrtLLM W31 swap + BlockMajorK - # shuffle), and selects the kernel. - # Quantize weights using production helper, EP-slice, then - # convert to TrtLLM format. + # Mirror production weight processing: quantize, EP-slice, then + # convert to the TrtLLM BlockMajorK format. from tests.kernels.moe.test_moe_layer import _quantize_fp8_halves from vllm.model_executor.layers.fused_moe.experts.trtllm_fp8_moe import ( TrtLlmFp8ExpertsModular, @@ -411,14 +432,32 @@ def _deep_ep_v2_moe_cudagraph( block_shape = [128, 128] qw = _quantize_fp8_halves(w1_ref, w2_ref, block_shape) + assert qw.w13_weight_scale is not None + assert qw.w2_weight_scale is not None + + # Reference MoE using the same blockwise FP8 quantization scheme as + # the production kernel. torch_experts quantizes activations before + # both GEMMs and dequantizes the operands for the reference matmuls. + reference_topk_weights = test_tensors.topk_weights.to(torch.bfloat16).to( + torch.float32 + ) + torch_combined = torch_experts( + test_tensors.rank_tokens, + qw.w13_weight, + qw.w2_weight, + reference_topk_weights, + test_tensors.topk, + w1_scale=qw.w13_weight_scale, + w2_scale=qw.w2_weight_scale, + quant_dtype=torch.float8_e4m3fn, + block_shape=block_shape, + ) # EP-slice before format conversion e_start = num_local_experts * pgi.rank e_end = e_start + num_local_experts w1_ep = qw.w13_weight[e_start:e_end] w2_ep = qw.w2_weight[e_start:e_end] - assert qw.w13_weight_scale is not None - assert qw.w2_weight_scale is not None w1_scale_ep = qw.w13_weight_scale[e_start:e_end] w2_scale_ep = qw.w2_weight_scale[e_start:e_end] @@ -452,11 +491,23 @@ def _deep_ep_v2_moe_cudagraph( w2_scale=w2_scale_ep, ) moe_config = make_dummy_moe_config( - num_experts=num_local_experts, + num_experts=config.num_experts, + num_local_experts=num_local_experts, experts_per_token=config.topk, hidden_dim=hidden_size, intermediate_size=config.n, ) + moe_parallel_config = dataclasses.replace( + moe_config.moe_parallel_config, + ep_size=pgi.world_size, + ep_rank=pgi.rank, + use_ep=True, + all2all_backend="deepep_v2", + ) + moe_config = dataclasses.replace( + moe_config, + moe_parallel_config=moe_parallel_config, + ) fused_experts = TrtLlmFp8ExpertsModular( moe_config=moe_config, quant_config=quant_config, @@ -480,21 +531,21 @@ def _deep_ep_v2_moe_cudagraph( mk_kernel = FusedMoEKernel( prepare_finalize=a2a, fused_experts=fused_experts, - inplace=False, ) - for _ in range(3): - out = mk_kernel.apply( - hidden_states=test_tensors.rank_tokens, - w1=w1_ep, - w2=w2_ep, - topk_weights=test_tensors.topk_weights, - topk_ids=test_tensors.topk, - activation=MoEActivation.SILU, - global_num_experts=config.num_experts, - expert_map=None, - apply_router_weight_on_input=False, - ) + with set_forward_context(None, vllm_cfg): + for _ in range(3): + out = mk_kernel.apply( + hidden_states=test_tensors.rank_tokens, + w1=w1_ep, + w2=w2_ep, + topk_weights=test_tensors.topk_weights, + topk_ids=test_tensors.topk, + activation=MoEActivation.SILU, + global_num_experts=config.num_experts, + expert_map=None, + apply_router_weight_on_input=False, + ) torch.testing.assert_close( torch_combined, diff --git a/tools/ep_kernels/README.md b/tools/ep_kernels/README.md index b4eabe18ca..04bc3ef276 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 diff --git a/vllm/distributed/device_communicators/all2all.py b/vllm/distributed/device_communicators/all2all.py index 33841de306..679764f6a8 100644 --- a/vllm/distributed/device_communicators/all2all.py +++ b/vllm/distributed/device_communicators/all2all.py @@ -978,6 +978,7 @@ class DeepEPV2All2AllManager(All2AllManagerBase): self._device_group = device_group self.handle_cache = Cache() self._num_sms: int | None = None + self._gin_checked = False def _make_all2all_kwargs( self, @@ -1000,11 +1001,37 @@ class DeepEPV2All2AllManager(All2AllManagerBase): explicitly_destroy=True, ) + def _check_gin_support(self, group) -> None: + from vllm.utils.nccl import query_nccl_gin_type + + # ProcessGroupNCCL creates communicators lazily. Initialize this exact + # group before querying so a null comm pointer is not mistaken for + # missing GIN support. + probe = torch.zeros(1, device="cuda") + torch.distributed.all_reduce(probe, group=group) + + gin_type = query_nccl_gin_type(group) + if gin_type is None: + raise RuntimeError( + "DeepEPv2 communicator properties query failed; " + "networking capability could not be determined." + ) + if gin_type == 0: + raise RuntimeError( + "DeepEPv2 requires NCCL GIN (GPU-Initiated Networking). " + "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_allocator.py b/vllm/distributed/device_communicators/pynccl_allocator.py index a5d7faceb2..4898c6212f 100644 --- a/vllm/distributed/device_communicators/pynccl_allocator.py +++ b/vllm/distributed/device_communicators/pynccl_allocator.py @@ -14,7 +14,7 @@ from vllm import envs 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, diff --git a/vllm/distributed/device_communicators/pynccl_wrapper.py b/vllm/distributed/device_communicators/pynccl_wrapper.py index 5ca8cc7c77..78ee81e7b5 100644 --- a/vllm/distributed/device_communicators/pynccl_wrapper.py +++ b/vllm/distributed/device_communicators/pynccl_wrapper.py @@ -51,6 +51,26 @@ class ncclUniqueId(ctypes.Structure): _fields_ = [("internal", ctypes.c_byte * 128)] +# NCCL 2.30+ ncclCommProperties_t. Only fields through ginType are read; +# trailing fields keep the layout aligned with NCCL's versioned structure. +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), + ] + + cudaStream_t = ctypes.c_void_p buffer_type = ctypes.c_void_p @@ -317,6 +337,12 @@ class NCCLLibrary: # ncclResult_t ncclCommWindowDeregister( # ncclComm_t comm, ncclWindow_t win); Function("ncclCommWindowDeregister", ncclResult_t, [ncclComm_t, ncclWindow_t]), + # Query runtime properties of a specific initialized communicator. + Function( + "ncclCommQueryProperties", + ncclResult_t, + [ncclComm_t, ctypes.POINTER(ncclCommProperties)], + ), ] # class attribute to store the mapping from the path to the library @@ -375,6 +401,9 @@ class NCCLLibrary: # Having an exception here on ROCm platform is # not allowed during graph capturing continue + elif func.name == "ncclCommQueryProperties": + # Optional on NCCL versions older than 2.29. + 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 4807bc076f..aeadaf5d22 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 @@ -62,3 +63,64 @@ 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 + + +def query_nccl_gin_type(group: torch.distributed.ProcessGroup) -> int | None: + """Return the GIN type for an initialized group, or ``None`` on failure.""" + from vllm.distributed.device_communicators.pynccl_wrapper import ( + NCCLLibrary, + ncclCommProperties, + ) + + try: + backend = group._get_backend(torch.device("cuda")) + # GIN is a property of this initialized communicator, not just the + # NCCL version. ncclCommQueryProperties requires its ncclComm_t. + comm_ptr = backend._comm_ptr() + if comm_ptr == 0: + return None + except Exception: + logger.warning( + "Failed to extract NCCL comm pointer from process group", + exc_info=True, + ) + return None + + try: + nccl = NCCLLibrary() + query_fn = nccl._funcs.get("ncclCommQueryProperties") + if query_fn is None: + return None + + props = ncclCommProperties() + ctypes.memset(ctypes.addressof(props), 0, ctypes.sizeof(props)) + props.size = ctypes.sizeof(props) + props.magic = 0xCAFEBEEF + props.version = nccl.ncclGetRawVersion() + result = query_fn(ctypes.c_void_p(comm_ptr), ctypes.byref(props)) + except Exception: + logger.warning("Failed to query NCCL communicator properties", exc_info=True) + return None + + if result != 0: + logger.warning("ncclCommQueryProperties returned error %d", result) + return None + return props.ginType