[ROCm][CI] Stage B gating (#42025 )

Signed-off-by: Andreas Karatzas <akaratza@amd.com>
DeepSeekV4-Pro enable cuda graph full and piecewise mode (#42604 )
2026-05-15 01:49:27 -07:00 · 2026-05-15 01:45:30 -07:00 · 2026-05-15 08:05:57 +00:00 · 2026-05-15 15:59:18 +08:00 · 2026-05-15 15:20:42 +08:00 · 2026-05-15 00:13:32 -07:00
289 changed files with 13987 additions and 5234 deletions
@@ -61,7 +61,7 @@ steps:
  source_file_dependencies:
  - csrc/cpu/
  - vllm/model_executor/layers/quantization/cpu_wna16.py
-  - vllm/model_executor/layers/quantization/gptq_marlin.py
+  - vllm/model_executor/layers/quantization/auto_gptq.py
  - vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w8a8_int8.py
  - vllm/model_executor/layers/quantization/kernels/scaled_mm/cpu.py
  - vllm/model_executor/layers/quantization/kernels/mixed_precision/cpu.py
@@ -36,6 +36,7 @@ steps:
        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --attention-backend=TRITON_ATTN &&
        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --quantization fp8 &&
        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8 &&
+        python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096 &&
        python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 &&
        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel'
@@ -62,4 +63,4 @@ steps:
        pytest -v -s v1/structured_output &&
        pytest -v -s v1/test_serial_utils.py &&
        pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py &&
-        pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py'
+        pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py'
@@ -2,7 +2,7 @@
 # We can use this script to compute baseline accuracy on chartqa for vllm.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.11"
+#   pip install "lm-eval[api]>=0.4.12"

 usage() {
    echo``
@@ -2,7 +2,7 @@
 # We can use this script to compute baseline accuracy on GSM for transformers.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.11"
+#   pip install "lm-eval[api]>=0.4.12"

 usage() {
    echo``
@@ -3,7 +3,7 @@
 # We use this for fp8, which HF does not support.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.11"
+#   pip install "lm-eval[api]>=0.4.12"

 usage() {
    echo``
@@ -3,7 +3,7 @@
 # We use this for fp8, which HF does not support.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.11"
+#   pip install "lm-eval[api]>=0.4.12"

 usage() {
    echo``
@@ -31,30 +31,9 @@
    }
  },
  "tests": [
-    {
-      "test_name": "serving_llama8B_tp1_sharegpt",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "sharegpt",
-        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp2_sharegpt",
-      "server_parameters": {
-        "tensor_parallel_size": 2
-      },
-      "client_parameters": {
-        "dataset_name": "sharegpt",
-        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
-      }
-    },
    {
      "test_name": "serving_llama8B_tp1_random_128_128",
      "server_parameters": {
-        "tensor_parallel_size": 1
      },
      "client_parameters": {
        "dataset_name": "random",
@@ -62,158 +41,10 @@
        "random-output-len": 128
      }
    },
-    {
-      "test_name": "serving_llama8B_tp2_random_128_128",
-      "server_parameters": {
-        "tensor_parallel_size": 2
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp4_random_128_128",
-      "server_parameters": {
-        "tensor_parallel_size": 4
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp1_random_128_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 2048
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp2_random_128_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 2
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 2048
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp4_random_128_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 4
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 2048
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp1_random_2048_128",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp2_random_2048_128",
-      "server_parameters": {
-        "tensor_parallel_size": 2
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp4_random_2048_128",
-      "server_parameters": {
-        "tensor_parallel_size": 4
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp1_random_2048_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 2048
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp2_random_2048_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 2
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 2048
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp4_random_2048_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 4
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 2048
-      }
-    },
    {
      "test_name": "serving_llama8B_int4_tp1_random_128_128",
      "server_parameters": {
-        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_int4_tp2_random_128_128",
-      "server_parameters": {
-        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
-        "tensor_parallel_size": 2
-      },
-      "client_parameters": {
-        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_int4_tp4_random_128_128",
-      "server_parameters": {
-        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
-        "tensor_parallel_size": 4
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4"
      },
      "client_parameters": {
        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
@@ -225,8 +56,7 @@
    {
      "test_name": "serving_llama8B_int8_tp1_random_128_128",
      "server_parameters": {
-        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
-        "tensor_parallel_size": 1
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8"
      },
      "client_parameters": {
        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
@@ -236,26 +66,12 @@
      }
    },
    {
-      "test_name": "serving_llama8B_int8_tp2_random_128_128",
+      "test_name": "serving_llama1B_tp1_random_128_128",
      "server_parameters": {
-        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
-        "tensor_parallel_size": 2
+        "model": "meta-llama/Llama-3.2-1B"
      },
      "client_parameters": {
-        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_int8_tp4_random_128_128",
-      "server_parameters": {
-        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
-        "tensor_parallel_size": 4
-      },
-      "client_parameters": {
-        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "model": "meta-llama/Llama-3.2-1B",
        "dataset_name": "random",
        "random-input-len": 128,
        "random-output-len": 128
@@ -264,8 +80,7 @@
    {
      "test_name": "serving_llama3B_tp1_random_128_128",
      "server_parameters": {
-        "model": "meta-llama/Llama-3.2-3B-Instruct",
-        "tensor_parallel_size": 1
+        "model": "meta-llama/Llama-3.2-3B-Instruct"
      },
      "client_parameters": {
        "model": "meta-llama/Llama-3.2-3B-Instruct",
@@ -274,11 +89,22 @@
        "random-output-len": 128
      }
    },
+    {
+      "test_name": "serving_llama70B_tp1_random_128_128",
+      "server_parameters": {
+        "model": "meta-llama/Llama-3.3-70B-Instruct"
+      },
+      "client_parameters": {
+        "model": "meta-llama/Llama-3.3-70B-Instruct",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
    {
      "test_name": "serving_granite2B_tp1_random_128_128",
      "server_parameters": {
-        "model": "ibm-granite/granite-3.2-2b-instruct",
-        "tensor_parallel_size": 1
+        "model": "ibm-granite/granite-3.2-2b-instruct"
      },
      "client_parameters": {
        "model": "ibm-granite/granite-3.2-2b-instruct",
@@ -290,8 +116,7 @@
    {
      "test_name": "serving_qwen1.7B_tp1_random_128_128",
      "server_parameters": {
-        "model": "Qwen/Qwen3-1.7B",
-        "tensor_parallel_size": 1
+        "model": "Qwen/Qwen3-1.7B"
      },
      "client_parameters": {
        "model": "Qwen/Qwen3-1.7B",
@@ -303,8 +128,7 @@
    {
      "test_name": "serving_qwen4B_tp1_random_128_128",
      "server_parameters": {
-        "model": "Qwen/Qwen3-4B",
-        "tensor_parallel_size": 1
+        "model": "Qwen/Qwen3-4B"
      },
      "client_parameters": {
        "model": "Qwen/Qwen3-4B",
@@ -316,8 +140,7 @@
    {
      "test_name": "serving_qwen8B_tp1_random_128_128",
      "server_parameters": {
-        "model": "Qwen/Qwen3-8B",
-        "tensor_parallel_size": 1
+        "model": "Qwen/Qwen3-8B"
      },
      "client_parameters": {
        "model": "Qwen/Qwen3-8B",
@@ -326,11 +149,34 @@
        "random-output-len": 128
      }
    },
+    {
+      "test_name": "serving_qwen14B_tp1_random_128_128",
+      "server_parameters": {
+        "model": "Qwen/Qwen3-14B"
+      },
+      "client_parameters": {
+        "model": "Qwen/Qwen3-14B",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_qwen30B_tp1_random_128_128",
+      "server_parameters": {
+        "model": "Qwen/Qwen3-30B-A3B"
+      },
+      "client_parameters": {
+        "model": "Qwen/Qwen3-30B-A3B",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
    {
      "test_name": "serving_glm9B_tp1_random_128_128",
      "server_parameters": {
-        "model": "zai-org/glm-4-9b-hf",
-        "tensor_parallel_size": 1
+        "model": "zai-org/glm-4-9b-hf"
      },
      "client_parameters": {
        "model": "zai-org/glm-4-9b-hf",
@@ -342,8 +188,7 @@
    {
      "test_name": "serving_gemma7B_tp1_random_128_128",
      "server_parameters": {
-        "model": "google/gemma-7b",
-        "tensor_parallel_size": 1
+        "model": "google/gemma-7b"
      },
      "client_parameters": {
        "model": "google/gemma-7b",
@@ -351,6 +196,94 @@
        "random-input-len": 128,
        "random-output-len": 128
      }
+    },
+    {
+      "test_name": "serving_gemma3-4b_tp1_random_128_128",
+      "server_environment_variables": {
+	"VLLM_CPU_SGL_KERNEL": 0
+      },
+      "server_parameters": {
+        "model": "google/gemma-3-4b-it"
+      },
+      "client_parameters": {
+        "model": "google/gemma-3-4b-it",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_gemma3-12b_tp1_random_128_128",
+      "server_environment_variables": {
+	"VLLM_CPU_SGL_KERNEL": 0
+      },
+      "server_parameters": {
+        "model": "google/gemma-3-12b-it"
+      },
+      "client_parameters": {
+        "model": "google/gemma-3-12b-it",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_gemma4-4b_tp1_random_128_128",
+      "server_environment_variables": {
+	"VLLM_CPU_SGL_KERNEL": 0
+      },
+      "server_parameters": {
+        "model": "google/gemma-4-E4B-it"
+      },
+      "client_parameters": {
+        "model": "google/gemma-4-E4B-it",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_gemma4-2b_tp1_random_128_128",
+      "server_environment_variables": {
+	"VLLM_CPU_SGL_KERNEL": 0
+      },
+      "server_parameters": {
+        "model": "google/gemma-4-E2B-it"
+      },
+      "client_parameters": {
+        "model": "google/gemma-4-E2B-it",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_gemma4-26b_tp1_random_128_128",
+      "server_environment_variables": {
+	"VLLM_CPU_SGL_KERNEL": 0,
+        "VLLM_CPU_ATTN_SPLIT_KV": 0
+      },
+      "server_parameters": {
+        "model": "google/gemma-4-26B-A4B-it"
+      },
+      "client_parameters": {
+        "model": "google/gemma-4-26B-A4B-it",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_phi4_tp1_random_128_128",
+      "server_parameters": {
+        "model": "microsoft/Phi-4-reasoning"
+      },
+      "client_parameters": {
+        "model": "microsoft/Phi-4-reasoning",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
    }
  ]
 }
@@ -61,7 +61,7 @@ echo "Results will be stored in: $RESULTS_DIR"
 echo "--- Installing Python dependencies ---"
 python3 -m pip install --progress-bar off git+https://github.com/thuml/depyf.git \
    && python3 -m pip install --progress-bar off pytest pytest-asyncio tpu-info \
-    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.11" \
+    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.12" \
    && python3 -m pip install --progress-bar off hf-transfer tblib==3.1.0
 echo "--- Python dependencies installed ---"

@@ -61,7 +61,7 @@ echo "Results will be stored in: $RESULTS_DIR"
 echo "--- Installing Python dependencies ---"
 python3 -m pip install --progress-bar off git+https://github.com/thuml/depyf.git \
    && python3 -m pip install --progress-bar off pytest pytest-asyncio tpu-info \
-    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.11" \
+    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.12" \
    && python3 -m pip install --progress-bar off hf-transfer tblib==3.1.0
 echo "--- Python dependencies installed ---"

@@ -798,32 +798,6 @@ steps:
  commands:
  - pytest -v -s v1/attention

- label: V1 Core + KV + Metrics # TBD
-  timeout_in_minutes: 60
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
-  agent_pool: mi250_1
-  optional: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/v1/core
-  - tests/v1/executor
-  - tests/v1/kv_offload
-  - tests/v1/worker
-  - tests/v1/kv_connector/unit
-  - tests/v1/metrics
-  - tests/entrypoints/openai/correctness/test_lmeval.py
-  commands:
-  - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - pytest -v -s -m 'not cpu_test' v1/core
-  - pytest -v -s v1/executor
-  - pytest -v -s v1/kv_offload
-  - pytest -v -s v1/worker
-  - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
-  - pytest -v -s -m 'not cpu_test' v1/metrics
-  - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
-  - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
-
 - label: V1 Sample + Logits # TBD
  timeout_in_minutes: 60
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -1279,6 +1253,7 @@ steps:
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
+  optional: true
  fast_check: true
  torch_nightly: true
  working_dir: "/vllm-workspace/tests"
@@ -1299,6 +1274,7 @@ steps:
  agent_pool: mi300_1
  fast_check: true
  torch_nightly: true
+  optional: true
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -1312,6 +1288,7 @@ steps:
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
+  optional: true
  fast_check: true
  torch_nightly: true
  working_dir: "/vllm-workspace/tests"
@@ -1328,6 +1305,7 @@ steps:
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
+  optional: true
  fast_check: true
  torch_nightly: true
  working_dir: "/vllm-workspace/tests"
@@ -1728,6 +1706,7 @@ steps:
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
+  optional: true
  parallelism: 2
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -1786,6 +1765,7 @@ steps:
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_1
+  optional: true
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -51,7 +51,7 @@ steps:

 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
  key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -36,6 +36,12 @@ steps:
  commands:
    - pytest -v -s v1/engine/test_preprocess_error_handling.py
    - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
+  mirror:
+    amd:
+      device: mi300_1
+      timeout_in_minutes: 40
+      depends_on:
+      - image-build-amd

 - label: e2e Scheduling (1 GPU)
  key: e2e-scheduling-1-gpu
@@ -46,6 +52,12 @@ steps:
    - tests/v1/e2e/general/
  commands:
    - pytest -v -s v1/e2e/general/test_async_scheduling.py
+  mirror:
+    amd:
+      device: mi250_1
+      timeout_in_minutes: 60
+      depends_on:
+      - image-build-amd

 - label: e2e Core (1 GPU)
  key: e2e-core-1-gpu
@@ -43,6 +43,12 @@ steps:
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
+  mirror:
+    amd:
+      device: mi300_1
+      timeout_in_minutes: 80
+      depends_on:
+      - image-build-amd

 - label: Entrypoints Integration (API Server openai - Part 2)
  key: entrypoints-integration-api-server-openai-part-2
@@ -55,6 +61,12 @@ steps:
  commands:
  - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
  - pytest -v -s entrypoints/test_chat_utils.py
+  mirror:
+    amd:
+      device: mi300_1
+      timeout_in_minutes: 60
+      depends_on:
+      - image-build-amd

 - label: Entrypoints Integration (API Server openai - Part 3)
  key: entrypoints-integration-api-server-openai-part-3
@@ -67,7 +79,13 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion  --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/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
+  mirror:
+    amd:
+      device: mi300_1
+      timeout_in_minutes: 60
+      depends_on:
+      - image-build-amd

 - label: Entrypoints Integration (API Server 2)
  key: entrypoints-integration-api-server-2
@@ -83,6 +101,11 @@ steps:
  - pytest -v -s entrypoints/serve/instrumentator
  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
  - pytest -v -s tool_use
+  mirror:
+    amd:
+      device: mi300_1
+      depends_on:
+      - image-build-amd

 - label: Entrypoints Integration (Speech to Text)
  key: entrypoints-integration-speech_to_text
@@ -62,6 +62,19 @@ steps:
  commands:
    - pytest -v -s kernels/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
  parallelism: 2
+  mirror:
+    amd:
+      device: mi300_1
+      source_file_dependencies:
+      - csrc/quantization/
+      - vllm/model_executor/layers/quantization
+      - tests/kernels/quantization
+      - tests/kernels/quantization/test_rocm_skinny_gemms.py
+      - vllm/_aiter_ops.py
+      - vllm/platforms/rocm.py
+      - vllm/model_executor/kernels/
+      depends_on:
+      - image-build-amd

 - label: Kernels MoE Test %N
  key: kernels-moe-test
@@ -129,6 +129,19 @@ steps:
    - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
    - pytest -v -s -m 'cpu_test' v1/metrics

+- label: Extract Hidden States Integration
+  key: extract-hidden-states-integration
+  timeout_in_minutes: 20
+  device: h200_18gb
+  source_file_dependencies:
+    - vllm/v1/spec_decode/extract_hidden_states.py
+    - vllm/model_executor/models/extract_hidden_states.py
+    - vllm/transformers_utils/configs/extract_hidden_states.py
+    - tests/v1/kv_connector/extract_hidden_states_integration
+  commands:
+    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+    - pytest -v -s v1/kv_connector/extract_hidden_states_integration
+
 - label: Regression
  key: regression
  timeout_in_minutes: 20
@@ -305,6 +318,20 @@ steps:
  - pytest -v -s transformers_utils
  - pytest -v -s config

+- label: Batch Invariance (A100)
+  key: batch-invariance-a100
+  timeout_in_minutes: 30
+  device: a100
+  source_file_dependencies:
+    - vllm/v1/attention
+    - vllm/model_executor/layers
+    - tests/v1/determinism/
+  commands:
+    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+    - pip install pytest-timeout pytest-forked
+    - pytest -v -s v1/determinism/test_batch_invariance.py
+    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
+
 - label: Batch Invariance (H100)
  key: batch-invariance-h100
  timeout_in_minutes: 30
@@ -92,6 +92,12 @@ steps:
  - tests/models/language/pooling
  commands:
    - pytest -v -s models/language/pooling -m 'not core_model'
+  mirror:
+    amd:
+      device: mi300_1
+      timeout_in_minutes: 100
+      depends_on:
+      - image-build-amd

 - label: Language Models Test (MTEB)
  key: language-models-test-mteb
@@ -1222,6 +1222,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
      CUDA_ARCHS "${DSV3_ROUTER_GEMM_ARCHS}")
    list(APPEND VLLM_MOE_EXT_SRC "${DSV3_ROUTER_GEMM_SRC}")
    message(STATUS "Building DSV3 router GEMM kernel for archs: ${DSV3_ROUTER_GEMM_ARCHS}")
+
+    # DeepSeek V4 fused RMSNorm + router GEMV - same arch gating as DSV3.
+    set(DSV4_NORM_ROUTER_GEMM_SRC
+      "csrc/moe/dsv4_norm_router_gemm_entry.cu"
+      "csrc/moe/dsv4_norm_router_gemm_kernel.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${DSV4_NORM_ROUTER_GEMM_SRC}"
+      CUDA_ARCHS "${DSV3_ROUTER_GEMM_ARCHS}")
+    list(APPEND VLLM_MOE_EXT_SRC "${DSV4_NORM_ROUTER_GEMM_SRC}")
+    message(STATUS "Building DSV4 norm+router GEMV kernel for archs: ${DSV3_ROUTER_GEMM_ARCHS}")
  else()
    message(STATUS "Not building DSV3 router GEMM kernel as no compatible archs found"
                   " (requires SM90+ and CUDA >= 12.0)")
@@ -53,6 +53,7 @@ backends:
  - FLASHINFER_MLA
  - FLASH_ATTN_MLA  # Hopper only
  - FLASHMLA        # Hopper only
+  - TOKENSPEED_MLA  # Blackwell + R1 dims + FP8 KV (use --kv-cache-dtype fp8)

 device: "cuda:0"
 repeats: 100
@@ -3,6 +3,7 @@
 # Compares all available MLA prefill backends:
 #   FA backends:  fa2, fa3, fa4 (FlashAttention versions)
 #   Non-FA:       flashinfer, cudnn, trtllm (Blackwell-only, require flashinfer)
+#   CuTe DSL:     tokenspeed (Blackwell + R1 dims, requires tokenspeed_mla)
 #
 # Uses cutlass_mla as the decode backend for impl construction
 # (only the prefill path is exercised).
@@ -120,6 +121,7 @@ prefill_backends:
  - flashinfer
  - cudnn
  - trtllm
+  - tokenspeed

 device: "cuda:0"
 repeats: 20
@@ -29,6 +29,7 @@ from vllm.config import (
    VllmConfig,
    set_current_vllm_config,
 )
+from vllm.v1.attention.backends.mla.prefill.registry import MLAPrefillBackendEnum

 # ============================================================================
 # VllmConfig Creation
@@ -79,8 +80,8 @@ def create_minimal_vllm_config(
        index_topk: Optional topk value for sparse MLA backends. If provided,
                    the config will include index_topk for sparse attention.
        prefill_backend: Prefill backend name (e.g., "fa3", "fa4", "flashinfer",
-                        "cudnn", "trtllm"). Configures the attention config to
-                        force the specified prefill backend.
+                        "trtllm"). Configures the attention config to force
+                        the specified prefill backend.

    Returns:
        VllmConfig for benchmarking
@@ -179,19 +180,13 @@ def create_minimal_vllm_config(

    if prefill_backend is not None:
        prefill_cfg = get_prefill_backend_config(prefill_backend)
+        vllm_config.attention_config.mla_prefill_backend = prefill_cfg[
+            "mla_prefill_backend"
+        ]
        if prefill_cfg["flash_attn_version"] is not None:
            vllm_config.attention_config.flash_attn_version = prefill_cfg[
                "flash_attn_version"
            ]
-        vllm_config.attention_config.disable_flashinfer_prefill = prefill_cfg[
-            "disable_flashinfer_prefill"
-        ]
-        vllm_config.attention_config.use_cudnn_prefill = prefill_cfg[
-            "use_cudnn_prefill"
-        ]
-        vllm_config.attention_config.use_trtllm_ragged_deepseek_prefill = prefill_cfg[
-            "use_trtllm_ragged_deepseek_prefill"
-        ]

    return vllm_config

@@ -206,39 +201,27 @@ def create_minimal_vllm_config(
 _PREFILL_BACKEND_CONFIG: dict[str, dict] = {
    "fa2": {
        "flash_attn_version": 2,
-        "disable_flashinfer_prefill": True,
-        "use_cudnn_prefill": False,
-        "use_trtllm_ragged_deepseek_prefill": False,
+        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
    },
    "fa3": {
        "flash_attn_version": 3,
-        "disable_flashinfer_prefill": True,
-        "use_cudnn_prefill": False,
-        "use_trtllm_ragged_deepseek_prefill": False,
+        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
    },
    "fa4": {
        "flash_attn_version": 4,
-        "disable_flashinfer_prefill": True,
-        "use_cudnn_prefill": False,
-        "use_trtllm_ragged_deepseek_prefill": False,
+        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
    },
    "flashinfer": {
        "flash_attn_version": None,
-        "disable_flashinfer_prefill": False,
-        "use_cudnn_prefill": False,
-        "use_trtllm_ragged_deepseek_prefill": False,
-    },
-    "cudnn": {
-        "flash_attn_version": None,
-        "disable_flashinfer_prefill": True,
-        "use_cudnn_prefill": True,
-        "use_trtllm_ragged_deepseek_prefill": False,
+        "mla_prefill_backend": MLAPrefillBackendEnum.FLASHINFER,
    },
    "trtllm": {
        "flash_attn_version": None,
-        "disable_flashinfer_prefill": True,
-        "use_cudnn_prefill": False,
-        "use_trtllm_ragged_deepseek_prefill": True,
+        "mla_prefill_backend": MLAPrefillBackendEnum.TRTLLM_RAGGED,
+    },
+    "tokenspeed": {
+        "flash_attn_version": None,
+        "mla_prefill_backend": MLAPrefillBackendEnum.TOKENSPEED_MLA,
    },
 }

@@ -625,6 +608,21 @@ def _create_backend_impl(
    # Create mock layer
    layer = MockLayer(device, impl=impl, kv_cache_spec=kv_cache_spec)

+    # Attach a prefill backend (MLAAttention does this in __init__; the metadata
+    # builder reads layer.prefill_backend from static_forward_context).
+    from vllm.v1.attention.backends.mla.prefill import get_mla_prefill_backend
+
+    prefill_backend_cls = get_mla_prefill_backend(vllm_config)
+    layer.prefill_backend = prefill_backend_cls(
+        num_heads=mla_dims["num_q_heads"],
+        scale=(mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"]) ** -0.5,
+        kv_lora_rank=mla_dims["kv_lora_rank"],
+        qk_nope_head_dim=mla_dims["qk_nope_head_dim"],
+        qk_rope_head_dim=mla_dims["qk_rope_head_dim"],
+        v_head_dim=mla_dims["v_head_dim"],
+        vllm_config=vllm_config,
+    )
+
    # Create builder instance if needed
    builder_instance = None
    if builder_class:
@@ -961,19 +959,6 @@ def _run_mla_benchmark_batched(
    results = []

    with set_current_vllm_config(vllm_config):
-        # Clear cached prefill backend detection functions so they re-evaluate
-        # with the current VllmConfig. These are @functools.cache decorated and
-        # would otherwise return stale results from a previous backend's config.
-        from vllm.model_executor.layers.attention.mla_attention import (
-            use_cudnn_prefill,
-            use_flashinfer_prefill,
-            use_trtllm_ragged_deepseek_prefill,
-        )
-
-        use_flashinfer_prefill.cache_clear()
-        use_cudnn_prefill.cache_clear()
-        use_trtllm_ragged_deepseek_prefill.cache_clear()
-
        # Create backend impl, layer, builder, and indexer (reused across benchmarks)
        impl, layer, builder_instance, indexer = _create_backend_impl(
            backend_cfg,
@@ -985,38 +970,36 @@ def _run_mla_benchmark_batched(
            kv_cache_dtype=kv_cache_dtype,
        )

-        # Verify the actual prefill backend matches what was requested
+        # Verify the actual prefill backend matches what was requested. The
+        # selector + impl construction already raise on misuse; here we just
+        # check the resolved class against the requested name as a sanity guard.
        if prefill_backend is not None:
-            prefill_cfg = get_prefill_backend_config(prefill_backend)
-            fa_version = prefill_cfg["flash_attn_version"]
-
-            if fa_version is not None:
-                # FA backend: verify the impl's FA version
-                actual_fa_version = getattr(impl, "vllm_flash_attn_version", None)
+            expected_class = {
+                "fa2": "FlashAttnPrefillBackend",
+                "fa3": "FlashAttnPrefillBackend",
+                "fa4": "FlashAttnPrefillBackend",
+                "flashinfer": "FlashInferPrefillBackend",
+                "trtllm": "TrtllmRaggedPrefillBackend",
+                "tokenspeed": "TokenspeedMLAPrefillBackend",
+            }.get(prefill_backend)
+            actual_class = type(getattr(layer, "prefill_backend", None)).__name__
+            if expected_class and actual_class != expected_class:
+                raise RuntimeError(
+                    f"Prefill backend '{prefill_backend}' requested "
+                    f"{expected_class}, got {actual_class}. Check "
+                    f"attention_config plumbing or installed deps."
+                )
+            if prefill_backend in {"fa2", "fa3", "fa4"}:
+                fa_version = int(prefill_backend[2:])
+                actual_fa_version = getattr(
+                    layer.prefill_backend, "vllm_flash_attn_version", None
+                )
                if actual_fa_version != fa_version:
                    raise RuntimeError(
                        f"Prefill backend '{prefill_backend}' requested FA "
-                        f"version {fa_version}, but the impl is using FA "
-                        f"version {actual_fa_version}. Check "
-                        f"vllm/v1/attention/backends/fa_utils.py."
+                        f"version {fa_version}, got "
+                        f"{actual_fa_version} on {actual_class}."
                    )
-            else:
-                # Non-FA backend: verify the builder picked the right path
-                expected_flags = {
-                    "flashinfer": "_use_fi_prefill",
-                    "cudnn": "_use_cudnn_prefill",
-                    "trtllm": "_use_trtllm_ragged_prefill",
-                }
-                flag_name = expected_flags.get(prefill_backend)
-                if flag_name and not getattr(builder_instance, flag_name, False):
-                    raise RuntimeError(
-                        f"Prefill backend '{prefill_backend}' was requested "
-                        f"but the metadata builder did not enable it. This "
-                        f"usually means a dependency is missing (e.g., "
-                        f"flashinfer not installed) or the platform doesn't "
-                        f"support it."
-                    )
-
        # Run each benchmark with the shared impl
        for config, threshold, num_splits in configs_with_params:
            # Set threshold for this benchmark (FlashAttn/FlashMLA only)
@@ -0,0 +1,183 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Benchmark and correctness check for ``ops.dsv4_norm_router_gemm``.
+
+Two implementations are compared:
+
+  1. ``unfused``   — ``vllm_ops.rms_norm`` then ``ops.dsv3_router_gemm``,
+                     i.e. the current vLLM hot path (two kernel launches).
+  2. ``fused``     — ``ops.dsv4_norm_router_gemm``, the new single-kernel
+                     fused path.
+
+Both produce ``(normed_x: bf16, router_logits: fp32)``.  The correctness
+check verifies that ``fused`` and ``unfused`` agree to within ~1 bf16
+ULP — that is the precision floor for this op.
+"""
+
+import argparse
+
+import torch
+
+from vllm import _custom_ops as vllm_ops
+from vllm.triton_utils import triton
+
+# The fused dsv4_norm_router_gemm kernel is templated only for DSV4-Pro
+# (hidden_size=7168, num_experts=384).  Other shapes fall back to the
+# unfused path on the Python side (NormGatedLinear), so benchmark only
+# the configuration that the fused kernel actually targets.
+HIDDEN_SIZE = 7168
+NUM_EXPERTS_CHOICES = (384,)
+RMS_EPS = 1e-6
+
+
+def unfused_norm_router_gemm(
+    x: torch.Tensor,
+    norm_weight: torch.Tensor,
+    gate_weight: torch.Tensor,
+    eps: float,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    # Call ``_C::rms_norm`` directly (mirroring ``_dsv4_pro_norm_gate``'s
+    # fallback path) so the benchmarked baseline doesn't inherit any
+    # Python wrapper overhead or risk falling through to the native
+    # eager-primitive ``RMSNorm.forward_native`` path.
+    normed = torch.empty_like(x)
+    torch.ops._C.rms_norm(normed, x, norm_weight, eps)
+    logits = vllm_ops.dsv3_router_gemm(normed, gate_weight, torch.float32)
+    return normed, logits
+
+
+def fused_norm_router_gemm(
+    x: torch.Tensor,
+    norm_weight: torch.Tensor,
+    gate_weight: torch.Tensor,
+    eps: float,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    return vllm_ops.dsv4_norm_router_gemm(x, norm_weight, gate_weight, eps)
+
+
+def _make_inputs(num_tokens: int, num_experts: int, hidden_size: int, seed: int = 0):
+    torch.manual_seed(seed)
+    device = "cuda"
+    x = torch.randn(num_tokens, hidden_size, dtype=torch.bfloat16, device=device)
+    norm_w = torch.randn(hidden_size, dtype=torch.bfloat16, device=device)
+    gate_w = torch.randn(num_experts, hidden_size, dtype=torch.bfloat16, device=device)
+    # Down-scale gate_w so the GEMV output stays in a representable range.
+    gate_w = gate_w / float(hidden_size) ** 0.5
+    norm_w = (norm_w * 0.1) + 1.0
+    return x, norm_w, gate_w
+
+
+def calculate_diff(
+    num_tokens: int,
+    num_experts: int,
+    hidden_size: int = HIDDEN_SIZE,
+    normed_atol: float = 2e-3,
+    logits_atol: float = 1e-2,
+    rtol: float = 1e-2,
+) -> None:
+    x, norm_w, gate_w = _make_inputs(num_tokens, num_experts, hidden_size)
+
+    normed_unfused, logits_unfused = unfused_norm_router_gemm(
+        x.clone(), norm_w, gate_w, RMS_EPS
+    )
+    normed_fused, logits_fused = fused_norm_router_gemm(
+        x.clone(), norm_w, gate_w, RMS_EPS
+    )
+
+    def _max_abs(a, b):
+        return (a.float() - b.float()).abs().max().item()
+
+    print(f"\n=== M={num_tokens} E={num_experts} H={hidden_size} ===")
+    print(f"normed_x  |fused - unfused| = {_max_abs(normed_fused, normed_unfused):.3e}")
+    print(f"logits    |fused - unfused| = {_max_abs(logits_fused, logits_unfused):.3e}")
+
+    ok_normed = torch.allclose(
+        normed_fused.float(),
+        normed_unfused.float(),
+        atol=normed_atol,
+        rtol=rtol,
+    )
+    ok_logits = torch.allclose(
+        logits_fused.float(),
+        logits_unfused.float(),
+        atol=logits_atol,
+        rtol=rtol,
+    )
+    if ok_normed and ok_logits:
+        print(
+            f"OK   fused vs unfused within "
+            f"normed_atol={normed_atol:.0e} logits_atol={logits_atol:.0e} "
+            f"rtol={rtol:.0e}"
+        )
+    else:
+        print(
+            f"FAIL normed_ok={ok_normed} logits_ok={ok_logits}; "
+            f"see max-abs values above"
+        )
+
+
+def get_benchmark():
+    # Only num_tokens varies (DSV4-Pro hard-codes E=384); single-axis
+    # sweep yields a clean line plot with M on the x-axis.
+    num_experts = NUM_EXPERTS_CHOICES[0]
+
+    @triton.testing.perf_report(
+        triton.testing.Benchmark(
+            x_names=["num_tokens"],
+            x_vals=list(range(1, 17)),
+            line_arg="provider",
+            line_vals=["unfused", "fused"],
+            line_names=["unfused (rms+dsv3)", "fused (dsv4)"],
+            styles=[("green", "-"), ("red", "-")],
+            ylabel="us",
+            plot_name=f"norm-router-gemm-E{num_experts}-H{HIDDEN_SIZE}",
+            args={},
+        )
+    )
+    def benchmark(num_tokens, provider):
+        x, norm_w, gate_w = _make_inputs(num_tokens, num_experts, HIDDEN_SIZE)
+
+        quantiles = [0.5, 0.2, 0.8]
+        if provider == "unfused":
+            fn = lambda: unfused_norm_router_gemm(  # noqa: E731
+                x, norm_w, gate_w, RMS_EPS
+            )
+        else:
+            fn = lambda: fused_norm_router_gemm(  # noqa: E731
+                x, norm_w, gate_w, RMS_EPS
+            )
+
+        ms, min_ms, max_ms = triton.testing.do_bench(fn, quantiles=quantiles)
+        return 1000 * ms, 1000 * max_ms, 1000 * min_ms
+
+    return benchmark
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument(
+        "--save-path",
+        type=str,
+        default="./configs/norm_router_gemm/",
+    )
+    parser.add_argument(
+        "--skip-bench",
+        action="store_true",
+        help="Run only the correctness check, not the perf sweep.",
+    )
+    args = parser.parse_args()
+
+    # Correctness sweep over the full fast-path range M=1..16.
+    for m in range(1, 17):
+        for e in NUM_EXPERTS_CHOICES:
+            calculate_diff(num_tokens=m, num_experts=e, hidden_size=HIDDEN_SIZE)
+
+    if args.skip_bench:
+        return
+
+    benchmark = get_benchmark()
+    benchmark.run(print_data=True, save_path=args.save_path)
+
+
+if __name__ == "__main__":
+    main()
@@ -12,7 +12,6 @@ from vllm._custom_ops import (
    cpu_attn_get_scheduler_metadata,
    cpu_attn_reshape_and_cache,
 )
-from vllm.platforms import CpuArchEnum, current_platform
 from vllm.utils.argparse_utils import FlexibleArgumentParser
 from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
 from vllm.v1.attention.backends.cpu_attn import CPUAttentionBackend, _get_attn_isa
@@ -22,15 +21,14 @@ def get_attn_isa(
    block_size: int | None = None,
    dtype: torch.dtype | None = None,
 ):
-    if block_size and dtype:
-        return _get_attn_isa(dtype, block_size)
-    else:
-        if current_platform.get_cpu_architecture() == CpuArchEnum.ARM:
-            return "neon"
-        elif torch.cpu._is_amx_tile_supported():
-            return "amx"
-        else:
-            return "vec"
+    # Delegate to _get_attn_isa so the fallback path applies the same arch
+    # gating (e.g. RISC-V RVV is only chosen when the build's hardcoded
+    # VLEN=128 kernel is actually present; on VLEN=256 / scalar hosts it
+    # correctly falls through to vec/vec16).
+    return _get_attn_isa(
+        dtype if dtype is not None else torch.bfloat16,
+        block_size if block_size else 32,
+    )


 # rand number generation takes too much time, cache rand tensors
@@ -235,7 +233,7 @@ if __name__ == "__main__":
    )
    parser.add_argument("--use-sink", action="store_true")
    parser.add_argument(
-        "--isa", type=str, choices=["vec", "neon", "amx", "vec16"], default=None
+        "--isa", type=str, choices=["vec", "neon", "amx", "vec16", "rvv"], default=None
    )
    parser.add_argument("--seed", type=int, default=0)
    parser.add_argument("--iters", type=int, default=20)
@@ -195,10 +195,12 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
    endif()
    if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
        message(STATUS "RISC-V RVV VLEN=${VLLM_RVV_VLEN}")
+        # Sources gate FP16/BF16 paths on the compiler-provided
+        # __riscv_zvfh / __riscv_zvfbfmin macros, which GCC and clang
+        # define automatically when those extensions appear in -march.
        if(RVV_BF16_FOUND)
            message(STATUS "BF16 extension detected")
            set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
-            add_compile_definitions(RISCV_BF16_SUPPORT)
        elseif(RVV_FP16_FOUND)
            message(WARNING "BF16 functionality is not available")
            set(MARCH_FLAGS -march=rv64gcv_zvfh_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
@@ -156,4 +156,20 @@

 #endif  // __aarch64__

+// RISC-V RVV
+#ifdef __riscv_v
+  #include <riscv_vector.h>
+
+  #ifdef __riscv_zihintpause
+    #define FAST_SPINNING __riscv_pause();
+  #endif
+
+  // FP32Vec16::exp() in cpu_types_riscv.hpp already implements the full
+  // polynomial approximation for RVV, so we simply delegate to it.
+  #define DEFINE_FAST_EXP                             \
+    auto fast_exp = [&](const vec_op::FP32Vec16& vec) \
+                        __attribute__((always_inline)) { return vec.exp(); };
+
+#endif  // __riscv_v
+
 #endif
@@ -29,6 +29,8 @@ torch::Tensor get_scheduler_metadata(
    isa = cpu_attention::ISA::NEON;
  } else if (isa_hint == "vxe") {
    isa = cpu_attention::ISA::VXE;
+  } else if (isa_hint == "rvv") {
+    isa = cpu_attention::ISA::RVV;
  } else if (isa_hint == "vsx") {
    isa = cpu_attention::ISA::VSX;
  } else {
@@ -131,6 +133,8 @@ void cpu_attn_reshape_and_cache(
      return cpu_attention::ISA::NEON;
    } else if (isa == "vxe") {
      return cpu_attention::ISA::VXE;
+    } else if (isa == "rvv") {
+      return cpu_attention::ISA::RVV;
    } else if (isa == "vsx") {
      return cpu_attention::ISA::VSX;
    } else {
@@ -12,7 +12,7 @@
 #include "cpu/utils.hpp"

 namespace cpu_attention {
-enum class ISA { AMX, VEC, VEC16, NEON, VXE, VSX };
+enum class ISA { AMX, VEC, VEC16, NEON, VXE, RVV, VSX };

 // Mirrors csrc/attention/dtype_fp8.cuh Fp8KVCacheDataType exactly.
 enum class Fp8KVCacheDataType {
@@ -164,6 +164,9 @@ struct AttentionMetadata {
      case ISA::VXE:
        ss << "VXE, ";
        break;
+      case ISA::RVV:
+        ss << "RVV, ";
+        break;
      case ISA::VSX:
        ss << "VSX, ";
        break;
@@ -0,0 +1,445 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+#ifndef CPU_ATTN_RVV_HPP
+#define CPU_ATTN_RVV_HPP
+
+// This kernel is currently hardcoded to VLEN=128 (m1/m2 intrinsics, vl=8).
+// The fixed-width typedefs below use `riscv_rvv_vector_bits(128)`, which
+// only matches `vfloat16m1_t`/`vuint16m1_t` register layout when VLEN==128;
+// at VLEN>=256 those typedefs fail to compile.  Scalar RISC-V builds
+// (-march=rv64gc) additionally don't have <riscv_vector.h>.  For both
+// cases we omit the file entirely and let the dispatcher fall back to the
+// scalar VEC / VEC16 implementations.  TODO: migrate to RVVI() macros +
+// semantic names in cpu_types_riscv_defs.hpp to support VLEN>=256 natively.
+#if defined(__riscv_v_min_vlen) && __riscv_v_min_vlen == 128
+
+  #include "cpu_attn_impl.hpp"
+  #include <riscv_vector.h>
+  #include <type_traits>
+
+namespace cpu_attention {
+
+namespace {
+
+// File-local concrete-LMUL typedefs.  The shared _defs.hpp exposes
+// VLEN-independent semantic names (fixed_fp32x8_t, fixed_fp16x8_t, ...),
+// but this kernel is currently hardcoded to VLEN=128 (m1/m2 intrinsics),
+// so keep the legacy concrete aliases scoped to this file.
+typedef vfloat16m1_t fixed_vfloat16m1_t
+    __attribute__((riscv_rvv_vector_bits(128)));
+typedef vfloat32m2_t fixed_vfloat32m2_t
+    __attribute__((riscv_rvv_vector_bits(256)));
+typedef vuint16m1_t fixed_vuint16m1_t
+    __attribute__((riscv_rvv_vector_bits(128)));
+typedef vuint32m2_t fixed_vuint32m2_t
+    __attribute__((riscv_rvv_vector_bits(256)));
+  #ifdef __riscv_zvfbfmin
+typedef vbfloat16m1_t fixed_vbfloat16m1_t
+    __attribute__((riscv_rvv_vector_bits(128)));
+  #endif
+
+  #define BLOCK_SIZE_ALIGNMENT 32
+  #define HEAD_SIZE_ALIGNMENT 32
+  #define MAX_Q_HEAD_NUM_PER_ITER 16
+
+// ============================================================================
+// B-matrix row loading: load 8 elements as FP32 (using m2 LMUL at VLEN=128)
+// ============================================================================
+
+template <typename kv_cache_t>
+FORCE_INLINE fixed_vfloat32m2_t load_row8_B_as_f32(const kv_cache_t* p);
+
+template <>
+FORCE_INLINE fixed_vfloat32m2_t load_row8_B_as_f32<float>(const float* p) {
+  return __riscv_vle32_v_f32m2(p, 8);
+}
+
+template <>
+FORCE_INLINE fixed_vfloat32m2_t
+load_row8_B_as_f32<c10::Half>(const c10::Half* p) {
+  #ifdef __riscv_zvfh
+  fixed_vfloat16m1_t h =
+      __riscv_vle16_v_f16m1(reinterpret_cast<const _Float16*>(p), 8);
+  return __riscv_vfwcvt_f_f_v_f32m2(h, 8);
+  #else
+  // Fallback for hardware without Zvfh: scalar half->float conversion.
+  // c10::Half provides operator float() so this is correct on any RVV CPU
+  // that has only the base V extension.  Slower than the Zvfh path, but
+  // keeps the kernel buildable on Zvfhmin-only / no-fp16 hardware.
+  alignas(16) float tmp[8];
+  for (int i = 0; i < 8; ++i) {
+    tmp[i] = static_cast<float>(p[i]);
+  }
+  return __riscv_vle32_v_f32m2(tmp, 8);
+  #endif
+}
+
+template <>
+FORCE_INLINE fixed_vfloat32m2_t
+load_row8_B_as_f32<c10::BFloat16>(const c10::BFloat16* p) {
+  #ifdef __riscv_zvfbfmin
+  fixed_vbfloat16m1_t bf =
+      __riscv_vle16_v_bf16m1(reinterpret_cast<const __bf16*>(p), 8);
+  return __riscv_vfwcvtbf16_f_f_v_f32m2(bf, 8);
+  #else
+  // Fallback: load as uint16, zero-extend to uint32, shift left by 16
+  fixed_vuint16m1_t raw =
+      __riscv_vle16_v_u16m1(reinterpret_cast<const uint16_t*>(p), 8);
+  fixed_vuint32m2_t wide = __riscv_vzext_vf2_u32m2(raw, 8);
+  fixed_vuint32m2_t shifted = __riscv_vsll_vx_u32m2(wide, 16, 8);
+  return __riscv_vreinterpret_v_u32m2_f32m2(shifted);
+  #endif
+}
+
+// ============================================================================
+// Micro kernel: Mx8 tile, K unrolled by 4, RVV scalar-broadcast FMA
+// ============================================================================
+//
+// NEON uses vfmaq_laneq_f32 (lane-indexed FMA from a preloaded A vector).
+// RVV has no lane-indexed FMA; instead we load A elements as scalars and
+// use __riscv_vfmacc_vf (scalar * vector + accumulator), which is equally
+// efficient and avoids the need for vrgather/vslidedown.
+//
+// At VLEN=128, m2 holds 8 x FP32, matching the 8-column tile width.
+// Register budget: M accumulators (m2 each) + 1 B temp = 2M+2 regs.
+// M=8 => 18 regs out of 32 available — no spills.
+
+template <int32_t M, typename kv_cache_t>
+FORCE_INLINE void gemm_micro_rvv_fma_Mx8_Ku4(
+    const float* __restrict A,       // [M x K]
+    const kv_cache_t* __restrict B,  // [K x 8]
+    float* __restrict C,             // [M x 8]
+    int64_t lda, int64_t ldb, int64_t ldc, int32_t K, bool accumulate) {
+  static_assert(1 <= M && M <= 8, "M must be in [1,8]");
+
+  constexpr size_t vl = 8;
+
+  // helpers for per-M codegen
+  #define ROWS_APPLY(OP) OP(0) OP(1) OP(2) OP(3) OP(4) OP(5) OP(6) OP(7)
+  #define IF_M(i) if constexpr (M > (i))
+
+    // A row base pointers
+  #define DECL_A(i) const float* a##i = A + (i) * lda;
+  ROWS_APPLY(DECL_A)
+  #undef DECL_A
+
+    // declare one m2 accumulator per row
+  #define DECL_ACC(i) fixed_vfloat32m2_t acc##i;
+  ROWS_APPLY(DECL_ACC)
+  #undef DECL_ACC
+
+    // initialize accumulators
+  #define INIT_ACC(i)                                      \
+    IF_M(i) {                                              \
+      if (accumulate) {                                    \
+        acc##i = __riscv_vle32_v_f32m2(C + (i) * ldc, vl); \
+      } else {                                             \
+        acc##i = __riscv_vfmv_v_f_f32m2(0.f, vl);          \
+      }                                                    \
+    }
+  ROWS_APPLY(INIT_ACC)
+  #undef INIT_ACC
+
+  int32_t k = 0;
+
+  // K unrolled by 4
+  for (; k + 3 < K; k += 4) {
+    // k + 0
+    {
+      fixed_vfloat32m2_t b =
+          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 0) * ldb);
+  #define STEP_K0(i)                                                    \
+    IF_M(i) {                                                           \
+      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 0), b, vl); \
+    }
+      ROWS_APPLY(STEP_K0)
+  #undef STEP_K0
+    }
+    // k + 1
+    {
+      fixed_vfloat32m2_t b =
+          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 1) * ldb);
+  #define STEP_K1(i)                                                    \
+    IF_M(i) {                                                           \
+      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 1), b, vl); \
+    }
+      ROWS_APPLY(STEP_K1)
+  #undef STEP_K1
+    }
+    // k + 2
+    {
+      fixed_vfloat32m2_t b =
+          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 2) * ldb);
+  #define STEP_K2(i)                                                    \
+    IF_M(i) {                                                           \
+      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 2), b, vl); \
+    }
+      ROWS_APPLY(STEP_K2)
+  #undef STEP_K2
+    }
+    // k + 3
+    {
+      fixed_vfloat32m2_t b =
+          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 3) * ldb);
+  #define STEP_K3(i)                                                    \
+    IF_M(i) {                                                           \
+      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 3), b, vl); \
+    }
+      ROWS_APPLY(STEP_K3)
+  #undef STEP_K3
+    }
+  }
+
+  // K tail
+  for (; k < K; ++k) {
+    fixed_vfloat32m2_t b = load_row8_B_as_f32<kv_cache_t>(B + (int64_t)k * ldb);
+  #define TAIL_ROW(i) \
+    IF_M(i) { acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k), b, vl); }
+    ROWS_APPLY(TAIL_ROW)
+  #undef TAIL_ROW
+  }
+
+    // store accumulators to C
+  #define STORE_ROW(i) \
+    IF_M(i) { __riscv_vse32_v_f32m2(C + (i) * ldc, acc##i, vl); }
+  ROWS_APPLY(STORE_ROW)
+  #undef STORE_ROW
+
+  #undef ROWS_APPLY
+  #undef IF_M
+}
+
+// ============================================================================
+// Macro kernel: dispatch M tiles of {8,4,2,1}, step N by 8
+// ============================================================================
+
+template <int32_t N, typename kv_cache_t>
+FORCE_INLINE void gemm_macro_rvv_fma_Mx8_Ku4(const float* __restrict A,
+                                             const kv_cache_t* __restrict B,
+                                             float* __restrict C, int32_t M,
+                                             int32_t K, int64_t lda,
+                                             int64_t ldb, int64_t ldc,
+                                             bool accumulate) {
+  static_assert(N % 8 == 0, "N must be a multiple of 8");
+  for (int32_t m = 0; m < M;) {
+    int32_t mb = (M - m >= 8) ? 8 : (M - m >= 4) ? 4 : (M - m >= 2) ? 2 : 1;
+    const float* Ab = A + m * lda;
+    float* Cb = C + m * ldc;
+
+    for (int32_t n = 0; n < N; n += 8) {
+      const kv_cache_t* Bn = B + n;
+      float* Cn = Cb + n;
+      switch (mb) {
+        case 8:
+          gemm_micro_rvv_fma_Mx8_Ku4<8, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
+                                                    K, accumulate);
+          break;
+        case 4:
+          gemm_micro_rvv_fma_Mx8_Ku4<4, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
+                                                    K, accumulate);
+          break;
+        case 2:
+          gemm_micro_rvv_fma_Mx8_Ku4<2, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
+                                                    K, accumulate);
+          break;
+        default:
+          gemm_micro_rvv_fma_Mx8_Ku4<1, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
+                                                    K, accumulate);
+          break;
+      }
+    }
+    m += mb;
+  }
+}
+
+// ============================================================================
+// TileGemm wrapper — plugs into AttentionMainLoop
+// ============================================================================
+
+template <typename kv_cache_t>
+class TileGemmRVV {
+ public:
+  template <AttentionGemmPhase phase, int32_t k_size>
+  FORCE_INLINE static void gemm(const int32_t m_size,
+                                float* __restrict__ a_tile,
+                                kv_cache_t* __restrict__ b_tile,
+                                float* __restrict__ c_tile, const int64_t lda,
+                                const int64_t ldb, const int64_t ldc,
+                                const int32_t block_size,
+                                const int32_t dynamic_k_size,
+                                const bool accum_c) {
+    if constexpr (phase == AttentionGemmPhase::QK) {
+      gemm_macro_rvv_fma_Mx8_Ku4<BLOCK_SIZE_ALIGNMENT, kv_cache_t>(
+          a_tile, b_tile, c_tile, m_size, k_size, lda, ldb, ldc, accum_c);
+    } else {
+      gemm_macro_rvv_fma_Mx8_Ku4<HEAD_SIZE_ALIGNMENT, kv_cache_t>(
+          a_tile, b_tile, c_tile, m_size, dynamic_k_size, lda, ldb, ldc,
+          accum_c);
+    }
+  }
+};
+
+}  // namespace
+
+// ============================================================================
+// AttentionImpl<ISA::RVV> — mirrors ISA::NEON specialization
+// ============================================================================
+
+template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
+class AttentionImpl<ISA::RVV, scalar_t, head_dim, kv_cache_scalar_t> {
+ public:
+  using query_t = scalar_t;
+  using q_buffer_t = float;
+  using kv_cache_t = scalar_t;
+  using logits_buffer_t = float;
+  using partial_output_buffer_t = float;
+  using prob_buffer_t = float;
+
+  constexpr static int64_t BlockSizeAlignment = BLOCK_SIZE_ALIGNMENT;
+  constexpr static int64_t HeadDimAlignment = HEAD_SIZE_ALIGNMENT;
+  constexpr static int64_t MaxQHeadNumPerIteration = MAX_Q_HEAD_NUM_PER_ITER;
+  constexpr static int64_t HeadDim = head_dim;
+  constexpr static ISA ISAType = ISA::RVV;
+  constexpr static bool scale_on_logits = false;
+
+  static_assert(HeadDim % HeadDimAlignment == 0);
+  static_assert(HeadDimAlignment % 8 == 0);
+  static_assert(BlockSizeAlignment % 8 == 0);
+
+ public:
+  template <template <typename tile_gemm_t> typename attention>
+  FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
+    attention<TileGemmRVV<kv_cache_t>> attention_iteration;
+    attention_iteration(CPU_ATTENTION_PARAMS);
+  }
+
+  constexpr static int64_t k_cache_token_group_stride(
+      const int32_t block_size) {
+    return BlockSizeAlignment;
+  }
+
+  constexpr static int64_t v_cache_token_group_stride(
+      const int32_t block_size) {
+    return head_dim * BlockSizeAlignment;
+  }
+
+  constexpr static int64_t v_cache_head_group_stride(const int32_t block_size) {
+    return HeadDimAlignment;
+  }
+
+  static void copy_q_heads_tile(scalar_t* __restrict__ src,
+                                float* __restrict__ q_buffer,
+                                const int32_t q_num,
+                                const int32_t q_heads_per_kv,
+                                const int64_t q_num_stride,
+                                const int64_t q_head_stride, float scale) {
+    static_assert(head_dim % 16 == 0);
+    constexpr int32_t unroll_size = head_dim / 16;
+    using load_vec_t = typename VecTypeTrait<scalar_t>::vec_t;
+
+    vec_op::FP32Vec16 scale_vec(scale);
+    for (int32_t q_num_idx = 0; q_num_idx < q_num; ++q_num_idx) {
+      for (int32_t q_head_idx = 0; q_head_idx < q_heads_per_kv; ++q_head_idx) {
+        scalar_t* __restrict__ curr_q =
+            src + q_num_idx * q_num_stride + q_head_idx * q_head_stride;
+        float* __restrict__ curr_q_buffer =
+            q_buffer + q_num_idx * q_heads_per_kv * head_dim +
+            q_head_idx * head_dim;
+
+        vec_op::unroll_loop<int32_t, unroll_size>([&](int32_t i) {
+          load_vec_t vec(curr_q);
+          vec_op::FP32Vec16 fp32_vec(vec);
+          fp32_vec = fp32_vec * scale_vec;
+          fp32_vec.save(curr_q_buffer);
+
+          curr_q += 16;
+          curr_q_buffer += 16;
+        });
+      }
+    }
+  }
+
+  static void reshape_and_cache(
+      const scalar_t* __restrict__ key, const scalar_t* __restrict__ value,
+      scalar_t* __restrict__ key_cache, scalar_t* __restrict__ value_cache,
+      const int64_t* __restrict__ slot_mapping, const int64_t token_num,
+      const int64_t key_token_num_stride, const int64_t value_token_num_stride,
+      const int64_t head_num, const int64_t key_head_num_stride,
+      const int64_t value_head_num_stride, const int64_t num_blocks,
+      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
+      const int64_t block_size, const int64_t block_size_stride,
+      const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
+  #pragma omp parallel for collapse(2)
+    for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
+      for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
+        const int64_t pos = slot_mapping[token_idx];
+        if (pos < 0) {
+          continue;
+        }
+
+        const int64_t block_idx = pos / block_size;
+        const int64_t block_offset = pos % block_size;
+        {
+          // Write Key (transpose to column-major: [head_dim, block_size])
+          const scalar_t* key_start_ptr = key +
+                                          token_idx * key_token_num_stride +
+                                          head_idx * key_head_num_stride;
+          scalar_t* key_cache_start_ptr =
+              key_cache + block_idx * num_blocks_stride +
+              head_idx * cache_head_num_stride + block_offset;
+
+          // Strided vector store for efficient transpose.
+          // Load contiguous key elements, store with stride = block_size.
+          {
+            const ptrdiff_t byte_stride = block_size * sizeof(scalar_t);
+            int64_t i = 0;
+            for (; i < head_dim;) {
+              size_t vl;
+              if constexpr (std::is_same_v<scalar_t, float>) {
+                vl = __riscv_vsetvl_e32m2(head_dim - i);
+                vfloat32m2_t v = __riscv_vle32_v_f32m2(
+                    reinterpret_cast<const float*>(key_start_ptr + i), vl);
+                __riscv_vsse32_v_f32m2(
+                    reinterpret_cast<float*>(key_cache_start_ptr +
+                                             i * block_size),
+                    byte_stride, v, vl);
+              } else {
+                // Half and BFloat16 are both 16-bit types
+                vl = __riscv_vsetvl_e16m1(head_dim - i);
+                vuint16m1_t v = __riscv_vle16_v_u16m1(
+                    reinterpret_cast<const uint16_t*>(key_start_ptr + i), vl);
+                __riscv_vsse16_v_u16m1(
+                    reinterpret_cast<uint16_t*>(key_cache_start_ptr +
+                                                i * block_size),
+                    byte_stride, v, vl);
+              }
+              i += vl;
+            }
+          }
+        }
+        {
+          // Write Value (row-major: [block_size, head_dim])
+          const scalar_t* value_start_ptr = value +
+                                            token_idx * value_token_num_stride +
+                                            head_idx * value_head_num_stride;
+          scalar_t* value_cache_start_ptr =
+              value_cache + block_idx * num_blocks_stride +
+              head_idx * cache_head_num_stride + block_offset * head_dim;
+          std::memcpy(value_cache_start_ptr, value_start_ptr,
+                      sizeof(scalar_t) * head_dim);
+        }
+      }
+    }
+  }
+};
+
+}  // namespace cpu_attention
+
+  #undef BLOCK_SIZE_ALIGNMENT
+  #undef HEAD_SIZE_ALIGNMENT
+  #undef MAX_Q_HEAD_NUM_PER_ITER
+
+#endif  // __riscv_v_min_vlen == 128
+
+#endif  // CPU_ATTN_RVV_HPP
@@ -72,7 +72,7 @@ typedef RVVTYPE(vuint16, LMUL_512, _t) fixed_u16x32_t
    __attribute__((riscv_rvv_vector_bits(512)));

 // bfloat16
-#ifdef RISCV_BF16_SUPPORT
+#ifdef __riscv_zvfbfmin
 typedef RVVTYPE(vbfloat16, LMUL_128, _t) fixed_bf16x8_t
    __attribute__((riscv_rvv_vector_bits(128)));
 typedef RVVTYPE(vbfloat16, LMUL_256, _t) fixed_bf16x16_t
@@ -15,8 +15,17 @@
 #include <torch/all.h>
 namespace vec_op {

-// BFloat16 is always supported on RISC-V: natively when RISCV_BF16_SUPPORT
-// is defined, otherwise via the FP32-simulation fallback path.
+// FP8 KV cache is not supported on RISC-V. These tag types and the
+// corresponding BF16Vec32 stub constructors below exist solely so that
+// templates referencing vec_op::fp8_*_tag in their bodies (e.g. in
+// cpu_attn_vec.hpp) compile under GCC's -Wtemplate-body lookup. The
+// stubs are never instantiated by CPU_ATTN_DISPATCH on __riscv.
+struct fp8_e4m3_tag {};
+struct fp8_e5m2_tag {};
+
+// BFloat16 is always supported on RISC-V: natively when __riscv_zvfbfmin
+// is defined (compiler-provided when -march includes zvfbfmin), otherwise
+// via the FP32-simulation fallback path.
 #define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
  AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)  \
@@ -106,7 +115,7 @@ struct FP16Vec16 : public Vec<FP16Vec16> {
 // BF16 Implementation
 // ============================================================================

-#ifdef RISCV_BF16_SUPPORT
+#ifdef __riscv_zvfbfmin

 FORCE_INLINE fixed_u16x8_t bf16_to_u16(fixed_bf16x8_t v) {
  return RVVI4(__riscv_vreinterpret_v_bf16, LMUL_128, _u16, LMUL_128)(v);
@@ -185,6 +194,13 @@ struct BF16Vec32 : public Vec<BF16Vec32> {

  explicit BF16Vec32(fixed_bf16x32_t data) : reg(data) {};

+  // FP8 KV cache stubs: never instantiated on RISC-V (CPU_ATTN_DISPATCH
+  // omits FP8 cases on __riscv); exist only so name lookup succeeds.
+  explicit BF16Vec32(const uint8_t* ptr, fp8_e4m3_tag)
+      : BF16Vec32(static_cast<const void*>(ptr)) {}
+  explicit BF16Vec32(const uint8_t* ptr, fp8_e5m2_tag)
+      : BF16Vec32(static_cast<const void*>(ptr)) {}
+
  explicit BF16Vec32(const BF16Vec8& v) {
    fixed_u16x8_t u16_val = bf16_to_u16(v.reg);
    fixed_u16x32_t u16_combined =
@@ -323,6 +339,13 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
    reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_1024)(tmp, 32);
  }

+  // FP8 KV cache stubs: never instantiated on RISC-V (CPU_ATTN_DISPATCH
+  // omits FP8 cases on __riscv); exist only so name lookup succeeds.
+  explicit BF16Vec32(const uint8_t* ptr, fp8_e4m3_tag)
+      : BF16Vec32(static_cast<const void*>(ptr)) {}
+  explicit BF16Vec32(const uint8_t* ptr, fp8_e5m2_tag)
+      : BF16Vec32(static_cast<const void*>(ptr)) {}
+
  explicit BF16Vec32(const BF16Vec8& v) {
    float tmp_small[8];
    RVVI(__riscv_vse32_v_f32, LMUL_256)(tmp_small, v.reg_fp32, 8);
@@ -410,7 +433,7 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
  explicit FP32Vec8(fixed_fp16x8_t v)
      : reg(RVVI(__riscv_vfwcvt_f_f_v_f32, LMUL_256)(v, VEC_ELEM_NUM)) {};

-#ifdef RISCV_BF16_SUPPORT
+#ifdef __riscv_zvfbfmin
  explicit FP32Vec8(fixed_bf16x8_t v)
      : reg(RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_256)(v, VEC_ELEM_NUM)) {};
  explicit FP32Vec8(const BF16Vec8& v)
@@ -608,7 +631,7 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
  explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
  explicit FP32Vec16(const FP16Vec16& v);

-#ifdef RISCV_BF16_SUPPORT
+#ifdef __riscv_zvfbfmin
  explicit FP32Vec16(fixed_bf16x16_t v)
      : reg(RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_512)(v, VEC_ELEM_NUM)) {};
  explicit FP32Vec16(const BF16Vec16& v)
@@ -868,7 +891,7 @@ inline void fma(FP32Vec16& acc, const FP32Vec16& a, const FP32Vec16& b) {
  acc = acc.fma(a, b);
 }

-#ifdef RISCV_BF16_SUPPORT
+#ifdef __riscv_zvfbfmin
 template <>
 inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
  *ptr = static_cast<__bf16>(v);
@@ -20,7 +20,8 @@ ISA_TYPES = {
    "VEC16": 2,
    "NEON": 3,
    "VXE": 4,
-    "VSX": 5,
+    "RVV": 5,
+    "VSX": 6,
 }

 # KV cache index: 0 = auto (same as scalar_t), 1 = fp8_e4m3, 2 = fp8_e5m2
@@ -38,7 +39,7 @@ KV_CACHE_CPP_TYPES = {
 }

 # ISAs supported for head_dims divisible by 32
-ISA_FOR_32 = ["AMX", "NEON", "VEC", "VEC16", "VXE", "VSX"]
+ISA_FOR_32 = ["AMX", "NEON", "VEC", "VEC16", "VXE", "RVV", "VSX"]

 # ISAs supported for head_dims divisible by 16 only
 ISA_FOR_16 = ["VEC16"]
@@ -149,6 +150,15 @@ def generate_header_file() -> str:
  #include "cpu_attn_vxe.hpp"
 #endif

+// cpu_attn_rvv.hpp is hardcoded to VLEN==128 (m1/m2 intrinsics, vl=8) and
+// itself includes <riscv_vector.h>, which is unavailable on scalar
+// (-march=rv64gc) builds. Gate the include the same way as the dispatch
+// macro below, so non-128 / scalar RISC-V builds skip it entirely.
+#if defined(__riscv) && defined(__riscv_v_min_vlen) && \
+    __riscv_v_min_vlen == 128
+  #include "cpu_attn_rvv.hpp"
+#endif
+
 #ifdef __powerpc__
  #include "cpu_attn_vsx.hpp"
 #endif
@@ -212,6 +222,23 @@ def generate_header_file() -> str:
        ["VXE", "VEC", "VEC16"],
        fp8=False,
    )
+    # RISC-V with RVV.  cpu_attn_rvv.hpp is hardcoded to VLEN==128
+    # (riscv_rvv_vector_bits(128) typedefs + vl=8 m1/m2 intrinsics), so
+    # we split the dispatch into two top-level branches: VLEN==128 builds
+    # get the full RVV+VEC+VEC16 case set, other VLEN builds get a
+    # VEC/VEC16-only fallback.  Preprocessor directives cannot appear
+    # inside a #define body, so this duplication is necessary.
+    header += _macro_block(
+        "#elif defined(__riscv) && defined(__riscv_v_min_vlen) "
+        "&& __riscv_v_min_vlen == 128",
+        ["RVV", "VEC", "VEC16"],
+        fp8=False,
+    )
+    header += _macro_block(
+        "#elif defined(__riscv)",
+        ["VEC", "VEC16"],
+        fp8=False,
+    )
    header += _macro_block(
        "#elif defined(__powerpc__)",
        ["VSX", "VEC", "VEC16"],
@@ -233,8 +260,8 @@ def generate_header_file() -> str:
        fp8=False,
    )
    header += (
-        "#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / "
-        "__s390x__ / __powerpc__ */\n\n"
+        "#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / __s390x__ /"
+        " __riscv / __powerpc__ */\n\n"
        "#endif  // CPU_ATTN_DISPATCH_GENERATED_H\n"
    )

@@ -0,0 +1,30 @@
+/*
+ * Fused RMSNorm + router GEMV for DeepSeek V4.
+ *
+ * Computes in a single kernel:
+ *   normed_x[m,k]   = x[m,k] * rsqrt(mean(x[m]^2) + eps) * norm_weight[k]
+ *   router_logits[m,n] = sum_k(normed_x[m,k] * gate_weight[n,k])
+ *
+ * The GEMV body mirrors the algorithm in csrc/moe/dsv3_router_gemm_*.cu
+ * (warp butterfly + smem cross-warp reduction, fp32 accumulation, PDL on
+ * SM90+).  Blocks 0..kNumTokens-1 each materialize one token's normed_x
+ * row to global memory using the algebraic identity
+ *      logits[m,n] = rsqrt[m] * sum_k(x[m,k] * nw[k] * gw[n,k])
+ * which lets every block produce its column of logits before normed_x
+ * exists in gmem.
+ *
+ * Logits output is fp32 only — DeepSeek V4 router gate is hard-coded to
+ * fp32 (vllm/model_executor/models/deepseek_v4.py:749).
+ */
+
+#pragma once
+
+#include <cuda_bf16.h>
+#include <cuda_runtime.h>
+
+#include "dsv3_router_gemm_utils.h"
+
+template <typename T, int kNumTokens, int kNumExperts, int kHiddenDim>
+void invokeNormRouterGemm(float* logits, __nv_bfloat16* normed_x, T const* x,
+                          T const* norm_weight, T const* gate_weight, float eps,
+                          cudaStream_t stream);
@@ -0,0 +1,130 @@
+/*
+ * TORCH op entry for the fused RMSNorm + router GEMV kernel
+ * (DeepSeek V4 Pro).  This op is DSV4-Pro-specific: the kernel is
+ * instantiated only for ``num_experts == 384`` and ``hidden_dim ==
+ * 7168``.  Other configurations (e.g. DSV4-Flash with H=4096) must
+ * fall back to the unfused ``rms_norm`` + ``dsv3_router_gemm`` path.
+ */
+
+#include <ATen/ATen.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <torch/all.h>
+
+#include <cuda_bf16.h>
+#include <cuda_runtime.h>
+
+#include "core/registration.h"
+#include "dsv4_norm_router_gemm.h"
+
+namespace {
+
+// DSV4-Pro hard-coded shape constants.  Renamed from the earlier
+// ``kKimiK2NumExperts`` to avoid the misleading impression that this
+// kernel targets Kimi K2 — 384 happens to match Kimi K2's gate but the
+// intent here is DSV4-Pro.
+constexpr int kDsv4NumExperts = 384;
+constexpr int kDsv4HiddenDim = 7168;
+
+template <int kBegin, int kEnd>
+struct LoopUnroller {
+  static void unroll(int num_tokens, float* logits, __nv_bfloat16* normed_x,
+                     __nv_bfloat16 const* x, __nv_bfloat16 const* norm_weight,
+                     __nv_bfloat16 const* gate_weight, float eps,
+                     cudaStream_t stream) {
+    if (num_tokens == kBegin) {
+      invokeNormRouterGemm<__nv_bfloat16, kBegin, kDsv4NumExperts,
+                           kDsv4HiddenDim>(logits, normed_x, x, norm_weight,
+                                           gate_weight, eps, stream);
+    } else {
+      LoopUnroller<kBegin + 1, kEnd>::unroll(num_tokens, logits, normed_x, x,
+                                             norm_weight, gate_weight, eps,
+                                             stream);
+    }
+  }
+};
+
+template <int kEnd>
+struct LoopUnroller<kEnd, kEnd> {
+  static void unroll(int num_tokens, float* logits, __nv_bfloat16* normed_x,
+                     __nv_bfloat16 const* x, __nv_bfloat16 const* norm_weight,
+                     __nv_bfloat16 const* gate_weight, float eps,
+                     cudaStream_t stream) {
+    if (num_tokens == kEnd) {
+      invokeNormRouterGemm<__nv_bfloat16, kEnd, kDsv4NumExperts,
+                           kDsv4HiddenDim>(logits, normed_x, x, norm_weight,
+                                           gate_weight, eps, stream);
+    } else {
+      throw std::invalid_argument(
+          "Invalid num_tokens, only supports 1 to 16 for "
+          "dsv4_norm_router_gemm");
+    }
+  }
+};
+
+}  // namespace
+
+void dsv4_norm_router_gemm(at::Tensor& logits,    // [num_tokens, E] fp32
+                           at::Tensor& normed_x,  // [num_tokens, H] bf16
+                           at::Tensor const& x,   // [num_tokens, H] bf16
+                           at::Tensor const& norm_weight,  // [H] bf16
+                           at::Tensor const& gate_weight,  // [E, H] bf16
+                           double eps) {
+  TORCH_CHECK(x.dim() == 2 && norm_weight.dim() == 1 && gate_weight.dim() == 2,
+              "x must be 2D, norm_weight 1D, gate_weight 2D");
+  TORCH_CHECK(logits.dim() == 2 && normed_x.dim() == 2,
+              "logits and normed_x must be 2D");
+
+  int const num_tokens = x.size(0);
+  int const hidden_dim = x.size(1);
+  int const num_experts = gate_weight.size(0);
+
+  TORCH_CHECK(hidden_dim == kDsv4HiddenDim,
+              "Expected hidden_dim=", kDsv4HiddenDim,
+              " (DSV4-Pro), but got hidden_dim=", hidden_dim);
+  TORCH_CHECK(gate_weight.size(1) == hidden_dim,
+              "gate_weight.shape[1] must equal x.shape[1]");
+  TORCH_CHECK(norm_weight.size(0) == hidden_dim,
+              "norm_weight.shape[0] must equal x.shape[1]");
+  TORCH_CHECK(num_experts == kDsv4NumExperts,
+              "Expected num_experts=", kDsv4NumExperts,
+              " (DSV4-Pro), but got num_experts=", num_experts);
+  TORCH_CHECK(num_tokens >= 1 && num_tokens <= 16,
+              "num_tokens must be in [1, 16] for dsv4_norm_router_gemm");
+
+  TORCH_CHECK(x.dtype() == at::kBFloat16, "x must be bf16");
+  TORCH_CHECK(norm_weight.dtype() == at::kBFloat16, "norm_weight must be bf16");
+  TORCH_CHECK(gate_weight.dtype() == at::kBFloat16, "gate_weight must be bf16");
+  TORCH_CHECK(normed_x.dtype() == at::kBFloat16, "normed_x must be bf16");
+  TORCH_CHECK(logits.dtype() == at::kFloat,
+              "logits must be float32 (DSV4 router output is hard-coded fp32)");
+
+  TORCH_CHECK(normed_x.size(0) == num_tokens && normed_x.size(1) == hidden_dim,
+              "normed_x must be [num_tokens, hidden_dim]");
+  TORCH_CHECK(logits.size(0) == num_tokens && logits.size(1) == num_experts,
+              "logits must be [num_tokens, num_experts]");
+
+  TORCH_CHECK(x.is_contiguous() && norm_weight.is_contiguous() &&
+                  gate_weight.is_contiguous() && normed_x.is_contiguous() &&
+                  logits.is_contiguous(),
+              "all tensors must be contiguous");
+
+  auto const sm = getSMVersion();
+  TORCH_CHECK(sm >= 90 && sm <= 103,
+              "dsv4_norm_router_gemm requires SM_90 <= CUDA ARCH <= SM_103");
+
+  cudaStream_t const stream = at::cuda::getCurrentCUDAStream();
+
+  auto* logits_ptr = reinterpret_cast<float*>(logits.mutable_data_ptr());
+  auto* nx_ptr = reinterpret_cast<__nv_bfloat16*>(normed_x.mutable_data_ptr());
+  auto* x_ptr = reinterpret_cast<__nv_bfloat16 const*>(x.data_ptr());
+  auto* nw_ptr = reinterpret_cast<__nv_bfloat16 const*>(norm_weight.data_ptr());
+  auto* gw_ptr = reinterpret_cast<__nv_bfloat16 const*>(gate_weight.data_ptr());
+  float const eps_f = static_cast<float>(eps);
+
+  LoopUnroller<1, 16>::unroll(num_tokens, logits_ptr, nx_ptr, x_ptr, nw_ptr,
+                              gw_ptr, eps_f, stream);
+}
+
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+  m.impl("dsv4_norm_router_gemm", &dsv4_norm_router_gemm);
+}
@@ -0,0 +1,249 @@
+/*
+ * Fused RMSNorm + router GEMV for DeepSeek V4 (logits are fp32; bf16
+ * output is unsupported because DSV4 hard-codes fp32 logits).  See
+ * dsv4_norm_router_gemm.h for the math.
+ *
+ * The GEMV body mirrors csrc/moe/dsv3_router_gemm_float_out.cu (warp
+ * butterfly reduction + smem cross-warp reduction, fp32 accumulation,
+ * 128-thread block, PDL on SM90+).  RMSNorm is folded into the same
+ * pass via the identity
+ *   logits[m,n] = rsqrt[m] * sum_k(x[m,k] * nw[k] * gw[n,k])
+ * so x is read exactly once per block during the GEMV phase.  Blocks
+ * 0..kNumTokens-1 each materialize one row of normed_x for downstream
+ * experts / shared_experts to consume.
+ */
+
+#include <ATen/ATen.h>
+#include <ATen/cuda/CUDAContext.h>
+
+#include <cuda_bf16.h>
+#include <cuda_runtime.h>
+
+#include "dsv4_norm_router_gemm.h"
+
+namespace {
+
+// Convert 8 bf16 values packed in uint4 into 8 floats. Mirrors the helper
+// in dsv3_router_gemm_float_out.cu (kept local so the dsv3 file stays
+// untouched).
+template <int VPT>
+__device__ __forceinline__ void bf16_uint4_to_float8(uint4 const& vec,
+                                                     float* dst) {
+  __nv_bfloat16* bf16_ptr =
+      reinterpret_cast<__nv_bfloat16*>(const_cast<uint4*>(&vec));
+#pragma unroll
+  for (int i = 0; i < VPT; i++) {
+    dst[i] = __bfloat162float(bf16_ptr[i]);
+  }
+}
+
+template <typename T, int kBlockSize, int VPT, int kNumTokens, int kNumExperts,
+          int kHiddenDim>
+__global__ __launch_bounds__(128, 1) void norm_router_gemm_kernel(
+    float* __restrict__ logits, __nv_bfloat16* __restrict__ normed_x,
+    T const* __restrict__ x, T const* __restrict__ norm_weight,
+    T const* __restrict__ gate_weight, float eps) {
+  static_assert(kBlockSize == 128, "kernel assumes blockDim.x == 128");
+  static_assert(kHiddenDim % (VPT * kBlockSize) == 0,
+                "kHiddenDim must be a multiple of VPT * kBlockSize");
+
+  int const n_idx = blockIdx.x;
+  int const tid = threadIdx.x;
+  constexpr int kWarpSize = 32;
+  constexpr int kNumWarps = kBlockSize / kWarpSize;
+  constexpr int k_elems_per_iter = VPT * kBlockSize;
+  constexpr int k_iterations = kHiddenDim / k_elems_per_iter;
+
+  T const* gw_col = gate_weight + n_idx * kHiddenDim;
+
+  // Per-thread accumulators — fp32 throughout, matching dsv3 / layernorm.
+  float partial[kNumTokens] = {};
+  float ss[kNumTokens] = {};
+
+  // Cross-warp reduction scratch.
+  __shared__ float sm_partial[kNumTokens][kNumWarps];
+  __shared__ float sm_ss[kNumTokens][kNumWarps];
+  __shared__ float s_rsqrt[kNumTokens];
+
+  int k_bases[k_iterations];
+#pragma unroll
+  for (int ki = 0; ki < k_iterations; ki++) {
+    k_bases[ki] = ki * k_elems_per_iter + tid * VPT;
+  }
+
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+  asm volatile("griddepcontrol.wait;");
+#endif
+
+  // ---- Phase 1: single pass over x, accumulate partial GEMV and ss. ----
+#pragma unroll
+  for (int ki = 0; ki < k_iterations; ki++) {
+    int const k_base = k_bases[ki];
+
+    uint4 nw_vec = *reinterpret_cast<uint4 const*>(norm_weight + k_base);
+    float nw_f[VPT];
+    bf16_uint4_to_float8<VPT>(nw_vec, nw_f);
+
+    uint4 b_vec = *reinterpret_cast<uint4 const*>(gw_col + k_base);
+    float b_f[VPT];
+    bf16_uint4_to_float8<VPT>(b_vec, b_f);
+
+#pragma unroll
+    for (int m = 0; m < kNumTokens; m++) {
+      uint4 a_vec =
+          *reinterpret_cast<uint4 const*>(x + m * kHiddenDim + k_base);
+      float a_f[VPT];
+      bf16_uint4_to_float8<VPT>(a_vec, a_f);
+
+#pragma unroll
+      for (int k = 0; k < VPT; k++) {
+        float a = a_f[k];
+        ss[m] += a * a;
+        partial[m] += a * nw_f[k] * b_f[k];
+      }
+    }
+  }
+
+  // ---- Phase 2: warp butterfly reduction for both ss[] and partial[]. ----
+  int const warpId = tid / kWarpSize;
+  int const laneId = tid % kWarpSize;
+
+#pragma unroll
+  for (int m = 0; m < kNumTokens; m++) {
+    float p = partial[m];
+    float s = ss[m];
+
+    p += __shfl_xor_sync(0xffffffff, p, 16);
+    s += __shfl_xor_sync(0xffffffff, s, 16);
+    p += __shfl_xor_sync(0xffffffff, p, 8);
+    s += __shfl_xor_sync(0xffffffff, s, 8);
+    p += __shfl_xor_sync(0xffffffff, p, 4);
+    s += __shfl_xor_sync(0xffffffff, s, 4);
+    p += __shfl_xor_sync(0xffffffff, p, 2);
+    s += __shfl_xor_sync(0xffffffff, s, 2);
+    p += __shfl_xor_sync(0xffffffff, p, 1);
+    s += __shfl_xor_sync(0xffffffff, s, 1);
+
+    if (laneId == 0) {
+      sm_partial[m][warpId] = p;
+      sm_ss[m][warpId] = s;
+    }
+  }
+
+  __syncthreads();
+
+  // ---- Phase 3: tid 0 finalises the reduction, writes logits, stashes
+  //               rsqrt[m] in smem for phase 4. ----
+  if (tid == 0) {
+#pragma unroll
+    for (int m = 0; m < kNumTokens; m++) {
+      float p_sum = 0.0f;
+      float s_sum = 0.0f;
+#pragma unroll
+      for (int w = 0; w < kNumWarps; w++) {
+        p_sum += sm_partial[m][w];
+        s_sum += sm_ss[m][w];
+      }
+      // Order matches layernorm_kernels.cu: rsqrtf(variance / H + eps).
+      // Use division (not multiply-by-reciprocal) to avoid an extra ULP
+      // mismatch with the reference RMSNorm.
+      float rs = rsqrtf(s_sum / static_cast<float>(kHiddenDim) + eps);
+      s_rsqrt[m] = rs;
+      logits[m * kNumExperts + n_idx] = p_sum * rs;
+    }
+  }
+
+  __syncthreads();
+
+  // ---- Phase 4: spread normed_x writes across blocks 0..kNumTokens-1.
+  //              Each writer block handles exactly one token row,
+  //              avoiding the long tail of block 0 doing all M rows.
+  //              Every block has every token's rsqrt[] in s_rsqrt
+  //              already (computed independently in phase 3), so no
+  //              cross-block synchronization is required. ----
+  if (n_idx < kNumTokens) {
+    int const m_writer = n_idx;
+    float const rs = s_rsqrt[m_writer];
+    __nv_bfloat16 const* x_row = x + m_writer * kHiddenDim;
+    __nv_bfloat16* normed_row = normed_x + m_writer * kHiddenDim;
+
+#pragma unroll
+    for (int ki = 0; ki < k_iterations; ki++) {
+      int const k_base = k_bases[ki];
+
+      uint4 nw_vec = *reinterpret_cast<uint4 const*>(norm_weight + k_base);
+      float nw_f[VPT];
+      bf16_uint4_to_float8<VPT>(nw_vec, nw_f);
+
+      uint4 a_vec = *reinterpret_cast<uint4 const*>(x_row + k_base);
+      float a_f[VPT];
+      bf16_uint4_to_float8<VPT>(a_vec, a_f);
+
+      uint4 normed_vec;
+      __nv_bfloat16* np = reinterpret_cast<__nv_bfloat16*>(&normed_vec);
+#pragma unroll
+      for (int k = 0; k < VPT; k++) {
+        np[k] = __float2bfloat16(a_f[k] * rs * nw_f[k]);
+      }
+      *reinterpret_cast<uint4*>(normed_row + k_base) = normed_vec;
+    }
+  }
+
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+  asm volatile("griddepcontrol.launch_dependents;");
+#endif
+}
+
+}  // namespace
+
+template <typename T, int kNumTokens, int kNumExperts, int kHiddenDim>
+void invokeNormRouterGemm(float* logits, __nv_bfloat16* normed_x, T const* x,
+                          T const* norm_weight, T const* gate_weight, float eps,
+                          cudaStream_t stream) {
+  constexpr int VPT = 16 / sizeof(T);
+  constexpr int kBlockSize = 128;
+
+  cudaLaunchConfig_t config;
+  config.gridDim = kNumExperts;
+  config.blockDim = kBlockSize;
+  config.dynamicSmemBytes = 0;
+  config.stream = stream;
+
+  cudaLaunchAttribute attrs[1];
+  attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
+  attrs[0].val.programmaticStreamSerializationAllowed = 1;
+  config.numAttrs = 1;
+  config.attrs = attrs;
+
+  cudaLaunchKernelEx(&config,
+                     norm_router_gemm_kernel<T, kBlockSize, VPT, kNumTokens,
+                                             kNumExperts, kHiddenDim>,
+                     logits, normed_x, x, norm_weight, gate_weight, eps);
+}
+
+// Template instantiations — DSV4-Pro is the only supported configuration:
+// num_experts=384, hidden_dim=7168.  Other shapes (e.g. DSV4-Flash with
+// hidden_dim=4096) fall back to the unfused path on the Python side.
+#define INSTANTIATE(M)                                                    \
+  template void invokeNormRouterGemm<__nv_bfloat16, M, 384, 7168>(        \
+      float*, __nv_bfloat16*, __nv_bfloat16 const*, __nv_bfloat16 const*, \
+      __nv_bfloat16 const*, float, cudaStream_t);
+
+INSTANTIATE(1)
+INSTANTIATE(2)
+INSTANTIATE(3)
+INSTANTIATE(4)
+INSTANTIATE(5)
+INSTANTIATE(6)
+INSTANTIATE(7)
+INSTANTIATE(8)
+INSTANTIATE(9)
+INSTANTIATE(10)
+INSTANTIATE(11)
+INSTANTIATE(12)
+INSTANTIATE(13)
+INSTANTIATE(14)
+INSTANTIATE(15)
+INSTANTIATE(16)
+
+#undef INSTANTIATE
@@ -75,4 +75,12 @@ void shuffle_rows(const torch::Tensor& input_tensor,
 // Supports num_tokens in [1, 16], num_experts in {256, 384}, hidden_dim = 7168
 void dsv3_router_gemm(torch::Tensor& output, const torch::Tensor& mat_a,
                      const torch::Tensor& mat_b);
+
+// Fused RMSNorm + router GEMV for DeepSeek V4. Produces both:
+//   normed_x[m,k]      = x[m,k] * rsqrt(mean(x[m]^2) + eps) * norm_weight[k]
+//   logits[m,n]        = sum_k(normed_x[m,k] * gate_weight[n,k])
+// in a single kernel launch. Same dim/dtype constraints as dsv3_router_gemm.
+void dsv4_norm_router_gemm(at::Tensor& logits, at::Tensor& normed_x,
+                           at::Tensor const& x, at::Tensor const& norm_weight,
+                           at::Tensor const& gate_weight, double eps);
 #endif
@@ -298,130 +298,131 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
    asm volatile("griddepcontrol.launch_dependents;");
 #endif
    return;
-  }
-
+  } else {
 #pragma unroll
-  for (int ii = 0; ii < VPT; ++ii) {
-    float val = row_chunk[ii];
-    float val_b = val * beta;
-    // Compute softplus: log(1 + exp(val)) with numerical stability
-    // When val > threshold, softplus(x) ≈ x to avoid exp overflow
-    val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
-    val = sqrtf(val);
-    if (correction_bias) {
-      const int group_id = ii / ELTS_PER_LDG;
-      const int local_id = ii % ELTS_PER_LDG;
-      const int expert_idx = first_elt_read_by_thread +
-                             group_id * THREADS_PER_ROW * ELTS_PER_LDG +
-                             local_id;
-      val = val + correction_bias[expert_idx];
+    for (int ii = 0; ii < VPT; ++ii) {
+      float val = row_chunk[ii];
+      float val_b = val * beta;
+      // Compute softplus: log(1 + exp(val)) with numerical stability
+      // When val > threshold, softplus(x) ≈ x to avoid exp overflow
+      val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
+      val = sqrtf(val);
+      if (correction_bias) {
+        const int group_id = ii / ELTS_PER_LDG;
+        const int local_id = ii % ELTS_PER_LDG;
+        const int expert_idx = first_elt_read_by_thread +
+                               group_id * THREADS_PER_ROW * ELTS_PER_LDG +
+                               local_id;
+        val = val + correction_bias[expert_idx];
+      }
+      row_chunk[ii] = val;
    }
-    row_chunk[ii] = val;
-  }

-  // Original TopK path: find top-k experts by score
-  // Now, sigmoid_res contains the sigmoid of the row chunk. Now, I want to find
-  // the topk elements in each row, along with the max index.
-  int start_col = first_elt_read_by_thread;
-  static constexpr int COLS_PER_GROUP_LDG = ELTS_PER_LDG * THREADS_PER_ROW;
+    // Original TopK path: find top-k experts by score
+    // Now, sigmoid_res contains the sigmoid of the row chunk. Now, I want to
+    // find the topk elements in each row, along with the max index.
+    int start_col = first_elt_read_by_thread;
+    static constexpr int COLS_PER_GROUP_LDG = ELTS_PER_LDG * THREADS_PER_ROW;

-  float selected_sum = 0.f;
-  for (int k_idx = 0; k_idx < k; ++k_idx) {
-    // First, each thread does the local argmax
-    float max_val = row_chunk[0];
-    int expert = start_col;
+    float selected_sum = 0.f;
+    for (int k_idx = 0; k_idx < k; ++k_idx) {
+      // First, each thread does the local argmax
+      float max_val = row_chunk[0];
+      int expert = start_col;
 #pragma unroll
-    for (int ldg = 0, col = start_col; ldg < LDG_PER_THREAD;
-         ++ldg, col += COLS_PER_GROUP_LDG) {
+      for (int ldg = 0, col = start_col; ldg < LDG_PER_THREAD;
+           ++ldg, col += COLS_PER_GROUP_LDG) {
 #pragma unroll
-      for (int ii = 0; ii < ELTS_PER_LDG; ++ii) {
-        float val = row_chunk[ldg * ELTS_PER_LDG + ii];
+        for (int ii = 0; ii < ELTS_PER_LDG; ++ii) {
+          float val = row_chunk[ldg * ELTS_PER_LDG + ii];

-        // No check on the experts here since columns with the smallest index
-        // are processed first and only updated if > (not >=)
-        if (val > max_val) {
-          max_val = val;
-          expert = col + ii;
+          // No check on the experts here since columns with the smallest index
+          // are processed first and only updated if > (not >=)
+          if (val > max_val) {
+            max_val = val;
+            expert = col + ii;
+          }
        }
      }
-    }

 // Now, we perform the argmax reduce. We use the butterfly pattern so threads
 // reach consensus about the max. This will be useful for K > 1 so that the
 // threads can agree on "who" had the max value. That thread can then blank out
 // their max with -inf and the warp can run more iterations...
 #pragma unroll
-    for (int mask = THREADS_PER_ROW / 2; mask > 0; mask /= 2) {
-      float other_max =
-          VLLM_SHFL_XOR_SYNC_WIDTH(max_val, mask, THREADS_PER_ROW);
-      int other_expert =
-          VLLM_SHFL_XOR_SYNC_WIDTH(expert, mask, THREADS_PER_ROW);
+      for (int mask = THREADS_PER_ROW / 2; mask > 0; mask /= 2) {
+        float other_max =
+            VLLM_SHFL_XOR_SYNC_WIDTH(max_val, mask, THREADS_PER_ROW);
+        int other_expert =
+            VLLM_SHFL_XOR_SYNC_WIDTH(expert, mask, THREADS_PER_ROW);

-      // We want lower indices to "win" in every thread so we break ties this
-      // way
-      if (other_max > max_val ||
-          (other_max == max_val && other_expert < expert)) {
-        max_val = other_max;
-        expert = other_expert;
+        // We want lower indices to "win" in every thread so we break ties this
+        // way
+        if (other_max > max_val ||
+            (other_max == max_val && other_expert < expert)) {
+          max_val = other_max;
+          expert = other_expert;
+        }
+      }
+
+      // Write the max for this k iteration to global memory.
+      if (thread_group_idx == 0) {
+        // Add a guard to ignore experts not included by this node
+        const bool node_uses_expert =
+            expert >= start_expert && expert < end_expert;
+        const bool should_process_row = row_is_active && node_uses_expert;
+
+        // The lead thread from each sub-group will write out the final results
+        // to global memory. (This will be a single) thread per row of the
+        // input/output matrices.
+        const int idx = k * thread_row + k_idx;
+        if (correction_bias != nullptr) {
+          max_val -= correction_bias[expert];
+        }
+        output[idx] = max_val;
+        indices[idx] =
+            should_process_row ? (expert - start_expert) : NUM_EXPERTS;
+        source_rows[idx] = k_idx * num_rows + thread_row;
+        if (renormalize) {
+          selected_sum += max_val;
+        }
+      }
+
+      // Finally, we clear the value in the thread with the current max if there
+      // is another iteration to run.
+      if (k_idx + 1 < k) {
+        const int ldg_group_for_expert = expert / COLS_PER_GROUP_LDG;
+        const int thread_to_clear_in_group =
+            (expert / ELTS_PER_LDG) % THREADS_PER_ROW;
+
+        // Only the thread in the group which produced the max will reset the
+        // "winning" value to -inf.
+        if (thread_group_idx == thread_to_clear_in_group) {
+          const int offset_for_expert = expert % ELTS_PER_LDG;
+          // Safe to set to any negative value since row_chunk values must be
+          // between 0 and 1.
+          row_chunk[ldg_group_for_expert * ELTS_PER_LDG + offset_for_expert] =
+              -10000.f;
+        }
      }
    }

-    // Write the max for this k iteration to global memory.
+    // Apply renormalization and routed scaling factor to final weights.
    if (thread_group_idx == 0) {
-      // Add a guard to ignore experts not included by this node
-      const bool node_uses_expert =
-          expert >= start_expert && expert < end_expert;
-      const bool should_process_row = row_is_active && node_uses_expert;
-
-      // The lead thread from each sub-group will write out the final results to
-      // global memory. (This will be a single) thread per row of the
-      // input/output matrices.
-      const int idx = k * thread_row + k_idx;
-      if (correction_bias != nullptr) {
-        max_val -= correction_bias[expert];
-      }
-      output[idx] = max_val;
-      indices[idx] = should_process_row ? (expert - start_expert) : NUM_EXPERTS;
-      source_rows[idx] = k_idx * num_rows + thread_row;
+      float scale = static_cast<float>(routed_scaling_factor);
      if (renormalize) {
-        selected_sum += max_val;
+        const float denom = selected_sum > 0.f ? selected_sum : 1.f;
+        scale /= denom;
+      }
+      for (int k_idx = 0; k_idx < k; ++k_idx) {
+        const int idx = k * thread_row + k_idx;
+        output[idx] = output[idx] * scale;
      }
    }
-
-    // Finally, we clear the value in the thread with the current max if there
-    // is another iteration to run.
-    if (k_idx + 1 < k) {
-      const int ldg_group_for_expert = expert / COLS_PER_GROUP_LDG;
-      const int thread_to_clear_in_group =
-          (expert / ELTS_PER_LDG) % THREADS_PER_ROW;
-
-      // Only the thread in the group which produced the max will reset the
-      // "winning" value to -inf.
-      if (thread_group_idx == thread_to_clear_in_group) {
-        const int offset_for_expert = expert % ELTS_PER_LDG;
-        // Safe to set to any negative value since row_chunk values must be
-        // between 0 and 1.
-        row_chunk[ldg_group_for_expert * ELTS_PER_LDG + offset_for_expert] =
-            -10000.f;
-      }
-    }
-  }
-
-  // Apply renormalization and routed scaling factor to final weights.
-  if (thread_group_idx == 0) {
-    float scale = static_cast<float>(routed_scaling_factor);
-    if (renormalize) {
-      const float denom = selected_sum > 0.f ? selected_sum : 1.f;
-      scale /= denom;
-    }
-    for (int k_idx = 0; k_idx < k; ++k_idx) {
-      const int idx = k * thread_row + k_idx;
-      output[idx] = output[idx] * scale;
-    }
-  }
 #if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
-  asm volatile("griddepcontrol.launch_dependents;");
+    asm volatile("griddepcontrol.launch_dependents;");
 #endif
+  }
 }

 namespace detail {
@@ -135,6 +135,12 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, m) {
  // DeepSeek V3 optimized router GEMM for SM90+
  m.def("dsv3_router_gemm(Tensor! output, Tensor mat_a, Tensor mat_b) -> ()");
  // conditionally compiled so impl registration is in source file
+
+  // DeepSeek V4 fused RMSNorm + router GEMV for SM90+
+  m.def(
+      "dsv4_norm_router_gemm(Tensor! logits, Tensor! normed_x, Tensor x, "
+      "Tensor norm_weight, Tensor gate_weight, float eps) -> ()");
+  // conditionally compiled so impl registration is in source file
 #endif
 }

@@ -845,9 +845,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
                libcublas-dev-${CUDA_VERSION_DASH} \
                libcusolver-dev-${CUDA_VERSION_DASH}"; \
    if [ "$INSTALL_KV_CONNECTORS" = "true" ]; then \
-        if [ "$CUDA_MAJOR" -ge 13 ]; then \
-            uv pip install --system nixl-cu13; \
-        fi; \
        uv pip install --system -r /tmp/kv_connectors.txt --no-build || ( \
            # if the above fails, install from source
            apt-get update -y && \
@@ -857,6 +854,8 @@ RUN --mount=type=cache,target=/root/.cache/uv \
            # clean up -dev packages, keep runtime libraries
            rm -rf /var/lib/apt/lists/* \
        ); \
+        # Force-reinstall the matching CUDA wheel so the correct nixl_ep_cpp.so is installed.
+        uv pip install --system --force-reinstall --no-deps nixl-cu${CUDA_MAJOR}; \
    fi

 ENV VLLM_USAGE_SOURCE production-docker-image
@@ -2,10 +2,18 @@
 ARG REMOTE_VLLM="0"
 ARG COMMON_WORKDIR=/app
 ARG BASE_IMAGE=rocm/vllm-dev:base
-# AMD NIC backend
-ARG NIC_BACKEND=none
+# NIC backend for MoRI RDMA support.
+# By default (all), drivers and userspace libraries for all supported NIC types
+# (ainic and bnxt) are installed; MoRI selects the appropriate one at runtime.
+# To install drivers for a single NIC type only, set NIC_BACKEND explicitly:
+#   --build-arg NIC_BACKEND=ainic   # AMD AINIC (Pensando) only
+#   --build-arg NIC_BACKEND=bnxt    # Broadcom Thor-2 only
+#   --build-arg NIC_BACKEND=none    # Install nothing.
+ARG NIC_BACKEND=all
 # AMD AINIC apt repo settings
-ARG AINIC_VERSION=1.117.5
+# Users can specify a custom version compatible with their host drivers.
+# The default version has been tested with ioinic-dkms=25.11.1.001
+ARG AINIC_VERSION=1.117.3-hydra
 ARG UBUNTU_CODENAME=jammy

 # Sccache configuration (only used in release pipeline)
@@ -231,39 +239,52 @@ ARG NIC_BACKEND
 ARG AINIC_VERSION
 ARG UBUNTU_CODENAME
 RUN /bin/bash -lc 'set -euo pipefail; \
+ \
+ install_ainic() { \
+   apt-get update && apt-get install -y --no-install-recommends ca-certificates curl gnupg apt-transport-https; \
+   rm -rf /var/lib/apt/lists/*; \
+   mkdir -p /etc/apt/keyrings; \
+   curl -fsSL https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor > /etc/apt/keyrings/amdainic.gpg; \
+   echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/amdainic.gpg] https://repo.radeon.com/amdainic/pensando/ubuntu/${AINIC_VERSION} ${UBUNTU_CODENAME} main" \
+     > /etc/apt/sources.list.d/amdainic.list; \
+   apt-get update && apt-get install -y --no-install-recommends \
+     libionic-dev \
+     ionic-common \
+   ; \
+   rm -rf /var/lib/apt/lists/*; \
+ }; \
+ \
+ # NOTE: requires FW 235.2.86.0 and kernel drivers on the host: \
+ #   bnxt-en-dkms=1.10.3.235.2.86.0 bnxt-re-dkms=235.2.86.0 (from packages.broadcom.com PPA) \
+ install_bnxt() { \
+   install -m 0755 -d /etc/apt/keyrings; \
+   curl -fsSL https://packages.broadcom.com/artifactory/api/security/keypair/PackagesKey/public \
+     -o /etc/apt/keyrings/broadcom-nic.asc; \
+   chmod a+r /etc/apt/keyrings/broadcom-nic.asc; \
+   echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/broadcom-nic.asc] https://packages.broadcom.com/artifactory/ethernet-nic-debian-public jammy main" \
+     > /etc/apt/sources.list.d/broadcom-nic.list; \
+   apt-get update && apt-get install -y --no-install-recommends \
+     bnxt-rocelib=235.2.86.0 \
+   ; \
+   cp -a /usr/local/lib/x86_64-linux-gnu/libbnxt_re* /usr/local/lib/; \
+   ldconfig; \
+   rm -rf /var/lib/apt/lists/*; \
+ }; \
+ \
 echo "[MORI] Install MoRI proxy deps"; \
 pip install --quiet --ignore-installed blinker && \
 pip install --quiet quart msgpack aiohttp pyzmq; \
 echo "[MORI] NIC_BACKEND=${NIC_BACKEND}"; \
-  \
-  # NIC backend deps — mori auto-detects NIC at runtime (MORI_DEVICE_NIC env var override).
-  # Only vendor packages are installed here for dlopen (e.g. libionic.so); no compile-time flags needed.
-  case "${NIC_BACKEND}" in \
-    # default: mlx5
-    none) \
-      ;; \
-    # AMD NIC
-    ainic) \
-      apt-get update && apt-get install -y --no-install-recommends ca-certificates curl gnupg apt-transport-https && \
-      rm -rf /var/lib/apt/lists/* && mkdir -p /etc/apt/keyrings; \
-      curl -fsSL https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor > /etc/apt/keyrings/amdainic.gpg; \
-      echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/amdainic.gpg] https://repo.radeon.com/amdainic/pensando/ubuntu/${AINIC_VERSION} ${UBUNTU_CODENAME} main" \
-        > /etc/apt/sources.list.d/amdainic.list; \
-      apt-get update && apt-get install -y --no-install-recommends \
-          libionic-dev \
-          ionic-common \
-      ; \
-      rm -rf /var/lib/apt/lists/*; \
-      ;; \
-    # TODO: Add Broadcom bnxt packages/repos here later.
-    # bnxt) \
-    #   echo "[MORI] Add Broadcom bnxt packages/repos here later."; \
-    #   ;; \
-    *) \
-      echo "ERROR: unknown NIC_BACKEND=${NIC_BACKEND}. Use one of: none, ainic"; \
-      exit 2; \
-      ;; \
-  esac;'
+ \
+ # NIC backend deps — mori auto-detects NIC at runtime (MORI_DEVICE_NIC env var override). \
+ # Only vendor packages are installed here for dlopen; no compile-time flags needed. \
+ case "${NIC_BACKEND}" in \
+   none)  ;; \
+   all)   install_ainic; install_bnxt ;; \
+   ainic) install_ainic ;; \
+   bnxt)  install_bnxt ;; \
+   *)     echo "ERROR: unknown NIC_BACKEND=${NIC_BACKEND}. Use one of: none, ainic, bnxt, all"; exit 2 ;; \
+ esac'

 # -----------------------
 # vLLM wheel release build stage (for building distributable wheels)
@@ -125,12 +125,13 @@ Priority is **1 = highest** (tried first).
 | Priority | Backend |
 | -------- | ------- |
 | 1 | `FLASHINFER_MLA` |
-| 2 | `CUTLASS_MLA` |
-| 3 | `FLASH_ATTN_MLA` |
-| 4 | `FLASHMLA` |
-| 5 | `TRITON_MLA` |
-| 6 | `FLASHINFER_MLA_SPARSE`**\*** |
-| 7 | `FLASHMLA_SPARSE` |
+| 2 | `TOKENSPEED_MLA` |
+| 3 | `CUTLASS_MLA` |
+| 4 | `FLASH_ATTN_MLA` |
+| 5 | `FLASHMLA` |
+| 6 | `TRITON_MLA` |
+| 7 | `FLASHINFER_MLA_SPARSE`**\*** |
+| 8 | `FLASHMLA_SPARSE` |

 **Ampere/Hopper (SM 8.x-9.x):**

@@ -202,6 +203,7 @@ hardware and configuration.
 | `FLASH_ATTN`‡ | FlashAttention varlen (FA2/FA3/FA4) | fp16, bf16 | Any | FA4 on SM100+, FA3 on SM90, FA2 otherwise |
 | `TRTLLM_RAGGED` | TensorRT-LLM ragged attention | fp16, bf16 | 10.x | DeepSeek R1 dims only |
 | `FLASHINFER` | FlashInfer CUTLASS backend | fp16, bf16 | 10.x | DeepSeek R1 dims only |
+| `TOKENSPEED_MLA` | | fp16, bf16 | 10.x | DeepSeek R1 dims only |

 > **‡** TRT-LLM Ragged is the default on Blackwell (SM100).
 > On other GPUs, FlashAttention is used as the default.
@@ -222,5 +224,6 @@ MLA decode backends are selected using the standard
 | `ROCM_AITER_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %1 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | N/A |
 | `ROCM_AITER_MLA_SPARSE` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 1, 64 | Any | ❌ | ❌ | ✅ | ❌ | ❌ | Decoder | N/A |
 | `ROCM_AITER_TRITON_MLA` | fp16, bf16 | `auto` | Any | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | N/A |
+| `TOKENSPEED_MLA` | fp16, bf16 | `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | 10.x |
 | `TRITON_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | %16 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | Any |
 | `XPU_MLA_SPARSE` | fp16, bf16 | `auto`, `float16`, `bfloat16` | Any | 576 | ❌ | ❌ | ✅ | ❌ | ❌ | Decoder | Any |
@@ -85,13 +85,14 @@ Models opt-in to encoder CUDA Graphs by implementing the [SupportsEncoderCudaGra

 | Architecture | Models | CG for Image | CG for Video |
 | ------------ | ------ | ------------ | ------------ |
+| `Qwen2VLForConditionalGeneration` | `Qwen2-VL` | ✅︎ | ✅︎ |
 | `Qwen2_5_VLForConditionalGeneration` | `Qwen2.5-VL` | ✅︎ | ✅︎ |
 | `Qwen3VLForConditionalGeneration` | `Qwen3-VL` | ✅︎ | ✅︎ |
 | `Qwen3_5ForConditionalGeneration` | `Qwen3.5` | ✅︎ | ✅︎ |

 !!! note
    Encoder CUDA Graphs have currently been tested with `--mm-encoder-attn-backend=FLASH_ATTN` and `--mm-encoder-attn-backend=FLASHINFER` on Blackwell GPUs.
-    For Qwen2.5-VL only FA2 and FA3 has been tested.
+    For Qwen2-VL and Qwen2.5-VL only FA2 and FA3 has been tested.

 ## Configuration

@@ -84,7 +84,7 @@ Since simple RTN does not require data for weight quantization and the activatio
 Install `vllm` and `lm-evaluation-harness` for evaluation:

 ```bash
-pip install vllm "lm-eval[api]>=0.4.11"
+pip install vllm "lm-eval[api]>=0.4.12"
 ```

 Load and run the model in `vllm`:
@@ -18,7 +18,7 @@ pip install llmcompressor
 Additionally, install `vllm` and `lm-evaluation-harness` for evaluation:

 ```bash
-pip install vllm "lm-eval[api]>=0.4.11"
+pip install vllm "lm-eval[api]>=0.4.12"
 ```

 ## Quantization Process
@@ -23,7 +23,7 @@ pip install llmcompressor
 Additionally, install `vllm` and `lm-evaluation-harness` for evaluation:

 ```bash
-pip install vllm "lm-eval[api]>=0.4.11"
+pip install vllm "lm-eval[api]>=0.4.12"
 ```

 ## Quantization Process
@@ -42,18 +42,59 @@ vllm serve meta-llama/Llama-3.1-8B --quantization mxfp8

 For fine-grained control, use a `quantization_config` dictionary.

+### Schema
+
+```yaml
+quantization_config:
+  linear:
+    weight: <name>      # see QUANT_KEY_NAMES in vllm/config/quantization.py
+    activation: <name>
+  moe:
+    weight: <name>
+    activation: <name>
+  ignore: [<layer-name-or-regex>, ...]
+```
+
+`linear` and `moe` accept a full `{weight, activation}` dict, or a bare
+string. A string resolves first against the `--quantization` shorthands
+(taking the matching layer-kind slot), then against `QUANT_KEY_NAMES` as a
+weight name. Unset fields fall back to the `--quantization` shorthand's
+defaults, or for already-quantized checkpoints to whatever the checkpoint
+declares.
+
+The CLI accepts the same shape as JSON or as dotted keys:
+
+```bash
+vllm serve <model> --quantization-config '{"moe":{"activation":"mxfp8"}}'
+vllm serve <model> --quantization-config.moe.activation mxfp8
+```
+
+### Activation overrides on already-quantized checkpoints
+
+For checkpoint-quantized models, `quantization_config` lets you pick an
+activation format independently of the baked-in weights. The supported
+overrides are checkpoint-specific; today this is wired up for MXFP4 MoE
+checkpoints (gpt-oss) where you can opt into FP8 activations:
+
+```bash
+vllm serve openai/gpt-oss-20b --quantization-config.moe.activation mxfp8
+```
+
+Combine with `--moe-backend` to pin a specific kernel family.
+
 ### Separate Schemes for Dense and MoE Layers

-You can apply different quantization schemes to dense linear layers and MoE expert layers:
+You can apply different quantization schemes to dense linear layers and MoE expert layers via the `linear` and `moe` fields. Each accepts either a full spec dict, or a bare string naming an online shorthand (e.g. `"fp8_per_block"`) or weight format (e.g. `"fp8_per_block_static"`); fields not set fall back to the shorthand defaults.

 ```python
 from vllm import LLM

+# Linear: per-block FP8; MoE: per-tensor FP8 (inherited from the shorthand)
 llm = LLM(
    "ibm-granite/granite-3.0-1b-a400m-base",
    quantization="fp8_per_tensor",
    quantization_config={
-        "linear_scheme_override": "fp8_per_block",
+        "linear": "fp8_per_block",
    },
 )
 ```
@@ -63,11 +104,12 @@ Or,
 ```python
 from vllm import LLM

+# Linear: per-tensor FP8 (inherited); MoE: per-block FP8
 llm = LLM(
    "ibm-granite/granite-3.0-1b-a400m-base",
    quantization="fp8_per_tensor",
    quantization_config={
-        "moe_scheme_override": "fp8_per_block",
+        "moe": "fp8_per_block",
    },
 )
 ```
@@ -20,7 +20,7 @@ for more installation details.
 Additionally, install `vllm` and `lm-evaluation-harness` for evaluation:

 ```bash
-pip install vllm "lm-eval[api]>=0.4.11"
+pip install vllm "lm-eval[api]>=0.4.12"
 ```

 ## Quantization Process
@@ -15,6 +15,7 @@ vLLM supports a variety of methods of speculative decoding. Model-based methods
 - [Multi-Layer Perceptron](mlp.md)
 - [N-Gram](n_gram.md)
 - [Suffix Decoding](suffix.md)
+- [Custom Proposer Backend (Experimental)](#custom-proposer-backend-experimental)

 ## Method Selection at a Glance

@@ -30,11 +31,22 @@ depend on your model family, traffic pattern, hardware, and sampling settings.
 | MLP speculator | Medium to high gain | Medium gain | Good when compatible MLP speculators are available. |
 | N-gram | Low to medium gain | Medium gain | Lightweight and easy to enable. |
 | Suffix decoding | Low to medium gain | Medium gain | No extra draft model; dynamic speculation depth. |
+| Custom Proposer | Varies | Varies | Bring your own proposer class (experimental). |

 For reproducible measurements in your environment, use
 [`examples/features/speculative_decoding/spec_decode_offline.py`](../../../examples/features/speculative_decoding/spec_decode_offline.py)
 or the [benchmark CLI guide](../../benchmarking/cli.md).

+## Custom Proposer Backend (Experimental)
+
+You can plug in your own custom proposer class for speculative decoding by setting the method to `custom_class` and providing the full module path to your class.
+Your custom class must accept a `VllmConfig` upon instantiation and implement a `propose` method.
+
+**Example configuration:**
+
+- `speculative_config.method = "custom_class"`
+- `speculative_config.model = "your_module.YourCustomProposerClass"`
+
 ## `--speculative-config` schema

 Use `--speculative-config` to pass speculative decoding settings as a JSON
@@ -11,24 +11,50 @@

 ### Text-only Language Models

-| Model                                | Architecture                             | Supported |
+| Model | Architecture | Supported |
 | ------------------------------------ | ---------------------------------------- | --------- |
-| meta-llama/Llama-3.1-8B-Instruct     | LlamaForCausalLM                         | ✅        |
-| meta-llama/Llama-3.2-3B-Instruct     | LlamaForCausalLM                         | ✅        |
-| ibm-granite/granite-3.2-2b-instruct  | GraniteForCausalLM                       | ✅        |
-| Qwen/Qwen3-1.7B                      | Qwen3ForCausalLM                         | ✅        |
-| Qwen/Qwen3-4B                        | Qwen3ForCausalLM                         | ✅        |
-| Qwen/Qwen3-8B                        | Qwen3ForCausalLM                         | ✅        |
-| zai-org/glm-4-9b-hf                  | GLMForCausalLM                           | ✅        |
-| google/gemma-7b                      | GemmaForCausalLM                         | ✅        |
+| unsloth/gpt-oss-20b | GptOssForCausalLM | ✅ |
+| meta-llama/Llama-3.1-8B-Instruct | LlamaForCausalLM | ✅ |
+| meta-llama/Llama-3.2-1B | LlamaForCausalLM | ✅ |
+| meta-llama/Llama-3.2-3B-Instruct | LlamaForCausalLM | ✅ |
+| meta-llama/Llama-3.3-70B-Instruct | LlamaForCausalLM | ✅ |
+| RedHatAI/Meta-Llama-3.1-8B-quantized.w8a8 | LlamaForCausalLM | ✅ |
+| RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8 | LlamaForCausalLM | ✅ |
+| RedHatAI/Llama-3.2-1B-Instruct-quantized.w8a8 | LlamaForCausalLM | ✅ |
+| RedHatAI/Llama-3.2-3B-Instruct-quantized.w8a8 | LlamaForCausalLM | ✅ |
+| RedHatAI/DeepSeek-R1-Distill-Llama-70B-quantized.w8a8 | LlamaForCausalLM | ✅ |
+| hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4 | LlamaForCausalLM | ✅ |
+| AMead10/Llama-3.2-1B-Instruct-AWQ | LlamaForCausalLM | ✅ |
+| AMead10/Llama-3.2-3B-Instruct-AWQ | LlamaForCausalLM | ✅ |
+| TheBloke/TinyLlama-1.1B-Chat-v1.0-AWQ | LlamaForCausalLM | ✅ |
+| TheBloke/TinyLlama-1.1B-Chat-v1.0-GPTQ | LlamaForCausalLM | ✅ |
+| ibm-granite/granite-3.2-2b-instruct | GraniteForCausalLM | ✅ |
+| Qwen/Qwen3-1.7B | Qwen3ForCausalLM | ✅ |
+| Qwen/Qwen3-4B | Qwen3ForCausalLM | ✅ |
+| Qwen/Qwen3-8B | Qwen3ForCausalLM | ✅ |
+| Qwen/Qwen3-14B | Qwen3ForCausalLM | ✅ |
+| Qwen/Qwen3-14B-AWQ | Qwen3ForCausalLM | ✅ |
+| Qwen/Qwen3-30B-A3B | Qwen3MoeForCausalLM | ✅ |
+| Qwen/QwQ-32B-AWQ | Qwen2ForCausalLM | ✅ |
+| Qwen/Qwen1.5-0.5B-Chat-GPTQ-Int4 | Qwen2ForCausalLM | ✅ |
+| RedHatAI/QwQ-32B-quantized.w8a8 | Qwen2ForCausalLM | ✅ |
+| zai-org/glm-4-9b-hf | GLMForCausalLM | ✅ |
+| google/gemma-7b | GemmaForCausalLM | ✅ |
+| microsoft/Phi-4-reasoning | Phi3ForCausalLM | ✅ |
+| TheBloke/Mistral-7B-Instruct-v0.2-AWQ | MistralForCausalLM | ✅ |

 ### Multimodal Language Models

-| Model                                | Architecture                             | Supported |
+| Model | Architecture | Supported |
 | ------------------------------------ | ---------------------------------------- | --------- |
-| Qwen/Qwen2.5-VL-7B-Instruct          | Qwen2VLForConditionalGeneration          | ✅        |
-| openai/whisper-large-v3              | WhisperForConditionalGeneration          | ✅        |
+| meta-llama/Llama-4-Scout-17B-16E-Instruct | Llama4ForConditionalGeneration | ✅ |
+| google/gemma-3-4b-it | Gemma3ForConditionalGeneration | ✅ |
+| google/gemma-3-12b-it | Gemma3ForConditionalGeneration | ✅ |
+| google/gemma-4-E4B-it | Gemma4ForConditionalGeneration | ✅ |
+| google/gemma-4-E2B-it | Gemma4ForConditionalGeneration | ✅ |
+| google/gemma-4-26B-A4B-it | Gemma4ForConditionalGeneration | ✅ |
+| microsoft/Phi-4-multimodal-instruct | Phi4MMForCausalLM | ✅ |
+| Qwen/Qwen2.5-VL-7B-Instruct | Qwen2VLForConditionalGeneration | ✅ |
+| openai/whisper-large-v3 | WhisperForConditionalGeneration | ✅ |

-✅ Runs and optimized.  
-🟨 Runs and correct but not optimized to green yet.  
-❌ Does not pass accuracy test or does not run.  
+✅ Runs and optimized.
@@ -131,24 +131,24 @@ enabling the corresponding APIs.

 ### `LLM.classify`

-The [classify][vllm.LLM.classify] method outputs a probability vector for each prompt.
+The [classify][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.classify] method outputs a probability vector for each prompt.
 It is primarily designed for [classification models](classify.md).
 For more information about `LLM.embed`, see [this page](classify.md#offline-inference).

 ### `LLM.embed`

-The [embed][vllm.LLM.embed] method outputs an embedding vector for each prompt.
+The [embed][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.embed] method outputs an embedding vector for each prompt.
 It is primarily designed for [embedding models](embed.md).
 For more information about `LLM.embed`, see [this page](embed.md#offline-inference).

 ### `LLM.score`

-The [score][vllm.LLM.score] method outputs similarity scores between sentence pairs.
+The [score][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.score] method outputs similarity scores between sentence pairs.
 It is primarily designed for [score models](scoring.md).

 ### `LLM.encode`

-The [encode][vllm.LLM.encode] method is available to all pooling models in vLLM.
+The [encode][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.encode] method is available to all pooling models in vLLM.

 Please use one of the more specific methods or set the task directly when using `LLM.encode`, refer to the [table above](#offline-apis-corresponding-to-pooling-usages).

@@ -77,7 +77,7 @@ The following [pooling parameters][vllm.PoolingParams] are supported.

 ### `LLM.classify`

-The [classify][vllm.LLM.classify] method outputs a probability vector for each prompt.
+The [classify][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.classify] method outputs a probability vector for each prompt.

 ```python
 from vllm import LLM
@@ -93,7 +93,7 @@ A code example can be found here: [examples/basic/offline_inference/classify.py]

 ### `LLM.encode`

-The [encode][vllm.LLM.encode] method is available to all pooling models in vLLM.
+The [encode][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.encode] method is available to all pooling models in vLLM.

 Set `pooling_task="classify"` when using `LLM.encode` for classification Models:

@@ -120,7 +120,7 @@ The following [pooling parameters][vllm.PoolingParams] are supported.

 ### `LLM.embed`

-The [embed][vllm.LLM.embed] method outputs an embedding vector for each prompt.
+The [embed][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.embed] method outputs an embedding vector for each prompt.

 ```python
 from vllm import LLM
@@ -136,7 +136,7 @@ A code example can be found here: [examples/basic/offline_inference/embed.py](..

 ### `LLM.encode`

-The [encode][vllm.LLM.encode] method is available to all pooling models in vLLM.
+The [encode][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.encode] method is available to all pooling models in vLLM.

 Set `pooling_task="embed"` when using `LLM.encode` for embedding Models:

@@ -152,7 +152,7 @@ print(f"Data: {data!r}")

 ### `LLM.score`

-The [score][vllm.LLM.score] method outputs similarity scores between sentence pairs.
+The [score][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.score] method outputs similarity scores between sentence pairs.

 All models that support embedding task also support using the score API to compute similarity scores by calculating the cosine similarity of two input prompt's embeddings.

@@ -87,7 +87,7 @@ The following [pooling parameters][vllm.PoolingParams] are supported.

 ### `LLM.encode`

-The [encode][vllm.LLM.encode] method is available to all pooling models in vLLM.
+The [encode][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.encode] method is available to all pooling models in vLLM.

 - Reward Models

@@ -130,7 +130,7 @@ The following [pooling parameters][vllm.PoolingParams] are only supported by cro

 ### `LLM.score`

-The [score][vllm.LLM.score] method outputs similarity scores between sentence pairs.
+The [score][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.score] method outputs similarity scores between sentence pairs.

 ```python
 from vllm import LLM
@@ -87,7 +87,7 @@ The following [pooling parameters][vllm.PoolingParams] are supported.

 ### `LLM.encode`

-The [encode][vllm.LLM.encode] method is available to all pooling models in vLLM.
+The [encode][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.encode] method is available to all pooling models in vLLM.

 Set `pooling_task="token_classify"` when using `LLM.encode` for token classification Models:

@@ -94,7 +94,7 @@ The following [pooling parameters][vllm.PoolingParams] are supported.

 ### `LLM.encode`

-The [encode][vllm.LLM.encode] method is available to all pooling models in vLLM.
+The [encode][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.encode] method is available to all pooling models in vLLM.

 Set `pooling_task="token_embed"` when using `LLM.encode` for token embedding Models:

@@ -110,7 +110,7 @@ print(f"Data: {data!r}")

 ### `LLM.score`

-The [score][vllm.LLM.score] method outputs similarity scores between sentence pairs.
+The [score][vllm.entrypoints.pooling.offline.PoolingOfflineMixin.score] method outputs similarity scores between sentence pairs.

 All models that support token embedding task also support using the score API to compute similarity scores by calculating the late interaction of two input prompts.

@@ -575,6 +575,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `IsaacForConditionalGeneration` | Isaac | T + I<sup>+</sup> | `PerceptronAI/Isaac-0.1` | ✅︎ | ✅︎ |
 | `InternS1ForConditionalGeneration` | Intern-S1 | T + I<sup>E+</sup> + V<sup>E+</sup> | `internlm/Intern-S1`, `internlm/Intern-S1-mini`, etc. | ✅︎ | ✅︎ |
 | `InternS1ProForConditionalGeneration` | Intern-S1-Pro | T + I<sup>E+</sup> + V<sup>E+</sup> | `internlm/Intern-S1-Pro`, etc. | ✅︎ | ✅︎ |
+| `InternS2PreviewForConditionalGeneration` | Intern-S2-Preview | T + I<sup>E+</sup> + V<sup>E+</sup> | `internlm/Intern-S2-Preview`, etc. | ✅︎ | ✅︎ |
 | `InternVLChatModel` | InternVL 3.5, InternVL 3.0, InternVideo 2.5, InternVL 2.5, Mono-InternVL, InternVL 2.0 | T + I<sup>E+</sup> + (V<sup>E+</sup>) | `OpenGVLab/InternVL3_5-14B`, `OpenGVLab/InternVL3-9B`, `OpenGVLab/InternVideo2_5_Chat_8B`, `OpenGVLab/InternVL2_5-4B`, `OpenGVLab/Mono-InternVL-2B`, `OpenGVLab/InternVL2-4B`, etc. | ✅︎ | ✅︎ |
 | `InternVLForConditionalGeneration` | InternVL 3.0 (HF format) | T + I<sup>E+</sup> + V<sup>E+</sup> | `OpenGVLab/InternVL3-1B-hf`, etc. | ✅︎ | ✅︎ |
 | `KananaVForConditionalGeneration` | Kanana-V | T + I<sup>+</sup> | `kakaocorp/kanana-1.5-v-3b-instruct`, etc. | | ✅︎ |
@@ -1,289 +0,0 @@
-# Routed Experts Replay
-
-## Overview
-
-Routed experts replay captures which MoE (Mixture of Experts) experts process each token during inference and returns this information alongside the generated text. This is essential for **reinforcement learning (RL) training pipelines** (such as GRPO and RLHF) where the training step needs to reconstruct expert routing decisions from the inference pass.
-
-When enabled, each API response includes:
-
- **`prompt_routed_experts`**: A `[prompt_len, num_moe_layers, top_k]` array of expert IDs for the prompt tokens (at the response level, shared across completions).
- **`routed_experts`**: A `[gen_len, num_moe_layers, top_k]` array of expert IDs for the generated tokens (per completion).
-
-For example, a model with 40 MoE layers and top-22 routing that processes a 100-token prompt and generates 50 tokens would return:
-
- `prompt_routed_experts`: shape `[100, 40, 22]`
- `routed_experts`: shape `[50, 40, 22]`
-
-Each value is an int16 expert ID in the range `[0, num_experts)`.
-
-## Quickstart
-
-### OpenAI API Server
-
-```bash
-vllm serve <MODEL> \
-    --enable-return-routed-experts \
-    --tensor-parallel-size 4 \
-    --enable-expert-parallel
-```
-
-Then query the `/v1/completions` endpoint as usual. The response includes routing data:
-
-```python
-import requests
-
-resp = requests.post("http://localhost:8000/v1/completions", json={
-    "model": "<MODEL>",
-    "prompt": "Explain quantum computing.",
-    "max_tokens": 64,
-    "temperature": 0.0,
-}).json()
-
-# Generation routing (per completion choice)
-gen_routing = resp["choices"][0]["routed_experts"]  # [gen_len, layers, top_k]
-
-# Prompt routing (shared across all choices)
-prompt_routing = resp["prompt_routed_experts"]      # [prompt_len, layers, top_k]
-
-print(f"Prompt routing shape: [{len(prompt_routing)}, "
-      f"{len(prompt_routing[0])}, {len(prompt_routing[0][0])}]")
-print(f"Gen routing shape: [{len(gen_routing)}, "
-      f"{len(gen_routing[0])}, {len(gen_routing[0][0])}]")
-```
-
-### Python SDK (Offline Inference)
-
-```python
-from vllm import LLM, SamplingParams
-
-llm = LLM(
-    model="<MODEL>",
-    enable_return_routed_experts=True,
-    tensor_parallel_size=4,
-    enable_expert_parallel=True,
-)
-
-outputs = llm.generate(
-    ["Explain quantum computing."],
-    SamplingParams(temperature=0, max_tokens=64),
-)
-
-result = outputs[0]
-
-# Prompt routing: numpy array, shape [prompt_len, num_moe_layers, top_k]
-prompt_routing = result.prompt_routed_experts
-print(f"Prompt routing: {prompt_routing.shape}, dtype={prompt_routing.dtype}")
-
-# Generation routing: numpy array, shape [gen_len, num_moe_layers, top_k]
-gen_routing = result.outputs[0].routed_experts
-print(f"Gen routing: {gen_routing.shape}, dtype={gen_routing.dtype}")
-```
-
-## Output Format
-
-### `CompletionOutput.routed_experts`
-
- **Type**: `numpy.ndarray` (Python SDK) or `list[list[list[int]]]` (JSON API)
- **Shape**: `[gen_len, num_moe_layers, top_k]`
- **Dtype**: `int16`
- **Content**: Expert IDs for **generated tokens only**. `gen_len` matches the number of generated tokens (i.e., `usage.completion_tokens` or fewer).
-
-### `RequestOutput.prompt_routed_experts`
-
- **Type**: `numpy.ndarray` (Python SDK) or `list[list[list[int]]]` (JSON API)
- **Shape**: `[prompt_len, num_moe_layers, top_k]`
- **Dtype**: `int16`
- **Content**: Expert IDs for **prompt tokens only**. `prompt_len` matches `usage.prompt_tokens`. This field lives on the request-level response (not per-choice), because prompt routing is shared across all completions when `n > 1`.
-
-### Why Separate Prompt and Generation Routing?
-
-When a request has multiple completions (`n > 1`), each completion shares the same prompt but produces different generated text. Storing prompt routing once on the `RequestOutput` (rather than duplicating it on every `CompletionOutput`) avoids redundant data. For RL training, the consumer typically needs:
-
-1. The prompt routing (once) to reconstruct the forward pass for the shared prefix.
-2. The per-completion generation routing to reconstruct each completion's forward pass.
-
-## Architecture
-
-### Data Flow
-
-```text
-Forward Pass                    Async D2H Pipeline              Output
-─────────────                   ──────────────────              ──────
-FusedMoE layer                  After forward pass:             On request finish:
-writes topk_ids    ──────►     D2H copy to pinned     ──────►  Extract from host cache
-to device buffer               staging buffer                   Split at prompt_len
-(L, N, K) int16                 (via CUDA stream)                Trim gen to output len
-                                Scatter to per-request           Serialize to API response
-                                host cache (numpy)
-```
-
-### Device Cache
-
-A pre-allocated GPU buffer with layout `(L, N, K)` where:
-
- `L` = number of MoE layers
- `N` = `max_num_batched_tokens`
- `K` = `num_experts_per_tok` (top-k)
-
-The `(L, N, K)` layout ensures that `buffer[layer_id]` gives a contiguous `(N, K)` view per layer. Each `FusedMoE` layer gets a persistent reference to its slice via `module._routing_replay_out = buffer[layer_id]`.
-
-**Dtype**: `int16` — sufficient for expert IDs (max ~512 experts in practice) and half the memory of `int32`.
-
-### Host Cache
-
-Per-request numpy arrays for accumulating routing data across decode steps. Each request gets a lazily allocated `(seq_len, L, K)` int16 buffer that grows as the sequence lengthens. Buffers are freed when a request completes.
-
-### Async D2H Pipeline
-
-After each forward pass, the model runner issues a non-blocking device-to-host copy on a dedicated CUDA stream:
-
-1. **Copy**: `pinned_staging[:, :total_tokens, :].copy_(device_buffer[:, :total_tokens, :])` on a separate stream, recorded with a CUDA event.
-2. **Scatter** (deferred to next step): On the *next* forward pass, synchronize the event (effectively free — an entire forward pass has elapsed) and scatter the staging data into per-request host cache buffers using the token positions.
-
-This design ensures the D2H copy overlaps with the next forward pass, minimizing GPU stall time.
-
-### CUDA Graph Compatibility
-
-CUDA graph compatibility requires two mechanisms:
-
-1. **Persistent tensor attribute**: Each `FusedMoE` layer stores a reference to its buffer slice as `module._routing_replay_out`. Because `torch.compile` captures module attributes by reference, graph replay always writes to the live buffer — not a stale snapshot.
-
-2. **Static marking**: Both the full `(L, N, K)` buffer and each per-layer `(N, K)` view are marked with `cudagraph_mark_tensor_static()`. This prevents CUDA graphs from snapshot/restore behavior that would zero the buffer on replay.
-
-### Multi-Node Support
-
-On multi-node tensor-parallel setups, all TP ranks allocate a device buffer (required for symmetric CUDA graph structure), but only TP rank 0 runs the D2H pipeline and host cache. Routing data flows from the model runner through `ModelRunnerOutput` via Ray DAG to the scheduler — no shared memory or file locks needed.
-
-### Routing Capture Path
-
-For the **non-monolithic (Triton) kernel path** (e.g., BF16 MoE), routing is captured after `select_experts()` in the MoE runner:
-
-```python
-routing_replay_out = getattr(layer, "_routing_replay_out", None)
-topk_weights, topk_ids = self.router.select_experts(...)
-
-if routing_replay_out is not None:
-    routing_replay_out[:topk_ids.shape[0]].copy_(topk_ids.to(torch.int16))
-```
-
-For the **monolithic kernel path** (e.g., FP8/MXFP8 via FlashInfer), `routing_replay_out` is threaded through the `apply_monolithic()` call chain and FlashInfer writes expert IDs directly during routing inside the fused kernel.
-
-### MTP (Multi-Token Prediction) Handling
-
-With MTP speculative decoding, the model captures routing for all tokens including speculative ones that may later be rejected. When a request finishes, the generation routing is trimmed to match the actual number of accepted output tokens:
-
-```python
-num_gen = self.detokenizer.num_output_tokens()
-if gen_routed_experts.shape[0] > num_gen and num_gen > 0:
-    gen_routed_experts = gen_routed_experts[:num_gen]
-```
-
-This ensures the routing array length always matches the token IDs in the response.
-
-## Design Decisions
-
-### Why Replace SharedMemory with Device Cache?
-
-The previous implementation used `multiprocessing.SharedMemory` with `fcntl` file locking to transfer routing data from GPU workers to the scheduler. This approach had fundamental problems:
-
- **Multi-node**: `SharedMemory` is node-local. On multi-node TP setups (required for 400B+ parameter models), the scheduler on node 0 cannot read shared memory from workers on other nodes.
- **Performance**: Synchronous `.cpu().numpy()` D2H transfers block the GPU. File-based locking adds further overhead.
- **CUDA graphs**: The callback-based capture mechanism bakes tensor references at trace time, causing stale data on graph replay.
-
-The device cache approach solves all three: data flows through Ray DAG (works multi-node), D2H is async (non-blocking), and persistent tensor attributes work with CUDA graphs.
-
-### Why `(L, N, K)` Layout Instead of `(N, L, K)`?
-
-FlashInfer's `routing_replay_out` parameter expects a contiguous `(N, K)` tensor per layer. With `(L, N, K)` layout, `buffer[layer_id]` gives a contiguous `(N, K)` view with zero-copy slicing. The previous `(N, L, K)` layout would require non-contiguous indexing or an explicit copy.
-
-### Why int16 Instead of int32?
-
-Expert IDs are small integers (typically 0-255 for models with up to 256 experts). `int16` supports up to 32,767 experts — far more than any current model — while halving GPU memory usage and D2H bandwidth compared to `int32`.
-
-### Why Split Prompt and Generation Routing?
-
-RL training pipelines process prompt and generation routing separately:
-
- Prompt routing reconstructs the shared forward pass for the input.
- Generation routing reconstructs each sampled trajectory.
-
-With `n > 1` completions, all completions share the same prompt routing. Duplicating it per completion would waste memory proportional to `n * prompt_len * L * K`. Instead, `prompt_routed_experts` is stored once on `RequestOutput` and shared.
-
-### Why Async D2H Instead of Synchronous Copy?
-
-A synchronous `.cpu()` call forces the GPU to drain its command queue before the copy can begin, stalling the pipeline. The async approach:
-
-1. Issues the copy on a separate CUDA stream (non-blocking to the main compute stream).
-2. Defers the host-side scatter to the *next* step, by which time the copy has finished.
-
-This means the D2H transfer overlaps entirely with the next forward pass, adding near-zero latency to the critical path.
-
-### Why All TP Ranks Get a Device Buffer?
-
-CUDA graph capture records the exact sequence of kernel calls and their arguments. If only rank 0 had a device buffer, the `FusedMoE` layer would take a different code path on rank 0 vs. other ranks (one writes to a buffer, others don't). This asymmetry causes different CUDA graph structures across ranks, which can lead to NCCL deadlocks during collective operations inside the graph. Giving all ranks a real buffer ensures symmetric graph structure. Only rank 0 does the D2H copy and host cache management.
-
-## Performance
-
-Routing replay adds a small overhead from the device buffer writes and async D2H copies. On tested configurations:
-
- **Throughput overhead** (random data, ISL=1024, OSL=1024): **~2%**
- **Memory overhead** (int16 buffer, 40 layers, 8192 tokens, top-22): **~14 MB per GPU**
- **Accuracy impact** (GSM8K): **Zero** (pass@1 identical with and without routing replay)
-
-The overhead is dominated by the per-layer `.copy_()` during the forward pass. The async D2H pipeline runs entirely in the background.
-
-## Supported Configurations
-
-| Configuration                            | Supported                                                 |
-|------------------------------------------|-----------------------------------------------------------|
-| BF16 Triton MoE (non-monolithic)         | Yes                                                       |
-| FP8/MXFP8 FlashInfer MoE (monolithic)    | Yes (requires FlashInfer with `routing_replay_out`)       |
-| CUDA graphs                              | Yes                                                       |
-| Multi-node tensor parallelism            | Yes                                                       |
-| Data parallelism (DP)                    | Yes                                                       |
-| Expert parallelism (EP)                  | Yes                                                       |
-| Prefix caching                           | Yes (cached positions marked with `-1` sentinel)          |
-| MTP speculative decoding                 | Yes (gen routing trimmed to accepted tokens)              |
-| `n > 1` (multiple completions)           | Yes (prompt routing shared, gen routing per-completion)   |
-
-## Limitations
-
- **Streaming**: Routing data is only available when the request finishes (not streamed incrementally).
- **V1 engine only**: Routing replay is implemented for the vLLM V1 engine.
- **Preempted requests**: When a request is preempted by the scheduler (and later resumed via re-prefill), any routing already accumulated in the worker's host cache for that request is dropped without being emitted. The consumer sees `routed_experts=None` for the resumed request with no other signal. Partial-rollout and async-RL pipelines that rely on routing for preempted requests should either disable preemption (`--no-enable-chunked-prefill` / sufficient KV headroom) or reconstruct routing on the resumed prefill.
- **Async scheduling**: Not supported; rejected at config time. The worker-side stop predicate reads `req_state.output_token_ids[-1]`, which under async scheduling is the placeholder `-1` until `AsyncGPUModelRunnerOutput` resolves the real sampled token, so EOS / stop-token finishes would silently drop routing. Use sync scheduling (the default when `--enable-return-routed-experts` is set, or set explicitly with the appropriate scheduler config).
- **Sequence parallelism / naive DP MoE dispatch**: Not supported on the FusedMoE layer; rejected at bind time. SP shards `topk_ids` along dim 0 across the TP group so each rank only captures `1/sp_size` of the rows; naive DP dispatch all-gathers tokens across DP ranks before routing, so `topk_ids.shape[0]` exceeds the per-rank buffer size. Both raise `NotImplementedError` from `bind_routing_capture_to_model`.
- **Pipeline / prefill-context / decode-context parallelism**: Not yet validated; rejected at config time.
-
-## CLI Reference
-
-| Flag                               | Description                                                            |
-|------------------------------------|------------------------------------------------------------------------|
-| `--enable-return-routed-experts`   | Enable routing replay capture and return expert IDs in API responses.  |
-
-## API Reference
-
-### Completions (`/v1/completions`)
-
-**Response-level field:**
-
-| Field                     | Type                                | Description                                                                 |
-|---------------------------|-------------------------------------|-----------------------------------------------------------------------------|
-| `prompt_routed_experts`   | `list[list[list[int]]]` or `null`   | Expert IDs for prompt tokens. Shape: `[prompt_len, num_moe_layers, top_k]`. |
-
-**Choice-level field:**
-
-| Field              | Type                                | Description                                                                   |
-|--------------------|-------------------------------------|-------------------------------------------------------------------------------|
-| `routed_experts`   | `list[list[list[int]]]` or `null`   | Expert IDs for generated tokens. Shape: `[gen_len, num_moe_layers, top_k]`.   |
-
-### Chat Completions (`/v1/chat/completions`)
-
-Same fields as above on `ChatCompletionResponse` and `ChatCompletionResponseChoice`.
-
-### Python SDK
-
-| Object               | Field                     | Type                     | Description                 |
-|----------------------|---------------------------|--------------------------|-----------------------------|
-| `RequestOutput`      | `prompt_routed_experts`   | `np.ndarray` or `None`   | `[prompt_len, L, K]` i16    |
-| `CompletionOutput`   | `routed_experts`          | `np.ndarray` or `None`   | `[gen_len, L, K]` int16     |
@@ -2557,6 +2557,7 @@ MODELS_SUPPORT_VIT_CUDA_GRAPH = [
    "qwen2_5_vl",
    "qwen3_vl",
    "qwen3_vl_moe",
+    "qwen2_vl",
    "qwen3_5",
    "qwen3_5_moe",
 ]
@@ -1,7 +1,7 @@
 <|im_start|>system
 Judge whether the Document meets the requirements based on the Query and the Instruct provided. Note that the answer can only be "yes" or "no".<|im_end|>
 <|im_start|>user
-<Instruct>: {{ messages | selectattr("role", "eq", "system") | map(attribute="content") | first | default("Given a web search query, retrieve relevant passages that answer the query") }}
+<Instruct>: {{ instruction | default(instruct | default(messages | selectattr("role", "eq", "system") | map(attribute="content") | first | default("Given a web search query, retrieve relevant passages that answer the query", true), true), true) }}
 <Query>: {{ messages | selectattr("role", "eq", "query") | map(attribute="content") | first }}
 <Document>: {{ messages | selectattr("role", "eq", "document") | map(attribute="content") | first }}<|im_end|>
 <|im_start|>assistant
@@ -1,13 +1,7 @@
 <|im_start|>system
 Judge whether the Document meets the requirements based on the Query and the Instruct provided. Note that the answer can only be "yes" or "no".<|im_end|>
 <|im_start|>user
-<Instruct>: {{
-    messages
-    | selectattr("role", "eq", "system")
-    | map(attribute="content")
-    | first
-    | default("Given a search query, retrieve relevant candidates that answer the query.")
-}}<Query>:{{
+<Instruct>: {{ instruction | default(instruct | default(messages | selectattr("role", "eq", "system") | map(attribute="content") | first | default("Given a search query, retrieve relevant candidates that answer the query.", true), true), true) }}<Query>:{{
    messages
    | selectattr("role", "eq", "query")
    | map(attribute="content")
@@ -19,7 +19,7 @@ pillow  # Required for image processing
 prometheus-fastapi-instrumentator >= 7.0.0
 tiktoken >= 0.6.0  # Required for DBRX tokenizer
 lm-format-enforcer == 0.11.3
-llguidance >= 1.3.0, < 1.4.0; platform_machine == "x86_64" or platform_machine == "arm64" or platform_machine == "aarch64" or platform_machine == "ppc64le"
+llguidance >= 1.7.0, < 1.8.0; platform_machine == "x86_64" or platform_machine == "arm64" or platform_machine == "aarch64" or platform_machine == "ppc64le"
 outlines_core == 0.2.14
 # required for outlines backend disk cache
 diskcache == 5.6.3
@@ -21,5 +21,8 @@ nvidia-cudnn-frontend>=1.13.0,<1.19.0
 fastsafetensors >= 0.2.2

 # QuACK and Cutlass DSL for FA4 (cute-DSL implementation)
-nvidia-cutlass-dsl[cu13]>=4.4.2
+nvidia-cutlass-dsl[cu13]==4.5.0
 quack-kernels>=0.3.3
+
+# Tokenspeed_MLA for faster mla with spec decode
+tokenspeed-mla==0.1.2
@@ -21,6 +21,3 @@ timm>=1.0.17
 # amd-quark: required for Quark quantization on ROCm 
 # To be consistent with test_quark.py
 amd-quark>=0.8.99
-# tilelang has to be installed for mhc module to be
-# imported correctly.
-tilelang==0.1.9
@@ -36,7 +36,7 @@ num2words # required for smolvlm test
 open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_common.py
 opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
-lm-eval[api]>=0.4.11 # required for model evaluation test
+lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
 transformers==5.5.3
 tokenizers==0.22.2
@@ -41,7 +41,6 @@ attrs==24.2.0
    # via
    #   aiohttp
    #   hypothesis
-    #   jsonlines
    #   jsonschema
    #   pytest-subtests
    #   referencing
@@ -357,8 +356,6 @@ joblib==1.4.2
    #   librosa
    #   nltk
    #   scikit-learn
-jsonlines==4.0.0
-    # via lm-eval
 jsonpointer==3.0.0
    # via jsonschema
 jsonschema==4.23.0
@@ -387,7 +384,7 @@ librosa==0.10.2.post1
    # via -r requirements/test/cuda.in
 llvmlite==0.47.0
    # via numba
-lm-eval==0.4.11
+lm-eval==0.4.12
    # via -r requirements/test/cuda.in
 lxml==5.3.0
    # via
@@ -1141,5 +1138,3 @@ yarl==1.17.1
    #   schemathesis
 zipp==3.23.0
    # via importlib-metadata
-zstandard==0.23.0
-    # via lm-eval
@@ -27,7 +27,7 @@ mistral_common[image,audio] >= 1.11.2 # required for voxtral test
 num2words # required for smolvlm test
 opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
-lm-eval[api]>=0.4.11 # required for model evaluation test
+lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
 transformers==5.5.3
 tokenizers==0.22.2
@@ -35,7 +35,7 @@ num2words # required for smolvlm test
 open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_common.py
 opencv-python-headless>=4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
-lm-eval[api]>=0.4.11 # required for model evaluation test
+lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
 transformers==5.5.3
 tokenizers==0.22.2
@@ -55,7 +55,6 @@ astor==0.8.1
 attrs==26.1.0
    # via
    #   aiohttp
-    #   jsonlines
    #   jsonschema
    #   pytest-subtests
    #   referencing
@@ -433,8 +432,6 @@ joblib==1.5.3
    #   librosa
    #   nltk
    #   scikit-learn
-jsonlines==4.0.0
-    # via lm-eval
 jsonpointer==3.1.0
    # via jsonschema
 jsonschema==4.26.0
@@ -466,13 +463,13 @@ libnacl==2.1.0
    # via tensorizer
 librosa==0.10.2.post1
    # via -r requirements/test/rocm.in
-llguidance==1.3.0
+llguidance==1.7.5
    # via
    #   -c requirements/common.txt
    #   -r requirements/test/../common.txt
 llvmlite==0.47.0
    # via numba
-lm-eval==0.4.11
+lm-eval==0.4.12
    # via -r requirements/test/rocm.in
 lm-format-enforcer==0.11.3
    # via
@@ -1362,8 +1359,6 @@ yarl==1.23.0
    #   schemathesis
 zipp==3.23.0
    # via importlib-metadata
-zstandard==0.25.0
-    # via lm-eval

 # The following packages were excluded from the output:
 # torch
@@ -15,4 +15,5 @@ torch==2.11.0+xpu
 torchaudio
 torchvision

+auto_round_lib>=0.13.0
 vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.8/vllm_xpu_kernels-0.1.8-cp38-abi3-manylinux_2_28_x86_64.whl
@@ -103,13 +103,6 @@ def run_e2e_fusion_test(monkeypatch, caplog_mp_spawn):
            if not supports_trtllm_attention():
                matches = matches._replace(attn_quant_fusion=0)

-        # TODO: remove this after finishing migration from envs to model kwargs
-        if model_name == "openai/gpt-oss-20b":
-            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.
        monkeypatch.setenv("VLLM_DISABLE_COMPILE_CACHE", "1")
@@ -214,4 +214,9 @@ gpt_oss_20b = ModelFusionInfo(
        sequence_parallel=n_layers * 2 + 1,
        async_tp=n_layers * 2,
    ),
+    model_kwargs=(
+        {"quantization_config": {"moe": {"activation": "mxfp8"}}}
+        if is_blackwell()
+        else {}
+    ),
 )
@@ -7,7 +7,7 @@ import torch

 import vllm.config
 from tests.compile.backend import TestBackend
-from vllm._aiter_ops import is_aiter_found_and_supported, rocm_aiter_ops
+from vllm._aiter_ops import rocm_aiter_ops
 from vllm.compilation.passes.utility.noop_elimination import NoOpEliminationPass
 from vllm.compilation.passes.utility.post_cleanup import PostCleanupPass
 from vllm.config import (
@@ -72,9 +72,8 @@ class TestModel(torch.nn.Module):
@pytest.mark.parametrize("hidden_size", [2880])
@pytest.mark.parametrize("num_local_experts", [128])
@pytest.mark.parametrize("x_pad_to_multiple", [256])
-@pytest.mark.skipif(
-    not is_aiter_found_and_supported(),
-    reason="Only test on ROCm with AITER installed and supported",
+@pytest.mark.skip(
+    reason="Skipping for now because of the accuracy issue. See: https://github.com/ROCm/aiter/issues/2614"
 )
 def test_fuse_act_padding(
    dtype: torch.dtype,
@@ -116,7 +115,6 @@ def test_fuse_act_padding(

        x = torch.rand(1, hidden_size)
        torch._dynamo.mark_dynamic(x, 0)
-
        outputs_unfused = model(x)

        model_fused = torch.compile(model, backend=backend)
@@ -40,7 +40,7 @@ from vllm.model_executor.kernels.linear import (
    TritonFp8BlockScaledMMKernel,
    _KernelT,
 )
-from vllm.model_executor.layers.layernorm import RMSNorm
+from vllm.model_executor.layers.layernorm import RMSNorm, RMSNormGated
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
    GroupShape,
    create_fp8_quant_key,
@@ -99,8 +99,6 @@ KERNEL_GROUPSHAPE_COMBINATIONS = (

 # For Aiter tests we toggle use_aiter_quant_op
 AITER_KERNEL_GROUPSHAPE_COMBINATIONS = [
-    # Per-token with ROCmFP8ScaledMMLinearKernel
-    (ROCmFP8ScaledMMLinearKernel, GroupShape.PER_TENSOR, False),
    # Per-token with RowWiseTorchFP8ScaledMMLinearKernel
    (RowWiseTorchFP8ScaledMMLinearKernel, GroupShape.PER_TOKEN, True),
    (RowWiseTorchFP8ScaledMMLinearKernel, GroupShape.PER_TOKEN, False),
@@ -441,3 +439,242 @@ def test_aiter_fusion_rmsnorm_quant(
        _run_fusion_test(
            model, fusion_pass, vllm_config, dtype, hidden_size, num_tokens
        )
+
+
+class TestGatedModel(torch.nn.Module):
+    """Model that uses RMSNormGated + reshape + group FP8 quant + linear.
+
+    Mimics GatedDeltaNetAttention's output projection path where:
+    - RMSNormGated operates on per-head tensors (N*H, D)
+    - Output is reshaped to (N, H*D) before group quantization + linear
+    """
+
+    def __init__(
+        self,
+        num_heads: int,
+        head_dim: int,
+        eps: float,
+        force_kernel: type[_KernelT],
+        group_shape: GroupShape,
+        dtype: torch.dtype,
+        use_aiter_quant: bool = True,
+    ):
+        super().__init__()
+        self.num_heads = num_heads
+        self.head_dim = head_dim
+        hidden_dim = num_heads * head_dim
+
+        self.norm = RMSNormGated(
+            head_dim,
+            eps=eps,
+            group_size=None,
+            norm_before_gate=True,
+        )
+
+        self.activation_quant_key = create_fp8_quant_key(
+            static=False, group_shape=group_shape
+        )
+        self.weight_quant_key = create_fp8_quant_key(
+            static=True, group_shape=GroupShape(group_shape.col, group_shape.col)
+        )
+
+        self.fp8_linear = TestFP8Layer(
+            weight_shape=(hidden_dim, hidden_dim),
+            activation_quant_key=self.activation_quant_key,
+            weight_quant_key=self.weight_quant_key,
+            force_kernel=force_kernel,
+            transpose_weights=True,
+            input_dtype=dtype,
+        )
+        self.fp8_linear.kernel.quant_fp8.use_aiter = use_aiter_quant
+
+    def forward(self, x, z):
+        num_heads = self.num_heads
+        head_dim = self.head_dim
+        hidden_dim = num_heads * head_dim
+        x = torch.relu(x)
+        z = torch.relu(z)
+        x_heads = x.reshape(-1, num_heads, head_dim).reshape(-1, head_dim)
+        z_heads = z.reshape(-1, num_heads, head_dim).reshape(-1, head_dim)
+        normed = self.norm(x_heads, z_heads)
+        merged = normed.reshape(-1, hidden_dim)
+        out = self.fp8_linear(merged)
+        return out
+
+    def ops_in_model_after(self):
+        from vllm.compilation.passes.fusion.rocm_aiter_fusion import (
+            AiterRMSNormGatedFp8GroupQuantPattern,
+        )
+
+        return [AiterRMSNormGatedFp8GroupQuantPattern.FUSED_OP]
+
+
+class _MockGDNLayer:
+    """Minimal mock to populate static_forward_context for pass discovery.
+
+    Uses __class__ assignment to pass isinstance checks against
+    GatedDeltaNetAttention without requiring a full config-based init.
+    """
+
+    def __init__(self, num_v_heads: int, head_v_dim: int, tp_size: int = 1):
+        self.num_v_heads = num_v_heads
+        self.head_v_dim = head_v_dim
+        self.tp_size = tp_size
+
+        from vllm.model_executor.layers.mamba.gdn_linear_attn import (
+            GatedDeltaNetAttention,
+        )
+
+        self.__class__ = GatedDeltaNetAttention
+
+
+@pytest.mark.parametrize("dtype", [torch.bfloat16])
+@pytest.mark.parametrize("num_heads", [2])
+@pytest.mark.parametrize("head_dim", [128])
+@pytest.mark.parametrize("num_tokens", [8])
+@pytest.mark.parametrize("eps", [1e-5, 1e-6])
+@pytest.mark.skipif(
+    (not current_platform.is_rocm() or not IS_AITER_FOUND),
+    reason="Only test on ROCm with aiter package installed",
+)
+def test_aiter_fusion_rmsnorm_gated_quant(
+    dtype: torch.dtype,
+    num_heads: int,
+    head_dim: int,
+    num_tokens: int,
+    eps: float,
+    monkeypatch: pytest.MonkeyPatch,
+):
+    group_shape = GroupShape(1, 128)
+    vllm_config = VllmConfig(
+        model_config=ModelConfig(dtype=dtype),
+        compilation_config=CompilationConfig(
+            mode=CompilationMode.VLLM_COMPILE,
+            custom_ops=["-rms_norm", "-silu_and_mul", "-quant_fp8"],
+            pass_config=PassConfig(fuse_norm_quant=True, eliminate_noops=True),
+        ),
+    )
+
+    with vllm.config.set_current_vllm_config(vllm_config), monkeypatch.context() as m:
+        from vllm.compilation.passes.fusion.rocm_aiter_fusion import (
+            RocmAiterRMSNormQuantFusionPass,
+        )
+
+        m.setenv("VLLM_ROCM_USE_AITER", "1")
+        rocm_aiter_ops.refresh_env_variables()
+
+        # Register a mock GDN layer so the pass discovers num_heads/head_dim
+        mock_gdn = _MockGDNLayer(num_v_heads=num_heads, head_v_dim=head_dim, tp_size=1)
+        vllm_config.compilation_config.static_forward_context["mock_gdn_layer"] = (
+            mock_gdn
+        )
+
+        torch.set_default_device("cuda")
+        torch.set_default_dtype(dtype)
+        torch.manual_seed(1)
+
+        fusion_pass = RocmAiterRMSNormQuantFusionPass(vllm_config)
+
+        model = TestGatedModel(
+            num_heads=num_heads,
+            head_dim=head_dim,
+            eps=eps,
+            force_kernel=AiterFp8BlockScaledMMKernel,
+            group_shape=group_shape,
+            dtype=dtype,
+            use_aiter_quant=True,
+        )
+
+        noop_pass = NoOpEliminationPass(vllm_config)
+        cleanup_pass = PostCleanupPass(vllm_config)
+
+        backend = TestBackend(noop_pass, fusion_pass, cleanup_pass)
+        backend2 = TestBackend(noop_pass, cleanup_pass)
+
+        hidden_dim = num_heads * head_dim
+        x = torch.rand(num_tokens, hidden_dim)
+        z = torch.rand(num_tokens, hidden_dim)
+        torch._dynamo.mark_dynamic(x, 0)
+        torch._dynamo.mark_dynamic(z, 0)
+
+        model_fused = torch.compile(model, backend=backend)
+        result_fused = model_fused(x, z)
+
+        model_unfused = torch.compile(model, backend=backend2)
+        result_unfused = model_unfused(x, z)
+
+        torch.testing.assert_close(result_fused, result_unfused, atol=1e-2, rtol=1e-2)
+
+        assert fusion_pass.matched_count == 1
+        backend.check_after_ops(model.ops_in_model_after())
+
+
+@pytest.mark.parametrize("dtype", [torch.bfloat16])
+@pytest.mark.parametrize("num_heads", [2])
+@pytest.mark.parametrize("head_dim", [128])
+@pytest.mark.parametrize("num_tokens", [8])
+@pytest.mark.parametrize("eps", [1e-6])
+@pytest.mark.skipif(
+    (not current_platform.is_rocm() or not IS_AITER_FOUND),
+    reason="Only test on ROCm with aiter package installed",
+)
+def test_aiter_fusion_rmsnorm_gated_quant_no_gdn_layers(
+    dtype: torch.dtype,
+    num_heads: int,
+    head_dim: int,
+    num_tokens: int,
+    eps: float,
+    monkeypatch: pytest.MonkeyPatch,
+):
+    """Verify that without GDN layers in static_forward_context,
+    the gated pattern is not registered and no matches occur."""
+    group_shape = GroupShape(1, 128)
+    vllm_config = VllmConfig(
+        model_config=ModelConfig(dtype=dtype),
+        compilation_config=CompilationConfig(
+            mode=CompilationMode.VLLM_COMPILE,
+            custom_ops=["-rms_norm", "-silu_and_mul", "-quant_fp8"],
+            pass_config=PassConfig(fuse_norm_quant=True, eliminate_noops=True),
+        ),
+    )
+
+    with vllm.config.set_current_vllm_config(vllm_config), monkeypatch.context() as m:
+        from vllm.compilation.passes.fusion.rocm_aiter_fusion import (
+            RocmAiterRMSNormQuantFusionPass,
+        )
+
+        m.setenv("VLLM_ROCM_USE_AITER", "1")
+        rocm_aiter_ops.refresh_env_variables()
+
+        torch.set_default_device("cuda")
+        torch.set_default_dtype(dtype)
+        torch.manual_seed(1)
+
+        # No mock GDN layer registered -- pass should not register gated pattern
+        fusion_pass = RocmAiterRMSNormQuantFusionPass(vllm_config)
+
+        model = TestGatedModel(
+            num_heads=num_heads,
+            head_dim=head_dim,
+            eps=eps,
+            force_kernel=AiterFp8BlockScaledMMKernel,
+            group_shape=group_shape,
+            dtype=dtype,
+            use_aiter_quant=True,
+        )
+
+        noop_pass = NoOpEliminationPass(vllm_config)
+        cleanup_pass = PostCleanupPass(vllm_config)
+
+        backend = TestBackend(noop_pass, fusion_pass, cleanup_pass)
+
+        hidden_dim = num_heads * head_dim
+        x = torch.rand(num_tokens, hidden_dim)
+        z = torch.rand(num_tokens, hidden_dim)
+        torch._dynamo.mark_dynamic(x, 0)
+        torch._dynamo.mark_dynamic(z, 0)
+
+        model_fused = torch.compile(model, backend=backend)
+        model_fused(x, z)
+
+        assert fusion_pass.matched_count == 0
@@ -224,19 +224,28 @@ def init_test_http_connection():
 def dist_init():
    from tests.utils import ensure_current_vllm_config

-    temp_file = tempfile.mkstemp()[1]
+    # Close the fd returned by mkstemp; FileStore opens the path itself.
+    # Leaving it open leaks one FD per test and eventually exhausts the
+    # ulimit, causing FileStore's destructor to throw c10::DistStoreError
+    # ("Too many open files") during gc and abort the process.
+    fd, temp_file = tempfile.mkstemp()
+    os.close(fd)

-    with ensure_current_vllm_config():
-        init_distributed_environment(
-            world_size=1,
-            rank=0,
-            distributed_init_method=f"file://{temp_file}",
-            local_rank=0,
-            backend="nccl",
-        )
-        initialize_model_parallel(1, 1)
-        yield
-    cleanup_dist_env_and_memory()
+    try:
+        with ensure_current_vllm_config():
+            init_distributed_environment(
+                world_size=1,
+                rank=0,
+                distributed_init_method=f"file://{temp_file}",
+                local_rank=0,
+                backend="nccl",
+            )
+            initialize_model_parallel(1, 1)
+            yield
+        cleanup_dist_env_and_memory()
+    finally:
+        with contextlib.suppress(OSError):
+            os.unlink(temp_file)


@pytest.fixture
@@ -785,10 +794,15 @@ class HfRunner:
        audios: PromptAudioInput | None = None,
        videos: PromptVideoInput | None = None,
        use_cache: bool = True,
+        tokenization_kwargs: dict[str, Any] | None = None,
        **kwargs: Any,
    ) -> list[TokensTextLogprobs]:
        all_inputs = self.get_inputs(
-            prompts, images=images, videos=videos, audios=audios
+            prompts,
+            images=images,
+            videos=videos,
+            audios=audios,
+            tokenization_kwargs=tokenization_kwargs,
        )

        all_logprobs: list[list[dict[int, float]]] = []
@@ -368,12 +368,8 @@ def test_attention_config():
            "true",
            "--attention-config.flash_attn_max_num_splits_for_cuda_graph",
            "16",
-            "--attention-config.use_trtllm_ragged_deepseek_prefill",
-            "true",
            "--attention-config.use_trtllm_attention",
            "true",
-            "--attention-config.disable_flashinfer_prefill",
-            "true",
            "--attention-config.disable_flashinfer_q_quantization",
            "true",
        ]
@@ -385,9 +381,7 @@ def test_attention_config():
    assert engine_args.attention_config.flash_attn_version == 3
    assert engine_args.attention_config.use_prefill_decode_attention is True
    assert engine_args.attention_config.flash_attn_max_num_splits_for_cuda_graph == 16
-    assert engine_args.attention_config.use_trtllm_ragged_deepseek_prefill is True
    assert engine_args.attention_config.use_trtllm_attention is True
-    assert engine_args.attention_config.disable_flashinfer_prefill is True
    assert engine_args.attention_config.disable_flashinfer_q_quantization is True

    # set to string form of a dict with all fields
@@ -397,10 +391,7 @@ def test_attention_config():
            '{"backend": "FLASHINFER", "flash_attn_version": 2, '
            '"use_prefill_decode_attention": false, '
            '"flash_attn_max_num_splits_for_cuda_graph": 8, '
-            '"use_cudnn_prefill": false, '
-            '"use_trtllm_ragged_deepseek_prefill": false, '
            '"use_trtllm_attention": false, '
-            '"disable_flashinfer_prefill": false, '
            '"disable_flashinfer_q_quantization": false}',
        ]
    )
@@ -411,10 +402,7 @@ def test_attention_config():
    assert engine_args.attention_config.flash_attn_version == 2
    assert engine_args.attention_config.use_prefill_decode_attention is False
    assert engine_args.attention_config.flash_attn_max_num_splits_for_cuda_graph == 8
-    assert engine_args.attention_config.use_cudnn_prefill is False
-    assert engine_args.attention_config.use_trtllm_ragged_deepseek_prefill is False
    assert engine_args.attention_config.use_trtllm_attention is False
-    assert engine_args.attention_config.disable_flashinfer_prefill is False
    assert engine_args.attention_config.disable_flashinfer_q_quantization is False

    # test --attention-backend flows into VllmConfig.attention_config
@@ -199,12 +199,107 @@ class TestExtractHarmonyStreamingDelta:
        assert delta_message.content == delta_text
        assert tools_streamed is False

+    @pytest.mark.parametrize("channel", ["commentary", "analysis"])
+    @patch("vllm.entrypoints.openai.chat_completion.stream_harmony.make_tool_call_id")
+    def test_new_tool_call_without_functions_prefix(
+        self, mock_make_tool_call_id, channel
+    ):
+        mock_make_tool_call_id.return_value = "call_bare123"
+        parser = MockStreamableParser()
+
+        token_states = [TokenState(channel=channel, recipient="get_weather", text="")]
+
+        delta_message, tools_streamed = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient=None,
+            include_reasoning=False,
+        )
+
+        assert delta_message is not None
+        assert len(delta_message.tool_calls) == 1
+        tool_call = delta_message.tool_calls[0]
+        assert tool_call.id == "call_bare123"
+        assert tool_call.type == "function"
+        assert tool_call.function.name == "get_weather"
+        assert tool_call.function.arguments == ""
+        assert tool_call.index == 0
+        assert tools_streamed is True
+
+    @pytest.mark.parametrize("channel", ["commentary", "analysis"])
+    def test_tool_call_argument_streaming_without_functions_prefix(self, channel):
+        parser = MockStreamableParser()
+        args_text = '{"location": "Paris"}'
+
+        token_states = [
+            TokenState(channel=channel, recipient="get_weather", text=args_text)
+        ]
+
+        delta_message, tools_streamed = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient="get_weather",
+            include_reasoning=False,
+        )
+
+        assert delta_message is not None
+        tool_call = delta_message.tool_calls[0]
+        assert tool_call.id is None
+        assert tool_call.function.arguments == args_text
+        assert tool_call.index == 0
+        assert tools_streamed is True
+
+    def test_tool_call_index_from_previous_messages_without_functions_prefix(self):
+        messages = [
+            MockMessage(channel="commentary", recipient="tool1"),
+        ]
+        parser = MockStreamableParser(messages=messages)
+
+        token_states = [
+            TokenState(channel="commentary", recipient="tool2", text="args")
+        ]
+
+        delta_message, _ = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient="tool2",
+            include_reasoning=False,
+        )
+
+        assert delta_message.tool_calls[0].index == 1
+
+    @pytest.mark.parametrize("channel", ["commentary", "analysis"])
+    @patch("vllm.entrypoints.openai.chat_completion.stream_harmony.make_tool_call_id")
+    def test_new_tool_call_dotted_function_name(self, mock_make_tool_call_id, channel):
+        mock_make_tool_call_id.return_value = "call_dotted123"
+        parser = MockStreamableParser()
+
+        token_states = [TokenState(channel=channel, recipient="math.sum", text="")]
+
+        delta_message, tools_streamed = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient=None,
+            include_reasoning=False,
+        )
+
+        assert delta_message is not None
+        assert len(delta_message.tool_calls) == 1
+        tool_call = delta_message.tool_calls[0]
+        assert tool_call.id == "call_dotted123"
+        assert tool_call.type == "function"
+        assert tool_call.function.name == "math.sum"
+        assert tool_call.function.arguments == ""
+        assert tool_call.index == 0
+        assert tools_streamed is True
+
    @pytest.mark.parametrize(
        "channel,recipient",
        [
            (None, None),
            ("unknown_channel", None),
            ("commentary", "browser.search"),
+            ("commentary", "assistant"),
        ],
    )
    def test_returns_none_for_invalid_inputs(self, channel, recipient):
@@ -348,3 +443,92 @@ class TestExtractHarmonyStreamingDelta:
        assert tool_c_args.function.arguments == '{"key_c": "val_c"}'

        assert delta_message.content == "Thinking... Thinking again..."
+
+
+class TestToolCallsOnNonStandardChannels:
+    """Tool calls are detected by recipient, not channel.
+
+    Models sometimes emit tool calls on unexpected channels (e.g. ``comment``
+    instead of ``commentary``).  These tests verify that the streaming delta
+    extraction is channel-agnostic for tool call detection.
+    """
+
+    @patch("vllm.entrypoints.openai.chat_completion.stream_harmony.make_tool_call_id")
+    def test_prefixed_tool_call_on_comment_channel(self, mock_make_tool_call_id):
+        mock_make_tool_call_id.return_value = "call_comment_chan"
+        parser = MockStreamableParser()
+
+        token_states = [
+            TokenState(channel="comment", recipient="functions.get_weather", text="")
+        ]
+
+        delta_message, tools_streamed = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient=None,
+            include_reasoning=False,
+        )
+
+        assert delta_message is not None
+        assert len(delta_message.tool_calls) == 1
+        assert delta_message.tool_calls[0].function.name == "get_weather"
+        assert tools_streamed is True
+
+    @patch("vllm.entrypoints.openai.chat_completion.stream_harmony.make_tool_call_id")
+    def test_bare_tool_call_on_comment_channel(self, mock_make_tool_call_id):
+        mock_make_tool_call_id.return_value = "call_bare_comment"
+        parser = MockStreamableParser()
+
+        token_states = [TokenState(channel="comment", recipient="get_weather", text="")]
+
+        delta_message, tools_streamed = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient=None,
+            include_reasoning=False,
+        )
+
+        assert delta_message is not None
+        assert len(delta_message.tool_calls) == 1
+        assert delta_message.tool_calls[0].function.name == "get_weather"
+        assert tools_streamed is True
+
+    def test_tool_call_arguments_on_comment_channel(self):
+        parser = MockStreamableParser()
+        args_text = '{"location": "Paris"}'
+
+        token_states = [
+            TokenState(
+                channel="comment", recipient="functions.get_weather", text=args_text
+            )
+        ]
+
+        delta_message, tools_streamed = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient="functions.get_weather",
+            include_reasoning=False,
+        )
+
+        assert delta_message is not None
+        assert delta_message.tool_calls[0].function.arguments == args_text
+        assert tools_streamed is True
+
+    def test_base_index_counts_tool_calls_on_comment_channel(self):
+        messages = [
+            MockMessage(channel="comment", recipient="functions.tool1"),
+        ]
+        parser = MockStreamableParser(messages=messages)
+
+        token_states = [
+            TokenState(channel="commentary", recipient="functions.tool2", text="args")
+        ]
+
+        delta_message, _ = extract_harmony_streaming_delta(
+            harmony_parser=parser,
+            token_states=token_states,
+            prev_recipient="functions.tool2",
+            include_reasoning=False,
+        )
+
+        assert delta_message.tool_calls[0].index == 1
@@ -7,9 +7,11 @@ from openai_harmony import Message, Role
 from tests.entrypoints.openai.utils import verify_harmony_messages
 from vllm.entrypoints.openai.parser.harmony_utils import (
    auto_drop_analysis_messages,
+    extract_function_from_recipient,
    get_encoding,
    get_system_message,
    has_custom_tools,
+    is_function_recipient,
    parse_chat_input_to_harmony_message,
    parse_chat_output,
 )
@@ -19,6 +21,182 @@ from vllm.entrypoints.openai.responses.harmony import (
 )


+class TestIsFunctionRecipient:
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "functions.get_weather",
+            "functions.search_web",
+            "functions.math.sum",
+        ],
+    )
+    def test_functions_prefix_accepted(self, recipient):
+        assert is_function_recipient(recipient) is True
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "get_weather",
+            "search_web",
+            "calculator",
+            "my-tool",
+        ],
+    )
+    def test_bare_function_name_accepted(self, recipient):
+        assert is_function_recipient(recipient) is True
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "assistant",
+        ],
+    )
+    def test_assistant_rejected(self, recipient):
+        assert is_function_recipient(recipient) is False
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "math.sum",
+            "code.run",
+            "namespace.tool_name",
+            "my.deeply.nested.tool",
+        ],
+    )
+    def test_dotted_function_names_accepted(self, recipient):
+        assert is_function_recipient(recipient) is True
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "python",
+            "browser",
+            "container",
+        ],
+    )
+    def test_builtin_tool_names_rejected(self, recipient):
+        assert is_function_recipient(recipient) is False
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "python.run",
+            "python.execute",
+            "browser.search",
+            "browser.open",
+            "container.exec",
+        ],
+    )
+    def test_builtin_dotted_variants_rejected(self, recipient):
+        assert is_function_recipient(recipient) is False
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "",
+            "functions.",
+        ],
+    )
+    def test_empty_recipients_rejected(self, recipient):
+        assert is_function_recipient(recipient) is False
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "<|start|>",
+            "<|end|>",
+            "<|channel|>",
+        ],
+    )
+    def test_harmony_tokens_rejected(self, recipient):
+        assert is_function_recipient(recipient) is False
+
+
+class TestIsFunctionRecipientWithAllowedNames:
+    """Tests for is_function_recipient with allowed_function_tool_names."""
+
+    def test_prefixed_always_accepted(self):
+        """functions. prefix is always accepted regardless of allowed names."""
+        fn_names = frozenset({"other_tool"})
+        assert is_function_recipient("functions.get_weather", fn_names) is True
+
+    def test_bare_name_accepted_when_in_allowed_names(self):
+        fn_names = frozenset({"get_weather", "search_web"})
+        assert is_function_recipient("get_weather", fn_names) is True
+        assert is_function_recipient("search_web", fn_names) is True
+
+    def test_bare_name_rejected_when_not_in_allowed_names(self):
+        fn_names = frozenset({"get_weather"})
+        assert is_function_recipient("unknown_tool", fn_names) is False
+
+    def test_dotted_name_accepted_when_in_allowed_names(self):
+        fn_names = frozenset({"math.sum", "namespace.tool_name"})
+        assert is_function_recipient("math.sum", fn_names) is True
+        assert is_function_recipient("namespace.tool_name", fn_names) is True
+
+    def test_dotted_name_rejected_when_not_in_allowed_names(self):
+        fn_names = frozenset({"get_weather"})
+        assert is_function_recipient("custom_server.search", fn_names) is False
+
+    def test_empty_allowed_names_rejects_bare_names(self):
+        """Empty frozenset means no function tools — bare names are not functions."""
+        fn_names: frozenset[str] = frozenset()
+        assert is_function_recipient("get_weather", fn_names) is False
+        assert is_function_recipient("math.sum", fn_names) is False
+
+    def test_builtin_tools_always_rejected(self):
+        fn_names = frozenset({"python", "browser", "container"})
+        assert is_function_recipient("python", fn_names) is False
+        assert is_function_recipient("browser", fn_names) is False
+        assert is_function_recipient("container", fn_names) is False
+
+    def test_builtin_dotted_always_rejected(self):
+        fn_names = frozenset({"python.run", "browser.search"})
+        assert is_function_recipient("python.run", fn_names) is False
+        assert is_function_recipient("browser.search", fn_names) is False
+
+    def test_none_allowed_names_uses_heuristic(self):
+        """When allowed names is None (Chat Completions), use heuristic."""
+        assert is_function_recipient("get_weather", None) is True
+        assert is_function_recipient("math.sum", None) is True
+        assert is_function_recipient("python", None) is False
+
+
+class TestExtractFunctionFromRecipient:
+    @pytest.mark.parametrize(
+        "recipient,expected",
+        [
+            ("functions.get_weather", "get_weather"),
+            ("functions.search_web", "search_web"),
+            ("functions.", ""),
+        ],
+    )
+    def test_strips_functions_prefix(self, recipient, expected):
+        assert extract_function_from_recipient(recipient) == expected
+
+    @pytest.mark.parametrize(
+        "recipient",
+        [
+            "get_weather",
+            "calculator",
+            "my-tool",
+        ],
+    )
+    def test_bare_name_returned_as_is(self, recipient):
+        assert extract_function_from_recipient(recipient) == recipient
+
+    @pytest.mark.parametrize(
+        "recipient,expected",
+        [
+            ("functions.math.sum", "math.sum"),
+            ("math.sum", "math.sum"),
+            ("namespace.tool_name", "namespace.tool_name"),
+        ],
+    )
+    def test_dotted_function_name_extraction(self, recipient, expected):
+        assert extract_function_from_recipient(recipient) == expected
+
+
 class TestCommonParseInputToHarmonyMessage:
    """
    Tests for scenarios that are common to both Chat Completion
@@ -246,7 +246,8 @@ class TestHarmonyToResponseOutput:
        message = message.with_channel("commentary")
        message = message.with_recipient("custom_tool")

-        output_items = harmony_to_response_output(message)
+        fn_names = frozenset({"other_tool"})
+        output_items = harmony_to_response_output(message, fn_names)

        assert len(output_items) == 1
        assert isinstance(output_items[0], McpCall)
@@ -286,13 +287,179 @@ class TestHarmonyToResponseOutput:
        assert len(output_items) == 0


+class TestHarmonyToResponseOutputWithFunctionToolNames:
+    """Tests for bare function name handling with function_tool_names."""
+
+    def test_bare_name_creates_function_call_when_in_tool_names(self):
+        """Bare function name matching a known tool creates function call."""
+        message = Message.from_role_and_content(
+            Role.ASSISTANT, '{"location": "San Francisco"}'
+        )
+        message = message.with_channel("commentary")
+        message = message.with_recipient("get_weather")
+
+        fn_names = frozenset({"get_weather"})
+        output_items = harmony_to_response_output(message, fn_names)
+
+        assert len(output_items) == 1
+        assert isinstance(output_items[0], ResponseFunctionToolCall)
+        assert output_items[0].type == "function_call"
+        assert output_items[0].name == "get_weather"
+        assert output_items[0].arguments == '{"location": "San Francisco"}'
+
+    def test_bare_name_creates_mcp_call_when_not_in_tool_names(self):
+        """Bare name not matching any known tool creates MCP call."""
+        message = Message.from_role_and_content(Role.ASSISTANT, '{"arg": "value"}')
+        message = message.with_channel("commentary")
+        message = message.with_recipient("custom_tool")
+
+        fn_names = frozenset({"get_weather"})
+        output_items = harmony_to_response_output(message, fn_names)
+
+        assert len(output_items) == 1
+        assert isinstance(output_items[0], McpCall)
+        assert output_items[0].type == "mcp_call"
+
+    def test_dotted_function_name_creates_function_call(self):
+        """Dotted function name in tool names creates function call."""
+        message = Message.from_role_and_content(Role.ASSISTANT, '{"a": 1, "b": 2}')
+        message = message.with_channel("commentary")
+        message = message.with_recipient("math.sum")
+
+        fn_names = frozenset({"math.sum"})
+        output_items = harmony_to_response_output(message, fn_names)
+
+        assert len(output_items) == 1
+        assert isinstance(output_items[0], ResponseFunctionToolCall)
+        assert output_items[0].name == "math.sum"
+
+    def test_empty_tool_names_defaults_to_mcp(self):
+        """With empty function_tool_names, bare names become MCP calls."""
+        message = Message.from_role_and_content(Role.ASSISTANT, '{"arg": "value"}')
+        message = message.with_channel("commentary")
+        message = message.with_recipient("get_weather")
+
+        output_items = harmony_to_response_output(message, frozenset())
+
+        assert len(output_items) == 1
+        assert isinstance(output_items[0], McpCall)
+
+    def test_prefixed_name_always_function_call(self):
+        """functions. prefix always creates function call even with empty tool names."""
+        message = Message.from_role_and_content(Role.ASSISTANT, '{"arg": "value"}')
+        message = message.with_channel("commentary")
+        message = message.with_recipient("functions.get_weather")
+
+        output_items = harmony_to_response_output(message, frozenset())
+
+        assert len(output_items) == 1
+        assert isinstance(output_items[0], ResponseFunctionToolCall)
+        assert output_items[0].name == "get_weather"
+
+
+class TestParserStateWithFunctionToolNames:
+    """Tests for parser_state_to_response_output with function_tool_names."""
+
+    def test_bare_name_creates_function_call(self):
+        from unittest.mock import Mock
+
+        parser = Mock()
+        parser.current_content = '{"arg": "value"}'
+        parser.current_role = Role.ASSISTANT
+        parser.current_channel = "commentary"
+        parser.current_recipient = "get_weather"
+
+        fn_names = frozenset({"get_weather"})
+        items = parser_state_to_response_output(parser, fn_names)
+
+        assert len(items) == 1
+        assert isinstance(items[0], ResponseFunctionToolCall)
+        assert items[0].name == "get_weather"
+        assert items[0].status == "in_progress"
+
+    def test_bare_name_creates_mcp_when_not_in_tool_names(self):
+        from unittest.mock import Mock
+
+        parser = Mock()
+        parser.current_content = '{"arg": "value"}'
+        parser.current_role = Role.ASSISTANT
+        parser.current_channel = "commentary"
+        parser.current_recipient = "unknown_tool"
+
+        fn_names = frozenset({"get_weather"})
+        items = parser_state_to_response_output(parser, fn_names)
+
+        assert len(items) == 1
+        assert isinstance(items[0], McpCall)
+        assert items[0].name == "unknown_tool"
+
+
+class TestToolCallsOnNonStandardChannels:
+    """Tests verifying tool calls are detected regardless of channel."""
+
+    def test_function_call_on_comment_channel(self):
+        message = Message.from_role_and_content(Role.ASSISTANT, '{"query": "weather"}')
+        message = message.with_channel("comment")
+        message = message.with_recipient("functions.get_weather")
+
+        output_items = harmony_to_response_output(message)
+
+        assert len(output_items) == 1
+        assert isinstance(output_items[0], ResponseFunctionToolCall)
+        assert output_items[0].type == "function_call"
+        assert output_items[0].name == "get_weather"
+
+    def test_bare_function_on_comment_channel(self):
+        message = Message.from_role_and_content(Role.ASSISTANT, '{"query": "weather"}')
+        message = message.with_channel("comment")
+        message = message.with_recipient("get_weather")
+
+        fn_names = frozenset({"get_weather"})
+        output_items = harmony_to_response_output(message, fn_names)
+
+        assert len(output_items) == 1
+        assert isinstance(output_items[0], ResponseFunctionToolCall)
+        assert output_items[0].name == "get_weather"
+
+    def test_parser_state_comment_channel_function(self):
+        from unittest.mock import Mock
+
+        parser = Mock()
+        parser.current_content = '{"arg": "value"}'
+        parser.current_role = Role.ASSISTANT
+        parser.current_channel = "comment"
+        parser.current_recipient = "functions.get_weather"
+
+        items = parser_state_to_response_output(parser)
+
+        assert len(items) == 1
+        assert isinstance(items[0], ResponseFunctionToolCall)
+        assert items[0].name == "get_weather"
+
+    def test_parser_state_comment_channel_mcp(self):
+        from unittest.mock import Mock
+
+        parser = Mock()
+        parser.current_content = '{"arg": "value"}'
+        parser.current_role = Role.ASSISTANT
+        parser.current_channel = "comment"
+        parser.current_recipient = "mcp.server.tool"
+
+        fn_names: frozenset[str] = frozenset()
+        items = parser_state_to_response_output(parser, fn_names)
+
+        assert len(items) == 1
+        assert isinstance(items[0], McpCall)
+
+
 def test_parse_mcp_call_basic() -> None:
    """Test that MCP calls are parsed with correct type and server_label."""
    message = Message.from_role_and_content(Role.ASSISTANT, '{"path": "/tmp"}')
    message = message.with_recipient("filesystem")
    message = message.with_channel("commentary")

-    output_items = harmony_to_response_output(message)
+    fn_names: frozenset[str] = frozenset()
+    output_items = harmony_to_response_output(message, fn_names)

    assert len(output_items) == 1
    assert isinstance(output_items[0], McpCall)
@@ -309,7 +476,8 @@ def test_parse_mcp_call_dotted_recipient() -> None:
    message = message.with_recipient("repo_browser.list")
    message = message.with_channel("commentary")

-    output_items = harmony_to_response_output(message)
+    fn_names: frozenset[str] = frozenset()
+    output_items = harmony_to_response_output(message, fn_names)

    assert len(output_items) == 1
    assert isinstance(output_items[0], McpCall)
@@ -371,7 +539,8 @@ def test_parser_state_to_response_output_commentary_channel() -> None:
    parser_mcp.current_channel = "commentary"
    parser_mcp.current_recipient = "filesystem"

-    mcp_items = parser_state_to_response_output(parser_mcp)
+    fn_names: frozenset[str] = frozenset()
+    mcp_items = parser_state_to_response_output(parser_mcp, fn_names)

    assert len(mcp_items) == 1
    assert isinstance(mcp_items[0], McpCall)
@@ -438,7 +607,8 @@ def test_parser_state_to_response_output_analysis_channel() -> None:
    parser_mcp.current_channel = "analysis"
    parser_mcp.current_recipient = "database"

-    mcp_items = parser_state_to_response_output(parser_mcp)
+    fn_names: frozenset[str] = frozenset()
+    mcp_items = parser_state_to_response_output(parser_mcp, fn_names)

    assert len(mcp_items) == 1
    assert isinstance(mcp_items[0], McpCall)
@@ -377,3 +377,135 @@ async def test_score_api_queries_list_documents_list(
        backend,
        "paired[3]_text_vs_text_plus_image",
    )
+
+
+INSTRUCTION = (
+    "Given a multimodal retrieval query, retrieve candidates that "
+    "visually or textually match the requested scene, object, or action."
+)
+
+
+@pytest.mark.asyncio
+async def test_score_api_instruction_field(
+    server: tuple[RemoteOpenAIServer, str],
+):
+    remote_server, _ = server
+
+    default_response = requests.post(
+        remote_server.url_for("score"),
+        json={
+            "model": MODEL_NAME,
+            "queries": query,
+            "documents": document,
+        },
+    )
+    default_response.raise_for_status()
+    default_score = ScoreResponse.model_validate(default_response.json())
+
+    instruction_response = requests.post(
+        remote_server.url_for("score"),
+        json={
+            "model": MODEL_NAME,
+            "queries": query,
+            "documents": document,
+            "instruction": INSTRUCTION,
+        },
+    )
+    instruction_response.raise_for_status()
+    instruction_score = ScoreResponse.model_validate(instruction_response.json())
+
+    assert instruction_score.id is not None
+    assert instruction_score.data is not None
+    assert len(instruction_score.data) == 1
+    assert instruction_score.usage.prompt_tokens > default_score.usage.prompt_tokens
+
+
+@pytest.mark.asyncio
+async def test_rerank_api_instruction_field(
+    server: tuple[RemoteOpenAIServer, str],
+):
+    remote_server, _ = server
+
+    doc_list = [
+        document,
+        {"content": [documents[0]]},
+        {"content": [documents[1]]},
+        {"content": [documents[0], documents[1]]},
+    ]
+
+    default_response = requests.post(
+        remote_server.url_for("rerank"),
+        json={
+            "model": MODEL_NAME,
+            "query": query,
+            "documents": doc_list,
+        },
+    )
+    default_response.raise_for_status()
+    default_rerank = RerankResponse.model_validate(default_response.json())
+
+    instruction_response = requests.post(
+        remote_server.url_for("rerank"),
+        json={
+            "model": MODEL_NAME,
+            "query": query,
+            "documents": doc_list,
+            "instruction": INSTRUCTION,
+        },
+    )
+    instruction_response.raise_for_status()
+    instruction_rerank = RerankResponse.model_validate(instruction_response.json())
+
+    assert instruction_rerank.id is not None
+    assert instruction_rerank.model is not None
+    assert instruction_rerank.usage is not None
+    assert len(instruction_rerank.results) == len(default_rerank.results)
+    assert instruction_rerank.usage.prompt_tokens > default_rerank.usage.prompt_tokens
+
+
+@pytest.mark.asyncio
+async def test_rerank_api_instruction_field_matches_chat_template_kwargs(
+    server: tuple[RemoteOpenAIServer, str],
+):
+    remote_server, _ = server
+
+    doc_list = [
+        document,
+        {"content": [documents[0]]},
+        {"content": [documents[1]]},
+        {"content": [documents[0], documents[1]]},
+    ]
+
+    field_response = requests.post(
+        remote_server.url_for("rerank"),
+        json={
+            "model": MODEL_NAME,
+            "query": query,
+            "documents": doc_list,
+            "instruction": INSTRUCTION,
+        },
+    )
+    field_response.raise_for_status()
+    field_rerank = RerankResponse.model_validate(field_response.json())
+
+    kwargs_response = requests.post(
+        remote_server.url_for("rerank"),
+        json={
+            "model": MODEL_NAME,
+            "query": query,
+            "documents": doc_list,
+            "chat_template_kwargs": {"instruction": INSTRUCTION},
+        },
+    )
+    kwargs_response.raise_for_status()
+    kwargs_rerank = RerankResponse.model_validate(kwargs_response.json())
+
+    assert kwargs_rerank.usage.prompt_tokens == field_rerank.usage.prompt_tokens
+
+    field_scores = [
+        r.relevance_score for r in sorted(field_rerank.results, key=lambda x: x.index)
+    ]
+    kwargs_scores = [
+        r.relevance_score for r in sorted(kwargs_rerank.results, key=lambda x: x.index)
+    ]
+    assert field_scores == pytest.approx(kwargs_scores)
@@ -3,6 +3,4 @@
 model_name: "openai/gpt-oss-20b"
 metric_threshold: 0.568
 reasoning_effort: "low"
-server_args: "--tensor-parallel-size 2"
-env:
-  VLLM_USE_FLASHINFER_MOE_MXFP4_MXFP8_CUTLASS: "1"
+server_args: "--tensor-parallel-size 2 --moe-backend flashinfer_cutlass --quantization-config.moe.activation mxfp8"
@@ -3,6 +3,4 @@
 model_name: "openai/gpt-oss-20b"
 metric_threshold: 0.568
 reasoning_effort: "low"
-server_args: "--tensor-parallel-size 2"
-env:
-  VLLM_USE_FLASHINFER_MOE_MXFP4_MXFP8: "1"
+server_args: "--tensor-parallel-size 2 --quantization-config.moe.activation mxfp8"
@@ -49,15 +49,14 @@ def get_attn_isa(
    block_size: int | None = None,
    dtype: torch.dtype | None = None,
 ):
-    if block_size and dtype:
-        return _get_attn_isa(dtype, block_size)
-    else:
-        if current_platform.get_cpu_architecture() == CpuArchEnum.ARM:
-            return "neon"
-        elif torch.cpu._is_amx_tile_supported():
-            return "amx"
-        else:
-            return "vec"
+    # Delegate to _get_attn_isa so the fallback path applies the same arch
+    # gating (e.g. RISC-V RVV is only chosen when the build's hardcoded
+    # VLEN=128 kernel is actually present; on VLEN=256 / scalar hosts it
+    # correctly falls through to vec/vec16).
+    return _get_attn_isa(
+        dtype if dtype is not None else torch.bfloat16,
+        block_size if block_size else 32,
+    )


 # rand number generation takes too much time, cache rand tensors
@@ -583,6 +582,52 @@ def test_varlen_with_paged_kv_normal_neon(
    )


+@pytest.mark.parametrize("kv_cache_dtype", ["auto", "fp8_e4m3"])
+@pytest.mark.parametrize("seq_lens", SEQ_LENS)
+@pytest.mark.parametrize("num_heads", NUM_HEADS)
+@pytest.mark.parametrize("head_size", HEAD_SIZES)
+@pytest.mark.parametrize("block_size", [96, 128])
+@pytest.mark.parametrize("sliding_window", SLIDING_WINDOWS)
+@pytest.mark.parametrize("dtype", QTYPES)
+@pytest.mark.parametrize("soft_cap", [None])
+@pytest.mark.parametrize("num_blocks", NUM_BLOCKS)
+@pytest.mark.parametrize("use_alibi", [False])
+@pytest.mark.parametrize("use_sink", [False])
+@pytest.mark.parametrize("isa", ["rvv"])
+@pytest.mark.skipif(
+    current_platform.get_cpu_architecture() != CpuArchEnum.RISCV,
+    reason="Not a RISC-V CPU.",
+)
+def test_varlen_with_paged_kv_normal_rvv(
+    seq_lens: list[tuple[int, int]],
+    num_heads: tuple[int, int],
+    head_size: int,
+    sliding_window: int | None,
+    dtype: torch.dtype,
+    block_size: int,
+    soft_cap: float | None,
+    num_blocks: int,
+    use_alibi: bool,
+    use_sink: bool,
+    isa: str,
+    kv_cache_dtype: str,
+) -> None:
+    varlen_with_paged_kv(
+        seq_lens=seq_lens,
+        num_heads=num_heads,
+        head_size=head_size,
+        sliding_window=sliding_window,
+        dtype=dtype,
+        block_size=block_size,
+        soft_cap=soft_cap,
+        num_blocks=num_blocks,
+        use_alibi=use_alibi,
+        use_sink=use_sink,
+        isa=isa,
+        kv_cache_dtype=kv_cache_dtype,
+    )
+
+
@pytest.mark.parametrize("kv_cache_dtype", ["auto", "fp8_e4m3"])
@pytest.mark.parametrize("seq_lens", SEQ_LENS)
@pytest.mark.parametrize("num_heads", NUM_HEADS)
@@ -66,6 +66,24 @@ def reference(x: torch.Tensor, use_ue8m0: bool) -> tuple[torch.Tensor, torch.Ten
    return reference_quant(ref_act_out, use_ue8m0)


+def reference_with_clamp(
+    x: torch.Tensor, use_ue8m0: bool, clamp_limit: float
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Pre-clamp inputs (gate from above, up symmetric) at the input dtype to
+    match the C++ compute() template, then run the standard silu_and_mul +
+    quant reference."""
+    N_2 = x.size(1) // 2
+    dtype = x.dtype
+    gate = x[..., :N_2].to(torch.float32).clamp(max=clamp_limit).to(dtype)
+    up = (
+        x[..., N_2:]
+        .to(torch.float32)
+        .clamp(min=-clamp_limit, max=clamp_limit)
+        .to(dtype)
+    )
+    return reference(torch.cat([gate, up], dim=-1), use_ue8m0)
+
+
@pytest.mark.parametrize("T", [128, 256, 512])
@pytest.mark.parametrize("N", [128 * 2, 256 * 2, 768 * 2, 2048 * 2, 7168 * 2])
@pytest.mark.skipif(
@@ -89,3 +107,32 @@ def test_silu_mul_fp8_quant_deep_gemm(T: int, N: int):

    torch.testing.assert_close(output.to(torch.float32), ref_output.to(torch.float32))
    torch.testing.assert_close(output_scales, ref_output_scales)
+
+
+@pytest.mark.parametrize("T", [128, 256, 512])
+@pytest.mark.parametrize("N", [128 * 2, 256 * 2, 768 * 2, 2048 * 2, 7168 * 2])
+@pytest.mark.parametrize("clamp_limit", [7.0, 10.0])
+@pytest.mark.skipif(
+    current_platform.is_rocm(),
+    reason="ROCm does not support DeepGemm.",
+)
+def test_silu_mul_fp8_quant_deep_gemm_clamp(T: int, N: int, clamp_limit: float):
+    set_random_seed(42)
+
+    # Use a wide distribution so values routinely exceed both clamp limits and
+    # the clamp branch is actually exercised (uniform [0, 1) inputs would never
+    # trigger it).
+    input = torch.randn((T, N), dtype=torch.bfloat16, device="cuda") * 8.0
+
+    use_ue8m0 = is_deep_gemm_e8m0_used()
+
+    # Test
+    output, output_scales = silu_mul_per_token_group_quant_fp8_colmajor(
+        input, use_ue8m0=use_ue8m0, clamp_limit=clamp_limit
+    )
+
+    # Reference
+    ref_output, ref_output_scales = reference_with_clamp(input, use_ue8m0, clamp_limit)
+
+    torch.testing.assert_close(output.to(torch.float32), ref_output.to(torch.float32))
+    torch.testing.assert_close(output_scales, ref_output_scales)
@@ -3,7 +3,7 @@
 import pytest
 import torch

-import vllm.model_executor.layers.mhc as mhc_ops  # noqa: F401
+import vllm.model_executor.kernels.mhc  # noqa: F401
 from vllm.platforms import current_platform
 from vllm.utils.torch_utils import set_random_seed

@@ -74,6 +74,23 @@ def mhc_post_ref(
    return (x.float().unsqueeze(-2) * post_layer_mix + term2).bfloat16()


+def hc_head_ref(
+    residual: torch.Tensor,
+    fn: torch.Tensor,
+    hc_scale: torch.Tensor,
+    hc_base: torch.Tensor,
+    rms_eps: float,
+    hc_eps: float,
+) -> torch.Tensor:
+    residual_flat = residual.flatten(-2).float()
+    residual_norm = residual_flat * torch.rsqrt(
+        residual_flat.square().mean(dim=-1, keepdim=True) + rms_eps
+    )
+    pre_mix = torch.nn.functional.linear(residual_norm, fn)
+    pre_mix = torch.sigmoid(pre_mix * hc_scale + hc_base) + hc_eps
+    return torch.sum(pre_mix.unsqueeze(-1) * residual.float(), dim=-2).bfloat16()
+
+
@pytest.mark.skipif(
    not current_platform.is_cuda(),
    reason="CUDA required",
@@ -121,7 +138,7 @@ def test_mhc_fused_post_pre(num_tokens, hidden_size, hc_mult):

    residual_ref, post_mix_ref, res_mix_ref, layer_input_ref = run_ref()

-    residual, post_mix, res_mix, x = torch.ops.vllm.mhc_fused_post_pre(
+    residual, post_mix, res_mix, x = torch.ops.vllm.mhc_fused_post_pre_tilelang(
        x,
        residual,
        post_layer_mix,
@@ -140,3 +157,42 @@ def test_mhc_fused_post_pre(num_tokens, hidden_size, hc_mult):
    torch.testing.assert_close(post_mix, post_mix_ref, atol=1e-2, rtol=1e-2)
    torch.testing.assert_close(res_mix, res_mix_ref, atol=1e-2, rtol=1e-2)
    torch.testing.assert_close(x, layer_input_ref, atol=1e-2, rtol=1e-2)
+
+
+@pytest.mark.skipif(
+    not current_platform.is_rocm(),
+    reason="ROCm required",
+)
+@pytest.mark.parametrize("num_tokens", [1, 4, 8, 128])
+@pytest.mark.parametrize("hidden_size", [4096, 7168])
+@pytest.mark.parametrize("hc_mult", [4])
+def test_hc_head_triton(num_tokens, hidden_size, hc_mult):
+    torch.set_default_device(DEVICE)
+    set_random_seed(0)
+
+    residual = torch.randn((num_tokens, hc_mult, hidden_size), dtype=torch.bfloat16)
+    fn = torch.randn((hc_mult, hc_mult * hidden_size), dtype=torch.float32) * 1e-4
+    hc_scale = torch.randn((1,), dtype=torch.float32) * 0.1
+    hc_base = torch.randn((hc_mult,), dtype=torch.float32) * 0.1
+    rms_eps = hc_eps = 1e-6
+
+    out = torch.empty((num_tokens, hidden_size), dtype=torch.bfloat16)
+    out.fill_(float("nan"))
+
+    result = torch.ops.vllm.hc_head_triton(
+        residual,
+        fn,
+        hc_scale,
+        hc_base,
+        out,
+        hidden_size,
+        rms_eps,
+        hc_eps,
+        hc_mult,
+    )
+
+    assert result is None
+    assert not torch.isnan(out).any()
+
+    out_ref = hc_head_ref(residual, fn, hc_scale, hc_base, rms_eps, hc_eps)
+    torch.testing.assert_close(out, out_ref, atol=5e-2, rtol=1e-2)
@@ -1,162 +1,249 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import types
+from types import SimpleNamespace
+from unittest.mock import patch

 import pytest
 import torch

+from vllm.distributed.eplb.eplb_state import EplbLayerState
+from vllm.model_executor.layers.fused_moe.config import RoutingMethodType
+from vllm.model_executor.layers.fused_moe.routed_experts_capturer import (
+    RoutedExpertsCapturer,
+)
+from vllm.model_executor.layers.fused_moe.router.base_router import BaseRouter
+
 pytestmark = pytest.mark.cpu_test

+_REC_MODULE = "vllm.model_executor.layers.fused_moe.routed_experts_capturer"

-def test_bind_routing_capture_to_model_sets_layer_view(monkeypatch):
-    import vllm.model_executor.layers.fused_moe.layer as fused_moe_layer
-    import vllm.model_executor.layers.fused_moe.routed_experts_capturer as rec_mod

-    class _DummyMoEConfig:
-        is_sequence_parallel = False
-        dp_size = 1
+def _capturer_with_buffer(
+    *,
+    max_tokens: int = 8,
+    num_layers: int = 4,
+    num_experts_per_tok: int = 2,
+    dp_rank: int = 0,
+    tp_size: int = 1,
+) -> RoutedExpertsCapturer:
+    # Bypass __init__ so the test can use a CPU buffer and skip the
+    # VllmConfig dependency. The CUDA device-tensor allocation in the
+    # real constructor is not what we are exercising here.
+    c = RoutedExpertsCapturer.__new__(RoutedExpertsCapturer)
+    c.dp_rank = dp_rank
+    c.tp_size = tp_size
+    c.device_buffer = torch.full(
+        (max_tokens, num_layers, num_experts_per_tok),
+        -1,
+        dtype=torch.int32,
+    )
+    return c

-    class _DummyQuantMethod:
-        supports_internal_mk = True
+
+class DummyRouter(BaseRouter):
+    @property
+    def routing_method_type(self) -> RoutingMethodType:
+        return RoutingMethodType.FUSED_TOPK
+
+    def _compute_routing(
+        self, hidden_states, router_logits, indices_type, *, input_ids=None
+    ):
+        topk_ids = torch.tensor([[1, 2], [3, 4]], dtype=torch.int64)
+        topk_weights = torch.ones_like(topk_ids, dtype=torch.float32)
+        return topk_weights, topk_ids
+
+    def _apply_eplb_mapping(self, topk_ids: torch.Tensor) -> torch.Tensor:
+        # Make mapping observable without requiring CUDA EPLB path.
+        return topk_ids + 10
+
+
+def _make_router(eplb_state: EplbLayerState | None = None) -> DummyRouter:
+    return DummyRouter(
+        top_k=2,
+        global_num_experts=16,
+        eplb_state=eplb_state,
+        indices_type_getter=None,
+    )
+
+
+def test_base_router_capture_pre_eplb_mapping():
+    router = _make_router()
+    captured = []
+
+    def capture_fn(ids):
+        captured.append(ids.clone())
+
+    router.set_capture_fn(capture_fn)
+    topk_weights, topk_ids = router.select_experts(
+        hidden_states=torch.empty(1),
+        router_logits=torch.empty(1),
+    )
+
+    assert topk_weights.shape == topk_ids.shape
+    assert len(captured) == 1
+    assert torch.equal(captured[0], torch.tensor([[1, 2], [3, 4]]))
+    assert torch.equal(topk_ids, torch.tensor([[11, 12], [13, 14]]))
+
+
+def test_base_router_capture_with_eplb_enabled():
+    eplb_state = EplbLayerState()
+    eplb_state.expert_load_view = torch.zeros(32, dtype=torch.int64)
+    eplb_state.logical_to_physical_map = torch.arange(32).view(32, 1)
+    eplb_state.logical_replica_count = torch.ones(32, dtype=torch.int64)
+    eplb_state.should_record_tensor = torch.ones((), dtype=torch.bool)
+    router = _make_router(eplb_state=eplb_state)
+
+    captured = []
+
+    def capture_fn(ids):
+        captured.append(ids.clone())
+
+    router.set_capture_fn(capture_fn)
+    _, topk_ids = router.select_experts(
+        hidden_states=torch.empty(1),
+        router_logits=torch.empty(1),
+    )
+
+    assert len(captured) == 1
+    # Capture should see logical ids pre-EPLB mapping.
+    assert torch.equal(captured[0], torch.tensor([[1, 2], [3, 4]]))
+    # Our DummyRouter mapping adds +10.
+    assert torch.equal(topk_ids, torch.tensor([[11, 12], [13, 14]]))
+
+
+def test_gpu_model_runner_binds_router_capture(monkeypatch):
+    from vllm.v1.worker import gpu_model_runner as gmr

    class DummyFusedMoE:
-        _routing_replay_out: torch.Tensor
+        def __init__(self):
+            self.layer_id = 7
+            self.router = _make_router()

-        def __init__(self, moe_layer_id):
-            self.moe_layer_id = moe_layer_id
-            self.moe_config = _DummyMoEConfig()
-            self.quant_method = _DummyQuantMethod()
+    class DummyCapturer:
+        def __init__(self):
+            self.calls = []
+
+        def capture(self, layer_id, topk_ids):
+            self.calls.append((layer_id, topk_ids))
+
+    dummy_module = DummyFusedMoE()
+
+    # Patch the runtime import inside _bind_routed_experts_capturer.
+    import vllm.model_executor.layers.fused_moe.layer as fused_moe_layer

    monkeypatch.setattr(fused_moe_layer, "FusedMoE", DummyFusedMoE)

-    num_layers, num_tokens, top_k = 4, 8, 2
-    buffer = torch.zeros((num_layers, num_tokens, top_k), dtype=torch.int16)
+    dummy_self = types.SimpleNamespace(
+        compilation_config=types.SimpleNamespace(
+            static_forward_context={"dummy": dummy_module}
+        )
+    )

-    class DummyDeviceCache:
-        def __init__(self, buf):
-            self.buffer = buf
+    capturer = DummyCapturer()
+    gmr.GPUModelRunner._bind_routed_experts_capturer(dummy_self, capturer)
+
+    assert dummy_module.router.capture_fn is not None
+    dummy_module.router.capture_fn(torch.tensor([[5, 6]]))
+
+    assert len(capturer.calls) == 1
+    layer_id, topk_ids = capturer.calls[0]
+    assert layer_id == 7
+    assert torch.equal(topk_ids, torch.tensor([[5, 6]]))
+
+
+def test_gpu_model_runner_binding_stage(monkeypatch):
+    from vllm.v1.worker import gpu_model_runner as gmr
+
+    class DummyFusedMoE:
+        def __init__(self):
+            self.layer_id = 11
+            self.router = _make_router()

    class DummyCapturer:
-        def get_device_cache(self):
-            return DummyDeviceCache(buffer)
+        def __init__(self):
+            self.calls = []

-    monkeypatch.setattr(rec_mod, "get_global_experts_capturer", lambda: DummyCapturer())
+        def capture(self, layer_id, topk_ids):
+            self.calls.append((layer_id, topk_ids))

-    m0 = DummyFusedMoE(moe_layer_id=0)
-    m2 = DummyFusedMoE(moe_layer_id=2)
+    dummy_module = DummyFusedMoE()

-    class DummyModel:
-        def modules(self):
-            return iter([m0, m2])
+    import vllm.model_executor.layers.fused_moe.layer as fused_moe_layer

-    rec_mod.bind_routing_capture_to_model(DummyModel())
+    monkeypatch.setattr(fused_moe_layer, "FusedMoE", DummyFusedMoE)

-    assert torch.equal(m0._routing_replay_out, buffer[0])
-    assert torch.equal(m2._routing_replay_out, buffer[2])
-
-
-def test_bind_routing_capture_to_model_noop_when_disabled(monkeypatch):
-    import vllm.model_executor.layers.fused_moe.routed_experts_capturer as rec_mod
-
-    class DummyCapturer:
-        def get_device_cache(self):
-            return None
-
-    monkeypatch.setattr(rec_mod, "get_global_experts_capturer", lambda: DummyCapturer())
-
-    class DummyModel:
-        def modules(self):
-            return iter([])
-
-    rec_mod.bind_routing_capture_to_model(DummyModel())
-
-
-# =========================================================================
-# Tests for device-cache routing replay architecture
-# =========================================================================
-
-
-class TestRoutedExpertsDeviceCache:
-    """Tests for _RoutedExpertsDeviceCache (GPU buffer for routing data)."""
-
-    def test_allocation_shape_and_dtype(self):
-        """Device cache allocates (L, N, K) int16 buffer."""
-        from vllm.model_executor.layers.fused_moe.routed_experts_capturer import (
-            _RoutedExpertsDeviceCache,
+    dummy_self = types.SimpleNamespace(
+        compilation_config=types.SimpleNamespace(
+            static_forward_context={"dummy": dummy_module}
        )
+    )

-        cache = _RoutedExpertsDeviceCache(
-            num_hidden_layers=40,
-            max_num_batched_tokens=8192,
-            num_experts_per_tok=8,
-            device="cpu",
-        )
-        assert cache.buffer.shape == (40, 8192, 8)
-        assert cache.buffer.dtype == torch.int16
+    # Before binding, no capture hook.
+    assert dummy_module.router.capture_fn is None

-    def test_per_layer_view_is_contiguous(self):
-        """buffer[layer_id] gives contiguous (N, K) view for FlashInfer."""
-        from vllm.model_executor.layers.fused_moe.routed_experts_capturer import (
-            _RoutedExpertsDeviceCache,
-        )
+    capturer = DummyCapturer()
+    gmr.GPUModelRunner._bind_routed_experts_capturer(dummy_self, capturer)

-        cache = _RoutedExpertsDeviceCache(
-            num_hidden_layers=40,
-            max_num_batched_tokens=8192,
-            num_experts_per_tok=8,
-            device="cpu",
-        )
-        layer_view = cache.buffer[0]
-        assert layer_view.is_contiguous()
-        assert layer_view.shape == (8192, 8)
+    # After binding, hook should exist and be callable.
+    assert callable(dummy_module.router.capture_fn)
+    dummy_module.router.capture_fn(torch.tensor([[9, 10]]))
+    assert len(capturer.calls) == 1


-class TestRoutedExpertsHostCache:
-    """Tests for _RoutedExpertsHostCache (per-request numpy buffer)."""
+def test_routed_experts_capturer_single_dp_no_metadata():
+    """dp_metadata is None: capture writes the full topk_ids rows."""
+    capturer = _capturer_with_buffer(dp_rank=0)
+    topk = torch.tensor([[1, 2], [3, 4], [5, 6]], dtype=torch.int32)
+    ctx = SimpleNamespace(dp_metadata=None)
+    with patch(f"{_REC_MODULE}.get_forward_context", return_value=ctx):
+        capturer.capture(layer_id=0, topk_ids=topk)
+    assert torch.equal(capturer.device_buffer[:3, 0, :], topk)
+    assert capturer.device_buffer[3, 0, 0].item() == -1

-    def test_sentinel_initialization(self):
-        """Host cache initializes with zeros by default."""
-        import numpy as np

-        from vllm.model_executor.layers.fused_moe.routed_experts_capturer import (
-            _RoutedExpertsHostCache,
-        )
+def test_routed_experts_capturer_dp_naive_concatenated_all_ranks():
+    """n == sum(num_tokens_dp): slice this rank's segment from concatenated topk."""
+    capturer = _capturer_with_buffer(dp_rank=1)
+    num_tokens_dp = torch.tensor([2, 3], dtype=torch.int32)
+    ctx = SimpleNamespace(
+        dp_metadata=SimpleNamespace(num_tokens_across_dp_cpu=num_tokens_dp)
+    )
+    # Concatenated order: rank0 rows then rank1 rows.
+    topk = torch.tensor(
+        [[0, 1], [2, 3], [10, 11], [12, 13], [14, 15]], dtype=torch.int32
+    )
+    with patch(f"{_REC_MODULE}.get_forward_context", return_value=ctx):
+        capturer.capture(layer_id=0, topk_ids=topk)
+    want = topk[2:5]
+    assert torch.equal(capturer.device_buffer[:3, 0, :], want)

-        cache = _RoutedExpertsHostCache(
-            num_hidden_layers=40,
-            num_experts_per_tok=8,
-            max_model_len=1024,
-        )
-        buf = cache.get_or_grow_buffer("req1", max_pos=100)
-        assert buf.dtype == np.int16
-        assert (buf == 0).all(), "Host cache must initialize with zeros"

-    def test_grow_preserves_existing_data(self):
-        """Growing the buffer preserves previously written data."""
-        from vllm.model_executor.layers.fused_moe.routed_experts_capturer import (
-            _RoutedExpertsHostCache,
-        )
+def test_routed_experts_capturer_dp_modular_local_tokens():
+    """n == token_num_per_dp: topk is already local to this DP rank."""
+    capturer = _capturer_with_buffer(dp_rank=1)
+    num_tokens_dp = torch.tensor([2, 3], dtype=torch.int32)
+    ctx = SimpleNamespace(
+        dp_metadata=SimpleNamespace(num_tokens_across_dp_cpu=num_tokens_dp)
+    )
+    topk = torch.tensor([[10, 11], [12, 13], [14, 15]], dtype=torch.int32)
+    with patch(f"{_REC_MODULE}.get_forward_context", return_value=ctx):
+        capturer.capture(layer_id=0, topk_ids=topk)
+    assert torch.equal(capturer.device_buffer[:3, 0, :], topk)

-        cache = _RoutedExpertsHostCache(
-            num_hidden_layers=40,
-            num_experts_per_tok=8,
-            max_model_len=1024,
-        )
-        buf = cache.get_or_grow_buffer("req1", max_pos=50)
-        buf[0, 0, 0] = 42
-        buf2 = cache.get_or_grow_buffer("req1", max_pos=200)
-        assert buf2[0, 0, 0] == 42, "Data lost during buffer grow"

-    def test_free_request_removes_buffer(self):
-        """Freeing a request removes its buffer."""
-        from vllm.model_executor.layers.fused_moe.routed_experts_capturer import (
-            _RoutedExpertsHostCache,
-        )
-
-        cache = _RoutedExpertsHostCache(
-            num_hidden_layers=40,
-            num_experts_per_tok=8,
-            max_model_len=1024,
-        )
-        cache.get_or_grow_buffer("req1", max_pos=50)
-        cache.free_request("req1")
-        assert cache.get_buffer("req1") is None
+def test_routed_experts_capturer_dp_unexpected_batch_raises():
+    """Mismatch between topk batch dim and DP layout: fail fast."""
+    capturer = _capturer_with_buffer(dp_rank=0)
+    num_tokens_dp = torch.tensor([2, 3], dtype=torch.int32)
+    ctx = SimpleNamespace(
+        dp_metadata=SimpleNamespace(num_tokens_across_dp_cpu=num_tokens_dp)
+    )
+    # total=5, local=2: n=1 matches neither naive (5) nor modular (2).
+    topk = torch.tensor([[1, 2]], dtype=torch.int32)
+    with (
+        patch(f"{_REC_MODULE}.get_forward_context", return_value=ctx),
+        pytest.raises(AssertionError, match="unexpected topk_ids batch dim"),
+    ):
+        capturer.capture(layer_id=0, topk_ids=topk)
+    assert capturer.device_buffer[0, 0, 0].item() == -1
@@ -316,7 +316,15 @@ VLM_TEST_SETTINGS = {
        stop_str=["<|im_end|>"],
        image_size_factors=[(0.10, 0.15)],
        max_tokens=64,
-        marks=[large_gpu_mark(min_gb=64)],
+        marks=[
+            pytest.mark.skip(
+                reason="Aria needs to update for latest transformers, "
+                "must have a vision_processor.py."
+                "An issue has been filed:"
+                "https://huggingface.co/rhymes-ai/Aria/discussions/23"
+            ),
+            large_gpu_mark(min_gb=64),
+        ],
    ),
    "aya_vision": VLMTestInfo(
        models=["CohereLabs/aya-vision-8b"],
@@ -468,7 +476,14 @@ VLM_TEST_SETTINGS = {
        max_tokens=8,
        num_logprobs=10,
        auto_cls=AutoModelForCausalLM,
-        marks=[large_gpu_mark(min_gb=32)],
+        marks=[
+            pytest.mark.skip(
+                reason="The code for this model has a bug."
+                "Please see the issue here:"
+                "https://huggingface.co/zai-org/glm-4v-9b/discussions/46."
+            ),
+            large_gpu_mark(min_gb=32),
+        ],
    ),
    "glm4_1v": VLMTestInfo(
        models=["zai-org/GLM-4.1V-9B-Thinking"],
@@ -513,7 +528,14 @@ VLM_TEST_SETTINGS = {
        num_logprobs=10,
        image_size_factors=[(0.25,), (0.25, 0.25, 0.25), (0.25, 0.2, 0.15)],
        auto_cls=AutoModelForImageTextToText,
-        marks=[large_gpu_mark(min_gb=32)],
+        marks=[
+            pytest.mark.skip(
+                reason="This test fails on both AMD and NV"
+                "hardware. please see the issue:"
+                "https://github.com/vllm-project/vllm/issues/42016"
+            ),
+            large_gpu_mark(min_gb=32),
+        ],
    ),
    "granite4_vision": VLMTestInfo(
        models=["ibm-granite/granite-vision-4.1-4b"],
@@ -14,7 +14,9 @@ from vllm.tokenizers import TokenizerLike
 from ....conftest import HfRunner, PromptImageInput, VllmRunner

 IMAGE = ImageAsset("paper-11").pil_image_ext(ext="png").convert("RGB")
-PROMPT = "</s><s><predict_bbox><predict_classes><output_markdown>"
+PROMPT = (
+    "</s><s><predict_bbox><predict_classes><output_markdown><predict_no_text_in_pic>"
+)


 class DummyLogprobs(dict[int, Logprob]):
@@ -85,7 +87,7 @@ def run_test(
                max_tokens,
                num_logprobs=num_logprobs,
                images=images,
-                use_cache=False,  # HF Nemotron Parse crashes here without this
+                tokenization_kwargs={"add_special_tokens": False},
            )
            for prompts, images in inputs
        ]
@@ -103,11 +105,7 @@ def run_test(
        )


-@pytest.mark.skip(
-    reason="Model's custom MBart decoder has head count mismatch with "
-    "transformers v5's GQA-aware cross-attention (8 vs 16 heads)"
-)
-@pytest.mark.parametrize("model", ["nvidia/NVIDIA-Nemotron-Parse-v1.1"])
+@pytest.mark.parametrize("model", ["nvidia/NVIDIA-Nemotron-Parse-v1.2"])
@pytest.mark.parametrize("dtype", ["bfloat16"])
@pytest.mark.parametrize("num_logprobs", [5])
 def test_models(
@@ -78,6 +78,18 @@ MODEL_CONFIGS: dict[str, VitCudagraphTestConfig] = {
        needs_video_metadata=True,
        marks=[pytest.mark.core_model],
    ),
+    "qwen2_vl": VitCudagraphTestConfig(
+        model="Qwen/Qwen2-VL-2B-Instruct",
+        image_prompt=qwen_vl_chat_template(
+            "<|vision_start|><|image_pad|><|vision_end|>What is in this image?"
+        ),
+        video_prompt=qwen_vl_chat_template(
+            "<|vision_start|><|video_pad|><|vision_end|>"
+            "Describe this video in one sentence."
+        ),
+        needs_video_metadata=False,
+        marks=[pytest.mark.core_model],
+    ),
 }


@@ -1,10 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Compares the outputs of gptq vs gptq_marlin.
+"""Tests AutoGPTQ (GPTQ with Marlin kernels) output correctness.

-Note: GPTQ and Marlin do not have bitwise correctness.
-As a result, in this test, we just confirm that the top selected tokens of the
-Marlin/GPTQ models are in the top 5 selections of each other.
 Note: Marlin internally uses locks to synchronize the threads. This can
 result in very slight nondeterminism for Marlin. As a result, we re-run the test
 up to 3 times to see if we pass.
@@ -36,10 +33,10 @@ MODELS = [

@pytest.mark.flaky(reruns=3)
@pytest.mark.skipif(
-    not is_quant_method_supported("gptq_marlin")
+    not is_quant_method_supported("auto_gptq")
    or current_platform.is_rocm()
    or not current_platform.is_cuda(),
-    reason="gptq_marlin is not supported on this GPU type.",
+    reason="auto_gptq is not supported on this GPU type.",
 )
@pytest.mark.parametrize("model", MODELS)
@pytest.mark.parametrize("dtype", ["half", "bfloat16"])
@@ -988,6 +988,11 @@ _MULTIMODAL_EXAMPLE_MODELS = {
        "internlm/Intern-S1-Pro",
        trust_remote_code=True,
    ),
+    "InternS2PreviewForConditionalGeneration": _HfExamplesInfo(
+        "internlm/Intern-S2-Preview",
+        trust_remote_code=True,
+        is_available_online=False,
+    ),
    "InternVLChatModel": _HfExamplesInfo(
        "OpenGVLab/InternVL2-1B",
        extras={
@@ -1334,7 +1339,16 @@ _MULTIMODAL_EXAMPLE_MODELS = {
    ),
    "RForConditionalGeneration": _HfExamplesInfo("YannQi/R-4B", trust_remote_code=True),
    "SkyworkR1VChatModel": _HfExamplesInfo(
-        "Skywork/Skywork-R1V-38B", trust_remote_code=True
+        "Skywork/Skywork-R1V-38B",
+        trust_remote_code=True,
+        max_transformers_version="4.57",
+        transformers_version_reason={
+            "hf": (
+                "SkyworkChatModel.__init__ does not call self.post_init(), so "
+                "`all_tied_weights_keys` is never set; Transformers v5 requires "
+                "this attribute in _move_missing_keys_from_meta_to_device."
+            )
+        },
    ),
    "SmolVLMForConditionalGeneration": _HfExamplesInfo(
        "HuggingFaceTB/SmolVLM2-2.2B-Instruct"
@@ -1379,7 +1393,7 @@ _MULTIMODAL_EXAMPLE_MODELS = {
        "CohereLabs/cohere-transcribe-03-2026", trust_remote_code=True
    ),
    "NemotronParseForConditionalGeneration": _HfExamplesInfo(
-        "nvidia/NVIDIA-Nemotron-Parse-v1.1", trust_remote_code=True
+        "nvidia/NVIDIA-Nemotron-Parse-v1.2", trust_remote_code=True
    ),
    "WhisperForConditionalGeneration": _HfExamplesInfo(
        "openai/whisper-large-v3-turbo",
@@ -13,7 +13,7 @@ from vllm.multimodal.video import (
    VideoLoader,
 )

-from .utils import create_video_from_image
+from .utils import create_long_gop_video, create_video_from_image

 pytestmark = pytest.mark.cpu_test

@@ -364,6 +364,49 @@ def test_pyav_dynamic_backend_loads_frames(
        assert metadata["video_backend"] == "pyav_dynamic"


+def test_pyav_backend_returns_target_frames_not_keyframes():
+    """Regression test: PyAV must decode forward past the seek keyframe.
+
+    container.seek() snaps backward to the nearest keyframe. With a long GOP
+    (here: one keyframe at frame 0), a decoder that does not advance forward
+    to the target PTS collapses every sampled slot onto the keyframe. This
+    test encodes a per-frame marker on the green channel and verifies the
+    returned frames are distinct, ordered, and match the requested indices.
+    """
+    num_frames = 50
+    num_sampled = 4
+    height, width = 64, 64
+
+    video_bytes = create_long_gop_video(
+        num_frames=num_frames, width=width, height=height
+    )
+
+    loader = VIDEO_LOADER_REGISTRY.load("opencv")
+    frames, metadata = loader.load_bytes(
+        video_bytes, num_frames=num_sampled, backend="pyav"
+    )
+    assert frames.shape == (num_sampled, height, width, 3)
+
+    requested = list(metadata["frames_indices"])
+    assert len(requested) == num_sampled
+
+    actual = [int(f[height // 2, width // 2, 1]) for f in frames]
+
+    assert len(set(actual)) == num_sampled, (
+        f"PyAV returned only {len(set(actual))} distinct frames for "
+        f"{num_sampled} requested indices: markers={actual}, "
+        f"requested={requested}. Keyframe-snap regression."
+    )
+
+    assert actual == sorted(actual), f"Returned frames out of order: markers={actual}"
+
+    for marker, want_idx in zip(actual, requested):
+        assert abs(marker - want_idx) <= 10, (
+            f"Frame mismatch: requested index {want_idx}, "
+            f"got marker {marker} (tolerance ±10)"
+        )
+
+
@pytest.mark.parametrize(
    "loader_key, kwargs, expected_num_frames",
    [
@@ -66,6 +66,43 @@ def create_video_from_image(
    return video_path


+def create_long_gop_video(
+    num_frames: int = 50,
+    fps: int = 30,
+    width: int = 64,
+    height: int = 64,
+) -> bytes:
+    """Encode an H.264 clip with one keyframe and green-channel = frame index.
+
+    The marker lets a test recover which frame the decoder actually returned,
+    independent of any metadata label.
+    """
+    import io
+
+    import av
+
+    buf = io.BytesIO()
+    with av.open(buf, mode="w", format="mp4") as container:
+        stream = container.add_stream("h264", rate=fps)
+        stream.width = width
+        stream.height = height
+        stream.pix_fmt = "yuv420p"
+        stream.codec_context.gop_size = num_frames
+        stream.codec_context.max_b_frames = 0
+        stream.codec_context.options = {
+            "x264-params": (f"scenecut=0:keyint={num_frames}:min-keyint={num_frames}")
+        }
+        for i in range(num_frames):
+            img = np.zeros((height, width, 3), dtype=np.uint8)
+            img[:, :, 1] = i
+            frame = av.VideoFrame.from_ndarray(img, format="rgb24")
+            for packet in stream.encode(frame):
+                container.mux(packet)
+        for packet in stream.encode():
+            container.mux(packet)
+    return buf.getvalue()
+
+
 def cosine_similarity(A: npt.NDArray, B: npt.NDArray, axis: int = -1) -> npt.NDArray:
    """Compute cosine similarity between two vectors."""
    return np.sum(A * B, axis=axis) / (
@@ -0,0 +1,56 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests that the auto_gptq quantization method works correctly.
+
+Run `pytest tests/quantization/test_auto_gptq.py -v -s`.
+"""
+
+import pytest
+import torch
+
+from tests.quantization.utils import is_quant_method_supported
+from vllm.model_executor.layers.quantization.auto_gptq import (
+    AutoGPTQConfig,
+    AutoGPTQLinearMethod,
+)
+
+PROMPT = "On the surface of Mars, we found"
+
+MODELS = [
+    "TheBloke/TinyLlama-1.1B-Chat-v1.0-GPTQ",
+]
+
+
+@pytest.mark.skipif(
+    not is_quant_method_supported("auto_gptq"),
+    reason="auto_gptq is not supported on this GPU type.",
+)
+@pytest.mark.parametrize("model_id", MODELS)
+def test_auto_gptq_quantization_method(vllm_runner, model_id: str, monkeypatch):
+    """Test that quantization='auto_gptq' loads and runs correctly."""
+    monkeypatch.setenv("VLLM_ALLOW_INSECURE_SERIALIZATION", "1")
+
+    with vllm_runner(
+        model_id,
+        dtype=torch.float16,
+        quantization="auto_gptq",
+        max_model_len=2048,
+        enforce_eager=True,
+    ) as llm:
+
+        def check_model(model):
+            for name, submodule in model.named_modules():
+                if name == "model.layers.0.self_attn.qkv_proj":
+                    assert isinstance(submodule.quant_method, AutoGPTQLinearMethod)
+                    break
+
+        llm.apply_model(check_model)
+
+        outputs = llm.generate_greedy([PROMPT], max_tokens=8)
+        assert outputs
+        assert len(outputs[0][1]) > 0
+
+
+def test_auto_gptq_config_get_name():
+    """Test that AutoGPTQConfig.get_name() returns 'auto_gptq'."""
+    assert AutoGPTQConfig.get_name() == "auto_gptq"
@@ -190,13 +190,27 @@ def test_gptoss_mxfp4bf16_moe_flashinfer(monkeypatch: pytest.MonkeyPatch):


 def test_gptoss_mxfp4mxfp8_moe_flashinfer_cutlass(monkeypatch: pytest.MonkeyPatch):
-    monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_MXFP4_MXFP8_CUTLASS", "1")
-    can_initialize("openai/gpt-oss-20b", hf_overrides=HF_OVERRIDE_TEXT)
+    can_initialize(
+        "openai/gpt-oss-20b",
+        hf_overrides=HF_OVERRIDE_TEXT,
+        extra_args=[
+            "--moe-backend",
+            "flashinfer_cutlass",
+            "--quantization-config.moe.activation",
+            "mxfp8",
+        ],
+    )


 def test_gptoss_mxfp4mxfp8_moe_flashinfer_trtllm(monkeypatch: pytest.MonkeyPatch):
-    monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_MXFP4_MXFP8", "1")
-    can_initialize("openai/gpt-oss-20b", hf_overrides=HF_OVERRIDE_TEXT)
+    can_initialize(
+        "openai/gpt-oss-20b",
+        hf_overrides=HF_OVERRIDE_TEXT,
+        extra_args=[
+            "--quantization-config.moe.activation",
+            "mxfp8",
+        ],
+    )


 def test_gptoss_eager(monkeypatch: pytest.MonkeyPatch):
@@ -24,31 +24,23 @@ MODEL_ARG_EXPTYPES = [
    # AUTOGPTQ
    # compat: autogptq <=0.7.1 is_marlin_format: bool
    # Model Serialized in Exllama Format.
-    (
-        "TheBloke/Llama-2-7B-Chat-GPTQ",
-        None,
-        "gptq_marlin" if current_platform.is_cuda() else "gptq",
-    ),
+    ("TheBloke/Llama-2-7B-Chat-GPTQ", None, "auto_gptq"),
    (
        "TheBloke/Llama-2-7B-Chat-GPTQ",
        "marlin",
-        "gptq_marlin" if current_platform.is_cuda() else "ERROR",
+        "auto_gptq" if current_platform.is_cuda() else "ERROR",
    ),
-    ("TheBloke/Llama-2-7B-Chat-GPTQ", "gptq", "gptq"),
+    ("TheBloke/Llama-2-7B-Chat-GPTQ", "gptq", "auto_gptq"),
    ("TheBloke/Llama-2-7B-Chat-GPTQ", "awq", "ERROR"),
    # compat: autogptq >=0.8.0 use checkpoint_format: str
    # Model Serialized in Exllama Format.
-    (
-        "LnL-AI/TinyLlama-1.1B-Chat-v1.0-GPTQ-4bit",
-        None,
-        "gptq_marlin" if current_platform.is_cuda() else "gptq",
-    ),
+    ("LnL-AI/TinyLlama-1.1B-Chat-v1.0-GPTQ-4bit", None, "auto_gptq"),
    (
        "LnL-AI/TinyLlama-1.1B-Chat-v1.0-GPTQ-4bit",
        "marlin",
-        "gptq_marlin" if current_platform.is_cuda() else "ERROR",
+        "auto_gptq" if current_platform.is_cuda() else "ERROR",
    ),
-    ("LnL-AI/TinyLlama-1.1B-Chat-v1.0-GPTQ-4bit", "gptq", "gptq"),
+    ("LnL-AI/TinyLlama-1.1B-Chat-v1.0-GPTQ-4bit", "gptq", "auto_gptq"),
    ("LnL-AI/TinyLlama-1.1B-Chat-v1.0-GPTQ-4bit", "awq", "ERROR"),
    # AUTOAWQ
    (
@@ -3,46 +3,36 @@
 """Tests whether gptq models with dynamic quantized can be loaded.

 Run `pytest tests/quantization/test_gptq_dynamic.py --forked`.
+
+Note: Only symmetric GPTQ models are supported after consolidation to Marlin.
 """

 import pytest
 import torch

 from vllm.model_executor.layers.linear import UnquantizedLinearMethod
-from vllm.model_executor.layers.quantization.gptq import GPTQLinearMethod
-from vllm.model_executor.layers.quantization.gptq_marlin import GPTQMarlinLinearMethod
+from vllm.model_executor.layers.quantization.auto_gptq import AutoGPTQLinearMethod
 from vllm.model_executor.layers.quantization.utils.gptq_utils import (
    get_dynamic_override,
 )
-from vllm.platforms import current_platform

 PROMPT = "On the surface of Mars, we found"

 # The first layer is quantized using bits=4, group_size=128
 # The second layer is quantized using bits=8, group_size=32
 # All other layers (layer index >= 2) are not quantized
-MODEL_QUANT = [
-    (
-        "ModelCloud/Qwen1.5-1.8B-Chat-GPTQ-4bits-dynamic-cfg-with-lm_head-symTrue",
-        current_platform.is_cuda(),
-    ),
-    (
-        "ModelCloud/Qwen1.5-1.8B-Chat-GPTQ-4bits-dynamic-cfg-with-lm_head-symFalse",
-        False,
-    ),
+# Note: Only symmetric models are supported with Marlin kernels
+MODELS = [
+    "ModelCloud/Qwen1.5-1.8B-Chat-GPTQ-4bits-dynamic-cfg-with-lm_head-symTrue",
 ]


-@pytest.mark.parametrize("model_id, use_marlin_kernel", MODEL_QUANT)
-def test_gptq_with_dynamic(
-    vllm_runner, model_id: str, use_marlin_kernel: bool, monkeypatch
-):
+@pytest.mark.parametrize("model_id", MODELS)
+def test_gptq_with_dynamic(vllm_runner, model_id: str, monkeypatch):
    # `LLM.apply_model` requires pickling a function.
    monkeypatch.setenv("VLLM_ALLOW_INSECURE_SERIALIZATION", "1")

-    linear_method_cls = (
-        GPTQMarlinLinearMethod if use_marlin_kernel else (GPTQLinearMethod)
-    )
+    linear_method_cls = AutoGPTQLinearMethod

    with vllm_runner(
        model_id, dtype=torch.float16, max_model_len=2048, enforce_eager=True
@@ -3,6 +3,9 @@
 """Tests whether vllm correctly load and run gptq_v2 format checkpoints.

 Run `pytest tests/quantization/test_gptq_v2.py --forked`.
+
+Note: 2/3-bit GPTQ models are no longer supported after the consolidation
+to Marlin kernels. Only 4/8-bit symmetric GPTQ models are supported.
 """

 import pytest
@@ -10,9 +13,10 @@ import torch
 from transformers import AutoTokenizer

 from vllm import SamplingParams
-from vllm.model_executor.layers.quantization.gptq import GPTQLinearMethod
+from vllm.model_executor.layers.quantization.auto_gptq import AutoGPTQLinearMethod

 # A dummy small model quantized by GPTQModel, stored in GPTQ v2 format
+# Note: This is a 2-bit model which is no longer supported with Marlin kernels
 MODELS = ["XXXXyu/Qwen3-1.7B-w2g64-gptq_v2"]

 # Generate multiple sequences for testing, because an 1.7B 2-bit model
@@ -20,27 +24,19 @@ MODELS = ["XXXXyu/Qwen3-1.7B-w2g64-gptq_v2"]
 N_SEQ = 5


+@pytest.mark.skip(reason="2-bit GPTQ is no longer supported after Marlin consolidation")
@pytest.mark.parametrize("model_id", MODELS)
 def test_model_load(vllm_runner, model_id, monkeypatch):
    # `LLM.apply_model` requires pickling a function.
    monkeypatch.setenv("VLLM_ALLOW_INSECURE_SERIALIZATION", "1")

-    # Only check the default GPTQ linear method (used for 2/3-bit models).
-    # 4/8-bit linear methods like Marlin already support gptq_v2.
-    linear_method_cls = GPTQLinearMethod
-
    with vllm_runner(model_id, dtype=torch.float16, max_model_len=512) as llm:

        def check_model(model_id):
            for name, submodule in model_id.named_modules():
                # Could check more modules if necessary
                if name == "model_id.layers.0.self_attn.qkv_proj":
-                    assert isinstance(submodule.quant_method, linear_method_cls)
-
-                    config = submodule.quant_method.quant_config
-                    assert config.checkpoint_format == "gptq_v2"
-                    assert submodule.quant_method.use_v2_format
-
+                    assert isinstance(submodule.quant_method, AutoGPTQLinearMethod)
                    # Just break since currently we only check 1 module
                    break

@@ -48,6 +44,7 @@ def test_model_load(vllm_runner, model_id, monkeypatch):
        llm.apply_model(check_model)


+@pytest.mark.skip(reason="2-bit GPTQ is no longer supported after Marlin consolidation")
@pytest.mark.parametrize("model_id", MODELS)
 def test_model_inference(vllm_runner, model_id):
    # Prepare prompt to test the model's generation result.
@@ -8,8 +8,7 @@ Run `pytest tests/quantization/test_quant_lm_head_true.py --forked`.
 import pytest
 import torch

-from vllm.model_executor.layers.quantization.gptq import GPTQLinearMethod
-from vllm.model_executor.layers.quantization.gptq_marlin import GPTQMarlinLinearMethod
+from vllm.model_executor.layers.quantization.auto_gptq import AutoGPTQLinearMethod
 from vllm.model_executor.layers.vocab_parallel_embedding import (
    UnquantizedEmbeddingMethod,
 )
@@ -40,7 +39,7 @@ def test_lm_head(
            if lm_head_quantized:
                assert isinstance(
                    lm_head_layer.quant_method,
-                    (GPTQLinearMethod, GPTQMarlinLinearMethod),
+                    AutoGPTQLinearMethod,
                )
            else:
                assert isinstance(
@@ -40,13 +40,12 @@ from vllm.platforms import current_platform
            Fp8PerBlockOnlineLinearMethod,
            Fp8PerBlockOnlineMoEMethod,
        ),
-        # quantization='online with linear_scheme_override and
-        # moe_scheme_override
+        # quantization='online' with per-layer-kind overrides
        (
            "online",
            {
-                "linear_scheme_override": "fp8_per_block",
-                "moe_scheme_override": "fp8_per_tensor",
+                "linear": "fp8_per_block",
+                "moe": "fp8_per_tensor",
            },
            Fp8PerBlockOnlineLinearMethod,
            Fp8PerTensorOnlineMoEMethod,
@@ -0,0 +1,134 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Unit tests for QuantizationConfigArgs parsing."""
+
+import pytest
+
+from vllm.config.quantization import (
+    QUANT_KEY_NAMES,
+    QuantizationConfigArgs,
+    QuantSpec,
+    resolve_quantization_config,
+)
+from vllm.model_executor.layers.quantization.utils.quant_utils import (
+    kFp8Dynamic128Sym,
+    kFp8DynamicTokenSym,
+    kFp8Static128BlockSym,
+    kFp8StaticTensorSym,
+    kInt8StaticChannelSym,
+    kMxfp8Dynamic,
+)
+
+# ---- QuantSpec ------------------------------------------------------------
+
+
+def test_quant_spec_resolves_string_to_quant_key():
+    spec = QuantSpec(weight="mxfp8", activation="fp8_per_token")
+    assert spec.weight == kMxfp8Dynamic
+    assert spec.activation == kFp8DynamicTokenSym
+
+
+def test_quant_spec_accepts_quant_key_directly():
+    spec = QuantSpec(weight=kFp8StaticTensorSym)
+    assert spec.weight is kFp8StaticTensorSym
+    assert spec.activation is None
+
+
+def test_quant_spec_rejects_unknown_name():
+    with pytest.raises(ValueError, match="unknown quantization name"):
+        QuantSpec(weight="not_a_real_format")
+
+
+# ---- QuantizationConfigArgs string shorthand on linear/moe ----------------
+
+
+def test_args_linear_string_resolves_via_quant_key_names():
+    # A bare QUANT_KEY_NAMES entry desugars to QuantSpec(weight=<key>).
+    args = QuantizationConfigArgs(linear="fp8_per_block_static")
+    assert args.linear == QuantSpec(weight=kFp8Static128BlockSym)
+    assert args.moe is None
+
+
+def test_args_moe_string_resolves_via_online_shorthand():
+    # An online-shorthand name pulls the matching slot from _ONLINE_SHORTHANDS
+    # (so `linear: "fp8_per_block"` and `moe: "fp8_per_block"` produce the
+    # same per-layer-kind spec the `--quantization fp8_per_block` shorthand
+    # would).
+    args = QuantizationConfigArgs(moe="fp8_per_block")
+    assert args.moe == QuantSpec(weight=kFp8Static128BlockSym)
+
+
+def test_args_string_shorthand_missing_slot_raises():
+    # int8_per_channel_weight_only sets only `moe`; using it on `linear`
+    # has no defined spec and should raise rather than silently no-op.
+    with pytest.raises(ValueError, match="does not define a linear spec"):
+        QuantizationConfigArgs(linear="int8_per_channel_weight_only")
+
+
+def test_args_accepts_dict_form():
+    args = QuantizationConfigArgs(moe={"activation": "mxfp8"})
+    assert args.moe == QuantSpec(weight=None, activation=kMxfp8Dynamic)
+
+
+# ---- resolve_quantization_config -----------------------------------------
+
+
+def test_resolve_shorthand_only_populates_both_slots():
+    args = resolve_quantization_config("fp8_per_block", None)
+    assert args.linear == QuantSpec(weight=kFp8Static128BlockSym)
+    assert args.moe == QuantSpec(weight=kFp8Static128BlockSym)
+
+
+def test_resolve_int8_shorthand_leaves_linear_unset():
+    # int8_per_channel_weight_only is MoE-only; linear stays None so that
+    # OnlineQuantizationConfig leaves Linear layers in full precision.
+    args = resolve_quantization_config("int8_per_channel_weight_only", None)
+    assert args.linear is None
+    assert args.moe == QuantSpec(weight=kInt8StaticChannelSym)
+
+
+def test_resolve_quantization_config_only():
+    # When only `quantization_config` is given (e.g. for an already-quantized
+    # checkpoint that needs an activation override), it's returned as-is.
+    args = resolve_quantization_config(None, {"moe": {"activation": "mxfp8"}})
+    assert args.linear is None
+    assert args.moe == QuantSpec(weight=None, activation=kMxfp8Dynamic)
+
+
+def test_resolve_merges_explicit_over_shorthand():
+    # Explicit linear in quantization_config wins; moe falls back to the
+    # shorthand's slot.
+    args = resolve_quantization_config(
+        "fp8_per_tensor",
+        {"linear": "fp8_per_block"},
+    )
+    assert args.linear == QuantSpec(weight=kFp8Static128BlockSym)
+    assert args.moe == QuantSpec(weight=kFp8StaticTensorSym)
+
+
+def test_resolve_rejects_quantization_config_with_non_shorthand_quant():
+    # If --quantization names something other than an online shorthand,
+    # quantization_config is not allowed via this path (checkpoint quant
+    # paths read it directly off ModelConfig instead).
+    with pytest.raises(ValueError, match="quantization_config is only supported"):
+        resolve_quantization_config("gptq", {"linear": "fp8_per_block"})
+
+
+# ---- QUANT_KEY_NAMES coverage --------------------------------------------
+
+
+def test_quant_key_names_round_trip():
+    # Every advertised name should round-trip through QuantSpec without error
+    # and produce the same QuantKey it maps to.
+    for name, expected in QUANT_KEY_NAMES.items():
+        assert QuantSpec(weight=name).weight == expected, name
+        assert QuantSpec(activation=name).activation == expected, name
+
+
+def test_static_block_weight_paired_with_dynamic_block_activation():
+    # The block-FP8 shorthand pair: 128x128 static weights + 1x128 dynamic
+    # activations. Pinning this so renames in QUANT_KEY_NAMES don't quietly
+    # rewire the kernel dispatch.
+    spec = QuantSpec(weight="fp8_per_block_static", activation="fp8_per_block_dynamic")
+    assert spec.weight == kFp8Static128BlockSym
+    assert spec.activation == kFp8Dynamic128Sym
@@ -240,8 +240,13 @@ WIKITEXT_ACCURACY_CONFIGS = [
    not QUARK_MXFP4_AVAILABLE,
    reason=f"amd-quark>={QUARK_MXFP4_MIN_VERSION} is not available",
 )
-@pytest.mark.parametrize("config", WIKITEXT_ACCURACY_CONFIGS)
-@pytest.mark.parametrize("tp_size", [1, 2])
+@pytest.mark.parametrize(
+    "config",
+    [pytest.param(val, id=f"config:{val}") for val in WIKITEXT_ACCURACY_CONFIGS],
+)
+@pytest.mark.parametrize(
+    "tp_size", [pytest.param(val, id=f"tp_size:{val}") for val in [1, 2]]
+)
 def test_ocp_mx_wikitext_correctness(config: AccuracyTestConfig, tp_size: int):
    device_count = torch.accelerator.device_count()
    if device_count < tp_size:
@@ -268,6 +273,53 @@ def test_ocp_mx_wikitext_correctness(config: AccuracyTestConfig, tp_size: int):
    ), f"Expected: {EXPECTED_VALUE} |  Measured: {measured_value}"


+@pytest.mark.skipif(
+    not QUARK_MXFP4_AVAILABLE,
+    reason=f"amd-quark>={QUARK_MXFP4_MIN_VERSION} is not available",
+)
+@pytest.mark.parametrize("tp_size", [1, 2])
+def test_nvfp4_wikitext_correctness(tp_size: int):
+    device_count = torch.accelerator.device_count()
+    if device_count < tp_size:
+        pytest.skip(f"This test requires >={tp_size} gpus, got only {device_count}")
+
+    # NOTE: expected_value from nvidia/Qwen3-30B-A3B-NVFP4
+    expected_value = 11.2391
+
+    model_name = "amd-quark/Qwen3-30B-A3B-nvfp4-quark"
+    task = "wikitext"
+
+    rtol = 0.25
+
+    config = AccuracyTestConfig(
+        model_name=model_name,
+        excepted_value=expected_value,
+    )
+
+    model_args = config.get_model_args(
+        tp_size=tp_size,
+        kwargs={
+            "cudagraph_capture_sizes": [16],
+        },
+    )
+    model_args.pop("add_bos_token")
+
+    # Smaller cudagraph_capture_sizes to speed up the test.
+    results = lm_eval.simple_evaluate(
+        model="vllm",
+        model_args=model_args,
+        tasks=task,
+        batch_size=64,
+    )
+
+    EXPECTED_VALUE = config.excepted_value
+    measured_value = results["results"][task]["word_perplexity,none"]
+    assert (
+        measured_value < EXPECTED_VALUE + rtol
+        and measured_value > EXPECTED_VALUE - rtol
+    ), f"Expected: {EXPECTED_VALUE} |  Measured: {measured_value}"
+
+
@pytest.mark.parametrize("config", GSM8K_ACCURACY_CONFIGS)
@pytest.mark.skipif(
    not QUARK_MXFP4_AVAILABLE,
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Andreas KaratzasandGitHub	d735968f6d	[ROCm][CI] Stage B gating (#42025 ) Signed-off-by: Andreas Karatzas <akaratza@amd.com>	2026-05-15 01:49:27 -07:00
bobofang_amd GitHub TJian	ccde9540be	DeepSeekV4-Pro enable cuda graph full and piecewise mode (#42604 ) Co-authored-by: TJian <tunjian.tan@embeddedllm.com>	2026-05-15 01:45:30 -07:00
wang.yuqiandGitHub	75fd68c7a5	[Entrypoints] Split the pooling offline API into PoolingOfflineMixin. (#42267 ) Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>	2026-05-15 08:05:57 +00:00
Yifan QiaoandGitHub	4b364f810e	[Core][DSV4] Skip caching SWA blocks that can never serve a prefix-cache hit (#42258 ) Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>	2026-05-15 15:59:18 +08:00
yupeng GitHub Jee Jee Li	31fa757cf9	[Misc] Make it simpler to replace out-of-tree layer classes with related LoRA layers. (#42306 ) Signed-off-by: paulyu12 <507435917@qq.com> Co-authored-by: Jee Jee Li <pandaleefree@gmail.com>	2026-05-15 15:20:42 +08:00
Cyrus LeungandGitHub	2676ab1e0b	[Deprecation] Remove old locations of `get_tokenizer` and `resolve_hf_chat_template` (#35024 ) Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk>	2026-05-15 00:13:32 -07:00
Daniel W GitHubLi, Jiang <jiang1.li@intel.com>	27b85d2084	[Bugfix] Clarify CPU backend memory error messages reference shared flag (#42479 ) Signed-off-by: daniel-devlab <282598346+daniel-devlab@users.noreply.github.com> Co-authored-by: Li, Jiang <jiang1.li@intel.com>	2026-05-15 06:35:05 +00:00
Louie TsaiandGitHub	e30f39c4f1	Update Intel Xeon model list and vLLM Benchmark Suite BKMs (#42607 ) Signed-off-by: louie-tsai <louie.tsai@intel.com>	2026-05-15 05:14:03 +00:00
Keyi Li GitHub Keyi Limergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	bf610c2f56	[Bugfix] Fix inverted condition causing thinking_token_budget to be silently ignored (#41674 ) Signed-off-by: Keyi Li <likey6688@gmail.com> Co-authored-by: Keyi Li <likey6688@gmail.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-15 12:48:49 +08:00
Isotr0py GitHub zxy	faa4b76afa	[Model] Support InternS2 Preview (#42705 ) Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: zxy <46674730+CUHKSZzxy@users.noreply.github.com>	2026-05-14 21:30:26 -07:00
lyd1992 GitHub Claude	f351455f0f	[CPU][RISC-V] Add RVV-optimized attention kernels for RISC-V Vector Extension (#40119 ) Signed-off-by: liuyudong <liuyudong@iscas.ac.cn> Co-authored-by: Claude <noreply@anthropic.com>	2026-05-15 12:08:23 +08:00
Cyrus LeungandGitHub	56434e8651	[Bugfix] Fix incorrect chat template format for Qwen3.5 (#42660 ) Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk>	2026-05-14 20:52:52 -07:00
Chao-Ju Chen GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	0d4d334eaa	Bump llguidance to 1.7 (#42150 ) Signed-off-by: RickyChen / 陳昭儒 <ricky.chen@infinirc.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-14 20:35:27 -04:00
Chengyi Nie GitHub Chengyi Nie Cursor	fa2a33b893	[Quant] Consolidate GPTQ: rename gptq_marlin.py to auto_gptq.py (#38288 ) Signed-off-by: Chengyi Nie <cnie@roblox.com> Co-authored-by: Chengyi Nie <cnie@roblox.com> Co-authored-by: Cursor <cursoragent@cursor.com>	2026-05-15 08:25:52 +08:00
Giancarlo DelfinandGitHub	3b6a204789	[Model Runner V2][Bug Fix][DSV4] Ensure lazy attention state initializations happen during cudagraph capture (#42444 ) Signed-off-by: Giancarlo Delfin <gdelfin@inferact.ai>	2026-05-14 16:16:17 -07:00
Yongye Zhu GitHub Claude	f8848b2f2d	[Bugfix] Add swiglu limits to deepgemm fp8 methods (#41986 ) Signed-off-by: Yongye Zhu <zyy1102000@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-05-14 15:43:13 -07:00
Charlie FuandGitHub	4cfcc0866f	[CI][ROCm] Remove unsupported cases in test_fusion.py (#38680 ) Signed-off-by: charlifu <charlifu@amd.com>	2026-05-14 17:37:18 -04:00
Tres GitHub Cursor	f887aa1a53	[Aiter][ROCm] RMSNormGated+GroupedQuantFP8 fusion (#40710 ) Signed-off-by: Tres Popp <tres.popp@amd.com> Signed-off-by: Tres Popp <trespopp@gmail.com> Co-authored-by: Cursor <cursoragent@cursor.com>	2026-05-14 15:37:09 -04:00
Matthew BonanniandGitHub	9898f94abe	[Attention] Remove deprecated MLA prefill arguments (#42555 ) Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>	2026-05-14 10:34:06 -07:00
Wentao Ye GitHub Nick Hill	ae4f59f0ec	[Model Runner v2] Oracle for model runner v2 - qwen3 dense model by default [1/N] (#39337 ) Signed-off-by: yewentao256 <zhyanwentao@126.com> Signed-off-by: Nick Hill <nickhill123@gmail.com> Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: Nick Hill <nickhill123@gmail.com>	2026-05-14 10:02:33 -07:00
RanranandGitHub	f3d5360591	[Bugfix][Multimodal] PyAV video backend returns keyframes labeled as targets (#42586 ) Signed-off-by: Ranran <hzz5361@psu.edu>	2026-05-14 08:56:59 -07:00
Baorun (Lauren) MuandGitHub	a7737cb4f3	[Fix] Misc Fixes in ViT CUDA Graph (#38040 ) Signed-off-by: Baorun Mu <bmu@nvidia.com>	2026-05-14 23:49:06 +08:00
Cyrus LeungandGitHub	b8a25d0e12	[Bugfix] Fix LM detection for Nemotron Parse (#42641 ) Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk>	2026-05-14 23:42:10 +08:00
frida-anderssonandGitHub	f07b1da797	[ROCm] Enable gluon paged MQA logits on gfx950 (MI355X) (#42062 ) Signed-off-by: Frida Andersson <fanderss@amd.com>	2026-05-14 15:39:26 +00:00
Vadim Gimpelson GitHub Nick Hill	f60c6b33a5	[V1][DP][LB] Publish request counts at the start of each engine step (#41626 ) Signed-off-by: Vadim Gimpelson <vadim.gimpelson@gmail.com> Signed-off-by: Vadim Gimpelson <156319763+vadiklyutiy@users.noreply.github.com> Co-authored-by: Nick Hill <nickhill123@gmail.com>	2026-05-14 15:39:24 +00:00
zhanqiuhu GitHub Nicolò Lucchesi	24337fb860	PD disagg with NIXL Connector: GDN support (Qwen3.5) (#41869 ) Signed-off-by: Zhanqiu Hu <zhu@redhat.com> Co-authored-by: Nicolò Lucchesi <nlucches@redhat.com>	2026-05-14 16:33:01 +02:00
Hongxin Xu GitHub Claude Opus 4.6 arlenxu Junjie Zhang	c7560af424	[RFC] Replace shared-memory routed experts with ModelRunnerOutput transfer and HTTP support (#39568 ) Signed-off-by: xhx1022 <1737006628@qq.com> Signed-off-by: arlenxu <arlenxu@tencent.com> Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-authored-by: arlenxu <arlenxu@tencent.com> Co-authored-by: Junjie Zhang <junj.jay.zhang@gmail.com>	2026-05-14 14:12:30 +00:00
Mohammad Miadh AngkadandGitHub	2317682f95	[Bugfix] Fix TRTLLM ragged MLA prefill workspace warmup (#42112 ) Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>	2026-05-14 09:48:56 -04:00
Martin Hickey GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>Or Ozeri	5bd8c71e79	[kv_offload] Implement `reset_cache()` for the offloading connector (#41956 ) Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> Co-authored-by: Or Ozeri <or@ozery.com>	2026-05-14 16:00:10 +03:00
Wentao YeandGitHub	6548560496	[Compile] Fix compile warning with topk softplus sqrt (#41261 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-14 05:12:50 -07:00
Jee Jee Li GitHub jeejeeleemergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	0a65d46628	[DSV4] Fuse norm and router for low latency scenario (#41263 ) Signed-off-by: Jee Jee Li <pandaleefree@gmail.com> Signed-off-by: jeejeelee <jeejeelee@verda-b300-05.datacrunch.io> Co-authored-by: jeejeelee <jeejeelee@verda-b300-05.datacrunch.io> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-14 05:11:02 -07:00
Zhenzhong XuandGitHub	1ea9401364	[Quantization][Autoround][Toolkit] Add W4A16 Support (#39778 ) Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com> Signed-off-by: Zhenzhong Xu <zhenzhong.xu@intel.com>	2026-05-14 19:18:49 +08:00
Libin Tang GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	9946c38b7f	[XPU] Fix double-transpose in XPUFP8ScaledMMLinearKernel for W8A8 quant method (#41689 ) Signed-off-by: Libin Tang <libin.tang@intel.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-14 17:17:39 +08:00
Wentao Ye GitHub Roger Wang	23c85343fb	[Bug] Fix DeepSeek V4 `AttributeError: module 'cutlass.cute.nvgpu' has no attribute 'LoadCacheMode'` (#42342 ) Signed-off-by: yewentao256 <zhyanwentao@126.com> Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: Roger Wang <hey@rogerw.io>	2026-05-14 02:00:20 -07:00
rasmithandGitHub	768f4a6f26	[CI][AMD][BugFix] Prevent triton compiler error when running test_moe_layer with use_ep = True on ROCm (#40857 ) Signed-off-by: Randall Smith <Randall.Smith@amd.com>	2026-05-14 08:44:22 +00:00
rasmithandGitHub	addef3299c	[CI][AMD] Skip tests where models have problems or fails on both HW types (#42126 ) Signed-off-by: Randall Smith <Randall.Smith@amd.com>	2026-05-14 08:21:06 +00:00
haic0 GitHub root Jhao-Ting Chen simondanielsson	ce29c26b31	Update Dockerfile.rocm for AINIC & Thor NIC (#40453 ) Signed-off-by: root <root@gbt350-odcdh5-wbb3.png-odc.dcgpu> Signed-off-by: Jhao-Ting Chen <jhaotingc@nvidia.com> Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com> Co-authored-by: root <root@gbt350-odcdh5-wbb3.png-odc.dcgpu> Co-authored-by: Jhao-Ting Chen <jhaotingc@nvidia.com> Co-authored-by: simondanielsson <simon.danielsson99@hotmail.com>	2026-05-14 15:24:27 +08:00
aoshen02andGitHub	8c79ad6580	Revert "[Core] Replace routing replay with device cache and async D2H pipeline" (#39917 ) (#42434 ) Signed-off-by: aoshen02 <aoshen@inferact.ai>	2026-05-13 23:49:01 -07:00
Yongye Zhu GitHub Roger Wang	0d2732dd91	[MLA Attention Backend] Add TOKENSPEED_MLA backend for DSR1/Kimi K25 prefill + decode on Blackwell (#41778 ) Signed-off-by: Yongye Zhu <zyy1102000@gmail.com> Signed-off-by: Roger Wang <hey@rogerw.io> Co-authored-by: Roger Wang <hey@rogerw.io>	2026-05-13 23:48:02 -07:00
Rebecca LeeandGitHub	fd7d858c8a	Use hidden_pad and intermediate_pad from vLLM #34301 (#42098 ) Signed-off-by: Rebecca Lee <Rebecca.Lee@amd.com>	2026-05-14 14:21:04 +08:00
liuzhenweiandGitHub	b26558d4a3	[CI][XPU] skip ut of offload connector (#42598 ) Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com>	2026-05-14 13:13:53 +08:00
Sarah SalahandGitHub	bf0d2dc6d7	[Misc] Fix mypy error in parser_manager type narrowing (#42441 ) Signed-off-by: Sarah-Salah <11881117+Sarah-Salah@users.noreply.github.com>	2026-05-14 02:48:59 +00:00
Harsh Shah GitHub Harsh Shah	ca60a4e84f	[Fix] Weight loading for qwen3_5 using runai_streamer (#42521 ) Signed-off-by: Harsh Shah <iharsh@google.com> Co-authored-by: Harsh Shah <iharsh@google.com>	2026-05-14 10:36:20 +08:00
Roy WangandGitHub	77e1421a68	[Bugfix] Fix EPLB initialization for VLM wrapper models (#39805 ) Signed-off-by: esmeetu <jasonailu87@gmail.com>	2026-05-14 02:26:15 +00:00
Kunshang JiandGitHub	751b9f14bd	[XPU][CT] Support mxfp8 moe model (#41918 ) Signed-off-by: Kunshang Ji <kunshang.ji@intel.com>	2026-05-14 09:47:10 +08:00
Krish GuptaandGitHub	70c00163ff	[Feature] Add instruction support for score/rerank chat templates (#42412 ) Signed-off-by: KrxGu <krishom70@gmail.com>	2026-05-14 09:41:22 +08:00
Siddharth BedekarandGitHub	f51f6844f9	[Bugfix][Spec Decode] Wire draft_probs into probabilistic draft_model rejection (#40269 )	2026-05-13 21:04:03 -04:00
longguoandGitHub	665f9c4253	[Bugfix] Fix Gemma4ToolParser streaming float corruption (#42128 ) Signed-off-by: abinggo <107740309+abinggo@users.noreply.github.com>	2026-05-13 18:03:30 -07:00
Flora FengandGitHub	1087676a90	[Refactor] Use shared utils in hermes tool parser (#42570 ) Signed-off-by: sfeng33 <4florafeng@gmail.com>	2026-05-13 20:35:45 -04:00
Yiyang "Ian" Liu GitHub Copilot sfeng33 Roger Wang	63cc8a55a9	fix(tool-parser): preserve "none"/"nil" strings as valid enum values in minimax_m2 (#39599 ) Signed-off-by: Yiyang Liu <yiyangliu@microsoft.com> Signed-off-by: Yiyang Liu <37043548+ianliuy@users.noreply.github.com> Signed-off-by: sfeng33 <4florafeng@gmail.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: sfeng33 <4florafeng@gmail.com> Co-authored-by: Roger Wang <hey@rogerw.io>	2026-05-13 20:35:34 -04:00
Divakar VermaandGitHub	ca7e4546da	[CI] set max transformers version for skywork model (#42104 ) Signed-off-by: Divakar Verma <divakar.verma@amd.com>	2026-05-13 16:53:49 -07:00
Yufeng He GitHub Matthew Bonanni	b2198670b1	[Bugfix] V1: support tuple model outputs in ubatch wrapper (dbo + spec decode) (#40789 ) Signed-off-by: Yufeng He <40085740+he-yufeng@users.noreply.github.com> Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>	2026-05-13 15:47:51 -07:00
Mohammad Miadh AngkadandGitHub	f1cc7aad3c	[Bugfix] Fix DeepSeek V4 MTP HC state handling (#42320 ) Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>	2026-05-13 15:44:52 -07:00
Lukas GeigerandGitHub	597ed13803	[Core][MM] Do not use urllib3 to parse data URLs (#42535 ) Signed-off-by: Lukas Geiger <lukas.geiger94@gmail.com>	2026-05-13 22:21:01 +00:00
liangel-02andGitHub	6b5c389ee3	expose flex block size for batch invariant mode (#41252 ) Signed-off-by: Angel Li <liangel@meta.com>	2026-05-13 14:11:57 -07:00
Michael GoinandGitHub	8efd508204	[Quantization] Rework quantization_config to use QuantKey and allow for activation override (#41566 )	2026-05-13 16:58:32 -04:00
ovidiusmandGitHub	cca32d55a2	[PD] Fix broken NIXL EP installation (#42542 ) Signed-off-by: Ovidiu Mara <ovidium@nvidia.com>	2026-05-13 13:55:51 -07:00
Walter Beller-MoralesandGitHub	873910d608	[Frontend] add support for thinking_token_budget in completions (#42116 )	2026-05-13 16:01:52 -04:00
Wentao YeandGitHub	3f611f6106	[CI] Fix pre-commit issue (#42563 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-13 12:37:26 -07:00
Nick HillandGitHub	a505cf807e	[ModelRunner V2] Share identical MTP weights (#42538 ) Signed-off-by: Nick Hill <nickhill123@gmail.com>	2026-05-13 18:57:04 +00:00
fxmarty-amd GitHub Kyle Sayers	40330967ab	[Quark] Support loading Quark NVFP4 checkpoints in vLLM (#35859 ) Signed-off-by: Felix Marty <Felix.Marty@amd.com> Signed-off-by: fxmarty-amd <felmarty@amd.com> Co-authored-by: Kyle Sayers <kylesayrs@gmail.com>	2026-05-13 11:17:36 -07:00
Fynn Schmitt-UlmsandGitHub	ab1ad0d7a9	Remove verifier model type check in speculative config (#42536 ) Signed-off-by: Fynn Schmitt-Ulms <fschmitt@redhat.com>	2026-05-13 18:14:39 +00:00
Ben BrowningandGitHub	0f69128a37	[Bugfix] Handle real-world gpt-oss tool call output in Harmony parsing (#42454 ) Signed-off-by: Ben Browning <bbrownin@redhat.com>	2026-05-13 17:54:46 +00:00
John Calderon GitHub Roger Wang	b3c69595a6	[MM][CG] Support ViT CG for Qwen2-VL (#41736 ) Signed-off-by: John Calderon <jcalderon@nvidia.com> Co-authored-by: Roger Wang <hey@rogerw.io>	2026-05-14 01:52:35 +08:00
Michael Goin GitHub Claude Opus 4.6	2f821faeae	[Spec Decode] Support hybrid attention models in extract_hidden_states (#39949 ) Signed-off-by: mgoin <mgoin64@gmail.com> Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-05-13 10:45:53 -07:00
Noelia Bentancor GitHub Copilot Autofix powered by AI Wentao Yemergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	5794c65f8c	[Bugfix][Model] Gemma4 MoE routing closure captures per_expert_scale, breaking functional_call substitution (#42250 ) Signed-off-by: Noelia <noeliabentancor1@gmail.com> Signed-off-by: Noelia Bentancor <71080743+NoeliaBentancor@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-13 17:43:12 +00:00
CynicDoraandGitHub	256dbcaabf	[Feature] Support custom callable proposer backend for speculative decoding (#39487 ) Signed-off-by: 524031910363 <hyzhyzsh@sjtu.edu.cn> Signed-off-by: CynicDora <hyzhyzsh@sjtu.edu.cn>	2026-05-13 16:53:01 +00:00
Wentao YeandGitHub	e35c0d4c63	[Feature] Support compile mode for batch invariance on SM80 (#42456 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-13 11:02:39 -04:00
Ronen SchafferandGitHub	11f6b545d4	[kv_offload] Add multi-tier KV cache offloading framework (#40020 ) Signed-off-by: Ronen Schaffer <ronen.schaffer@ibm.com>	2026-05-13 17:21:43 +03:00
TJian GitHub Yongye Zhu	a8887c208f	[Bugfix] [ROCm] [DSV4] [Perf] Add aiter mhc support (#41946 ) Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com> Signed-off-by: Yongye Zhu <zyy1102000@gmail.com> Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>	2026-05-13 21:43:15 +08:00
zofia GitHub Kunshang Ji	0ddaf6dffa	[XPU] [CT] Enable CT W4A4MxFp4 path and add xpu kernel (#38896 ) Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com> Signed-off-by: zofia <110436990+zufangzhu@users.noreply.github.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-05-13 06:43:00 -07:00
Marek WawrzosandGitHub	67671692ac	[CI] Re-enable Nemotron Parse parity test and switch testing to nemotron-parse v1.2 (#42498 ) Signed-off-by: <mwawrzos@nvidia.com>	2026-05-13 21:05:27 +08:00