Compare commits

...
Author SHA1 Message Date
Isotr0pyandkhluu 5a532c7e0b Revert "[Misc] Move pyav and soundfile to common requirements" (#40276)
Co-authored-by: Roger Wang <hey@rogerw.io>
(cherry picked from commit 67eb6083e3)
2026-04-22 02:00:47 -07:00
Wentao Yeandkhluu 2f3555bf53 [Deprecation] Deprecate cprofile and cprofile_context (#39100)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
(cherry picked from commit 301024aa9c)
2026-04-22 02:00:12 -07:00
Jhao-Ting Chenandkhluu 4699f1bf8b fix: clamp NaN/Inf in topk_softmax to prevent duplicate expert IDs (#39391)
Signed-off-by: Jhao-Ting Chen <jhaotingc@nvidia.com>
(cherry picked from commit 28c222157b)
2026-04-22 01:59:52 -07:00
Nicolò Lucchesiandkhluu 681a6371cc [Bugfix][CI] Fix tests/distributed/test_torchrun_example_moe.py (#40349)
Signed-off-by: NickLucche <nlucches@redhat.com>
(cherry picked from commit 304d5ba1a0)
2026-04-22 01:59:37 -07:00
TJianandkhluu e86d349053 [ROCm] [Wheel] [Bugfix] [Critical] Remove any packages installed from github from rocm.txt e.g fastsafetensors as it is incompatible with uv pip (#40461)
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
(cherry picked from commit 583e6f2226)
2026-04-22 01:58:51 -07:00
10 changed files with 129 additions and 16 deletions
+19 -2
View File
@@ -126,7 +126,9 @@ __launch_bounds__(TPB) __global__
{
const int idx = thread_row_offset + ii;
const float val = toFloat(input[idx]);
const float softmax_val = expf(val - float_max) * normalizing_factor;
float softmax_val = expf(val - float_max) * normalizing_factor;
// Clamp NaN/Inf to 0 to prevent duplicate expert IDs downstream.
if (isnan(softmax_val) || isinf(softmax_val)) softmax_val = 0.f;
output[idx] = softmax_val;
}
}
@@ -147,7 +149,9 @@ __launch_bounds__(TPB) __global__
{
const int idx = thread_row_offset + ii;
const float val = toFloat(input[idx]);
const float sigmoid_val = 1.0f / (1.0f + __expf(-val));
float sigmoid_val = 1.0f / (1.0f + __expf(-val));
// Clamp NaN/Inf to 0 to prevent duplicate expert IDs downstream.
if (isnan(sigmoid_val) || isinf(sigmoid_val)) sigmoid_val = 0.f;
output[idx] = sigmoid_val;
}
}
@@ -442,6 +446,19 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
}
}
// Fix: clamp NaN/Inf values to 0 to prevent duplicate expert IDs.
// NaN gating (from degenerate hidden states in CUDA graph padding) causes
// softmax to produce all-NaN, which makes the argmax loop always pick
// expert 0 for every top-k slot, producing duplicate expert IDs that
// crash FlashInfer's three-step MoE sort.
// With 0s, the argmax uses index tie-breaking to pick [0,1,2,...,k-1].
#pragma unroll
for (int ii = 0; ii < VPT; ++ii) {
if (isnan(row_chunk[ii]) || isinf(row_chunk[ii])) {
row_chunk[ii] = 0.f;
}
}
static constexpr int COLS_PER_GROUP_LDG = ELTS_PER_LDG * THREADS_PER_ROW;
// If bias is not null, use biased value for selection
+24
View File
@@ -261,6 +261,30 @@ RUN --mount=type=bind,source=.git,target=vllm/.git \
&& echo "Detected vLLM version: ${VLLM_VERSION}" \
&& echo "${VLLM_VERSION}" > /tmp/vllm_version.txt
# Fail if git-based package dependencies are found in requirements files
# (uv doesn't handle git+ URLs well, and packages should be distributed on PyPI)
# Extra notes: pip install is able to handle git+ URLs, but uv doesn't.
RUN echo "Checking for git-based packages in requirements files..." \
&& echo "Checking common.txt for git-based packages:" \
&& if grep -q 'git+' ${COMMON_WORKDIR}/vllm/requirements/common.txt; then \
echo "ERROR: Git-based packages found in common.txt:"; \
grep 'git+' ${COMMON_WORKDIR}/vllm/requirements/common.txt; \
echo "Please publish these packages to PyPI instead of using git dependencies."; \
exit 1; \
else \
echo " ✓ No git-based packages found in common.txt"; \
fi \
&& echo "Checking rocm.txt for git-based packages:" \
&& if grep -q 'git+' ${COMMON_WORKDIR}/vllm/requirements/rocm.txt; then \
echo "ERROR: Git-based packages found in rocm.txt:"; \
grep 'git+' ${COMMON_WORKDIR}/vllm/requirements/rocm.txt; \
echo "Please publish these packages to PyPI instead of using git dependencies."; \
exit 1; \
else \
echo " ✓ No git-based packages found in rocm.txt"; \
fi \
&& echo "All requirements files are clean - no git-based packages found"
# Pin vLLM dependencies to exact versions of custom ROCm wheels
# This ensures 'pip install vllm' automatically installs correct torch/triton/torchvision/amdsmi
COPY tools/vllm-rocm/pin_rocm_dependencies.py /tmp/pin_rocm_dependencies.py
+2 -2
View File
@@ -206,8 +206,8 @@ Both the `vllm.utils.profiling.cprofile` and `vllm.utils.profiling.cprofile_cont
used to profile a section of code.
!!! note
The legacy import paths `vllm.utils.cprofile` and `vllm.utils.cprofile_context` are deprecated.
Please use `vllm.utils.profiling.cprofile` and `vllm.utils.profiling.cprofile_context` instead.
The `vllm.utils.profiling` helpers are deprecated and will be removed in
`v0.21`. Please use Python's `cProfile` module directly instead.
### Example usage - decorator
-2
View File
@@ -32,9 +32,7 @@ pyzmq >= 25.0.0
msgspec
gguf >= 0.17.0
mistral_common[image] >= 1.11.0
av # required for audio in video IO
opencv-python-headless >= 4.13.0 # required for video IO
soundfile # required for audio IO
pyyaml
six>=1.16.0; python_version > '3.11' # transitive dependency of pandas that needs to be the latest version for python 3.12
setuptools>=77.0.3,<81.0.0; python_version > '3.11' # Setuptools is used by triton, we need to ensure a modern version is installed for 3.12+ so that it does not try to import distutils, which was removed in 3.12
+1 -3
View File
@@ -20,6 +20,4 @@ conch-triton-kernels==1.2.1
timm>=1.0.17
# amd-quark: required for Quark quantization on ROCm
# To be consistent with test_quark.py
amd-quark>=0.8.99
# Required for faster safetensors model loading
fastsafetensors @ git+https://github.com/foundation-model-stack/fastsafetensors.git@0.2.2
amd-quark>=0.8.99
+2 -7
View File
@@ -76,9 +76,7 @@ attrs==26.1.0
audioread==3.0.1
# via librosa
av==16.1.0
# via
# -r requirements/test/../common.txt
# -r requirements/test/rocm.in
# via -r requirements/test/rocm.in
azure-core==1.39.0
# via
# azure-identity
@@ -278,9 +276,7 @@ fastar==0.10.0
fastparquet==2026.3.0
# via genai-perf
fastsafetensors @ git+https://github.com/foundation-model-stack/fastsafetensors.git@65d80088fca7a8f567fba30415fbcc80f7d2259c
# via
# -c requirements/rocm.txt
# -r requirements/test/rocm.in
# via -r requirements/test/rocm.in
filelock==3.25.2
# via
# -c requirements/common.txt
@@ -1333,7 +1329,6 @@ sortedcontainers==2.4.0
# via hypothesis
soundfile==0.13.1
# via
# -r requirements/test/../common.txt
# -r requirements/test/rocm.in
# genai-perf
# librosa
+2
View File
@@ -1094,7 +1094,9 @@ setup(
"instanttensor": ["instanttensor >= 0.1.5"],
"runai": ["runai-model-streamer[s3,gcs,azure] >= 0.15.7"],
"audio": [
"av",
"scipy",
"soundfile",
"mistral_common[audio]",
], # Required for audio processing
"video": [], # Kept for backwards compatibility
@@ -38,6 +38,8 @@ llm = LLM(
distributed_executor_backend="external_launcher",
gpu_memory_utilization=random.uniform(0.7, 0.9),
seed=0,
max_model_len=1024,
max_num_seqs=16,
)
outputs = llm.generate(prompts, sampling_params)
+67
View File
@@ -135,3 +135,70 @@ def test_fused_topk_bias(
topk_weights_ref.to(torch.float32), topk_weights, atol=1e-2, rtol=1e-2
)
torch.testing.assert_close(topk_ids_ref.to(torch.int32), topk_ids, atol=0, rtol=0)
@pytest.mark.skipif(
not current_platform.is_cuda(), reason="This test is skipped on non-CUDA platform."
)
@pytest.mark.parametrize("num_experts", [6, 8, 16])
@pytest.mark.parametrize("topk", [3, 4])
@pytest.mark.parametrize("scoring_func", ["softmax", "sigmoid"])
@pytest.mark.parametrize("bad_value", [float("nan"), float("inf")])
@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.half, torch.float32])
def test_fused_topk_nan_inf_clamp(
num_experts: int,
topk: int,
scoring_func: str,
bad_value: float,
dtype: torch.dtype,
):
"""Regression test for the NaN/Inf clamp in topk_softmax_kernels.cu.
Degenerate hidden states (e.g., from CUDA graph padding) can produce
NaN/Inf gating logits. Without the clamp, softmax/sigmoid outputs are
NaN and the argmax loop picks expert 0 for every top-k slot (since
"NaN > NaN" is false per IEEE 754), yielding duplicate expert IDs that
crash downstream MoE sort kernels. The fix clamps NaN/Inf to 0 before
argmax so index tie-breaking selects unique experts [0, 1, ..., k-1].
"""
torch.manual_seed(0)
num_tokens = 4
hidden_size = 1024
hidden_states = torch.randn((num_tokens, hidden_size), dtype=dtype, device="cuda")
# Row 0: all normal. Rows 1-3: fully poisoned with NaN or Inf.
gating_output = torch.randn((num_tokens, num_experts), dtype=dtype, device="cuda")
gating_output[1:, :] = bad_value
topk_weights, topk_ids, _ = fused_topk(
hidden_states=hidden_states,
gating_output=gating_output,
topk=topk,
renormalize=False,
scoring_func=scoring_func,
)
# Normal row must still match the torch reference.
ref_weights, ref_ids = torch_topk(
gating_output=gating_output[:1],
topk=topk,
renormalize=False,
scoring_func=scoring_func,
)
torch.testing.assert_close(
ref_weights.to(torch.float32), topk_weights[:1], atol=1e-2, rtol=1e-2
)
torch.testing.assert_close(ref_ids.to(torch.int32), topk_ids[:1], atol=0, rtol=0)
# Poisoned rows: IDs must be unique (no duplicates) and weights must be
# finite (no NaN/Inf propagation into downstream MoE kernels).
for row in range(1, num_tokens):
row_ids = topk_ids[row]
assert row_ids.unique().numel() == topk, (
f"Row {row} has duplicate expert IDs {row_ids.tolist()} "
f"(bad_value={bad_value}, scoring_func={scoring_func})"
)
assert torch.isfinite(topk_weights[row]).all(), (
f"Row {row} has non-finite weights {topk_weights[row].tolist()} "
f"(bad_value={bad_value}, scoring_func={scoring_func})"
)
+10
View File
@@ -8,7 +8,13 @@ from collections.abc import Callable
from functools import wraps
from typing import Any
from typing_extensions import deprecated
@deprecated(
"vllm.utils.profiling.cprofile_context() is deprecated and will be removed "
"in v0.21. Use Python's cProfile module directly instead."
)
@contextlib.contextmanager
def cprofile_context(save_file: str | None = None):
"""Run a cprofile
@@ -32,6 +38,10 @@ def cprofile_context(save_file: str | None = None):
prof.print_stats(sort="cumtime")
@deprecated(
"vllm.utils.profiling.cprofile() is deprecated and will be removed in "
"v0.21. Use Python's cProfile module directly instead."
)
def cprofile(save_file: str | None = None, enabled: bool = True):
"""Decorator to profile a Python method using cProfile.