[ROCm] Remove stale SDPA and skinny GEMM workarounds (#50907)

Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com>
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: OpenAI Codex <codex@openai.com>

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++) {

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"

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",

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
 
 

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)

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
 

tests/entrypoints/pooling/embed/conftest.py

@@ -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,
-    )

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
 

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
 

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

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
 

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
 

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)

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
 

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
 

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)

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
 

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"
 

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)

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)

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
 

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
 
 

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
 
 

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()

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:

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(

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):

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)

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
 

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,

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()]

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)

tests/models/language/generation/conftest.py

@@ -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,
-    )

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:

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")

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(

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 = []

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"

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(

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)
 

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,

tests/models/multimodal/pooling/conftest.py

@@ -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

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)

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.

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 = (

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)

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}"
+    )

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()

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

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")

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.
 

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)

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

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

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

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