diff --git a/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu b/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu index 78aed496bb..98a04de783 100644 --- a/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu +++ b/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu @@ -532,7 +532,9 @@ __global__ void moe_lora_align_block_size_kernel( // Populate the token_mask based on the token-LoRA mapping int num_tokens = numel / topk_num; - if (threadIdx.x == 0) { + // Only the even counting block owns per-LoRA metadata. The odd block only + // initializes sorted_token_ids and must not race the final count write. + if (blockIdx.x % 2 == 0 && threadIdx.x == 0) { total_tokens_post_pad[lora_id] = 0; for (int i = 0; i < num_tokens; i++) { diff --git a/examples/rl/rlhf_async_new_apis.py b/examples/rl/rlhf_async_new_apis.py index acd1a42128..27d3f4b4e6 100644 --- a/examples/rl/rlhf_async_new_apis.py +++ b/examples/rl/rlhf_async_new_apis.py @@ -173,8 +173,6 @@ ray_env_vars = {} if current_platform.is_rocm(): # Workaround for RCCL bug. See https://github.com/ROCm/rocm-systems/issues/5756 ray_env_vars["RAY_EXPERIMENTAL_NOSET_HIP_VISIBLE_DEVICES"] = "1" - # For ROCm, BATCH_INVARIANT vllm is not supported - ray_env_vars["VLLM_ROCM_USE_SKINNY_GEMM"] = "0" else: # Enable batch invariance for deterministic outputs on NVIDIA ray_env_vars["VLLM_BATCH_INVARIANT"] = "1" diff --git a/tests/entrypoints/multimodal/openai/chat_completion/test_vision.py b/tests/entrypoints/multimodal/openai/chat_completion/test_vision.py index b33311f8af..e6b8dc23e9 100644 --- a/tests/entrypoints/multimodal/openai/chat_completion/test_vision.py +++ b/tests/entrypoints/multimodal/openai/chat_completion/test_vision.py @@ -9,7 +9,7 @@ import pytest_asyncio from transformers import AutoProcessor from tests.entrypoints.multimodal.conftest import TEST_IMAGE_ASSETS -from tests.utils import ROCM_ENV_OVERRIDES, ROCM_EXTRA_ARGS, RemoteOpenAIServer +from tests.utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer from vllm.multimodal.media import MediaWithBytes from vllm.multimodal.utils import encode_image_url, fetch_image from vllm.platforms import current_platform @@ -92,7 +92,6 @@ def server(): # ROCm: Increase timeouts to handle potential network delays and slower # video processing when downloading multiple videos from external sources env_overrides = { - **ROCM_ENV_OVERRIDES, **( { "VLLM_VIDEO_FETCH_TIMEOUT": "120", diff --git a/tests/entrypoints/openai/chat_completion/test_completion_with_function_calling.py b/tests/entrypoints/openai/chat_completion/test_completion_with_function_calling.py index 33cb576f35..3e59e7402e 100644 --- a/tests/entrypoints/openai/chat_completion/test_completion_with_function_calling.py +++ b/tests/entrypoints/openai/chat_completion/test_completion_with_function_calling.py @@ -10,7 +10,7 @@ import pytest import pytest_asyncio # downloading lora to test lora requests -from tests.utils import ROCM_ENV_OVERRIDES, ROCM_EXTRA_ARGS, RemoteOpenAIServer +from tests.utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer # any model with a chat template should work here MODEL_NAME = "Qwen/Qwen3-0.6B" @@ -144,9 +144,7 @@ def server(): "--enforce-eager", ] + ROCM_EXTRA_ARGS - with RemoteOpenAIServer( - MODEL_NAME, args, env_dict=ROCM_ENV_OVERRIDES - ) as remote_server: + with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/basic/test_encode.py b/tests/entrypoints/pooling/basic/test_encode.py index ab3a0610c3..ca85d2758f 100644 --- a/tests/entrypoints/pooling/basic/test_encode.py +++ b/tests/entrypoints/pooling/basic/test_encode.py @@ -7,7 +7,6 @@ import pytest from vllm import LLM, PoolingParams from vllm.distributed import cleanup_dist_env_and_memory -from vllm.platforms import current_platform MODEL_NAME = "intfloat/multilingual-e5-small" @@ -30,12 +29,6 @@ TOKEN_IDS = [ @pytest.fixture(scope="module") def llm(): - # ROCm: Use FLEX_ATTENTION backend as it's the only attention backend - # that supports encoder-only models on ROCm. - attention_config = None - if current_platform.is_rocm(): - attention_config = {"backend": "FLEX_ATTENTION"} - # pytest caches the fixture so we use weakref.proxy to # enable garbage collection llm = LLM( @@ -45,7 +38,6 @@ def llm(): gpu_memory_utilization=0.75, enforce_eager=True, seed=0, - attention_config=attention_config, ) yield weakref.proxy(llm) diff --git a/tests/entrypoints/pooling/basic/test_truncation.py b/tests/entrypoints/pooling/basic/test_truncation.py index 54cf083490..fa6094074b 100644 --- a/tests/entrypoints/pooling/basic/test_truncation.py +++ b/tests/entrypoints/pooling/basic/test_truncation.py @@ -7,7 +7,6 @@ import pytest import pytest_asyncio from tests.utils import RemoteOpenAIServer -from vllm.platforms import current_platform MODEL_NAME = "sentence-transformers/all-MiniLM-L12-v2" max_model_len = 128 @@ -39,10 +38,6 @@ def server(): str(max_model_len), ] - # ROCm: Use Flex Attention to support encoder-only self-attention. - if current_platform.is_rocm(): - args.extend(["--attention-backend", "FLEX_ATTENTION"]) - with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/embed/conftest.py b/tests/entrypoints/pooling/embed/conftest.py deleted file mode 100644 index 002b858740..0000000000 --- a/tests/entrypoints/pooling/embed/conftest.py +++ /dev/null @@ -1,28 +0,0 @@ -# SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project -"""Pytest configuration for vLLM pooling embed tests.""" - -import warnings - -import torch - -from vllm.platforms import current_platform - - -def pytest_collection_modifyitems(config, items): - """Configure ROCm-specific settings based on collected tests.""" - if not current_platform.is_rocm(): - return - - # Disable Flash/MemEfficient SDP on ROCm to avoid HF Transformers - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - # TODO: Remove once ROCm SDP accuracy issues are resolved on HuggingFace - torch.backends.cuda.enable_flash_sdp(False) - torch.backends.cuda.enable_mem_efficient_sdp(False) - torch.backends.cuda.enable_math_sdp(True) - warnings.warn( - "ROCm: Disabled flash_sdp and mem_efficient_sdp, enabled math_sdp " - "to avoid HuggingFace Transformers accuracy issues", - UserWarning, - stacklevel=1, - ) diff --git a/tests/entrypoints/pooling/embed/test_correctness_mteb.py b/tests/entrypoints/pooling/embed/test_correctness_mteb.py index 4c8d9f0d82..7018ae4b82 100644 --- a/tests/entrypoints/pooling/embed/test_correctness_mteb.py +++ b/tests/entrypoints/pooling/embed/test_correctness_mteb.py @@ -11,7 +11,6 @@ from tests.models.language.pooling_mteb_test.mteb_embed_utils import ( run_mteb_embed_task, ) from tests.utils import RemoteOpenAIServer -from vllm.platforms import current_platform os.environ["VLLM_LOGGING_LEVEL"] = "WARNING" @@ -23,10 +22,6 @@ MAIN_SCORE = 0.7422994752439667 def server(): args = ["--runner", "pooling", "--enforce-eager", "--disable-uvicorn-access-log"] - # ROCm: Use Flex Attention to support encoder-only self-attention. - if current_platform.is_rocm(): - args.extend(["--attention-backend", "FLEX_ATTENTION"]) - with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/embed/test_offline.py b/tests/entrypoints/pooling/embed/test_offline.py index c19e13075e..6dd2955ead 100644 --- a/tests/entrypoints/pooling/embed/test_offline.py +++ b/tests/entrypoints/pooling/embed/test_offline.py @@ -8,7 +8,6 @@ import torch.nn.functional as F from vllm import LLM, EmbeddingRequestOutput, PoolingParams from vllm.distributed import cleanup_dist_env_and_memory -from vllm.platforms import current_platform from vllm.tasks import PoolingTask MODEL_NAME = "intfloat/multilingual-e5-small" @@ -20,12 +19,6 @@ embedding_size = 384 @pytest.fixture(scope="module") def llm(): - # ROCm: Use FLEX_ATTENTION backend as it's the only attention backend - # that supports encoder-only models on ROCm. - attention_config = None - if current_platform.is_rocm(): - attention_config = {"backend": "FLEX_ATTENTION"} - # pytest caches the fixture so we use weakref.proxy to # enable garbage collection llm = LLM( @@ -35,7 +28,6 @@ def llm(): gpu_memory_utilization=0.75, enforce_eager=True, seed=0, - attention_config=attention_config, ) assert embedding_size == llm.model_config.embedding_size diff --git a/tests/entrypoints/pooling/embed/test_online.py b/tests/entrypoints/pooling/embed/test_online.py index 96555ee363..be600cdc88 100644 --- a/tests/entrypoints/pooling/embed/test_online.py +++ b/tests/entrypoints/pooling/embed/test_online.py @@ -50,14 +50,6 @@ input_tokens = [ 2, ] -if current_platform.is_rocm(): - # Disable Flash/MemEfficient SDP on ROCm to avoid HF Transformers - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - # TODO: Remove once ROCm SDP accuracy issues are resolved on HuggingFace - torch.backends.cuda.enable_flash_sdp(False) - torch.backends.cuda.enable_mem_efficient_sdp(False) - torch.backends.cuda.enable_math_sdp(True) - # On ROCm, floating-point reductions in attention and GEMM kernels are # non-associative and sensitive to batch geometry. Force LLM instances # into an identical, deterministic execution mode: @@ -507,11 +499,17 @@ async def test_base64_embedding(hf_model, client: openai.AsyncOpenAI, model_name "The best thing about vLLM is that it supports many different models", ] + def check_outputs(outputs: list[list[float]]) -> None: + # The embeddings endpoint executes list elements as separate scheduler + # requests, so use the same per-request geometry for the HF reference. + for input_text, output in zip(input_texts, outputs, strict=True): + run_embedding_correctness_test(hf_model, [input_text], [output]) + responses_float = await client.embeddings.create( input=input_texts, model=model_name, encoding_format="float" ) float_data = [d.embedding for d in responses_float.data] - run_embedding_correctness_test(hf_model, input_texts, float_data) + check_outputs(float_data) responses_base64 = await client.embeddings.create( input=input_texts, model=model_name, encoding_format="base64" @@ -522,14 +520,14 @@ async def test_base64_embedding(hf_model, client: openai.AsyncOpenAI, model_name np.frombuffer(base64.b64decode(data.embedding), dtype="float32").tolist() ) - run_embedding_correctness_test(hf_model, input_texts, base64_data) + check_outputs(base64_data) # Default response is float32 decoded from base64 by OpenAI Client responses_default = await client.embeddings.create( input=input_texts, model=model_name ) default_data = [d.embedding for d in responses_default.data] - run_embedding_correctness_test(hf_model, input_texts, default_data) + check_outputs(default_data) @pytest.mark.asyncio diff --git a/tests/entrypoints/pooling/embed/test_online_dimensions.py b/tests/entrypoints/pooling/embed/test_online_dimensions.py index 638f521898..7abec99ec3 100644 --- a/tests/entrypoints/pooling/embed/test_online_dimensions.py +++ b/tests/entrypoints/pooling/embed/test_online_dimensions.py @@ -12,7 +12,6 @@ from tests.models.language.pooling.embed_utils import run_embedding_correctness_ from tests.models.utils import EmbedModelInfo from tests.utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer from vllm.entrypoints.pooling.embed.protocol import EmbeddingResponse -from vllm.platforms import current_platform MODELS = [ EmbedModelInfo("intfloat/multilingual-e5-small", is_matryoshka=False), @@ -57,10 +56,6 @@ def server(model_info, dtype: str): ["--trust_remote_code", "--hf_overrides", '{"matryoshka_dimensions":[256]}'] ) - # ROCm: Use Flex Attention to support encoder-only self-attention. - if current_platform.is_rocm(): - args.extend(["--attention-backend", "FLEX_ATTENTION"]) - with RemoteOpenAIServer(model_info.name, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/embed/test_online_long_text.py b/tests/entrypoints/pooling/embed/test_online_long_text.py index c1dc5bf1ed..60cacb17f8 100644 --- a/tests/entrypoints/pooling/embed/test_online_long_text.py +++ b/tests/entrypoints/pooling/embed/test_online_long_text.py @@ -16,7 +16,6 @@ import pytest_asyncio from tests.utils import RemoteOpenAIServer from vllm.entrypoints.pooling.embed.protocol import EmbeddingResponse -from vllm.platforms import current_platform def _generate_random_text(word_count: int) -> str: @@ -225,10 +224,6 @@ def server_with_chunked_processing(): "0.8", ] - # ROCm: Use Flex Attention to support encoder-only self-attention. - if current_platform.is_rocm(): - args.extend(["--attention-backend", "FLEX_ATTENTION"]) - with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/scoring/test_bi_encoder_offline.py b/tests/entrypoints/pooling/scoring/test_bi_encoder_offline.py index aba70293f2..e68fb61a99 100644 --- a/tests/entrypoints/pooling/scoring/test_bi_encoder_offline.py +++ b/tests/entrypoints/pooling/scoring/test_bi_encoder_offline.py @@ -8,7 +8,6 @@ import pytest from tests.entrypoints.pooling.scoring.util import EncoderScoringHfRunner from vllm import LLM from vllm.distributed import cleanup_dist_env_and_memory -from vllm.platforms import current_platform MODEL_NAME = "intfloat/multilingual-e5-small" PROMPT = "The chef prepared a delicious meal." @@ -29,12 +28,6 @@ DTYPE = "half" @pytest.fixture(scope="module") def llm(): - # ROCm: Use FLEX_ATTENTION backend as it's the only attention backend - # that supports encoder-only models on ROCm. - attention_config = None - if current_platform.is_rocm(): - attention_config = {"backend": "FLEX_ATTENTION"} - # pytest caches the fixture so we use weakref.proxy to # enable garbage collection llm = LLM( @@ -44,7 +37,6 @@ def llm(): gpu_memory_utilization=0.75, enforce_eager=True, seed=0, - attention_config=attention_config, ) yield weakref.proxy(llm) diff --git a/tests/entrypoints/pooling/scoring/test_bi_encoder_online.py b/tests/entrypoints/pooling/scoring/test_bi_encoder_online.py index 3925140566..dac7b6db6c 100644 --- a/tests/entrypoints/pooling/scoring/test_bi_encoder_online.py +++ b/tests/entrypoints/pooling/scoring/test_bi_encoder_online.py @@ -8,7 +8,6 @@ from tests.entrypoints.pooling.scoring.util import EncoderScoringHfRunner from tests.utils import RemoteOpenAIServer from vllm.entrypoints.pooling.pooling.protocol import PoolingResponse from vllm.entrypoints.pooling.scoring.protocol import RerankResponse, ScoreResponse -from vllm.platforms import current_platform MODEL_NAME = "BAAI/bge-base-en-v1.5" input_text = "This product was excellent and exceeded my expectations" @@ -30,10 +29,6 @@ TEXTS_2 = [ def server(): args = ["--enforce-eager", "--max-model-len", "100", "--dtype", DTYPE] - # ROCm: Use Flex Attention to support encoder-only self-attention. - if current_platform.is_rocm(): - args.extend(["--attention-backend", "FLEX_ATTENTION"]) - with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/scoring/test_cross_encoder_correctness_mteb.py b/tests/entrypoints/pooling/scoring/test_cross_encoder_correctness_mteb.py index 1ee45b4459..1e18eb60a3 100644 --- a/tests/entrypoints/pooling/scoring/test_cross_encoder_correctness_mteb.py +++ b/tests/entrypoints/pooling/scoring/test_cross_encoder_correctness_mteb.py @@ -13,7 +13,6 @@ from tests.models.language.pooling_mteb_test.mteb_score_utils import ( run_mteb_rerank, ) from tests.utils import RemoteOpenAIServer -from vllm.platforms import current_platform os.environ["VLLM_LOGGING_LEVEL"] = "WARNING" @@ -25,10 +24,6 @@ st_main_score = 0.33457 def server(): args = ["--runner", "pooling", "--enforce-eager", "--disable-uvicorn-access-log"] - # ROCm: Use Flex Attention to support encoder-only self-attention. - if current_platform.is_rocm(): - args.extend(["--attention-backend", "FLEX_ATTENTION"]) - with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/scoring/test_cross_encoder_offline.py b/tests/entrypoints/pooling/scoring/test_cross_encoder_offline.py index 56e83de3f7..3ad044987a 100644 --- a/tests/entrypoints/pooling/scoring/test_cross_encoder_offline.py +++ b/tests/entrypoints/pooling/scoring/test_cross_encoder_offline.py @@ -9,7 +9,6 @@ import torch from tests.models.utils import softmax from vllm import LLM, PoolingParams from vllm.distributed import cleanup_dist_env_and_memory -from vllm.platforms import current_platform MODEL_NAME = "tomaarsen/Qwen3-Reranker-0.6B-seq-cls" PROMPT = "The chef prepared a delicious meal." @@ -26,12 +25,6 @@ TEXTS_2 = [ @pytest.fixture(scope="module") def llm(): - # ROCm: Use FLEX_ATTENTION backend as it's the only attention backend - # that supports encoder-only models on ROCm. - attention_config = None - if current_platform.is_rocm(): - attention_config = {"backend": "FLEX_ATTENTION"} - # pytest caches the fixture so we use weakref.proxy to # enable garbage collection llm = LLM( @@ -41,7 +34,6 @@ def llm(): gpu_memory_utilization=0.75, enforce_eager=True, seed=0, - attention_config=attention_config, ) yield weakref.proxy(llm) diff --git a/tests/entrypoints/pooling/scoring/test_cross_encoder_online.py b/tests/entrypoints/pooling/scoring/test_cross_encoder_online.py index 54b4fe075b..401bf4ff22 100644 --- a/tests/entrypoints/pooling/scoring/test_cross_encoder_online.py +++ b/tests/entrypoints/pooling/scoring/test_cross_encoder_online.py @@ -9,7 +9,6 @@ import torch.nn.functional as F from tests.utils import RemoteOpenAIServer from vllm.entrypoints.pooling.pooling.protocol import PoolingResponse from vllm.entrypoints.pooling.scoring.protocol import RerankResponse, ScoreResponse -from vllm.platforms import current_platform MODEL_NAME = "BAAI/bge-reranker-base" DTYPE = "half" @@ -32,10 +31,6 @@ TEXTS_2 = [ def server(): args = ["--enforce-eager", "--max-model-len", "100", "--dtype", DTYPE] - # ROCm: Use Flex Attention to support encoder-only self-attention. - if current_platform.is_rocm(): - args.extend(["--attention-backend", "FLEX_ATTENTION"]) - with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py b/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py index aae69fe217..107de837fe 100644 --- a/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py +++ b/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py @@ -53,12 +53,6 @@ BACKEND_ABS_TOL: dict[str, float] = { "FLEX_ATTENTION": 0.006, } -# ROCm: disable skinny GEMM to avoid non-deterministic results from -# atomic reductions in wvSplitKrc kernel. -# See: https://github.com/vllm-project/vllm/pull/33493#issuecomment-3906083975 -ROCM_ENV_OVERRIDES = ( - {"VLLM_ROCM_USE_SKINNY_GEMM": "0"} if current_platform.is_rocm() else {} -) # ROCm: disable prefix caching and eliminate batch variance to reduce # test flakiness. ROCM_EXTRA_ARGS = ( @@ -132,7 +126,6 @@ def server(request): args += ["--attention-config", json.dumps({"backend": backend})] args += ROCM_EXTRA_ARGS - env = dict(ROCM_ENV_OVERRIDES) if backend != "ROCM_AITER_FA": env["VLLM_ROCM_USE_AITER"] = "0" diff --git a/tests/entrypoints/pooling/scoring/test_late_interaction_offline.py b/tests/entrypoints/pooling/scoring/test_late_interaction_offline.py index ea162d630f..9a4db3d3f8 100644 --- a/tests/entrypoints/pooling/scoring/test_late_interaction_offline.py +++ b/tests/entrypoints/pooling/scoring/test_late_interaction_offline.py @@ -7,7 +7,6 @@ import pytest from vllm import LLM from vllm.distributed import cleanup_dist_env_and_memory -from vllm.platforms import current_platform from .util import ColBERTScoringHfRunner @@ -32,12 +31,6 @@ DTYPE = "half" @pytest.fixture(scope="module") def llm(): - # ROCm: Use FLEX_ATTENTION backend as it's the only attention backend - # that supports encoder-only models on ROCm. - attention_config = None - if current_platform.is_rocm(): - attention_config = {"backend": "FLEX_ATTENTION"} - # pytest caches the fixture so we use weakref.proxy to # enable garbage collection llm = LLM( @@ -47,7 +40,6 @@ def llm(): gpu_memory_utilization=0.75, enforce_eager=True, seed=0, - attention_config=attention_config, ) yield weakref.proxy(llm) diff --git a/tests/entrypoints/pooling/scoring/test_late_interaction_offline_vision.py b/tests/entrypoints/pooling/scoring/test_late_interaction_offline_vision.py index f5bbb208b4..4812bd2ab3 100644 --- a/tests/entrypoints/pooling/scoring/test_late_interaction_offline_vision.py +++ b/tests/entrypoints/pooling/scoring/test_late_interaction_offline_vision.py @@ -7,7 +7,6 @@ import pytest from vllm import LLM from vllm.distributed import cleanup_dist_env_and_memory -from vllm.platforms import current_platform from .util import make_base64_image, make_image_mm_param @@ -16,12 +15,6 @@ MODEL_NAME = "vidore/colpali-v1.3-hf" @pytest.fixture(scope="module") def llm(): - # ROCm: Use FLEX_ATTENTION backend as it's the only attention backend - # that supports encoder-only models on ROCm. - attention_config = None - if current_platform.is_rocm(): - attention_config = {"backend": "FLEX_ATTENTION"} - # pytest caches the fixture so we use weakref.proxy to # enable garbage collection llm = LLM( @@ -31,7 +24,6 @@ def llm(): gpu_memory_utilization=0.75, enforce_eager=True, seed=0, - attention_config=attention_config, ) yield weakref.proxy(llm) diff --git a/tests/entrypoints/pooling/token_embed/test_offline.py b/tests/entrypoints/pooling/token_embed/test_offline.py index 5bf15052d1..7bc7be4866 100644 --- a/tests/entrypoints/pooling/token_embed/test_offline.py +++ b/tests/entrypoints/pooling/token_embed/test_offline.py @@ -7,7 +7,6 @@ import pytest from vllm import LLM, PoolingRequestOutput from vllm.config import PoolerConfig from vllm.distributed import cleanup_dist_env_and_memory -from vllm.platforms import current_platform from vllm.tasks import PoolingTask MODEL_NAME = "intfloat/multilingual-e5-small" @@ -19,12 +18,6 @@ embedding_size = 384 @pytest.fixture(scope="module") def llm(): - # ROCm: Use FLEX_ATTENTION backend as it's the only attention backend - # that supports encoder-only models on ROCm. - attention_config = None - if current_platform.is_rocm(): - attention_config = {"backend": "FLEX_ATTENTION"} - # pytest caches the fixture so we use weakref.proxy to # enable garbage collection llm = LLM( @@ -35,7 +28,6 @@ def llm(): gpu_memory_utilization=0.75, enforce_eager=True, seed=0, - attention_config=attention_config, ) assert embedding_size == llm.model_config.embedding_size diff --git a/tests/entrypoints/scale_out/token_in_token_out/test_serving_multimodal_tokens.py b/tests/entrypoints/scale_out/token_in_token_out/test_serving_multimodal_tokens.py index b4811666d4..60f5e9e420 100644 --- a/tests/entrypoints/scale_out/token_in_token_out/test_serving_multimodal_tokens.py +++ b/tests/entrypoints/scale_out/token_in_token_out/test_serving_multimodal_tokens.py @@ -9,8 +9,6 @@ processor in the pytest parent process to keep os.fork() in sibling tests (e.g. test_weight_transfer_llm.py) deadlock-free. """ -import os - import httpx import pytest import pytest_asyncio @@ -41,10 +39,7 @@ def server(): "--no-enable-prefix-caching", ] - envs = os.environ.copy() - envs["VLLM_ROCM_USE_SKINNY_GEMM"] = "0" - - with RemoteOpenAIServer(MODEL_NAME, args, env_dict=envs) as remote_server: + with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py b/tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py index 3425bdef63..4ed5923980 100644 --- a/tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py +++ b/tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py @@ -2,7 +2,6 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM project import json -import os import httpx import pytest @@ -79,11 +78,7 @@ def server(request): else [str(extra_args)] ) - envs = os.environ.copy() - # See: https://github.com/vllm-project/vllm/pull/33493#issuecomment-3888060787 - envs["VLLM_ROCM_USE_SKINNY_GEMM"] = "0" - - with RemoteOpenAIServer(MODEL_NAME, args, env_dict=envs) as remote_server: + with RemoteOpenAIServer(MODEL_NAME, args) as remote_server: yield remote_server diff --git a/tests/entrypoints/speech_to_text/conftest.py b/tests/entrypoints/speech_to_text/conftest.py index 098a9a7232..94505bd054 100644 --- a/tests/entrypoints/speech_to_text/conftest.py +++ b/tests/entrypoints/speech_to_text/conftest.py @@ -16,19 +16,6 @@ def add_attention_backend(server_args, attention_config): server_args.extend(["--attention-backend", attention_config["backend"]]) -@pytest.fixture(scope="module") -def rocm_aiter_fa_attention(): - """Return attention config for transcription/translation tests on ROCm. - - On ROCm, audio tests require ROCM_AITER_FA attention backend. - """ - from vllm.platforms import current_platform - - if current_platform.is_rocm(): - return {"backend": "ROCM_AITER_FA"} - return None - - @pytest.fixture def mary_had_lamb(): path = AudioAsset("mary_had_lamb").get_local_path() diff --git a/tests/entrypoints/speech_to_text/realtime/test_realtime_validation.py b/tests/entrypoints/speech_to_text/realtime/test_realtime_validation.py index 675922f588..7fb7a5522d 100644 --- a/tests/entrypoints/speech_to_text/realtime/test_realtime_validation.py +++ b/tests/entrypoints/speech_to_text/realtime/test_realtime_validation.py @@ -10,15 +10,13 @@ import pybase64 as base64 import pytest import websockets -from tests.entrypoints.speech_to_text.conftest import add_attention_backend -from tests.utils import ROCM_ENV_OVERRIDES, ROCM_EXTRA_ARGS, RemoteOpenAIServer +from tests.utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer from vllm.assets.audio import AudioAsset from vllm.multimodal.media.audio import load_audio # Increase engine iteration timeout for ROCm where first-use JIT compilation # can exceed the default 60s, causing a silent deadlock in feed_tokens. REALTIME_ENV_OVERRIDES = { - **ROCM_ENV_OVERRIDES, "VLLM_ENGINE_ITERATION_TIMEOUT_S": "600", } @@ -72,17 +70,13 @@ def mary_had_lamb_audio_chunks() -> list[str]: @pytest.mark.asyncio @pytest.mark.parametrize("model_name", [MODEL_NAME]) -async def test_multi_chunk_streaming( - model_name, mary_had_lamb_audio_chunks, rocm_aiter_fa_attention -): +async def test_multi_chunk_streaming(model_name, mary_had_lamb_audio_chunks): """Test streaming multiple audio chunks before committing.""" server_args = ["--enforce-eager", "--max-model-len", "2048"] if model_name.startswith("mistralai"): server_args += MISTRAL_FORMAT_ARGS - add_attention_backend(server_args, rocm_aiter_fa_attention) - with RemoteOpenAIServer( model_name, server_args, env_dict=REALTIME_ENV_OVERRIDES ) as remote_server: @@ -174,7 +168,7 @@ async def test_multi_chunk_streaming( @pytest.mark.asyncio @pytest.mark.parametrize("model_name", [MODEL_NAME]) async def test_empty_commit_does_not_crash_engine( - model_name, mary_had_lamb_audio_chunks, rocm_aiter_fa_attention + model_name, mary_had_lamb_audio_chunks ): """Test that committing without audio does not crash the engine. @@ -189,8 +183,6 @@ async def test_empty_commit_does_not_crash_engine( if model_name.startswith("mistralai"): server_args += MISTRAL_FORMAT_ARGS - add_attention_backend(server_args, rocm_aiter_fa_attention) - with RemoteOpenAIServer( model_name, server_args, env_dict=REALTIME_ENV_OVERRIDES ) as remote_server: @@ -273,17 +265,13 @@ async def test_empty_commit_does_not_crash_engine( @pytest.mark.asyncio @pytest.mark.parametrize("model_name", [MODEL_NAME]) -async def test_session_update_invalid_model_returns_error( - model_name, rocm_aiter_fa_attention -): +async def test_session_update_invalid_model_returns_error(model_name): """Test that session.update with an invalid model returns an error.""" server_args = ["--enforce-eager", "--max-model-len", "2048"] if model_name.startswith("mistralai"): server_args += MISTRAL_FORMAT_ARGS - add_attention_backend(server_args, rocm_aiter_fa_attention) - with RemoteOpenAIServer( model_name, server_args, env_dict=REALTIME_ENV_OVERRIDES ) as remote_server: @@ -305,9 +293,7 @@ async def test_session_update_invalid_model_returns_error( @pytest.mark.asyncio @pytest.mark.parametrize("model_name", [MODEL_NAME]) -async def test_commit_without_session_update_returns_error( - model_name, rocm_aiter_fa_attention -): +async def test_commit_without_session_update_returns_error(model_name): """Test that committing before validating the model returns an error and does not fall through to processing.""" server_args = ["--enforce-eager", "--max-model-len", "2048"] @@ -315,8 +301,6 @@ async def test_commit_without_session_update_returns_error( if model_name.startswith("mistralai"): server_args += MISTRAL_FORMAT_ARGS - add_attention_backend(server_args, rocm_aiter_fa_attention) - with RemoteOpenAIServer( model_name, server_args, env_dict=REALTIME_ENV_OVERRIDES ) as remote_server: diff --git a/tests/entrypoints/speech_to_text/transcription/test_transcription_validation.py b/tests/entrypoints/speech_to_text/transcription/test_transcription_validation.py index 5ea218406b..2237a091ce 100644 --- a/tests/entrypoints/speech_to_text/transcription/test_transcription_validation.py +++ b/tests/entrypoints/speech_to_text/transcription/test_transcription_validation.py @@ -6,8 +6,7 @@ import json import pytest -from tests.entrypoints.speech_to_text.conftest import add_attention_backend -from tests.utils import ROCM_ENV_OVERRIDES, ROCM_EXTRA_ARGS, RemoteOpenAIServer +from tests.utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer MISTRAL_FORMAT_ARGS = [ "--tokenizer_mode", @@ -66,18 +65,14 @@ async def transcribe_and_check( @pytest.mark.parametrize( "model_name", ["mistralai/Voxtral-Mini-3B-2507", "Qwen/Qwen3-ASR-0.6B"] ) -async def test_basic_audio(mary_had_lamb, model_name, rocm_aiter_fa_attention): +async def test_basic_audio(mary_had_lamb, model_name): server_args = ["--enforce-eager", *ROCM_EXTRA_ARGS] if model_name.startswith("mistralai"): server_args += MISTRAL_FORMAT_ARGS - add_attention_backend(server_args, rocm_aiter_fa_attention) - # Based on https://github.com/openai/openai-cookbook/blob/main/examples/Whisper_prompting_guide.ipynb. - with RemoteOpenAIServer( - model_name, server_args, env_dict=ROCM_ENV_OVERRIDES - ) as remote_server: + with RemoteOpenAIServer(model_name, server_args) as remote_server: client = remote_server.get_async_client() await transcribe_and_check( client, @@ -90,7 +85,7 @@ async def test_basic_audio(mary_had_lamb, model_name, rocm_aiter_fa_attention): @pytest.mark.asyncio -async def test_basic_audio_with_lora(mary_had_lamb, rocm_aiter_fa_attention): +async def test_basic_audio_with_lora(mary_had_lamb): """Ensure STT (transcribe) requests can pass LoRA through to generate.""" # ROCm SPECIFIC CONFIGURATION: # To ensure the test passes on ROCm, we modify the max model length to 512. @@ -112,12 +107,8 @@ async def test_basic_audio_with_lora(mary_had_lamb, rocm_aiter_fa_attention): "1", ] - add_attention_backend(server_args, rocm_aiter_fa_attention) - # Based on https://github.com/openai/openai-cookbook/blob/main/examples/Whisper_prompting_guide.ipynb. - with RemoteOpenAIServer( - model_name, server_args, env_dict=ROCM_ENV_OVERRIDES - ) as remote_server: + with RemoteOpenAIServer(model_name, server_args) as remote_server: client = remote_server.get_async_client() await transcribe_and_check( client, @@ -133,18 +124,15 @@ async def test_basic_audio_with_lora(mary_had_lamb, rocm_aiter_fa_attention): @pytest.mark.parametrize( "model_name", ["google/gemma-3n-E2B-it", "Qwen/Qwen3-ASR-0.6B"] ) -async def test_basic_audio_foscolo(foscolo, rocm_aiter_fa_attention, model_name): +async def test_basic_audio_foscolo(foscolo, model_name): # Gemma accuracy on some of the audio samples we use is particularly bad, # hence we use a different one here. WER is evaluated separately. server_args = ["--enforce-eager", *ROCM_EXTRA_ARGS] - add_attention_backend(server_args, rocm_aiter_fa_attention) - with RemoteOpenAIServer( model_name, server_args, max_wait_seconds=480, - env_dict=ROCM_ENV_OVERRIDES, ) as remote_server: client = remote_server.get_async_client() await transcribe_and_check( diff --git a/tests/entrypoints/speech_to_text/translation/test_translation_validation.py b/tests/entrypoints/speech_to_text/translation/test_translation_validation.py index 25f26404d0..5cc1a2eff1 100644 --- a/tests/entrypoints/speech_to_text/translation/test_translation_validation.py +++ b/tests/entrypoints/speech_to_text/translation/test_translation_validation.py @@ -28,27 +28,33 @@ def _get_rocm_attention_config(model_name): Whisper uses cross-attention (ENCODER_DECODER) which ROCM_AITER_FA does not support. For Whisper we use ROCM_AITER_UNIFIED_ATTN (or TRITON_ATTN - as fallback); other models can use ROCM_AITER_FA. + as fallback). Other models use ROCM_AITER_FA only on supported CDNA3+ + hardware and otherwise use the platform default. """ from vllm.platforms import current_platform if not current_platform.is_rocm(): return None - if "whisper" in model_name.lower(): - try: - from vllm.platforms.rocm import get_cdna_version + try: + from vllm.platforms.rocm import get_cdna_version - if get_cdna_version() > 2: - return {"backend": "ROCM_AITER_UNIFIED_ATTN"} - except ImportError: - logger.warning( - "Could not check cdna version from rocm platform, " - "falling back to TRITON_ATTN for Whisper." - ) + cdna_version = get_cdna_version() + except ImportError: + logger.warning( + "Could not check cdna version from rocm platform; " + "using a compatible fallback attention backend." + ) + cdna_version = -1 + + if "whisper" in model_name.lower(): + if cdna_version > 2: + return {"backend": "ROCM_AITER_UNIFIED_ATTN"} return {"backend": "TRITON_ATTN"} - return {"backend": "ROCM_AITER_FA"} + if cdna_version > 2: + return {"backend": "ROCM_AITER_FA"} + return None def _get_server_args(attention_config): diff --git a/tests/kernels/attention/test_prefix_prefill.py b/tests/kernels/attention/test_prefix_prefill.py index e63d2c0adf..e2e9c4e8c2 100644 --- a/tests/kernels/attention/test_prefix_prefill.py +++ b/tests/kernels/attention/test_prefix_prefill.py @@ -5,12 +5,10 @@ import math import random import time from collections.abc import Callable -from contextlib import nullcontext import pytest import torch import torch.nn.functional as F -from torch.nn.attention import SDPBackend, sdpa_kernel from vllm.platforms import current_platform from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed @@ -895,21 +893,15 @@ def test_contexted_kv_attention_alibi( query_len, seq_len, alibi_slopes, device, dtype ) - # Compute attention. On ROCm we force use of the Math SDPA backend rather than - # the Flash or Mem-Efficient backends for increased numerical accuracy - if current_platform.is_rocm(): - sdpa_context = sdpa_kernel(SDPBackend.MATH) - else: - sdpa_context = nullcontext() - with sdpa_context: - out = F.scaled_dot_product_attention( - q_sdpa, - k_sdpa, - v_sdpa, - attn_mask=alibi_mask, - dropout_p=0.0, - scale=scale, - ) + # Compute attention + out = F.scaled_dot_product_attention( + q_sdpa, + k_sdpa, + v_sdpa, + attn_mask=alibi_mask, + dropout_p=0.0, + scale=scale, + ) # Reshape output back to [query_len, num_heads, head_size] out = out.view(num_heads, query_len, head_size).permute(1, 0, 2) diff --git a/tests/kernels/quantization/test_rocm_skinny_gemms.py b/tests/kernels/quantization/test_rocm_skinny_gemms.py index e137231dc9..0c1cbe10bf 100644 --- a/tests/kernels/quantization/test_rocm_skinny_gemms.py +++ b/tests/kernels/quantization/test_rocm_skinny_gemms.py @@ -140,7 +140,7 @@ def test_rocm_wvsplitkrc_kernel(xnorm, n, k, m, dtype, seed, padded_a, bias_mode GrpsShrB = min(N_p2 // 16, 4) # Given the above, how many CUs would we need? CuNeeded = rndup_cus * GrpsShrB - # candidate for atomic reduce count splitk? + # Deterministic reduction stores one float workspace value per K shard. fits_wvsplitkrc = (N_p2 * m * ((k + 512 - 1) // 512)) <= 128 * 1024 * 12 fits_wvsplitkrc &= CuNeeded <= cu_count diff --git a/tests/lora/test_gptoss_tp.py b/tests/lora/test_gptoss_tp.py index 0e7fe60b36..70b94ea2eb 100644 --- a/tests/lora/test_gptoss_tp.py +++ b/tests/lora/test_gptoss_tp.py @@ -50,6 +50,20 @@ def reformat(text: str) -> str: return text +def _enable_deterministic_lora_shrink(monkeypatch: pytest.MonkeyPatch) -> None: + if not current_platform.is_rocm(): + return + + # These tests assert exact greedy outputs. Force the Triton LoRA shrink + # kernel to use SPLIT_K=1 so it stores the complete reduction directly + # instead of accumulating split-K partial results with atomic_add. + monkeypatch.setenv("VLLM_BATCH_INVARIANT", "1") + # The kernel configuration reads VLLM_BATCH_INVARIANT at import time. + # Spawn the engine process so it observes this setting even if the LoRA + # Triton utilities were already imported during test collection. + monkeypatch.setenv("VLLM_WORKER_MULTIPROC_METHOD", "spawn") + + def generate_and_test(llm: vllm.LLM, lora_path: str, lora_id: int) -> None: prompts = [ PROMPT_TEMPLATE.format( @@ -106,7 +120,10 @@ def test_gpt_oss_lora( gptoss20b_lora_files, mxfp4_use_marlin, specialize_active_lora, + monkeypatch: pytest.MonkeyPatch, ): + _enable_deterministic_lora_shrink(monkeypatch) + llm = vllm.LLM( MODEL_PATH, max_model_len=1024, @@ -146,7 +163,10 @@ def test_gpt_oss_lora_tp2( gptoss20b_lora_files, fully_sharded_loras, mxfp4_use_marlin, + monkeypatch: pytest.MonkeyPatch, ): + _enable_deterministic_lora_shrink(monkeypatch) + llm = vllm.LLM( MODEL_PATH, max_model_len=1024, diff --git a/tests/lora/test_moe_lora_align_sum.py b/tests/lora/test_moe_lora_align_sum.py index 77a44970ba..80ba0a56ca 100644 --- a/tests/lora/test_moe_lora_align_sum.py +++ b/tests/lora/test_moe_lora_align_sum.py @@ -40,6 +40,18 @@ def sample_data(num_experts, max_loras, num_tokens, topk_num): @pytest.mark.parametrize("block_size", [16]) def test_moe_lora_align_block_size( num_tokens, topk_num, num_experts, max_loras, block_size +): + _run_moe_lora_align_block_size_case( + num_tokens, topk_num, num_experts, max_loras, block_size + ) + + +def test_moe_lora_align_block_size_large_path_counts(): + _run_moe_lora_align_block_size_case(100, 12, 64, 4, 16) + + +def _run_moe_lora_align_block_size_case( + num_tokens, topk_num, num_experts, max_loras, block_size ): # sample data random.seed(1) @@ -94,8 +106,17 @@ def test_moe_lora_align_block_size( # verify values expert_ids = expert_ids.view(max_loras, -1) sorted_token_ids = sorted_token_ids.view(max_loras, -1, block_size) + topk_ids_cpu = topk_ids.cpu() + token_lora_mapping_cpu = token_lora_mapping.cpu() for lora_idx in range(max_loras): + expert_counts = torch.bincount( + topk_ids_cpu[token_lora_mapping_cpu == lora_idx].flatten(), + minlength=num_experts, + ) + expected_post_pad = round_up(expert_counts, block_size).sum().item() + assert num_tokens_post_pad[lora_idx].item() == expected_post_pad + for token_idx in range(sorted_token_ids.size(1)): block = sorted_token_ids[lora_idx][token_idx] indices = block[block != topk_ids.numel()] diff --git a/tests/model_executor/layers/test_rocm_unquantized_gemm.py b/tests/model_executor/layers/test_rocm_unquantized_gemm.py index f20be7cfb6..66b4af7727 100644 --- a/tests/model_executor/layers/test_rocm_unquantized_gemm.py +++ b/tests/model_executor/layers/test_rocm_unquantized_gemm.py @@ -42,6 +42,113 @@ def test_rocm_unquantized_gemm_gfx1x_wvsplitk_path(monkeypatch): assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3) +def test_rocm_unquantized_gemm_makes_skinny_activation_contiguous(monkeypatch): + x = torch.randn(64, 4, dtype=torch.float16).t() + weight = torch.randn(128, 64, dtype=torch.float16) + assert x.shape == (4, 64) + assert x.stride() == (1, 4) + + monkeypatch.setattr(utils, "use_aiter_triton_gemm", lambda *args: False) + monkeypatch.setattr(utils.envs, "VLLM_ROCM_USE_SKINNY_GEMM", True) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx1x", lambda: True) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx9", lambda: False) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx950", lambda: False) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx1250", lambda: False) + monkeypatch.setattr(utils, "num_compute_units", lambda: 120) + + wvsplitk_mock = MagicMock(side_effect=lambda w, x_view, _, __: x_view @ w.t()) + monkeypatch.setattr(utils.ops, "wvSplitK", wvsplitk_mock) + + out = utils.rocm_unquantized_gemm_impl(x, weight, None) + ref = torch.nn.functional.linear(x, weight, None) + + wvsplitk_mock.assert_called_once() + x_view = wvsplitk_mock.call_args.args[1] + assert x_view.is_contiguous() + assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3) + + +def test_rocm_unquantized_gemm_makes_llmm1_activation_contiguous(monkeypatch): + x = torch.randn(1, 128, dtype=torch.float16)[:, ::2] + weight = torch.randn(4, 64, dtype=torch.float16) + assert x.shape == (1, 64) + assert x.stride() == (128, 2) + + monkeypatch.setattr(utils, "use_aiter_triton_gemm", lambda *args: False) + monkeypatch.setattr(utils.envs, "VLLM_ROCM_USE_SKINNY_GEMM", True) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx1x", lambda: True) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx9", lambda: False) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx950", lambda: False) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx1250", lambda: False) + monkeypatch.setattr(utils, "num_compute_units", lambda: 120) + + llmm1_mock = MagicMock(side_effect=lambda w, x_view, _: x_view @ w.t()) + monkeypatch.setattr(utils.ops, "LLMM1", llmm1_mock) + + out = utils.rocm_unquantized_gemm_impl(x, weight, None) + ref = torch.nn.functional.linear(x, weight, None) + + llmm1_mock.assert_called_once() + x_view = llmm1_mock.call_args.args[1] + assert x_view.is_contiguous() + assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3) + + +@pytest.mark.parametrize("noncontiguous_operand", ["weight", "bias"]) +def test_rocm_unquantized_gemm_rejects_unsupported_skinny_layouts( + monkeypatch, noncontiguous_operand +): + x = torch.randn(4, 64, dtype=torch.float16) + weight = torch.randn(128, 64, dtype=torch.float16) + bias = torch.randn(128, dtype=torch.float16) + if noncontiguous_operand == "weight": + weight = torch.randn(64, 128, dtype=torch.float16).t() + assert not weight.is_contiguous() + else: + bias = torch.randn(256, dtype=torch.float16)[::2] + assert not bias.is_contiguous() + + monkeypatch.setattr(utils, "use_aiter_triton_gemm", lambda *args: False) + monkeypatch.setattr(utils.rocm_aiter_ops, "is_tgemm_enabled", lambda: False) + monkeypatch.setattr(utils.envs, "VLLM_ROCM_USE_SKINNY_GEMM", True) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx1x", lambda: True) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx9", lambda: False) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx950", lambda: False) + monkeypatch.setattr("vllm.platforms.rocm.on_gfx1250", lambda: False) + monkeypatch.setattr(utils, "num_compute_units", lambda: 120) + + wvsplitk_mock = MagicMock() + monkeypatch.setattr(utils.ops, "wvSplitK", wvsplitk_mock) + llmm1_mock = MagicMock() + monkeypatch.setattr(utils.ops, "LLMM1", llmm1_mock) + + out = utils.rocm_unquantized_gemm_impl(x, weight, bias) + ref = torch.nn.functional.linear(x, weight, bias) + + wvsplitk_mock.assert_not_called() + llmm1_mock.assert_not_called() + assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3) + + +@pytest.mark.skipif(not current_platform.is_rocm(), reason="ROCm-only kernel test") +@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16]) +def test_rocm_unquantized_gemm_noncontiguous_activation_real_kernel(monkeypatch, dtype): + x = torch.randn(64, 4, device="cuda", dtype=dtype).t() + weight = torch.randn(128, 64, device="cuda", dtype=dtype) + assert x.stride() == (1, 4) + + monkeypatch.setattr(utils.envs, "VLLM_ROCM_USE_SKINNY_GEMM", True) + original_wvsplitk = utils.ops.wvSplitK + wvsplitk_mock = MagicMock(side_effect=original_wvsplitk) + monkeypatch.setattr(utils.ops, "wvSplitK", wvsplitk_mock) + + out = utils.rocm_unquantized_gemm_impl(x, weight, None) + ref = torch.nn.functional.linear(x, weight, None) + + wvsplitk_mock.assert_called_once() + torch.testing.assert_close(out, ref, atol=1e-2, rtol=1e-2) + + def test_rocm_unquantized_gemm_gfx1x_n_gt_5_falls_back(monkeypatch): # wvSplitK skinny GEMM handles n in [1, 5] (see PR #40687); n > 5 must # fall back to torch.nn.functional.linear. @@ -70,8 +177,9 @@ def test_rocm_unquantized_gemm_gfx1x_n_gt_5_falls_back(monkeypatch): def test_rocm_unquantized_gemm_gfx950_wvsplitkrc_path(monkeypatch): - x = torch.randn(16, 1024, dtype=torch.float16) + x = torch.randn(1024, 16, dtype=torch.float16).t() weight = torch.randn(256, 1024, dtype=torch.float16) + assert x.stride() == (1, 16) monkeypatch.setattr(utils, "use_aiter_triton_gemm", lambda *args: False) monkeypatch.setattr(utils.envs, "VLLM_ROCM_USE_SKINNY_GEMM", True) @@ -91,4 +199,6 @@ def test_rocm_unquantized_gemm_gfx950_wvsplitkrc_path(monkeypatch): wvsplitkrc_mock.assert_called_once() wvsplitk_mock.assert_not_called() + x_view = wvsplitkrc_mock.call_args.args[0] + assert x_view.is_contiguous() assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3) diff --git a/tests/models/language/generation/conftest.py b/tests/models/language/generation/conftest.py deleted file mode 100644 index aeb13bde46..0000000000 --- a/tests/models/language/generation/conftest.py +++ /dev/null @@ -1,46 +0,0 @@ -# SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project -"""Pytest configuration for vLLM language generation tests.""" - -import os -import warnings - -import torch - -from vllm.platforms import current_platform - - -def pytest_configure(config): - """Early ROCm configuration that must happen before test collection.""" - if not current_platform.is_rocm(): - return - - # Disable skinny GEMM on ROCm to avoid non-deterministic results - # from atomic reductions in wvSplitKrc kernel. - # See: https://github.com/vllm-project/vllm/pull/33493#issuecomment-3906083975 - os.environ["VLLM_ROCM_USE_SKINNY_GEMM"] = "0" - warnings.warn( - "ROCm: Set VLLM_ROCM_USE_SKINNY_GEMM=0 to avoid non-deterministic " - "results from skinny GEMM atomic reductions", - UserWarning, - stacklevel=1, - ) - - -def pytest_sessionstart(session): - """Configure ROCm-specific settings before test session starts.""" - if not current_platform.is_rocm(): - return - - # Disable Flash/MemEfficient SDP on ROCm to avoid HF Transformers - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - # TODO: Remove once ROCm SDP accuracy issues are resolved on HuggingFace - torch.backends.cuda.enable_flash_sdp(False) - torch.backends.cuda.enable_mem_efficient_sdp(False) - torch.backends.cuda.enable_math_sdp(True) - warnings.warn( - "ROCm: Disabled flash_sdp and mem_efficient_sdp, enabled math_sdp " - "to avoid HuggingFace Transformers accuracy issues", - UserWarning, - stacklevel=1, - ) diff --git a/tests/models/language/generation/test_common.py b/tests/models/language/generation/test_common.py index db6b74ee61..33ff4e7ccb 100644 --- a/tests/models/language/generation/test_common.py +++ b/tests/models/language/generation/test_common.py @@ -126,6 +126,12 @@ def test_models( model_info.check_available_online(on_fail="skip") model_info.check_transformers_version(on_fail="skip") + if current_platform.is_rocm() and model == "TitanML/tiny-mixtral": + # Its single-token router selects LLMM1, whose low-precision + # accumulation can change the top-2 experts. Keep the optimized kernel + # enabled generally, but use the reference GEMM for this accuracy test. + monkeypatch.setenv("VLLM_ROCM_USE_SKINNY_GEMM", "0") + if use_rocm_aiter and (model in AITER_MODEL_LIST): monkeypatch.setenv("VLLM_ROCM_USE_AITER", "1") elif use_rocm_aiter and model not in AITER_MODEL_LIST: diff --git a/tests/models/language/pooling/conftest.py b/tests/models/language/pooling/conftest.py index 6348d49cd7..0571f1cd6d 100644 --- a/tests/models/language/pooling/conftest.py +++ b/tests/models/language/pooling/conftest.py @@ -2,8 +2,6 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM project """Pytest configuration for vLLM language generation tests.""" -import warnings - import torch from vllm.platforms import current_platform @@ -11,19 +9,5 @@ from vllm.platforms import current_platform def pytest_sessionstart(session): """Configure ROCm-specific settings before test session starts.""" - if not current_platform.is_rocm(): - return - - # Disable Flash/MemEfficient SDP on ROCm to avoid HF Transformers - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - # TODO: Remove once ROCm SDP accuracy issues are resolved on HuggingFace - torch.backends.cuda.enable_flash_sdp(False) - torch.backends.cuda.enable_mem_efficient_sdp(False) - torch.backends.cuda.enable_math_sdp(True) - torch.set_float32_matmul_precision("high") - warnings.warn( - "ROCm: Disabled flash_sdp and mem_efficient_sdp, enabled math_sdp " - "to avoid HuggingFace Transformers accuracy issues", - UserWarning, - stacklevel=1, - ) + if current_platform.is_rocm(): + torch.set_float32_matmul_precision("high") diff --git a/tests/models/language/pooling/test_gritlm.py b/tests/models/language/pooling/test_gritlm.py index 7b6c176fd0..64265d6a0c 100644 --- a/tests/models/language/pooling/test_gritlm.py +++ b/tests/models/language/pooling/test_gritlm.py @@ -7,7 +7,7 @@ from scipy.spatial.distance import cosine from vllm import LLM, SamplingParams from vllm.config import ModelConfig -from ....utils import ROCM_ENV_OVERRIDES, ROCM_EXTRA_ARGS, RemoteOpenAIServer +from ....utils import ROCM_EXTRA_ARGS, RemoteOpenAIServer from .embed_utils import run_client_embeddings MODEL_NAME = "parasail-ai/GritLM-7B-vllm" @@ -134,7 +134,7 @@ async def test_gritlm_api_server_embedding(): *ROCM_EXTRA_ARGS, ] - with RemoteOpenAIServer(MODEL_NAME, args, env_dict=ROCM_ENV_OVERRIDES) as server: + with RemoteOpenAIServer(MODEL_NAME, args) as server: client_embedding = server.get_async_client() d_rep = await run_client_embeddings( diff --git a/tests/models/language/pooling/test_token_classification.py b/tests/models/language/pooling/test_token_classification.py index 5b18bdf1a5..d8d813f404 100644 --- a/tests/models/language/pooling/test_token_classification.py +++ b/tests/models/language/pooling/test_token_classification.py @@ -48,17 +48,10 @@ def test_bert_model_runner_v2( vllm_model.token_classify(prompts) for prompts in prompt_batches ] - # Use eager attention on ROCm to avoid HF Transformers flash attention - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - hf_model_kwargs = {} - if current_platform.is_rocm(): - hf_model_kwargs["attn_implementation"] = "eager" - with hf_runner( model, dtype=dtype, auto_cls=AutoModelForTokenClassification, - model_kwargs=hf_model_kwargs, ) as hf_model: tokenizer = hf_model.tokenizer hf_outputs = [] @@ -100,17 +93,10 @@ def test_modernbert_models( with vllm_runner(model, max_model_len=None, dtype=dtype) as vllm_model: vllm_outputs = vllm_model.token_classify(example_prompts) - # Use eager attention on ROCm to avoid HF Transformers flash attention - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - hf_model_kwargs = {} - if current_platform.is_rocm(): - hf_model_kwargs["attn_implementation"] = "eager" - with hf_runner( model, dtype=dtype, auto_cls=AutoModelForTokenClassification, - model_kwargs=hf_model_kwargs, ) as hf_model: tokenizer = hf_model.tokenizer hf_outputs = [] @@ -140,17 +126,10 @@ def test_xlm_roberta_models( with vllm_runner(model, max_model_len=None, dtype=dtype) as vllm_model: vllm_outputs = vllm_model.token_classify(example_prompts) - # Use eager attention on ROCm to avoid HF Transformers flash attention - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - hf_model_kwargs = {} - if current_platform.is_rocm(): - hf_model_kwargs["attn_implementation"] = "eager" - with hf_runner( model, dtype=dtype, auto_cls=AutoModelForTokenClassification, - model_kwargs=hf_model_kwargs, ) as hf_model: tokenizer = hf_model.tokenizer hf_outputs = [] @@ -197,15 +176,10 @@ def test_openai_privacy_filter( with vllm_runner(model, max_model_len=None, dtype=dtype) as vllm_model: vllm_outputs = vllm_model.token_classify(PRIVACY_FILTER_PROMPTS) - hf_model_kwargs = {} - if current_platform.is_rocm(): - hf_model_kwargs["attn_implementation"] = "eager" - with hf_runner( model, dtype=dtype, auto_cls=AutoModelForTokenClassification, - model_kwargs=hf_model_kwargs, ) as hf_model: tokenizer = hf_model.tokenizer hf_outputs = [] @@ -276,12 +250,6 @@ def test_bert_for_masked_lm( with vllm_runner(model, max_model_len=None, dtype=dtype) as vllm_model: vllm_outputs = vllm_model.token_classify(example_prompts) - # Use eager attention on ROCm to avoid HF Transformers flash attention - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - hf_model_kwargs = {} - if current_platform.is_rocm(): - hf_model_kwargs["attn_implementation"] = "eager" - # Run hf_runner reference with "highest" fp32 precision to match # default behvior of vLLM. This is needed on ROCm since the # pooling tests set matmul precision to "high" in conftest.py @@ -291,7 +259,6 @@ def test_bert_for_masked_lm( model, dtype=dtype, auto_cls=AutoModelForMaskedLM, - model_kwargs=hf_model_kwargs, ) as hf_model: tokenizer = hf_model.tokenizer hf_outputs = [] diff --git a/tests/models/multimodal/conftest.py b/tests/models/multimodal/conftest.py index 70d4eee9b3..54e545d6bf 100644 --- a/tests/models/multimodal/conftest.py +++ b/tests/models/multimodal/conftest.py @@ -2,14 +2,8 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM project """Pytest configuration for vLLM multimodal tests.""" -import os -import warnings - -import torch import torch.nn.functional as F -from vllm.platforms import current_platform - # openbmb/MiniCPM-Llama3-V-2_5's remote ``resampler.py`` annotates with ``List`` # without importing it, relying on the alias leaking from # ``from torch.nn.functional import *``. Recent torch uses builtin generics and no @@ -17,62 +11,3 @@ from vllm.platforms import current_platform # TODO: Remove once the checkpoint fixes its imports. if not hasattr(F, "List"): F.List = list - - -def pytest_configure(config): - """Early ROCm configuration that must happen before test collection.""" - if not current_platform.is_rocm(): - return - - # Disable skinny GEMM on ROCm to avoid non-deterministic results - # from atomic reductions in wvSplitKrc kernel. - # See: https://github.com/vllm-project/vllm/pull/33493#issuecomment-3906083975 - os.environ["VLLM_ROCM_USE_SKINNY_GEMM"] = "0" - warnings.warn( - "ROCm: Set VLLM_ROCM_USE_SKINNY_GEMM=0 to avoid non-deterministic " - "results from skinny GEMM atomic reductions", - UserWarning, - stacklevel=1, - ) - - -def pytest_collection_modifyitems(config, items): - """Configure ROCm-specific settings based on collected tests.""" - if not current_platform.is_rocm(): - return - - skip_patterns = ["test_granite_speech.py"] - if any(pattern in str(arg) for arg in config.args for pattern in skip_patterns): - return - - # Disable Flash/MemEfficient SDP on ROCm to avoid HF Transformers - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - # TODO: Remove once ROCm SDP accuracy issues are resolved on HuggingFace - torch.backends.cuda.enable_flash_sdp(False) - torch.backends.cuda.enable_mem_efficient_sdp(False) - torch.backends.cuda.enable_math_sdp(True) - warnings.warn( - "ROCm: Disabled flash_sdp and mem_efficient_sdp, enabled math_sdp " - "to avoid HuggingFace Transformers accuracy issues", - UserWarning, - stacklevel=1, - ) - - -def patch_hf_vision_attn_for_rocm(model): - """Force SDPA for HF vision encoders on ROCm. - - HF's flash_attention_2 has accuracy issues on ROCm that bypass - torch.backends.cuda settings. This forces SDPA which then uses - math_sdp via the pytest_collection_modifyitems settings. - """ - if not current_platform.is_rocm(): - return - - inner = getattr(model, "model", model) - - if hasattr(inner, "vision_embedding"): - vit = inner.vision_embedding[0] - for layer in vit.encoder.layers: - if hasattr(layer, "self_attn"): - layer.self_attn.vision_config._attn_implementation = "sdpa" diff --git a/tests/models/multimodal/generation/test_common.py b/tests/models/multimodal/generation/test_common.py index 476cc44d9c..b367767c6b 100644 --- a/tests/models/multimodal/generation/test_common.py +++ b/tests/models/multimodal/generation/test_common.py @@ -249,10 +249,7 @@ VLM_TEST_SETTINGS = { vllm_runner_kwargs={ "model_impl": "transformers", }, - marks=[ - pytest.mark.core_model, - *([large_gpu_mark(min_gb=80)] if current_platform.is_rocm() else []), - ], + marks=[pytest.mark.core_model], ), "idefics3-transformers": VLMTestInfo( models=["HuggingFaceTB/SmolVLM-256M-Instruct"], @@ -282,19 +279,8 @@ VLM_TEST_SETTINGS = { image_size_factors=[(0.25, 0.2, 0.15)], vllm_runner_kwargs={ "model_impl": "transformers", - # TODO: [ROCm] Revert this once issue #30167 is resolved - **( - { - "mm_processor_kwargs": { - "min_pixels": 256 * 28 * 28, - "max_pixels": 1280 * 28 * 28, - }, - } - if current_platform.is_rocm() - else {} - ), }, - marks=[large_gpu_mark(min_gb=80 if current_platform.is_rocm() else 32)], + marks=[large_gpu_mark(min_gb=32)], ), #### Extended model tests "aria": VLMTestInfo( diff --git a/tests/models/multimodal/generation/test_whisper.py b/tests/models/multimodal/generation/test_whisper.py index 93634760a5..310c7fe0f5 100644 --- a/tests/models/multimodal/generation/test_whisper.py +++ b/tests/models/multimodal/generation/test_whisper.py @@ -129,7 +129,6 @@ def check_model_available(model: str) -> None: @pytest.mark.parametrize("max_tokens", [64]) @pytest.mark.parametrize("beam_width", [1, 2]) def test_beam_search_encoder_decoder( - monkeypatch, hf_runner, vllm_runner, dtype: str, @@ -138,9 +137,6 @@ def test_beam_search_encoder_decoder( resampled_assets, ) -> None: """Test beam search with encoder-decoder models (Whisper).""" - if current_platform.is_rocm(): - monkeypatch.setenv("VLLM_ROCM_USE_SKINNY_GEMM", "0") - model = "openai/whisper-large-v3-turbo" check_model_available(model) diff --git a/tests/models/multimodal/generation/vlm_utils/model_utils.py b/tests/models/multimodal/generation/vlm_utils/model_utils.py index a4c58c1215..e85a49311a 100644 --- a/tests/models/multimodal/generation/vlm_utils/model_utils.py +++ b/tests/models/multimodal/generation/vlm_utils/model_utils.py @@ -5,9 +5,7 @@ for manipulating the input / output of HF & vLLM test runners, which are typically specific to a small subset of models. """ -import logging import types -import warnings from pathlib import PosixPath import numpy as np @@ -34,8 +32,6 @@ from vllm.utils.collection_utils import is_list_of from .....conftest import HfRunner, ImageAsset, ImageTestAssets from .types import RunnerOutput -logger = logging.getLogger(__name__) - ####### vLLM output processors functions def blip2_vllm_to_hf_output(vllm_output: RunnerOutput, model: str) -> RunnerOutput: @@ -559,32 +555,6 @@ def isaac_patch_hf_runner(hf_model: HfRunner) -> HfRunner: # ---------------------------- isaac_model = hf_model.model.model - # [ROCm] Disable Flash/MemEfficient SDP on ROCm to avoid HF Transformers - # accuracy issues: https://github.com/vllm-project/vllm/issues/30167 - # TODO: Remove once ROCm SDP accuracy issues are resolved on HuggingFace - # ---------------------------- - from ...conftest import patch_hf_vision_attn_for_rocm - - try: - patch_hf_vision_attn_for_rocm(hf_model.model) - except AttributeError as e: - if "vision_config" in str(e): - warnings.warn( - f"Skipping ROCm vision attention patch for Isaac model: {e}. " - "This is expected for models without vision_config in " - "attention layers (e.g., Siglip2VariableLengthAttention).", - stacklevel=2, - ) - else: - logger.error( - "Unexpected AttributeError during ROCm vision attention patch: %s. " - "Model type: %s. Inner model type: %s.", - e, - type(hf_model.model).__name__, - type(getattr(hf_model.model, "model", None)).__name__, - ) - raise - def patched_forward( self, input_ids=None, diff --git a/tests/models/multimodal/pooling/conftest.py b/tests/models/multimodal/pooling/conftest.py deleted file mode 100644 index 401bc39b4b..0000000000 --- a/tests/models/multimodal/pooling/conftest.py +++ /dev/null @@ -1,18 +0,0 @@ -# SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project -"""Pytest configuration for vLLM pooling tests.""" - -import pytest - -from vllm.platforms import current_platform - - -@pytest.fixture -def siglip_attention_config(): - """Return attention config for SigLIP tests on ROCm. - - On ROCm, SigLIP tests require FLEX_ATTENTION backend. - """ - if current_platform.is_rocm(): - return {"backend": "FLEX_ATTENTION"} - return None diff --git a/tests/models/multimodal/pooling/test_siglip.py b/tests/models/multimodal/pooling/test_siglip.py index 4617250e38..bca598b42c 100644 --- a/tests/models/multimodal/pooling/test_siglip.py +++ b/tests/models/multimodal/pooling/test_siglip.py @@ -39,7 +39,6 @@ def _run_test( *, dtype: str, tokenization_kwargs: dict[str, Any] | None = None, - attention_config: dict[str, Any] | None = None, ) -> None: if tokenization_kwargs is None: tokenization_kwargs = {} @@ -51,7 +50,6 @@ def _run_test( enforce_eager=True, max_model_len=64, gpu_memory_utilization=0.7, - attention_config=attention_config, ) as vllm_model: vllm_outputs = vllm_model.embed( input_texts, images=input_images, tokenization_kwargs=tokenization_kwargs @@ -96,7 +94,6 @@ def test_models_text( hf_runner, vllm_runner, image_assets, - siglip_attention_config, model: str, dtype: str, ) -> None: @@ -115,7 +112,6 @@ def test_models_text( "padding": "max_length", "max_length": 64, }, # siglip2 was trained with this padding setting. - attention_config=siglip_attention_config, ) @@ -125,7 +121,6 @@ def test_models_image( hf_runner, vllm_runner, image_assets, - siglip_attention_config, model: str, dtype: str, ) -> None: @@ -142,7 +137,6 @@ def test_models_image( input_images, model, dtype=dtype, - attention_config=siglip_attention_config, ) @@ -151,7 +145,6 @@ def test_models_image( def test_models_text_image_no_crash( vllm_runner, image_assets, - siglip_attention_config, model: str, dtype: str, ) -> None: @@ -165,7 +158,6 @@ def test_models_text_image_no_crash( enforce_eager=True, max_model_len=64, gpu_memory_utilization=0.7, - attention_config=siglip_attention_config, ) as vllm_model: with pytest.raises(ValueError, match="not both"): vllm_model.embed(texts, images=images) diff --git a/tests/samplers/test_beam_search.py b/tests/samplers/test_beam_search.py index 5104469663..91dee5634b 100644 --- a/tests/samplers/test_beam_search.py +++ b/tests/samplers/test_beam_search.py @@ -50,7 +50,6 @@ MODELS = ["TinyLlama/TinyLlama-1.1B-Chat-v1.0"] @pytest.mark.parametrize("max_tokens", MAX_TOKENS) @pytest.mark.parametrize("beam_width", BEAM_WIDTHS) def test_beam_search_single_input( - monkeypatch, hf_runner, vllm_runner, example_prompts, @@ -59,9 +58,6 @@ def test_beam_search_single_input( max_tokens: int, beam_width: int, ) -> None: - if current_platform.is_rocm(): - monkeypatch.setenv("VLLM_ROCM_USE_SKINNY_GEMM", "0") - example_prompts = example_prompts[:1] with hf_runner(model, dtype=dtype) as hf_model: hf_outputs = hf_model.generate_beam_search( @@ -95,7 +91,6 @@ def test_beam_search_single_input( @pytest.mark.parametrize("max_tokens", MAX_TOKENS) @pytest.mark.parametrize("beam_width", BEAM_WIDTHS) def test_beam_search_with_concurrency_limit( - monkeypatch, hf_runner, vllm_runner, example_prompts, @@ -104,9 +99,6 @@ def test_beam_search_with_concurrency_limit( max_tokens: int, beam_width: int, ) -> None: - if current_platform.is_rocm(): - monkeypatch.setenv("VLLM_ROCM_USE_SKINNY_GEMM", "0") - # example_prompts[1]&[3]&[7] fails due to unknown reason even without # concurrency limit. skip them for now. example_prompts = example_prompts[:8] @@ -156,7 +148,6 @@ def test_beam_search_with_concurrency_limit( @pytest.mark.parametrize("max_tokens", MAX_TOKENS) @pytest.mark.parametrize("beam_width", MM_BEAM_WIDTHS) def test_beam_search_passes_multimodal_data( - monkeypatch, hf_runner, vllm_runner, dtype: str, @@ -164,9 +155,6 @@ def test_beam_search_passes_multimodal_data( beam_width: int, ) -> None: """Ensure that beam search passes multimodal data through correctly.""" - if current_platform.is_rocm(): - monkeypatch.setenv("VLLM_ROCM_USE_SKINNY_GEMM", "0") - # NOTE - this test is primarily to check that mm data is passed to beams # correctly. As such, we just need to check one extra modality to make # sure things pass through properly. diff --git a/tests/utils.py b/tests/utils.py index fc5ede2d99..fac455ab73 100644 --- a/tests/utils.py +++ b/tests/utils.py @@ -131,12 +131,6 @@ else: VLLM_PATH = Path(__file__).parent.parent """Path to root of the vLLM repository.""" -# ROCm: disable skinny GEMM to avoid non-deterministic results from -# atomic reductions in wvSplitKrc kernel. -# See: https://github.com/vllm-project/vllm/pull/33493#issuecomment-3906083975 -ROCM_ENV_OVERRIDES = ( - {"VLLM_ROCM_USE_SKINNY_GEMM": "0"} if current_platform.is_rocm() else {} -) # ROCm: disable prefix caching and eliminate batch variance to reduce # test flakiness. ROCM_EXTRA_ARGS = ( diff --git a/tests/v1/distributed/test_internal_lb_dp.py b/tests/v1/distributed/test_internal_lb_dp.py index 31859ffba4..70e416787c 100644 --- a/tests/v1/distributed/test_internal_lb_dp.py +++ b/tests/v1/distributed/test_internal_lb_dp.py @@ -12,7 +12,7 @@ import pytest import pytest_asyncio import requests -from tests.utils import ROCM_ENV_OVERRIDES, RemoteOpenAIServer +from tests.utils import RemoteOpenAIServer from tests.v1.utils import check_request_balancing from vllm.platforms import current_platform @@ -186,7 +186,6 @@ class MultinodeInternalLBServerManager: auto_port=False, env_dict={ "VLLM_SERVER_DEV_MODE": "1", - **ROCM_ENV_OVERRIDES, current_platform.device_control_env_var: ",".join( str(current_platform.device_id_to_physical_device_id(i)) for i in range(r, r + gpus_per_node) @@ -308,7 +307,6 @@ class APIOnlyServerManager: auto_port=False, env_dict={ "VLLM_SERVER_DEV_MODE": "1", - **ROCM_ENV_OVERRIDES, # No GPUs needed for API-only server }, ) @@ -329,7 +327,6 @@ class APIOnlyServerManager: engines_server_args, auto_port=False, env_dict={ - **ROCM_ENV_OVERRIDES, current_platform.device_control_env_var: ",".join( str(current_platform.device_id_to_physical_device_id(i)) for i in range(self.dp_size * self.tp_size) diff --git a/tests/v1/kv_connector/nixl_integration/test_accuracy.py b/tests/v1/kv_connector/nixl_integration/test_accuracy.py index 469ed5f9bc..76d7b8d9c3 100644 --- a/tests/v1/kv_connector/nixl_integration/test_accuracy.py +++ b/tests/v1/kv_connector/nixl_integration/test_accuracy.py @@ -38,6 +38,19 @@ SIMPLE_PROMPT = ( MODEL_NAME = os.environ.get("TEST_MODEL", "Qwen/Qwen3-0.6B") +def _meets_accuracy_threshold(measured_value: float, expected_value: float) -> bool: + return measured_value >= expected_value - RTOL + + +def test_accuracy_threshold_is_one_sided(): + expected_value = 0.77 + minimum_value = expected_value - RTOL + + assert _meets_accuracy_threshold(minimum_value, expected_value) + assert _meets_accuracy_threshold(expected_value + RTOL + 0.01, expected_value) + assert not _meets_accuracy_threshold(minimum_value - 0.01, expected_value) + + def run_simple_prompt(): client = openai.OpenAI(api_key="EMPTY", base_url=BASE_URL) completion = client.completions.create(model=MODEL_NAME, prompt=SIMPLE_PROMPT) @@ -93,7 +106,7 @@ def test_accuracy(): ) return - assert ( - measured_value - RTOL < expected_value - and measured_value + RTOL > expected_value - ), f"Expected: {expected_value} | Measured: {measured_value}" + minimum_value = expected_value - RTOL + assert _meets_accuracy_threshold(measured_value, expected_value), ( + f"Expected at least: {minimum_value} | Measured: {measured_value}" + ) diff --git a/tests/v1/kv_offload/cpu/test_gpu_worker.py b/tests/v1/kv_offload/cpu/test_gpu_worker.py index 9a90dedb3d..1b6e3a2586 100644 --- a/tests/v1/kv_offload/cpu/test_gpu_worker.py +++ b/tests/v1/kv_offload/cpu/test_gpu_worker.py @@ -18,6 +18,7 @@ from vllm.v1.kv_offload.base import ( CanonicalKVCaches, CanonicalKVCacheTensor, GPULoadStoreSpec, + TransferResult, ) from vllm.v1.kv_offload.cpu import gpu_worker from vllm.v1.kv_offload.cpu.common import CPULoadStoreSpec @@ -612,3 +613,72 @@ def test_transfer_multi_group( ) worker.shutdown() + + +@pytest.mark.skipif( + not current_platform.is_cuda_alike(), + reason="stream ordering test requires a CUDA-like platform", +) +@torch.inference_mode() +def test_load_waits_for_pending_compute_stream_writes(default_vllm_config) -> None: + """A CPU load must land after pending writes to its GPU destination.""" + device = DEVICES[0] + page_size_bytes = 128 * 1024 + num_blocks = 64 + loaded_blocks = list(range(32)) + sentinel = 0x5A + + gpu_tensor = torch.zeros( + (num_blocks, page_size_bytes), dtype=torch.int8, device=device + ) + loaded_block_ids = torch.tensor(loaded_blocks, dtype=torch.long, device=device) + worker = CPUOffloadingWorker( + kv_caches=CanonicalKVCaches( + tensors=[ + CanonicalKVCacheTensor( + tensor=gpu_tensor, page_size_bytes=page_size_bytes + ) + ], + group_data_refs=[ + [CanonicalKVCacheRef(tensor_idx=0, page_size_bytes=page_size_bytes)] + ], + ), + blocks_per_chunk=1, + num_cpu_blocks=num_blocks, + ) + worker._load_handler.src_tensors[0].fill_(sentinel) + expected = torch.full((page_size_bytes,), sentinel, dtype=torch.int8) + + try: + for trial in range(3): + gpu_tensor.fill_(0x11) + torch.accelerator.synchronize() + + # Model a delayed zero of a freshly allocated KV block. Without a + # compute-stream dependency, the DMA can finish during the sleep + # and this later fill wipes out the loaded cache contents. + torch.cuda._sleep(50_000_000) + gpu_tensor.index_fill_(0, loaded_block_ids, 0) + + assert worker.submit_load( + trial + 1, + CPULoadStoreSpec(loaded_blocks), + GPULoadStoreSpec( + loaded_blocks, + group_sizes=(len(loaded_blocks),), + block_indices=(0,), + ), + ) + deadline = time.time() + 10 + finished: list[TransferResult] = [] + while time.time() < deadline and not finished: + finished = worker.get_finished() + if not finished: + time.sleep(0.001) + assert finished and finished[0].success, f"load {trial} did not finish" + + torch.accelerator.synchronize() + for block_id in loaded_blocks: + torch.testing.assert_close(gpu_tensor[block_id].cpu(), expected) + finally: + worker.shutdown() diff --git a/tests/v1/sample/test_logprobs.py b/tests/v1/sample/test_logprobs.py index d643b0b6fd..05a1e93618 100644 --- a/tests/v1/sample/test_logprobs.py +++ b/tests/v1/sample/test_logprobs.py @@ -1114,7 +1114,6 @@ def test_correct_decoded_token_preserves_valid_tokens(): def test_spec_decode_logprobs( logprobs_mode: LogprobsMode, model_setup: tuple[str, str, dict, int], - monkeypatch, ): """Spec decode logprobs should match those of the base model. @@ -1127,20 +1126,9 @@ def test_spec_decode_logprobs( logprobs_mode: logprobs mode. model_setup: Tuple of (method, base model name, speculative_config dict, top_logprobs). - monkeypatch: pytest fixture for setting env vars. """ from vllm import LLM - # The ROCm skinny GEMM kernels (gemm_kernels.cu) are - # non-deterministic across LLM instantiations due to persistent - # workgroup scheduling and wave-level shuffle reductions, which - # causes logprob differences that get misattributed to spec decode. - # Disable them so this test isolates spec decode correctness only. - # TODO(akaratza): Remove this workaround once the follow-up to - # https://github.com/vllm-project/vllm/pull/33493#issuecomment-3906083975 - # lands with a determinism fix for wvSplitK kernels. - monkeypatch.setenv("VLLM_ROCM_USE_SKINNY_GEMM", "0") - method, model_name, spec_config, top_logprobs = model_setup prompt = "Hello world " * 50 diff --git a/tests/v1/worker/test_gpu_model_runner.py b/tests/v1/worker/test_gpu_model_runner.py index 814d060767..2db024dc6b 100644 --- a/tests/v1/worker/test_gpu_model_runner.py +++ b/tests/v1/worker/test_gpu_model_runner.py @@ -1,6 +1,7 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +import gc from types import SimpleNamespace from unittest.mock import Mock @@ -126,6 +127,34 @@ def get_vllm_config(): return vllm_config +@pytest.mark.parametrize("gc_initially_enabled", [True, False]) +def test_freeze_gc_disables_and_restores_automatic_gc( + monkeypatch: pytest.MonkeyPatch, + gc_initially_enabled: bool, +): + monkeypatch.setenv("VLLM_ENABLE_CUDAGRAPH_GC", "0") + original_gc_state = gc.isenabled() + try: + if gc_initially_enabled: + gc.enable() + else: + gc.disable() + + with ( + pytest.raises(RuntimeError, match="capture failed"), + GPUModelRunner._freeze_gc(), + ): + assert not gc.isenabled() + raise RuntimeError("capture failed") + + assert gc.isenabled() is gc_initially_enabled + finally: + if original_gc_state: + gc.enable() + else: + gc.disable() + + @pytest.fixture def model_runner(): vllm_config = get_vllm_config() @@ -896,6 +925,37 @@ def test_sample_passes_reordered_draft_probs_to_rejection_sampler(): assert torch.equal(passed_draft_probs, expected_draft_probs) +def test_dummy_sampler_run_warms_all_greedy_rejection_sampler(monkeypatch): + runner = object.__new__(GPUModelRunner) + runner.device = torch.device("cpu") + runner.vllm_config = SimpleNamespace( + model_config=SimpleNamespace(multimodal_config=None) + ) + runner.model = SimpleNamespace(compute_logits=Mock(return_value=torch.randn(3, 8))) + runner.sampler = Mock(return_value="sampler_output") + runner.sampler.logprobs_mode = "processed_logprobs" + runner.speculative_config = SimpleNamespace( + rejection_sample_method="standard", + draft_sample_method="greedy", + ) + runner.rejection_sampler = Mock() + synchronize = Mock() + monkeypatch.setattr(torch.accelerator, "synchronize", synchronize) + + output = GPUModelRunner._dummy_sampler_run(runner, torch.randn(3, 4)) + + assert output == "sampler_output" + assert runner.rejection_sampler.call_count == 2 + mixed_metadata = runner.rejection_sampler.call_args_list[0].args[3] + all_greedy_metadata = runner.rejection_sampler.call_args_list[1].args[3] + assert not mixed_metadata.all_greedy + assert mixed_metadata.temperature is not None + assert all_greedy_metadata.all_greedy + assert not all_greedy_metadata.all_random + assert all_greedy_metadata.temperature is None + synchronize.assert_called_once_with() + + def test_invalid_draft_suffixes_remain_rejected_in_metadata(): runner = object.__new__(GPUModelRunner) runner.device = torch.device("cpu") diff --git a/vllm/_aiter_ops.py b/vllm/_aiter_ops.py index 7b1403b762..982de32b0d 100644 --- a/vllm/_aiter_ops.py +++ b/vllm/_aiter_ops.py @@ -1,12 +1,15 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +import ctypes import functools +import os from collections.abc import Callable import torch from torch._ops import OpOverload import vllm.envs as envs +from vllm.logger import init_logger from vllm.platforms import current_platform from vllm.utils.import_utils import PlaceholderModule from vllm.utils.torch_utils import direct_register_custom_op @@ -15,6 +18,8 @@ from vllm.v1.attention.ops.rocm_aiter_mla_sparse import ( rocm_aiter_sparse_attn_indexer_fake, ) +logger = init_logger(__name__) + try: import pandas as pd except ImportError: @@ -49,6 +54,83 @@ def is_aiter_found() -> bool: IS_AITER_FOUND = is_aiter_found() +class _DlInfo(ctypes.Structure): + _fields_ = [ + ("dli_fname", ctypes.c_char_p), + ("dli_fbase", ctypes.c_void_p), + ("dli_sname", ctypes.c_char_p), + ("dli_saddr", ctypes.c_void_p), + ] + + +_HIP_RUNTIME_HANDLE: ctypes.CDLL | None = None + + +def _find_torch_hip_runtime_path() -> str | None: + """Find the HIP runtime already selected by PyTorch's extension graph.""" + try: + torch_extension = ctypes.CDLL(torch._C.__file__) + hip_symbol = torch_extension.hipGetDeviceCount + dladdr = ctypes.CDLL(None).dladdr + dladdr.argtypes = [ctypes.c_void_p, ctypes.POINTER(_DlInfo)] + dladdr.restype = ctypes.c_int + info = _DlInfo() + if not dladdr(ctypes.cast(hip_symbol, ctypes.c_void_p), ctypes.byref(info)): + return None + except (AttributeError, OSError, TypeError, ctypes.ArgumentError): + return None + if info.dli_fname is None: + return None + + runtime_path = os.fsdecode(info.dli_fname) + if "libamdhip64.so" not in os.path.basename(runtime_path): + return None + return runtime_path + + +def _promote_torch_hip_runtime() -> None: + """Make PyTorch's HIP runtime symbols visible to AITER JIT extensions. + + PyTorch loads its native extension with ``RTLD_LOCAL``. Some AITER JIT + extensions rely on HIP registration and launch symbols being available in + the process-global lookup scope. Reopening the exact runtime that provides + PyTorch's ``hipGetDeviceCount`` symbol promotes the existing DSO instead of + selecting or initializing a second HIP runtime. + """ + global _HIP_RUNTIME_HANDLE + if _HIP_RUNTIME_HANDLE is not None: + return + + runtime_path = _find_torch_hip_runtime_path() + if runtime_path is None: + logger.warning_once( + "Could not locate PyTorch's HIP runtime; AITER JIT extensions may " + "fail to resolve HIP symbols." + ) + return + + try: + # Retain the handle for the process lifetime. Besides preventing a + # future dlclose, this documents that the RTLD_GLOBAL promotion is + # intentional and must outlive every lazily imported AITER extension. + _HIP_RUNTIME_HANDLE = ctypes.CDLL(runtime_path, mode=ctypes.RTLD_GLOBAL) + except OSError as err: + logger.warning_once( + "Could not promote PyTorch's HIP runtime at %s; AITER JIT " + "extensions may fail to resolve HIP symbols: %s", + runtime_path, + err, + ) + + +def _maybe_promote_torch_hip_runtime() -> None: + if current_platform.is_rocm() and IS_AITER_FOUND: + _promote_torch_hip_runtime() + + +_maybe_promote_torch_hip_runtime() + + def is_aiter_found_and_supported() -> bool: """Check if AITER library is available and platform supports it. diff --git a/vllm/model_executor/layers/utils.py b/vllm/model_executor/layers/utils.py index 1eec3ba47b..16539bdf65 100644 --- a/vllm/model_executor/layers/utils.py +++ b/vllm/model_executor/layers/utils.py @@ -134,12 +134,16 @@ def rocm_unquantized_gemm_impl( GrpsShrB = min(N_p2 // 16, 4) # Given the above, how many CUs would we need? CuNeeded = rndup_cus * GrpsShrB - # candidate for atomic reduce count splitk? + # Deterministic reduction stores one float workspace value per K shard. fits_wvsplitkrc = ( N_p2 * m * ((k + 512 - 1) // 512) ) <= 128 * 1024 * 12 # deterministic fits_wvsplitkrc &= CuNeeded <= cu_count + skinny_operands_compatible = weight.is_contiguous() and ( + bias is None or bias.is_contiguous() + ) + use_skinny_reduce_counting = ( envs.VLLM_ROCM_USE_SKINNY_GEMM and on_gfx950() @@ -151,12 +155,13 @@ def rocm_unquantized_gemm_impl( and k > 512 and m % 16 == 0 and fits_wvsplitkrc - and weight.is_contiguous() + and skinny_operands_compatible ) ) if use_skinny_reduce_counting: - return ops.wvSplitKrc(x, weight, cu_count, bias) + x_view = x.reshape(-1, x.size(-1)).contiguous() + return ops.wvSplitKrc(x_view, weight, cu_count, bias) # gfx1250's aiter gemm_a16w16 uses the gluon backend, which requires # K % 256 == 0 (it walks K with fixed-size descriptors and won't pad a @@ -174,10 +179,13 @@ def rocm_unquantized_gemm_impl( # TODO GFX1250: Include once skinny GEMM is supported on gfx1250 and x.dtype in [torch.float16, torch.bfloat16] and k % 8 == 0 + and skinny_operands_compatible ) if use_skinny: - x_view = x.reshape(-1, x.size(-1)) + # The skinny kernels assume contiguous K elements. A shape-preserving + # reshape can retain a transposed activation's non-contiguous strides. + x_view = x.reshape(-1, x.size(-1)).contiguous() if m > 8 and 0 < n <= 5: cu_count = num_compute_units() out = ops.wvSplitK(weight, x_view, cu_count, bias) diff --git a/vllm/model_executor/models/transformers/multimodal.py b/vllm/model_executor/models/transformers/multimodal.py index f073e8c186..84d6813c8e 100644 --- a/vllm/model_executor/models/transformers/multimodal.py +++ b/vllm/model_executor/models/transformers/multimodal.py @@ -18,7 +18,7 @@ from collections import defaultdict from collections.abc import Mapping, Sequence -from contextlib import ExitStack, contextmanager, nullcontext +from contextlib import ExitStack, contextmanager from typing import TYPE_CHECKING, Any import torch @@ -53,7 +53,6 @@ from vllm.multimodal.processing import ( PromptUpdate, TimingContext, ) -from vllm.platforms import current_platform from vllm.sequence import IntermediateTensors if TYPE_CHECKING: @@ -727,15 +726,9 @@ class MultiModalMixin(SupportsMultiModal, SupportsMRoPE): kwargs.pop("token_type_ids", None) kwargs.pop("mm_token_type_ids", None) - context = nullcontext() - if current_platform.is_rocm(): - context = torch.nn.attention.sdpa_kernel( - backends=[torch.nn.attention.SDPBackend.MATH] - ) - with context: - audio_output = self.model.get_audio_features( - input_features, return_dict=True, **kwargs - ) + audio_output = self.model.get_audio_features( + input_features, return_dict=True, **kwargs + ) audio_embeddings = audio_output.pooler_output split_sizes = num_audio_tokens.flatten().tolist() @@ -756,23 +749,7 @@ class MultiModalMixin(SupportsMultiModal, SupportsMRoPE): num_image_patches = kwargs.pop("num_image_patches") - context = nullcontext() - if current_platform.is_rocm(): - # ROCm: Force math SDP backend for vision encoder to avoid accuracy issues - # with flash_sdp and mem_efficient_sdp - # TODO: [ROCm] Fix accuracy issues with flash backend - logger.debug( - "ROCm platform detected. Forcing math SDP backend " - "for vision encoder. Currently ROCm platform has " - "accuracy issues with `flash_sdp` and" - "`mem_efficient_sdp` backends. See issue: " - "https://github.com/vllm-project/vllm/issues/30167" - ) - context = torch.nn.attention.sdpa_kernel( - backends=[torch.nn.attention.SDPBackend.MATH] - ) - with context: - vision_embeddings = self.model.get_image_features(pixel_values, **kwargs) + vision_embeddings = self.model.get_image_features(pixel_values, **kwargs) # Transformers `v5`, `self.get_image_features` returns a tuple # containing the features and optionally attentions/hidden_states diff --git a/vllm/v1/core/kv_cache_utils.py b/vllm/v1/core/kv_cache_utils.py index 770dffa4f2..7640093f6e 100644 --- a/vllm/v1/core/kv_cache_utils.py +++ b/vllm/v1/core/kv_cache_utils.py @@ -619,8 +619,7 @@ def resolve_kv_cache_block_sizes( ``cache_config.prefix_match_unit`` override if set, else the GCD of group block sizes; every group's block size must be divisible by it. Returns the scheduler block size (i.e. disables finer hashing) if block - hashing is inactive or a mamba group's block size diverges from the - cache block size (mamba_cache_mode != "align"). + hashing is inactive or a mamba group is not using cache mode "align". """ cache_config = vllm_config.cache_config dcp = vllm_config.parallel_config.decode_context_parallel_size diff --git a/vllm/v1/kv_offload/cpu/gpu_worker.py b/vllm/v1/kv_offload/cpu/gpu_worker.py index ab4fedcbe0..d2625c2376 100644 --- a/vllm/v1/kv_offload/cpu/gpu_worker.py +++ b/vllm/v1/kv_offload/cpu/gpu_worker.py @@ -634,9 +634,11 @@ class SingleDirectionOffloadingHandler: else torch.Event(enable_timing=True) ) - if self.gpu_to_cpu: - # wait for model computation to finish before offloading - stream.wait_stream(current_platform.current_stream()) + # Stores must wait for the model to finish writing the KV they read. + # Loads must wait for pending writes (including zeroing) to their + # destination blocks; otherwise an earlier transfer can be overwritten + # by compute-stream work that was already queued when the load began. + stream.wait_stream(current_platform.current_stream()) if self._transfers: last_transfer: Transfer = self._transfers[-1] last_event = last_transfer.end_event diff --git a/vllm/v1/worker/gpu_model_runner.py b/vllm/v1/worker/gpu_model_runner.py index d11cb858e9..2f9c4f811e 100644 --- a/vllm/v1/worker/gpu_model_runner.py +++ b/vllm/v1/worker/gpu_model_runner.py @@ -6420,6 +6420,21 @@ class GPUModelRunner( logits, dummy_metadata, ) + # All-greedy is a separate Triton specialization from the + # mixed-sampling path above. Compile and exercise it before serving. + all_greedy_metadata = replace( + dummy_metadata, + temperature=None, + all_greedy=True, + all_random=False, + ) + self.rejection_sampler( + dummy_spec_decode_metadata, + draft_probs, + logits, + all_greedy_metadata, + ) + torch.accelerator.synchronize() return sampler_output def _dummy_pooler_run_task( @@ -6615,16 +6630,26 @@ class GPUModelRunner( @staticmethod @contextmanager def _freeze_gc(): + gc_was_enabled = gc.isenabled() gc.collect() should_freeze = not envs.VLLM_ENABLE_CUDAGRAPH_GC if should_freeze: gc.freeze() + # A Triton kernel finalized during stream capture unloads its + # module and invalidates the captured graph. + gc.disable() try: yield finally: if should_freeze: - gc.unfreeze() - gc.collect() + try: + gc.unfreeze() + gc.collect() + finally: + if gc_was_enabled: + gc.enable() + else: + gc.disable() def shutdown(self) -> None: """Release GPU tensors (model weights, KV caches, workspace) so that