forked from Karylab-cklius/vllm
Compare commits
5
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
5a532c7e0b | ||
|
|
2f3555bf53 | ||
|
|
4699f1bf8b | ||
|
|
681a6371cc | ||
|
|
e86d349053 |
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
@@ -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
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -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})"
|
||||
)
|
||||
|
||||
@@ -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.
|
||||
|
||||
|
||||
Reference in New Issue
Block a user