forked from Karylab-cklius/vllm
Compare 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 int idx = thread_row_offset + ii;
|
||||||
const float val = toFloat(input[idx]);
|
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;
|
output[idx] = softmax_val;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
@@ -147,7 +149,9 @@ __launch_bounds__(TPB) __global__
|
|||||||
{
|
{
|
||||||
const int idx = thread_row_offset + ii;
|
const int idx = thread_row_offset + ii;
|
||||||
const float val = toFloat(input[idx]);
|
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;
|
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;
|
static constexpr int COLS_PER_GROUP_LDG = ELTS_PER_LDG * THREADS_PER_ROW;
|
||||||
|
|
||||||
// If bias is not null, use biased value for selection
|
// 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 "Detected vLLM version: ${VLLM_VERSION}" \
|
||||||
&& echo "${VLLM_VERSION}" > /tmp/vllm_version.txt
|
&& 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
|
# Pin vLLM dependencies to exact versions of custom ROCm wheels
|
||||||
# This ensures 'pip install vllm' automatically installs correct torch/triton/torchvision/amdsmi
|
# 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
|
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.
|
used to profile a section of code.
|
||||||
|
|
||||||
!!! note
|
!!! note
|
||||||
The legacy import paths `vllm.utils.cprofile` and `vllm.utils.cprofile_context` are deprecated.
|
The `vllm.utils.profiling` helpers are deprecated and will be removed in
|
||||||
Please use `vllm.utils.profiling.cprofile` and `vllm.utils.profiling.cprofile_context` instead.
|
`v0.21`. Please use Python's `cProfile` module directly instead.
|
||||||
|
|
||||||
### Example usage - decorator
|
### Example usage - decorator
|
||||||
|
|
||||||
|
|||||||
@@ -32,9 +32,7 @@ pyzmq >= 25.0.0
|
|||||||
msgspec
|
msgspec
|
||||||
gguf >= 0.17.0
|
gguf >= 0.17.0
|
||||||
mistral_common[image] >= 1.11.0
|
mistral_common[image] >= 1.11.0
|
||||||
av # required for audio in video IO
|
|
||||||
opencv-python-headless >= 4.13.0 # required for video IO
|
opencv-python-headless >= 4.13.0 # required for video IO
|
||||||
soundfile # required for audio IO
|
|
||||||
pyyaml
|
pyyaml
|
||||||
six>=1.16.0; python_version > '3.11' # transitive dependency of pandas that needs to be the latest version for python 3.12
|
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
|
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
|
timm>=1.0.17
|
||||||
# amd-quark: required for Quark quantization on ROCm
|
# amd-quark: required for Quark quantization on ROCm
|
||||||
# To be consistent with test_quark.py
|
# To be consistent with test_quark.py
|
||||||
amd-quark>=0.8.99
|
amd-quark>=0.8.99
|
||||||
# Required for faster safetensors model loading
|
|
||||||
fastsafetensors @ git+https://github.com/foundation-model-stack/fastsafetensors.git@0.2.2
|
|
||||||
@@ -76,9 +76,7 @@ attrs==26.1.0
|
|||||||
audioread==3.0.1
|
audioread==3.0.1
|
||||||
# via librosa
|
# via librosa
|
||||||
av==16.1.0
|
av==16.1.0
|
||||||
# via
|
# via -r requirements/test/rocm.in
|
||||||
# -r requirements/test/../common.txt
|
|
||||||
# -r requirements/test/rocm.in
|
|
||||||
azure-core==1.39.0
|
azure-core==1.39.0
|
||||||
# via
|
# via
|
||||||
# azure-identity
|
# azure-identity
|
||||||
@@ -278,9 +276,7 @@ fastar==0.10.0
|
|||||||
fastparquet==2026.3.0
|
fastparquet==2026.3.0
|
||||||
# via genai-perf
|
# via genai-perf
|
||||||
fastsafetensors @ git+https://github.com/foundation-model-stack/fastsafetensors.git@65d80088fca7a8f567fba30415fbcc80f7d2259c
|
fastsafetensors @ git+https://github.com/foundation-model-stack/fastsafetensors.git@65d80088fca7a8f567fba30415fbcc80f7d2259c
|
||||||
# via
|
# via -r requirements/test/rocm.in
|
||||||
# -c requirements/rocm.txt
|
|
||||||
# -r requirements/test/rocm.in
|
|
||||||
filelock==3.25.2
|
filelock==3.25.2
|
||||||
# via
|
# via
|
||||||
# -c requirements/common.txt
|
# -c requirements/common.txt
|
||||||
@@ -1333,7 +1329,6 @@ sortedcontainers==2.4.0
|
|||||||
# via hypothesis
|
# via hypothesis
|
||||||
soundfile==0.13.1
|
soundfile==0.13.1
|
||||||
# via
|
# via
|
||||||
# -r requirements/test/../common.txt
|
|
||||||
# -r requirements/test/rocm.in
|
# -r requirements/test/rocm.in
|
||||||
# genai-perf
|
# genai-perf
|
||||||
# librosa
|
# librosa
|
||||||
|
|||||||
@@ -1094,7 +1094,9 @@ setup(
|
|||||||
"instanttensor": ["instanttensor >= 0.1.5"],
|
"instanttensor": ["instanttensor >= 0.1.5"],
|
||||||
"runai": ["runai-model-streamer[s3,gcs,azure] >= 0.15.7"],
|
"runai": ["runai-model-streamer[s3,gcs,azure] >= 0.15.7"],
|
||||||
"audio": [
|
"audio": [
|
||||||
|
"av",
|
||||||
"scipy",
|
"scipy",
|
||||||
|
"soundfile",
|
||||||
"mistral_common[audio]",
|
"mistral_common[audio]",
|
||||||
], # Required for audio processing
|
], # Required for audio processing
|
||||||
"video": [], # Kept for backwards compatibility
|
"video": [], # Kept for backwards compatibility
|
||||||
|
|||||||
@@ -38,6 +38,8 @@ llm = LLM(
|
|||||||
distributed_executor_backend="external_launcher",
|
distributed_executor_backend="external_launcher",
|
||||||
gpu_memory_utilization=random.uniform(0.7, 0.9),
|
gpu_memory_utilization=random.uniform(0.7, 0.9),
|
||||||
seed=0,
|
seed=0,
|
||||||
|
max_model_len=1024,
|
||||||
|
max_num_seqs=16,
|
||||||
)
|
)
|
||||||
|
|
||||||
outputs = llm.generate(prompts, sampling_params)
|
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
|
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)
|
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 functools import wraps
|
||||||
from typing import Any
|
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
|
@contextlib.contextmanager
|
||||||
def cprofile_context(save_file: str | None = None):
|
def cprofile_context(save_file: str | None = None):
|
||||||
"""Run a cprofile
|
"""Run a cprofile
|
||||||
@@ -32,6 +38,10 @@ def cprofile_context(save_file: str | None = None):
|
|||||||
prof.print_stats(sort="cumtime")
|
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):
|
def cprofile(save_file: str | None = None, enabled: bool = True):
|
||||||
"""Decorator to profile a Python method using cProfile.
|
"""Decorator to profile a Python method using cProfile.
|
||||||
|
|
||||||
|
|||||||
Reference in New Issue
Block a user