Upstream sync 13/N: merge 1ab2801dde [ROCm] Remove stale SDPA and skinny GEMM workarounds (#50907) (conflict)

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 @@ def generate(self, token_ids: list[int], max_new_tokens: int) -> list[int]:
 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 @@
 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_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 @@
 
 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 @@
 
 @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_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/test_correctness_mteb.py

@@ -11,7 +11,6 @@
     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 @@
 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 @@
 
 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 @@
 
 @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 @@
     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.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 @@
 
 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 @@
 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 @@
 
 @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.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 @@
 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 @@
     run_mteb_rerank,
 )
 from tests.utils import RemoteOpenAIServer
-from vllm.platforms import current_platform
 
 os.environ["VLLM_LOGGING_LEVEL"] = "WARNING"
 
@@ -25,10 +24,6 @@
 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 @@
 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 @@
 
 @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 @@
 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 @@
 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