[WideEP] Update NCCL to 2.30.7 to enable DeepEPv2 in the vllm/vllm-openai image (#45321)
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com> Signed-off-by: Tyler Michael Smith <tyler@vllm.ai> Signed-off-by: Tyler Michael Smith <tyler@tylermsmith.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Ilya Markov <ilmarkov@users.noreply.github.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Codex <noreply@openai.com>
This commit is contained in:
co-authored by
Claude
Ilya Markov
OpenAI Codex
Codex
parent
e222c33f2f
commit
5d8e90a966
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
+25
-8
@@ -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)
|
||||
# ----------------------------------------------------------------------
|
||||
|
||||
@@ -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"
|
||||
|
||||
@@ -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:
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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]
|
||||
|
||||
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user