[Hybrid] calling get_mamba_groups() once at MambaCopyBuffers.create() (#37318 )

Signed-off-by: Francesco Fusco <ffu@zurich.ibm.com>
[ROCm] Enable DeepEP ROCm as all2allbackend for AMD GPUs. (#34692 )
2026-03-21 09:29:43 +00:00 · 2026-03-21 00:32:31 -07:00 · 2026-03-21 15:27:12 +08:00 · 2026-03-21 13:48:54 +08:00 · 2026-03-21 12:57:25 +08:00 · 2026-03-20 21:35:49 -07:00
179 changed files with 4171 additions and 2741 deletions
@@ -127,7 +127,7 @@ run_and_track_test() {

 # --- Actual Test Execution ---
 run_and_track_test 1 "test_struct_output_generate.py" \
-    "python3 -m pytest -s -v /workspace/vllm/tests/v1/entrypoints/llm/test_struct_output_generate.py -k \"not test_structured_output_with_reasoning_matrices\""
+    "python3 -m pytest -s -v /workspace/vllm/tests/entrypoints/llm/test_struct_output_generate.py -k \"not test_structured_output_with_reasoning_matrices\""
 run_and_track_test 2 "test_moe_pallas.py" \
    "python3 -m pytest -s -v /workspace/vllm/tests/tpu/test_moe_pallas.py"
 run_and_track_test 3 "test_lora.py" \
@@ -39,8 +39,7 @@
 #####################################################################################################################################
 #                                                                                                                                   #
 # IMPORTANT:                                                                                                                        #
-#   * Currently AMD CI has MI300 agents, MI325 agents, and MI355 agents. Of those, AMD is using mostly MI325 and MI355. AMD team    #
-#     is actively working on enabling more MI300 machines. All upcoming feature improvements are tracked in:                        #
+#   * Currently AMD CI has MI250 agents, MI325 agents, and MI355 agents. All upcoming feature improvements are tracked in:          #
 #         https://github.com/vllm-project/vllm/issues/34994                                                                         #
 #                                                                                                                                   #
 #-----------------------------------------------------------------------------------------------------------------------------------#
@@ -49,13 +48,15 @@
 #   * [Pytorch Nightly Dependency Override Check]: if this test fails, it means the nightly torch version is not compatible with    #
 #                                                  some of the dependencies. Please check the error message and add the package to  #
 #                                                  whitelist in `/vllm/tools/pre_commit/generate_nightly_torch_test.py`.            #
-#   * [Entrypoints Integration Test (LLM)]:                                                                                         #
+#   * [Entrypoints Integration (LLM)]:                                                                                              #
 #     - {`pytest -v -s entrypoints/llm/test_generate.py`}: It needs a clean process                                                 #
 #     - {`pytest -v -s entrypoints/offline_mode`}: Needs to avoid interference with other tests                                     #
-#   * [V1 Test e2e + engine]: The test uses 4 GPUs, but we schedule it on 8-GPU machines for stability. See discussion here:        #
-#                             https://github.com/vllm-project/vllm/pull/31040                                                       #
-#   * [V1 others]:                                                                                                                  #
-#     - Split the tests to avoid interference                                                                                       #
+#   * [Engine / Engine (1 GPU) / e2e Scheduling / e2e Core / V1 e2e / Spec Decode / V1 Sample + Logits / V1 Core + KV + Metrics]:   #
+#     - Previously a single "V1 Test e2e + engine" step, now split across multiple groups.                                          #
+#     - V1 e2e (2/4 GPUs) uses 4 GPUs but is scheduled on 8-GPU machines for stability. See:                                        #
+#       https://github.com/vllm-project/vllm/pull/31040                                                                             #
+#   * [V1 Sample + Logits / V1 Core + KV + Metrics / V1 others (CPU)]:                                                              #
+#     - Previously a single "V1 others" step, now split to avoid interference.                                                      #
 #     - Integration test for streaming correctness (requires special branch for __harness__ lib).                                   #
 #   * [V1 others (CPU)]: Split the tests to avoid interference                                                                      #
 #   * [PyTorch Compilation Unit Tests]: Run unit tests defined directly under `compile/`, not including subdirectories, which       #
@@ -83,9 +84,9 @@
 #                                          run plamo2 model in vLLM.                                                                #
 #   * [Language Models Test (Extended Generation)]: Install fast path packages for testing against transformers (mamba, conv1d)     #
 #                                                   and to run plamo2 model in vLLM.                                                #
-#   * [Multi-Modal Models (Standard)]:                                                                                              #
+#   * [Multi-Modal Models (Standard) 1-4]:                                                                                          #
 #     - Do NOT remove `VLLM_WORKER_MULTIPROC_METHOD=spawn` setting as ROCm requires this for certain models to function.            #
-#   * [Transformers Nightly Models Test]: Whisper needs `VLLM_WORKER_MULTIPROC_METHOD=spawn` to avoid deadlock.                     #
+#   * [Transformers Nightly Models]: Whisper needs `VLLM_WORKER_MULTIPROC_METHOD=spawn` to avoid deadlock.                          #
 #   * [Plugin Tests (2 GPUs)]:                                                                                                      #
 #     - {`pytest -v -s entrypoints/openai/test_oot_registration.py`}: It needs a clean process                                      #
 #     - {`pytest -v -s models/test_oot_registration.py`}: It needs a clean process                                                  #
@@ -94,11 +95,11 @@
 #     - There is some Tensor Parallelism related processing logic in LoRA that requires multi-GPU testing for validation.           #
 #     - {`pytest -v -s -x lora/test_gptoss_tp.py`}: Disabled for now because MXFP4 backend on non-cuda platform doesn't support     #
 #                                                   LoRA yet.                                                                       #
-#   * [Distributed Tests (GPU_TAG)]: Don't test llama model here, it seems hf implementation is buggy. See:                         #
-#                                    https://github.com/vllm-project/vllm/pull/5689                                                 #
-#   * [Distributed Tests (GPU_TAG)]: Some old E2E tests were removed in https://github.com/vllm-project/vllm/pull/33293 in          #
-#                                    favor of new tests in fusions_e2e. We avoid replicating the new jobs in                        #
-#                                    this file as it's deprecated.                                                                  #
+#   * [Distributed Tests (NxGPUs)(HW-TAG)]: Don't test llama model here, it seems hf implementation is buggy. See:                  #
+#                                           https://github.com/vllm-project/vllm/pull/5689                                          #
+#   * [Distributed Tests (NxGPUs)(HW-TAG)]: Some old E2E tests were removed in https://github.com/vllm-project/vllm/pull/33293     #
+#                                           in favor of new tests in fusions_e2e. We avoid replicating the new jobs in              #
+#                                           this file as it's deprecated.                                                           #
 #                                                                                                                                   #
 #####################################################################################################################################

@@ -223,7 +224,7 @@ steps:
  - vllm/platforms/rocm.py
  commands:
  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling
+  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling


 - label: Entrypoints Integration (LLM) # TBD
@@ -254,11 +255,11 @@ steps:
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/rpc
-  - tests/entrypoints/instrumentator
+  - tests/entrypoints/serve/instrumentator
  - tests/tool_use
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/instrumentator
+  - pytest -v -s entrypoints/serve/instrumentator
  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
  - pytest -v -s tool_use

@@ -483,19 +484,6 @@ steps:
    - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"


- label: Entrypoints V1 # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
-  agent_pool: mi250_1
-  optional: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/v1
-  commands:
-  - pytest -v -s v1/entrypoints
-
-
 - label: V1 Sample + Logits # TBD
  timeout_in_minutes: 60
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -1173,14 +1161,14 @@ steps:
  - vllm/v1/engine/
  - vllm/v1/worker/
  - tests/v1/distributed
-  - tests/v1/entrypoints/openai/test_multi_api_servers.py
+  - tests/entrypoints/openai/test_multi_api_servers.py
  - vllm/platforms/rocm.py
  commands:
  - export TORCH_NCCL_BLOCKING_WAIT=1
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
-  - DP_SIZE=2 pytest -v -s v1/entrypoints/openai/test_multi_api_servers.py
+  - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py


 - label: Distributed Compile + RPC Tests (2 GPUs) # TBD
@@ -1402,7 +1390,7 @@ steps:
 - label: Distributed Tests (2 GPUs)(H100-MI250) # TBD
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
-  agent_pool: mi250_2
+  agent_pool: mi325_2
  num_gpus: 2
  working_dir: "/vllm-workspace/"
  source_file_dependencies:
@@ -1412,7 +1400,6 @@ steps:
  - vllm/v1/attention/backends/
  - vllm/v1/attention/selector.py
  - tests/distributed/test_context_parallel.py
-  - tests/v1/distributed/test_dbo.py
  - examples/offline_inference/data_parallel.py
  - vllm/_aiter_ops.py
  - vllm/platforms/rocm.py
@@ -1420,7 +1407,6 @@ steps:
  - export TORCH_NCCL_BLOCKING_WAIT=1
  - pytest -v -s tests/distributed/test_context_parallel.py
  - VLLM_LOGGING_LEVEL=DEBUG python3 examples/offline_inference/data_parallel.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=allgather_reducescatter --disable-nccl-for-dp-synchronization
-  - pytest -v -s tests/v1/distributed/test_dbo.py


 #####################################################################################################################################
@@ -1477,11 +1463,11 @@ steps:
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/rpc
-  - tests/entrypoints/instrumentator
+  - tests/entrypoints/serve/instrumentator
  - tests/tool_use
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/instrumentator
+  - pytest -v -s entrypoints/serve/instrumentator
  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
  - pytest -v -s tool_use

@@ -1760,6 +1746,7 @@ steps:
  timeout_in_minutes: 106
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_4
+  optional: true
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -1768,19 +1755,6 @@ steps:
    - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle_correctness_heavy"


- label: Entrypoints V1 # 25.7m
-  timeout_in_minutes: 45
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
-  agent_pool: mi325_1
-  optional: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/v1
-  commands:
-  - pytest -v -s v1/entrypoints
-
-
 - label: V1 Spec Decode # TBD
  timeout_in_minutes: 40
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
@@ -2393,14 +2367,14 @@ steps:
  - vllm/v1/engine/
  - vllm/v1/worker/
  - tests/v1/distributed
-  - tests/v1/entrypoints/openai/test_multi_api_servers.py
+  - tests/entrypoints/openai/test_multi_api_servers.py
  - vllm/platforms/rocm.py
  commands:
  - export TORCH_NCCL_BLOCKING_WAIT=1
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
-  - DP_SIZE=2 pytest -v -s v1/entrypoints/openai/test_multi_api_servers.py
+  - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py


 - label: Distributed Compile + RPC Tests (2 GPUs) # 56.1m
@@ -2578,6 +2552,7 @@ steps:
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_4
  num_gpus: 4
+  optional: true
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -2594,21 +2569,16 @@ steps:
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_2
  num_gpus: 2
-  optional: true
  working_dir: "/vllm-workspace/"
  source_file_dependencies:
  - vllm/distributed/
  - vllm/v1/distributed/
  - vllm/model_executor/layers/fused_moe/
-  - tests/distributed/test_context_parallel.py
  - tests/v1/distributed/test_dbo.py
-  - examples/offline_inference/data_parallel.py
  - vllm/_aiter_ops.py
  - vllm/platforms/rocm.py
  commands:
  - export TORCH_NCCL_BLOCKING_WAIT=1
-  - pytest -v -s tests/distributed/test_context_parallel.py
-  - VLLM_LOGGING_LEVEL=DEBUG python3 examples/offline_inference/data_parallel.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
  - pytest -v -s tests/v1/distributed/test_dbo.py


@@ -2667,7 +2637,7 @@ steps:
  - vllm/_aiter_ops.py
  - vllm/platforms/rocm.py
  commands:
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8.txt
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt


 - label: LM Eval Large Models (H200-MI325) # TBD
@@ -2698,6 +2668,7 @@ steps:
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_4
  num_gpus: 4
+  optional: true
  working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
  source_file_dependencies:
  - csrc/
@@ -2718,6 +2689,7 @@ steps:
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_4
  num_gpus: 4
+  optional: true
  working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
  source_file_dependencies:
  - csrc/
@@ -2783,6 +2755,7 @@ steps:
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_4
  num_gpus: 4
+  optional: true
  working_dir: "/vllm-workspace"
  source_file_dependencies:
  - vllm/model_executor/models/
@@ -2825,6 +2798,7 @@ steps:
  timeout_in_minutes: 11
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_4
+  optional: true
  working_dir: "/vllm-workspace"
  source_file_dependencies:
  - vllm/model_executor/models/
@@ -2846,6 +2820,7 @@ steps:
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_4
  num_gpus: 4
+  optional: true
  working_dir: "/vllm-workspace"
  source_file_dependencies:
  - vllm/model_executor/models/
@@ -2988,11 +2963,11 @@ steps:
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/rpc
-  - tests/entrypoints/instrumentator
+  - tests/entrypoints/serve/instrumentator
  - tests/tool_use
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/instrumentator
+  - pytest -v -s entrypoints/serve/instrumentator
  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
  - pytest -v -s tool_use

@@ -3597,7 +3572,7 @@ steps:
  - vllm/_aiter_ops.py
  - vllm/platforms/rocm.py
  commands:
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8.txt
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt


 - label: LM Eval Large Models (4 GPUs)(FP8) # TBD
@@ -27,14 +27,14 @@ steps:
  - vllm/v1/engine/
  - vllm/v1/worker/
  - tests/v1/distributed
-  - tests/v1/entrypoints/openai/test_multi_api_servers.py
+  - tests/entrypoints/openai/test_multi_api_servers.py
  commands:
  # https://github.com/NVIDIA/nccl/issues/1838
  - export NCCL_CUMEM_HOST_ENABLE=0
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
-  - DP_SIZE=2 pytest -v -s v1/entrypoints/openai/test_multi_api_servers.py
+  - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py

 - label: Distributed Compile + RPC Tests (2 GPUs)
  timeout_in_minutes: 20
@@ -70,3 +70,15 @@ steps:
      device: mi325_4
      depends_on:
      - image-build-amd
+
+- label: V1 e2e (4xH100)
+  timeout_in_minutes: 60
+  device: h100
+  num_devices: 4
+  optional: true
+  source_file_dependencies:
+    - vllm/v1/attention/backends/utils.py
+    - vllm/v1/worker/gpu_model_runner.py
+    - tests/v1/e2e/test_hybrid_chunked_prefill.py
+  commands:
+    - pytest -v -s v1/e2e/test_hybrid_chunked_prefill.py
@@ -10,7 +10,7 @@ steps:
  - tests/entrypoints/
  commands:
  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling
+  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling

 - label: Entrypoints Integration (LLM)
  timeout_in_minutes: 40
@@ -34,7 +34,7 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py --ignore=entrypoints/openai/correctness/  --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py --ignore=entrypoints/openai/correctness/  --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
  - pytest -v -s entrypoints/test_chat_utils.py
  mirror:
    amd:
@@ -48,11 +48,11 @@ steps:
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/rpc
-  - tests/entrypoints/instrumentator
+  - tests/entrypoints/serve/instrumentator
  - tests/tool_use
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/instrumentator
+  - pytest -v -s entrypoints/serve/instrumentator
  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
  - pytest -v -s tool_use

@@ -75,19 +75,6 @@ steps:
  commands:
  - pytest -v -s entrypoints/openai/responses

- label: Entrypoints V1
-  timeout_in_minutes: 50
-  source_file_dependencies:
-    - vllm/
-    - tests/v1
-  commands:
-    - pytest -v -s v1/entrypoints
-  mirror:
-    amd:
-      device: mi325_1
-      depends_on:
-      - image-build-amd
-
 - label: OpenAI API Correctness
  timeout_in_minutes: 30
  source_file_dependencies:
@@ -45,6 +45,22 @@ steps:
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt

+- label: LM Eval Qwen3.5 Models (B200)
+  timeout_in_minutes: 120
+  device: b200
+  optional: true
+  num_devices: 2
+  source_file_dependencies:
+  - vllm/model_executor/models/qwen3_5.py
+  - vllm/model_executor/models/qwen3_5_mtp.py
+  - vllm/transformers_utils/configs/qwen3_5.py
+  - vllm/transformers_utils/configs/qwen3_5_moe.py
+  - vllm/model_executor/models/qwen3_next.py
+  - vllm/model_executor/models/qwen3_next_mtp.py
+  - vllm/model_executor/layers/fla/ops/
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt
+
 - label: LM Eval Large Models (H200)
  timeout_in_minutes: 60
  device: h200
@@ -11,7 +11,7 @@ steps:
  - vllm/v1/attention/
  - tests/v1/engine/test_llm_engine.py
  - tests/v1/e2e/
-  - tests/v1/entrypoints/llm/test_struct_output_generate.py
+  - tests/entrypoints/llm/test_struct_output_generate.py
  commands:
  - set -x
  - export VLLM_USE_V2_MODEL_RUNNER=1
@@ -22,7 +22,7 @@ steps:
  - pytest -v -s v1/e2e/general/test_context_length.py
  - pytest -v -s v1/e2e/general/test_min_tokens.py
  # Temporary hack filter to exclude ngram spec decoding based tests.
-  - pytest -v -s v1/entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
+  - pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"

 - label: Model Runner V2 Examples
  timeout_in_minutes: 45
@@ -75,7 +75,7 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
 /tests/multimodal @DarkLight1337 @ywang96 @NickLucche
 /tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety
 /tests/test_inputs.py @DarkLight1337 @ywang96
-/tests/v1/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
+/tests/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
 /tests/v1/structured_output @mgoin @russellb @aarnphm
 /tests/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery
 /tests/weight_loading @mgoin @youkaichao @yewentao256
@@ -260,7 +260,7 @@ pull_request_rules:
      - files=examples/offline_inference/structured_outputs.py
      - files=examples/online_serving/structured_outputs/structured_outputs.py
      - files~=^tests/v1/structured_output/
-      - files=tests/v1/entrypoints/llm/test_struct_output_generate.py
+      - files=tests/entrypoints/llm/test_struct_output_generate.py
      - files~=^vllm/v1/structured_output/
  actions:
    label:
@@ -26,6 +26,16 @@
  #define __HIP__GFX9__
 #endif

+#if defined(__HIPCC__) &&                                                    \
+    (defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1150__) || \
+     defined(__gfx1151__) || defined(__gfx1200__) || defined(__gfx1201__))
+  #define __HIP__GFX1X__
+#endif
+
+#if defined(__HIPCC__) && (defined(__gfx1200__) || defined(__gfx1201__))
+  #define __HIP__GFX12__
+#endif
+
 #if defined(__HIPCC__) && (defined(__gfx942__) || defined(__gfx950__))
  #define __HIP__MI3XX__
 #endif
@@ -37,15 +47,31 @@
 #endif

 int get_lds_size() {
-  static bool is_cached = false;
-  static int result;
-  if (is_cached == false) {
-    auto dprops = at::cuda::getCurrentDeviceProperties();
-    std::string device_arch = dprops->gcnArchName;
-    size_t substring = device_arch.find("gfx95");
-    result = (substring == std::string::npos ? 64 * 1024 : 160 * 1024);
-    is_cached = true;
-  }
+  static const int result = [] {
+    const auto* dprops = at::cuda::getCurrentDeviceProperties();
+    const std::string device_arch = dprops->gcnArchName;
+    return device_arch.find("gfx95") == std::string::npos ? 64 * 1024
+                                                          : 160 * 1024;
+  }();
+  return result;
+}
+
+bool on_gfx1x() {
+  static const bool result = [] {
+    const auto* dprops = at::cuda::getCurrentDeviceProperties();
+    const std::string device_arch = dprops->gcnArchName;
+    return device_arch.find("gfx11") != std::string::npos ||
+           device_arch.find("gfx12") != std::string::npos;
+  }();
+  return result;
+}
+
+bool on_gfx12() {
+  static const bool result = [] {
+    const auto* dprops = at::cuda::getCurrentDeviceProperties();
+    const std::string device_arch = dprops->gcnArchName;
+    return device_arch.find("gfx12") != std::string::npos;
+  }();
  return result;
 }

@@ -286,21 +312,35 @@ torch::Tensor LLMM1(at::Tensor& in_a, at::Tensor& in_b,
  return out_c;
 }

-#define DOT2C(V0, V2, V3)                                                     \
-  if constexpr (std::is_same_v<scalar_t, half>) {                             \
-    asm("v_dot2c_f32_f16 %0, %2, %3" : "=v"(V0) : "0"(V0), "v"(V2), "v"(V3)); \
-  } else if constexpr (std::is_same_v<scalar_t, __hip_bfloat16>) {            \
-    float2 s = __bfloat1622float2(*((__hip_bfloat162*)(&(V2)))) *             \
-               __bfloat1622float2(*((__hip_bfloat162*)(&(V3))));              \
-    V0 += (s.x + s.y);                                                        \
-  }
+#if defined(__HIP__GFX9__) && !defined(__HIP__GFX1X__)
+  #define DOT2C(V0, V2, V3)                                          \
+    if constexpr (std::is_same_v<scalar_t, half>) {                  \
+      asm("v_dot2c_f32_f16 %0, %2, %3"                               \
+          : "=v"(V0)                                                 \
+          : "0"(V0), "v"(V2), "v"(V3));                              \
+    } else if constexpr (std::is_same_v<scalar_t, __hip_bfloat16>) { \
+      float2 s = __bfloat1622float2(*((__hip_bfloat162*)(&(V2)))) *  \
+                 __bfloat1622float2(*((__hip_bfloat162*)(&(V3))));   \
+      V0 += (s.x + s.y);                                             \
+    }
+#elif defined(__HIP__GFX1X__)
+  // gfx1x: v_dot2_f32_f16 (VOP3-P, dot10-insts, available on gfx11+gfx12)
+  #define DOT2C(V0, V2, V3)                                               \
+    if constexpr (std::is_same_v<scalar_t, half>) {                       \
+      asm("v_dot2_f32_f16 %0, %1, %2, %0" : "+v"(V0) : "v"(V2), "v"(V3)); \
+    } else if constexpr (std::is_same_v<scalar_t, __hip_bfloat16>) {      \
+      float2 s = __bfloat1622float2(*((__hip_bfloat162*)(&(V2)))) *       \
+                 __bfloat1622float2(*((__hip_bfloat162*)(&(V3))));        \
+      V0 += (s.x + s.y);                                                  \
+    }
+#endif

 // To avoid LLVM silently upcasting to double
 __device__ inline unsigned int min__(uint32_t a, uint32_t b) {
  return min(a, b);
 }

-#if defined(__HIP__GFX9__)  // TODO: Add NAVI support
+#if defined(__HIP__GFX9__) || defined(__HIP__GFX1X__)
 // This version targets cases where A[] fits LDS capacity
 template <typename scalar_t, int THRDS, int YTILE, int WvPrGrp, int A_CHUNK,
          int UNRL, int N>
@@ -442,14 +482,18 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
                                                1);  // row_shr2
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x111, 0xf, 0xf,
                                                1);  // row_shr1
+  #if defined(__HIP__GFX9__)
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x142, 0xf, 0xf,
                                                1);  // ROW_BCAST15
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x143, 0xf, 0xf,
                                                1);  // ROW_BCAST31
+  #else
+          sum[n][y] += __shfl_xor(sum[n][y], 16);
+  #endif
        }
      }

-      if (threadIdx.x == 63) {
+      if (threadIdx.x == (THRDS - 1)) {
        scalar_t biases[N][YTILE] = {};
        if (BIAS)
          for (int n = 0; n < N; n++) {
@@ -469,9 +513,10 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
        }
      }
    } else {
-  #pragma unroll
+  #ifdef __HIP__GFX9__
+    #pragma unroll
      for (int n = 0; n < N; n++) {
-  #pragma unroll
+    #pragma unroll
        for (int y = 0; y < YTILE; y++) {
          /*float accm1 = 0;
           for (int i=0; i<64; i++)
@@ -498,7 +543,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
          sum4[n][y][0] = accm;
        }
      }
-      if (threadIdx.x == 63) {
+      if (threadIdx.x == (THRDS - 1)) {
        scalar_t biases[N][YTILE] = {};
        if (BIAS)
          for (int n = 0; n < N; n++) {
@@ -513,11 +558,12 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
          }
        }
      }
+  #endif  // __HIP__GFX9__ (MFMA path)
    }
    m += CuCount * _WvPrGrp * YTILE;
  }
 }
-#else   // !defined(__HIP__GFX9__) TODO: Add NAVI support
+#else
 template <typename scalar_t, int THRDS, int YTILE, int WvPrGrp, int A_CHUNK,
          int UNRL, int N>
 __global__ void wvSplitK_hf_sml_(const int K, const int Kbp, const int Kap,
@@ -528,9 +574,9 @@ __global__ void wvSplitK_hf_sml_(const int K, const int Kbp, const int Kap,
                                 const int _WvPrGrp, const int CuCount) {
  UNREACHABLE_CODE
 }
-#endif  // defined(__HIP__GFX9__) TODO: Add NAVI support
+#endif

-#if defined(__HIP__GFX9__)  // TODO: Add NAVI support
+#if defined(__HIP__GFX9__) || defined(__HIP__GFX1X__)
 // This version targets cases where A[] marginally exceeds LDS capacity
 template <typename scalar_t, int THRDS, int YTILE, int WvPrGrp, int A_CHUNK,
          int UNRL, int N>
@@ -657,14 +703,18 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
                                                1);  // row_shr2
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x111, 0xf, 0xf,
                                                1);  // row_shr1
+  #if defined(__HIP__GFX9__)
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x142, 0xf, 0xf,
                                                1);  // ROW_BCAST15
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x143, 0xf, 0xf,
                                                1);  // ROW_BCAST31
+  #else
+          sum[n][y] += __shfl_xor(sum[n][y], 16);
+  #endif
        }
      }

-      if (threadIdx.x == 63) {
+      if (threadIdx.x == (THRDS - 1)) {
        scalar_t biases[N][YTILE] = {};
        if (BIAS)
          for (int n = 0; n < N; n++) {
@@ -686,9 +736,10 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
        }
      }
    } else {
-  #pragma unroll
+  #ifdef __HIP__GFX9__
+    #pragma unroll
      for (int n = 0; n < N; n++) {
-  #pragma unroll
+    #pragma unroll
        for (int y = 0; y < YTILE; y++) {
          // float accm1 = 0;
          // for (int i=0; i<64; i++)
@@ -713,7 +764,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
          sum4[n][y][0] = accm;
        }
      }
-      if (threadIdx.x == 63) {
+      if (threadIdx.x == (THRDS - 1)) {
        scalar_t biases[N][YTILE] = {};
        if (BIAS)
          for (int n = 0; n < N; n++) {
@@ -730,6 +781,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
          }
        }
      }
+  #endif  // __HIP__GFX9__ (MFMA path)
    }

    m += CuCount * _WvPrGrp * YTILE;
@@ -746,7 +798,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
  }
 }

-#else   // !defined(__HIP__GFX9__) TODO: Add NAVI support
+#else
 template <typename scalar_t, int THRDS, int YTILE, int WvPrGrp, int A_CHUNK,
          int UNRL, int N>
 __global__ void wvSplitK_hf_(const int K, const int Kbp, const int Kap,
@@ -756,9 +808,9 @@ __global__ void wvSplitK_hf_(const int K, const int Kbp, const int Kap,
                             const int _WvPrGrp, const int CuCount) {
  UNREACHABLE_CODE
 }
-#endif  // defined(__HIP__GFX9__) TODO: Add NAVI support
+#endif

-#if defined(__HIP__GFX9__)  // TODO: Add NAVI support
+#if defined(__HIP__GFX9__) || defined(__HIP__GFX1X__)
 // This version targets big A[] cases, where it is much larger than LDS capacity
 template <typename scalar_t, int THRDS, int YTILE, int WvPrGrp, int A_CHUNK,
          int UNRL, int N>
@@ -1004,14 +1056,18 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
                                                1);  // row_shr2
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x111, 0xf, 0xf,
                                                1);  // row_shr1
+  #if defined(__HIP__GFX9__)
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x142, 0xf, 0xf,
                                                1);  // ROW_BCAST15
          sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x143, 0xf, 0xf,
                                                1);  // ROW_BCAST31
+  #else
+          sum[n][y] += __shfl_xor(sum[n][y], 16);
+  #endif
        }
      }

-      if (threadIdx.x == 63) {
+      if (threadIdx.x == (THRDS - 1)) {
        scalar_t biases[N][YTILE] = {};
        if (BIAS)
          for (int n = 0; n < N; n++) {
@@ -1033,9 +1089,10 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
        }
      }
    } else {
-  #pragma unroll
+  #ifdef __HIP__GFX9__
+    #pragma unroll
      for (int n = 0; n < N; n++) {
-  #pragma unroll
+    #pragma unroll
        for (int y = 0; y < YTILE; y++) {
          float accm = sum4[n][y][0];
          accm += __builtin_amdgcn_mov_dpp(sum4[n][y][1], 0x101, 0xf, 0xf,
@@ -1057,7 +1114,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
          sum4[n][y][0] = accm;
        }
      }
-      if (threadIdx.x == 63) {
+      if (threadIdx.x == (THRDS - 1)) {
        scalar_t biases[N][YTILE] = {};
        if (BIAS)
          for (int n = 0; n < N; n++) {
@@ -1074,6 +1131,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
          }
        }
      }
+  #endif  // __HIP__GFX9__ (MFMA path)
    }

    m += CuCount * _WvPrGrp * YTILE;
@@ -1090,7 +1148,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
    }
  }
 }
-#else   // !defined(__HIP__GFX9__) TODO: Add NAVI support
+#else
 template <typename scalar_t, int THRDS, int YTILE, int WvPrGrp, int A_CHUNK,
          int UNRL, int N>
 __global__ void wvSplitK_hf_big_(const int K, const int Kbp, const int Kap,
@@ -1101,7 +1159,7 @@ __global__ void wvSplitK_hf_big_(const int K, const int Kbp, const int Kap,
                                 const int _WvPrGrp, const int CuCount) {
  UNREACHABLE_CODE
 }
-#endif  // defined(__HIP__GFX9__) TODO: Add NAVI support
+#endif

 // Find the min val of div2 that doesn't increase N/(div1*div2)
 int mindiv(int N, int div1, int div2) {
@@ -1148,40 +1206,40 @@ torch::Tensor wvSplitK(const at::Tensor& in_a, const at::Tensor& in_b,
  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  const int max_lds_len = get_lds_size() / 2;

-#define WVSPLITK(_YTILE, _UNRL, _N)                                           \
+#define WVSPLITK_CFG(_THRDS, _WVPRGRP, _YTILE, _UNRL, _N)                     \
  {                                                                           \
-    dim3 block(64, 16);                                                       \
-    int __wvPrGrp = mindiv(M_in, CuCount * _YTILE, 16);                       \
+    dim3 block(_THRDS, _WVPRGRP);                                             \
+    int __wvPrGrp = mindiv(M_in, CuCount * _YTILE, _WVPRGRP);                 \
    if ((Kbp_in * N_in <= max_lds_len) && (M_in % _YTILE == 0))               \
-      wvSplitK_hf_sml_<fptype, 64, _YTILE, 16, 8, _UNRL, _N>                  \
+      wvSplitK_hf_sml_<fptype, _THRDS, _YTILE, _WVPRGRP, 8, _UNRL, _N>        \
          <<<grid, block, 0, stream>>>(K_in, Kap_in, Kbp_in, M_in, Bx_in,     \
                                       By_in, af4, bf4, biasf4, c, __wvPrGrp, \
                                       CuCount);                              \
    else if (Kbp_in * N_in <= max_lds_len * 1.2)                              \
-      wvSplitK_hf_<fptype, 64, _YTILE, 16, 8, _UNRL, _N>                      \
+      wvSplitK_hf_<fptype, _THRDS, _YTILE, _WVPRGRP, 8, _UNRL, _N>            \
          <<<grid, block, 0, stream>>>(K_in, Kap_in, Kbp_in, M_in, Bx_in,     \
                                       By_in, af4, bf4, biasf4, c, __wvPrGrp, \
                                       CuCount);                              \
    else                                                                      \
-      wvSplitK_hf_big_<fptype, 64, _YTILE, 16, 8, _UNRL, _N>                  \
+      wvSplitK_hf_big_<fptype, _THRDS, _YTILE, _WVPRGRP, 8, _UNRL, _N>        \
          <<<grid, block, 0, stream>>>(K_in, Kap_in, Kbp_in, M_in, Bx_in,     \
                                       By_in, af4, bf4, biasf4, c, __wvPrGrp, \
                                       CuCount);                              \
  }

-#define WVSPLIT_TILE(_sYT, __N)                           \
+#define WVSPLIT_TILE_CFG(_THRDS, _WVPRGRP, _sYT, __N)     \
  {                                                       \
    bool fit_lds = (Kbp_in * N_in <= max_lds_len);        \
    if (_sYT <= 1)                                        \
-      WVSPLITK(1, 4, __N)                                 \
+      WVSPLITK_CFG(_THRDS, _WVPRGRP, 1, 4, __N)           \
    else if ((__N == 1) || (!fit_lds) || (_sYT <= 4 * 2)) \
-      WVSPLITK(2, 2, __N)                                 \
+      WVSPLITK_CFG(_THRDS, _WVPRGRP, 2, 2, __N)           \
    else if (_sYT <= 4 * 3)                               \
-      WVSPLITK(3, 2, __N)                                 \
+      WVSPLITK_CFG(_THRDS, _WVPRGRP, 3, 2, __N)           \
    else if (__N == 4)                                    \
-      WVSPLITK(4, 1, __N)                                 \
+      WVSPLITK_CFG(_THRDS, _WVPRGRP, 4, 1, __N)           \
    else                                                  \
-      WVSPLITK(4, 2, __N)                                 \
+      WVSPLITK_CFG(_THRDS, _WVPRGRP, 4, 2, __N)           \
  }

  AT_DISPATCH_REDUCED_FLOATING_TYPES(in_b.scalar_type(), "wvSplitK", [&] {
@@ -1198,18 +1256,31 @@ torch::Tensor wvSplitK(const at::Tensor& in_a, const at::Tensor& in_b,
    // then cut the active waves to balance their distribution...
    int sYT = (M_in + CuCount * 4 - 1) / (CuCount * 4);

+    const bool use_wave32 = on_gfx1x();
    switch (N_in) {
      case 1:
-        WVSPLIT_TILE(sYT, 1)
+        if (use_wave32)
+          WVSPLIT_TILE_CFG(32, 16, sYT, 1)
+        else
+          WVSPLIT_TILE_CFG(64, 16, sYT, 1)
        break;
      case 2:
-        WVSPLIT_TILE(sYT, 2)
+        if (use_wave32)
+          WVSPLIT_TILE_CFG(32, 16, sYT, 2)
+        else
+          WVSPLIT_TILE_CFG(64, 16, sYT, 2)
        break;
      case 3:
-        WVSPLIT_TILE(sYT, 3)
+        if (use_wave32)
+          WVSPLIT_TILE_CFG(32, 16, sYT, 3)
+        else
+          WVSPLIT_TILE_CFG(64, 16, sYT, 3)
        break;
      case 4:
-        WVSPLIT_TILE(sYT, 4)
+        if (use_wave32)
+          WVSPLIT_TILE_CFG(32, 16, sYT, 4)
+        else
+          WVSPLIT_TILE_CFG(64, 16, sYT, 4)
        break;
      default:
        throw std::runtime_error(
@@ -1653,7 +1724,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
  #endif
  }
 }
-#else   // !defined(__HIP__GFX9__) TODO: Add NAVI support
+#else
 template <typename scalar_t, int THRDS, int YTILE, int WvPrGrp, int A_CHUNK,
          int UNRL, int N, int GrpsShrB, int CHUNKK, int DTRMNSTC>
 __global__ void wvSplitKrc_(const int actlN, const int K, const int Kap,
@@ -1688,6 +1759,8 @@ torch::Tensor wvSplitKrc(const at::Tensor& in_a, const at::Tensor& in_b,
  TORCH_CHECK(in_a.dtype() == torch::kFloat16 ||
              in_a.dtype() == torch::kBFloat16);

+  const at::cuda::OptionalCUDAGuard device_guard(device_of(in_a));
+
  auto out_c = torch::empty(
      {N_in, M_in},
      torch::TensorOptions().dtype(in_a.dtype()).device(in_a.device()));
@@ -1696,7 +1769,6 @@ torch::Tensor wvSplitKrc(const at::Tensor& in_a, const at::Tensor& in_b,

  dim3 grid(CuCount);

-  const at::cuda::OptionalCUDAGuard device_guard(device_of(in_a));
  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  // const int max_lds_len = get_lds_size() / 2;

@@ -1773,7 +1845,7 @@ torch::Tensor wvSplitKrc(const at::Tensor& in_a, const at::Tensor& in_b,
  return out_c;
 }

-#if defined(__HIP__MI3XX__)  // TODO: Add NAVI support
+#if defined(__HIP__MI3XX__) || defined(__HIP__GFX12__)
 template <typename scalar_t, typename fp8_t, int THRDS, int YTILE, int WvPrGrp,
          int A_CHUNK, int UNRL, int N>
 __global__ void __launch_bounds__(WvPrGrp* THRDS)
@@ -1817,12 +1889,17 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)

  uint32_t m = (blockIdx.x * _WvPrGrp + (threadIdx.y % _WvPrGrp)) * YTILE;

-  using floatx16 = __attribute__((__vector_size__(16 * sizeof(float)))) float;
  float sA = *s_A;
  float sB = *s_B;

  while (m < M) {
+  #ifdef __HIP__GFX12__
+    // gfx12: per-lane scalar accumulation via v_dot4_f32_fp8_fp8
+    float sum[N][YTILE] = {};
+  #else
+    // gfx9: MFMA accumulation
    scalar8 sum[N][YTILE] = {};
+  #endif
    for (uint32_t k1 = 0; k1 < K; k1 += THRDS * A_CHUNK * UNRL) {
      bigType bigA[N][UNRL] = {};
      bigType bigB[YTILE][UNRL];
@@ -1854,6 +1931,17 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
  #pragma unroll
      for (uint32_t k2 = 0; k2 < UNRL; k2++) {
        for (uint32_t n = 0; n < N; n++) {
+  #ifdef __HIP__GFX12__
+          // gfx12: 4 x dot4 per A_CHUNK=16 bytes (4 FP8 per dot4)
+          for (int y = 0; y < YTILE; ++y) {
+    #pragma unroll
+            for (int i = 0; i < A_CHUNK / 4; i++) {
+              sum[n][y] = __builtin_amdgcn_dot4_f32_fp8_fp8(
+                  bigA[n][k2].i[i], bigB[y][k2].i[i], sum[n][y]);
+            }
+          }
+  #else
+          // gfx9: MFMA path
          for (int i = 0; i < A_CHUNK; i += 8) {
            for (int y = 0; y < YTILE; ++y) {
              sum[n][y] = __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(
@@ -1861,11 +1949,33 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
                  0);
            }
          }
+  #endif
        }
      }
    }

    // Final reduction
+  #ifdef __HIP__GFX12__
+    // gfx12 wave32: DPP row_shr within 16-lane rows + cross-row shuffle
+    for (int n = 0; n < N; n++) {
+      for (int y = 0; y < YTILE; y++) {
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:8 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:4 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:2 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:1 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        sum[n][y] += __shfl_xor(sum[n][y], 16);
+      }
+    }
+  #else
+    // gfx9 MFMA reduction
    for (int n = 0; n < N; n++) {
      for (int y = 0; y < YTILE; y++) {
        float accm0 = sum[n][y][0];
@@ -1880,8 +1990,15 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
        sum[n][y][0] = accm0;
      }
    }
+  #endif

-    if (threadIdx.x == 0) {
+    const bool writeback_lane =
+  #ifdef __HIP__GFX12__
+        threadIdx.x == (THRDS - 1);
+  #else
+        threadIdx.x == 0;
+  #endif
+    if (writeback_lane) {
      scalar_t biases[N][YTILE] = {};
      if (BIAS)
        for (int n = 0; n < N; n++) {
@@ -1892,13 +2009,17 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
      for (int n = 0; n < N; n++) {
        for (int y = 0; y < YTILE; y++) {
          if (y + m >= M) break;  // To avoid mem access fault.
-          sum[n][y][0] *= sA * sB;
+  #ifdef __HIP__GFX12__
+          float result = sum[n][y] * sA * sB;
+  #else
+          float result = sum[n][y][0] * sA * sB;
+  #endif
          if constexpr (std::is_same_v<scalar_t, half>) {
-            sum[n][y][0] += __half2float(biases[n][y]);
+            result += __half2float(biases[n][y]);
          } else if constexpr (std::is_same_v<scalar_t, __hip_bfloat16>) {
-            sum[n][y][0] += __bfloat162float(biases[n][y]);
+            result += __bfloat162float(biases[n][y]);
          }
-          C[m + y + n * M] = __float2s<scalar_t>(sum[n][y][0]);
+          C[m + y + n * M] = __float2s<scalar_t>(result);
        }
      }
    }
@@ -1906,7 +2027,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
    m += CuCount * _WvPrGrp * YTILE;
  }
 }
-#else   // !defined(__HIP__MI3XX__) TODO: Add NAVI support
+#else   // !defined(__HIP__MI3XX__) && !defined(__HIP__GFX12__)
 template <typename scalar_t, typename fp8_t, int THRDS, int YTILE, int WvPrGrp,
          int A_CHUNK, int UNRL, int N>
 __global__ void wvSplitKQ_hf_sml_(const int K, const int Kap, const int Kbp,
@@ -1918,9 +2039,9 @@ __global__ void wvSplitKQ_hf_sml_(const int K, const int Kap, const int Kbp,
                                  const int _WvPrGrp, const int CuCount) {
  UNREACHABLE_CODE
 }
-#endif  // defined(__HIP__MI3XX__) TODO: Add NAVI support
+#endif  // defined(__HIP__MI3XX__) || defined(__HIP__GFX12__)

-#if defined(__HIP__MI3XX__)  // TODO: Add NAVI support
+#if defined(__HIP__MI3XX__) || defined(__HIP__GFX12__)
 template <typename scalar_t, typename fp8_t, int THRDS, int YTILE, int WvPrGrp,
          int A_CHUNK, int UNRL, int N>
 __global__ void __launch_bounds__(WvPrGrp* THRDS)
@@ -1963,12 +2084,17 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)

  uint32_t m = (blockIdx.x * _WvPrGrp + (threadIdx.y % _WvPrGrp)) * YTILE;

-  using floatx16 = __attribute__((__vector_size__(16 * sizeof(float)))) float;
  float sA = *s_A;
  float sB = *s_B;

  while (m < M) {
+  #ifdef __HIP__GFX12__
+    // gfx12: per-lane scalar accumulation via v_dot4_f32_fp8_fp8
+    float sum[N][YTILE] = {};
+  #else
+    // gfx9: MFMA accumulation
    scalar8 sum[N][YTILE] = {};
+  #endif
    for (uint32_t k1 = 0; k1 < K; k1 += THRDS * A_CHUNK * UNRL) {
      bigType bigA[N][UNRL] = {};
      bigType bigB[YTILE][UNRL];
@@ -2002,6 +2128,17 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
  #pragma unroll
      for (uint32_t k2 = 0; k2 < UNRL; k2++) {
        for (uint32_t n = 0; n < N; n++) {
+  #ifdef __HIP__GFX12__
+          // gfx12: 4 x dot4 per A_CHUNK=16 bytes (4 FP8 per dot4)
+          for (int y = 0; y < YTILE; ++y) {
+    #pragma unroll
+            for (int i = 0; i < A_CHUNK / 4; i++) {
+              sum[n][y] = __builtin_amdgcn_dot4_f32_fp8_fp8(
+                  bigA[n][k2].i[i], bigB[y][k2].i[i], sum[n][y]);
+            }
+          }
+  #else
+          // gfx9: MFMA path
          for (int i = 0; i < A_CHUNK; i += 8) {
            for (int y = 0; y < YTILE; ++y) {
              sum[n][y] = __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(
@@ -2009,11 +2146,33 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
                  0);
            }
          }
+  #endif
        }
      }
    }

    // Final reduction
+  #ifdef __HIP__GFX12__
+    // gfx12 wave32: DPP row_shr within 16-lane rows + cross-row shuffle
+    for (int n = 0; n < N; n++) {
+      for (int y = 0; y < YTILE; y++) {
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:8 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:4 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:2 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        asm("s_nop 0\n\tv_add_f32 %0, %2, %3 row_shr:1 bound_ctrl:0 "
+            : "=v"(sum[n][y])
+            : "0"(sum[n][y]), "v"(sum[n][y]), "v"(sum[n][y]));
+        sum[n][y] += __shfl_xor(sum[n][y], 16);
+      }
+    }
+  #else
+    // gfx9 MFMA reduction
    for (int n = 0; n < N; n++) {
      for (int y = 0; y < YTILE; y++) {
        float accm0 = sum[n][y][0];
@@ -2028,8 +2187,15 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
        sum[n][y][0] = accm0;
      }
    }
+  #endif

-    if (threadIdx.x == 0) {
+    const bool writeback_lane =
+  #ifdef __HIP__GFX12__
+        threadIdx.x == (THRDS - 1);
+  #else
+        threadIdx.x == 0;
+  #endif
+    if (writeback_lane) {
      scalar_t biases[N][YTILE] = {};
      if (BIAS)
        for (int n = 0; n < N; n++) {
@@ -2040,13 +2206,17 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
      for (int n = 0; n < N; n++) {
        for (int y = 0; y < YTILE; y++) {
          if (y + m >= M) break;  // To avoid mem access fault.
-          sum[n][y][0] *= sA * sB;
+  #ifdef __HIP__GFX12__
+          float result = sum[n][y] * sA * sB;
+  #else
+          float result = sum[n][y][0] * sA * sB;
+  #endif
          if constexpr (std::is_same_v<scalar_t, half>) {
-            sum[n][y][0] += __half2float(biases[n][y]);
+            result += __half2float(biases[n][y]);
          } else if constexpr (std::is_same_v<scalar_t, __hip_bfloat16>) {
-            sum[n][y][0] += __bfloat162float(biases[n][y]);
+            result += __bfloat162float(biases[n][y]);
          }
-          C[m + y + n * M] = __float2s<scalar_t>(sum[n][y][0]);
+          C[m + y + n * M] = __float2s<scalar_t>(result);
        }
      }
    }
@@ -2054,7 +2224,7 @@ __global__ void __launch_bounds__(WvPrGrp* THRDS)
    m += CuCount * _WvPrGrp * YTILE;
  }
 }
-#else   // !defined(__HIP__MI3XX__) TODO: Add NAVI support
+#else   // !defined(__HIP__MI3XX__) && !defined(__HIP__GFX12__)
 template <typename scalar_t, typename fp8_t, int THRDS, int YTILE, int WvPrGrp,
          int A_CHUNK, int UNRL, int N>
 __global__ void wvSplitKQ_hf_(const int K, const int Kap, const int Kbp,
@@ -2066,7 +2236,7 @@ __global__ void wvSplitKQ_hf_(const int K, const int Kap, const int Kbp,
                              const int CuCount) {
  UNREACHABLE_CODE
 }
-#endif  // defined(__HIP__MI3XX__) TODO: Add NAVI support
+#endif  // defined(__HIP__MI3XX__) || defined(__HIP__GFX12__)

 void wvSplitKQ(const at::Tensor& in_b, const at::Tensor& in_a,
               const std::optional<at::Tensor>& in_bias, at::Tensor& out_c,
@@ -2099,24 +2269,30 @@ void wvSplitKQ(const at::Tensor& in_b, const at::Tensor& in_a,
  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  const int max_lds_len = get_lds_size();

-#define WVSPLITKQ(_WvPrGrp, _YTILEs, _YTILEm, _UNRLs, _UNRLm, _N)             \
-  {                                                                           \
-    dim3 block(64, _WvPrGrp);                                                 \
-    if ((Kap_in * N_in <= max_lds_len) && (M_in % _YTILEs == 0)) {            \
-      int __wvPrGrp = min(_WvPrGrp, mindiv(M_in, CuCount * _YTILEs, 16));     \
-      wvSplitKQ_hf_sml_<fptype, fp8_t, 64, _YTILEs, _WvPrGrp, 16, _UNRLs, _N> \
-          <<<grid, block, 0, stream>>>(K_in, Kap_in, Kbp_in, M_in, Bx_in,     \
-                                       By_in, b_ptr, a_ptr, bias_ptr, c_ptr,  \
-                                       s_a, s_b, __wvPrGrp, CuCount);         \
-    } else {                                                                  \
-      int __wvPrGrp = min(_WvPrGrp, mindiv(M_in, CuCount * _YTILEm, 16));     \
-      wvSplitKQ_hf_<fptype, fp8_t, 64, _YTILEm, _WvPrGrp, 16, _UNRLm, _N>     \
-          <<<grid, block, 0, stream>>>(K_in, Kap_in, Kbp_in, M_in, Bx_in,     \
-                                       By_in, b_ptr, a_ptr, bias_ptr, c_ptr,  \
-                                       s_a, s_b, __wvPrGrp, CuCount);         \
-    }                                                                         \
+#define WVSPLITKQ_IMPL(_THRDS, _WvPrGrp, _YTILEs, _YTILEm, _UNRLs, _UNRLm, _N) \
+  {                                                                            \
+    dim3 block(_THRDS, _WvPrGrp);                                              \
+    if ((Kap_in * N_in <= max_lds_len) && (M_in % _YTILEs == 0)) {             \
+      int __wvPrGrp = min(_WvPrGrp, mindiv(M_in, CuCount * _YTILEs, 16));      \
+      wvSplitKQ_hf_sml_<fptype, fp8_t, _THRDS, _YTILEs, _WvPrGrp, 16, _UNRLs,  \
+                        _N><<<grid, block, 0, stream>>>(                       \
+          K_in, Kap_in, Kbp_in, M_in, Bx_in, By_in, b_ptr, a_ptr, bias_ptr,    \
+          c_ptr, s_a, s_b, __wvPrGrp, CuCount);                                \
+    } else {                                                                   \
+      int __wvPrGrp = min(_WvPrGrp, mindiv(M_in, CuCount * _YTILEm, 16));      \
+      wvSplitKQ_hf_<fptype, fp8_t, _THRDS, _YTILEm, _WvPrGrp, 16, _UNRLm, _N>  \
+          <<<grid, block, 0, stream>>>(K_in, Kap_in, Kbp_in, M_in, Bx_in,      \
+                                       By_in, b_ptr, a_ptr, bias_ptr, c_ptr,   \
+                                       s_a, s_b, __wvPrGrp, CuCount);          \
+    }                                                                          \
  }

+#define WVSPLITKQ(_WvPrGrp, _YTILEs, _YTILEm, _UNRLs, _UNRLm, _N)      \
+  if (on_gfx12())                                                      \
+    WVSPLITKQ_IMPL(32, _WvPrGrp, _YTILEs, _YTILEm, _UNRLs, _UNRLm, _N) \
+  else                                                                 \
+    WVSPLITKQ_IMPL(64, _WvPrGrp, _YTILEs, _YTILEm, _UNRLs, _UNRLm, _N)
+
  AT_DISPATCH_REDUCED_FLOATING_TYPES(out_c.scalar_type(), "wvSplitKQ", [&] {
    using fptype = typename scalar<scalar_t>::type;
    auto c_ptr = reinterpret_cast<fptype*>(out_c.data_ptr());
@@ -2136,10 +2312,10 @@ void wvSplitKQ(const at::Tensor& in_b, const at::Tensor& in_a,
          WVSPLITKQ(16, 2, 2, 2, 2, 2)
          break;
        case 3:
-          WVSPLITKQ(16, 2, 2, 2, 2, 3)
+          WVSPLITKQ(16, 2, 2, 1, 1, 3)
          break;
        case 4:
-          WVSPLITKQ(16, 2, 2, 2, 2, 4)
+          WVSPLITKQ(16, 2, 2, 1, 1, 4)
          break;
        default:
          throw std::runtime_error(
@@ -44,7 +44,7 @@ ENV DEBIAN_FRONTEND=noninteractive

 # Install Python and other dependencies
 RUN apt-get update -y \
-    && apt-get install -y software-properties-common git curl sudo vim less libgfortran5 libopenmpi-dev libpci-dev \
+    && apt-get install -y software-properties-common git curl sudo vim less libgfortran5 libopenmpi-dev libpci-dev liblzma-dev pkg-config \
    && for i in 1 2 3; do \
        add-apt-repository -y ppa:deadsnakes/ppa && break || \
        { echo "Attempt $i failed, retrying in 5s..."; sleep 5; }; \
@@ -3,6 +3,10 @@
 !!! warning
    Profiling is only intended for vLLM developers and maintainers to understand the proportion of time spent in different parts of the codebase. **vLLM end-users should never turn on profiling** as it will significantly slow down the inference.

+!!! tip "Choosing a profiler"
+    - Use **Nsight Systems** for low-overhead, performance-critical profiling.
+    - Use **PyTorch Profiler** for medium-overhead profiling with richer debugging information (e.g., stack traces, memory, shapes). Note that enabling these features adds overhead and is not recommended for benchmarking.
+
 ## Profile with PyTorch Profiler

 We support tracing vLLM workers using different profilers. You can enable profiling by setting the `--profiler-config` flag when launching the server.
@@ -88,8 +88,8 @@ To be used with a particular `FusedMoEPrepareAndFinalizeModular` subclass, MoE k
 | flashinfer | standard | nvfp4,</br>fp8 | T | <sup>5</sup> | N | Y | [`FlashInferExperts`][vllm.model_executor.layers.fused_moe.flashinfer_cutlass_moe.FlashInferExperts] |
 | gpt oss triton | standard | N/A | N/A | <sup>5</sup> | Y | Y | [`triton_kernel_fused_experts`][vllm.model_executor.layers.fused_moe.gpt_oss_triton_kernels_moe.triton_kernel_fused_experts],</br>[`OAITritonExperts`][vllm.model_executor.layers.fused_moe.gpt_oss_triton_kernels_moe.OAITritonExperts] |
 | marlin | standard,</br>batched | <sup>3</sup> / N/A | <sup>3</sup> / N/A | silu,</br>swigluoai | Y | Y | [`fused_marlin_moe`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.fused_marlin_moe],</br>[`MarlinExperts`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.MarlinExperts],</br>[`BatchedMarlinExperts`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.BatchedMarlinExperts] |
-| trtllm | standard | mxfp4,</br>nvfp4 | G(16),G(32) | <sup>5</sup> | N | Y | [`TrtLlmGenExperts`][vllm.model_executor.layers.fused_moe.trtllm_moe.TrtLlmGenExperts] |
-| rocm aiter moe | standard | fp8 | G(128),A,T | silu, gelu | Y | N | [`rocm_aiter_fused_experts`][vllm.model_executor.layers.fused_moe.rocm_aiter_fused_moe.rocm_aiter_fused_experts] |
+| trtllm | standard | mxfp4,</br>nvfp4 | G(16),G(32) | <sup>5</sup> | N | Y | [`TrtLlmMxfp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsMonolithic],</br>[`TrtLlmMxfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsModular],</br>[`TrtLlmNvFp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsMonolithic],</br>[`TrtLlmNvfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsModular] |
+| rocm aiter moe | standard | mxfp4,</br>fp8 | G(32),G(128),A,T | silu, gelu,</br>swigluoai | Y | N | `rocm_aiter_fused_experts`,</br>`AiterExperts` |
 | cpu_fused_moe | standard | N/A | N/A | silu | N | N | [`CPUFusedMOE`][vllm.model_executor.layers.fused_moe.cpu_fused_moe.CPUFusedMOE] |
 | naive batched<sup>4</sup> | batched | int8,</br>fp8 | G,A,T | silu, gelu | <sup>6</sup> | Y | [`NaiveBatchedExperts`][vllm.model_executor.layers.fused_moe.fused_batched_moe.NaiveBatchedExperts] |

@@ -31,28 +31,29 @@ Of course, we also have "plugin" tasks that allow users to customize input and o

 ### Pooling Tasks

-| Pooling Tasks      | Granularity   | Outputs                                         |
-|--------------------|---------------|-------------------------------------------------|
-| `classify`         | Sequence-wise | probability vector of classes for each sequence |
-| `score` (see note) | Sequence-wise | reranker score for each sequence                |
-| `embed`            | Sequence-wise | vector representations for each sequence        |
-| `token_classify`   | Token-wise    | probability vector of classes for each token    |
-| `token_embed`      | Token-wise    | vector representations for each token           |
+| Pooling Tasks         | Granularity   | Outputs                                         |
+|-----------------------|---------------|-------------------------------------------------|
+| `classify` (see note) | Sequence-wise | probability vector of classes for each sequence |
+| `embed`               | Sequence-wise | vector representations for each sequence        |
+| `token_classify`      | Token-wise    | probability vector of classes for each token    |
+| `token_embed`         | Token-wise    | vector representations for each token           |

 !!! note
    Within classification tasks, there is a specialized subcategory: Cross-encoder (aka reranker) models. These models are a subset of classification models that accept two prompts as input and output num_labels equal to 1.

 ### Score Types

-| Pooling Tasks      | Granularity   | Outputs                                         | Score Types        | scoring function         |
-|--------------------|---------------|-------------------------------------------------|--------------------|--------------------------|
-| `classify`         | Sequence-wise | probability vector of classes for each sequence | nan                | nan                      |
-| `score` (see note) | Sequence-wise | reranker score for each sequence                | `cross-encoder`    | linear classifier        |
-| `embed`            | Sequence-wise | vector representations for each sequence        | `bi-encoder`       | cosine similarity        |
-| `token_classify`   | Token-wise    | probability vector of classes for each token    | nan                | nan                      |
-| `token_embed`      | Token-wise    | vector representations for each token           | `late-interaction` | late interaction(MaxSim) |
+The scoring models is designed to compute similarity scores between two input prompts. It supports three model types (aka `score_type`): `cross-encoder`, `late-interaction`, and `bi-encoder`.

-The score models is designed to compute similarity scores between two input prompts. It supports three model types (aka `score_type`): `cross-encoder`, `late-interaction`, and `bi-encoder`.
+| Pooling Tasks         | Granularity   | Outputs                                      | Score Types        | scoring function         |
+|-----------------------|---------------|----------------------------------------------|--------------------|--------------------------|
+| `classify` (see note) | Sequence-wise | reranker score for each sequence             | `cross-encoder`    | linear classifier        |
+| `embed`               | Sequence-wise | vector representations for each sequence     | `bi-encoder`       | cosine similarity        |
+| `token_classify`      | Token-wise    | probability vector of classes for each token | nan                | nan                      |
+| `token_embed`         | Token-wise    | vector representations for each token        | `late-interaction` | late interaction(MaxSim) |
+
+!!! note
+    Only when a classification model outputs num_labels equal to 1 can it be used as a scoring model and have its scoring API enabled.

 ### Pooling Usages

@@ -85,14 +86,16 @@ enabling the corresponding APIs.

 ### Offline APIs corresponding to pooling tasks

-| Task             | APIs                                                                       |
-|------------------|----------------------------------------------------------------------------|
-| `embed`          | `LLM.embed(...)`,`LLM.encode(..., pooling_task="embed")`, `LLM.score(...)` |
-| `classify`       | `LLM.classify(...)`, `LLM.encode(..., pooling_task="classify")`            |
-| `score`          | `LLM.score(...)`                                                           |
-| `token_classify` | `LLM.reward(...)`, `LLM.encode(..., pooling_task="token_classify")`        |
-| `token_embed`    | `LLM.encode(..., pooling_task="token_embed")`, `LLM.score(...)`            |
-| `plugin`         | `LLM.encode(..., pooling_task="plugin")`                                   |
+| Task             | APIs                                                                                  |
+|------------------|---------------------------------------------------------------------------------------|
+| `embed`          | `LLM.embed(...)`, `LLM.encode(..., pooling_task="embed")`, `LLM.score(...)`(see note) |
+| `classify`       | `LLM.classify(...)`, `LLM.encode(..., pooling_task="classify")`, `LLM.score(...)`     |
+| `token_classify` | `LLM.reward(...)`, `LLM.encode(..., pooling_task="token_classify")`                   |
+| `token_embed`    | `LLM.encode(..., pooling_task="token_embed")`, `LLM.score(...)`                       |
+| `plugin`         | `LLM.encode(..., pooling_task="plugin")`                                              |
+
+!!! note
+    Only when a classification model outputs num_labels equal to 1 can it be used as a scoring model and have its scoring API enabled.

 ### `LLM.classify`

@@ -206,11 +209,11 @@ If `--runner pooling` has been set (manually or automatically) but the model doe
 vLLM will attempt to automatically convert the model according to the architecture names
 shown in the table below.

-| Architecture                                    | `--convert` | Supported pooling tasks               |
-| ----------------------------------------------- | ----------- | ------------------------------------- |
-| `*ForTextEncoding`, `*EmbeddingModel`, `*Model` | `embed`     | `token_embed`, `embed`                |
-| `*ForRewardModeling`, `*RewardModel`            | `embed`     | `token_embed`, `embed`                |
-| `*For*Classification`, `*ClassificationModel`   | `classify`  | `token_classify`, `classify`, `score` |
+| Architecture                                    | `--convert` | Supported pooling tasks      |
+|-------------------------------------------------|-------------|------------------------------|
+| `*ForTextEncoding`, `*EmbeddingModel`, `*Model` | `embed`     | `token_embed`, `embed`       |
+| `*ForRewardModeling`, `*RewardModel`            | `embed`     | `token_embed`, `embed`       |
+| `*For*Classification`, `*ClassificationModel`   | `classify`  | `token_classify`, `classify` |

 !!! tip
    You can explicitly set `--convert <type>` to specify how to convert the model.
@@ -251,3 +254,7 @@ Pooling models now default support all pooling, you can use it without any setti

 - Extracting hidden states prefers using `token_embed` task.
 - Named Entity Recognition (NER) and reward models prefers using `token_classify` task.
+
+### Score task
+
+`score` task is deprecated and will be removed in v0.20. Please use `classify` instead. Only when a classification model outputs num_labels equal to 1 can it be used as a scoring model and have its scoring API enabled.
@@ -17,6 +17,8 @@ The key distinction between (sequence) classification and token classification l

 Many classification models support both (sequence) classification and token classification. For further details on token classification, please refer to [this page](token_classify.md).

+Only when a classification model outputs num_labels equal to 1 can it be used as a scoring model and have its scoring API enabled, please refer to [this page](scoring.md).
+
 ## Typical Use Cases

 ### Classification
@@ -54,7 +56,7 @@ If your model is not in the above list, we will try to automatically convert the

 Cross-encoder (aka reranker) models are a subset of classification models that accept two prompts as input and output num_labels equal to 1. Most classification models can also be used as [cross-encoder models](scoring.md#cross-encoder-models). For more information on cross-encoder models, please refer to [this page](scoring.md).

--8<-- "docs/models/pooling_models/scoring.md:supported-score-models"
+--8<-- "docs/models/pooling_models/scoring.md:supported-cross-encoder-models"

 ### Reward Models

@@ -10,11 +10,11 @@ The score models is designed to compute similarity scores between two input prom
 - Model Usage: Scoring
 - Pooling Task:

-| Score Types        | Pooling Tasks | scoring function         |
-|--------------------|---------------|--------------------------|
-| `cross-encoder`    | `score`       | linear classifier        |
-| `late-interaction` | `token_embed` | late interaction(MaxSim) |
-| `bi-encoder`       | `embed`       | cosine similarity        |
+| Score Types        | Pooling Tasks         | scoring function         |
+|--------------------|-----------------------|--------------------------|
+| `cross-encoder`    | `classify` (see note) | linear classifier        |
+| `late-interaction` | `token_embed`         | late interaction(MaxSim) |
+| `bi-encoder`       | `embed`               | cosine similarity        |

 - Offline APIs:
    - `LLM.score`
@@ -22,13 +22,16 @@ The score models is designed to compute similarity scores between two input prom
    - [Score API](scoring.md#score-api) (`/score`)
    - [Rerank API](scoring.md#rerank-api) (`/rerank`, `/v1/rerank`, `/v2/rerank`)

+!!! note
+    Only when a classification model outputs num_labels equal to 1 can it be used as a scoring model and have its scoring API enabled.
+
 ## Supported Models

 ### Cross-encoder models

 [Cross-encoder](https://www.sbert.net/examples/applications/cross-encoder/README.html) (aka reranker) models are a subset of classification models that accept two prompts as input and output num_labels equal to 1.

--8<-- [start:supported-score-models]
+--8<-- [start:supported-cross-encoder-models]

 #### Text-only Models

@@ -99,7 +102,7 @@ The score models is designed to compute similarity scores between two input prom
    vllm serve Qwen/Qwen3-VL-Reranker-2B --hf_overrides '{"architectures": ["Qwen3VLForSequenceClassification"],"classifier_from_token": ["no", "yes"],"is_original_qwen3_reranker": true}'
    ```

--8<-- [end:supported-score-models]
+--8<-- [end:supported-cross-encoder-models]

 ### Late-interaction models

@@ -11,6 +11,7 @@ vLLM supports ColBERT models with multiple encoder backbones:
 | `HF_ColBERT` | BERT | `answerdotai/answerai-colbert-small-v1`, `colbert-ir/colbertv2.0` |
 | `ColBERTModernBertModel` | ModernBERT | `lightonai/GTE-ModernColBERT-v1` |
 | `ColBERTJinaRobertaModel` | Jina XLM-RoBERTa | `jinaai/jina-colbert-v2` |
+| `ColBERTLfm2Model` | LFM2 | `LiquidAI/LFM2-ColBERT-350M` |

 **BERT-based ColBERT** models work out of the box:

@@ -29,6 +30,10 @@ vllm serve lightonai/GTE-ModernColBERT-v1 \
 vllm serve jinaai/jina-colbert-v2 \
    --hf-overrides '{"architectures": ["ColBERTJinaRobertaModel"]}' \
    --trust-remote-code
+
+# LFM2 backbone
+vllm serve LiquidAI/LFM2-ColBERT-350M \
+    --hf-overrides '{"architectures": ["ColBERTLfm2Model"]}'
 ```

 Then you can use the rerank API:
@@ -39,6 +39,7 @@ Models of any architecture can be converted into embedding models using `--conve

 | Architecture | Models | Example HF Models | [LoRA](../../features/lora.md) | [PP](../../serving/parallelism_scaling.md) |
 | ------------ | ------ | ----------------- | -------------------- | ------------------------- |
+| `ColBERTLfm2Model` | LFM2 | `LiquidAI/LFM2-ColBERT-350M` | | |
 | `ColBERTModernBertModel` | ModernBERT | `lightonai/GTE-ModernColBERT-v1` | | |
 | `ColBERTJinaRobertaModel` | Jina XLM-RoBERTa | `jinaai/jina-colbert-v2` | | |
 | `HF_ColBERT` | BERT | `answerdotai/answerai-colbert-small-v1`, `colbert-ir/colbertv2.0` | | |
@@ -21,6 +21,7 @@ vocos # required for minicpmo_26 test
 peft>=0.15.0 # required for phi-4-mm test
 pqdm
 ray[cgraph,default]>=2.48.0 # Ray Compiled Graph, required by pipeline parallelism tests
+resampy # required for audio tests
 sentence-transformers>=5.2.0 # required for embedding tests
 soundfile # required for audio tests
 jiwer # required for audio tests
@@ -544,6 +544,7 @@ numba==0.61.2
    # via
    #   -r requirements/test.in
    #   librosa
+    #   resampy
 numpy==2.2.6
    # via
    #   -r requirements/test.in
@@ -584,6 +585,7 @@ numpy==2.2.6
    #   pyogrio
    #   pywavelets
    #   rasterio
+    #   resampy
    #   rioxarray
    #   rouge-score
    #   runai-model-streamer
@@ -995,6 +997,8 @@ requests==2.32.3
    #   tiktoken
    #   transformers
    #   wandb
+resampy==0.4.3
+    # via -r requirements/test.in
 responses==0.25.3
    # via genai-perf
 rfc3339-validator==0.1.4
@@ -987,11 +987,11 @@ setup(
        "instanttensor": ["instanttensor >= 0.1.5"],
        "runai": ["runai-model-streamer[s3,gcs,azure] >= 0.15.7"],
        "audio": [
-            "librosa",
+            "av",
+            "resampy",
            "scipy",
            "soundfile",
            "mistral_common[audio]",
-            "av",
        ],  # Required for audio processing
        "video": [],  # Kept for backwards compatibility
        "flashinfer": [],  # Kept for backwards compatibility
@@ -84,7 +84,10 @@ def run_e2e_fusion_test(monkeypatch, caplog_mp_spawn):

        # TODO: remove this after finishing migration from envs to model kwargs
        if model_name == "openai/gpt-oss-20b":
-            monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_MXFP4_MXFP8", "1")
+            from .common import is_blackwell
+
+            if is_blackwell():
+                monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_MXFP4_MXFP8", "1")

        # Disable, compile cache to make sure custom passes run.
        # Otherwise, we can't verify fusion happened through the logs.
@@ -14,6 +14,7 @@ from unittest.mock import Mock, patch
 import pytest
 import torch

+import vllm.envs as envs
 import vllm.model_executor.layers.activation
 from vllm.compilation.backends import VllmBackend
 from vllm.compilation.caching import (
@@ -162,6 +163,9 @@ def test_save_and_load(monkeypatch: pytest.MonkeyPatch):

@pytest.mark.skipif(not is_torch_equal_or_newer("2.10.0"), reason="requires torch 2.10")
 def test_save_and_load_slice(monkeypatch: pytest.MonkeyPatch):
+    from torch._subclasses import FakeTensorMode
+    from torch.fx.experimental.symbolic_shapes import ShapeEnv
+
    def foo(x: torch.Tensor):
        return x[slice(0, x.shape[0])]

@@ -172,12 +176,13 @@ def test_save_and_load_slice(monkeypatch: pytest.MonkeyPatch):
    gm = torch.fx.symbolic_trace(foo)
    assert "getitem_1 = x[slice(0, getitem, None)]" in gm.code
    with use_vllm_config(vllm_config):
-        payload = VllmSerializableFunction.serialize_compile_artifacts(
-            VllmSerializableFunction(gm, (example_input,), "", foo)
+        payload = VllmSerializableFunction.serialize_graph_module(gm)
+        fake_mode = FakeTensorMode(shape_env=ShapeEnv())
+        loaded_gm = VllmSerializableFunction.deserialize_graph_module(
+            payload, fake_mode
        )
-        fn = VllmSerializableFunction.deserialize_compile_artifacts(payload)

-    assert gm.code == fn.graph_module.code
+    assert gm.code == loaded_gm.code


@pytest.mark.skipif(not is_torch_equal_or_newer("2.10.0"), reason="requires torch 2.10")
@@ -725,6 +730,10 @@ class TestStandaloneCompiledArtifactsIntegration:
        ]:
            assert cache.get(submod, shape) == shared_data

+    @pytest.mark.skipif(
+        envs.VLLM_USE_MEGA_AOT_ARTIFACT,
+        reason="There's no AOT Autograd run with mega artifact",
+    )
    def test_functorch_config(self):
        vllm_config = make_vllm_config()
        example_inputs = (torch.randn(10, 10),)
@@ -9,11 +9,15 @@ then runs in the parent with clean in-memory state but populated caches.

 import multiprocessing as mp

+import pytest
 from torch._dynamo.utils import counters

+import vllm.envs as envs
 from vllm.compilation.counter import compilation_counter
 from vllm.config import CompilationConfig, CompilationMode, CUDAGraphMode

+from ..utils import fork_new_process_for_each_test
+
 MODEL = "microsoft/Phi-tiny-MoE-instruct"


@@ -45,8 +49,11 @@ def _cold_start(vllm_runner):
    assert counters["aot_autograd"]["autograd_cache_hit"] == 0


-def test_moe_startup(monkeypatch, vllm_runner, fresh_vllm_cache):
+@fork_new_process_for_each_test
+@pytest.mark.parametrize("mega_aot_artifact", ["0", "1"])
+def test_moe_startup(monkeypatch, vllm_runner, fresh_vllm_cache, mega_aot_artifact):
    monkeypatch.setenv("VLLM_ENABLE_V1_MULTIPROCESSING", "0")
+    monkeypatch.setenv("VLLM_USE_MEGA_AOT_ARTIFACT", mega_aot_artifact)

    # Cold start in a forked child (must fork before CUDA init).
    # This model has 32 identical transformer layers which produce
@@ -64,7 +71,12 @@ def test_moe_startup(monkeypatch, vllm_runner, fresh_vllm_cache):
        num_compiled_artifacts_saved=0,
    ):
        _run_vllm(vllm_runner)
-    assert counters["aot_autograd"]["total"] == 30
+    if envs.VLLM_USE_MEGA_AOT_ARTIFACT:
+        # MEGA_AOT_ARTIFACT is enabled, so we expect no aot_autograd running on
+        # subgraphs.
+        assert counters["aot_autograd"]["total"] == 0
+    else:
+        assert counters["aot_autograd"]["total"] == 30
    assert counters["aot_autograd"]["autograd_cache_miss"] == 0
    assert (
        counters["aot_autograd"]["autograd_cache_hit"] == 0
@@ -6,9 +6,6 @@ from copy import deepcopy

 from tblib import pickling_support

-# Import fixture
-from tests.v1.entrypoints.conftest import sample_json_schema  # noqa
-
 # ruff: noqa

 # Install support for pickling exceptions so that we can nicely propagate
@@ -81,6 +78,55 @@ if TYPE_CHECKING:

 logger = init_logger(__name__)

+
+@pytest.fixture
+def sample_json_schema():
+    return {
+        "type": "object",
+        "properties": {
+            "name": {"type": "string"},
+            "age": {"type": "integer"},
+            "skills": {
+                "type": "array",
+                "items": {
+                    "type": "string",
+                },
+            },
+            "grade": {
+                "type": "string",
+                "pattern": "^[A-D]$",
+            },
+            "email": {
+                "type": "string",
+                "pattern": "^[a-zA-Z0-9._%+-]+@[a-zA-Z0-9.-]+\\.[a-zA-Z]{2,}$",
+            },
+            "work_history": {
+                "type": "array",
+                "items": {
+                    "type": "object",
+                    "properties": {
+                        "company": {"type": "string"},
+                        "duration": {
+                            "type": "number",
+                            "minimum": 0.0,
+                            "maximum": 100.0,
+                        },
+                        "position": {"type": "string"},
+                    },
+                    "required": ["company", "duration", "position"],
+                    "additionalProperties": False,
+                },
+                "minItems": 0,
+                "maxItems": 3,
+            },
+        },
+        "required": ["name", "age", "skills", "grade", "email", "work_history"],
+        "additionalProperties": False,
+        "minProperties": 1,
+        "maxProperties": 10,
+    }
+
+
 _TEST_DIR = os.path.dirname(__file__)
 _TEST_PROMPTS = [os.path.join(_TEST_DIR, "prompts", "example.txt")]
 _LONG_PROMPTS = [os.path.join(_TEST_DIR, "prompts", "summary.txt")]
@@ -24,6 +24,108 @@ from vllm.sampling_params import (
    StructuredOutputsParams,
 )

+SAMPLE_REGEX = (
+    r"((25[0-5]|(2[0-4]|1\d|[1-9]|)\d)\.){3}"
+    r"(25[0-5]|(2[0-4]|1\d|[1-9]|)\d)"
+)
+
+# Note: Ensure this only uses attributes compatible with xgrammar
+SAMPLE_JSON_SCHEMA = {
+    "type": "object",
+    "properties": {
+        "name": {"type": "string"},
+        "age": {"type": "integer"},
+        "skills": {
+            "type": "array",
+            "items": {
+                "type": "string",
+            },
+        },
+        "grade": {
+            "type": "string",
+            "pattern": "^[A-D]$",  # Regex pattern
+        },
+        "email": {
+            "type": "string",
+            "pattern": "^[a-zA-Z0-9._%+-]+@[a-zA-Z0-9.-]+\\.[a-zA-Z]{2,}$",
+        },
+        "work_history": {
+            "type": "array",
+            "items": {
+                "type": "object",
+                "properties": {
+                    "company": {"type": "string"},
+                    "duration": {
+                        "type": "number",
+                        "minimum": 0.0,
+                        "maximum": 100.0,  # Numeric range
+                    },
+                    "position": {"type": "string"},
+                },
+                "required": ["company", "duration", "position"],
+                "additionalProperties": False,
+            },
+            "minItems": 0,
+            "maxItems": 3,
+        },
+    },
+    "required": ["name", "age", "skills", "grade", "email", "work_history"],
+    "additionalProperties": False,
+    "minProperties": 1,
+    "maxProperties": 10,
+}
+
+# A schema unsupported by xgrammar
+UNSUPPORTED_JSON_SCHEMA = {
+    "type": "object",
+    "properties": {
+        "score": {
+            "type": "integer",
+            "multipleOf": 5,  # Numeric multiple
+        },
+        "tags": {
+            "type": "array",
+            "items": {"type": "string", "minLength": 10, "maxLength": 20},
+        },
+    },
+    "required": ["score", "tags"],
+    "additionalProperties": False,
+    "patternProperties": {
+        "^score$": {"type": "integer"},
+    },
+}
+
+SAMPLE_STRUCTURED_OUTPUTS_CHOICES = [
+    "Python",
+    "Java",
+    "JavaScript",
+    "C++",
+    "C#",
+    "PHP",
+    "TypeScript",
+    "Ruby",
+    "Swift",
+    "Kotlin",
+]
+
+SAMPLE_SQL_EBNF = """
+root ::= select_statement
+select_statement ::= "SELECT" column "from" table "where" condition
+column ::= "col_1" | "col_2"
+table ::= "table_1" | "table_2"
+condition ::= column "=" number
+number ::= "1" | "2"
+"""
+
+SAMPLE_SQL_LARK = """
+start: select_statement
+select_statement: "SELECT" column "from" table "where" condition
+column: "col_1" | "col_2"
+table: "table_1" | "table_2"
+condition: column "=" number
+number: "1" | "2"
+"""
+
 NGRAM_SPEC_CONFIG = {
    "model": "[ngram]",
    "num_speculative_tokens": 5,
@@ -110,17 +212,17 @@ class CarDescription(BaseModel):
    PARAMS_MODELS_BACKENDS_TOKENIZER_MODE,
 )
 def test_structured_output(
-    sample_json_schema: dict[str, Any],
-    unsupported_json_schema: dict[str, Any],
-    sample_sql_ebnf: str,
-    sample_sql_lark: str,
-    sample_regex: str,
-    sample_structured_outputs_choices: str,
    backend: str,
    tokenizer_mode: str,
    model_name: str,
    speculative_config: dict[str, Any],
 ):
+    sample_json_schema = SAMPLE_JSON_SCHEMA
+    unsupported_json_schema = UNSUPPORTED_JSON_SCHEMA
+    sample_sql_ebnf = SAMPLE_SQL_EBNF
+    sample_sql_lark = SAMPLE_SQL_LARK
+    sample_regex = SAMPLE_REGEX
+    sample_structured_outputs_choices = SAMPLE_STRUCTURED_OUTPUTS_CHOICES
    if current_platform.is_tpu() and speculative_config:
        pytest.skip("TPU does not support speculative decoding")

@@ -702,10 +804,10 @@ def test_structured_output_with_reasoning_matrices(

@pytest.mark.parametrize("model_name, tokenizer_mode", PARAMS_MODELS_TOKENIZER_MODE)
 def test_structured_output_auto_mode(
-    unsupported_json_schema: dict[str, Any],
    model_name: str,
    tokenizer_mode: str,
 ):
+    unsupported_json_schema = UNSUPPORTED_JSON_SCHEMA
    llm = LLM(
        model=model_name,
        max_model_len=1024,
@@ -808,9 +910,9 @@ def test_guidance_no_additional_properties():

@pytest.mark.parametrize("backend", ["guidance", "xgrammar", "outlines"])
 def test_structured_output_batched_with_non_structured_outputs_requests(
-    sample_json_schema: dict[str, Any],
    backend: str,
 ):
+    sample_json_schema = SAMPLE_JSON_SCHEMA
    # Don't use eager execution on TPUs because we want to test for no
    # recompilation at runtime
    enforce_eager = bool(not current_platform.is_tpu())
@@ -231,13 +231,14 @@ def k2_server():
        "--gpu-memory-utilization",
        "0.4",
    ] + ROCM_EXTRA_ARGS
-    # hack to test kimi_k2 tool use tool_id format.
-    # avoid error in is_deepseek_mla check by setting kv_lora_rank=null
+    # Test kimi_k2 tool use tool_id format by overriding model_type.
+    # is_deepseek_mla safely returns False via getattr when kv_lora_rank
+    # is absent from the underlying config.
    with RemoteOpenAIServer(
        MODEL_NAME,
        args,
        env_dict=ROCM_ENV_OVERRIDES,
-        override_hf_configs={"model_type": "kimi_k2", "kv_lora_rank": None},
+        override_hf_configs={"model_type": "kimi_k2"},
    ) as remote_server:
        yield remote_server

@@ -152,5 +152,5 @@ async def test_basic_audio_foscolo(foscolo, rocm_aiter_fa_attention, model_name)
            model_name,
            foscolo,
            language="it",
-            expected_text="ove il mio corpo fanciulletto giacque",
+            expected_text="ove il mio corpo fanciulletto",
        )
@@ -275,7 +275,7 @@ INPUT_REASONING_BATCH = "\n".join(
    ]
 )

-MINIMAL_WAV_BASE64 = "UklGRiQAAABXQVZFZm10IBAAAAABAAEAQB8AAEAfAAABAAgAZGF0YQAAAAA="
+MINIMAL_WAV_BASE64 = "UklGRigAAABXQVZFZm10IBAAAAABAAEAgD4AAAB9AAACABAAZGF0YQQAAAAAAP9/"
 INPUT_TRANSCRIPTION_BATCH = (
    json.dumps(
        {
@@ -11,11 +11,10 @@ import pytest_asyncio
 import requests
 from fastapi import Request

+from tests.utils import RemoteOpenAIServer
 from vllm.v1.engine.exceptions import EngineDeadError
 from vllm.version import __version__ as VLLM_VERSION

-from ...utils import RemoteOpenAIServer
-
 MODEL_NAME = "Qwen/Qwen3-0.6B"


@@ -10,7 +10,7 @@ from http import HTTPStatus
 import pytest
 import requests

-from ...utils import RemoteOpenAIServer
+from tests.utils import RemoteOpenAIServer

 # Use a small embeddings model for faster startup and smaller memory footprint.
 # Since we are not testing any chat functionality,
@@ -5,7 +5,7 @@ import openai
 import pytest
 import pytest_asyncio

-from ...utils import RemoteOpenAIServer
+from tests.utils import RemoteOpenAIServer

 # any model with a chat template should work here
 MODEL_NAME = "Qwen/Qwen3-0.6B"
@@ -0,0 +1,8 @@
+model_name: "Qwen/Qwen3.5-35B-A3B"
+accuracy_threshold: 0.86
+num_questions: 1319
+num_fewshot: 5
+server_args: >-
+  --max-model-len 4096
+  --data-parallel-size 2
+  --enable-expert-parallel
@@ -0,0 +1,9 @@
+model_name: "Qwen/Qwen3.5-35B-A3B-FP8"
+accuracy_threshold: 0.86
+num_questions: 1319
+num_fewshot: 5
+server_args: >-
+  --max-model-len 4096
+  --data-parallel-size 2
+  --enable-expert-parallel
+  --kv-cache-dtype fp8
@@ -0,0 +1 @@
+Qwen3.5-35B-A3B-DEP2.yaml
@@ -14,8 +14,19 @@ from vllm.config import (
 )
 from vllm.platforms import current_platform
 from vllm.platforms.cpu import CpuPlatform
-from vllm.platforms.cuda import CudaPlatform
-from vllm.platforms.rocm import RocmPlatform
+
+# CudaPlatform and RocmPlatform import their respective compiled C extensions
+# at module level, raising ModuleNotFoundError on incompatible builds.
+try:
+    from vllm.platforms.cuda import CudaPlatform
+except (ImportError, ModuleNotFoundError):
+    CudaPlatform = None
+
+try:
+    from vllm.platforms.rocm import RocmPlatform
+except (ImportError, ModuleNotFoundError):
+    RocmPlatform = None
+
 from vllm.v1.attention.backends.registry import AttentionBackendEnum
 from vllm.v1.attention.selector import _cached_get_attn_backend, get_attn_backend

@@ -101,6 +112,8 @@ def test_backend_selection(
            assert backend.get_name() == "CPU_ATTN"

        elif device == "hip":
+            if RocmPlatform is None:
+                pytest.skip("RocmPlatform not available")
            with patch("vllm.platforms.current_platform", RocmPlatform()):
                if use_mla:
                    # ROCm MLA backend logic:
@@ -126,6 +139,8 @@ def test_backend_selection(
                    assert backend.get_name() == expected

        elif device == "cuda":
+            if CudaPlatform is None:
+                pytest.skip("CudaPlatform not available")
            with patch("vllm.platforms.current_platform", CudaPlatform()):
                capability = torch.cuda.get_device_capability()
                if use_mla:
@@ -214,7 +229,7 @@ def test_backend_selection(
                    assert backend.get_name() == expected


-@pytest.mark.parametrize("device", ["cpu", "cuda"])
+@pytest.mark.parametrize("device", ["cpu", "cuda", "hip"])
 def test_fp32_fallback(device: str):
    """Test attention backend selection with fp32."""
    # Use default config (no backend specified)
@@ -227,10 +242,25 @@ def test_fp32_fallback(device: str):
            assert backend.get_name() == "CPU_ATTN"

        elif device == "cuda":
+            if CudaPlatform is None:
+                pytest.skip("CudaPlatform not available")
            with patch("vllm.platforms.current_platform", CudaPlatform()):
                backend = get_attn_backend(16, torch.float32, None)
            assert backend.get_name() == "FLEX_ATTENTION"

+        elif device == "hip":
+            if RocmPlatform is None:
+                pytest.skip("RocmPlatform not available")
+            # ROCm backends do not support head_size=16 (minimum is 32).
+            # No known HuggingFace transformer model uses head_size=16.
+            # Revisit if a real model with this head size is identified
+            # and accuracy-tested.
+            with (
+                patch("vllm.platforms.current_platform", RocmPlatform()),
+                pytest.raises(ValueError, match="No valid attention backend"),
+            ):
+                get_attn_backend(16, torch.float32, None)
+

 def test_flash_attn(monkeypatch: pytest.MonkeyPatch):
    """Test FlashAttn validation."""
@@ -367,6 +397,8 @@ def test_per_head_quant_scales_backend_selection(
        attention_config=attention_config, cache_config=cache_config
    )

+    if CudaPlatform is None:
+        pytest.skip("CudaPlatform not available")
    with (
        set_current_vllm_config(vllm_config),
        patch("vllm.platforms.current_platform", CudaPlatform()),
@@ -6,6 +6,7 @@ import pytest
 import torch
 import torch.nn.functional as F

+from vllm.platforms import current_platform
 from vllm.utils.import_utils import has_triton_kernels

 if not has_triton_kernels():
@@ -14,6 +15,7 @@ if not has_triton_kernels():
        allow_module_level=True,
    )

+import triton_kernels.matmul_ogs_details.opt_flags as opt_flags
 import triton_kernels.swiglu
 from triton_kernels.matmul_ogs import FlexCtx, PrecisionConfig
 from triton_kernels.numerics import InFlexData
@@ -21,12 +23,16 @@ from triton_kernels.numerics_details.mxfp import downcast_to_mxfp, upcast_from_m
 from triton_kernels.tensor import FP4, convert_layout, wrap_torch_tensor
 from triton_kernels.tensor_details import layout
 from triton_kernels.testing import assert_close
+from triton_kernels.topk import topk as topk_fn

 from vllm.model_executor.layers.fused_moe.config import mxfp4_w4a16_moe_quant_config
 from vllm.model_executor.layers.fused_moe.gpt_oss_triton_kernels_moe import (
+    legacy_routing,
+    make_routing_data,
    triton_kernel_moe_forward,
 )
 from vllm.utils.math_utils import round_up
+from vllm.utils.torch_utils import set_random_seed

 from .utils import shuffle_weight

@@ -299,6 +305,12 @@ def test_equiv(num_token, a_dtype, w_dtype, tp, workspace_init):
        pc2,
    ) = init_compute_data(M, K, N, E, a_dtype, w_dtype, num_warps=8)

+    if current_platform.is_device_capability_family(100):
+        constraints = {
+            "is_persistent": True,
+        }
+        opt_flags.update_opt_flags_constraints(constraints)
+
    if a_dtype == "bf16" and w_dtype == "mx4":
        quant_config = mxfp4_w4a16_moe_quant_config(
            w1_scale=pc1,
@@ -355,3 +367,43 @@ def test_unit_shuffle():
    )

    assert_close(ref=out_ref, tri=out)
+
+
+@pytest.mark.parametrize("num_tokens", [2, 8, 64])
+@pytest.mark.parametrize("num_experts", [32, 128])
+@pytest.mark.parametrize("topk", [1, 4])
+@pytest.mark.parametrize("renormalize", [True, False])
+@pytest.mark.parametrize("dtype", [torch.bfloat16])
+def test_legacy_routing(
+    num_tokens: int, num_experts: int, topk: int, renormalize: bool, dtype: torch.dtype
+):
+    set_random_seed(0)
+    gating_output = torch.randn(num_tokens, num_experts, device="cuda", dtype=dtype)
+
+    sm_first = not renormalize
+    logits = gating_output
+    if sm_first:
+        logits = torch.softmax(logits, dim=-1)
+    sparse_logits = topk_fn(logits, topk, apply_softmax=not sm_first)
+    topk_ids = sparse_logits.indx.to(torch.long)
+    topk_weights = sparse_logits.vals
+    routing_data_ref, gather_indx_ref, scatter_indx_ref = make_routing_data(
+        topk_ids, topk_weights, num_experts
+    )
+
+    routing_data, gather_indx, scatter_indx = legacy_routing(
+        gating_output, topk, sm_first=sm_first
+    )
+
+    assert_close(
+        ref=gather_indx_ref.src_indx, tri=gather_indx.src_indx, maxtol=0, rmstol=0
+    )
+    assert_close(
+        ref=gather_indx_ref.dst_indx, tri=gather_indx.dst_indx, maxtol=0, rmstol=0
+    )
+    assert_close(
+        ref=scatter_indx_ref.src_indx, tri=scatter_indx.src_indx, maxtol=0, rmstol=0
+    )
+    assert_close(
+        ref=scatter_indx_ref.dst_indx, tri=scatter_indx.dst_indx, maxtol=0, rmstol=0
+    )
@@ -82,7 +82,7 @@ def test_mxfp4_loading_and_execution_moe(vllm_runner, model_case: ModelCase):
        model_case.model_id,
        tensor_parallel_size=model_case.tp,
        load_format="dummy",
-        cudagraph_capture_sizes=[16],
+        compilation_config={"cudagraph_capture_sizes": [16]},
    ) as llm:
        # Disabled as check_model is broken: https://github.com/vllm-project/vllm/pull/18465#issuecomment-3329880562
        # def check_model(model):
@@ -17,89 +17,6 @@ from unittest.mock import MagicMock, patch
 import pytest
 import torch

-from vllm.model_executor.layers.quantization.mxfp4 import (
-    Mxfp4Backend,
-    Mxfp4MoEMethod,
-)
-
-
-def _make_mock_moe_config(ep_size: int = 1) -> MagicMock:
-    """Create a mock FusedMoEConfig with the given EP size."""
-    parallel_config = MagicMock()
-    parallel_config.ep_size = ep_size
-
-    moe_config = MagicMock()
-    moe_config.ep_size = ep_size
-    moe_config.is_lora_enabled = False
-    moe_config.moe_parallel_config = parallel_config
-    return moe_config
-
-
-class TestMxfp4TritonIsMonolithic:
-    """Verify that is_monolithic is always True for the TRITON backend,
-    regardless of EP size, since triton_kernel_moe_forward now handles
-    expert_map remapping internally."""
-
-    @pytest.mark.parametrize(
-        "backend,ep_size,expected_monolithic",
-        [
-            # TRITON is always monolithic (handles EP via expert_map remapping)
-            (Mxfp4Backend.TRITON, 1, True),
-            (Mxfp4Backend.TRITON, 2, True),
-            (Mxfp4Backend.TRITON, 4, True),
-            # SM100 backends are always monolithic
-            (Mxfp4Backend.SM100_FI_MXFP4_MXFP8_TRTLLM, 1, True),
-            (Mxfp4Backend.SM100_FI_MXFP4_MXFP8_TRTLLM, 2, True),
-            (Mxfp4Backend.SM100_FI_MXFP4_BF16, 1, True),
-            (Mxfp4Backend.SM100_FI_MXFP4_BF16, 2, True),
-            # MARLIN is never monolithic
-            (Mxfp4Backend.MARLIN, 1, False),
-            (Mxfp4Backend.MARLIN, 2, False),
-        ],
-        ids=[
-            "triton-no-ep",
-            "triton-ep2",
-            "triton-ep4",
-            "sm100-trtllm-no-ep",
-            "sm100-trtllm-ep2",
-            "sm100-bf16-no-ep",
-            "sm100-bf16-ep2",
-            "marlin-no-ep",
-            "marlin-ep2",
-        ],
-    )
-    @patch(
-        "vllm.model_executor.layers.quantization.mxfp4.get_mxfp4_backend",
-    )
-    @patch(
-        "vllm.model_executor.layers.quantization.mxfp4.get_current_vllm_config",
-    )
-    def test_is_monolithic(
-        self,
-        mock_get_config,
-        mock_get_backend,
-        backend,
-        ep_size,
-        expected_monolithic,
-    ):
-        """is_monolithic should be True for TRITON regardless of EP size."""
-        mock_get_backend.return_value = backend
-
-        mock_compilation_config = MagicMock()
-        mock_compilation_config.max_cudagraph_capture_size = 1024
-        mock_vllm_config = MagicMock()
-        mock_vllm_config.compilation_config = mock_compilation_config
-        mock_get_config.return_value = mock_vllm_config
-
-        moe_config = _make_mock_moe_config(ep_size=ep_size)
-        method = Mxfp4MoEMethod(moe_config)
-
-        assert method.is_monolithic == expected_monolithic, (
-            f"Expected is_monolithic={expected_monolithic} for "
-            f"backend={backend.name}, ep_size={ep_size}, "
-            f"but got {method.is_monolithic}."
-        )
-

 class TestTritonMoeForwardExpertMap:
    """Test that triton_kernel_moe_forward applies expert_map remapping
@@ -160,6 +160,8 @@ def test_rocm_wvsplitkrc_kernel(xnorm, n, k, m, dtype, seed, padded_a, bias_mode
        BIAS = torch.rand(m, dtype=dtype, device="cuda") * 2 - 1
    elif bias_mode == 2:
        BIAS = torch.rand(n, m, dtype=dtype, device="cuda") * 2 - 1
+    elif bias_mode == 3:
+        BIAS = torch.rand(1, m, dtype=dtype, device="cuda") * 2 - 1

    ref_out = torch.nn.functional.linear(A, B, BIAS)
    out = ops.wvSplitKrc(A, B, cu_count, BIAS)
@@ -224,10 +226,9 @@ def test_rocm_wvsplitk_kernel(
    ref_out = torch.nn.functional.linear(A, B, BIAS)
    out = ops.wvSplitK(B, A.view(-1, A.size(-1)), cu_count, BIAS)

-    if xnorm:
-        assert torch.allclose(out, ref_out, atol=1e-3, rtol=1e-8)
-    else:
-        assert torch.allclose(out, ref_out, atol=1e-3, rtol=1e-2)
+    # Accumulation error in fp16 GEMM scales with sqrt(K)
+    atol = torch.finfo(dtype).eps * math.sqrt(k)
+    torch.testing.assert_close(out, ref_out, atol=atol, rtol=1e-2)


@pytest.mark.parametrize("xnorm", [False, True])
@@ -0,0 +1,89 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+from unittest.mock import MagicMock
+
+import pytest
+import torch
+
+from vllm.platforms import current_platform
+
+if current_platform.is_cuda():
+    pytest.skip(
+        "ROCm skinny GEMM tests are not supported on CUDA.",
+        allow_module_level=True,
+    )
+
+from vllm.model_executor.layers import utils
+
+
+def test_rocm_unquantized_gemm_gfx1x_wvsplitk_path(monkeypatch):
+    x = torch.randn(1, 64, dtype=torch.float16)
+    weight = torch.randn(128, 64, dtype=torch.float16)
+
+    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(utils, "get_cu_count", lambda: 120)
+
+    wvsplitk_mock = MagicMock(side_effect=lambda w, x_view, _, __: x_view @ w.t())
+    monkeypatch.setattr(utils.ops, "wvSplitK", wvsplitk_mock)
+    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)
+
+    wvsplitk_mock.assert_called_once()
+    llmm1_mock.assert_not_called()
+    assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3)
+
+
+def test_rocm_unquantized_gemm_gfx1x_n_gt_4_falls_back(monkeypatch):
+    x = torch.randn(5, 64, dtype=torch.float16)
+    weight = torch.randn(128, 64, dtype=torch.float16)
+
+    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(utils, "get_cu_count", lambda: 120)
+
+    wvsplitk_mock = MagicMock(side_effect=lambda w, x_view, _, __: x_view @ w.t())
+    monkeypatch.setattr(utils.ops, "wvSplitK", wvsplitk_mock)
+    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)
+
+    wvsplitk_mock.assert_not_called()
+    llmm1_mock.assert_not_called()
+    assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3)
+
+
+def test_rocm_unquantized_gemm_gfx950_wvsplitkrc_path(monkeypatch):
+    x = torch.randn(16, 1024, dtype=torch.float16)
+    weight = torch.randn(256, 1024, dtype=torch.float16)
+
+    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: False)
+    monkeypatch.setattr("vllm.platforms.rocm.on_gfx9", lambda: False)
+    monkeypatch.setattr("vllm.platforms.rocm.on_gfx950", lambda: True)
+    monkeypatch.setattr(utils, "get_cu_count", lambda: 120)
+
+    wvsplitkrc_mock = MagicMock(side_effect=lambda w, x_view, _, __: x_view @ w.t())
+    monkeypatch.setattr(utils.ops, "wvSplitKrc", wvsplitkrc_mock)
+    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)
+
+    wvsplitkrc_mock.assert_called_once()
+    wvsplitk_mock.assert_not_called()
+    assert torch.allclose(out, ref, atol=1e-3, rtol=1e-3)
@@ -59,6 +59,22 @@ COLBERT_MODELS = {
            "model_cls": "AutoModel",
        },
    },
+    "lfm2": {
+        "model": "LiquidAI/LFM2-ColBERT-350M",
+        "colbert_dim": 128,
+        "max_model_len": 511,
+        "extra_kwargs": {
+            "hf_overrides": {
+                "architectures": ["ColBERTLfm2Model"],
+            },
+        },
+        "hf_comparison": {
+            "weights_file": "1_Dense/model.safetensors",
+            "weights_key": "linear.weight",
+            "trust_remote_code": False,
+            "model_cls": "AutoModel",
+        },
+    },
 }


@@ -220,7 +220,10 @@ VLM_TEST_SETTINGS = {
        vllm_runner_kwargs={
            "model_impl": "transformers",
        },
-        marks=[pytest.mark.core_model],
+        marks=[
+            pytest.mark.core_model,
+            *([large_gpu_mark(min_gb=80)] if current_platform.is_rocm() else []),
+        ],
    ),
    "idefics3-transformers": VLMTestInfo(
        models=["HuggingFaceTB/SmolVLM-256M-Instruct"],
@@ -39,7 +39,11 @@ models = [MODEL_NAME]
 def granite_speech_attention_config():
    """Return attention config for Granite Speech tests on ROCm."""
    if current_platform.is_rocm():
-        return {"backend": "ROCM_AITER_FA"}
+        from vllm.platforms.rocm import on_mi3xx
+
+        if on_mi3xx():
+            return {"backend": "ROCM_AITER_FA"}
+        return {"backend": "TRITON_ATTN"}
    return None


@@ -323,10 +323,7 @@ def build_audio_inputs_from_test_info(
        test_info.audio_idx_to_prompt,
        test_info.prompt_formatter,
    )
-    resampler = AudioResampler(
-        target_sr=16000,
-        method="librosa",
-    )
+    resampler = AudioResampler(target_sr=16000)
    audios = [asset.audio_and_sample_rate for asset in audio_assets]
    resampled_audios = [
        (
@@ -24,6 +24,7 @@ from transformers import (
    GenerationConfig,
    GenerationMixin,
 )
+from transformers.masking_utils import create_causal_mask
 from transformers.video_utils import VideoMetadata

 from vllm.logprobs import SampleLogprobs
@@ -680,10 +681,14 @@ def isaac_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
        sin = sin.to(inputs_embeds.dtype)

        # Prepare attention mask
-        if attention_mask is not None:
-            attention_mask = self._update_causal_mask(
-                attention_mask, inputs_embeds, cache_position, past_key_values, False
-            )
+        attention_mask = create_causal_mask(
+            config=self.config,
+            input_embeds=inputs_embeds,
+            attention_mask=attention_mask,
+            past_key_values=past_key_values,
+            position_ids=position_ids,
+            cache_position=cache_position,
+        )

        # Initialize and collect hidden states
        hidden_states = inputs_embeds
@@ -22,8 +22,10 @@ from vllm.entrypoints.chat_utils import (
    ChatCompletionContentPartTextParam,
 )
 from vllm.entrypoints.pooling.score.utils import ScoreMultiModalParam
+from vllm.platforms import current_platform

 from ....conftest import IMAGE_ASSETS, HfRunner, PromptImageInput, VllmRunner
+from ....utils import ROCM_ENGINE_KWARGS
 from ...utils import check_embeddings_close

 # Prefixes used by the model API
@@ -70,6 +72,7 @@ def _run_test(
        max_model_len=2048,
        enforce_eager=True,
        trust_remote_code=True,
+        **ROCM_ENGINE_KWARGS,
    ) as vllm_model:
        vllm_outputs = vllm_model.embed(input_texts, images=input_images)

@@ -250,6 +253,7 @@ def _run_vllm_reranker(
        max_model_len=2048,
        enforce_eager=True,
        trust_remote_code=True,
+        **ROCM_ENGINE_KWARGS,
    ) as vllm_model:
        has_images = any(img is not None for _, img in docs)

@@ -322,8 +326,11 @@ def _run_reranker_test(
    assert len(hf_scores) == len(vllm_scores), (
        f"Output length mismatch: HF={len(hf_scores)}, vLLM={len(vllm_scores)}"
    )
+    # NOTE: ROCm shows slightly higher numerical variance dues to different attention
+    # backend between vLLM and HF; use a marginally looser tolerance
+    rel_tol = 0.022 if current_platform.is_rocm() else 0.02
    for i, (hf_score, vllm_score) in enumerate(zip(hf_scores, vllm_scores)):
-        assert hf_score == pytest.approx(vllm_score, rel=0.02), (
+        assert hf_score == pytest.approx(vllm_score, rel=rel_tol), (
            f"Score mismatch at index {i}: HF={hf_score:.4f}, vLLM={vllm_score:.4f}"
        )

@@ -628,6 +628,11 @@ _LATE_INTERACTION_EXAMPLE_MODELS = {
        trust_remote_code=True,
        hf_overrides={"architectures": ["ColBERTJinaRobertaModel"]},
    ),
+    "ColBERTLfm2Model": _HfExamplesInfo(
+        "LiquidAI/LFM2-ColBERT-350M",
+        trust_remote_code=True,
+        hf_overrides={"architectures": ["ColBERTLfm2Model"]},
+    ),
    # [Multimodal]
    "ColModernVBertForRetrieval": _HfExamplesInfo(
        "ModernVBERT/colmodernvbert-merged",
@@ -10,6 +10,8 @@ import pytest

 from vllm.multimodal.media import AudioMediaIO

+from ...conftest import AudioTestAssets
+
 pytestmark = pytest.mark.cpu_test

 ASSETS_DIR = Path(__file__).parent.parent / "assets"
@@ -22,40 +24,32 @@ def dummy_audio():


@pytest.fixture
-def dummy_audio_bytes():
-    return b"FAKEAUDIOBYTES"
+def dummy_audio_bytes(audio_assets: AudioTestAssets):
+    with open(audio_assets[0].get_local_path(), "rb") as f:
+        return f.read()


 def test_audio_media_io_load_bytes(dummy_audio_bytes):
    audio_io = AudioMediaIO()
-    with patch("librosa.load") as mock_load:
-        mock_load.return_value = (np.array([0.1, 0.2]), 16000)
-        out = audio_io.load_bytes(dummy_audio_bytes)
-        mock_load.assert_called_once()
-        assert isinstance(out[0], np.ndarray)
-        assert out[1] == 16000
+    out = audio_io.load_bytes(dummy_audio_bytes)
+    assert isinstance(out[0], np.ndarray)
+    assert out[1] == 16000


 def test_audio_media_io_load_base64(dummy_audio_bytes):
    audio_io = AudioMediaIO()
    encoded = base64.b64encode(dummy_audio_bytes).decode("utf-8")
-    with patch.object(AudioMediaIO, "load_bytes") as mock_load_bytes:
-        mock_load_bytes.return_value = (np.array([0.1, 0.2]), 16000)
-        out = audio_io.load_base64("audio/wav", encoded)
-        mock_load_bytes.assert_called_once()
-        assert isinstance(out[0], np.ndarray)
-        assert out[1] == 16000
+    out = audio_io.load_base64("audio/wav", encoded)
+    assert isinstance(out[0], np.ndarray)
+    assert out[1] == 16000


-def test_audio_media_io_load_file():
+def test_audio_media_io_load_file(audio_assets: AudioTestAssets):
    audio_io = AudioMediaIO()
-    path = Path("/fake/path.wav")
-    with patch("librosa.load") as mock_load:
-        mock_load.return_value = (np.array([0.1, 0.2]), 16000)
-        out = audio_io.load_file(path)
-        mock_load.assert_called_once_with(path, sr=None)
-        assert isinstance(out[0], np.ndarray)
-        assert out[1] == 16000
+    path = audio_assets[0].get_local_path()
+    out = audio_io.load_file(path)
+    assert isinstance(out[0], np.ndarray)
+    assert out[1] == 16000


 def test_audio_media_io_encode_base64(dummy_audio):
@@ -14,7 +14,7 @@ from vllm.multimodal.audio import (
    AudioSpec,
    ChannelReduction,
    normalize_audio,
-    resample_audio_librosa,
+    resample_audio_pyav,
    resample_audio_scipy,
    split_audio,
 )
@@ -25,14 +25,14 @@ def dummy_audio():
    return np.array([0.0, 0.1, 0.2, 0.3, 0.4], dtype=float)


-def test_resample_audio_librosa(dummy_audio):
-    with patch("vllm.multimodal.audio.librosa.resample") as mock_resample:
-        mock_resample.return_value = dummy_audio * 2
-        out = resample_audio_librosa(dummy_audio, orig_sr=44100, target_sr=22050)
-        mock_resample.assert_called_once_with(
-            dummy_audio, orig_sr=44100, target_sr=22050
-        )
-        assert np.all(out == dummy_audio * 2)
+def test_resample_audio_pyav(dummy_audio):
+    out_down = resample_audio_pyav(dummy_audio, orig_sr=4, target_sr=2)
+    out_up = resample_audio_pyav(dummy_audio, orig_sr=2, target_sr=4)
+    out_same = resample_audio_pyav(dummy_audio, orig_sr=4, target_sr=4)
+
+    assert len(out_down) == 3
+    assert len(out_up) == 10
+    assert np.all(out_same == dummy_audio)


 def test_resample_audio_scipy(dummy_audio):
@@ -56,9 +56,9 @@ def test_resample_audio_scipy_non_integer_ratio(dummy_audio):
    assert np.isfinite(out).all()


-def test_audio_resampler_librosa_calls_resample(dummy_audio):
-    resampler = AudioResampler(target_sr=22050, method="librosa")
-    with patch("vllm.multimodal.audio.resample_audio_librosa") as mock_resample:
+def test_audio_resampler_pyav_calls_resample(dummy_audio):
+    resampler = AudioResampler(target_sr=22050, method="pyav")
+    with patch("vllm.multimodal.audio.resample_audio_pyav") as mock_resample:
        mock_resample.return_value = dummy_audio
        out = resampler.resample(dummy_audio, orig_sr=44100)
        mock_resample.assert_called_once_with(
@@ -423,13 +423,13 @@ class TestAudioPipelineE2E:
        # Verify channel averaging: mean of [0.5, -0.5] = 0.0
        np.testing.assert_array_almost_equal(audio_output, np.zeros(16000), decimal=5)

-    def test_librosa_mono_passthrough_e2e(self):
-        """Full pipeline: librosa mono format → preserved as mono."""
+    def test_pyav_mono_passthrough_e2e(self):
+        """Full pipeline: pyav mono format → preserved as mono."""
        from vllm.multimodal.parse import MultiModalDataParser

-        # Simulate librosa output: already mono (time,) format
-        mono_librosa = np.random.randn(16000).astype(np.float32)
-        assert mono_librosa.shape == (16000,)
+        # Simulate pyav output: already mono (time,) format
+        mono_pyav = np.random.randn(16000).astype(np.float32)
+        assert mono_pyav.shape == (16000,)

        # Create parser with mono normalization
        parser = MultiModalDataParser(
@@ -438,7 +438,7 @@ class TestAudioPipelineE2E:
        )

        # Process audio through the parser
-        result = parser._parse_audio_data((mono_librosa, 16000))
+        result = parser._parse_audio_data((mono_pyav, 16000))
        audio_output = result.get(0)

        # Verify output is still mono 1D
@@ -446,7 +446,7 @@ class TestAudioPipelineE2E:
        assert audio_output.shape == (16000,)

        # Verify audio content is preserved
-        np.testing.assert_array_almost_equal(audio_output, mono_librosa)
+        np.testing.assert_array_almost_equal(audio_output, mono_pyav)

    def test_multichannel_5_1_surround_to_mono_e2e(self):
        """Full pipeline: 5.1 surround (6 channels) → mono output."""
@@ -74,7 +74,7 @@ def test_embed_dimensions(model_info: EmbedModelInfo):
        pooling_params.verify(model_config)


-@pytest.mark.parametrize("task", ["score", "classify"])
+@pytest.mark.parametrize("task", ["classify"])
 def test_classify(task):
    model_config = MockModelConfig(pooler_config=PoolerConfig(seq_pooling_type="CLS"))

@@ -10,9 +10,10 @@ from vllm.v1.attention.backends.utils import reorder_batch_to_split_decodes_and_


 class MockInputBatch:
-    def __init__(self, req_ids, num_computed_tokens_cpu):
+    def __init__(self, req_ids, num_computed_tokens_cpu, num_prompt_tokens):
        self.req_ids = req_ids
        self.num_computed_tokens_cpu = num_computed_tokens_cpu
+        self.num_prompt_tokens = num_prompt_tokens

    def swap_states(self, i, j):
        self.req_ids[i], self.req_ids[j] = self.req_ids[j], self.req_ids[i]
@@ -20,6 +21,10 @@ class MockInputBatch:
            self.num_computed_tokens_cpu[j],
            self.num_computed_tokens_cpu[i],
        )
+        self.num_prompt_tokens[i], self.num_prompt_tokens[j] = (
+            self.num_prompt_tokens[j],
+            self.num_prompt_tokens[i],
+        )


 class MockSchedulerOutput:
@@ -29,96 +34,139 @@ class MockSchedulerOutput:

@dataclass
 class ReorderTestCase:
-    requests: list[tuple[int, int]]  # (num_scheduled_tokens, num_computed_tokens)
+    # (num_scheduled_tokens, num_computed_tokens, num_prompt_tokens)
+    requests: list[tuple[int, int, int]]
    expected_order: list[int]
    expected_modified: bool
    decode_threshold: int = 1


 # Test cases for batch reordering
+# Format: (num_scheduled, num_computed, num_prompt)
 REORDER_TEST_CASES = {
    "all_decodes": ReorderTestCase(
-        requests=[(1, 10), (1, 20), (1, 30)],
+        requests=[(1, 10, 10), (1, 20, 20), (1, 30, 30)],
        expected_order=[0, 1, 2],
        expected_modified=False,
    ),
-    "all_prefills": ReorderTestCase(
-        requests=[(100, 100), (200, 200), (300, 300)],
+    "all_long_extends": ReorderTestCase(
+        requests=[(100, 100, 100), (200, 200, 200), (300, 300, 300)],
        expected_order=[0, 1, 2],
        expected_modified=False,
    ),
-    "mixed_interleaved": ReorderTestCase(
-        requests=[(100, 100), (1, 10), (200, 200), (1, 20)],
-        expected_order=[3, 1, 2, 0],  # Only swap 0↔3, keep 1 and 2 in place
+    "mixed_decodes_long_extends": ReorderTestCase(
+        requests=[(100, 100, 100), (1, 10, 10), (200, 200, 200), (1, 20, 20)],
+        expected_order=[3, 1, 2, 0],
        expected_modified=True,
    ),
    "already_ordered": ReorderTestCase(
-        requests=[(1, 10), (1, 20), (100, 100), (200, 0)],
+        requests=[(1, 10, 10), (1, 20, 20), (100, 100, 100), (200, 0, 200)],
        expected_order=[0, 1, 2, 3],
        expected_modified=False,
    ),
    "single_request": ReorderTestCase(
-        requests=[(1, 10)],
+        requests=[(1, 10, 10)],
        expected_order=[0],
        expected_modified=False,
    ),
    "higher_threshold": ReorderTestCase(
-        requests=[(2, 10), (3, 20), (5, 30), (6, 40)],
+        requests=[(2, 10, 10), (3, 20, 20), (5, 30, 30), (6, 40, 40)],
        expected_order=[0, 1, 2, 3],
        expected_modified=False,
        decode_threshold=4,
    ),
    "decodes_at_end": ReorderTestCase(
-        requests=[(100, 100), (200, 200), (1, 10), (1, 20)],
+        requests=[(100, 100, 100), (200, 200, 200), (1, 10, 10), (1, 20, 20)],
        expected_order=[2, 3, 0, 1],
        expected_modified=True,
    ),
-    "decode_extend_prefill": ReorderTestCase(
-        requests=[(100, 0), (10, 50), (1, 10)],
+    "decode_long_extend_prefill": ReorderTestCase(
+        requests=[(100, 0, 100), (10, 50, 50), (1, 10, 10)],
        expected_order=[2, 1, 0],
        expected_modified=True,
    ),
-    "extend_prefill_only": ReorderTestCase(
-        requests=[(100, 0), (10, 50), (200, 0), (20, 75)],
-        expected_order=[3, 1, 2, 0],  # Only swap 0↔3, keep 1 and 2 in place
+    "long_extend_prefill_only": ReorderTestCase(
+        requests=[(100, 0, 100), (10, 50, 50), (200, 0, 200), (20, 75, 75)],
+        expected_order=[3, 1, 2, 0],
        expected_modified=True,
    ),
-    "complicated_mixed_interleaved": ReorderTestCase(
+    "complicated_mixed": ReorderTestCase(
        requests=[
-            (1, 20),
-            (1, 50),
-            (374, 0),
-            (300, 20),
-            (1, 20),
-            (256, 0),
-            (1, 5),
-            (27, 0),
-            (1, 4),
+            (1, 20, 20),  # decode
+            (1, 50, 50),  # decode
+            (374, 0, 374),  # prefill
+            (300, 20, 20),  # long_extend
+            (1, 20, 20),  # decode
+            (256, 0, 256),  # prefill
+            (1, 5, 5),  # decode
+            (27, 0, 27),  # prefill
+            (1, 4, 4),  # decode
        ],
        expected_order=[0, 1, 6, 8, 4, 3, 2, 7, 5],
        expected_modified=True,
    ),
    "new_request_single_token_prefill": ReorderTestCase(
        requests=[
-            (100, 0),
-            (1, 0),  # New request with only 1 token (STILL prefill)
-            (50, 100),
-            (1, 10),
+            (100, 0, 100),  # prefill
+            (1, 0, 1),  # prefill (single token, still prefill)
+            (50, 100, 100),  # long_extend
+            (1, 10, 10),  # decode
        ],
-        # Only index 3 is a true decode (has num_computed_tokens > 0)
        expected_order=[3, 2, 0, 1],
        expected_modified=True,
    ),
    "multiple_new_requests_single_token_prefill": ReorderTestCase(
        requests=[
-            (1, 0),  # New prefill (1 token, no computed)
-            (1, 0),  # New prefill (1 token, no computed)
-            (1, 50),
-            (200, 0),
+            (1, 0, 1),  # prefill
+            (1, 0, 1),  # prefill
+            (1, 50, 50),  # decode
+            (200, 0, 200),  # prefill
        ],
        expected_order=[2, 1, 0, 3],
        expected_modified=True,
    ),
+    "four_way_already_ordered": ReorderTestCase(
+        requests=[
+            (1, 100, 100),  # decode
+            (1, 50, 100),  # short_extend
+            (10, 50, 100),  # long_extend
+            (100, 0, 100),  # prefill
+        ],
+        expected_order=[0, 1, 2, 3],
+        expected_modified=False,
+    ),
+    "four_way_needs_reorder": ReorderTestCase(
+        requests=[
+            (100, 0, 100),  # prefill
+            (1, 50, 100),  # short_extend
+            (1, 100, 100),  # decode
+            (10, 50, 100),  # long_extend
+        ],
+        expected_order=[2, 1, 3, 0],
+        expected_modified=True,
+    ),
+    "four_way_multiple_short_extends": ReorderTestCase(
+        requests=[
+            (2, 100, 100),  # decode
+            (2, 50, 200),  # short_extend
+            (2, 75, 150),  # short_extend
+            (2, 200, 200),  # decode
+        ],
+        expected_order=[0, 3, 2, 1],
+        expected_modified=True,
+        decode_threshold=2,
+    ),
+    "four_way_spec_decode_threshold": ReorderTestCase(
+        requests=[
+            (5, 100, 100),  # decode
+            (5, 50, 100),  # short_extend
+            (5, 0, 100),  # prefill
+            (10, 50, 100),  # long_extend
+        ],
+        expected_order=[0, 1, 3, 2],
+        expected_modified=True,
+        decode_threshold=5,
+    ),
 }


@@ -129,8 +177,9 @@ def test_reorder_batch_to_split_decodes_and_prefills(test_case: ReorderTestCase)
    req_ids = [f"r{i}" for i in range(len(test_case.requests))]
    num_computed_tokens = np.array([r[1] for r in test_case.requests], dtype=np.int32)
    num_scheduled_tokens = {f"r{i}": r[0] for i, r in enumerate(test_case.requests)}
+    num_prompt_tokens = np.array([r[2] for r in test_case.requests], dtype=np.int32)

-    input_batch = MockInputBatch(req_ids, num_computed_tokens)
+    input_batch = MockInputBatch(req_ids, num_computed_tokens, num_prompt_tokens)
    scheduler_output = MockSchedulerOutput(num_scheduled_tokens)

    modified = reorder_batch_to_split_decodes_and_prefills(
@@ -43,7 +43,7 @@ MESSAGES = [
        pytest.param("Qwen/Qwen3.5-4B", marks=[large_gpu_mark(min_gb=40)]),
        pytest.param(
            "nvidia/NVIDIA-Nemotron-3-Super-120B-A12B-FP8",
-            marks=[large_gpu_mark(min_gb=80)] + multi_gpu_marks(num_gpus=2),
+            marks=[large_gpu_mark(min_gb=80)] + multi_gpu_marks(num_gpus=4),
        ),
    ],
 )
@@ -68,7 +68,7 @@ def test_mtp_speculative_mixed_batch_short_prefill(
        max_num_batched_tokens=chunk_size,
        max_model_len=512,
        enforce_eager=True,
-        tensor_parallel_size=2,
+        tensor_parallel_size=4,
        trust_remote_code=True,
        enable_chunked_prefill=True,
        enable_prefix_caching=enable_prefix_caching,
@@ -1,173 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import pytest
-
-
-@pytest.fixture
-def sample_prompts():
-    return [
-        "Hello, my name is",
-        "The president of the United States is",
-        "The capital of France is",
-        "The future of AI is",
-    ]
-
-
-@pytest.fixture
-def sample_token_ids():
-    return [
-        [0],
-        [0, 1],
-        [0, 2, 1],
-        [0, 3, 1, 2],
-    ]
-
-
-@pytest.fixture
-def sample_regex():
-    return (
-        r"((25[0-5]|(2[0-4]|1\d|[1-9]|)\d)\.){3}"
-        r"(25[0-5]|(2[0-4]|1\d|[1-9]|)\d)"
-    )
-
-
-# Note: Ensure this only uses attributes compatible with xgrammar
-@pytest.fixture
-def sample_json_schema():
-    return {
-        "type": "object",
-        "properties": {
-            "name": {"type": "string"},
-            "age": {"type": "integer"},
-            "skills": {
-                "type": "array",
-                "items": {
-                    "type": "string",
-                },
-            },
-            "grade": {
-                "type": "string",
-                "pattern": "^[A-D]$",  # Regex pattern
-            },
-            "email": {
-                "type": "string",
-                "pattern": "^[a-zA-Z0-9._%+-]+@[a-zA-Z0-9.-]+\\.[a-zA-Z]{2,}$",
-            },
-            "work_history": {
-                "type": "array",
-                "items": {
-                    "type": "object",
-                    "properties": {
-                        "company": {"type": "string"},
-                        "duration": {
-                            "type": "number",
-                            "minimum": 0.0,
-                            "maximum": 100.0,  # Numeric range
-                        },
-                        "position": {"type": "string"},
-                    },
-                    "required": ["company", "duration", "position"],
-                    "additionalProperties": False,
-                },
-                "minItems": 0,
-                "maxItems": 3,
-            },
-        },
-        "required": ["name", "age", "skills", "grade", "email", "work_history"],
-        "additionalProperties": False,
-        "minProperties": 1,
-        "maxProperties": 10,
-    }
-
-
-# A schema unsupported by xgrammar
-@pytest.fixture
-def unsupported_json_schema():
-    return {
-        "type": "object",
-        "properties": {
-            "score": {
-                "type": "integer",
-                "multipleOf": 5,  # Numeric multiple
-            },
-            "tags": {
-                "type": "array",
-                "items": {"type": "string", "minLength": 10, "maxLength": 20},
-            },
-        },
-        "required": ["score", "tags"],
-        "additionalProperties": False,
-        "patternProperties": {
-            "^score$": {"type": "integer"},
-        },
-    }
-
-
-@pytest.fixture
-def sample_definition_json_schema():
-    return {
-        "$defs": {
-            "Step": {
-                "properties": {
-                    "explanation": {"title": "Explanation", "type": "string"},
-                    "output": {"title": "Output", "type": "string"},
-                },
-                "required": ["explanation", "output"],
-                "title": "Step",
-                "type": "object",
-            }
-        },
-        "properties": {
-            "steps": {
-                "items": {"$ref": "#/$defs/Step"},
-                "title": "Steps",
-                "type": "array",
-            },
-            "final_answer": {"title": "Final Answer", "type": "string"},
-        },
-        "required": ["steps", "final_answer"],
-        "title": "MathReasoning",
-        "type": "object",
-        "additionalProperties": False,
-    }
-
-
-@pytest.fixture
-def sample_structured_outputs_choices():
-    return [
-        "Python",
-        "Java",
-        "JavaScript",
-        "C++",
-        "C#",
-        "PHP",
-        "TypeScript",
-        "Ruby",
-        "Swift",
-        "Kotlin",
-    ]
-
-
-@pytest.fixture
-def sample_sql_ebnf():
-    return """
-root ::= select_statement
-select_statement ::= "SELECT" column "from" table "where" condition
-column ::= "col_1" | "col_2"
-table ::= "table_1" | "table_2"
-condition ::= column "=" number
-number ::= "1" | "2"
-"""
-
-
-@pytest.fixture
-def sample_sql_lark():
-    return """
-start: select_statement
-select_statement: "SELECT" column "from" table "where" condition
-column: "col_1" | "col_2"
-table: "table_1" | "table_2"
-condition: column "=" number
-number: "1" | "2"
-"""
@@ -0,0 +1,207 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""Unit tests for MRv2 GPUModelRunner.add_requests streaming input support."""
+
+from unittest.mock import Mock
+
+import pytest
+import torch
+
+from vllm.v1.core.sched.output import (
+    CachedRequestData,
+    NewRequestData,
+    SchedulerOutput,
+)
+from vllm.v1.worker.gpu.model_runner import GPUModelRunner
+from vllm.v1.worker.gpu.states import RequestState
+
+pytestmark = pytest.mark.cpu_test
+
+
+@pytest.fixture
+def mock_model_runner_with_req_states():
+    """Create a mock MRv2 GPUModelRunner with a real RequestState."""
+
+    runner = Mock(spec=GPUModelRunner)
+    runner.req_states = RequestState(
+        max_num_reqs=10,
+        max_model_len=1024,
+        max_num_batched_tokens=1024,
+        num_speculative_steps=0,
+        vocab_size=32000,
+        device=torch.device("cpu"),
+        model_dtype=torch.float32,
+        cache_draft_logits=False,
+    )
+    runner.encoder_cache = None
+    runner.model_state = Mock()
+    runner.block_tables = Mock()
+    runner.lora_state = Mock()
+    runner.sampler = None
+    runner.prompt_logprobs_worker = None
+    runner.is_last_pp_rank = False
+
+    # Mock staged writes — they use Triton kernels that require GPU
+    runner.req_states.apply_staged_writes = Mock()
+
+    # Bind the real methods to our mock
+    runner._remove_request = GPUModelRunner._remove_request.__get__(runner)
+    runner.add_requests = GPUModelRunner.add_requests.__get__(runner)
+    return runner
+
+
+def _make_scheduler_output(new_reqs):
+    return SchedulerOutput(
+        scheduled_new_reqs=new_reqs,
+        scheduled_cached_reqs=CachedRequestData.make_empty(),
+        num_scheduled_tokens={},
+        total_num_scheduled_tokens=0,
+        scheduled_spec_decode_tokens={},
+        scheduled_encoder_inputs={},
+        num_common_prefix_blocks=[],
+        finished_req_ids=set(),
+        free_encoder_mm_hashes=[],
+    )
+
+
+def test_e2e_streaming_request_update_basic_flow(
+    mock_model_runner_with_req_states,
+):
+    """Test that streaming sessions are updated correctly.
+
+    This test validates that when a streaming session is updated with new
+    prompt tokens:
+    1. The old request state is removed (no free_indices leak)
+    2. The new state is written with updated prefill_token_ids
+    3. model_state and block_tables are re-registered for the new state
+    """
+    runner = mock_model_runner_with_req_states
+    req_states = runner.req_states
+    req_id = "streaming_req_0"
+    initial_free = len(req_states.free_indices)
+
+    # Step 1: Add initial request with 3 prompt tokens, all computed
+    initial_req_data = NewRequestData(
+        req_id=req_id,
+        prompt_token_ids=[1, 2, 3],
+        prefill_token_ids=[1, 2, 3],
+        mm_features=[],
+        sampling_params=None,
+        pooling_params=None,
+        block_ids=([0],),
+        num_computed_tokens=3,
+        lora_request=None,
+    )
+    runner.add_requests(_make_scheduler_output([initial_req_data]))
+    assert req_id in req_states.req_id_to_index
+    assert len(req_states.free_indices) == initial_free - 1
+
+    # Step 2: Create streaming update with extended prompt
+    # The scheduler has already set prefill_token_ids to the full sequence
+    # (original prompt + intermediate output + new prompt tokens)
+    updated_req_data = NewRequestData(
+        req_id=req_id,
+        prompt_token_ids=[1, 2, 3],
+        prefill_token_ids=[1, 2, 3, 10, 4, 5],
+        mm_features=[],
+        sampling_params=None,
+        pooling_params=None,
+        block_ids=([0, 1],),
+        num_computed_tokens=4,  # 3 original prompt + 1 intermediate output
+        lora_request=None,
+    )
+    runner.add_requests(_make_scheduler_output([updated_req_data]))
+
+    # Step 3: Verify no free_indices leak (old slot recycled)
+    assert len(req_states.free_indices) == initial_free - 1
+
+    # Verify the request is still tracked with exactly one index
+    assert req_id in req_states.req_id_to_index
+    assert sum(1 for v in req_states.index_to_req_id.values() if v == req_id) == 1
+
+    # Verify state was updated with new values
+    new_idx = req_states.req_id_to_index[req_id]
+    assert req_states.prompt_len.np[new_idx] == 3
+    assert req_states.prefill_len.np[new_idx] == 6
+    assert req_states.num_computed_prefill_tokens[new_idx] == 4
+
+    # Verify model_state and block_tables were re-registered
+    runner.model_state.add_request.assert_called_with(new_idx, updated_req_data)
+    runner.block_tables.append_block_ids.assert_called_with(
+        new_idx, ([0, 1],), overwrite=True
+    )
+
+
+def test_e2e_streaming_with_multimodal_features(
+    mock_model_runner_with_req_states,
+):
+    """Test that streaming sessions with multimodal features are updated.
+
+    This test validates that when a streaming session with mm features
+    is updated:
+    1. The old request state is removed (no free_indices leak)
+    2. encoder_cache is cleaned up and re-registered with new mm_features
+    3. model_state is re-registered (recomputes M-RoPE positions etc.)
+    """
+    runner = mock_model_runner_with_req_states
+    req_states = runner.req_states
+    req_id = "streaming_mm_req_0"
+    initial_free = len(req_states.free_indices)
+
+    # Enable encoder_cache for multimodal
+    runner.encoder_cache = Mock()
+
+    # Step 1: Add initial request with one audio feature
+    mm_feature_1 = Mock()
+    initial_req_data = NewRequestData(
+        req_id=req_id,
+        prompt_token_ids=[1, 2] + [0] * 10 + [3, 4],
+        prefill_token_ids=[1, 2] + [0] * 10 + [3, 4],
+        mm_features=[mm_feature_1],
+        sampling_params=None,
+        pooling_params=None,
+        block_ids=([0],),
+        num_computed_tokens=14,
+        lora_request=None,
+    )
+    runner.add_requests(_make_scheduler_output([initial_req_data]))
+    assert req_id in req_states.req_id_to_index
+
+    # Reset mocks to track only the streaming update calls
+    runner.encoder_cache.reset_mock()
+    runner.model_state.reset_mock()
+
+    # Step 2: Create streaming update with additional multimodal feature
+    # The scheduler has folded the intermediate output (100) into
+    # prefill_token_ids and added a new audio chunk
+    mm_feature_2 = Mock()
+    updated_req_data = NewRequestData(
+        req_id=req_id,
+        prompt_token_ids=[1, 2] + [0] * 10 + [3, 4],
+        prefill_token_ids=[1, 2] + [0] * 10 + [3, 4, 100] + [0] * 5 + [5],
+        mm_features=[mm_feature_1, mm_feature_2],
+        sampling_params=None,
+        pooling_params=None,
+        block_ids=([0, 1],),
+        num_computed_tokens=14,
+        lora_request=None,
+    )
+    runner.add_requests(_make_scheduler_output([updated_req_data]))
+
+    # Step 3: Verify no free_indices leak
+    assert len(req_states.free_indices) == initial_free - 1
+    assert sum(1 for v in req_states.index_to_req_id.values() if v == req_id) == 1
+
+    # Verify encoder_cache was cleaned up and re-registered
+    runner.encoder_cache.remove_request.assert_called_once_with(req_id)
+    runner.encoder_cache.add_request.assert_called_once_with(
+        req_id, [mm_feature_1, mm_feature_2]
+    )
+
+    # Verify model_state was re-registered with new data
+    new_idx = req_states.req_id_to_index[req_id]
+    runner.model_state.add_request.assert_called_once_with(new_idx, updated_req_data)
+
+    # Verify updated prefill length
+    assert req_states.prefill_len.np[new_idx] == 21
@@ -36,6 +36,7 @@ def test_resumed_req_ids_cleared_from_mamba_state_idx():
    spec = MagicMock(block_size=64, num_speculative_blocks=0)
    cache_config = MagicMock(enable_prefix_caching=True)
    input_batch = MagicMock(req_ids=[])
+    copy_bufs = MagicMock(mamba_group_ids=[0], mamba_spec=spec)

    mamba_state_idx = {
        "finished": 1,
@@ -62,7 +63,7 @@ def test_resumed_req_ids_cleared_from_mamba_state_idx():
            {},
            {},
            (),
-            MagicMock(),
+            copy_bufs,
        )

    assert mamba_state_idx == {"keep": 99}
@@ -8,15 +8,10 @@ from urllib.parse import urljoin

 import numpy.typing as npt

-from vllm.utils.import_utils import PlaceholderModule
+from vllm.multimodal.media.audio import load_audio

 from .base import VLLM_S3_BUCKET_URL, get_vllm_public_assets

-try:
-    import librosa
-except ImportError:
-    librosa = PlaceholderModule("librosa")  # type: ignore[assignment]
-
 ASSET_DIR = "multimodal_asset"

 AudioAssetName = Literal["winning_call", "mary_had_lamb"]
@@ -33,7 +28,7 @@ class AudioAsset:
    @property
    def audio_and_sample_rate(self) -> tuple[npt.NDArray, float]:
        audio_path = get_vllm_public_assets(filename=self.filename, s3_prefix=ASSET_DIR)
-        return librosa.load(audio_path, sr=None)
+        return load_audio(audio_path, sr=None)

    def get_local_path(self) -> Path:
        return get_vllm_public_assets(filename=self.filename, s3_prefix=ASSET_DIR)
@@ -10,15 +10,10 @@ import numpy.typing as npt
 from huggingface_hub import hf_hub_download
 from PIL import Image

-from vllm.utils.import_utils import PlaceholderModule
+from vllm.multimodal.media.audio import load_audio_pyav

 from .base import get_cache_dir

-try:
-    import librosa
-except ImportError:
-    librosa = PlaceholderModule("librosa")  # type: ignore[assignment]
-

@lru_cache
 def download_video_asset(filename: str) -> str:
@@ -146,4 +141,4 @@ class VideoAsset:

        See also: examples/offline_inference/qwen2_5_omni/only_thinker.py
        """
-        return librosa.load(self.video_path, sr=sampling_rate)[0]
+        return load_audio_pyav(self.video_path, sr=sampling_rate)[0]
@@ -38,6 +38,7 @@ from typing_extensions import deprecated
 from vllm.lora.request import LoRARequest
 from vllm.lora.utils import get_adapter_absolute_path
 from vllm.multimodal import MultiModalDataDict
+from vllm.multimodal.audio import get_audio_duration
 from vllm.multimodal.image import convert_image_mode
 from vllm.tokenizers import TokenizerLike
 from vllm.utils.argparse_utils import FlexibleArgumentParser
@@ -54,10 +55,6 @@ try:
 except ImportError:
    pd = PlaceholderModule("pandas")

-try:
-    import librosa
-except ImportError:
-    librosa = PlaceholderModule("librosa")

 logger = logging.getLogger(__name__)

@@ -3253,7 +3250,7 @@ class ASRDataset(HuggingFaceDataset):
                break
            audio = item["audio"]
            y, sr = audio["array"], audio["sampling_rate"]
-            duration_s = librosa.get_duration(y=y, sr=sr)
+            duration_s = get_audio_duration(y=y, sr=sr)
            if duration_s < asr_min_audio_len_sec or duration_s > asr_max_audio_len_sec:
                skipped += 1
                continue
@@ -11,10 +11,13 @@ from typing import Any, Literal
 from unittest.mock import patch

 import torch
+from torch._subclasses import FakeTensorMode
+from torch.fx._graph_pickler import GraphPickler, Options
 from torch.utils import _pytree as pytree

 import vllm.envs as envs
 from vllm.compilation.compiler_interface import get_inductor_factors
+from vllm.compilation.counter import compilation_counter
 from vllm.config import VllmConfig, get_current_vllm_config
 from vllm.config.utils import hash_factors
 from vllm.logger import init_logger
@@ -59,6 +62,7 @@ class StandaloneCompiledArtifacts:
        self.submodule_bytes[f"{submod_name}_{shape}"] = hex_digest
        if hex_digest not in self.submodule_bytes_store:
            self.submodule_bytes_store[hex_digest] = entry
+            compilation_counter.num_compiled_artifacts_saved += 1
            logger.debug(
                "inserting new artifact for submod %s with shape %s "
                "(%s bytes) at hash %s",
@@ -122,6 +126,7 @@ class StandaloneCompiledArtifacts:

        def _load_entry(entry_bytes: bytes) -> AOTCompiledArtifact:
            entry = pickle.loads(entry_bytes)
+            compilation_counter.num_compiled_artifacts_loaded += 1
            return AOTCompiledArtifact.deserialize(entry)

        with concurrent.futures.ThreadPoolExecutor() as executor:
@@ -206,26 +211,8 @@ class VllmSerializableFunction(SerializableCallable):  # type: ignore[misc]
        return self.optimized_call(*args, **kwargs)

    @classmethod
-    def serialize_compile_artifacts(
-        cls, compiled_fn: "VllmSerializableFunction"
-    ) -> bytes:
+    def serialize_graph_module(cls, graph_module: torch.fx.GraphModule) -> bytes:
        import sympy
-        from torch._subclasses import FakeTensorMode
-        from torch.fx._graph_pickler import GraphPickler, Options
-
-        state = compiled_fn.__dict__.copy()
-        state.pop("optimized_call")
-        state.pop("shape_env")
-        state.pop("vllm_backend", None)
-        state.pop("_fake_mode", None)
-        for node in state["graph_module"].graph.nodes:
-            node.meta.pop("source_fn_stack", None)
-            node.meta.pop("nn_module_stack", None)
-        for name, submod in state["graph_module"].named_children():
-            if hasattr(submod, "graph"):
-                for node in submod.graph.nodes:
-                    node.meta.pop("source_fn_stack", None)
-                    node.meta.pop("nn_module_stack", None)

        graph_reducer_override = GraphPickler.reducer_override

@@ -242,6 +229,37 @@ class VllmSerializableFunction(SerializableCallable):  # type: ignore[misc]
                return type(None), ()
            return graph_reducer_override(self, obj)

+        with (
+            patch.object(GraphPickler, "reducer_override", _graph_reducer_override),
+            patch_pytree_map_over_slice(),
+        ):
+            return GraphPickler.dumps(graph_module, Options(ops_filter=None))
+
+    @classmethod
+    def deserialize_graph_module(
+        cls, data: bytes, fake_mode: FakeTensorMode
+    ) -> torch.fx.GraphModule:
+        with patch_pytree_map_over_slice():
+            return GraphPickler.loads(data, fake_mode)
+
+    @classmethod
+    def serialize_compile_artifacts(
+        cls, compiled_fn: "VllmSerializableFunction"
+    ) -> bytes:
+        state = compiled_fn.__dict__.copy()
+        state.pop("optimized_call")
+        state.pop("shape_env")
+        state.pop("vllm_backend", None)
+        state.pop("_fake_mode", None)
+        for node in state["graph_module"].graph.nodes:
+            node.meta.pop("source_fn_stack", None)
+            node.meta.pop("nn_module_stack", None)
+        for name, submod in state["graph_module"].named_children():
+            if hasattr(submod, "graph"):
+                for node in submod.graph.nodes:
+                    node.meta.pop("source_fn_stack", None)
+                    node.meta.pop("nn_module_stack", None)
+
        if state.get("sym_tensor_indices"):
            # put tensor inputs on meta device since their data
            # isn't needed, yet we need the meta for make_copy_and_call
@@ -257,14 +275,9 @@ class VllmSerializableFunction(SerializableCallable):  # type: ignore[misc]
                lambda inp: torch.empty_like(inp, device="meta"),
                state["example_inputs"],
            )
-        with (
-            patch.object(GraphPickler, "reducer_override", _graph_reducer_override),
-            patch_pytree_map_over_slice(),
-        ):
-            state["graph_module"] = GraphPickler.dumps(
-                state["graph_module"], Options(ops_filter=None)
-            )
-            state["example_inputs"] = GraphPickler.dumps(state["example_inputs"])
+
+        state["graph_module"] = cls.serialize_graph_module(state["graph_module"])
+        state["example_inputs"] = GraphPickler.dumps(state["example_inputs"])

        if compiled_fn.vllm_backend:
            (
@@ -280,14 +293,14 @@ class VllmSerializableFunction(SerializableCallable):  # type: ignore[misc]
    @classmethod
    def deserialize_compile_artifacts(cls, data: bytes) -> "VllmSerializableFunction":
        from torch._guards import TracingContext, tracing
-        from torch._subclasses import FakeTensorMode
-        from torch.fx._graph_pickler import GraphPickler
        from torch.fx.experimental.symbolic_shapes import ShapeEnv

        state = pickle.loads(data)
        fake_mode = FakeTensorMode(shape_env=ShapeEnv())
-        with patch_pytree_map_over_slice():
-            state["graph_module"] = GraphPickler.loads(state["graph_module"], fake_mode)
+
+        state["graph_module"] = cls.deserialize_graph_module(
+            state["graph_module"], fake_mode
+        )
        state["graph_module"].recompile()
        state["example_inputs"] = GraphPickler.loads(state["example_inputs"], fake_mode)

@@ -1435,10 +1435,10 @@ class ModelConfig:
    @property
    def score_type(self) -> ScoreType:
        """
-        Score API handles score/rerank for:
-        - "score" task (score_type: cross-encoder models)
-        - "embed" task (score_type: bi-encoder models)
-        - "token_embed" task (score_type: late interaction models)
+        Scoring API handles score/rerank for:\n
+        - "classify" task (score_type: cross-encoder models)\n
+        - "embed" task (score_type: bi-encoder models)\n
+        - "token_embed" task (score_type: late interaction models)\n
        """
        # fixme: self._model_info.score_type is the score type before
        #  as_seq_cls_model, which is "bi-encoder", rather than the
@@ -45,10 +45,10 @@ class ProfilerConfig:
    worker's traces (CPU & GPU) will be saved under this directory. Note that
    it must be an absolute path."""

-    torch_profiler_with_stack: bool = False
-    """If `True`, enables stack tracing in the torch profiler. Disabled by default
-    to reduce overhead. Can be enabled via VLLM_TORCH_PROFILER_WITH_STACK=1 env var
-    or --profiler-config.torch_profiler_with_stack=true CLI flag."""
+    torch_profiler_with_stack: bool = True
+    """If `True`, enables stack tracing in the torch profiler. Enabled by default
+    as it is useful for debugging. Can be disabled via 
+    --profiler-config.torch_profiler_with_stack=false CLI flag."""

    torch_profiler_with_flops: bool = False
    """If `True`, enables FLOPS counting in the torch profiler. Disabled by default."""
@@ -10,6 +10,7 @@ import vllm.envs as envs
 from vllm.distributed import get_dp_group, get_ep_group
 from vllm.forward_context import get_forward_context
 from vllm.logger import init_logger
+from vllm.platforms import current_platform
 from vllm.utils.flashinfer import (
    has_flashinfer_nvlink_one_sided,
    has_flashinfer_nvlink_two_sided,
@@ -325,14 +326,20 @@ class DeepEPHTAll2AllManager(DeepEPAll2AllManagerBase):

        assert num_rdma_bytes is not None
        assert num_qps_per_rank is not None
-        return dict(
+        # TODO: remove platform-specific logic
+        # once ROCm DeepEP is updated with the latest APIs.
+        kwargs = dict(
            group=self.cpu_group,
            num_nvl_bytes=num_nvl_bytes,
            num_rdma_bytes=num_rdma_bytes,
            low_latency_mode=False,
            num_qps_per_rank=num_qps_per_rank,
-            explicitly_destroy=True,
        )
+        if not current_platform.is_rocm():
+            kwargs.update(
+                explicitly_destroy=True,
+            )
+        return kwargs

    def get_handle(self, kwargs):
        assert len(kwargs) == 0, (
@@ -397,16 +404,22 @@ class DeepEPLLAll2AllManager(DeepEPAll2AllManagerBase):
        )

        assert num_rdma_bytes is not None
-        return dict(
+        # TODO: remove platform-specific logic
+        # once ROCm DeepEP is updated with the latest APIs.
+        kwargs = dict(
            group=self.cpu_group,
            num_nvl_bytes=num_nvl_bytes,
            num_rdma_bytes=num_rdma_bytes,
            low_latency_mode=True,
            num_qps_per_rank=num_qps_per_rank,
-            allow_nvlink_for_low_latency_mode=True,
-            allow_mnnvl=envs.VLLM_DEEPEP_LOW_LATENCY_USE_MNNVL,
-            explicitly_destroy=True,
        )
+        if not current_platform.is_rocm():
+            kwargs.update(
+                allow_nvlink_for_low_latency_mode=True,
+                allow_mnnvl=envs.VLLM_DEEPEP_LOW_LATENCY_USE_MNNVL,
+                explicitly_destroy=True,
+            )
+        return kwargs

    def get_handle(self, kwargs):
        """
@@ -338,6 +338,7 @@ class CudaCommunicator(DeviceCommunicatorBase):

    def destroy(self):
        if self.pynccl_comm is not None:
+            self.pynccl_comm.destroy()
            self.pynccl_comm = None
        if self.ca_comm is not None:
            self.ca_comm = None
@@ -145,6 +145,13 @@ class PyNcclCommunicator:
            stream.synchronize()
            del data

+    def destroy(self):
+        if self.available and not self.disabled:
+            with torch.accelerator.device_index(self.device.index):
+                self.nccl.ncclCommDestroy(self.comm)
+            self.available = False
+            self.disabled = True
+
    def all_reduce(
        self,
        in_tensor: torch.Tensor,
@@ -145,11 +145,37 @@ class ElasticEPScalingExecutor:
            raise ValueError(f"Unknown execute method: {execute_method}")
        return method(*args, **kwargs)

+    def _set_eplb_suppressed(self, suppressed: bool) -> None:
+        self.worker.model_runner.eep_eplb_suppressed = suppressed
+        ep_group = get_standby_ep_group() or get_ep_group()
+        if ep_group.rank == 0:
+            logger.info(
+                "[Elastic EP] EPLB %s elastic scaling transition",
+                "disabled during" if suppressed else "re-enabled after",
+            )
+
+    def load_model(self) -> None:
+        (
+            expanded_physical_to_logical,
+            num_logical_experts,
+            old_num_physical_experts,
+        ) = self.receive_expert_mapping()
+        num_physical_experts = expanded_physical_to_logical.shape[1]
+        self.worker.parallel_config.eplb_config.num_redundant_experts = (
+            num_physical_experts - num_logical_experts
+        )
+        self.worker.load_model(load_dummy_weights=True)
+        self.worker.model_runner.setup_eplb_from_mapping(
+            expanded_physical_to_logical, old_num_physical_experts
+        )
+        self._set_eplb_suppressed(True)
+
    def create_standby_groups(
        self, reconfig_request: ReconfigureDistributedRequest
    ) -> None:
        self.reconfig_request = reconfig_request
        new_dp_size = reconfig_request.new_data_parallel_size
+        old_dp_size = get_dp_group().world_size
        world_size = self.worker.vllm_config.parallel_config.world_size
        new_world_size_across_dp = world_size * new_dp_size
        updated_config = copy.copy(self.worker.vllm_config)
@@ -165,11 +191,8 @@ class ElasticEPScalingExecutor:
                coord_store_port=reconfig_request.coord_store_port,
                enable_eplb=updated_config.parallel_config.enable_eplb,
            )
-        self.worker.model_runner.eep_eplb_suppressed = True
-        standby_ep_group = get_standby_ep_group()
-        assert standby_ep_group is not None
-        if standby_ep_group.rank == 0:
-            logger.info("[Elastic EP] EPLB disabled during elastic scaling transition")
+        if new_dp_size > old_dp_size:
+            self._set_eplb_suppressed(True)

    def transfer_weights(self, old_dp_size: int, new_dp_size: int) -> None:
        standby_dp_group = get_standby_dp_group()
@@ -237,13 +260,31 @@ class ElasticEPScalingExecutor:
            device=self.worker.device,
        )

+    def _release_cuda_graphs(self) -> None:
+        if isinstance(self.worker.model_runner.model, CUDAGraphWrapper):
+            wrapper = self.worker.model_runner.model
+            wrapper.concrete_cudagraph_entries = {}
+
+        elif isinstance(self.worker.model_runner.model, UBatchWrapper):
+            raise RuntimeError("DBO is not yet supported in elastic EP")
+
+        torch.compiler.reset()
+        with set_current_vllm_config(self.worker.vllm_config):
+            reset_compile_wrapper(self.worker.model_runner.get_model())
+
+        gc.collect()
+        torch.accelerator.synchronize()
+        torch.accelerator.empty_cache()
+
    def switch_and_remove(self) -> None:
+        self._release_cuda_graphs()
        _replace_active_groups(world=None, dp=None, ep=None, eplb=None, node_count=None)

    def switch_and_prepare(self) -> None:
        old_dp_size = get_dp_group().world_size
        old_ep_size = get_ep_group().world_size

+        self._release_cuda_graphs()
        _replace_active_groups(**pop_standby_groups())

        parallel_config = self.worker.vllm_config.parallel_config
@@ -384,13 +425,6 @@ class ElasticEPScalingExecutor:
            compilation_counter.stock_torch_compile_count += 1
            self.worker.model_runner.model.compile(fullgraph=True, backend=backend)

-        # release all previously captured CUDA graphs
-        if isinstance(self.worker.model_runner.model, CUDAGraphWrapper):
-            wrapper = self.worker.model_runner.model
-            wrapper.concrete_cudagraph_entries = {}
-        elif isinstance(self.worker.model_runner.model, UBatchWrapper):
-            raise RuntimeError("DBO is not yet supported in elastic EP")
-
        multi_block_table = self.worker.model_runner.input_batch.block_table
        saved_block_tables: list[tuple[torch.Tensor, torch.Tensor]] = []
        for bt in multi_block_table.block_tables:
@@ -399,14 +433,6 @@ class ElasticEPScalingExecutor:
            )
        multi_block_table.clear()

-        # reset the compile wrapper
-        torch.compiler.reset()
-        with set_current_vllm_config(self.worker.vllm_config):
-            reset_compile_wrapper(self.worker.model_runner.get_model())
-
-        gc.collect()
-        torch.accelerator.synchronize()
-        torch.accelerator.empty_cache()
        unlock_workspace()
        self.worker.compile_or_warm_up_model()
        lock_workspace()
@@ -416,8 +442,12 @@ class ElasticEPScalingExecutor:
        ):
            bt.block_table.gpu.copy_(saved_gpu)
            bt.block_table.cpu.copy_(saved_cpu)
+        if new_dp_size < old_dp_size:
+            self._set_eplb_suppressed(False)

-    def perform_eplb_reshuffle(self, new_dp_size: int | None = None) -> None:
+    def _perform_eplb_reshuffle(
+        self, rank_mapping: dict[int, int] | None = None
+    ) -> None:
        if get_ep_group().rank == 0:
            logger.info("[Elastic EP] Starting expert resharding...")

@@ -428,20 +458,9 @@ class ElasticEPScalingExecutor:
        eplb_model_state = eplb_state.model_states[model_config.compute_hash()]
        is_async_enabled = eplb_state.is_async
        eplb_state.is_async = False
-        if new_dp_size is None:
+        if rank_mapping is None:
            eplb_state.rearrange()
        else:
-            # scale down
-            parallel_config = self.worker.vllm_config.parallel_config
-            tp_size = parallel_config.tensor_parallel_size
-            old_ep_size = parallel_config.data_parallel_size * tp_size
-            new_ep_size = new_dp_size * tp_size
-
-            rank_mapping = {
-                old_ep_rank: old_ep_rank if old_ep_rank < new_ep_size else -1
-                for old_ep_rank in range(old_ep_size)
-            }
-
            eplb_state.rearrange(rank_mapping=rank_mapping)
        # NOTE(yongji): check whether we need to synchronize here
        torch.accelerator.synchronize()
@@ -451,10 +470,25 @@ class ElasticEPScalingExecutor:
            eplb_model_state.physical_to_logical_map.shape[1]
        )
        eplb_state.is_async = is_async_enabled
-        self.worker.model_runner.eep_eplb_suppressed = False
        if get_ep_group().rank == 0:
            logger.info("[Elastic EP] Expert resharding completed")

+    def perform_eplb_reshuffle(self) -> None:
+        self._perform_eplb_reshuffle()
+        self._set_eplb_suppressed(False)
+
+    def perform_scale_down_eplb_reshuffle(self, new_dp_size: int) -> None:
+        self._set_eplb_suppressed(True)
+        parallel_config = self.worker.vllm_config.parallel_config
+        tp_size = parallel_config.tensor_parallel_size
+        old_ep_size = parallel_config.data_parallel_size * tp_size
+        new_ep_size = new_dp_size * tp_size
+        rank_mapping = {
+            old_ep_rank: old_ep_rank if old_ep_rank < new_ep_size else -1
+            for old_ep_rank in range(old_ep_size)
+        }
+        self._perform_eplb_reshuffle(rank_mapping=rank_mapping)
+
    def receive_weights(self) -> None:
        dp_group = get_dp_group()
        assert isinstance(dp_group, StatelessGroupCoordinator)
@@ -43,9 +43,10 @@ class ScaleUpExistingEngineState(enum.IntEnum):


 class ScaleUpNewEngineState(enum.IntEnum):
-    PREPARE = 0
-    EPLB_RESHUFFLE = 1
-    COMPLETE = 2
+    PRE_KV_INIT = 0
+    PREPARE = 1
+    EPLB_RESHUFFLE = 2
+    COMPLETE = 3


 class ScaleDownRemainingEngineState(enum.IntEnum):
@@ -104,7 +105,7 @@ class ElasticEPScalingState:
        self.state: EngineState
        if scale_type == "scale_up":
            self.state = (
-                ScaleUpNewEngineState.PREPARE
+                ScaleUpNewEngineState.PRE_KV_INIT
                if worker_type == "new"
                else ScaleUpExistingEngineState.WAIT_NEW_CORE_ENGINES_INIT
            )
@@ -142,6 +143,12 @@ class ElasticEPScalingState:
            else self._progress_remaining_engine()
        )

+    def run_pre_kv_init_states(self) -> None:
+        assert self.scale_type == "scale_up" and self.worker_type == "new"
+        assert self.state == ScaleUpNewEngineState.PRE_KV_INIT
+        assert self.progress()
+        assert self.state == ScaleUpNewEngineState.PREPARE
+
    def _execute_tcp_store_barrier(
        self, dp_store, group_rank, group_size, barrier_id, timeout=None
    ):
@@ -303,7 +310,23 @@ class ElasticEPScalingState:
        state = self.state
        assert self.new_dp_group is not None and self.new_dp_store is not None

-        if state == ScaleUpNewEngineState.PREPARE:
+        if state == ScaleUpNewEngineState.PRE_KV_INIT:
+            self.engine_core._eep_send_engine_core_notification(
+                EEPNotificationType.NEW_CORE_ENGINES_WEIGHTS_INIT_READY
+            )
+            self.model_executor.collective_rpc(
+                "elastic_ep_execute", args=("receive_weights",)
+            )
+            self.engine_core.available_gpu_memory_for_kv_cache = (
+                ParallelConfig.sync_kv_cache_memory_size(self.new_dp_group, -1)
+            )
+            self.model_executor.collective_rpc(
+                "elastic_ep_execute", args=("prepare_new_worker",)
+            )
+            self.state = ScaleUpNewEngineState.PREPARE
+            return True
+
+        elif state == ScaleUpNewEngineState.PREPARE:
            tensor = torch.tensor([0, 0, 0], dtype=torch.int32, device="cpu")
            torch.distributed.all_reduce(
                tensor,
@@ -403,7 +426,6 @@ class ElasticEPScalingState:
            self.engine_core._eep_send_engine_core_notification(
                EEPNotificationType.SHUTDOWN_COMPLETE
            )
-            self.engine_core.shutdown()
            return True

        else:
@@ -525,7 +547,7 @@ class ElasticEPScalingState:
        self.model_executor.collective_rpc(
            "elastic_ep_execute",
            args=(
-                "perform_eplb_reshuffle",
+                "perform_scale_down_eplb_reshuffle",
                self.reconfig_request.new_data_parallel_size,
            ),
        )
@@ -126,28 +126,17 @@ class KVConnectorPromMetrics:
        self._labelnames = labelnames
        self.per_engine_labelvalues = per_engine_labelvalues

-    def make_per_engine(self, metric: PromMetric) -> dict[int, PromMetric]:
-        """
-        Create a per-engine child of a prometheus_client.Metric with
-        the appropriate labels set. The parent metric must be created
-        using the labelnames list.
-        """
-        return {
-            idx: metric.labels(*labelvalues)
-            for idx, labelvalues in self.per_engine_labelvalues.items()
-        }
-
    def observe(self, transfer_stats_data: dict[str, Any], engine_idx: int = 0):
        """
        Record the supplied transfer statistics to Prometheus metrics. These
        statistics are engine-specific, and should be recorded to a metric
        with the appropriate 'engine' label. These metric instances can be
-        created using the make_per_engine() helper method.
+        created using the create_metric_per_engine() helper method.
        """
        raise NotImplementedError


-class KVConnectorPrometheus:
+class KVConnectorProm:
    """
    Support for registering per-connector Prometheus metrics, and
    recording transfer statistics to those metrics. Uses
@@ -1396,9 +1396,6 @@ class MoRIIOConnectorWorker:
            remote_ip=meta.remote_host,
        )

-    def _is_last_layer(self, layer_name):
-        return layer_name == list(self.kv_caches.keys())[-1]
-
    def merge_contiguous_blocks(
        self,
        offsets_local: list[int],
@@ -65,6 +65,7 @@ from vllm.v1.kv_cache_interface import (
    SlidingWindowSpec,
    UniformTypeKVCacheSpecs,
 )
+from vllm.v1.metrics.utils import create_metric_per_engine
 from vllm.v1.worker.block_table import BlockTable
 from vllm.v1.worker.utils import select_common_block_size

@@ -3057,7 +3058,9 @@ class NixlPromMetrics(KVConnectorPromMetrics):
            buckets=buckets[1:],
            labelnames=labelnames,
        )
-        self.nixl_histogram_xfer_time = self.make_per_engine(nixl_histogram_xfer_time)
+        self.nixl_histogram_xfer_time = create_metric_per_engine(
+            nixl_histogram_xfer_time, self.per_engine_labelvalues
+        )
        nixl_histogram_post_time = self._histogram_cls(
            name="vllm:nixl_post_time_seconds",
            documentation="Histogram of transfer post time for NIXL KV"
@@ -3065,7 +3068,9 @@ class NixlPromMetrics(KVConnectorPromMetrics):
            buckets=buckets,
            labelnames=labelnames,
        )
-        self.nixl_histogram_post_time = self.make_per_engine(nixl_histogram_post_time)
+        self.nixl_histogram_post_time = create_metric_per_engine(
+            nixl_histogram_post_time, self.per_engine_labelvalues
+        )
        # uniform 2kb to 16gb range
        buckets = [2 ** (10 + i) for i in range(1, 25, 2)]
        nixl_histogram_bytes_transferred = self._histogram_cls(
@@ -3074,8 +3079,8 @@ class NixlPromMetrics(KVConnectorPromMetrics):
            buckets=buckets,
            labelnames=labelnames,
        )
-        self.nixl_histogram_bytes_transferred = self.make_per_engine(
-            nixl_histogram_bytes_transferred
+        self.nixl_histogram_bytes_transferred = create_metric_per_engine(
+            nixl_histogram_bytes_transferred, self.per_engine_labelvalues
        )
        buckets = [
            10,
@@ -3100,24 +3105,24 @@ class NixlPromMetrics(KVConnectorPromMetrics):
            buckets=buckets,
            labelnames=labelnames,
        )
-        self.nixl_histogram_num_descriptors = self.make_per_engine(
-            nixl_histogram_num_descriptors
+        self.nixl_histogram_num_descriptors = create_metric_per_engine(
+            nixl_histogram_num_descriptors, self.per_engine_labelvalues
        )
        counter_nixl_num_failed_transfers = self._counter_cls(
            name="vllm:nixl_num_failed_transfers",
            documentation="Number of failed NIXL KV Cache transfers.",
            labelnames=labelnames,
        )
-        self.counter_nixl_num_failed_transfers = self.make_per_engine(
-            counter_nixl_num_failed_transfers
+        self.counter_nixl_num_failed_transfers = create_metric_per_engine(
+            counter_nixl_num_failed_transfers, self.per_engine_labelvalues
        )
        counter_nixl_num_failed_notifications = self._counter_cls(
            name="vllm:nixl_num_failed_notifications",
            documentation="Number of failed NIXL KV Cache notifications.",
            labelnames=labelnames,
        )
-        self.counter_nixl_num_failed_notifications = self.make_per_engine(
-            counter_nixl_num_failed_notifications
+        self.counter_nixl_num_failed_notifications = create_metric_per_engine(
+            counter_nixl_num_failed_notifications, self.per_engine_labelvalues
        )

        counter_nixl_num_kv_expired_reqs = self._counter_cls(
@@ -3126,8 +3131,8 @@ class NixlPromMetrics(KVConnectorPromMetrics):
            "NOTE: This metric is tracked on the P instance.",
            labelnames=labelnames,
        )
-        self.counter_nixl_num_kv_expired_reqs = self.make_per_engine(
-            counter_nixl_num_kv_expired_reqs
+        self.counter_nixl_num_kv_expired_reqs = create_metric_per_engine(
+            counter_nixl_num_kv_expired_reqs, self.per_engine_labelvalues
        )

    def observe(self, transfer_stats_data: dict[str, Any], engine_idx: int = 0):
@@ -1477,9 +1477,9 @@ class LLM:
            data_1 = data_1 * len(data_2)

        if pooling_params is None:
-            pooling_params = PoolingParams(task="score")
+            pooling_params = PoolingParams(task="classify")
        elif pooling_params.task is None:
-            pooling_params.task = "score"
+            pooling_params.task = "classify"

        pooling_params_list = list[PoolingParams]()

@@ -22,7 +22,7 @@ from fastapi.middleware.cors import CORSMiddleware
 from starlette.datastructures import State

 import vllm.envs as envs
-from vllm.config import VllmConfig
+from vllm.config import ModelConfig, VllmConfig
 from vllm.engine.arg_utils import AsyncEngineArgs
 from vllm.engine.protocol import EngineClient
 from vllm.entrypoints.chat_utils import load_chat_template
@@ -155,7 +155,9 @@ async def build_async_engine_client_from_engine_args(


 def build_app(
-    args: Namespace, supported_tasks: tuple["SupportedTask", ...] | None = None
+    args: Namespace,
+    supported_tasks: tuple["SupportedTask", ...] | None = None,
+    model_config: ModelConfig | None = None,
 ) -> FastAPI:
    if supported_tasks is None:
        warnings.warn(
@@ -191,7 +193,7 @@ def build_app(
        attach_router as register_sagemaker_api_router,
    )

-    register_sagemaker_api_router(app, supported_tasks)
+    register_sagemaker_api_router(app, supported_tasks, model_config)

    if "generate" in supported_tasks:
        from vllm.entrypoints.openai.generate.api_router import (
@@ -242,7 +244,7 @@ def build_app(
    if any(task in POOLING_TASKS for task in supported_tasks):
        from vllm.entrypoints.pooling import register_pooling_api_routers

-        register_pooling_api_routers(app, supported_tasks)
+        register_pooling_api_routers(app, supported_tasks, model_config)

    app.root_path = args.root_path
    app.add_middleware(
@@ -583,8 +585,10 @@ async def build_and_serve(
        uvicorn_kwargs["log_config"] = log_config

    supported_tasks = await engine_client.get_supported_tasks()
+    model_config = engine_client.model_config
+
    logger.info("Supported tasks: %s", supported_tasks)
-    app = build_app(args, supported_tasks)
+    app = build_app(args, supported_tasks, model_config)
    await init_app_state(engine_client, app.state, args, supported_tasks)

    logger.info("Starting vLLM server on %s", listen_address)
@@ -9,7 +9,7 @@ from abc import ABC, abstractmethod
 from collections.abc import Callable
 from contextlib import AsyncExitStack
 from dataclasses import replace
-from typing import TYPE_CHECKING, Final, Union
+from typing import TYPE_CHECKING, Any, Final, Union

 from openai.types.responses.response_function_tool_call_output_item import (
    ResponseFunctionToolCallOutputItem,
@@ -182,6 +182,7 @@ class SimpleContext(ConversationContext):
        self.all_turn_metrics = []

        self.input_messages: list[ResponseRawMessageAndToken] = []
+        self.kv_transfer_params: dict[str, Any] | None = None

    def append_output(self, output) -> None:
        self.last_output = output
@@ -190,6 +191,8 @@ class SimpleContext(ConversationContext):
        self.num_prompt_tokens = len(output.prompt_token_ids or [])
        self.num_cached_tokens = output.num_cached_tokens or 0
        self.num_output_tokens += len(output.outputs[0].token_ids or [])
+        if output.kv_transfer_params is not None:
+            self.kv_transfer_params = output.kv_transfer_params

        # Accumulate text, token_ids, and logprobs for streaming mode
        delta_output = output.outputs[0]
@@ -308,11 +311,14 @@ class ParsableContext(ConversationContext):
        self.input_messages: list[ResponseRawMessageAndToken] = []
        self.output_messages: list[ResponseRawMessageAndToken] = []
        self._accumulated_token_ids: list[int] = []
+        self.kv_transfer_params: dict[str, Any] | None = None

    def append_output(self, output: RequestOutput) -> None:
        self.num_prompt_tokens = len(output.prompt_token_ids or [])
        self.num_cached_tokens = output.num_cached_tokens or 0
        self.num_output_tokens += len(output.outputs[0].token_ids or [])
+        if output.kv_transfer_params is not None:
+            self.kv_transfer_params = output.kv_transfer_params
        self.parser.process(output.outputs[0])
        output_token_ids = output.outputs[0].token_ids or []
        self._accumulated_token_ids.extend(output_token_ids)
@@ -538,6 +544,7 @@ class HarmonyContext(ConversationContext):
        self.all_turn_metrics: list[TurnMetrics] = []
        self.is_first_turn = True
        self.first_tok_of_message = True  # For streaming support
+        self.kv_transfer_params: dict[str, Any] | None = None

    def _update_num_reasoning_tokens(self):
        channel = self.parser.current_channel
@@ -557,6 +564,8 @@ class HarmonyContext(ConversationContext):
            self._update_num_reasoning_tokens()
        self._update_prefill_token_usage(output)
        self._update_decode_token_usage(output)
+        if output.kv_transfer_params is not None:
+            self.kv_transfer_params = output.kv_transfer_params
        # Append current turn to all turn list for next turn's calculations
        self.all_turn_metrics.append(self.current_turn_metrics.copy())
        self.current_turn_metrics.reset()
@@ -868,6 +877,8 @@ class StreamingHarmonyContext(HarmonyContext):
        if last_delta_text:
            self.last_content_delta = last_delta_text
        self._update_decode_token_usage(output)
+        if output.kv_transfer_params is not None:
+            self.kv_transfer_params = output.kv_transfer_params

        # For streaming, update previous turn when message is complete
        if output.finished:
@@ -252,6 +252,10 @@ class ResponsesRequest(OpenAIBaseModel):
            "numeric values, used by custom extensions."
        ),
    )
+    kv_transfer_params: dict[str, Any] | None = Field(
+        default=None,
+        description="KVTransfer parameters used for disaggregated serving.",
+    )
    # --8<-- [end:responses-extra-params]

    def build_chat_params(
@@ -351,6 +355,10 @@ class ResponsesRequest(OpenAIBaseModel):
        if isinstance(stop, str):
            stop = [stop]

+        extra_args: dict[str, Any] = self.vllm_xargs if self.vllm_xargs else {}
+        if self.kv_transfer_params:
+            extra_args["kv_transfer_params"] = self.kv_transfer_params
+
        return SamplingParams.from_optional(
            temperature=temperature,
            top_p=top_p,
@@ -367,7 +375,7 @@ class ResponsesRequest(OpenAIBaseModel):
            ),
            structured_outputs=structured_outputs,
            logit_bias=self.logit_bias,
-            extra_args=self.vllm_xargs or {},
+            extra_args=extra_args,
            skip_clone=True,  # Created fresh per request, safe to skip clone
            skip_special_tokens=self.skip_special_tokens,
            include_stop_str_in_output=self.include_stop_str_in_output,
@@ -488,6 +496,11 @@ class ResponsesResponse(OpenAIBaseModel):
    usage: ResponseUsage | None = None
    user: str | None = None

+    # vLLM-specific fields that are not in OpenAI spec
+    kv_transfer_params: dict[str, Any] | None = Field(
+        default=None, description="KVTransfer parameters."
+    )
+
    # --8<-- [start:responses-response-extra-params]
    # These are populated when enable_response_messages is set to True
    # NOTE: custom serialization is needed
@@ -531,6 +544,7 @@ class ResponsesResponse(OpenAIBaseModel):
        usage: ResponseUsage | None = None,
        input_messages: ResponseInputOutputMessage | None = None,
        output_messages: ResponseInputOutputMessage | None = None,
+        kv_transfer_params: dict[str, Any] | None = None,
    ) -> "ResponsesResponse":
        incomplete_details: IncompleteDetails | None = None
        if status == "incomplete":
@@ -566,6 +580,7 @@ class ResponsesResponse(OpenAIBaseModel):
            truncation=request.truncation,
            user=request.user,
            usage=usage,
+            kv_transfer_params=kv_transfer_params,
        )


@@ -873,6 +873,7 @@ class OpenAIServingResponses(OpenAIServing):
            output=output,
            status=status,
            usage=usage,
+            kv_transfer_params=context.kv_transfer_params,
        )

        if request.store:
@@ -42,32 +42,13 @@ from vllm.inputs import EncoderDecoderInputs, ProcessorInputs
 from vllm.logger import init_logger
 from vllm.logprobs import FlatLogprobs, Logprob
 from vllm.model_executor.models import SupportsTranscription
-from vllm.multimodal.audio import split_audio
-from vllm.multimodal.media.audio import extract_audio_from_video_bytes
+from vllm.multimodal.audio import get_audio_duration, split_audio
+from vllm.multimodal.media.audio import load_audio
 from vllm.outputs import RequestOutput
 from vllm.renderers.inputs import DictPrompt, EncoderDecoderDictPrompt
 from vllm.renderers.inputs.preprocess import parse_enc_dec_prompt, parse_model_prompt
 from vllm.sampling_params import BeamSearchParams, SamplingParams
 from vllm.tokenizers import get_tokenizer
-from vllm.utils.import_utils import PlaceholderModule
-
-try:
-    import librosa
-except ImportError:
-    librosa = PlaceholderModule("librosa")  # type: ignore[assignment]
-
-try:
-    import soundfile as sf
-except ImportError:
-    sf = PlaceholderModule("soundfile")  # type: ignore[assignment]
-
-# Public libsndfile error codes exposed via `soundfile.LibsndfileError.code`, soundfile
-# being librosa's main backend. Used to validate if an audio loading error is due to a
-# server error vs a client error (invalid audio file).
-# 1 = unrecognised format      (file is not a supported audio container)
-# 3 = malformed file           (corrupt or structurally invalid audio)
-# 4 = unsupported encoding     (codec not supported by this libsndfile build)
-_BAD_SF_CODES = {1, 3, 4}

 SpeechToTextResponse: TypeAlias = TranscriptionResponse | TranslationResponse
 SpeechToTextResponseVerbose: TypeAlias = (
@@ -214,32 +195,13 @@ class OpenAISpeechToText(OpenAIServing):
        # pre-requisite for chunking, as it assumes Whisper SR.
        try:
            with io.BytesIO(audio_data) as buf:
-                y, sr = librosa.load(buf, sr=self.asr_config.sample_rate)  # type: ignore[return-value]
-        except sf.LibsndfileError as exc:
-            # Only fall back for known format-detection failures.
-            # Re-raise anything else (e.g. corrupt but recognised format).
-            if exc.code not in _BAD_SF_CODES:
-                raise
-            logger.debug(
-                "librosa/soundfile could not decode audio from BytesIO "
-                "(code=%s: %s); falling back to pyav in-process decode",
-                exc.code,
-                exc,
-            )
-            try:
-                native_y, native_sr = extract_audio_from_video_bytes(audio_data)
-                sr = self.asr_config.sample_rate
-                y = librosa.resample(native_y, orig_sr=native_sr, target_sr=sr)
-            except Exception as pyav_exc:
-                logger.debug(
-                    "pyAV fallback also failed: %s",
-                    pyav_exc,
-                )
-                raise ValueError("Invalid or unsupported audio file.") from pyav_exc
+                y, sr = load_audio(buf, sr=self.asr_config.sample_rate)
+        except Exception as exc:
+            raise ValueError("Invalid or unsupported audio file.") from exc

-        duration = librosa.get_duration(y=y, sr=sr)
-        do_split_audio = (
-            self.asr_config.allow_audio_chunking
+        duration = get_audio_duration(y=y, sr=sr)
+        do_split_audio = self.asr_config.allow_audio_chunking and (
+            self.asr_config.max_audio_clip_s is not None
            and duration > self.asr_config.max_audio_clip_s
        )

@@ -5,6 +5,9 @@ from typing import TYPE_CHECKING

 from fastapi import FastAPI

+from vllm.config import ModelConfig
+from vllm.logger import init_logger
+
 if TYPE_CHECKING:
    from argparse import Namespace

@@ -17,9 +20,30 @@ else:
    RequestLogger = object
    SupportedTask = object

+logger = init_logger(__name__)
+
+
+def enable_scoring_api(
+    supported_tasks: tuple["SupportedTask", ...],
+    model_config: ModelConfig | None = None,
+) -> bool:
+    if any(t in supported_tasks for t in ("embed", "token_embed")):
+        return True
+
+    if model_config is not None and "classify" in supported_tasks:
+        num_labels = getattr(model_config.hf_config, "num_labels", 0)
+        if num_labels != 1:
+            logger.debug_once("Score API is only enabled for num_labels == 1.")
+            return False
+        return True
+
+    return False
+

 def register_pooling_api_routers(
-    app: FastAPI, supported_tasks: tuple["SupportedTask", ...]
+    app: FastAPI,
+    supported_tasks: tuple["SupportedTask", ...],
+    model_config: ModelConfig | None = None,
 ):
    from vllm.entrypoints.pooling.pooling.api_router import router as pooling_router

@@ -37,11 +61,7 @@ def register_pooling_api_routers(

        app.include_router(embed_router)

-    # Score API handles score/rerank for:
-    # - "score" task (score_type: cross-encoder models)
-    # - "embed" task (score_type: bi-encoder models)
-    # - "token_embed" task (score_type: late interaction models)
-    if any(t in supported_tasks for t in ("score", "embed", "token_embed")):
+    if enable_scoring_api(supported_tasks, model_config):
        from vllm.entrypoints.pooling.score.api_router import router as score_router

        app.include_router(score_router)
@@ -61,6 +81,8 @@ def init_pooling_state(
    from vllm.entrypoints.pooling.score.serving import ServingScores
    from vllm.tasks import POOLING_TASKS

+    model_config = engine_client.model_config
+
    resolved_chat_template = load_chat_template(args.chat_template)

    state.serving_pooling = (
@@ -102,10 +124,6 @@ def init_pooling_state(
        if "classify" in supported_tasks
        else None
    )
-    # Score API handles score/rerank for:
-    # - "score" task (score_type: cross-encoder models)
-    # - "embed" task (score_type: bi-encoder models)
-    # - "token_embed" task (score_type: late interaction models)
    state.serving_scores = (
        ServingScores(
            engine_client,
@@ -114,6 +132,6 @@ def init_pooling_state(
            score_template=resolved_chat_template,
            log_error_stack=args.log_error_stack,
        )
-        if any(t in supported_tasks for t in ("embed", "score", "token_embed"))
+        if enable_scoring_api(supported_tasks, model_config)
        else None
    )
@@ -35,7 +35,7 @@ class ScoreRequestMixin(PoolingBasicRequestMixin, ClassifyRequestMixin):
            max_total_tokens_param="max_model_len",
        )

-    def to_pooling_params(self, task: PoolingTask = "score"):
+    def to_pooling_params(self, task: PoolingTask = "classify"):
        return PoolingParams(
            task=task,
            use_activation=self.use_activation,
@@ -111,7 +111,7 @@ class RerankRequest(PoolingBasicRequestMixin, ClassifyRequestMixin):
            max_total_tokens_param="max_model_len",
        )

-    def to_pooling_params(self, task: PoolingTask = "score"):
+    def to_pooling_params(self, task: PoolingTask = "classify"):
        return PoolingParams(
            task=task,
            use_activation=self.use_activation,
@@ -413,7 +413,7 @@ class ServingScores(OpenAIServing):
        # Schedule the request and get the result generator.
        generators: list[AsyncGenerator[PoolingRequestOutput, None]] = []

-        default_pooling_params = request.to_pooling_params("score")
+        default_pooling_params = request.to_pooling_params("classify")

        for i, engine_prompt in enumerate(engine_prompts):
            request_id_item = f"{request_id}-{i}"
@@ -10,9 +10,11 @@ import pydantic
 from fastapi import APIRouter, Depends, FastAPI, HTTPException, Request
 from fastapi.responses import JSONResponse, Response

+from vllm.config import ModelConfig
 from vllm.entrypoints.openai.engine.protocol import ErrorResponse
 from vllm.entrypoints.openai.engine.serving import OpenAIServing
 from vllm.entrypoints.openai.utils import validate_json_request
+from vllm.entrypoints.pooling import enable_scoring_api
 from vllm.entrypoints.pooling.base.serving import PoolingServing
 from vllm.entrypoints.serve.instrumentator.basic import base
 from vllm.entrypoints.serve.instrumentator.health import health
@@ -25,7 +27,10 @@ GetHandlerFn = Callable[[Request], OpenAIServing | PoolingServing | None]
 EndpointFn = Callable[[RequestType, Request], Awaitable[Any]]


-def get_invocation_types(supported_tasks: tuple["SupportedTask", ...]):
+def get_invocation_types(
+    supported_tasks: tuple["SupportedTask", ...],
+    model_config: ModelConfig | None = None,
+):
    # NOTE: Items defined earlier take higher priority
    INVOCATION_TYPES: list[tuple[RequestType, tuple[GetHandlerFn, EndpointFn]]] = []

@@ -70,7 +75,7 @@ def get_invocation_types(supported_tasks: tuple["SupportedTask", ...]):
            (ClassificationRequest, (classify, create_classify)),
        ]

-    if "score" in supported_tasks:
+    if enable_scoring_api(supported_tasks, model_config):
        from vllm.entrypoints.pooling.score.api_router import do_rerank, rerank
        from vllm.entrypoints.pooling.score.protocol import RerankRequest

@@ -78,7 +83,6 @@ def get_invocation_types(supported_tasks: tuple["SupportedTask", ...]):
            (RerankRequest, (rerank, do_rerank)),
        ]

-    if "score" in supported_tasks or "embed" in supported_tasks:
        from vllm.entrypoints.pooling.score.api_router import create_score, score
        from vllm.entrypoints.pooling.score.protocol import ScoreRequest

@@ -97,11 +101,15 @@ def get_invocation_types(supported_tasks: tuple["SupportedTask", ...]):
    return INVOCATION_TYPES


-def attach_router(app: FastAPI, supported_tasks: tuple["SupportedTask", ...]):
+def attach_router(
+    app: FastAPI,
+    supported_tasks: tuple["SupportedTask", ...],
+    model_config: ModelConfig | None = None,
+):
    router = APIRouter()

    # NOTE: Construct the TypeAdapters only once
-    INVOCATION_TYPES = get_invocation_types(supported_tasks)
+    INVOCATION_TYPES = get_invocation_types(supported_tasks, model_config)
    INVOCATION_VALIDATORS = [
        (pydantic.TypeAdapter(request_type), (get_handler, endpoint))
        for request_type, (get_handler, endpoint) in INVOCATION_TYPES
@@ -124,6 +124,14 @@ class ConchLinearKernel(MPLinearKernel):

        w_q, w_s, w_zp, _ = self._get_weight_params(layer)

+        # Map channelwise group_size=-1 to the actual input dimension K.
+        # The conch kernel computes stride_mul = block_k / group_size;
+        # passing -1 produces a negative stride that reads out-of-bounds
+        # scale values for all K-blocks after the first.
+        group_size = self.config.group_size
+        if group_size == -1:
+            group_size = x.shape[-1]
+
        output = mixed_precision_gemm(
            x=x,
            w_q_packed=w_q.data,
@@ -131,7 +139,7 @@ class ConchLinearKernel(MPLinearKernel):
            w_zp=w_zp.data if w_zp is not None else None,
            weight_size_bits=self.config.weight_type.size_bits,
            weight_bias=self.config.weight_type.bias,
-            group_size=self.config.group_size,
+            group_size=group_size,
        )

        if bias is not None:
@@ -59,6 +59,13 @@ class ExllamaLinearKernel(MPLinearKernel):
                f"{cls.SUPPORTED_QUANT_TYPES}",
            )

+        if c.group_size <= 0:
+            return (
+                False,
+                f"Group size ({c.group_size}) must be positive, "
+                "Exllama does not support channelwise quantization",
+            )
+
        if c.full_weight_shape[0] % c.group_size != 0:
            return (
                False,
@@ -346,7 +346,7 @@ class FusedMoEQuantConfig:

    @property
    def use_fp8_w8a8(self) -> bool:
-        return self.quant_dtype == torch.float8_e4m3fn
+        return self.quant_dtype == current_platform.fp8_dtype()

    @property
    def use_int8_w8a8(self) -> bool:
@@ -566,7 +566,7 @@ def fp8_w8a8_moe_quant_config(
    Construct a quant config for fp8 activations and fp8 weights.
    """
    return FusedMoEQuantConfig.make(
-        torch.float8_e4m3fn,
+        current_platform.fp8_dtype(),
        w1_scale=w1_scale,
        g1_alphas=g1_alphas,
        w2_scale=w2_scale,
@@ -16,6 +16,7 @@ from vllm.model_executor.layers.fused_moe.utils import (
    moe_kernel_quantize_input,
    normalize_batched_scales_shape,
 )
+from vllm.platforms import current_platform
 from vllm.v1.worker.ubatching import (
    dbo_current_ubatch_id,
    dbo_enabled,
@@ -158,11 +159,6 @@ class DeepEPLLPrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):
            return topk_ids
        return self.global_to_physical[topk_ids]

-    def _map_local_to_global_ids(self, expert_topk_ids: torch.Tensor) -> torch.Tensor:
-        if self.local_expert_global_ids is None:
-            return expert_topk_ids
-        return self.local_expert_global_ids[expert_topk_ids]
-
    def _do_quant(
        self,
        x: torch.Tensor | tuple[torch.Tensor, torch.Tensor],
@@ -295,23 +291,46 @@ class DeepEPLLPrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):

        # Dispatch
        dispatch_topk_ids = self._map_global_to_physical_ids(topk_ids)
-        expert_x, expert_num_tokens, handle, _, hook = self.buffer.low_latency_dispatch(
-            a1,
-            dispatch_topk_ids,
-            self.max_tokens_per_rank,
-            num_experts,
-            use_fp8=self.use_fp8_dispatch,
-            round_scale=self.use_ue8m0_dispatch,
-            use_ue8m0=self.use_ue8m0_dispatch,
-            **(dict(use_nvfp4=True) if use_nvfp4 else dict()),
-            **(
-                dict(x_global_scale=qc_a1_gscale_or_scale)
-                if qc_a1_gscale_or_scale is not None
-                else dict()
-            ),
-            async_finish=False,
-            return_recv_hook=True,
-        )
+        if current_platform.is_rocm():
+            (
+                expert_x,
+                expert_num_tokens,
+                handle,
+                _,
+                hook,
+            ) = self.buffer.low_latency_dispatch(
+                a1,
+                dispatch_topk_ids,
+                self.max_tokens_per_rank,
+                num_experts,
+                use_fp8=self.use_fp8_dispatch,
+                async_finish=False,
+                return_recv_hook=True,
+            )
+        else:
+            (
+                expert_x,
+                expert_num_tokens,
+                handle,
+                _,
+                hook,
+            ) = self.buffer.low_latency_dispatch(
+                a1,
+                dispatch_topk_ids,
+                self.max_tokens_per_rank,
+                num_experts,
+                use_fp8=self.use_fp8_dispatch,
+                round_scale=self.use_ue8m0_dispatch,
+                use_ue8m0=self.use_ue8m0_dispatch,
+                **(dict(use_nvfp4=True) if use_nvfp4 else dict()),
+                **(
+                    dict(x_global_scale=qc_a1_gscale_or_scale)
+                    if qc_a1_gscale_or_scale is not None
+                    else dict()
+                ),
+                async_finish=False,
+                return_recv_hook=True,
+            )
        self.handles[a2a_idx] = handle

        return (
@@ -269,9 +269,16 @@ class TrtLlmFp8ExpertsMonolithic(TrtLlmFp8ExpertsBase, mk.FusedMoEExpertsMonolit
        weight_key: QuantKey | None,
        activation_key: QuantKey | None,
    ) -> bool:
-        """Monolithic kernels need to express router support."""
+        """Monolithic kernels need to express router support.
+        Renormalize/RenormalizeNaive are excluded: the monolithic kernel's
+        internal routing for these methods produces output uncorrelated
+        with the modular kernel's output and with Triton kernel's output
+        for Qwen3.5-35B-A3B-FP8.
+        See: https://github.com/vllm-project/vllm/issues/37591
+        """
        # NOTE(dbari): TopK routing could also be enabled, but need to validate models
        # NOTE(dbari): Default is not implemented and should not be enabled until it is
+
        if (weight_key, activation_key) in [
            (kFp8Static128BlockSym, kFp8Dynamic128Sym),
            (kMxfp8Static, kMxfp8Dynamic),
@@ -279,16 +286,12 @@ class TrtLlmFp8ExpertsMonolithic(TrtLlmFp8ExpertsBase, mk.FusedMoEExpertsMonolit
            # NOTE(rob): potentially allow others here. This is a conservative list.
            return routing_method in [
                RoutingMethodType.DeepSeekV3,
-                RoutingMethodType.Renormalize,
-                RoutingMethodType.RenormalizeNaive,
            ]
        elif (weight_key, activation_key) == (kFp8StaticTensorSym, kFp8StaticTensorSym):
            # NOTE(dbari): as above, potentially allow others here.
            return routing_method in [
                RoutingMethodType.DeepSeekV3,
                RoutingMethodType.Llama4,
-                RoutingMethodType.Renormalize,
-                RoutingMethodType.RenormalizeNaive,
            ]
        else:
            raise ValueError("Unsupported quantization scheme.")
@@ -0,0 +1,352 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import torch
+
+import vllm.model_executor.layers.fused_moe.modular_kernel as mk
+from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+from vllm.model_executor.layers.fused_moe.config import (
+    FusedMoEConfig,
+    FusedMoEParallelConfig,
+    FusedMoEQuantConfig,
+    RoutingMethodType,
+)
+from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
+    TopKWeightAndReduceNoOP,
+)
+from vllm.model_executor.layers.quantization.utils.quant_utils import (
+    QuantKey,
+    kMxfp4Static,
+    kMxfp8Dynamic,
+)
+from vllm.platforms import current_platform
+from vllm.utils.flashinfer import has_flashinfer
+
+
+class TrtLlmMxfp4ExpertsBase:
+    """
+    MXFP4 TRTLLM-Gen MoE kernels. Shared base for modular and monolithic.
+    """
+
+    def __init__(
+        self,
+        moe_config: FusedMoEConfig,
+        quant_config: FusedMoEQuantConfig,
+    ):
+        # NOTE: FusedMoEExperts.__init__ is called by the concrete subclass
+        # (Monolithic/Modular) via MRO, not here, to avoid mypy issues with
+        # multiple inheritance. This matches the NvFP4 expert pattern.
+        self.moe_config = moe_config
+        self.quant_config = quant_config
+
+        self.routing_method_type = moe_config.routing_method
+        self.topk = moe_config.experts_per_token
+        self.intermediate_size_per_partition = (
+            moe_config.intermediate_size_per_partition
+        )
+        self.hidden_dim = moe_config.hidden_dim
+        self.local_num_experts = moe_config.num_local_experts
+        self.ep_rank = moe_config.moe_parallel_config.ep_rank
+
+        # MXFP4-specific TRTLLM parameters
+        device = torch.accelerator.current_device_index()
+        self.gemm1_alpha = torch.tensor(
+            [1.702] * self.local_num_experts,
+            dtype=torch.float32,
+            device=device,
+        )
+        self.gemm1_beta = torch.tensor(
+            [1.0] * self.local_num_experts,
+            dtype=torch.float32,
+            device=device,
+        )
+        self.gemm1_clamp_limit = torch.tensor(
+            [7.0] * self.local_num_experts,
+            dtype=torch.float32,
+            device=device,
+        )
+
+        from vllm.config import get_current_vllm_config
+
+        self.max_capture_size = (
+            get_current_vllm_config().compilation_config.max_cudagraph_capture_size
+        )
+
+        # P1-5 fix: use public quant_dtype property instead of private _a1
+        self.use_mxfp8_input = quant_config.quant_dtype == "mxfp8"
+
+    @staticmethod
+    def _supports_current_device() -> bool:
+        p = current_platform
+        return p.is_cuda() and p.is_device_capability_family(100) and has_flashinfer()
+
+    @staticmethod
+    def _supports_no_act_and_mul() -> bool:
+        return False
+
+    @staticmethod
+    def _supports_quant_scheme(
+        weight_key: QuantKey | None,
+        activation_key: QuantKey | None,
+    ) -> bool:
+        SUPPORTED_W_A = [
+            (kMxfp4Static, None),
+            (kMxfp4Static, kMxfp8Dynamic),
+        ]
+        return (weight_key, activation_key) in SUPPORTED_W_A
+
+    @staticmethod
+    def _supports_activation(activation: MoEActivation) -> bool:
+        return activation == MoEActivation.SWIGLUOAI
+
+    @staticmethod
+    def activation_format() -> mk.FusedMoEActivationFormat:
+        return mk.FusedMoEActivationFormat.Standard
+
+    def supports_chunking(self) -> bool:
+        return False
+
+    def supports_expert_map(self) -> bool:
+        return False
+
+    @property
+    def expects_unquantized_inputs(self) -> bool:
+        # Expert handles MXFP8 quantization internally if needed
+        return True
+
+
+class TrtLlmMxfp4ExpertsMonolithic(
+    TrtLlmMxfp4ExpertsBase, mk.FusedMoEExpertsMonolithic
+):
+    """
+    Monolithic version of the MXFP4 TRTLLM kernel (router + experts).
+    Wraps flashinfer.trtllm_fp4_block_scale_moe().
+    """
+
+    @staticmethod
+    def _supports_parallel_config(
+        moe_parallel_config: FusedMoEParallelConfig,
+    ) -> bool:
+        return (
+            not moe_parallel_config.use_all2all_kernels
+            and not moe_parallel_config.enable_eplb
+            and moe_parallel_config.dp_size <= 1
+        )
+
+    @staticmethod
+    def _supports_routing_method(
+        routing_method: RoutingMethodType,
+        weight_key: QuantKey | None,
+        activation_key: QuantKey | None,
+    ) -> bool:
+        return routing_method in [
+            RoutingMethodType.Renormalize,
+            RoutingMethodType.RenormalizeNaive,
+        ]
+
+    @staticmethod
+    def _supports_router_logits_dtype(
+        router_logits_dtype: torch.dtype | None,
+        routing_method: RoutingMethodType,
+    ) -> bool:
+        # Kernel converts to bfloat16 internally
+        return True
+
+    def apply(
+        self,
+        hidden_states: torch.Tensor,
+        w1: torch.Tensor,
+        w2: torch.Tensor,
+        router_logits: torch.Tensor,
+        activation: MoEActivation,
+        global_num_experts: int,
+        expert_map: torch.Tensor | None,
+        a1q_scale: torch.Tensor | None,
+        apply_router_weight_on_input: bool,
+        # grouped topk + fused topk bias parameters
+        num_expert_group: int | None = None,
+        e_score_correction_bias: torch.Tensor | None = None,
+        routed_scaling_factor: float | None = None,
+        topk_group: int | None = None,
+    ) -> torch.Tensor:
+        from flashinfer import trtllm_fp4_block_scale_moe
+
+        # Handle input quantization
+        if self.use_mxfp8_input:
+            from flashinfer import mxfp8_quantize
+
+            x_quant, x_scale = mxfp8_quantize(
+                hidden_states,
+                is_sf_swizzled_layout=False,
+                alignment=256,
+            )
+            x_scale = x_scale.view(torch.float8_e4m3fn).reshape(
+                *hidden_states.shape[:-1], -1
+            )
+        else:
+            assert hidden_states.dtype == torch.bfloat16
+            x_quant = hidden_states
+            x_scale = None
+
+        output = torch.empty_like(hidden_states)
+
+        return trtllm_fp4_block_scale_moe(
+            routing_logits=router_logits.to(torch.bfloat16),
+            routing_bias=None,
+            hidden_states=x_quant,
+            hidden_states_scale=x_scale,
+            gemm1_weights=w1,
+            gemm1_weights_scale=self.w1_scale,
+            gemm1_bias=self.w1_bias,
+            gemm1_alpha=self.gemm1_alpha,
+            gemm1_beta=self.gemm1_beta,
+            gemm1_clamp_limit=self.gemm1_clamp_limit,
+            gemm2_weights=w2,
+            gemm2_weights_scale=self.w2_scale,
+            gemm2_bias=self.w2_bias,
+            output1_scale_scalar=None,
+            output1_scale_gate_scalar=None,
+            output2_scale_scalar=None,
+            num_experts=global_num_experts,
+            top_k=self.topk,
+            n_group=None,
+            topk_group=None,
+            intermediate_size=self.intermediate_size_per_partition,
+            local_expert_offset=self.ep_rank * self.local_num_experts,
+            local_num_experts=self.local_num_experts,
+            routed_scaling_factor=None,
+            routing_method_type=self.routing_method_type,
+            do_finalize=True,
+            tune_max_num_tokens=max(self.max_capture_size, 1),
+            output=output,
+        )[0]
+
+
+class TrtLlmMxfp4ExpertsModular(TrtLlmMxfp4ExpertsBase, mk.FusedMoEExpertsModular):
+    """
+    Modular version of the MXFP4 TRTLLM kernel (just the experts).
+    Wraps flashinfer.trtllm_fp4_block_scale_routed_moe().
+    Moved from trtllm_moe.py.
+    """
+
+    @property
+    def expects_unquantized_inputs(self) -> bool:
+        return True
+
+    @staticmethod
+    def _supports_parallel_config(
+        moe_parallel_config: FusedMoEParallelConfig,
+    ) -> bool:
+        return True
+
+    def supports_expert_map(self) -> bool:
+        return True
+
+    def finalize_weight_and_reduce_impl(self) -> mk.TopKWeightAndReduce:
+        return TopKWeightAndReduceNoOP()
+
+    def workspace_shapes(
+        self,
+        M: int,
+        N: int,
+        K: int,
+        topk: int,
+        global_num_experts: int,
+        local_num_experts: int,
+        expert_tokens_meta: mk.ExpertTokensMetadata | None,
+        activation: MoEActivation,
+    ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
+        # The workspaces for this implementation are managed by flashinfer.
+        workspace1 = (0,)
+        workspace2 = (0,)
+        output = (M, K)
+        return (workspace1, workspace2, output)
+
+    def apply(
+        self,
+        output: torch.Tensor,
+        hidden_states: torch.Tensor,
+        w1: torch.Tensor,
+        w2: torch.Tensor,
+        topk_weights: torch.Tensor,
+        topk_ids: torch.Tensor,
+        activation: MoEActivation,
+        global_num_experts: int,
+        expert_map: torch.Tensor | None,
+        a1q_scale: torch.Tensor | None,
+        a2_scale: torch.Tensor | None,
+        workspace13: torch.Tensor,
+        workspace2: torch.Tensor,
+        expert_tokens_meta: mk.ExpertTokensMetadata | None,
+        apply_router_weight_on_input: bool,
+    ):
+        topk = topk_ids.size(-1)
+        local_num_experts = w1.size(0)
+        intermediate_size = w2.size(1)
+        local_expert_offset = self.moe_config.ep_rank * local_num_experts
+
+        # Handle input quantization
+        if self.use_mxfp8_input:
+            from flashinfer import mxfp8_quantize
+
+            x_quant, x_scale = mxfp8_quantize(
+                hidden_states,
+                is_sf_swizzled_layout=False,
+                alignment=256,
+            )
+            x_scale = x_scale.view(torch.float8_e4m3fn).reshape(
+                *hidden_states.shape[:-1], -1
+            )
+        else:
+            assert hidden_states.dtype == torch.bfloat16
+            x_quant = hidden_states
+            x_scale = None
+
+        packed_tensor = (topk_ids.to(torch.int32) << 16) | topk_weights.to(
+            torch.bfloat16
+        ).view(torch.int16)
+
+        assert self.w1_scale is not None
+        assert self.w2_scale is not None
+        kwargs = {
+            "topk_ids": packed_tensor,
+            "routing_bias": None,
+            "hidden_states": x_quant,
+            "hidden_states_scale": x_scale,
+            "gemm1_weights": w1,
+            "gemm1_weights_scale": self.w1_scale,
+            "gemm1_bias": self.w1_bias,
+            "gemm1_alpha": self.gemm1_alpha,
+            "gemm1_beta": self.gemm1_beta,
+            "gemm1_clamp_limit": self.gemm1_clamp_limit,
+            "gemm2_weights": w2,
+            "gemm2_weights_scale": self.w2_scale,
+            "gemm2_bias": self.w2_bias,
+            "output1_scale_scalar": None,
+            "output1_scale_gate_scalar": None,
+            "output2_scale_scalar": None,
+            "num_experts": global_num_experts,
+            "top_k": topk,
+            "n_group": None,
+            "topk_group": None,
+            "intermediate_size": intermediate_size,
+            "local_expert_offset": local_expert_offset,
+            "local_num_experts": local_num_experts,
+            "routed_scaling_factor": None,
+            "routing_method_type": self.routing_method_type,
+            "do_finalize": True,
+            "output": output,
+            "tune_max_num_tokens": max(self.max_capture_size, 1),
+        }
+
+        from flashinfer import trtllm_fp4_block_scale_routed_moe
+
+        from vllm.utils.flashinfer import autotune
+
+        with autotune(False):
+            # Enable autotune when,
+            # https://github.com/flashinfer-ai/flashinfer/issues/2023 is
+            # resolved.
+            trtllm_fp4_block_scale_routed_moe(**kwargs)
+
+        return output
@@ -1017,6 +1017,7 @@ class BatchedTritonExperts(mk.FusedMoEExpertsModular):
            torch.float16,
            torch.bfloat16,
            torch.float8_e4m3fn,
+            torch.float8_e4m3fnuz,
        ]
        assert expert_tokens_meta is not None

@@ -1046,7 +1047,7 @@ class BatchedTritonExperts(mk.FusedMoEExpertsModular):
            compute_type = tl.float16
        elif hidden_states.dtype == torch.float32:
            compute_type = tl.float32
-        elif hidden_states.dtype == torch.float8_e4m3fn:
+        elif hidden_states.dtype == current_platform.fp8_dtype():
            compute_type = tl.bfloat16
        else:
            raise ValueError(f"Unsupported compute_type: {hidden_states.dtype}")
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Francesco FuscoandGitHub	298e510848	[Hybrid] calling get_mamba_groups() once at MambaCopyBuffers.create() (#37318 ) Signed-off-by: Francesco Fusco <ffu@zurich.ibm.com>	2026-03-21 09:29:43 +00:00
Chaitanya Sri Krishna Lolla GitHub Tej Kiran	3982bc2cd0	[ROCm] Enable DeepEP ROCm as all2allbackend for AMD GPUs. (#34692 ) Signed-off-by: Tej Kiran <vpolamre@amd.com> Co-authored-by: Tej Kiran <vpolamre@amd.com>	2026-03-21 00:32:31 -07:00
Andreas KaratzasandGitHub	02eec7ecbe	[ROCm][CI] Update GSM8K eval config to use fp8-and-mixed models list (MI355) (#37721 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-21 15:27:12 +08:00
Bongwoo Bak GitHub Chauncey	17ee641c45	[Responses API] Add kv_transfer_params for PD disaggregation (#37424 ) Signed-off-by: bongwoobak <bongwoobak@gmail.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com>	2026-03-21 13:48:54 +08:00
Andreas KaratzasandGitHub	0d50fa1db6	[ROCm][CI] Mark gemma3 as large GPU test to avoid OOM on MI250 (#37610 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-21 12:57:25 +08:00
Simon MoandGitHub	1fa1e53a73	Revert "[compile] Initialize passes at VllmBackend init" (#37733 )	2026-03-20 21:35:49 -07:00
Andreas KaratzasandGitHub	3ffa52009f	[ROCm][CI] Guard CudaPlatform/RocmPlatform imports to fix test collection on cross-platform builds (#37617 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-21 11:58:58 +08:00
Yongye Zhu GitHub Claude Opus 4.6	87bd91892f	[MoE Refactor] Mxfp4 oracle rebased (#37128 ) Signed-off-by: Yongye Zhu <zyy1102000@gmail.com> Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-03-21 03:37:04 +00:00
Isotr0pyandGitHub	c7f98b4d0a	[Frontend] Remove librosa from audio dependency (#37058 ) Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>	2026-03-21 11:36:15 +08:00
tmm77andGitHub	1c472f8fe1	Add get_device_uuid for rocm (#37694 ) Signed-off-by: Tiffany Mintz <Tiffany.Mintz@amd.com>	2026-03-21 11:33:16 +08:00
Itay Alroy GitHub Ron Tourgeman	c57d38d603	elastic_ep: Fix issues with repeated scale up/down cycles (#37131 ) Signed-off-by: Itay Alroy <ialroy@nvidia.com> Co-authored-by: Ron Tourgeman <rtourgeman@nvidia.com>	2026-03-20 23:13:02 +00:00
Kaihang JiangandGitHub	e5ed6c6c13	[BugFix] Allow qk_nope_head_dim=192 in FlashInfer MLA backend checks (#37475 ) Signed-off-by: Kaihang Jiang <kaihangj@nvidia.com>	2026-03-20 16:14:55 -06:00
Wentao YeandGitHub	b3d0b37908	[Refactor] Remove unused dead code (#36171 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-03-20 16:12:51 -06:00
Santino RamosandGitHub	85f671b8e1	[Model Runner V2] Support Streaming Inputs (#37028 ) Signed-off-by: Santino Ramos <elsantinoramos@gmail.com>	2026-03-20 20:42:25 +00:00
Andreas KaratzasandGitHub	8bc6b5cdb0	[ROCm][CI] Setting some mi325_4 tests back to optional (in parity with upstream) (#37711 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-20 12:25:08 -07:00
Vadim GimpelsonandGitHub	4f16ebbbd3	[Bugfix] Disable monolithic TRTLLM MoE for Renormalize routing (#37591 ) (#37605 ) Signed-off-by: Vadim Gimpelson <vadim.gimpelson@gmail.com>	2026-03-20 12:19:26 -07:00
Angela YiandGitHub	12fd17eb51	[compile] Initialize passes at VllmBackend init (#35216 ) Signed-off-by: angelayi <yiangela7@gmail.com>	2026-03-20 11:40:33 -07:00
Cyrus LeungandGitHub	37aadf6237	[Model] Update Kimi-K25 and Isaac processors to fit HF-style (#37693 ) Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk>	2026-03-20 18:30:22 +00:00
Le YangandGitHub	d7d2b5e405	[Bugfix] Disable --calculate-kv-scales for hybrid GDN/Mamba+Attention… (#37565 ) Signed-off-by: Young-Leo <562593859@qq.com>	2026-03-20 18:28:34 +00:00
SherryC41 GitHub sherryC41	6ec5e9fd37	refactor: abstract deepgemm support into platform (#37519 ) Co-authored-by: sherryC41 <sherry.c.c41@gmail.com>	2026-03-20 17:54:08 +00:00
Lucas WilkinsonandGitHub	e1d85e5c24	[Attention] Support distinguishing between short extends and decodes (#37303 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>	2026-03-20 10:49:36 -07:00
Peter Pan GitHub Matthew Bonanni	79eb9369c5	fix CUDAGraph memory being counted twice (#37426 ) Signed-off-by: Peter Pan <Peter.Pan@daocloud.io> Signed-off-by: Peter Pan <peter.pan@daocloud.io> Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>	2026-03-20 17:36:32 +00:00
Woosuk KwonandGitHub	e80cfe575d	[MRV2] Avoid recompilation of _gather_block_tables_kernel (#37645 ) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>	2026-03-20 10:31:45 -07:00
Xin YangandGitHub	d0532bf38d	[Perf] Eliminate redundant SparseMatrix creation in gpt_oss_triton_kernels (#37683 ) Signed-off-by: Xin Yang <xyangx@amazon.com>	2026-03-20 11:28:41 -06:00
Andreas KaratzasandGitHub	fb4e8bf442	[ROCm][CI] Fix accuracy for llama-nemotron-vl pooling tests (#37613 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-20 10:16:59 -07:00
Harry MellorandGitHub	6ade4bc5a5	Fix various config related issues for Transformers v5 (#37681 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-03-20 16:30:12 +00:00
Zhengxu ChenandGitHub	2e089b96a8	[compile] Add compiled artifact counter for VLLM_USE_MEGA_AOT_ARTIFACT=1. (#37589 ) Signed-off-by: zhxchen17 <zhxchen17@fb.com>	2026-03-20 16:22:46 +00:00
Martin HickeyandGitHub	880be2b1b8	[Metrics] Some small refactoring for better maintainability (#33898 ) Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com>	2026-03-20 16:11:34 +00:00
Zhengxu ChenandGitHub	c0f5fae601	[compile] Fix aot test failures with torch 2.12. (#37604 ) Signed-off-by: zhxchen17 <zhxchen17@fb.com>	2026-03-20 16:06:29 +00:00
Rémi DelacourtandGitHub	aa84e43ccb	[Pixtral] Enable Pixtral language model support Eagle3 (#37182 ) Signed-off-by: remi <remi@mistral.ai>	2026-03-20 15:50:15 +00:00
Matthias GehreandGitHub	5e806bcf54	[Bugfix] Fix ConchLinearKernel channelwise quantization (group_size=-1) (#37329 ) Signed-off-by: Matthias Gehre <matthias.gehre@amd.com>	2026-03-20 10:32:21 -05:00
Matthias GehreandGitHub	56a62c310c	[Bugfix] Reject channelwise quantization (group_size <= 0) in ExllamaLinearKernel (#37331 ) Signed-off-by: Matthias Gehre <matthias.gehre@amd.com>	2026-03-20 10:31:57 -05:00
L.B.R.GitHub L.B.R.	1779c09898	[ROCm] Enable wvSplitK skinny GEMM kernel for RDNA4/gfx1x decode (#34709 ) Signed-off-by: L.B.R. <lbr@mmonad.com> Co-authored-by: L.B.R. <lbr@mmonad.com>	2026-03-20 10:11:23 -05:00
xuebwang-amdandGitHub	44eea10f68	[ROCm][Quantization] make quark ocp mx dtype parser robust for weight-only quantization (#36232 ) Signed-off-by: xuebwang-amd <xuebwang@amd.com>	2026-03-20 10:10:03 -05:00
Ilya BoytsovandGitHub	8b6c6b9505	[Model] Add LFM2-ColBERT-350M support (#37528 ) Signed-off-by: Ilya Boytsov <ilyaboytsov1805@gmail.com>	2026-03-20 14:57:57 +00:00
Harry MellorandGitHub	9f6d9dd371	Fix attribute error in `isaac_patch_hf_runner` (#37685 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-03-20 14:49:40 +00:00
Jee Jee LiandGitHub	dd20ee4e3e	[UX] Enable torch_profiler_with_stack (#37571 ) Signed-off-by: Jee Jee Li <pandaleefree@gmail.com>	2026-03-20 11:17:26 +00:00
ChaunceyandGitHub	0523449c9c	[Misc] Use logger.info_once for auto tool choice log message (#37661 ) Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>	2026-03-20 10:40:36 +00:00
Flora FengandGitHub	b4c1aef21c	[Refactor] Relocate tests from tests/v1/entrypoints/ to tests/entrypoints/ (#37500 ) Signed-off-by: sfeng33 <4florafeng@gmail.com>	2026-03-20 02:50:34 -07:00
Flora FengandGitHub	6050b93bed	[Refactor] Move serve entrypoint tests under tests/entrypoints/serve/ (#37595 ) Signed-off-by: sfeng33 <4florafeng@gmail.com>	2026-03-20 02:10:47 -07:00
Andreas KaratzasandGitHub	5a4a179591	[ROCm][CI] Fix granite_speech test for gfx90a by selecting compatible attention backend (#37611 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-20 17:07:26 +08:00
Andreas KaratzasandGitHub	37cd9fc107	[ROCm][CI] Remove deepep DBO tests on gfx90a (#37614 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-20 17:07:07 +08:00
Andreas KaratzasandGitHub	9cfd4ebb5e	[ROCm][CI] Update GSM8K eval config to use fp8-and-mixed models list (#37619 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-03-20 17:06:53 +08:00
wang.yuqiandGitHub	ed359c497a	[Model] Deprecate the score task (this will not affect users). (#37537 ) Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>	2026-03-20 08:07:56 +00:00