Remove broken allocator pre-warming (initializes CUDA before fork)

The pre-warming called torch.cuda.mem_get_info() in the parent process before the EngineCore subprocess forks, which initialized CUDA. This caused "Cannot re-initialize CUDA in forked subprocess" in the child process, breaking tests that previously passed. The pre-warming also wouldn't have helped — PyTorch's allocator pool is per-process, so the child's pool starts empty regardless. Signed-off-by: khluu <khluu000@gmail.com> Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
[CI] Migrate all remaining gpu_1_queue jobs to h200_18gb MIG
2026-05-13 02:26:43 -07:00 · 2026-05-12 18:50:08 -07:00 · 2026-05-12 15:57:19 -07:00 · 2026-05-12 02:00:37 -07:00 · 2026-05-12 08:45:51 +00:00 · 2026-05-12 07:47:47 +00:00
278 changed files with 9553 additions and 2541 deletions
@@ -460,7 +460,7 @@ steps:
  - tests/lora
  - vllm/platforms/rocm.py
  commands:
-  - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_llm_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py
+  - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_qwen3_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py

 #------------------------------------------------------  mi250 · model_executor  -------------------------------------------------------#

@@ -929,6 +929,7 @@ steps:
  - tests/tokenizers_
  - tests/reasoning
  - tests/tool_parsers
+  - tests/parser
  - tests/transformers_utils
  - tests/config
  commands:
@@ -942,6 +943,7 @@ steps:
  - pytest -v -s tokenizers_
  - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
  - pytest -v -s tool_parsers
+  - pytest -v -s parser
  - pytest -v -s transformers_utils
  - pytest -v -s config

@@ -1320,7 +1322,6 @@ steps:
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
  - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - pytest -v -s entrypoints/openai/speech_to_text/
  - pytest -v -s entrypoints/test_chat_utils.py

 - label: Entrypoints Integration (API Server openai - Part 3) # TBD
@@ -1336,7 +1337,21 @@ 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/speech_to_text/ --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
+
+- label: Entrypoints Integration (Speech to Text) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  fast_check: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/speech_to_text
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/speech_to_text

 - label: Entrypoints Integration (LLM) # TBD
  timeout_in_minutes: 180
@@ -1760,7 +1775,7 @@ steps:
  - export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
  - pytest -v -s -x lora/test_chatglm3_tp.py
  - pytest -v -s -x lora/test_llama_tp.py
-  - pytest -v -s -x lora/test_llm_with_multi_loras.py
+  - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
  - pytest -v -s -x lora/test_olmoe_tp.py
  - pytest -v -s -x lora/test_gptoss_tp.py
  - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
@@ -2766,7 +2781,6 @@ steps:
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
  - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - pytest -v -s entrypoints/openai/speech_to_text/
  - pytest -v -s entrypoints/test_chat_utils.py

 - label: Entrypoints Integration (API Server openai - Part 3) # TBD
@@ -2782,7 +2796,21 @@ 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/speech_to_text/ --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
+
+- label: Entrypoints Integration (Speech to Text) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi355]
+  agent_pool: mi355_1
+  fast_check: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/speech_to_text
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/speech_to_text

 - label: Entrypoints Integration (Pooling) # TBD
  timeout_in_minutes: 180
@@ -20,6 +20,7 @@ steps:
 - label: Cudagraph
  key: cudagraph
  timeout_in_minutes: 20
+  device: h200_18gb
  source_file_dependencies:
  - tests/v1/cudagraph
  - vllm/v1/cudagraph_dispatcher.py
@@ -30,6 +30,7 @@ steps:
 - label: Engine (1 GPU)
  key: engine-1-gpu
  timeout_in_minutes: 30
+  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/engine/
    - tests/v1/engine/
@@ -50,6 +51,7 @@ steps:
 - label: e2e Core (1 GPU)
  key: e2e-core-1-gpu
  timeout_in_minutes: 30
+  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/
    - tests/v1/e2e/general/
@@ -6,17 +6,19 @@ steps:
  key: entrypoints-unit-tests
  timeout_in_minutes: 10
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/entrypoints
  - tests/entrypoints/
  commands:
  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling
+  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text

 - label: Entrypoints Integration (LLM)
  key: entrypoints-integration-llm
  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/llm
@@ -36,6 +38,7 @@ steps:
  key: entrypoints-integration-api-server-openai-part-1
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/openai
@@ -44,18 +47,17 @@ steps:
  - 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

-
 - label: Entrypoints Integration (API Server openai - Part 2)
  key: entrypoints-integration-api-server-openai-part-2
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/openai
  - tests/entrypoints/test_chat_utils
  commands:
  - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - pytest -v -s entrypoints/openai/speech_to_text/
  - pytest -v -s entrypoints/test_chat_utils.py

 - label: Entrypoints Integration (API Server openai - Part 3)
@@ -69,12 +71,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/speech_to_text/ --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

 - label: Entrypoints Integration (API Server 2)
  key: entrypoints-integration-api-server-2
  timeout_in_minutes: 130
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/rpc
@@ -86,10 +89,23 @@ steps:
  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
  - pytest -v -s tool_use

+- label: Entrypoints Integration (Speech to Text)
+  key: entrypoints-integration-speech_to_text
+  timeout_in_minutes: 50
+  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/speech_to_text
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/speech_to_text
+
 - label: Entrypoints Integration (Pooling)
  key: entrypoints-integration-pooling
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/pooling
@@ -101,6 +117,7 @@ steps:
  key: entrypoints-integration-responses-api
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/entrypoints/openai/responses
@@ -115,5 +132,5 @@ steps:
  - csrc/
  - vllm/entrypoints/openai/
  - vllm/model_executor/models/whisper.py
-  commands: # LMEval+Transcription WER check
+  commands: # LMEval
  - pytest -s entrypoints/openai/correctness/
@@ -17,6 +17,7 @@ steps:
 - label: Kernels Core Operation Test
  key: kernels-core-operation-test
  timeout_in_minutes: 75
+  device: h200_18gb
  source_file_dependencies:
  - csrc/
  - tests/kernels/core
@@ -41,6 +42,7 @@ steps:
 - label: Kernels Attention Test %N
  key: kernels-attention-test
  timeout_in_minutes: 35
+  device: h200_18gb
  source_file_dependencies:
  - csrc/attention/
  - vllm/v1/attention
@@ -55,6 +57,7 @@ steps:
 - label: Kernels Quantization Test %N
  key: kernels-quantization-test
  timeout_in_minutes: 90
+  device: h200_18gb
  source_file_dependencies:
  - csrc/quantization/
  - vllm/model_executor/layers/quantization
@@ -66,6 +69,7 @@ steps:
 - label: Kernels MoE Test %N
  key: kernels-moe-test
  timeout_in_minutes: 25
+  device: h200_18gb
  source_file_dependencies:
  - csrc/quantization/cutlass_w8a8/moe/
  - csrc/moe/
@@ -82,6 +86,7 @@ steps:
 - label: Kernels Mamba Test
  key: kernels-mamba-test
  timeout_in_minutes: 45
+  device: h200_18gb
  source_file_dependencies:
  - csrc/mamba/
  - tests/kernels/mamba
@@ -91,6 +96,7 @@ steps:

 - label: Kernels KDA Test
  timeout_in_minutes: 20
+  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers/fla/ops/kda.py
  - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
@@ -242,7 +248,7 @@ steps:
 - label: Kernels FusedMoE Layer Test (2 B200s)
  key: kernels-fusedmoe-layer-test-2-b200s
  timeout_in_minutes: 90
-  device: b200
+  device: b200-k8s
  num_devices: 2
  source_file_dependencies:
  - csrc/quantization/cutlass_w8a8/moe/
@@ -5,6 +5,7 @@ steps:
 - label: LM Eval Small Models
  key: lm-eval-small-models
  timeout_in_minutes: 75
+  device: h200_18gb
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -40,7 +41,7 @@ steps:
 - label: LM Eval Small Models (B200)
  key: lm-eval-small-models-b200
  timeout_in_minutes: 120
-  device: b200
+  device: b200-k8s
  optional: true
  source_file_dependencies:
  - csrc/
@@ -48,6 +49,18 @@ steps:
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt

+- label: LM Eval Large Models (B200, EP)
+  key: lm-eval-large-models-b200-ep
+  timeout_in_minutes: 120
+  device: b200-k8s
+  optional: true
+  num_devices: 2
+  source_file_dependencies:
+  - csrc/
+  - vllm/model_executor/layers/quantization
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell-ep.txt
+
 - label: LM Eval Qwen3.5 Models (B200)
  key: lm-eval-qwen3-5-models-b200
  timeout_in_minutes: 120
@@ -92,7 +105,7 @@ steps:

 - label: MoE Refactor Integration Test (B200 DP - TEMPORARY)
  key: moe-refactor-integration-test-b200-dp-temporary
-  device: b200
+  device: b200-k8s
  optional: true
  num_devices: 2
  commands:
@@ -102,6 +115,7 @@ steps:
 - label: LM Eval TurboQuant KV Cache
  key: lm-eval-turboquant-kv-cache
  timeout_in_minutes: 75
+  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers/quantization/turboquant/
  - vllm/v1/attention/backends/turboquant_attn.py
@@ -127,7 +141,7 @@ steps:
 - label: GPQA Eval (GPT-OSS) (B200)
  key: gpqa-eval-gpt-oss-b200
  timeout_in_minutes: 120
-  device: b200
+  device: b200-k8s
  optional: true
  num_devices: 2
  source_file_dependencies:
@@ -140,6 +154,7 @@ steps:

 - label: MRCR Eval Small Models
  timeout_in_minutes: 30
+  device: h200_18gb
  source_file_dependencies:
  - tests/evals/mrcr/
  commands:
@@ -5,11 +5,12 @@ steps:
 - label: LoRA %N
  key: lora
  timeout_in_minutes: 30
+  device: h200_18gb
  source_file_dependencies:
  - vllm/lora
  - tests/lora
  commands:
-    - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_llm_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py 
+    - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_qwen3_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py 
  parallelism: 4


@@ -31,7 +32,7 @@ steps:
    # requires multi-GPU testing for validation.
    - pytest -v -s -x lora/test_chatglm3_tp.py
    - pytest -v -s -x lora/test_llama_tp.py
-    - pytest -v -s -x lora/test_llm_with_multi_loras.py
+    - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
    - pytest -v -s -x lora/test_olmoe_tp.py
    - pytest -v -s -x lora/test_gptoss_tp.py
    - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
@@ -5,6 +5,7 @@ steps:
 - label: V1 Spec Decode
  key: v1-spec-decode
  timeout_in_minutes: 30
+  device: h200_18gb
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -58,6 +59,7 @@ steps:
 - label: V1 Core + KV + Metrics
  key: v1-core-kv-metrics
  timeout_in_minutes: 30
+  device: h200_18gb
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -153,6 +155,7 @@ steps:
  key: examples
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/examples"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/entrypoints
  - vllm/multimodal
@@ -210,7 +213,9 @@ steps:
 - label: Python-only Installation
  key: python-only-installation
  depends_on: ~
+  optional: true
  timeout_in_minutes: 20
+  device: h200_18gb
  source_file_dependencies:
  - tests/standalone_tests/python_only_compile.sh
  - setup.py
@@ -220,6 +225,7 @@ steps:
 - label: Async Engine, Inputs, Utils, Worker
  key: async-engine-inputs-utils-worker
  timeout_in_minutes: 50
+  device: h200_18gb
  source_file_dependencies:
  - vllm/assets/
  - vllm/config/
@@ -282,6 +288,7 @@ steps:
  - tests/tokenizers_
  - tests/reasoning
  - tests/tool_parsers
+  - tests/parser
  - tests/transformers_utils
  - tests/config
  device: cpu-small
@@ -296,6 +303,7 @@ steps:
  - pytest -v -s tokenizers_
  - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
  - pytest -v -s tool_parsers
+  - pytest -v -s parser
  - pytest -v -s transformers_utils
  - pytest -v -s config

@@ -5,6 +5,7 @@ steps:
 - label: Model Executor
  key: model-executor
  timeout_in_minutes: 35
+  device: h200_18gb
  source_file_dependencies:
  - vllm/engine/arg_utils.py
  - vllm/config/model.py
@@ -5,6 +5,7 @@ steps:
 - label: Model Runner V2 Core Tests
  key: model-runner-v2-core-tests
  timeout_in_minutes: 45
+  device: h200_18gb
  source_file_dependencies:
  - vllm/v1/worker/gpu/
  - vllm/v1/worker/gpu_worker.py
@@ -29,6 +30,7 @@ steps:
  key: model-runner-v2-examples
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/examples"
+  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/worker/gpu/
    - vllm/v1/core/sched/
@@ -102,6 +104,7 @@ steps:
  key: model-runner-v2-spec-decode
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
+  device: h200_18gb
  source_file_dependencies:
  - vllm/v1/worker/gpu/
  - vllm/v1/worker/gpu_worker.py
@@ -6,6 +6,7 @@ steps:
  key: basic-models-tests-initialization
  timeout_in_minutes: 45
  torch_nightly: true
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/test_initialization.py
@@ -19,6 +20,7 @@ steps:
 - label: Basic Models Tests (Extra Initialization) %N
  key: basic-models-tests-extra-initialization
  timeout_in_minutes: 45
+  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/test_initialization.py
@@ -35,6 +37,7 @@ steps:
 - label: Basic Models Tests (Other)
  key: basic-models-tests-other
  timeout_in_minutes: 45
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/test_terratorch.py
@@ -61,6 +64,7 @@ steps:
  working_dir: "/vllm-workspace/"
  optional: true
  soft_fail: true
+  device: h200_18gb
  commands:
    - pip install --upgrade git+https://github.com/huggingface/transformers
    - pytest -v -s tests/models/test_initialization.py
@@ -77,6 +81,7 @@ steps:
  working_dir: "/vllm-workspace/"
  optional: true
  soft_fail: true
+  device: h200_18gb
  commands:
    - pip install transformers==4.57.5
    - pytest -v -s tests/models/test_initialization.py
@@ -5,6 +5,7 @@ steps:
 - label: Language Models Tests (Standard)
  key: language-models-tests-standard
  timeout_in_minutes: 25
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/language
@@ -18,6 +19,7 @@ steps:
 - label: Language Models Tests (Extra Standard) %N
  key: language-models-tests-extra-standard
  timeout_in_minutes: 45
+  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/language/pooling/test_embedding.py
@@ -35,6 +37,7 @@ steps:
 - label: Language Models Tests (Hybrid) %N
  key: language-models-tests-hybrid
  timeout_in_minutes: 75
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/language/generation
@@ -60,6 +63,7 @@ steps:
 - label: Language Models Test (Extended Generation) # 80min
  key: language-models-test-extended-generation
  timeout_in_minutes: 110
+  device: h200_18gb
  optional: true
  source_file_dependencies:
  - vllm/
@@ -85,6 +89,7 @@ steps:
 - label: Language Models Test (Extended Pooling)  # 36min
  key: language-models-test-extended-pooling
  timeout_in_minutes: 50
+  device: h200_18gb
  optional: true
  source_file_dependencies:
  - vllm/
@@ -40,6 +40,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
  key: multi-modal-models-standard-3-llava-qwen2-vl
  timeout_in_minutes: 45
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -56,6 +57,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
  key: multi-modal-models-standard-4-other-whisper
  timeout_in_minutes: 45
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -94,6 +96,7 @@ steps:
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
  key: multi-modal-accuracy-eval-small-models
  timeout_in_minutes: 70
+  device: h200_18gb
  working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
  source_file_dependencies:
  - vllm/multimodal/
@@ -105,6 +108,7 @@ steps:
 - label: Multi-Modal Models (Extended Generation 1)
  key: multi-modal-models-extended-generation-1
  optional: true
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal/generation
@@ -122,6 +126,7 @@ steps:
 - label: Multi-Modal Models (Extended Generation 2)
  key: multi-modal-models-extended-generation-2
  optional: true
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal/generation
@@ -132,6 +137,7 @@ steps:
 - label: Multi-Modal Models (Extended Generation 3)
  key: multi-modal-models-extended-generation-3
  optional: true
+  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal/generation
@@ -5,6 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
  key: pytorch-compilation-unit-tests
  timeout_in_minutes: 10
+  device: h200_18gb
  source_file_dependencies:
    - vllm/__init__.py
    - vllm/_aiter_ops.py
@@ -78,6 +79,7 @@ steps:
 - label: PyTorch Compilation Passes Unit Tests
  key: pytorch-compilation-passes-unit-tests
  timeout_in_minutes: 20
+  device: h200_18gb
  source_file_dependencies:
    - vllm/__init__.py
    - vllm/_aiter_ops.py
@@ -110,6 +112,7 @@ steps:
 - label: PyTorch Fullgraph Smoke Test
  key: pytorch-fullgraph-smoke-test
  timeout_in_minutes: 35
+  device: h200_18gb
  source_file_dependencies:
  - vllm/__init__.py
  - vllm/_aiter_ops.py
@@ -5,6 +5,7 @@ steps:
 - label: Quantization
  key: quantization
  timeout_in_minutes: 90
+  device: h200_18gb
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -42,6 +43,7 @@ steps:
 - label: Quantized Models Test
  key: quantized-models-test
  timeout_in_minutes: 60
+  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers/quantization
  - tests/models/quantization
@@ -5,6 +5,7 @@ steps:
 - label: Samplers Test
  key: samplers-test
  timeout_in_minutes: 75
+  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers
  - vllm/sampling_metadata.py
@@ -16,7 +16,7 @@ steps:
 - label: Spec Decode Eagle Nightly B200
  key: spec-decode-eagle-nightly-b200
  timeout_in_minutes: 30
-  device: b200
+  device: b200-k8s
  optional: true
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -40,7 +40,7 @@ steps:
 - label: Spec Decode Speculators + MTP Nightly B200
  key: spec-decode-speculators-mtp-nightly-b200
  timeout_in_minutes: 30
-  device: b200
+  device: b200-k8s
  optional: true
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -100,7 +100,7 @@ steps:

 - label: Spec Decode MTP hybrid (B200)
  timeout_in_minutes: 30
-  device: b200
+  device: b200-k8s
  optional: true
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -31,8 +31,8 @@
 /vllm/entrypoints/cli @hmellor @mgoin @DarkLight1337 @russellb
 /vllm/entrypoints/mcp @heheda12345
 /vllm/entrypoints/openai @aarnphm @chaunceyjiang @DarkLight1337 @russellb
-/vllm/entrypoints/openai/realtime @njhill
-/vllm/entrypoints/openai/speech_to_text @NickLucche
+/vllm/entrypoints/speech_to_text/realtime @njhill
+/vllm/entrypoints/speech_to_text @NickLucche
 /vllm/entrypoints/pooling @noooop
 /vllm/entrypoints/sagemaker @DarkLight1337
 /vllm/entrypoints/serve @njhill
@@ -109,6 +109,24 @@ else()
  set(CUDA_SUPPORTED_ARCHS "7.0;7.5;8.0;8.6;8.7;8.9;9.0")
 endif()

+#
+# spinloop extension (pure CXX; must stay above the non-CUDA device branch so
+# CPU builds define the target before the early return)
+#
+set(VLLM_SPINLOOP_EXT_SRC "csrc/spinloop.cpp")
+set(SPINLOOP_COMPILE_FLAGS "")
+if(CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|amd64")
+  list(APPEND SPINLOOP_COMPILE_FLAGS "-mmwaitx")
+endif()
+define_extension_target(
+  spinloop
+  DESTINATION vllm
+  LANGUAGE CXX
+  SOURCES ${VLLM_SPINLOOP_EXT_SRC}
+  COMPILE_FLAGS ${SPINLOOP_COMPILE_FLAGS}
+  USE_SABI 3.11
+  WITH_SOABI)
+
 #
 # Forward the non-CUDA device extensions to external CMake scripts.
 #
@@ -21,28 +21,33 @@ namespace vllm {

 // NOTE Be EXTRA careful with raw_kv_scalar_t, for __half and __nv_bfloat16 it's
 // using u16 as the backing type.
-template <typename qk_t, bool IS_NEOX, typename raw_kv_scalar_t,
-          typename cache_t, Fp8KVCacheDataType kv_dt>
+template <typename qk_t, typename cos_sin_t, bool IS_NEOX,
+          typename raw_kv_scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt>
 __global__ void concat_and_cache_mla_rope_fused_kernel(
    const int64_t* __restrict__ positions,  // [num_tokens]
    qk_t* __restrict__ q_pe,        // [num_tokens, num_q_heads, rot_dim]
    qk_t* __restrict__ k_pe,        // [num_tokens, rot_dim]
    const qk_t* __restrict__ kv_c,  // [num_tokens, kv_lora_rank]
-    const qk_t* __restrict__ rope_cos_sin_cache,  // [max_position, 2,
-                                                  // rot_dim // 2]
+    const cos_sin_t* __restrict__ rope_cos_sin_cache,  // [max_position, 2,
+                                                       // rot_dim // 2]
    const int rot_dim, const int64_t q_pe_stride_token,
    const int64_t q_pe_stride_head, const int64_t k_pe_stride,
    const int64_t kv_c_stride, const int num_q_heads,
    cache_t* __restrict__ kv_cache,  // [num_blocks, block_size, (kv_lora_rank +
                                     // rot_dim)]
-    const int64_t* __restrict__ kv_cache_slot_mapping,  // [num_tokens]
+    const int64_t* __restrict__ slot_mapping,  // [num_tokens]
    const int block_stride, const int entry_stride, const int kv_lora_rank,
    const int block_size, const float* kv_cache_quant_scale) {
  // Each thread block is responsible for one token.
  const int64_t token_idx = blockIdx.x;
+  const int64_t slot_idx = slot_mapping[token_idx];
+  // NOTE: slot_idx can be -1 if the token is padded
+  if (slot_idx < 0) {
+    return;
+  }
  const int64_t pos = positions[token_idx];

-  const qk_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;
+  const cos_sin_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;

  const int embed_dim = rot_dim / 2;

@@ -54,8 +59,8 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(

    // NOTE: Would be nice to have interleaved sin/cos so we could just load
    // both at the same time.
-    qk_t cos = VLLM_LDG(cos_sin_ptr + pair_idx);
-    qk_t sin = VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim);
+    qk_t cos = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx));
+    qk_t sin = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim));

    qk_t* q_pe_head_ptr =
        q_pe + token_idx * q_pe_stride_token + head_idx * q_pe_stride_head;
@@ -81,21 +86,15 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
    q_pe_head_ptr[pair_idx_y] = y_dst;
  }

-  const int64_t slot_idx = kv_cache_slot_mapping[token_idx];
  const int64_t block_idx = slot_idx / block_size;
  const int64_t entry_idx = slot_idx % block_size;

-  // NOTE: slot_idx can be -1 if the token is padded
-  if (slot_idx < 0) {
-    return;
-  }
-
  // K with 1 HEAD
  for (int i = threadIdx.x; i < embed_dim; i += blockDim.x) {
    int pair_idx = i;

-    qk_t cos = VLLM_LDG(cos_sin_ptr + pair_idx);
-    qk_t sin = VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim);
+    qk_t cos = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx));
+    qk_t sin = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim));

    qk_t* k_pe_head_ptr = k_pe + token_idx * k_pe_stride;

@@ -165,36 +164,43 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(

 }  // namespace vllm

-#define CALL_CONCAT_AND_CACHE_MLA_ROPE_FUSED(RAW_KV_T, CACHE_T, KV_DTYPE)      \
-  do {                                                                         \
-    VLLM_DISPATCH_FLOATING_TYPES(q_pe.scalar_type(), "qk_scalar_type", [&] {   \
-      using qk_t = scalar_t;                                                   \
-      if (rope_is_neox) {                                                      \
-        vllm::concat_and_cache_mla_rope_fused_kernel<qk_t, true, RAW_KV_T,     \
-                                                     CACHE_T, KV_DTYPE>        \
-            <<<grid, block, 0, stream>>>(                                      \
-                positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),          \
-                k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),                  \
-                rope_cos_sin_cache.data_ptr<qk_t>(), rot_dim,                  \
-                q_pe_stride_token, q_pe_stride_head, k_pe_stride, kv_c_stride, \
-                num_q_heads, reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),  \
-                kv_cache_slot_mapping.data_ptr<int64_t>(), block_stride,       \
-                entry_stride, kv_lora_rank, block_size,                        \
-                kv_cache_quant_scale.data_ptr<float>());                       \
-      } else {                                                                 \
-        vllm::concat_and_cache_mla_rope_fused_kernel<qk_t, false, RAW_KV_T,    \
-                                                     CACHE_T, KV_DTYPE>        \
-            <<<grid, block, 0, stream>>>(                                      \
-                positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),          \
-                k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),                  \
-                rope_cos_sin_cache.data_ptr<qk_t>(), rot_dim,                  \
-                q_pe_stride_token, q_pe_stride_head, k_pe_stride, kv_c_stride, \
-                num_q_heads, reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),  \
-                kv_cache_slot_mapping.data_ptr<int64_t>(), block_stride,       \
-                entry_stride, kv_lora_rank, block_size,                        \
-                kv_cache_quant_scale.data_ptr<float>());                       \
-      }                                                                        \
-    });                                                                        \
+#define CALL_CONCAT_AND_CACHE_MLA_ROPE_FUSED(RAW_KV_T, CACHE_T, KV_DTYPE)     \
+  do {                                                                        \
+    VLLM_DISPATCH_FLOATING_TYPES(q_pe.scalar_type(), "qk_scalar_type", [&] {  \
+      using qk_t = scalar_t;                                                  \
+      VLLM_DISPATCH_FLOATING_TYPES(                                           \
+          rope_cos_sin_cache.scalar_type(), "rope_cos_sin_cache_scalar_type", \
+          [&] {                                                               \
+            using cos_sin_t = scalar_t;                                       \
+            if (rope_is_neox) {                                               \
+              vllm::concat_and_cache_mla_rope_fused_kernel<                   \
+                  qk_t, cos_sin_t, true, RAW_KV_T, CACHE_T, KV_DTYPE>         \
+                  <<<grid, block, 0, stream>>>(                               \
+                      positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),   \
+                      k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),           \
+                      rope_cos_sin_cache.data_ptr<cos_sin_t>(), rot_dim,      \
+                      q_pe_stride_token, q_pe_stride_head, k_pe_stride,       \
+                      kv_c_stride, num_q_heads,                               \
+                      reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),        \
+                      slot_mapping.data_ptr<int64_t>(), block_stride,         \
+                      entry_stride, kv_lora_rank, block_size,                 \
+                      kv_cache_quant_scale.data_ptr<float>());                \
+            } else {                                                          \
+              vllm::concat_and_cache_mla_rope_fused_kernel<                   \
+                  qk_t, cos_sin_t, false, RAW_KV_T, CACHE_T, KV_DTYPE>        \
+                  <<<grid, block, 0, stream>>>(                               \
+                      positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),   \
+                      k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),           \
+                      rope_cos_sin_cache.data_ptr<cos_sin_t>(), rot_dim,      \
+                      q_pe_stride_token, q_pe_stride_head, k_pe_stride,       \
+                      kv_c_stride, num_q_heads,                               \
+                      reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),        \
+                      slot_mapping.data_ptr<int64_t>(), block_stride,         \
+                      entry_stride, kv_lora_rank, block_size,                 \
+                      kv_cache_quant_scale.data_ptr<float>());                \
+            }                                                                 \
+          });                                                                 \
+    });                                                                       \
  } while (false)

 // Executes RoPE on q_pe and k_pe, then writes k_pe and kv_c in the kv cache.
@@ -208,43 +214,52 @@ void concat_and_cache_mla_rope_fused(
    torch::Tensor& kv_c,                // [num_tokens, kv_lora_rank]
    torch::Tensor& rope_cos_sin_cache,  // [max_position, rot_dim]
    bool rope_is_neox,
-    torch::Tensor&
-        kv_cache_slot_mapping,  // [num_tokens] or [num_actual_tokens]
+    torch::Tensor& slot_mapping,  // [num_tokens] or [num_actual_tokens]
    torch::Tensor&
        kv_cache,  // [num_blocks, block_size, (kv_lora_rank + rot_dim)]
    const std::string& kv_cache_dtype, torch::Tensor& kv_cache_quant_scale) {
-  const int64_t num_tokens = q_pe.size(0);
+  // NOTE(woosuk): In vLLM V1, query/key/position.size(0) can be different from
+  // slot_mapping.size(0) because of padding for CUDA graphs.
+  // In vLLM V0, key.size(0) is always equal to slot_mapping.size(0) because
+  // both include padding.
+  // In vLLM V1, however, key.size(0) can be larger than slot_mapping.size(0)
+  // since key includes padding for CUDA graphs, while slot_mapping does not.
+  // In this case, slot_mapping.size(0) represents the actual number of tokens
+  // before padding.
+  // For compatibility with both cases, we use slot_mapping.size(0) as the
+  // number of tokens.
+  int num_tokens = slot_mapping.size(0);
+  int num_padded_tokens = q_pe.size(0);
+  TORCH_CHECK_GE(num_padded_tokens, num_tokens);

  const int num_q_heads = q_pe.size(1);
  const int rot_dim = q_pe.size(2);
  const int kv_lora_rank = kv_c.size(1);

-  TORCH_CHECK(positions.size(0) >=
-              num_tokens);  // CUDA Graphs might pad this for us
+  TORCH_CHECK_EQ(positions.size(0), num_padded_tokens);
  TORCH_CHECK_EQ(positions.dim(), 1);
  TORCH_CHECK_EQ(positions.scalar_type(), c10::ScalarType::Long);

-  TORCH_CHECK_EQ(q_pe.size(0), num_tokens);
+  TORCH_CHECK_EQ(q_pe.dim(), 3);
+  TORCH_CHECK_EQ(q_pe.size(0), num_padded_tokens);
  TORCH_CHECK_EQ(q_pe.size(1), num_q_heads);
  TORCH_CHECK_EQ(q_pe.size(2), rot_dim);
-  TORCH_CHECK_EQ(q_pe.dim(), 3);

-  TORCH_CHECK_EQ(k_pe.size(0), num_tokens);
-  TORCH_CHECK_EQ(k_pe.size(1), rot_dim);
  TORCH_CHECK_EQ(k_pe.dim(), 2);
+  TORCH_CHECK_EQ(k_pe.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(k_pe.size(1), rot_dim);
  TORCH_CHECK_EQ(k_pe.scalar_type(), q_pe.scalar_type());

-  TORCH_CHECK_EQ(kv_c.size(0), num_tokens);
-  TORCH_CHECK_EQ(kv_c.size(1), kv_lora_rank);
  TORCH_CHECK_EQ(kv_c.dim(), 2);
+  TORCH_CHECK_EQ(kv_c.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(kv_c.size(1), kv_lora_rank);
  TORCH_CHECK_EQ(kv_c.scalar_type(), q_pe.scalar_type());
  TORCH_CHECK_EQ(kv_c.dtype(), q_pe.dtype());

  TORCH_CHECK_EQ(rope_cos_sin_cache.size(1), rot_dim);
-  TORCH_CHECK_EQ(rope_cos_sin_cache.scalar_type(), q_pe.scalar_type());

-  TORCH_CHECK_EQ(kv_cache_slot_mapping.size(0), num_tokens);
-  TORCH_CHECK_EQ(kv_cache_slot_mapping.scalar_type(), c10::ScalarType::Long);
+  TORCH_CHECK_EQ(slot_mapping.size(0), num_tokens);
+  TORCH_CHECK_EQ(slot_mapping.scalar_type(), c10::ScalarType::Long);

  TORCH_CHECK_EQ(kv_cache.size(2), kv_lora_rank + rot_dim);
  TORCH_CHECK_EQ(kv_cache.dim(), 3);
@@ -1,7 +1,6 @@
 #pragma once
 #include <cstdlib>
 #include <string>
-#include <cctype>

 namespace vllm {

@@ -1,5 +1,6 @@
 #include "scaled_mm_kernels.hpp"
 #include "scaled_mm_sm100_fp8_dispatch.cuh"
+#include "core/batch_invariant.hpp"

 namespace vllm {

@@ -13,9 +14,17 @@ void cutlass_scaled_mm_sm100_fp8(
    STD_TORCH_CHECK(bias->scalar_type() == out.scalar_type(),
                    "currently bias dtype must match output dtype ",
                    out.scalar_type());
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm100_fp8_batch_invariant_epilogue<true>(
+          out, a, b, a_scales, b_scales, *bias);
+    }
    return cutlass_scaled_mm_sm100_fp8_epilogue<true>(out, a, b, a_scales,
                                                      b_scales, *bias);
  } else {
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm100_fp8_batch_invariant_epilogue<false>(
+          out, a, b, a_scales, b_scales);
+    }
    return cutlass_scaled_mm_sm100_fp8_epilogue<false>(out, a, b, a_scales,
                                                       b_scales);
  }
@@ -294,6 +294,34 @@ inline void cutlass_gemm_sm100_fp8_dispatch(
  }
 }

+template <typename InType, typename OutType, bool EnableBias,
+          typename... EpilogueArgs>
+inline void cutlass_gemm_sm100_fp8_batch_invariant_dispatch(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, torch::stable::Tensor const& a_scales,
+    torch::stable::Tensor const& b_scales, EpilogueArgs&&... args) {
+  static_assert(std::is_same<InType, cutlass::float_e4m3_t>());
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  using Cutlass3xGemmM64SwapAB =
+      typename sm100_fp8_config_M64_swap_ab<InType, OutType,
+                                            EnableBias>::Cutlass3xGemm;
+  using Cutlass3xGemmM64 =
+      typename sm100_fp8_config_M64<InType, OutType, EnableBias>::Cutlass3xGemm;
+
+  // keep the CUTLASS config independent of M for batch invariance
+  uint32_t const k = a.size(1);
+  if (k < 4096) {
+    return cutlass_gemm_caller_sm100_fp8<Cutlass3xGemmM64>(
+        out, a, b, a_scales, b_scales, std::forward<EpilogueArgs>(args)...);
+  }
+  return cutlass_gemm_caller_sm100_fp8<Cutlass3xGemmM64SwapAB>(
+      out, a, b, b_scales, a_scales, std::forward<EpilogueArgs>(args)...);
+}
+
 template <bool EnableBias, typename... EpilogueArgs>
 void cutlass_scaled_mm_sm100_fp8_epilogue(torch::stable::Tensor& out,
                                          torch::stable::Tensor const& a,
@@ -320,4 +348,28 @@ void cutlass_scaled_mm_sm100_fp8_epilogue(torch::stable::Tensor& out,
  }
 }

+template <bool EnableBias, typename... EpilogueArgs>
+void cutlass_scaled_mm_sm100_fp8_batch_invariant_epilogue(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, torch::stable::Tensor const& a_scales,
+    torch::stable::Tensor const& b_scales, EpilogueArgs&&... epilogue_args) {
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  if (out.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
+    return cutlass_gemm_sm100_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::bfloat16_t, EnableBias>(
+        out, a, b, a_scales, b_scales,
+        std::forward<EpilogueArgs>(epilogue_args)...);
+  } else {
+    STD_TORCH_CHECK(out.scalar_type() == torch::headeronly::ScalarType::Half);
+    return cutlass_gemm_sm100_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::half_t, EnableBias>(
+        out, a, b, a_scales, b_scales,
+        std::forward<EpilogueArgs>(epilogue_args)...);
+  }
+}
+
 }  // namespace vllm
@@ -1,5 +1,6 @@
 #include "scaled_mm_kernels.hpp"
 #include "scaled_mm_sm120_fp8_dispatch.cuh"
+#include "core/batch_invariant.hpp"
 #include "cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp"

 namespace vllm {
@@ -14,9 +15,17 @@ void cutlass_scaled_mm_sm120_fp8(
    STD_TORCH_CHECK(bias->scalar_type() == out.scalar_type(),
                    "currently bias dtype must match output dtype ",
                    out.scalar_type());
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm120_fp8_batch_invariant_epilogue<
+          c3x::ScaledEpilogueBias>(out, a, b, a_scales, b_scales, *bias);
+    }
    return cutlass_scaled_mm_sm120_fp8_epilogue<c3x::ScaledEpilogueBias>(
        out, a, b, a_scales, b_scales, *bias);
  } else {
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm120_fp8_batch_invariant_epilogue<
+          c3x::ScaledEpilogue>(out, a, b, a_scales, b_scales);
+    }
    return cutlass_scaled_mm_sm120_fp8_epilogue<c3x::ScaledEpilogue>(
        out, a, b, a_scales, b_scales);
  }
@@ -179,6 +179,26 @@ inline void cutlass_gemm_sm120_fp8_dispatch(torch::stable::Tensor& out,
      out, a, b, std::forward<EpilogueArgs>(args)...);
 }

+template <typename InType, typename OutType,
+          template <typename, typename, typename> typename Epilogue,
+          typename... EpilogueArgs>
+inline void cutlass_gemm_sm120_fp8_batch_invariant_dispatch(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, EpilogueArgs&&... args) {
+  static_assert(std::is_same<InType, cutlass::float_e4m3_t>());
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  using Cutlass3xGemmM64 =
+      typename sm120_fp8_config_M64<InType, OutType, Epilogue>::Cutlass3xGemm;
+
+  // keep the CUTLASS config independent of M for batch invariance
+  return cutlass_gemm_caller<Cutlass3xGemmM64>(
+      out, a, b, std::forward<EpilogueArgs>(args)...);
+}
+
 template <template <typename, typename, typename> typename Epilogue,
          typename... EpilogueArgs>
 void cutlass_scaled_mm_sm120_fp8_epilogue(torch::stable::Tensor& out,
@@ -202,4 +222,26 @@ void cutlass_scaled_mm_sm120_fp8_epilogue(torch::stable::Tensor& out,
  }
 }

+template <template <typename, typename, typename> typename Epilogue,
+          typename... EpilogueArgs>
+void cutlass_scaled_mm_sm120_fp8_batch_invariant_epilogue(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, EpilogueArgs&&... epilogue_args) {
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  if (out.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
+    return cutlass_gemm_sm120_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::bfloat16_t, Epilogue>(
+        out, a, b, std::forward<EpilogueArgs>(epilogue_args)...);
+  } else {
+    STD_TORCH_CHECK(out.scalar_type() == torch::headeronly::ScalarType::Half);
+    return cutlass_gemm_sm120_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::half_t, Epilogue>(
+        out, a, b, std::forward<EpilogueArgs>(epilogue_args)...);
+  }
+}
+
 }  // namespace vllm
@@ -1,5 +1,6 @@
 #include "scaled_mm_kernels.hpp"
 #include "scaled_mm_sm90_fp8_dispatch.cuh"
+#include "core/batch_invariant.hpp"

 namespace vllm {

@@ -13,9 +14,17 @@ void cutlass_scaled_mm_sm90_fp8(
    STD_TORCH_CHECK(bias->scalar_type() == out.scalar_type(),
                    "currently bias dtype must match output dtype ",
                    out.scalar_type());
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm90_fp8_batch_invariant_epilogue<true>(
+          out, a, b, a_scales, b_scales, *bias);
+    }
    return cutlass_scaled_mm_sm90_fp8_epilogue<true>(out, a, b, a_scales,
                                                     b_scales, *bias);
  } else {
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm90_fp8_batch_invariant_epilogue<false>(
+          out, a, b, a_scales, b_scales);
+    }
    return cutlass_scaled_mm_sm90_fp8_epilogue<false>(out, a, b, a_scales,
                                                      b_scales);
  }
@@ -349,6 +349,35 @@ inline void cutlass_gemm_sm90_fp8_dispatch(
  }
 }

+template <typename InType, typename OutType, bool EnableBias,
+          typename... EpilogueArgs>
+inline void cutlass_gemm_sm90_fp8_batch_invariant_dispatch(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, torch::stable::Tensor const& a_scales,
+    torch::stable::Tensor const& b_scales, EpilogueArgs&&... args) {
+  static_assert(std::is_same<InType, cutlass::float_e4m3_t>());
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  using Cutlass3xGemmM64_N1280 =
+      typename sm90_fp8_config_M64_N1280<InType, OutType,
+                                         EnableBias>::Cutlass3xGemm;
+  using Cutlass3xGemmM64_N8192 =
+      typename sm90_fp8_config_M64_N8192<InType, OutType,
+                                         EnableBias>::Cutlass3xGemm;
+
+  // keep the CUTLASS config independent of M for batch invariance
+  uint32_t const n = b.size(1);
+  if (n <= 1280) {
+    return cutlass_gemm_caller_sm90_fp8<Cutlass3xGemmM64_N1280>(
+        out, a, b, b_scales, a_scales, std::forward<EpilogueArgs>(args)...);
+  }
+  return cutlass_gemm_caller_sm90_fp8<Cutlass3xGemmM64_N8192>(
+      out, a, b, b_scales, a_scales, std::forward<EpilogueArgs>(args)...);
+}
+
 template <bool EnableBias, typename... EpilogueArgs>
 void cutlass_scaled_mm_sm90_fp8_epilogue(torch::stable::Tensor& out,
                                         torch::stable::Tensor const& a,
@@ -375,4 +404,28 @@ void cutlass_scaled_mm_sm90_fp8_epilogue(torch::stable::Tensor& out,
  }
 }

+template <bool EnableBias, typename... EpilogueArgs>
+void cutlass_scaled_mm_sm90_fp8_batch_invariant_epilogue(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, torch::stable::Tensor const& a_scales,
+    torch::stable::Tensor const& b_scales, EpilogueArgs&&... epilogue_args) {
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  if (out.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
+    return cutlass_gemm_sm90_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::bfloat16_t, EnableBias>(
+        out, a, b, a_scales, b_scales,
+        std::forward<EpilogueArgs>(epilogue_args)...);
+  } else {
+    STD_TORCH_CHECK(out.scalar_type() == torch::headeronly::ScalarType::Half);
+    return cutlass_gemm_sm90_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::half_t, EnableBias>(
+        out, a, b, a_scales, b_scales,
+        std::forward<EpilogueArgs>(epilogue_args)...);
+  }
+}
+
 }  // namespace vllm
@@ -9,6 +9,7 @@
 #include "scaled_mm_c2x_sm89_fp8_dispatch.cuh"
 #include "scaled_mm_c2x_sm89_int8_dispatch.cuh"

+#include "core/batch_invariant.hpp"
 #include "libtorch_stable/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp"

 using namespace vllm;
@@ -191,9 +192,17 @@ void cutlass_scaled_mm_sm89(torch::stable::Tensor& out,
    STD_TORCH_CHECK(bias->scalar_type() == out.scalar_type(),
                    "currently bias dtype must match output dtype ",
                    out.scalar_type());
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm89_fp8_batch_invariant_epilogue<
+          c2x::ScaledEpilogueBias>(out, a, b, a_scales, b_scales, *bias);
+    }
    return cutlass_scaled_mm_sm89_epilogue<c2x::ScaledEpilogueBias>(
        out, a, b, a_scales, b_scales, *bias);
  } else {
+    if (vllm_is_batch_invariant()) {
+      return cutlass_scaled_mm_sm89_fp8_batch_invariant_epilogue<
+          c2x::ScaledEpilogue>(out, a, b, a_scales, b_scales);
+    }
    return cutlass_scaled_mm_sm89_epilogue<c2x::ScaledEpilogue>(
        out, a, b, a_scales, b_scales);
  }
@@ -381,4 +381,43 @@ inline void cutlass_gemm_sm89_fp8_dispatch(torch::stable::Tensor& out,
  }
 }

+template <typename InType, typename OutType,
+          template <typename, typename> typename Epilogue,
+          typename... EpilogueArgs>
+inline void cutlass_gemm_sm89_fp8_batch_invariant_dispatch(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, EpilogueArgs&&... args) {
+  static_assert(std::is_same<InType, cutlass::float_e4m3_t>());
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  // keep the CUTLASS config independent of M for batch invariance
+  return sm89_fp8_config_M64::dispatch<InType, OutType, Epilogue>(
+      out, a, b, std::forward<EpilogueArgs>(args)...);
+}
+
+template <template <typename, typename> typename Epilogue,
+          typename... EpilogueArgs>
+void cutlass_scaled_mm_sm89_fp8_batch_invariant_epilogue(
+    torch::stable::Tensor& out, torch::stable::Tensor const& a,
+    torch::stable::Tensor const& b, EpilogueArgs&&... epilogue_args) {
+  STD_TORCH_CHECK(a.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+  STD_TORCH_CHECK(b.scalar_type() ==
+                  torch::headeronly::ScalarType::Float8_e4m3fn);
+
+  if (out.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
+    return cutlass_gemm_sm89_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::bfloat16_t, Epilogue>(
+        out, a, b, std::forward<EpilogueArgs>(epilogue_args)...);
+  } else {
+    STD_TORCH_CHECK(out.scalar_type() == torch::headeronly::ScalarType::Half);
+    return cutlass_gemm_sm89_fp8_batch_invariant_dispatch<
+        cutlass::float_e4m3_t, cutlass::half_t, Epilogue>(
+        out, a, b, std::forward<EpilogueArgs>(epilogue_args)...);
+  }
+}
+
 }  // namespace vllm
@@ -1,7 +1,7 @@
 #include <cstdint>

 /* Adapted from ./csrc/quantization/gguf/mmq.cuh
-   based on ./vllm/model_executor/layers/fused_moe/fused_moe.py */
+   based on ./vllm/model_executor/layers/fused_moe/experts/triton_moe.py */
 template <typename scalar_t, int qk, int qr, int qi, bool need_sum,
          typename block_q_t, int mmq_x, int mmq_y, int nwarps,
          allocate_tiles_cuda_t allocate_tiles, load_tiles_cuda_t load_tiles,
@@ -0,0 +1,204 @@
+#include <Python.h>
+
+extern "C" {
+
+#include <stdbool.h>
+#include <time.h>
+
+#if defined(__i386__) || defined(__x86_64__)
+  #include <cpuid.h>
+  #include <mwaitxintrin.h>
+#endif
+
+#if defined(CLOCK_MONOTONIC_RAW)
+  #define TIMEOUT_CLOCK CLOCK_MONOTONIC_RAW
+#else
+  #define TIMEOUT_CLOCK CLOCK_MONOTONIC
+#endif
+
+#define CPU_SUPPORT_NONE 0
+#define CPU_SUPPORT_MONITORX 1
+
+#define MWAITX_DEFAULT_TIMEOUT_CYCLES 1000000
+
+typedef struct {
+  unsigned int cpu_support;
+  unsigned int max_monitor_line_size;
+} spinloop_state_t;
+
+static void determine_cpu_support(spinloop_state_t* state) {
+  state->cpu_support = CPU_SUPPORT_NONE;
+  state->max_monitor_line_size = 0;
+
+#if defined(__i386__) || defined(__x86_64__)
+  unsigned int eax, ebx, ecx, edx;
+  if (__get_cpuid(0, &eax, &ebx, &ecx, &edx) == 1) {
+    // AMD CPU (possible monitorx/mwaitx support)
+    if (ebx == 0x68747541 && edx == 0x69746e65 && ecx == 0x444d4163) {
+      if (__get_cpuid(0x80000000, &eax, &ebx, &ecx, &edx) == 1 &&
+          eax >= 0x80000001 &&
+          __get_cpuid(0x80000001, &eax, &ebx, &ecx, &edx) == 1) {
+        if ((ecx & (1 << 29)) != 0) {
+          state->cpu_support = CPU_SUPPORT_MONITORX;
+        }
+      }
+    }
+  }
+
+  if (state->cpu_support == CPU_SUPPORT_MONITORX) {
+    if (__get_cpuid(5, &eax, &ebx, &ecx, &edx) == 1) {
+      state->max_monitor_line_size = ebx & 0xff;
+    }
+  }
+#endif
+}
+
+static PyObject* method_spinloop(PyObject* self, PyObject* args,
+                                 PyObject* kwargs) {
+  Py_buffer buffer;
+  PyObject* callback;
+  double timeout = 0.;
+
+  spinloop_state_t* state = (spinloop_state_t*)PyModule_GetState(self);
+  if (state == NULL) {
+    PyErr_SetString(PyExc_TypeError, "Failed to retrieve module state!");
+    return NULL;
+  }
+
+  static const char* keywords[] = {"buffer", "callback", "timeout", NULL};
+  if (!PyArg_ParseTupleAndKeywords(args, kwargs, "y*O|d", (char**)keywords,
+                                   &buffer, &callback, &timeout)) {
+    return NULL;
+  }
+
+  if (!PyCallable_Check(callback)) {
+    PyErr_SetString(PyExc_TypeError, "callback parameter must be callable!");
+    PyBuffer_Release(&buffer);
+    return NULL;
+  }
+
+  struct timespec t_start;
+  if (clock_gettime(TIMEOUT_CLOCK, &t_start) != 0) {
+    PyErr_SetString(PyExc_RuntimeError, "clock_gettime() failed!");
+    PyBuffer_Release(&buffer);
+    return NULL;
+  }
+
+  bool result = false;
+  bool error = false;
+  bool have_timeout = (timeout > 1e-9);
+  unsigned int iteration = 0;
+  const bool buffer_qualifies = (buffer.len <= state->max_monitor_line_size);
+
+  while (true) {
+    PyObject* res = PyObject_CallNoArgs(callback);
+    if (res == NULL) {
+      error = true;
+      break;
+    }
+    int ok = (res == Py_True);
+    Py_DECREF(res);
+
+    if (ok) {
+      result = true;
+      break;
+    }
+
+    // Check timeout at most every 16 iterations to avoid clock_gettime and
+    // comparison cost
+    if (have_timeout && (iteration & 15u) == 0) {
+      struct timespec t_now;
+      if (clock_gettime(TIMEOUT_CLOCK, &t_now) != 0) {
+        PyErr_SetString(PyExc_RuntimeError, "clock_gettime() failed!");
+        error = true;
+        break;
+      }
+
+      const double elapsed = (double)(t_now.tv_sec - t_start.tv_sec) +
+                             (t_now.tv_nsec - t_start.tv_nsec) * 1e-9;
+      if (elapsed >= timeout) {
+        result = false;
+        break;
+      }
+    }
+    ++iteration;
+
+#if defined(__i386__) || defined(__x86_64__)
+    // monitorx + mwaitx with qualified buffer
+    if (buffer_qualifies && state->cpu_support == CPU_SUPPORT_MONITORX) {
+      _mm_monitorx(buffer.buf, 0, 0);
+
+      // Check once more in case the buffer has been modified while we were
+      // arming the monitor hardware
+      res = PyObject_CallNoArgs(callback);
+      if (res == NULL) {
+        error = true;
+        break;
+      }
+      ok = (res == Py_True);
+      Py_DECREF(res);
+
+      if (ok) {
+        result = true;
+        break;
+      }
+
+      // Run mwaitx with enabled timeout (bit 1). The actual timeout value
+      // is not very important, we just want to ensure we don't lock up
+      // here for too long.
+      Py_BEGIN_ALLOW_THREADS _mm_mwaitx((1 << 1), 0,
+                                        MWAITX_DEFAULT_TIMEOUT_CYCLES);
+      Py_END_ALLOW_THREADS
+    }
+
+    // Fallback: Busy poll
+    else {
+#endif
+      // Give other threads a chance to be scheduled
+      Py_BEGIN_ALLOW_THREADS
+#if defined(__i386__) || defined(__x86_64__)
+      __builtin_ia32_pause();
+#elif defined(__aarch64__)
+        __asm__ volatile("yield" :: : "memory");
+#endif
+      Py_END_ALLOW_THREADS
+#if defined(__i386__) || defined(__x86_64__)
+    }
+#endif
+  }
+
+  PyBuffer_Release(&buffer);
+
+  if (error) {
+    return NULL;
+  }
+
+  if (result) {
+    Py_RETURN_TRUE;
+  }
+
+  Py_RETURN_FALSE;
+}
+
+static PyMethodDef spinloop_methods[] = {
+    {"spinloop", (PyCFunction)method_spinloop, METH_VARARGS | METH_KEYWORDS,
+     "Wait for store with callback"},
+    {NULL, NULL, 0, NULL}};
+
+static struct PyModuleDef spinloop_module = {
+    PyModuleDef_HEAD_INIT, "spinloop",
+    "Hardware-optimized spinloops for Python", sizeof(spinloop_state_t),
+    spinloop_methods};
+
+PyMODINIT_FUNC PyInit_spinloop(void) {
+  PyObject* m = PyModule_Create(&spinloop_module);
+  if (m != NULL) {
+    spinloop_state_t* state = (spinloop_state_t*)PyModule_GetState(m);
+    if (state != NULL) {
+      determine_cpu_support(state);
+    }
+  }
+  return m;
+}
+
+}  // extern "C"
@@ -39,7 +39,8 @@ th {
 | Spec Bench | ✅ | ✅ | `wget https://raw.githubusercontent.com/hemingkx/Spec-Bench/refs/heads/main/data/spec_bench/question.jsonl` |
 | SPEED-Bench | ✅ | ✅ | `curl -LsSf https://raw.githubusercontent.com/NVIDIA-NeMo/Skills/refs/heads/main/nemo_skills/dataset/speed-bench/prepare.py \| python3 -` |
 | Custom | ✅ | ✅ | Local file: `data.jsonl` |
-| Custom MM | ✅ | ✅ | Local file: `mm_data.jsonl` |
+| Custom Audio | ✅ | ✅ | Local file: `audio_data.jsonl` |
+| Custom Image | ✅ | ✅ | Local file: `image_data.jsonl` |

 Legend:

@@ -142,7 +143,7 @@ Example output: ![Dataset Statistics](../assets/contributing/vllm_bench_serve_da

 #### Custom Dataset

-If the dataset you want to benchmark is not supported yet in vLLM, even then you can benchmark on it using `CustomDataset`. Your data needs to be in `.jsonl` format and needs to have "prompt" field per entry, e.g., data.jsonl
+If the dataset you want to benchmark is not supported yet in vLLM, even then you can benchmark on it using `CustomDataset`. At inference time, use the option `--dataset-name custom`. Your data needs to be in the `.jsonl` format and needs to have "prompt" field per entry, e.g., data.jsonl

 ```json
 {"prompt": "What is the capital of India?"}
@@ -173,9 +174,62 @@ vllm bench serve --port 9001 --save-result --save-detailed \

 You can skip applying chat template if your data already has it by using `--custom-skip-chat-template`.

-#### Custom multimodal dataset
+#### Custom Audio Dataset

-If the multimodal dataset you want to benchmark is not supported yet in vLLM, then you can benchmark on it using `CustomMMDataset`. Your data needs to be in `.jsonl` format and needs to have "prompt" and "image_files" field per entry, e.g., `mm_data.jsonl`:
+If the audio dataset you want to benchmark is not supported yet in vLLM, then you can benchmark on it using `CustomAudioDataset`. At inference time, use the option `--dataset-name custom_audio`. Your data needs to be in the `.jsonl` format and needs to have "prompt" and "audio" fields per entry, e.g., `audio_data.jsonl`:
+
+```json
+{"prompt": "What does this audio say?", "audio": "/path/to/audio_1.wav"}
+{"prompt": "Transcribe the audio.", "audio": "/path/to/audio_2.wav"}
+```
+
+- **Supported models:** The `CustomAudioDataset` class supports two types of audio models: ASR models (e.g. Whisper) which do not require a "prompt" field; and multimodal audio-text chat models (e.g. Qwen2-Audio). Since these model types require different arguments at inference, we are giving two examples.
+
+- **Example 1: Whisper**
+
+Whisper is a dedicated ASR encoder-decoder model, so it uses `--backend openai-audio` and `--endpoint /v1/audio/transcriptions`.
+
+```bash
+# start server
+vllm serve openai/whisper-tiny
+```
+
+```bash
+vllm bench serve \
+  --model openai/whisper-tiny \
+  --backend openai-audio \
+  --endpoint /v1/audio/transcriptions \
+  --dataset-name custom_audio \
+  --dataset-path audio_data.jsonl \
+  --no-oversample \
+  --custom-output-len 256 \
+  --save-result \
+  --save-detailed \
+  --result-filename whisper_bench.json
+```
+
+- **Example 2: Qwen2-Audio**
+
+Qwen2-Audio is a multimodal chat model that can do ASR and speech analysis, so it uses `--backend openai-chat`, and `--endpoint /v1/chat/completions`. It also requires `--enable-multimodal-chat` to enable multimodal chat transformation.
+
+```bash
+vllm bench serve \
+  --model Qwen/Qwen2-Audio-7B-Instruct \
+  --backend openai-chat \
+  --endpoint /v1/chat/completions \
+  --dataset-name custom_audio \
+  --dataset-path audio_data.jsonl \
+  --no-oversample \
+  --custom-output-len 256 \
+  --enable-multimodal-chat \
+  --save-result \
+  --save-detailed \
+  --result-filename qwen_bench.json
+```
+
+#### Custom Image Dataset
+
+If the image dataset you want to benchmark is not supported yet in vLLM, then you can benchmark on it using `CustomImageDataset`. At inference time, use the option `--dataset-name custom_image`. Your data needs to be in the `.jsonl` format and needs to have "prompt" and "image_files" fields per entry, e.g., `image_data.jsonl`:

 ```json
 {"prompt": "How many animals are present in the given image?", "image_files": ["/path/to/image/folder/horsepony.jpg"]}
@@ -193,8 +247,8 @@ vllm bench serve--save-result --save-detailed \
  --backend openai-chat \
  --model Qwen/Qwen2-VL-7B-Instruct \
  --endpoint /v1/chat/completions \
-  --dataset-name custom_mm \
-  --dataset-path <path-to-your-mm-data-jsonl> \
+  --dataset-name custom_image \
+  --dataset-path <path-to-your-image-data-jsonl> \
  --allowed-local-media-path /path/to/image/folder
 ```

@@ -142,7 +142,7 @@ We use "mamba-like" to refer to layers that possess a state that is updated in-p
 For implementing new custom mamba-like layers, one should inherit from `MambaBase` and implement the methods `get_state_dtype`, `get_state_shape` to calculate the data types and state shapes at runtime, as well as `mamba_type` and `get_attn_backend`.
 It is also necessary to implement the "attention meta-data" class which handles the meta-data that is common across all layers.
 Please see [`LinearAttentionMetadata`](../../../vllm/v1/attention/backends/linear_attn.py) or [`ShortConvAttentionMetadata`](../../../vllm/v1/attention/backends/short_conv_attn.py) for examples of this.
-It is also worth noting that we should update `MAMBA_TYPE_TO_BACKEND_MAP` and `MambaAttentionBackendEnum` in [`registry.py`](../../../vllm/v1/attention/backends/registry.py) when adding a new mamba backend.
+It is also worth noting that we should update `MambaAttentionBackendEnum` in [`registry.py`](../../../vllm/v1/attention/backends/registry.py) when adding a new mamba backend.
 Finally, if one wants to support torch compile and CUDA graphs, it necessary to wrap the call to the mamba-like layer inside a custom op and register it.
 Please see the calls to `direct_register_custom_op` in [vllm/model_executor/models/minimax_text_01.py](../../../vllm/model_executor/models/minimax_text_01.py) or [vllm/model_executor/layers/mamba/short_conv.py](../../../vllm/model_executor/layers/mamba/short_conv.py) for examples of this.
 The new custom op should then be added to the list `_attention_ops` in [vllm/config/compilation.py](../../../vllm/config/compilation.py) to ensure that piecewise CUDA graphs works as intended.
@@ -138,7 +138,7 @@ For example:

 --8<-- "vllm/model_executor/models/transformers/moe.py:transformers_fused_moe"

--8<-- "vllm/model_executor/layers/fused_moe/fused_moe.py:grouped_topk"
+--8<-- "vllm/model_executor/layers/fused_moe/router/grouped_topk_router.py:grouped_topk"
 ```

 **9. Norm:**
@@ -80,14 +80,14 @@ To be used with a particular `FusedMoEPrepareAndFinalizeModular` subclass, MoE k

 | Kernel | Input act. format | Quant. types | Quant. format | Activation function | Apply Weight On Input | Modular | Source |
 | ------ | ----------------- | ------------ | ------------- | ------------------- | --------------------- | ------- | ------ |
-| triton | standard | all<sup>1</sup> | G,A,T | silu, gelu,</br>swigluoai,</br>silu_no_mul,</br>gelu_no_mul | Y | Y | [`fused_experts`][vllm.model_executor.layers.fused_moe.fused_moe.fused_experts],</br>[`TritonExperts`][vllm.model_executor.layers.fused_moe.fused_moe.TritonExperts] |
+| triton | standard | all<sup>1</sup> | G,A,T | silu, gelu,</br>swigluoai,</br>silu_no_mul,</br>gelu_no_mul | Y | Y | [`fused_experts`][vllm.model_executor.layers.fused_moe.fused_moe.fused_experts],</br>[`TritonExperts`][vllm.model_executor.layers.fused_moe.experts.triton_moe.TritonExperts] |
 | triton (batched) | batched | all<sup>1</sup> | G,A,T | silu, gelu | <sup>6</sup> | Y | [`BatchedTritonExperts`][vllm.model_executor.layers.fused_moe.fused_batched_moe.BatchedTritonExperts] |
 | deep gemm | standard,</br>batched | fp8 | G(128),A,T | silu, gelu | <sup>6</sup> | Y | </br>[`DeepGemmExperts`][vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe.DeepGemmExperts],</br>[`BatchedDeepGemmExperts`][vllm.model_executor.layers.fused_moe.experts.batched_deep_gemm_moe.BatchedDeepGemmExperts] |
 | cutlass_fp4 | standard,</br>batched | nvfp4 | A,T | silu | Y | Y | [`CutlassExpertsFp4`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassExpertsFp4] |
 | cutlass_fp8 | standard,</br>batched | fp8 | A,T | silu, gelu | Y | Y | [`CutlassExpertsFp8`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassExpertsFp8],</br>[`CutlasBatchedExpertsFp8`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassBatchedExpertsFp8] |
-| flashinfer | standard | nvfp4,</br>fp8 | T | <sup>5</sup> | N | Y | [`FlashInferExperts`][vllm.model_executor.layers.fused_moe.flashinfer_cutlass_moe.FlashInferExperts] |
+| flashinfer | standard | nvfp4,</br>fp8 | T | <sup>5</sup> | N | Y | [`FlashInferExperts`][vllm.model_executor.layers.fused_moe.experts.flashinfer_cutlass_moe.FlashInferExperts] |
 | gpt oss triton | standard | N/A | N/A | <sup>5</sup> | Y | Y | [`triton_kernel_fused_experts`][vllm.model_executor.layers.fused_moe.experts.gpt_oss_triton_kernels_moe.triton_kernel_fused_experts],</br>[`OAITritonExperts`][vllm.model_executor.layers.fused_moe.experts.gpt_oss_triton_kernels_moe.OAITritonExperts] |
-| marlin | standard,</br>batched | <sup>3</sup> / N/A | <sup>3</sup> / N/A | silu,</br>swigluoai | Y | Y | [`fused_marlin_moe`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.fused_marlin_moe],</br>[`MarlinExperts`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.MarlinExperts],</br>[`BatchedMarlinExperts`][vllm.model_executor.layers.fused_moe.fused_marlin_moe.BatchedMarlinExperts] |
+| marlin | standard,</br>batched | <sup>3</sup> / N/A | <sup>3</sup> / N/A | silu,</br>swigluoai | Y | Y | [`fused_marlin_moe`][vllm.model_executor.layers.fused_moe.experts.marlin_moe.fused_marlin_moe],</br>[`MarlinExperts`][vllm.model_executor.layers.fused_moe.experts.marlin_moe.MarlinExperts],</br>[`BatchedMarlinExperts`][vllm.model_executor.layers.fused_moe.experts.marlin_moe.BatchedMarlinExperts] |
 | trtllm | standard | mxfp4,</br>nvfp4 | G(16),G(32) | <sup>5</sup> | N | Y | [`TrtLlmMxfp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsMonolithic],</br>[`TrtLlmMxfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsModular],</br>[`TrtLlmNvFp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsMonolithic],</br>[`TrtLlmNvfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsModular] |
 | rocm aiter moe | standard | mxfp4,</br>fp8 | G(32),G(128),A,T | silu, gelu,</br>swigluoai | Y | N | `rocm_aiter_fused_experts`,</br>`AiterExperts` |
 | cpu_fused_moe | standard | N/A | N/A | silu | N | N | [`CPUFusedMOE`][vllm.model_executor.layers.fused_moe.cpu_fused_moe.CPUFusedMOE] |
@@ -126,9 +126,15 @@ python tests/v1/kv_connector/nixl_integration/toy_proxy_server.py \
    - Set when prefiller and decoder are on different machines
    - Connection info is passed via KVTransferParams from prefiller to decoder for handshake

- `VLLM_NIXL_ABORT_REQUEST_TIMEOUT`: Timeout (in seconds) for automatically releasing the prefiller’s KV cache for a particular request. (Optional)
+- `kv_lease_duration` (via `kv_connector_extra_config`): Lease duration (in seconds) for the prefiller's KV cache blocks. (Optional)
+    - Default: 30
+    - When a prefill request finishes, its KV blocks are held for this duration waiting for the decoder to read them. While the request is queued on the decoder, periodic heartbeats automatically extend the lease. If neither a heartbeat nor a read notification arrives before the lease expires, the blocks are freed. The heartbeat interval and extension amount are derived automatically from this value.
+    - Example: `--kv-transfer-config '{"kv_connector_extra_config": {"kv_lease_duration": 60}}'`
+
+- `decoder_kv_blocks_ttl` (via `kv_connector_extra_config`): TTL (in seconds) for KV blocks cached on the decoder in bidirectional transfer mode. (Optional)
    - Default: 480
-    - If a request is aborted and the decoder has not yet read the KV-cache blocks through the nixl channel, the prefill instance will release its KV-cache blocks after this timeout to avoid holding them indefinitely.
+    - In bidirectional mode, the decoder caches KV blocks for multi-turn conversations. This TTL controls how long those blocks are held before being released. Unlike the prefiller lease, this TTL is not renewed via heartbeats.
+    - Example: `--kv-transfer-config '{"kv_connector_extra_config": {"decoder_kv_blocks_ttl": 600}}'`

 ## Multi-Instance Setup

@@ -3,9 +3,10 @@
 vLLM supports the following hardware platforms:

 - [GPU](gpu.md)
-    - [NVIDIA CUDA](gpu.md#nvidia-cuda)
-    - [AMD ROCm](gpu.md#amd-rocm)
-    - [Intel XPU](gpu.md#intel-xpu)
+    - [NVIDIA CUDA](gpu.md)
+    - [AMD ROCm](gpu.md)
+    - [Intel XPU](gpu.md)
+    - [Apple Silicon](gpu.md) (via [vLLM-Metal](https://github.com/vllm-project/vllm-metal))
 - [CPU](cpu.md)
    - [Intel/AMD x86](cpu.md#intelamd-x86)
    - [ARM AArch64](cpu.md#arm-aarch64)
@@ -0,0 +1,125 @@
+<!-- markdownlint-disable MD041 -->
+--8<-- [start:installation]
+
+For GPU-accelerated inference on Apple Silicon, use [vLLM-Metal](https://github.com/vllm-project/vllm-metal), a community-maintained hardware plugin that uses MLX as the compute backend and provides native GPU acceleration via Apple's Metal framework.
+
+vLLM-Metal works with MLX-optimized models from the [mlx-community](https://huggingface.co/mlx-community) organization on Hugging Face, which provides quantized versions of popular models optimized for Apple Silicon.
+
+!!! tip
+    For installation and usage instructions, see the [Set up using vLLM-Metal](#set-up-using-vllm-metal) section below.
+
+--8<-- [end:installation]
+--8<-- [start:requirements]
+
+- OS: macOS Sonoma or later
+- Hardware: Apple Silicon
+- Metal support enabled
+
+!!! note
+    See the [Set up using vLLM-Metal](#set-up-using-vllm-metal) section below for installation instructions.
+
+--8<-- [end:requirements]
+--8<-- [start:set-up-using-python]
+
+## Set up using vLLM-Metal
+
+vLLM-Metal is distributed as a separate package that provides native GPU acceleration on Apple Silicon.
+
+To install vLLM-Metal, follow the installation instructions in the [vLLM-Metal documentation](https://github.com/vllm-project/vllm-metal#installation).
+
+The installation will:
+
+1. Set up the appropriate Python environment
+2. Install MLX and required dependencies
+3. Install the vLLM-Metal package
+
+After installation, you can start using vLLM with Metal GPU acceleration.
+
+!!! tip
+    When using vLLM-Metal, use models from the [mlx-community](https://huggingface.co/mlx-community) on Hugging Face for best performance. These models are optimized for MLX and often include quantized versions (4-bit, 8-bit) that run efficiently on Apple Silicon.
+
+    Example model: `mlx-community/Qwen2.5-0.5B-Instruct-4bit`
+
+### Using vLLM-Metal
+
+After installation, vLLM-Metal provides an easy-to-use CLI for running an OpenAI-compatible API server:
+
+```bash
+# Activate the vLLM-Metal environment
+source ~/.venv-vllm-metal/bin/activate
+
+# Start the API server (specify your mlx-community model or it will use default)
+vllm serve
+```
+
+Once the server is running, you have multiple options to interact with it:
+
+#### Option 1: Interactive chat
+
+Open a new terminal and start an interactive chat session:
+
+```bash
+source ~/.venv-vllm-metal/bin/activate
+vllm chat
+```
+
+#### Option 2: API requests with curl
+
+```bash
+curl http://localhost:8000/v1/chat/completions \
+  -H "Content-Type: application/json" \
+  -d '{
+    "messages": [{"role": "user", "content": "Hello!"}],
+    "max_tokens": 50
+  }'
+```
+
+#### Option 3: Python with OpenAI SDK
+
+```python
+from openai import OpenAI
+
+client = OpenAI(
+    base_url="http://localhost:8000/v1",
+    api_key="dummy"  # No auth required for local server
+)
+
+response = client.chat.completions.create(
+    model="mlx-community/Qwen2.5-0.5B-Instruct-4bit",
+    messages=[{"role": "user", "content": "Hello!"}]
+)
+
+print(response.choices[0].message.content)
+```
+
+For more details on the `vllm` CLI commands, see the [OpenAI-compatible server documentation](../../serving/openai_compatible_server.md).
+
+--8<-- [end:set-up-using-python]
+--8<-- [start:pre-built-wheels]
+
+vLLM-Metal is installed via the vLLM-Metal package. See the [Set up using vLLM-Metal](#set-up-using-vllm-metal) section above.
+
+--8<-- [end:pre-built-wheels]
+--8<-- [start:build-wheel-from-source]
+
+For build instructions from source, refer to the [vLLM-Metal documentation](https://github.com/vllm-project/vllm-metal#installation).
+
+--8<-- [end:build-wheel-from-source]
+--8<-- [start:pre-built-images]
+
+--8<-- [end:pre-built-images]
+--8<-- [start:build-image-from-source]
+
+--8<-- [end:build-image-from-source]
+--8<-- [start:supported-features]
+
+vLLM-Metal provides:
+
+- Native GPU acceleration using Metal
+- MLX-based compute backend optimized for Apple Silicon
+- OpenAI-compatible API server
+- Support for popular model architectures
+
+For specific feature support and limitations, refer to the [vLLM-Metal documentation](https://github.com/vllm-project/vllm-metal).
+
+--8<-- [end:supported-features]
@@ -18,6 +18,10 @@ vLLM is a Python library that supports the following GPU variants. Select your G

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:installation"

+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:installation"
+
 ## Requirements

 - OS: Linux
@@ -38,6 +42,10 @@ vLLM is a Python library that supports the following GPU variants. Select your G

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:requirements"

+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:requirements"
+
 ## Set up using Python

 ### Create a new Python environment
@@ -56,6 +64,10 @@ vLLM is a Python library that supports the following GPU variants. Select your G

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:set-up-using-python"

+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:set-up-using-python"
+
 ### Pre-built wheels {#pre-built-wheels}

 === "NVIDIA CUDA"
@@ -70,6 +82,10 @@ vLLM is a Python library that supports the following GPU variants. Select your G

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:pre-built-wheels"

+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:pre-built-wheels"
+
 ### Build wheel from source

 === "NVIDIA CUDA"
@@ -84,6 +100,10 @@ vLLM is a Python library that supports the following GPU variants. Select your G

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:build-wheel-from-source"

+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:build-wheel-from-source"
+
 ## Set up using Docker

 ### Pre-built images
@@ -102,6 +122,10 @@ vLLM is a Python library that supports the following GPU variants. Select your G

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:pre-built-images"

+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:pre-built-images"
+
 --8<-- [end:pre-built-images]

 ### Build image from source
@@ -120,6 +144,10 @@ vLLM is a Python library that supports the following GPU variants. Select your G

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:build-image-from-source"

+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:build-image-from-source"
+
 --8<-- [end:build-image-from-source]

 ## Supported features
@@ -135,3 +163,7 @@ vLLM is a Python library that supports the following GPU variants. Select your G
 === "Intel XPU"

    --8<-- "docs/getting_started/installation/gpu.xpu.inc.md:supported-features"
+
+=== "Apple Silicon"
+
+    --8<-- "docs/getting_started/installation/gpu.apple.inc.md:supported-features"
@@ -10,6 +10,9 @@ This guide will help you quickly get started with vLLM to perform:
 - OS: Linux
 - Python: 3.10 -- 3.13

+!!! note
+    vLLM also works on macOS with [vLLM-Metal](https://github.com/vllm-project/vllm-metal) for Apple Silicon GPU acceleration. See the [GPU installation guide](installation/gpu.md) and select the "Apple Silicon" tab.
+
 ## Installation

 === "NVIDIA CUDA"
@@ -73,6 +76,18 @@ This guide will help you quickly get started with vLLM to perform:
    !!! note
        For more detailed instructions, including Docker, installing from source, and troubleshooting, please refer to the [vLLM on TPU documentation](https://docs.vllm.ai/projects/tpu/en/latest/).

+=== "Apple Silicon (Mac)"
+
+    If you are using Apple Silicon Macs, you can use vLLM-Metal for GPU-accelerated inference via Apple's Metal framework.
+
+    Follow the installation instructions in the [vLLM-Metal documentation](https://github.com/vllm-project/vllm-metal#installation).
+
+    !!! note
+        vLLM-Metal uses MLX instead of PyTorch as the compute backend and requires MLX-optimized models from the [mlx-community](https://huggingface.co/mlx-community) on Hugging Face.
+
+    !!! tip
+        For more detailed instructions, please refer to the [GPU installation guide](installation/gpu.md) and select the "Apple Silicon" tab.
+
 !!! note
    For more detail and non-CUDA platforms, please refer to the [installation guide](installation/README.md) for specific instructions on how to install vLLM.

@@ -385,7 +385,7 @@ th {
 | `DeepseekForCausalLM` | DeepSeek | `deepseek-ai/deepseek-llm-67b-base`, `deepseek-ai/deepseek-llm-7b-chat`, etc. | ✅︎ | ✅︎ |
 | `DeepseekV2ForCausalLM` | DeepSeek-V2 | `deepseek-ai/DeepSeek-V2`, `deepseek-ai/DeepSeek-V2-Chat`, etc. | ✅︎ | ✅︎ |
 | `DeepseekV3ForCausalLM` | DeepSeek-V3 | `deepseek-ai/DeepSeek-V3`, `deepseek-ai/DeepSeek-R1`, `deepseek-ai/DeepSeek-V3.1`, etc. | ✅︎ | ✅︎ |
-| `DeepseekV4ForCausalLM` | DeepSeek-V4 | `deepseek-ai/DeepSeek-V4-Flash`, `deepseek-ai/DeepSeek-V4-Pro`, etc. | | |
+| `DeepseekV4ForCausalLM` | DeepSeek-V4 | `deepseek-ai/DeepSeek-V4-Flash`, `deepseek-ai/DeepSeek-V4-Pro`, etc. | | ✅︎ |
 | `Dots1ForCausalLM` | dots.llm1 | `rednote-hilab/dots.llm1.base`, `rednote-hilab/dots.llm1.inst`, etc. | | ✅︎ |
 | `DotsOCRForCausalLM` | dots_ocr | `rednote-hilab/dots.ocr` | ✅︎ | ✅︎ |
 | `Ernie4_5ForCausalLM` | Ernie4.5 | `baidu/ERNIE-4.5-0.3B-PT`, etc. | ✅︎ | ✅︎ |
@@ -598,7 +598,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `MiniMaxVL01ForConditionalGeneration` | MiniMax-VL | T + I<sup>E+</sup> | `MiniMaxAI/MiniMax-VL-01`, etc. | | ✅︎ |
 | `Mistral3ForConditionalGeneration` | Mistral3 (HF Transformers) | T + I<sup>+</sup> | `mistralai/Mistral-Small-3.1-24B-Instruct-2503`, etc. | ✅︎ | ✅︎ |
 | `MolmoForCausalLM` | Molmo | T + I<sup>+</sup> | `allenai/Molmo-7B-D-0924`, `allenai/Molmo-7B-O-0924`, etc. | ✅︎ | ✅︎ |
-| `Molmo2ForConditionalGeneration` | Molmo2 | T + I<sup>+</sup> / V | `allenai/Molmo2-4B`, `allenai/Molmo2-8B`, `allenai/Molmo2-O-7B` | ✅︎ | ✅︎ |
+| `Molmo2ForConditionalGeneration` | Molmo2 | T + I<sup>+</sup> / V | `allenai/Molmo2-4B`, `allenai/Molmo2-8B`, `allenai/Molmo2-O-7B`, `allenai/MolmoWeb-4B`<sup>^</sup>, `allenai/MolmoWeb-8B`<sup>^</sup> | ✅︎ | ✅︎ |
 | `Moondream3ForCausalLM` | Moondream3 | T + I | `moondream/moondream3-preview` | | ✅︎ |
 | `MusicFlamingoForConditionalGeneration` | MusicFlamingo | T + A | `nvidia/music-flamingo-2601-hf`, `nvidia/music-flamingo-think-2601-hf` | ✅︎ | ✅︎ |
 | `NVLM_D_Model` | NVLM-D 1.0 | T + I<sup>+</sup> | `nvidia/NVLM-D-72B`, etc. | | ✅︎ |
@@ -666,6 +666,11 @@ Some models are supported only via the [Transformers modeling backend](#transfor
 !!! note
    For `InternVLChatModel`, only InternVL2.5 with Qwen2.5 text backbone (`OpenGVLab/InternVL2.5-1B` etc.), InternVL3 and InternVL3.5 have video inputs support currently.

+!!! note
+    To use `allenai/MolmoWeb-4B` or `allenai/MolmoWeb-8B`, serve the checkpoint
+    with the Molmo2 architecture and disable multimodal-prefix attention:
+    `--hf-overrides '{"architectures": ["Molmo2ForConditionalGeneration"], "is_mm_prefix_lm": false}'`.
+
 !!! note
    `Moondream3ForCausalLM` uses task-specific prompt templates for `query`
    and `caption`. The native `detect` and `point` skills require custom
@@ -333,6 +333,43 @@ Most cache paths default to subdirectories under a single root. Changing `VLLM_C
 - **Do not copy cache contents from untrusted sources.** If you distribute cache artifacts between environments, ensure they originate from a trusted build pipeline.
 - **Container deployments:** If mounting cache directories into containers, ensure the volume source is trusted.

+## FIPS Compatibility
+
+FIPS compliance depends on many factors, so a vLLM deployment is not automatically FIPS compliant. Recent changes have improved vLLM's *tolerance* of FIPS-enabled hosts — that is, avoiding crashes when non-approved algorithms are blocked — but tolerance is not the same as compliance. Whether a deployment satisfies FIPS requirements depends on the host operating system, the OpenSSL provider backing Python's `hashlib` and `ssl` modules, and which optional dependencies are installed.
+
+### FIPS-relevant configuration
+
+Operators running vLLM on FIPS-enabled hosts should select FIPS-approved algorithms via the following knobs:
+
+- **Multimodal input hashing** — `VLLM_MM_HASHER_ALGORITHM` defaults to `blake3`, which is not FIPS-approved. Set it to `sha256` or `sha512` in FIPS-enabled environments.
+- **Prefix-cache hashing** — set `--prefix-caching-hash-algo` (config field `prefix_caching_hash_algo`) to `sha256` or `sha256_cbor`. The `xxhash` and `xxhash_cbor` options are not FIPS-approved.
+- **TLS ciphers** — use `--ssl-ciphers` to restrict the API server's TLS handshake to FIPS-approved cipher suites that match your environment's policy.
+
+### Automatic fallback for non-security MD5 use
+
+vLLM uses MD5 in a few places to derive non-security cache keys (for example, configuration hashes). These call sites pass `usedforsecurity=False` and additionally fall back to SHA-256 when the underlying OpenSSL provider refuses MD5 outright (see `safe_hash()` in `vllm/utils/hashing.py`). No user action is required; this behavior is documented so that auditors and security reviewers can identify the MD5 references and understand their purpose.
+
+### Dependencies that provide non-FIPS hash implementations
+
+Some dependencies expose hash implementations that are not FIPS-approved. vLLM only invokes them when the corresponding algorithm is selected, but operators with strict cryptographic controls may want to ensure the code paths are not exercised — and, where policy requires, that the packages themselves are absent:
+
+- `blake3` — currently listed in `requirements/common.txt`, so a standard install pulls it in. It is imported lazily and only used when `VLLM_MM_HASHER_ALGORITHM=blake3` (the default). Setting `VLLM_MM_HASHER_ALGORITHM` to `sha256` or `sha512` is sufficient to keep the non-FIPS code path dormant. If your policy additionally forbids the package being present, uninstall it after `pip install` (`pip uninstall blake3`); vLLM will continue to function as long as `VLLM_MM_HASHER_ALGORITHM` is set to a non-blake3 value.
+- `xxhash` — a true optional dependency (not in `requirements/common.txt`). It is only imported when an `xxhash`-based prefix-cache algorithm is selected. Leave it uninstalled and select a `sha256`-based prefix-cache algorithm.
+
+### Beyond hashing: other FIPS considerations
+
+Hashing is the area where vLLM has explicit FIPS-aware code, but a FIPS-compliant deployment depends on several factors that sit outside vLLM itself. Operators should evaluate the following with their platform and security teams:
+
+- **Host crypto provider.** Python's `hashlib` and `ssl` modules are FIPS-aware only when Python is linked against a FIPS-validated OpenSSL (or equivalent) provider supplied by the host OS. vLLM inherits whatever provider the host configures — it does not bundle one.
+- **API server TLS.** TLS termination for the OpenAI-compatible API server uses the host's OpenSSL via Python's `ssl` module. Restrict the cipher suite with `--ssl-ciphers` to match your environment's FIPS policy, and ensure server certificates are issued with FIPS-approved algorithms and key sizes.
+- **Outbound HTTPS.** Model and asset downloads (for example, via `huggingface_hub`) use the same host TLS stack. The same provider/cipher considerations apply.
+- **Inter-node communication is unencrypted by default.** As described in [Inter-Node Communication](#inter-node-communication), PyTorch Distributed, KV-cache transfer, and data-parallel channels do not encrypt traffic. FIPS environments that require FIPS-approved cryptography for data in transit must provide that protection externally — for example, via an mTLS sidecar or IPsec terminated by a FIPS-validated module — since vLLM's internal channels cannot satisfy the requirement on their own. Network isolation alone is not cryptography and does not meet a "FIPS-approved cryptography for data in transit" requirement, though it remains a useful defense-in-depth measure.
+- **Dependencies that bundle their own OpenSSL.** Some Python wheels statically link OpenSSL builds that fail the kernel FIPS self-test on FIPS-enabled hosts (`FATAL FIPS SELFTEST FAILURE`). `opencv-python-headless` is a known example; other manylinux wheels may behave similarly. Audit your installed wheels for bundled crypto libraries when troubleshooting FIPS startup failures.
+- **Accelerator and ML libraries.** PyTorch, CUDA, cuDNN, NCCL, and similar components have their own crypto and FIPS posture independent of vLLM. NVIDIA publishes FIPS-validated builds for some libraries; vLLM does not pin to those builds, so selecting and validating them is the operator's responsibility.
+- **What is *not* a FIPS concern in vLLM.** Random number generation used for token sampling (Python/NumPy/PyTorch RNGs) is not a cryptographic use and is out of scope for FIPS. Pickled cache artifacts are a separate security concern covered under [Cache Directory Security](#cache-directory-security).
+
+In short: the configuration knobs above let vLLM avoid non-approved algorithms, and the automatic fallbacks let it run without crashing on FIPS-enabled hosts. End-to-end FIPS compliance, however, is a property of the full deployment — host OS, crypto provider, transitive dependencies, and network architecture — not of vLLM alone.
+
 ## Reporting Security Vulnerabilities

 If you believe you have found a security vulnerability in vLLM, please report it following the project's security policy. For more information on how to report security issues and the project's security policy, please see the [vLLM Security Policy](https://github.com/vllm-project/vllm/blob/main/SECURITY.md).
@@ -263,7 +263,7 @@
            {%- if message.get('tool_responses') -%}
                {#- Legacy: tool_responses embedded on the assistant message (Google/Gemma native) -#}
                {%- for tool_response in message['tool_responses'] -%}
-                    {{- format_tool_response_block(tool_response['name'] | default('unknown'), tool_response['response']) -}}
+                    {{- format_tool_response_block(tool_response['name'] | default('unknown', true), tool_response['response']) -}}
                    {%- set ns_tr_out.flag = true -%}
                    {%- set ns.prev_message_type = 'tool_response' -%}
                {%- endfor -%}
@@ -277,7 +277,7 @@
                    {%- else -%}
                        {%- set follow = loop_messages[k] -%}
                        {#- Resolve tool_call_id to function name -#}
-                        {%- set ns_tname = namespace(name=follow.get('name') | default('unknown')) -%}
+                        {%- set ns_tname = namespace(name=follow.get('name') | default('unknown', true)) -%}
                        {%- for tc in message['tool_calls'] -%}
                            {%- if tc.get('id') == follow.get('tool_call_id') -%}
                                {%- set ns_tname.name = tc['function']['name'] -%}
@@ -123,7 +123,8 @@ python = "./.venv"
 extend-exclude = ["tests/models/fixtures/*", "tests/prompts/*", "tests/tokenizers_/*",
    "benchmarks/sonnet.txt", "tests/lora/data/*", "build/*",
    "examples/pooling/token_embed/*", "tests/models/language/pooling/*",
-    "vllm/third_party/*", "vllm/entrypoints/serve/instrumentator/static/*", "tests/entrypoints/openai/speech_to_text/test_transcription_validation.py",
+    "vllm/third_party/*", "vllm/entrypoints/serve/instrumentator/static/*",
+    "tests/entrypoints/speech_to_text/transcription/test_transcription_validation.py",
    "docs/governance/process.md", "docs/assets/contributing/vllm_bench_serve_timeline.html", 
    "tests/v1/engine/test_fast_incdec_prefix_err.py", ".git/*", "csrc/cpu/sgl-kernels/*"]
 ignore-hidden = false
@@ -686,6 +686,7 @@ class precompiled_wheel_utils:
                    "vllm/vllm_flash_attn/_vllm_fa2_C.abi3.so",
                    "vllm/vllm_flash_attn/_vllm_fa3_C.abi3.so",
                    "vllm/cumem_allocator.abi3.so",
+                    "vllm/spinloop.abi3.so",
                    # ROCm-specific libraries
                    "vllm/_rocm_C.abi3.so",
                ]
@@ -993,6 +994,8 @@ if _is_cuda() or _is_hip():
    # copying the relevant .py files from the source repository.
    ext_modules.append(CMakeExtension(name="vllm.triton_kernels", optional=True))

+ext_modules.append(CMakeExtension(name="vllm.spinloop"))
+
 if _is_hip():
    ext_modules.append(CMakeExtension(name="vllm._rocm_C"))

@@ -0,0 +1,166 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Unit tests for the DoubleQuant fan-out variants registered by
+``RocmAiterRMSNormQuantFusionPass``.
+
+Both variants target a 1-to-2 fan-out where one ``rms_norm`` output feeds
+two distinct ``rocm_aiter_group_fp8_quant`` consumers and rewrite it into
+two independent fused ``rms_norm + group_fp8_quant`` ops:
+
+* ``DoubleAiterRMSFp8GroupQuantPattern`` matches the un-viewed shape
+  (e.g. Kimi-K2.5 / DSR1).
+* ``DoubleAiterRMSFp8GroupQuantViewPattern`` (this PR) is the view-tolerant
+  sibling that additionally matches the
+  ``rms_norm -> view -> group_fp8_quant`` shape that DSv3.2's MLA indexer
+  q_c norm exposes through ``Fp8BlockScaledMMLinearKernel.apply_weights``'s
+  2D-flatten boilerplate.
+"""
+
+import pytest
+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.compilation.passes.utility.noop_elimination import NoOpEliminationPass
+from vllm.compilation.passes.utility.post_cleanup import PostCleanupPass
+from vllm.config import (
+    CompilationConfig,
+    CompilationMode,
+    ModelConfig,
+    PassConfig,
+    VllmConfig,
+)
+
+EPS = 1e-5
+HIDDEN_SIZE = 256
+GROUP_SIZE = 128
+
+
+class _NoViewDoubleQuantModel(torch.nn.Module):
+    """``rms_norm -> 2x group_fp8_quant`` fan-out (Kimi-K2.5 / DSR1 shape)."""
+
+    def __init__(self) -> None:
+        super().__init__()
+        self.weight = torch.nn.Parameter(torch.ones(HIDDEN_SIZE, dtype=torch.bfloat16))
+
+    def forward(
+        self, x: torch.Tensor
+    ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
+        # avoid graph input being a direct arg to a matched pattern node
+        x = torch.relu(x)
+        rms = torch.ops.vllm_ir.rms_norm(x, self.weight, EPS)
+        q1, s1 = torch.ops.vllm.rocm_aiter_group_fp8_quant.default(rms, GROUP_SIZE)
+        q2, s2 = torch.ops.vllm.rocm_aiter_group_fp8_quant.default(rms, GROUP_SIZE)
+        return q1, s1, q2, s2
+
+
+class _ViewDoubleQuantModel(torch.nn.Module):
+    """``rms_norm -> view -> 2x group_fp8_quant`` fan-out (DSv3.2 shape).
+
+    Reproduces the FX-graph shape produced by ``Fp8BlockScaledMMLinearKernel``'s
+    2D-flatten before the FP8 group quant op.
+    """
+
+    def __init__(self) -> None:
+        super().__init__()
+        self.weight = torch.nn.Parameter(torch.ones(HIDDEN_SIZE, dtype=torch.bfloat16))
+
+    def forward(
+        self, x: torch.Tensor
+    ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
+        x = torch.relu(x)
+        rms = torch.ops.vllm_ir.rms_norm(x, self.weight, EPS)
+        view = rms.view(-1, rms.shape[-1])
+        q1, s1 = torch.ops.vllm.rocm_aiter_group_fp8_quant.default(view, GROUP_SIZE)
+        q2, s2 = torch.ops.vllm.rocm_aiter_group_fp8_quant.default(view, GROUP_SIZE)
+        return q1, s1, q2, s2
+
+
+@pytest.mark.parametrize(
+    "model_cls",
+    [_NoViewDoubleQuantModel, _ViewDoubleQuantModel],
+    ids=["no_view", "with_view"],
+)
+@pytest.mark.skipif(
+    not is_aiter_found_and_supported(),
+    reason="Only test on ROCm with AITER installed and supported",
+)
+def test_double_aiter_rms_fp8_group_quant_fusion(
+    model_cls: type[torch.nn.Module],
+    monkeypatch: pytest.MonkeyPatch,
+) -> None:
+    """
+    Both fan-out shapes (with and without an intermediate view) must fuse
+    into ``rocm_aiter_rmsnorm_fp8_group_quant``: the no-view shape via
+    ``DoubleAiterRMSFp8GroupQuantPattern`` and the viewed shape via the
+    new ``DoubleAiterRMSFp8GroupQuantViewPattern`` sibling.
+
+    A failure on the ``with_view`` parametrization is a regression on the
+    DSv3.2 q_c norm path that this PR's view-tolerant pattern is intended
+    to cover.
+    """
+    torch._dynamo.reset()
+
+    vllm_config = VllmConfig(
+        model_config=ModelConfig(dtype=torch.bfloat16),
+        compilation_config=CompilationConfig(
+            mode=CompilationMode.VLLM_COMPILE,
+            custom_ops=["+rms_norm", "+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,
+        )
+
+        torch.set_default_device("cuda")
+        torch.set_default_dtype(torch.bfloat16)
+        torch.manual_seed(0)
+
+        m.setenv("VLLM_ROCM_USE_AITER", "1")
+        rocm_aiter_ops.refresh_env_variables()
+
+        fusion_pass = RocmAiterRMSNormQuantFusionPass(vllm_config)
+        passes = [
+            NoOpEliminationPass(vllm_config),
+            fusion_pass,
+            PostCleanupPass(vllm_config),
+        ]
+        backend = TestBackend(*passes)
+        model = model_cls()
+
+        x = torch.randn(8, HIDDEN_SIZE)
+        torch._dynamo.mark_dynamic(x, 0)
+
+        outputs_unfused = model(x)
+        model_fused = torch.compile(model, backend=backend)
+        outputs_fused = model_fused(x)
+
+        # Both consumers must be rewritten into the fused op (one
+        # ``register_replacement`` rewrite covers the whole 1-to-2 fan-out).
+        assert fusion_pass.matched_count == 1, (
+            f"Expected the {model_cls.__name__} fan-out to fuse via the "
+            f"DoubleQuant pattern (matched_count == 1), got "
+            f"{fusion_pass.matched_count}"
+        )
+
+        fused_op = rocm_aiter_ops.get_rmsnorm_group_fused_quant_op()
+        backend.check_after_ops([fused_op])
+
+        # Numerical parity sanity-check: the fused pair must match the
+        # unfused pair on FP8 outputs (exact byte-equality is the goal,
+        # but allow a tiny tolerance for any residual numeric noise).
+        for fused_t, unfused_t in zip(outputs_fused, outputs_unfused):
+            torch.testing.assert_close(
+                fused_t.to(torch.float32),
+                unfused_t.to(torch.float32),
+                atol=1e-2,
+                rtol=1e-2,
+            )
@@ -0,0 +1,413 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import torch
+
+import vllm.config
+from tests.compile.backend import TestBackend
+from tests.v1.attention.utils import BatchSpec, create_common_attn_metadata
+from vllm._aiter_ops import is_aiter_found_and_supported, rocm_aiter_ops
+from vllm.compilation.passes.fusion.mla_rope_kvcache_cat_fusion import (
+    MLARoPEKVCacheCatFusionPass,
+)
+from vllm.compilation.passes.utility.fix_functionalization import (
+    FixFunctionalizationPass,
+)
+from vllm.compilation.passes.utility.noop_elimination import NoOpEliminationPass
+from vllm.compilation.passes.utility.post_cleanup import PostCleanupPass
+from vllm.config import (
+    CacheConfig,
+    CompilationConfig,
+    CompilationMode,
+    ModelConfig,
+    PassConfig,
+    VllmConfig,
+)
+from vllm.forward_context import get_forward_context, set_forward_context
+from vllm.model_executor.layers.attention import MLAAttention
+from vllm.model_executor.layers.linear import ColumnParallelLinear
+from vllm.model_executor.layers.rotary_embedding import (
+    DeepseekScalingRotaryEmbedding,
+    RotaryEmbedding,
+)
+from vllm.platforms import current_platform
+from vllm.utils.torch_utils import _encode_layer_name
+from vllm.v1.attention.backend import (
+    AttentionBackend,
+    CommonAttentionMetadata,
+)
+from vllm.v1.attention.backends.fa_utils import flash_attn_supports_mla
+from vllm.v1.attention.backends.registry import AttentionBackendEnum
+
+INDEX_SELECT_OP = torch.ops.aten.index.Tensor
+VLLM_UNIFIED_MLA_KV_CACHE_UPDATE_OP = torch.ops.vllm.unified_mla_kv_cache_update
+FP8_DTYPE = current_platform.fp8_dtype()
+
+
+class MLARoPEKVCacheCatTestModel(torch.nn.Module):
+    def __init__(
+        self,
+        vllm_config: VllmConfig,
+        attn_backend: AttentionBackendEnum,
+        use_deepseek_scaling_rope: bool,
+        num_heads: int,
+        qk_nope_head_dim: int,
+        qk_rope_head_dim: int,
+        v_head_dim: int,
+        q_lora_rank: int,
+        kv_lora_rank: int,
+        is_neox: bool,
+        dtype: torch.dtype,
+        device: torch.device,
+        prefix: str = "model.layers.0.self_attn.attn",
+    ):
+        super().__init__()
+        self.num_heads = num_heads
+        self.qk_nope_head_dim = qk_nope_head_dim
+        self.qk_rope_head_dim = qk_rope_head_dim
+        self.qk_head_dim = qk_nope_head_dim + qk_rope_head_dim
+        self.v_head_dim = v_head_dim
+        self.q_lora_rank = q_lora_rank
+        self.kv_lora_rank = kv_lora_rank
+        self.dtype = dtype
+        self.device = device
+        self.layer_name = prefix
+
+        self.num_kv_heads = 1
+        self.head_size = kv_lora_rank + qk_rope_head_dim
+        self.block_size = vllm_config.cache_config.block_size
+        self.scale = self.qk_head_dim**-0.5
+
+        if use_deepseek_scaling_rope:
+            self.rotary_emb = DeepseekScalingRotaryEmbedding(
+                head_size=qk_rope_head_dim,
+                rotary_dim=qk_rope_head_dim,
+                max_position_embeddings=4096,
+                base=10000,
+                is_neox_style=is_neox,
+                scaling_factor=1.0,
+                dtype=dtype,
+            )
+        else:
+            self.rotary_emb = RotaryEmbedding(
+                head_size=qk_rope_head_dim,
+                rotary_dim=qk_rope_head_dim,
+                max_position_embeddings=4096,
+                base=10000,
+                is_neox_style=is_neox,
+                dtype=dtype,
+            )
+
+        # Initialize intermediate mm layers for unit test
+        self.q_b_proj = ColumnParallelLinear(
+            self.q_lora_rank,
+            self.num_heads * self.qk_head_dim,
+            bias=False,
+            prefix=f"{prefix}.q_b_proj",
+        ).to(device)
+        self.kv_b_proj = ColumnParallelLinear(
+            self.kv_lora_rank,
+            self.num_heads * (self.qk_nope_head_dim + self.v_head_dim),
+            bias=False,
+            prefix=f"{prefix}.kv_b_proj",
+        ).to(device)
+
+        # ColumnParallelLinear default init in bf16 with seed 0 produces
+        # near-zero weights (7/4.7M nonzero), making the GEMM output almost
+        # entirely zero and masking correctness bugs. Reinitialize to get
+        # dense outputs.
+        with torch.no_grad():
+            torch.nn.init.normal_(self.q_b_proj.weight, std=0.02)
+            torch.nn.init.normal_(self.kv_b_proj.weight, std=0.02)
+
+        # Register layer metadata for the fusion pass via MLAAttention
+        self.mla_attn = MLAAttention(
+            num_heads=self.num_heads,
+            scale=self.scale,
+            qk_nope_head_dim=self.qk_nope_head_dim,
+            qk_rope_head_dim=self.qk_rope_head_dim,
+            v_head_dim=self.v_head_dim,
+            q_lora_rank=self.q_lora_rank,
+            kv_lora_rank=self.kv_lora_rank,
+            kv_b_proj=self.kv_b_proj,
+            cache_config=vllm_config.cache_config,
+            quant_config=vllm_config.quant_config,
+            prefix=prefix,
+            attn_backend=attn_backend.get_class(),
+        )
+        self.attn_backend: type[AttentionBackend] = self.mla_attn.get_attn_backend()
+        self.mla_attn._k_scale = self.mla_attn._k_scale.to(device)
+        self.mla_attn._v_scale = self.mla_attn._v_scale.to(device)
+
+        # Keep both the string dtype (for ops) and torch dtype (for tensors)
+        self.kv_cache_dtype_str = vllm_config.cache_config.cache_dtype
+        self.kv_cache_dtype = (
+            FP8_DTYPE if self.kv_cache_dtype_str.startswith("fp8") else self.dtype
+        )
+
+        # Initialize attn MetadataBuilder
+        self.builder = self.attn_backend.get_builder_cls()(
+            kv_cache_spec=self.mla_attn.get_kv_cache_spec(vllm_config),
+            layer_names=[self.mla_attn.layer_name],
+            vllm_config=vllm_config,
+            device=device,
+        )
+
+    def build_attn_metadata(self, batch_size: int) -> CommonAttentionMetadata:
+        """Initialize attention metadata."""
+        # Create common attn metadata
+        batch_spec = BatchSpec(seq_lens=[1] * batch_size, query_lens=[1] * batch_size)
+        common_attn_metadata = create_common_attn_metadata(
+            batch_spec, self.block_size, self.device, arange_block_indices=True
+        )
+
+        max_blocks = (max(batch_spec.seq_lens) + self.block_size - 1) // self.block_size
+        num_blocks = batch_size * max_blocks
+
+        # Fetch the attention backend and kv cache shape and stride order
+        kv_cache_shape = self.attn_backend.get_kv_cache_shape(
+            num_blocks, self.block_size, self.num_kv_heads, self.head_size
+        )
+        try:
+            kv_cache_stride_order = self.attn_backend.get_kv_cache_stride_order()
+        except (AttributeError, NotImplementedError):
+            kv_cache_stride_order = tuple(range(len(kv_cache_shape)))
+
+        kv_cache_shape = tuple(kv_cache_shape[i] for i in kv_cache_stride_order)
+        inv_order = [
+            kv_cache_stride_order.index(i) for i in range(len(kv_cache_stride_order))
+        ]
+
+        raw_tensor = torch.zeros(
+            num_blocks * self.block_size * self.num_kv_heads * self.head_size,
+            dtype=self.kv_cache_dtype,
+            device=self.device,
+        )
+        raw_tensor = raw_tensor.view(kv_cache_shape)
+        kv_cache = raw_tensor.permute(*inv_order)
+
+        self.mla_attn.kv_cache = kv_cache
+
+        # Build attn metadata
+        attn_metadata = self.builder.build(
+            common_prefix_len=0, common_attn_metadata=common_attn_metadata
+        )
+
+        return attn_metadata
+
+    def forward(
+        self, qkv_lora: torch.Tensor, positions: torch.Tensor
+    ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
+        qkv_lora = qkv_lora.clone()
+        q_c, kv_lora = qkv_lora.split(
+            [self.q_lora_rank, self.kv_lora_rank + self.qk_rope_head_dim],
+            dim=-1,
+        )
+        q = self.q_b_proj(q_c)[0]
+        kv_c, k_pe = kv_lora.split([self.kv_lora_rank, self.qk_rope_head_dim], dim=-1)
+
+        q = q.view(-1, self.num_heads, self.qk_head_dim)
+        k_pe = k_pe.unsqueeze(1)
+
+        q[..., self.qk_nope_head_dim :], k_pe = self.rotary_emb(
+            positions, q[..., self.qk_nope_head_dim :], k_pe
+        )
+
+        dummy = torch.ops.vllm.unified_mla_kv_cache_update(
+            kv_c,
+            k_pe,
+            _encode_layer_name(self.layer_name),
+            self.kv_cache_dtype_str,
+            self.mla_attn._k_scale,
+        )
+        return q, kv_c, k_pe, dummy
+
+    def ops_in_model_before(self) -> list[torch._ops.OpOverload]:
+        ops = [
+            INDEX_SELECT_OP,
+            torch.ops.vllm.unified_mla_kv_cache_update.default,
+        ]
+        return ops
+
+    def ops_in_model_after(self) -> list[torch._ops.OpOverload]:
+        return [torch.ops.vllm.fused_rope_unified_mla_kv_cache_update.default]
+
+
+MLA_BACKENDS = [AttentionBackendEnum.TRITON_MLA]
+if flash_attn_supports_mla():
+    MLA_BACKENDS += [AttentionBackendEnum.FLASH_ATTN_MLA]
+if is_aiter_found_and_supported():
+    MLA_BACKENDS += [AttentionBackendEnum.ROCM_AITER_MLA]
+
+
+@pytest.mark.parametrize("attn_backend", MLA_BACKENDS)
+@pytest.mark.parametrize("use_deepseek_scaling_rope", [True])
+@pytest.mark.parametrize("num_heads", [16])
+@pytest.mark.parametrize("qk_nope_head_dim", [128])
+@pytest.mark.parametrize("qk_rope_head_dim", [64])
+@pytest.mark.parametrize("v_head_dim", [128])
+@pytest.mark.parametrize("q_lora_rank", [1536])
+@pytest.mark.parametrize("kv_lora_rank", [512])
+@pytest.mark.parametrize("block_size", [16])
+@pytest.mark.parametrize("is_neox", [True, False])
+@pytest.mark.parametrize("dtype", [torch.bfloat16])
+@pytest.mark.parametrize("kv_cache_dtype", ["auto", "fp8"])
+@pytest.mark.skipif(
+    not current_platform.is_cuda_alike(),
+    reason="MLA RoPE+KVCache+Cat fusion is only supported on CUDA and ROCm.",
+)
+def test_mla_rope_kvcache_cat_fusion(
+    attn_backend: AttentionBackendEnum,
+    use_deepseek_scaling_rope: bool,
+    num_heads: int,
+    qk_nope_head_dim: int,
+    qk_rope_head_dim: int,
+    v_head_dim: int,
+    q_lora_rank: int,
+    kv_lora_rank: int,
+    block_size: int,
+    is_neox: bool,
+    dtype: torch.dtype,
+    kv_cache_dtype: str,
+    monkeypatch: pytest.MonkeyPatch,
+):
+    torch.set_default_device("cuda")
+    torch.set_default_dtype(dtype)
+    torch.manual_seed(0)
+
+    vllm_config = VllmConfig(
+        model_config=ModelConfig(
+            model="deepseek-ai/DeepSeek-V2-Lite",
+            dtype=dtype,
+        ),
+        cache_config=CacheConfig(
+            block_size=block_size,
+            cache_dtype=kv_cache_dtype,
+        ),
+        compilation_config=CompilationConfig(
+            mode=CompilationMode.VLLM_COMPILE,
+            pass_config=PassConfig(
+                fuse_rope_kvcache_cat_mla=True,
+                eliminate_noops=True,
+            ),
+        ),
+    )
+
+    with vllm.config.set_current_vllm_config(vllm_config), monkeypatch.context() as m:
+        if not torch.distributed.is_initialized():
+            from vllm.distributed.parallel_state import (
+                init_distributed_environment,
+                initialize_model_parallel,
+            )
+            from vllm.utils.system_utils import update_environment_variables
+
+            update_environment_variables(
+                {
+                    "RANK": "0",
+                    "LOCAL_RANK": "0",
+                    "WORLD_SIZE": "1",
+                    "MASTER_ADDR": "localhost",
+                    "MASTER_PORT": "54321",
+                }
+            )
+            init_distributed_environment()
+            initialize_model_parallel()
+
+        if attn_backend == AttentionBackendEnum.ROCM_AITER_MLA:
+            m.setenv("VLLM_ROCM_USE_AITER", "1")
+            rocm_aiter_ops.refresh_env_variables()
+
+        model = MLARoPEKVCacheCatTestModel(
+            vllm_config=vllm_config,
+            attn_backend=attn_backend,
+            use_deepseek_scaling_rope=use_deepseek_scaling_rope,
+            num_heads=num_heads,
+            qk_nope_head_dim=qk_nope_head_dim,
+            qk_rope_head_dim=qk_rope_head_dim,
+            v_head_dim=v_head_dim,
+            q_lora_rank=q_lora_rank,
+            kv_lora_rank=kv_lora_rank,
+            is_neox=is_neox,
+            dtype=dtype,
+            device=torch.get_default_device(),
+        )
+
+        fusion_pass = MLARoPEKVCacheCatFusionPass(vllm_config)
+        # note: FixFunctionalizationPass is required to correctly lower
+        # the fused op to its inplace version with auto-functionalization v1.
+        # Without it, decompose_auto_functionalized calls clone_preserve_strides
+        # on the non-contiguous q_pe slice directly, and inductor's lowering
+        # of the resulting as_strided chain incorrectly drops the storage offset.
+        # auto-functionalization v2 avoids this: it clones the contiguous base
+        # tensor (_all_bases) and reconstructs the slice as a view, so the
+        # offset is never passed through as_strided lowering.
+        passes = [
+            NoOpEliminationPass(vllm_config),
+            fusion_pass,
+            PostCleanupPass(vllm_config),
+            FixFunctionalizationPass(vllm_config),
+        ]
+        backend = TestBackend(*passes)
+
+        T = 5
+
+        qkv_lora = torch.randn(
+            T,
+            q_lora_rank + kv_lora_rank + qk_rope_head_dim,
+            dtype=dtype,
+        )
+        pos = torch.arange(T, dtype=torch.long)
+
+        qkv_unfused = qkv_lora.clone()
+        pos_unfused = pos.clone()
+
+        # Run unfused version
+        with set_forward_context(None, vllm_config):
+            forward_context = get_forward_context()
+            attn_metadata = model.build_attn_metadata(T)
+            forward_context.slot_mapping = {
+                model.layer_name: attn_metadata.slot_mapping
+            }
+            q_unfused, kv_c_unfused, k_pe_unfused, dummy = model(
+                qkv_unfused, pos_unfused
+            )
+            attn_layer = forward_context.no_compile_layers[model.layer_name]
+            kv_cache_unfused = attn_layer.kv_cache.clone()
+        del dummy
+
+        # Run fused version (compiled)
+        torch._dynamo.mark_dynamic(qkv_lora, 0)
+        torch._dynamo.mark_dynamic(pos, 0)
+        with set_forward_context(None, vllm_config):
+            model_fused = torch.compile(model, backend=backend)
+            forward_context = get_forward_context()
+            attn_metadata = model.build_attn_metadata(T)
+            forward_context.slot_mapping = {
+                model.layer_name: attn_metadata.slot_mapping
+            }
+            q_fused, kv_c_fused, k_pe_fused, dummy = model_fused(qkv_lora, pos)
+            attn_layer = forward_context.no_compile_layers[model.layer_name]
+            kv_cache_fused = attn_layer.kv_cache
+        del dummy
+
+        assert fusion_pass.matched_count == 1
+
+        backend.check_before_ops(model.ops_in_model_before())
+        backend.check_after_ops(model.ops_in_model_after())
+
+        if dtype == torch.float16:
+            ATOL, RTOL = (2e-3, 2e-3)
+        else:
+            ATOL, RTOL = (1e-2, 1e-2)
+
+        torch.testing.assert_close(q_unfused, q_fused, atol=ATOL, rtol=RTOL)
+        torch.testing.assert_close(kv_c_unfused, kv_c_fused, atol=ATOL, rtol=RTOL)
+        torch.testing.assert_close(k_pe_unfused, k_pe_fused, atol=ATOL, rtol=RTOL)
+        # Cannot compare fp8_* directly here, cast to model dtype instead
+        torch.testing.assert_close(
+            kv_cache_unfused.view(dtype),
+            kv_cache_fused.view(dtype),
+            atol=ATOL,
+            rtol=RTOL,
+        )
@@ -34,7 +34,6 @@ from vllm.v1.attention.backend import (
    CommonAttentionMetadata,
 )
 from vllm.v1.attention.backends.registry import AttentionBackendEnum
-from vllm.v1.kv_cache_interface import AttentionSpec

 INDEX_SELECT_OP = torch.ops.aten.index.Tensor
 VLLM_UNIFIED_KV_CACHE_UPDATE_OP = torch.ops.vllm.unified_kv_cache_update
@@ -102,13 +101,8 @@ class QKRoPEKVCacheTestModel(torch.nn.Module):
        )

        # Initialize attn MetadataBuilder
-        self.builder = self.attn.attn_backend.get_builder_cls()(
-            kv_cache_spec=AttentionSpec(
-                block_size=self.block_size,
-                num_kv_heads=self.num_kv_heads,
-                head_size=head_size,
-                dtype=self.kv_cache_dtype,
-            ),
+        self.builder = self.attn_backend.get_builder_cls()(
+            kv_cache_spec=self.attn.get_kv_cache_spec(vllm_config),
            layer_names=[self.attn.layer_name],
            vllm_config=vllm_config,
            device=device,
@@ -126,12 +120,11 @@ class QKRoPEKVCacheTestModel(torch.nn.Module):
        num_blocks = batch_size * max_blocks

        # Fetch the attention backend and kv cache shape and stride order
-        attn_backend = self.attn.attn_backend
-        kv_cache_shape = attn_backend.get_kv_cache_shape(
+        kv_cache_shape = self.attn_backend.get_kv_cache_shape(
            num_blocks, self.block_size, self.num_kv_heads, self.head_size
        )
        try:
-            kv_cache_stride_order = attn_backend.get_kv_cache_stride_order()
+            kv_cache_stride_order = self.attn_backend.get_kv_cache_stride_order()
        except (AttributeError, NotImplementedError):
            kv_cache_stride_order = tuple(range(len(kv_cache_shape)))

@@ -67,58 +67,3 @@ async def test_chat_with_enable_force_include_usage(
            chunk.usage.prompt_tokens + chunk.usage.completion_tokens
        )
        last_completion_tokens = chunk.usage.completion_tokens
-
-
-@pytest.fixture(scope="module")
-def transcription_server_with_force_include_usage():
-    args = [
-        # use half precision for speed and memory savings in CI environment
-        "--dtype",
-        "bfloat16",
-        "--max-num-seqs",
-        "4",
-        "--enforce-eager",
-        "--enable-force-include-usage",
-        "--gpu-memory-utilization",
-        "0.2",
-    ]
-
-    with RemoteOpenAIServer("openai/whisper-large-v3-turbo", args) as remote_server:
-        yield remote_server
-
-
-@pytest_asyncio.fixture
-async def transcription_client_with_force_include_usage(
-    transcription_server_with_force_include_usage,
-):
-    async with (
-        transcription_server_with_force_include_usage.get_async_client() as async_client
-    ):
-        yield async_client
-
-
-@pytest.mark.asyncio
-async def test_transcription_with_enable_force_include_usage(
-    transcription_client_with_force_include_usage, winning_call
-):
-    res = (
-        await transcription_client_with_force_include_usage.audio.transcriptions.create(
-            model="openai/whisper-large-v3-turbo",
-            file=winning_call,
-            language="en",
-            temperature=0.0,
-            stream=True,
-            timeout=30,
-        )
-    )
-
-    async for chunk in res:
-        if not len(chunk.choices):
-            # final usage sent
-            usage = chunk.usage
-            assert isinstance(usage, dict)
-            assert usage["prompt_tokens"] > 0
-            assert usage["completion_tokens"] > 0
-            assert usage["total_tokens"] > 0
-        else:
-            assert not hasattr(chunk, "usage")
@@ -0,0 +1,492 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Cross-API render parity tests.
+
+Verifies that the chat completion input path (parse_chat_input_to_harmony_message)
+and the responses API input path (response_input_to_harmony) produce identical
+Harmony messages and identical rendered token sequences when given equivalent
+conversation representations.
+
+The chat completion API encodes reasoning and tool calls as fields on a single
+assistant message dict; the responses API encodes them as separate typed items
+in request.input. Both paths must converge on the same Harmony message list and
+therefore the same rendered prompt.
+
+Each test:
+  1. Builds Harmony messages from each path for a single message or sequence.
+  2. Asserts message-level properties (role, channel, recipient, content)
+     using verify_harmony_messages.
+  3. Asserts that render_for_completion produces identical token sequences.
+"""
+
+from openai.types.responses import ResponseFunctionToolCall
+
+from tests.entrypoints.openai.utils import verify_harmony_messages
+from vllm.entrypoints.openai.parser.harmony_utils import (
+    get_system_message,
+    parse_chat_input_to_harmony_message,
+    render_for_completion,
+)
+from vllm.entrypoints.openai.responses.harmony import response_input_to_harmony
+
+# Use a fixed date so the system message is deterministic across both paths.
+_DATE = "2025-01-01"
+
+
+def _system():
+    return get_system_message(start_date=_DATE)
+
+
+class TestResponseInputToHarmonyRenderParity:
+    """Each test drives the same conversation through both APIs and asserts
+    identical Harmony messages and rendered token sequences."""
+
+    # -----------------------------------------------------------------------
+    # Single-message cases
+    # -----------------------------------------------------------------------
+
+    def test_user_message(self):
+        chat_msgs = parse_chat_input_to_harmony_message(
+            {"role": "user", "content": "What's the weather in Paris?"}
+        )
+        resp_msgs = [
+            response_input_to_harmony(
+                {
+                    "type": "message",
+                    "role": "user",
+                    "content": "What's the weather in Paris?",
+                },
+                prev_responses=[],
+            )
+        ]
+
+        expected = [{"role": "user", "content": "What's the weather in Paris?"}]
+        verify_harmony_messages(chat_msgs, expected)
+        verify_harmony_messages(resp_msgs, expected)
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
+
+    def test_assistant_final_message(self):
+        chat_msgs = parse_chat_input_to_harmony_message(
+            {"role": "assistant", "content": "It is 18°C in Paris."}
+        )
+        resp_msgs = [
+            response_input_to_harmony(
+                {
+                    "type": "message",
+                    "role": "assistant",
+                    "content": "It is 18°C in Paris.",
+                },
+                prev_responses=[],
+            )
+        ]
+
+        expected = [
+            {"role": "assistant", "channel": "final", "content": "It is 18°C in Paris."}
+        ]
+        verify_harmony_messages(chat_msgs, expected)
+        verify_harmony_messages(resp_msgs, expected)
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
+
+    def test_reasoning_item(self):
+        # Chat path: assistant message with only a reasoning field and no content.
+        chat_msgs = parse_chat_input_to_harmony_message(
+            {
+                "role": "assistant",
+                "reasoning": "I should call get_weather.",
+                "content": "",
+            }
+        )
+        resp_msgs = [
+            response_input_to_harmony(
+                {
+                    "type": "reasoning",
+                    "content": [
+                        {"type": "reasoning_text", "text": "I should call get_weather."}
+                    ],
+                },
+                prev_responses=[],
+            )
+        ]
+
+        expected = [
+            {
+                "role": "assistant",
+                "channel": "analysis",
+                "content": "I should call get_weather.",
+            }
+        ]
+        verify_harmony_messages(chat_msgs, expected)
+        verify_harmony_messages(resp_msgs, expected)
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
+
+    def test_function_call(self):
+        chat_msgs = parse_chat_input_to_harmony_message(
+            {
+                "role": "assistant",
+                "tool_calls": [
+                    {
+                        "id": "call_1",
+                        "function": {
+                            "name": "get_weather",
+                            "arguments": '{"location": "Paris"}',
+                        },
+                    }
+                ],
+            }
+        )
+        resp_msgs = [
+            response_input_to_harmony(
+                {
+                    "type": "function_call",
+                    "name": "get_weather",
+                    "arguments": '{"location": "Paris"}',
+                },
+                prev_responses=[],
+            )
+        ]
+
+        expected = [
+            {
+                "role": "assistant",
+                "channel": "commentary",
+                "recipient": "functions.get_weather",
+                "content": '{"location": "Paris"}',
+                "content_type": "json",
+            }
+        ]
+        verify_harmony_messages(chat_msgs, expected)
+        verify_harmony_messages(resp_msgs, expected)
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
+
+    def test_tool_output(self):
+        prev_call = ResponseFunctionToolCall(
+            id="fc_1",
+            call_id="call_1",
+            name="get_weather",
+            arguments='{"location": "Paris"}',
+            type="function_call",
+        )
+
+        chat_msgs = parse_chat_input_to_harmony_message(
+            {"role": "tool", "tool_call_id": "call_1", "content": "18°C, clear skies."},
+            tool_id_names={"call_1": "get_weather"},
+        )
+        resp_msgs = [
+            response_input_to_harmony(
+                {
+                    "type": "function_call_output",
+                    "call_id": "call_1",
+                    "output": "18°C, clear skies.",
+                },
+                prev_responses=[prev_call],
+            )
+        ]
+
+        expected = [
+            {
+                "role": "tool",
+                "author_name": "functions.get_weather",
+                "channel": "commentary",
+                "recipient": "assistant",
+                "content": "18°C, clear skies.",
+            }
+        ]
+        verify_harmony_messages(chat_msgs, expected)
+        verify_harmony_messages(resp_msgs, expected)
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
+
+    # -----------------------------------------------------------------------
+    # Combined and multi-turn cases
+    # -----------------------------------------------------------------------
+
+    def test_reasoning_combined_with_function_call(self):
+        """Chat API packs reasoning + tool_calls into one dict; responses API
+        represents them as two separate items. Both must produce the same two
+        Harmony messages in the same order: analysis then commentary."""
+        chat_msgs = parse_chat_input_to_harmony_message(
+            {
+                "role": "assistant",
+                "reasoning": "I should get the weather for Paris.",
+                "tool_calls": [
+                    {
+                        "id": "call_1",
+                        "function": {
+                            "name": "get_weather",
+                            "arguments": '{"location": "Paris"}',
+                        },
+                    }
+                ],
+            }
+        )
+        resp_msgs = [
+            response_input_to_harmony(
+                {
+                    "type": "reasoning",
+                    "content": [
+                        {
+                            "type": "reasoning_text",
+                            "text": "I should get the weather for Paris.",
+                        }
+                    ],
+                },
+                prev_responses=[],
+            ),
+            response_input_to_harmony(
+                {
+                    "type": "function_call",
+                    "name": "get_weather",
+                    "arguments": '{"location": "Paris"}',
+                },
+                prev_responses=[],
+            ),
+        ]
+
+        expected = [
+            {
+                "role": "assistant",
+                "channel": "analysis",
+                "content": "I should get the weather for Paris.",
+            },
+            {
+                "role": "assistant",
+                "channel": "commentary",
+                "recipient": "functions.get_weather",
+                "content": '{"location": "Paris"}',
+                "content_type": "json",
+            },
+        ]
+        verify_harmony_messages(chat_msgs, expected)
+        verify_harmony_messages(resp_msgs, expected)
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
+
+    def test_full_multi_turn_tool_call_conversation(self):
+        """Full conversation: user -> reasoning + tool_call -> tool_output -> final.
+
+        Both APIs must render the complete conversation to identical token sequences.
+        This exercises the entire input pipeline including all message types and
+        the Rust harmony encoder.
+        """
+        prev_call = ResponseFunctionToolCall(
+            id="fc_1",
+            call_id="call_1",
+            name="get_weather",
+            arguments='{"location": "Paris"}',
+            type="function_call",
+        )
+
+        # --- Chat completion API path ---
+        tool_id_names = {"call_1": "get_weather"}
+        chat_msgs = []
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {"role": "user", "content": "What's the weather in Paris?"}
+        )
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {
+                "role": "assistant",
+                "reasoning": "I should call get_weather for Paris.",
+                "tool_calls": [
+                    {
+                        "id": "call_1",
+                        "function": {
+                            "name": "get_weather",
+                            "arguments": '{"location": "Paris"}',
+                        },
+                    }
+                ],
+            }
+        )
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {"role": "tool", "tool_call_id": "call_1", "content": "18°C, clear skies."},
+            tool_id_names=tool_id_names,
+        )
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {
+                "role": "assistant",
+                "content": "It is currently 18°C in Paris with clear skies.",
+            }
+        )
+
+        # --- Responses API path ---
+        resp_input = [
+            {
+                "type": "message",
+                "role": "user",
+                "content": "What's the weather in Paris?",
+            },
+            {
+                "type": "reasoning",
+                "content": [
+                    {
+                        "type": "reasoning_text",
+                        "text": "I should call get_weather for Paris.",
+                    }
+                ],
+            },
+            {
+                "type": "function_call",
+                "name": "get_weather",
+                "arguments": '{"location": "Paris"}',
+            },
+            {
+                "type": "function_call_output",
+                "call_id": "call_1",
+                "output": "18°C, clear skies.",
+            },
+            {
+                "type": "message",
+                "role": "assistant",
+                "content": "It is currently 18°C in Paris with clear skies.",
+            },
+        ]
+        resp_msgs = [
+            response_input_to_harmony(item, prev_responses=[prev_call])
+            for item in resp_input
+        ]
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
+
+    def test_multi_turn_two_tool_calls_with_reasoning_between(self):
+        """Validates parity for a chain of two tool calls, each with its own
+        reasoning trace. Reasoning traces in between commentary-channel tool
+        calls must survive as analysis-channel messages in both paths.
+        """
+        prev_call_1 = ResponseFunctionToolCall(
+            id="fc_1",
+            call_id="call_1",
+            name="get_weather",
+            arguments='{"location": "Paris"}',
+            type="function_call",
+        )
+        prev_call_2 = ResponseFunctionToolCall(
+            id="fc_2",
+            call_id="call_2",
+            name="get_forecast",
+            arguments='{"location": "Paris", "days": 7}',
+            type="function_call",
+        )
+
+        # --- Chat completion API path ---
+        tool_id_names = {"call_1": "get_weather", "call_2": "get_forecast"}
+        chat_msgs = []
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {"role": "user", "content": "What's the weather and forecast for Paris?"}
+        )
+        # First reasoning + tool call
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {
+                "role": "assistant",
+                "reasoning": "I need current weather first.",
+                "tool_calls": [
+                    {
+                        "id": "call_1",
+                        "function": {
+                            "name": "get_weather",
+                            "arguments": '{"location": "Paris"}',
+                        },
+                    }
+                ],
+            }
+        )
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {"role": "tool", "tool_call_id": "call_1", "content": "18°C, clear skies."},
+            tool_id_names=tool_id_names,
+        )
+        # Second reasoning + tool call
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {
+                "role": "assistant",
+                "reasoning": "Now I need the weekly forecast.",
+                "tool_calls": [
+                    {
+                        "id": "call_2",
+                        "function": {
+                            "name": "get_forecast",
+                            "arguments": '{"location": "Paris", "days": 7}',
+                        },
+                    }
+                ],
+            }
+        )
+        chat_msgs += parse_chat_input_to_harmony_message(
+            {
+                "role": "tool",
+                "tool_call_id": "call_2",
+                "content": "Mon 17°C, Tue 19°C, Wed 16°C",
+            },
+            tool_id_names=tool_id_names,
+        )
+
+        # --- Responses API path ---
+        prev_responses = [prev_call_1, prev_call_2]
+        resp_input = [
+            {
+                "type": "message",
+                "role": "user",
+                "content": "What's the weather and forecast for Paris?",
+            },
+            # First reasoning + tool call
+            {
+                "type": "reasoning",
+                "content": [
+                    {"type": "reasoning_text", "text": "I need current weather first."}
+                ],
+            },
+            {
+                "type": "function_call",
+                "name": "get_weather",
+                "arguments": '{"location": "Paris"}',
+            },
+            {
+                "type": "function_call_output",
+                "call_id": "call_1",
+                "output": "18°C, clear skies.",
+            },
+            # Second reasoning + tool call
+            {
+                "type": "reasoning",
+                "content": [
+                    {
+                        "type": "reasoning_text",
+                        "text": "Now I need the weekly forecast.",
+                    }
+                ],
+            },
+            {
+                "type": "function_call",
+                "name": "get_forecast",
+                "arguments": '{"location": "Paris", "days": 7}',
+            },
+            {
+                "type": "function_call_output",
+                "call_id": "call_2",
+                "output": "Mon 17°C, Tue 19°C, Wed 16°C",
+            },
+        ]
+        resp_msgs = [
+            response_input_to_harmony(item, prev_responses=prev_responses)
+            for item in resp_input
+        ]
+
+        assert render_for_completion([_system()] + chat_msgs) == render_for_completion(
+            [_system()] + resp_msgs
+        )
@@ -325,8 +325,12 @@ async def test_function_calling_with_streaming_expected_arguments(
    "tool_choice",
    ["auto", "required", {"type": "function", "name": "get_current_weather"}],
 )
+@pytest.mark.parametrize(
+    "enable_thinking",
+    [True, False],
+)
 async def test_function_calling_with_streaming_types(
-    client: openai.AsyncOpenAI, model_name: str, tool_choice
+    client: openai.AsyncOpenAI, model_name: str, tool_choice, enable_thinking: bool
 ):
    # this links the "done" type with the "start" type
    # so every "done" type should have a corresponding "start" type
@@ -436,6 +440,7 @@ async def test_function_calling_with_streaming_types(
        input=input_list,
        tools=tools,
        tool_choice=tool_choice,
+        extra_body={"chat_template_kwargs": {"enable_thinking": enable_thinking}},
        stream=True,
    )

@@ -0,0 +1,274 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Unit tests for response_input_to_harmony.
+
+Covers every type branch in the function and verifies that each produced
+Harmony Message has the correct role, channel, recipient, content_type,
+author name, and text content.
+"""
+
+import pytest
+from openai.types.responses import ResponseFunctionToolCall, ResponseReasoningItem
+from openai.types.responses.response_reasoning_item import (
+    Content as ReasoningTextContent,
+)
+from openai_harmony import Role
+
+from vllm.entrypoints.openai.responses.harmony import response_input_to_harmony
+
+# ---------------------------------------------------------------------------
+# Shared fixtures
+# ---------------------------------------------------------------------------
+
+_PREV_CALL = ResponseFunctionToolCall(
+    id="fc_test",
+    call_id="call_test",
+    name="get_weather",
+    arguments='{"location": "Paris"}',
+    type="function_call",
+)
+
+_REASONING_ITEM = ResponseReasoningItem(
+    id="rs_test",
+    type="reasoning",
+    content=[ReasoningTextContent(type="reasoning_text", text="Thinking hard.")],
+    summary=[],
+    status=None,
+)
+
+
+class TestResponseInputToHarmonyMessage:
+    """Unit tests for every message type handled by response_input_to_harmony."""
+
+    # -----------------------------------------------------------------------
+    # type="message" (or no type key)
+    # -----------------------------------------------------------------------
+
+    def test_user_message_string_content(self):
+        msg = response_input_to_harmony(
+            {"type": "message", "role": "user", "content": "Hello"},
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.USER
+        assert msg.content[0].text == "Hello"
+        assert msg.channel is None
+
+    def test_no_type_key_defaults_to_message_branch(self):
+        """Omitting 'type' should fall through to the message branch."""
+        msg = response_input_to_harmony(
+            {"role": "user", "content": "Hello"},
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.USER
+        assert msg.content[0].text == "Hello"
+
+    def test_system_message(self):
+        msg = response_input_to_harmony(
+            {"type": "message", "role": "system", "content": "Be helpful."},
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.SYSTEM
+        assert msg.content[0].text == "Be helpful."
+        assert msg.channel is None
+
+    def test_assistant_message_gets_final_channel(self):
+        msg = response_input_to_harmony(
+            {"type": "message", "role": "assistant", "content": "The answer is 42."},
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.ASSISTANT
+        assert msg.channel == "final"
+        assert msg.content[0].text == "The answer is 42."
+
+    def test_developer_message_gets_instructions_prefix(self):
+        msg = response_input_to_harmony(
+            {"type": "message", "role": "developer", "content": "Be concise."},
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.DEVELOPER
+        assert msg.content[0].text == "Instructions:\nBe concise."
+        assert msg.channel is None
+
+    def test_message_with_array_content(self):
+        msg = response_input_to_harmony(
+            {
+                "type": "message",
+                "role": "user",
+                "content": [
+                    {"type": "text", "text": "Part one. "},
+                    {"type": "text", "text": "Part two."},
+                ],
+            },
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.USER
+        assert len(msg.content) == 2
+        assert msg.content[0].text == "Part one. "
+        assert msg.content[1].text == "Part two."
+
+    def test_developer_message_array_content_gets_prefix_on_each_part(self):
+        msg = response_input_to_harmony(
+            {
+                "type": "message",
+                "role": "developer",
+                "content": [
+                    {"type": "text", "text": "Rule 1."},
+                    {"type": "text", "text": "Rule 2."},
+                ],
+            },
+            prev_responses=[],
+        )
+
+        assert msg.content[0].text == "Instructions:\nRule 1."
+        assert msg.content[1].text == "Instructions:\nRule 2."
+
+    # -----------------------------------------------------------------------
+    # type="reasoning"
+    # -----------------------------------------------------------------------
+
+    def test_reasoning_gets_analysis_channel(self):
+        msg = response_input_to_harmony(
+            {
+                "type": "reasoning",
+                "content": [
+                    {"type": "reasoning_text", "text": "I should call get_weather."}
+                ],
+            },
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.ASSISTANT
+        assert msg.channel == "analysis"
+        assert msg.content[0].text == "I should call get_weather."
+
+    def test_reasoning_pydantic_model_input(self):
+        """A Pydantic ResponseReasoningItem should be model_dump()'d before parsing."""
+        msg = response_input_to_harmony(_REASONING_ITEM, prev_responses=[])
+
+        assert msg.author.role == Role.ASSISTANT
+        assert msg.channel == "analysis"
+        assert msg.content[0].text == "Thinking hard."
+
+    # -----------------------------------------------------------------------
+    # type="function_call"
+    # -----------------------------------------------------------------------
+
+    def test_function_call_channel_recipient_and_content_type(self):
+        msg = response_input_to_harmony(
+            {
+                "type": "function_call",
+                "name": "get_weather",
+                "arguments": '{"location": "Paris"}',
+            },
+            prev_responses=[],
+        )
+
+        assert msg.author.role == Role.ASSISTANT
+        assert msg.channel == "commentary"
+        assert msg.recipient == "functions.get_weather"
+        assert msg.content_type == "json"
+        assert msg.content[0].text == '{"location": "Paris"}'
+
+    def test_function_call_empty_arguments(self):
+        msg = response_input_to_harmony(
+            {"type": "function_call", "name": "ping", "arguments": ""},
+            prev_responses=[],
+        )
+
+        assert msg.recipient == "functions.ping"
+        assert msg.content[0].text == ""
+
+    # -----------------------------------------------------------------------
+    # type="function_call_output"
+    # -----------------------------------------------------------------------
+
+    def test_function_call_output_channel_recipient_and_author_name(self):
+        msg = response_input_to_harmony(
+            {"type": "function_call_output", "call_id": "call_test", "output": "18°C"},
+            prev_responses=[_PREV_CALL],
+        )
+
+        assert msg.author.role == Role.TOOL
+        assert msg.author.name == "functions.get_weather"
+        assert msg.channel == "commentary"
+        assert msg.recipient == "assistant"
+        assert msg.content[0].text == "18°C"
+
+    def test_function_call_output_uses_most_recent_matching_call(self):
+        """When multiple prev_responses share a call_id, the last one wins
+        because the search is reversed."""
+        earlier = ResponseFunctionToolCall(
+            id="fc_old",
+            call_id="call_test",
+            name="old_func",
+            arguments="{}",
+            type="function_call",
+        )
+        later = ResponseFunctionToolCall(
+            id="fc_new",
+            call_id="call_test",
+            name="get_weather",
+            arguments="{}",
+            type="function_call",
+        )
+
+        msg = response_input_to_harmony(
+            {
+                "type": "function_call_output",
+                "call_id": "call_test",
+                "output": "result",
+            },
+            prev_responses=[earlier, later],
+        )
+
+        assert msg.author.name == "functions.get_weather"
+
+    def test_function_call_output_skips_non_function_call_items_in_prev_responses(
+        self,
+    ):
+        """ResponseReasoningItem entries in prev_responses should be ignored."""
+        msg = response_input_to_harmony(
+            {
+                "type": "function_call_output",
+                "call_id": "call_test",
+                "output": "18°C",
+            },
+            prev_responses=[_REASONING_ITEM, _PREV_CALL],
+        )
+
+        assert msg.author.name == "functions.get_weather"
+
+    def test_function_call_output_raises_if_no_matching_call(self):
+        with pytest.raises(ValueError, match="No call message found for"):
+            response_input_to_harmony(
+                {
+                    "type": "function_call_output",
+                    "call_id": "no_such_id",
+                    "output": "x",
+                },
+                prev_responses=[_PREV_CALL],
+            )
+
+    def test_function_call_output_raises_on_empty_prev_responses(self):
+        with pytest.raises(ValueError, match="No call message found for"):
+            response_input_to_harmony(
+                {"type": "function_call_output", "call_id": "call_test", "output": "x"},
+                prev_responses=[],
+            )
+
+    # -----------------------------------------------------------------------
+    # Error cases
+    # -----------------------------------------------------------------------
+
+    def test_unknown_type_raises_value_error(self):
+        with pytest.raises(ValueError, match="Unknown input type"):
+            response_input_to_harmony(
+                {"type": "image_url", "url": "https://example.com/img.png"},
+                prev_responses=[],
+            )
@@ -10,7 +10,7 @@ import pybase64 as base64
 import pytest
 import websockets

-from tests.entrypoints.openai.conftest import add_attention_backend
+from tests.entrypoints.speech_to_text.conftest import add_attention_backend
 from tests.utils import ROCM_ENV_OVERRIDES, ROCM_EXTRA_ARGS, RemoteOpenAIServer
 from vllm.assets.audio import AudioAsset
 from vllm.multimodal.media.audio import load_audio
@@ -163,6 +163,11 @@ async def test_multi_chunk_streaming(
                " A little piece of practical poetry. Mary had a little lamb,"
                " it sleeps with quite a flow, and everywhere that Mary went,"
                " the lamb was sure to go."
+            ) or full_text == (
+                " First words I spoke in the original phonograph."
+                " A little piece of practical poetry. Mary had a little lamb,"
+                " it squeaked with quite a flow, and everywhere that Mary went,"
+                " the lamb was sure to go."
            )


@@ -7,8 +7,8 @@ from unittest.mock import AsyncMock, Mock

 import pytest

-from vllm.entrypoints.openai.speech_to_text.protocol import TranscriptionResponse
-from vllm.entrypoints.openai.speech_to_text.speech_to_text import OpenAISpeechToText
+from vllm.entrypoints.speech_to_text.base.serving import OpenAISpeechToText
+from vllm.entrypoints.speech_to_text.transcription.protocol import TranscriptionResponse


 async def _never_finishes():
@@ -0,0 +1,61 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import pytest
+import pytest_asyncio
+
+from tests.utils import RemoteOpenAIServer
+
+
+@pytest.fixture(scope="module")
+def transcription_server_with_force_include_usage():
+    args = [
+        # use half precision for speed and memory savings in CI environment
+        "--dtype",
+        "bfloat16",
+        "--max-num-seqs",
+        "4",
+        "--enforce-eager",
+        "--enable-force-include-usage",
+        "--gpu-memory-utilization",
+        "0.2",
+    ]
+
+    with RemoteOpenAIServer("openai/whisper-large-v3-turbo", args) as remote_server:
+        yield remote_server
+
+
+@pytest_asyncio.fixture
+async def transcription_client_with_force_include_usage(
+    transcription_server_with_force_include_usage,
+):
+    async with (
+        transcription_server_with_force_include_usage.get_async_client() as async_client
+    ):
+        yield async_client
+
+
+@pytest.mark.asyncio
+async def test_transcription_with_enable_force_include_usage(
+    transcription_client_with_force_include_usage, winning_call
+):
+    res = (
+        await transcription_client_with_force_include_usage.audio.transcriptions.create(
+            model="openai/whisper-large-v3-turbo",
+            file=winning_call,
+            language="en",
+            temperature=0.0,
+            stream=True,
+            timeout=30,
+        )
+    )
+
+    async for chunk in res:
+        if not len(chunk.choices):
+            # final usage sent
+            usage = chunk.usage
+            assert isinstance(usage, dict)
+            assert usage["prompt_tokens"] > 0
+            assert usage["completion_tokens"] > 0
+            assert usage["total_tokens"] > 0
+        else:
+            assert not hasattr(chunk, "usage")
@@ -24,12 +24,14 @@ from vllm.entrypoints.openai.engine.protocol import (
    RequestResponseMetadata,
 )
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
-from vllm.entrypoints.openai.speech_to_text.protocol import TranscriptionRequest
-from vllm.entrypoints.openai.speech_to_text.serving import OpenAIServingTranscription
-from vllm.entrypoints.openai.speech_to_text.speech_to_text import (
+from vllm.entrypoints.speech_to_text.base.serving import (
    OpenAISpeechToText,
    asr_inter_chunk_separator,
 )
+from vllm.entrypoints.speech_to_text.transcription.protocol import TranscriptionRequest
+from vllm.entrypoints.speech_to_text.transcription.serving import (
+    OpenAIServingTranscription,
+)
 from vllm.model_executor.models.interfaces import SupportsTranscription
 from vllm.outputs import CompletionOutput, RequestOutput

@@ -6,7 +6,7 @@ import json

 import pytest

-from tests.entrypoints.openai.conftest import add_attention_backend
+from tests.entrypoints.speech_to_text.conftest import add_attention_backend
 from tests.utils import ROCM_ENV_OVERRIDES, ROCM_EXTRA_ARGS, RemoteOpenAIServer

 MISTRAL_FORMAT_ARGS = [
@@ -13,7 +13,7 @@ import pytest
 import pytest_asyncio
 import soundfile as sf

-from tests.entrypoints.openai.conftest import add_attention_backend
+from tests.entrypoints.speech_to_text.conftest import add_attention_backend
 from tests.utils import RemoteOpenAIServer
 from vllm.logger import init_logger
 from vllm.multimodal.media.audio import load_audio
@@ -0,0 +1,3 @@
+Qwen3-Next-80B-A3B-NVFP4-EP2.yaml
+Qwen3-Next-FP8-EP2.yaml
+Nemotron-3-Super-120B-A12B-NVFP4.yaml
@@ -3,6 +3,3 @@ Qwen2.5-VL-3B-Instruct-FP8-dynamic.yaml
 Qwen1.5-MoE-W4A16-CT.yaml
 DeepSeek-V2-Lite-Instruct-FP8.yaml
 Qwen3-30B-A3B-NVFP4.yaml
-Qwen3-Next-80B-A3B-NVFP4-EP2.yaml
-Qwen3-Next-FP8-EP2.yaml
-Nemotron-3-Super-120B-A12B-NVFP4.yaml
@@ -0,0 +1,377 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import torch
+
+from vllm.platforms import current_platform
+
+pytestmark = pytest.mark.skipif(
+    not current_platform.is_rocm(), reason="Only used by ROCm"
+)
+
+NOPE_HEAD_DIM = 448
+ROPE_HEAD_DIM = 64
+HEAD_DIM = NOPE_HEAD_DIM + ROPE_HEAD_DIM
+
+
+def _ref_global_topk_ragged(
+    topk_indices: torch.Tensor,
+    token_to_req_indices: torch.Tensor,
+    block_table: torch.Tensor,
+    block_size: int,
+    is_valid_token: torch.Tensor,
+) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
+    topk = topk_indices.reshape(topk_indices.shape[0], -1)
+    valid = (topk >= 0) & is_valid_token[:, None]
+    lens = valid.sum(dim=1, dtype=torch.int32)
+    indptr = torch.zeros(lens.shape[0] + 1, dtype=torch.int32, device=topk.device)
+    torch.cumsum(lens, dim=0, out=indptr[1:])
+
+    safe_topk = torch.clamp(topk, min=0)
+    block_indices = safe_topk // block_size
+    block_offsets = safe_topk % block_size
+    req_indices = token_to_req_indices[:, None].expand_as(topk)
+    slot_ids = block_table[req_indices, block_indices] * block_size + block_offsets
+
+    offsets = torch.arange(topk.shape[1], dtype=torch.int32, device=topk.device)
+    positions = indptr[:-1, None] + offsets[None, :]
+    return slot_ids[valid], positions[valid].to(torch.long), indptr, lens
+
+
+def _ref_sparse_prefill_ragged(
+    q: torch.Tensor,
+    kv: torch.Tensor,
+    rows: list[list[int]],
+    scale: float,
+    attn_sink: torch.Tensor | None,
+) -> torch.Tensor:
+    q_f32 = q.float()
+    kv_f32 = kv.float()
+    out = torch.empty_like(q_f32)
+
+    for query_idx in range(q.shape[0]):
+        row_indices = rows[query_idx]
+        for head_idx in range(q.shape[1]):
+            if row_indices:
+                selected_kv = kv_f32[row_indices]
+                scores = torch.mv(selected_kv, q_f32[query_idx, head_idx]) * scale
+                if attn_sink is not None:
+                    scores_with_sink = torch.cat(
+                        [scores, attn_sink[head_idx].float().reshape(1)]
+                    )
+                    probs = torch.softmax(scores_with_sink, dim=0)[:-1]
+                else:
+                    probs = torch.softmax(scores, dim=0)
+                out[query_idx, head_idx] = torch.sum(
+                    probs[:, None] * selected_kv, dim=0
+                )
+            else:
+                out[query_idx, head_idx] = 0
+    return out.to(torch.bfloat16)
+
+
+def _pack_fp8_ds_mla_cache(kv: torch.Tensor, block_size: int) -> torch.Tensor:
+    assert kv.shape[-1] == HEAD_DIM
+    num_tokens = kv.shape[0]
+    num_blocks = (num_tokens + block_size - 1) // block_size
+    cache = torch.zeros(
+        (num_blocks, block_size, 584),
+        dtype=torch.uint8,
+        device=kv.device,
+    )
+    cache_flat = cache.view(torch.uint8).flatten()
+    kv_nope_fp8 = (
+        kv[:, :NOPE_HEAD_DIM].to(current_platform.fp8_dtype()).view(torch.uint8)
+    )
+    kv_rope_u8 = kv[:, NOPE_HEAD_DIM:].contiguous().view(torch.uint8)
+
+    for slot in range(num_tokens):
+        block_idx = slot // block_size
+        pos = slot % block_size
+        block_base = block_idx * cache.stride(0)
+        token_base = block_base + pos * 576
+        scale_base = block_base + block_size * 576 + pos * 8
+        cache_flat[token_base : token_base + NOPE_HEAD_DIM].copy_(kv_nope_fp8[slot])
+        cache_flat[
+            token_base + NOPE_HEAD_DIM : token_base + NOPE_HEAD_DIM + ROPE_HEAD_DIM * 2
+        ].copy_(kv_rope_u8[slot])
+        cache_flat[scale_base : scale_base + 7].fill_(127)
+    return cache
+
+
+def _read_fp8_ds_mla_cache(
+    cache: torch.Tensor, slot: int, block_size: int
+) -> torch.Tensor:
+    cache_flat = cache.view(torch.uint8).flatten()
+    block_idx = slot // block_size
+    pos = slot % block_size
+    block_base = block_idx * cache.stride(0)
+    token_base = block_base + pos * 576
+
+    nope_u8 = cache_flat[token_base : token_base + NOPE_HEAD_DIM]
+    nope = nope_u8.view(current_platform.fp8_dtype()).to(torch.float32)
+    rope_u8 = cache_flat[
+        token_base + NOPE_HEAD_DIM : token_base + NOPE_HEAD_DIM + ROPE_HEAD_DIM * 2
+    ]
+    rope = rope_u8.view(torch.bfloat16).to(torch.float32)
+    return torch.cat([nope, rope])
+
+
+def _ref_sparse_decode_ragged(
+    q: torch.Tensor,
+    main_cache: torch.Tensor,
+    main_rows: list[list[int]],
+    scale: float,
+    attn_sink: torch.Tensor | None,
+    block_size: int,
+    extra_cache: torch.Tensor | None = None,
+    extra_rows: list[list[int]] | None = None,
+) -> torch.Tensor:
+    q_f32 = q.float()
+    out = torch.empty_like(q_f32)
+
+    for query_idx in range(q.shape[0]):
+        row_kv = [
+            _read_fp8_ds_mla_cache(main_cache, int(slot), block_size)
+            for slot in main_rows[query_idx]
+        ]
+        if extra_cache is not None and extra_rows is not None:
+            row_kv.extend(
+                _read_fp8_ds_mla_cache(extra_cache, int(slot), block_size)
+                for slot in extra_rows[query_idx]
+            )
+
+        kv = torch.stack(row_kv).to(q.device)
+        for head_idx in range(q.shape[1]):
+            scores = torch.mv(kv, q_f32[query_idx, head_idx]) * scale
+            if attn_sink is not None:
+                scores_with_sink = torch.cat(
+                    [scores, attn_sink[head_idx].float().reshape(1)]
+                )
+                probs = torch.softmax(scores_with_sink, dim=0)[:-1]
+            else:
+                probs = torch.softmax(scores, dim=0)
+            out[query_idx, head_idx] = torch.sum(probs[:, None] * kv, dim=0)
+    return out.to(torch.bfloat16)
+
+
+def _ref_combine_topk_swa_ragged(
+    device: torch.device,
+) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    expected_ragged = torch.tensor(
+        [
+            100,
+            101,
+            7,
+            8,
+            9,
+            110,
+            111,
+            8,
+            9,
+            10,
+            120,
+            121,
+            122,
+            9,
+            10,
+            11,
+            150,
+            27,
+            28,
+            29,
+            160,
+            161,
+            28,
+            29,
+            30,
+        ],
+        dtype=torch.int32,
+        device=device,
+    )
+    expected_lens = torch.tensor([5, 5, 6, 4, 5], dtype=torch.int32, device=device)
+    expected_indptr = torch.zeros(6, dtype=torch.int32, device=device)
+    torch.cumsum(expected_lens, dim=0, out=expected_indptr[1:])
+    return expected_ragged, expected_indptr, expected_lens
+
+
+@torch.inference_mode()
+def test_compute_global_topk_ragged_indices_and_indptr() -> None:
+    from vllm.v1.attention.backends.mla.rocm_aiter_mla_sparse_dsv4 import (
+        compute_global_topk_ragged_indices_and_indptr,
+    )
+
+    device = torch.device("cuda")
+    block_size = 4
+    topk_indices = torch.tensor(
+        [
+            [0, 3, 4, -1],
+            [5, 8, -1, -1],
+            [2, 7, 9, -1],
+        ],
+        dtype=torch.int32,
+        device=device,
+    )
+    token_to_req_indices = torch.tensor([0, 1, 1], dtype=torch.int32, device=device)
+    block_table = torch.tensor(
+        [
+            [10, 11, 12],
+            [20, 21, 22],
+        ],
+        dtype=torch.int32,
+        device=device,
+    )
+    is_valid_token = torch.tensor([True, False, True], dtype=torch.bool, device=device)
+
+    actual_ragged, actual_indptr, actual_lens = (
+        compute_global_topk_ragged_indices_and_indptr(
+            topk_indices,
+            token_to_req_indices,
+            block_table,
+            block_size,
+            is_valid_token,
+        )
+    )
+    expected_values, expected_positions, expected_indptr, expected_lens = (
+        _ref_global_topk_ragged(
+            topk_indices,
+            token_to_req_indices,
+            block_table,
+            block_size,
+            is_valid_token,
+        )
+    )
+
+    torch.testing.assert_close(actual_ragged[expected_positions], expected_values)
+    torch.testing.assert_close(actual_indptr, expected_indptr)
+    torch.testing.assert_close(actual_lens, expected_lens)
+
+
+@torch.inference_mode()
+def test_sparse_attn_prefill_ragged_kernel() -> None:
+    from vllm.v1.attention.ops.rocm_aiter_mla_sparse import (
+        _rocm_sparse_attn_prefill_ragged_triton,
+    )
+
+    device = torch.device("cuda")
+    torch.manual_seed(0)
+    q = torch.randn(3, 3, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
+    kv = torch.randn(5, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
+    indices = torch.tensor([0, 2, 1, 3, 4], dtype=torch.int32, device=device)
+    indptr = torch.tensor([0, 2, 5, 5], dtype=torch.int32, device=device)
+    attn_sink = torch.tensor([-0.25, 0.0, 0.25], dtype=torch.float32, device=device)
+    scale = HEAD_DIM**-0.5
+
+    actual = _rocm_sparse_attn_prefill_ragged_triton(
+        q=q,
+        kv=kv,
+        indices=indices,
+        indptr=indptr,
+        scale=scale,
+        attn_sink=attn_sink,
+        nope_head_dim=NOPE_HEAD_DIM,
+        rope_head_dim=ROPE_HEAD_DIM,
+    )
+    expected = _ref_sparse_prefill_ragged(
+        q, kv, [[0, 2], [1, 3, 4], []], scale, attn_sink
+    )
+
+    torch.testing.assert_close(actual, expected, atol=2e-2, rtol=2e-2)
+
+
+@torch.inference_mode()
+def test_sparse_attn_decode_ragged_kernel() -> None:
+    from vllm.v1.attention.ops.rocm_aiter_mla_sparse import (
+        _rocm_sparse_attn_decode_ragged_triton,
+    )
+
+    device = torch.device("cuda")
+    torch.manual_seed(1)
+    block_size = 4
+    q = torch.randn(2, 3, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
+    main_kv = torch.randn(6, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
+    extra_kv = torch.randn(5, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
+    main_cache = _pack_fp8_ds_mla_cache(main_kv, block_size)
+    extra_cache = _pack_fp8_ds_mla_cache(extra_kv, block_size)
+    main_indices = torch.tensor([0, 2, 4, 1], dtype=torch.int32, device=device)
+    main_indptr = torch.tensor([0, 2, 4], dtype=torch.int32, device=device)
+    extra_indices = torch.tensor([1, 3, 0], dtype=torch.int32, device=device)
+    extra_indptr = torch.tensor([0, 1, 3], dtype=torch.int32, device=device)
+    attn_sink = torch.tensor([-0.1, 0.0, 0.1], dtype=torch.float32, device=device)
+    scale = HEAD_DIM**-0.5
+
+    actual = _rocm_sparse_attn_decode_ragged_triton(
+        q=q,
+        main_cache=main_cache,
+        main_indices=main_indices,
+        main_indptr=main_indptr,
+        scale=scale,
+        attn_sink=attn_sink,
+        nope_head_dim=NOPE_HEAD_DIM,
+        rope_head_dim=ROPE_HEAD_DIM,
+        extra_cache=extra_cache,
+        extra_indices=extra_indices,
+        extra_indptr=extra_indptr,
+    )
+    expected = _ref_sparse_decode_ragged(
+        q=q,
+        main_cache=main_cache,
+        main_rows=[[0, 2], [4, 1]],
+        scale=scale,
+        attn_sink=attn_sink,
+        block_size=block_size,
+        extra_cache=extra_cache,
+        extra_rows=[[1], [3, 0]],
+    )
+
+    torch.testing.assert_close(actual, expected, atol=2e-2, rtol=2e-2)
+
+
+@torch.inference_mode()
+def test_combine_topk_swa_indices_ragged() -> None:
+    from vllm.v1.attention.backends.mla.rocm_aiter_mla_sparse_dsv4 import (
+        combine_topk_swa_indices_ragged,
+    )
+
+    device = torch.device("cuda")
+    topk_indices = torch.tensor(
+        [
+            [100, 101, 102, 103],
+            [110, 111, 112, 113],
+            [120, 121, 122, 123],
+            [130, 131, 132, 133],
+            [140, 141, 142, 143],
+        ],
+        dtype=torch.int32,
+        device=device,
+    )
+    query_start_loc = torch.tensor([0, 3, 5], dtype=torch.int32, device=device)
+    seq_lens = torch.tensor([6, 4], dtype=torch.int32, device=device)
+    gather_lens = torch.tensor([4, 3], dtype=torch.int32, device=device)
+    window_size = 3
+    compress_ratio = 2
+    topk = 4
+    M = 20
+    N = 8
+
+    actual_ragged, actual_indptr, actual_lens = combine_topk_swa_indices_ragged(
+        topk_indices,
+        query_start_loc,
+        seq_lens,
+        gather_lens,
+        window_size,
+        compress_ratio,
+        topk,
+        M,
+        N,
+    )
+    expected_ragged, expected_indptr, expected_lens = _ref_combine_topk_swa_ragged(
+        device
+    )
+
+    torch.testing.assert_close(
+        actual_ragged[: expected_ragged.numel()], expected_ragged
+    )
+    torch.testing.assert_close(actual_indptr, expected_indptr)
+    torch.testing.assert_close(actual_lens, expected_lens)
@@ -13,6 +13,7 @@ from vllm.model_executor.layers.mamba.ops.ssu_dispatch import (
    selective_state_update,
 )
 from vllm.utils.torch_utils import set_random_seed
+from vllm.v1.attention.backends.registry import MambaAttentionBackendEnum
 from vllm.v1.kv_cache_interface import (
    KVCacheConfig,
    KVCacheGroupSpec,
@@ -27,7 +28,9 @@ except ImportError:
    HAS_FLASHINFER = False


-def _kv_cache_config_with_ssu(mamba_type: str = "mamba2") -> KVCacheConfig:
+def _kv_cache_config_with_ssu(
+    mamba_type: MambaAttentionBackendEnum = MambaAttentionBackendEnum.MAMBA2,
+) -> KVCacheConfig:
    spec = MambaSpec(
        block_size=16,
        shapes=((16, 64),),
@@ -77,7 +80,12 @@ def test_uninitialized_backend_raises():


@pytest.mark.parametrize(
-    "mamba_type", ["linear_attention", "gdn_attention", "short_conv"]
+    "mamba_type",
+    [
+        MambaAttentionBackendEnum.LINEAR,
+        MambaAttentionBackendEnum.GDN_ATTN,
+        MambaAttentionBackendEnum.SHORT_CONV,
+    ],
 )
 def test_init_is_noop_for_non_ssu_mamba_type(mamba_type):
    import vllm.model_executor.layers.mamba.ops.ssu_dispatch as mod
@@ -237,7 +237,7 @@ if has_mori():
    )

 if has_flashinfer_cutlass_fused_moe() and current_platform.has_device_capability(100):
-    from vllm.model_executor.layers.fused_moe.flashinfer_cutlass_moe import (
+    from vllm.model_executor.layers.fused_moe.experts.flashinfer_cutlass_moe import (
        FlashInferExperts,
    )
    from vllm.model_executor.layers.fused_moe.prepare_finalize.flashinfer_nvlink_two_sided import (  # noqa: E501
@@ -298,7 +298,7 @@ if has_flashinfer_cutlass_fused_moe() and current_platform.has_device_capability
    )

 if has_aiter():
-    from vllm.model_executor.layers.fused_moe.rocm_aiter_fused_moe import (
+    from vllm.model_executor.layers.fused_moe.experts.rocm_aiter_moe import (
        AiterExperts,
    )

@@ -18,12 +18,12 @@ from vllm.model_executor.layers.fused_moe.config import (
    RoutingMethodType,
    fp8_w8a8_moe_quant_config,
 )
+from vllm.model_executor.layers.fused_moe.experts.flashinfer_cutlass_moe import (
+    FlashInferExperts,
+)
 from vllm.model_executor.layers.fused_moe.experts.trtllm_fp8_moe import (
    TrtLlmFp8ExpertsMonolithic,
 )
-from vllm.model_executor.layers.fused_moe.flashinfer_cutlass_moe import (
-    FlashInferExperts,
-)
 from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts
 from vllm.model_executor.layers.quantization.utils.flashinfer_utils import (
    rotate_weights_for_fi_trtllm_fp8_per_tensor_moe,
@@ -22,7 +22,7 @@ from vllm.model_executor.layers.fused_moe.config import (
    FusedMoEParallelConfig,
    RoutingMethodType,
 )
-from vllm.model_executor.layers.fused_moe.flashinfer_cutlass_moe import (
+from vllm.model_executor.layers.fused_moe.experts.flashinfer_cutlass_moe import (
    FlashInferExperts,
    is_valid_flashinfer_cutlass_fused_moe,
 )
@@ -5,7 +5,7 @@
 import pytest
 import torch

-from vllm.model_executor.layers.fused_moe.fused_marlin_moe import (
+from vllm.model_executor.layers.fused_moe.experts.marlin_moe import (
    fused_marlin_moe,
 )
 from vllm.model_executor.layers.fused_moe.router.grouped_topk_router import (
@@ -32,7 +32,7 @@ from vllm.model_executor.layers.fused_moe.config import (
    int4_w4a16_moe_quant_config,
    int8_w8a16_moe_quant_config,
 )
-from vllm.model_executor.layers.fused_moe.fused_marlin_moe import (
+from vllm.model_executor.layers.fused_moe.experts.marlin_moe import (
    batched_fused_marlin_moe,
    fused_marlin_moe,
 )
@@ -20,7 +20,7 @@ if not current_platform.is_rocm():
    pytest.skip("This test can only run on ROCm.", allow_module_level=True)

 # this import statement is needed to ensure the ops are registered
-import vllm.model_executor.layers.fused_moe.rocm_aiter_fused_moe  # noqa: F401
+import vllm.model_executor.layers.fused_moe.experts.rocm_aiter_moe  # noqa: F401

 # need to import once to ensure the ops are registered
 # Check if aiter package is installed
@@ -15,7 +15,7 @@ from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
    FUSED_MOE_UNQUANTIZED_CONFIG,
 )
-from vllm.model_executor.layers.fused_moe.fused_moe import TritonExperts
+from vllm.model_executor.layers.fused_moe.experts.triton_moe import TritonExperts
 from vllm.platforms import current_platform

 # Test parameters
@@ -151,7 +151,7 @@ def test_triton_experts_no_mul_activation(
@torch.inference_mode()
 def test_workspace_shapes_no_mul_vs_gated():
    """Test that workspace shapes differ correctly between gated and non-gated."""
-    from vllm.model_executor.layers.fused_moe.fused_moe import TritonExperts
+    from vllm.model_executor.layers.fused_moe.experts.triton_moe import TritonExperts

    M, N, K, topk = 64, 256, 128, 2

@@ -192,7 +192,7 @@ def test_workspace_shapes_no_mul_vs_gated():
@torch.inference_mode()
 def test_adjust_n_for_activation():
    """Test the adjust_N_for_activation method."""
-    from vllm.model_executor.layers.fused_moe.fused_moe import TritonExperts
+    from vllm.model_executor.layers.fused_moe.experts.triton_moe import TritonExperts

    experts = TritonExperts(
        moe_config=make_dummy_moe_config(),
@@ -158,7 +158,7 @@ def test_select_cuda_flashinfer_trtllm_backend(mock_is_supported_trtllm, monkeyp
    return_value=(False, None),
 )
@patch(
-    "vllm.model_executor.layers.fused_moe.flashinfer_cutlass_moe.FlashInferExperts.is_supported_config",
+    "vllm.model_executor.layers.fused_moe.experts.flashinfer_cutlass_moe.FlashInferExperts.is_supported_config",
    return_value=(True, None),
 )
@pytest.mark.skipif(
@@ -17,12 +17,14 @@ from vllm.model_executor.layers.fused_moe.config import (
    FusedMoEQuantConfig,
    RoutingMethodType,
 )
+from vllm.model_executor.layers.fused_moe.experts.triton_moe import (
+    TritonExperts,
+)
 from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
    BatchedTritonExperts,
    NaiveBatchedExperts,
 )
 from vllm.model_executor.layers.fused_moe.fused_moe import (
-    TritonExperts,
    fused_experts,
 )
 from vllm.model_executor.layers.fused_moe.modular_kernel import FusedMoEKernel
@@ -3,11 +3,12 @@
 """
 Round-trip tests for compressor → FP8 quant + KV cache insert → gather + dequant.

-Four test functions cover five paths:
+These tests cover:
  A) DeepseekV4 Attention: head_dim=512 (448 FP8 nope + 64 bf16 rope), quant_block=64
-  B) Indexer:       head_dim=128 (all FP8), quant_block=128
-  C) DeepseekV4 Attention magnitude range: correctness across small/large values
-  D) Indexer fused Triton kernel: compress+norm+rope+quant+insert
+  B) Fused dequant+gather K cache
+  C) Indexer:       head_dim=128 (all FP8), quant_block=128
+  D) DeepseekV4 Attention magnitude range: correctness across small/large values
+  E) Indexer fused Triton kernel: compress+norm+rope+quant+insert
 """

 import math
@@ -134,7 +135,140 @@ def test_deepseek_v4_attention_quant_cache_roundtrip(num_tokens: int, block_size
    )


-# ── Test B: Indexer path ────────────────────────────────────────────────────
+# ── Test B: Fused dequant+gather K cache ────────────────────────────────────
+
+
+def _dequantize_and_gather_k_cache_reference(
+    out: torch.Tensor,
+    k_cache: torch.Tensor,
+    seq_lens: torch.Tensor,
+    gather_lens: torch.Tensor | None,
+    block_table: torch.Tensor,
+    block_size: int,
+    offset: int,
+) -> None:
+    fp8_dim = 448
+    bf16_dim = 64
+    scale_dim = 8
+    quant_block = 64
+    token_data_size = fp8_dim + bf16_dim * 2
+
+    for req_id in range(seq_lens.shape[0]):
+        seq_len = seq_lens[req_id].item()
+        gather_len = gather_lens[req_id].item() if gather_lens is not None else seq_len
+        start_pos = seq_len - gather_len
+
+        for i in range(gather_len):
+            pos = start_pos + i
+            pos_in_block = pos % block_size
+            block_idx = block_table[req_id, pos // block_size].item()
+            cache_block = k_cache[block_idx].view(-1)
+
+            token_data_start = pos_in_block * token_data_size
+            fp8_bytes = cache_block[token_data_start : token_data_start + fp8_dim]
+            fp8_vals = fp8_bytes.view(torch.float8_e4m3fn).float()
+
+            scale_start = block_size * token_data_size + pos_in_block * scale_dim
+            encoded_scales = cache_block[scale_start : scale_start + scale_dim]
+            scales = torch.exp2(encoded_scales[:7].float() - 127.0)
+            dequant = fp8_vals * scales.repeat_interleave(quant_block)
+
+            bf16_start = token_data_start + fp8_dim
+            bf16_bytes = cache_block[bf16_start : bf16_start + bf16_dim * 2]
+            bf16_tail = bf16_bytes.view(torch.bfloat16)
+
+            out[req_id, offset + i, :fp8_dim] = dequant
+            out[req_id, offset + i, fp8_dim:] = bf16_tail
+
+
+@pytest.mark.parametrize(
+    ("seq_lens_host", "gather_lens_host", "offset"),
+    [
+        ([9, 23, 7], None, 0),
+        ([19, 8, 257], [6, 8, 129], 5),
+    ],
+)
+def test_dequantize_and_gather_k_cache(
+    seq_lens_host: list[int],
+    gather_lens_host: list[int] | None,
+    offset: int,
+):
+    block_size = 64
+    head_dim = 512
+    nope_dim = 448
+    scale_dim = 8
+    head_bytes = nope_dim + (head_dim - nope_dim) * 2 + scale_dim
+    device = "cuda"
+    num_reqs = len(seq_lens_host)
+    num_tokens = sum(seq_lens_host)
+    max_gather_len = max(gather_lens_host or seq_lens_host)
+    max_blocks_per_seq = math.ceil(max(seq_lens_host) / block_size)
+    num_blocks = sum(math.ceil(seq_len / block_size) for seq_len in seq_lens_host)
+
+    compressed_kv = torch.randn(
+        num_tokens, head_dim, dtype=torch.bfloat16, device=device
+    )
+
+    # Randomize physical pages so the test covers block-table translation.
+    # Keep padded block-table entries invalid to catch accidental reads.
+    physical_blocks = torch.randperm(num_blocks, device=device)
+    block_table = torch.full(
+        (num_reqs, max_blocks_per_seq), int(-1e6), dtype=torch.int32, device=device
+    )
+    start = 0
+    for req_id, seq_len in enumerate(seq_lens_host):
+        num_req_blocks = math.ceil(seq_len / block_size)
+        req_blocks = physical_blocks[start : start + num_req_blocks]
+        block_table[req_id, :num_req_blocks] = req_blocks
+        start += num_req_blocks
+
+    # Build slot_mapping for quantize_and_insert_k_cache.
+    slot_mapping = torch.empty(num_tokens, dtype=torch.int64, device=device)
+    start = 0
+    for req_id, seq_len in enumerate(seq_lens_host):
+        logical_pos = torch.arange(seq_len, dtype=torch.int64, device=device)
+        block_idx = block_table[req_id, logical_pos // block_size].to(torch.int64)
+        token_slots = block_idx * block_size + logical_pos % block_size
+        slot_mapping[start : start + seq_len] = token_slots
+        start += seq_len
+
+    # Insert compressed K into the paged cache layout used by the gather op.
+    k_cache = torch.empty(
+        num_blocks, block_size, head_bytes, dtype=torch.uint8, device=device
+    )
+    k_cache_2d = k_cache.view(num_blocks, -1)
+    quantize_and_insert_k_cache(compressed_kv, k_cache_2d, slot_mapping, block_size)
+
+    out_shape = (num_reqs, offset + max_gather_len + 3, head_dim)
+    ref_out = torch.empty(out_shape, dtype=torch.bfloat16, device=device)
+    actual_out = torch.empty_like(ref_out)
+    seq_lens = torch.tensor(seq_lens_host, dtype=torch.int32, device=device)
+    gather_lens = (
+        torch.tensor(gather_lens_host, dtype=torch.int32, device=device)
+        if gather_lens_host is not None
+        else None
+    )
+
+    # Compare production gather against a PyTorch reference for valid output rows.
+    _dequantize_and_gather_k_cache_reference(
+        ref_out, k_cache, seq_lens, gather_lens, block_table, block_size, offset
+    )
+    dequantize_and_gather_k_cache(
+        actual_out, k_cache, seq_lens, gather_lens, block_table, block_size, offset
+    )
+    torch.accelerator.synchronize()
+
+    # only check non-padded content
+    for req_id, seq_len in enumerate(seq_lens_host):
+        gather_len = (
+            gather_lens_host[req_id] if gather_lens_host is not None else seq_len
+        )
+        actual = actual_out[req_id, offset : offset + gather_len]
+        expected = ref_out[req_id, offset : offset + gather_len]
+        torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+
+
+# ── Test C: Indexer path ────────────────────────────────────────────────────


@pytest.mark.parametrize("num_tokens", [1, 4, 8, 17])
@@ -254,7 +388,7 @@ def test_indexer_gather_accepts_upper_bound_output():
    assert torch.all(dst_scale[valid_tokens:] == sentinel)


-# ── Test C: DeepseekV4 attention with values at different magnitudes ───────────
+# ── Test D: DeepseekV4 attention with values at different magnitudes ───────────


 def test_deepseek_v4_quant_magnitude_range():
@@ -316,7 +450,7 @@ def test_deepseek_v4_quant_magnitude_range():
            )


-# ── Test D: Indexer fused K-cache insert (Triton kernels) ────────────────────
+# ── Test E: Indexer fused K-cache insert (Triton kernels) ────────────────────
 #
 # Both kernels share the same Triton signature; use_fp4 selects between them.
 # Full pipeline: state-cache gather → softmax-weighted compress → RMSNorm →
@@ -0,0 +1,142 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import pytest
+import torch
+
+import vllm.model_executor.layers.mhc as mhc_ops  # noqa: F401
+from vllm.platforms import current_platform
+from vllm.utils.torch_utils import set_random_seed
+
+DEVICE = current_platform.device_type
+
+
+def sinkhorn_normalize_ref(x: torch.Tensor, repeat: int, eps: float) -> torch.Tensor:
+    x = x.softmax(-1) + eps
+    x = x / (x.sum(-2, keepdim=True) + eps)
+    for _ in range(repeat - 1):
+        x = x / (x.sum(-1, keepdim=True) + eps)
+        x = x / (x.sum(-2, keepdim=True) + eps)
+    return x
+
+
+def mhc_pre_ref(
+    residual: torch.Tensor,
+    fn: torch.Tensor,
+    hc_scale: torch.Tensor,
+    hc_base: torch.Tensor,
+    rms_eps: float,
+    hc_pre_eps: float,
+    hc_sinkhorn_eps: float,
+    hc_post_mult_value: float,
+    sinkhorn_repeat: int,
+) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    """mHC pre reference kernel from tilelang repo: https://github.com/tile-ai/tilelang/blob/d135bd1cd2d2eee74fbb41dd0a0831a427194c86/examples/deepseek_mhc/example_mhc_pre.py#L303"""
+    hc_mult = residual.shape[-2]
+
+    residual_flat = residual.flatten(-2, -1).float()
+    sqrsum = residual_flat.square().sum(-1)
+    mixes = (
+        residual_flat @ fn.T * (sqrsum.unsqueeze(-1) / fn.shape[-1] + rms_eps).rsqrt()
+    )
+
+    hc_scale = torch.cat(
+        [
+            hc_scale[0].expand(hc_mult),
+            hc_scale[1].expand(hc_mult),
+            hc_scale[2].expand(hc_mult * hc_mult),
+        ],
+    )
+    mixes = mixes * hc_scale + hc_base
+
+    pre_mix = mixes[:, :hc_mult].sigmoid().unsqueeze(-1) + hc_pre_eps
+    post_mix = (
+        mixes[:, hc_mult : 2 * hc_mult].sigmoid() * hc_post_mult_value
+    ).unsqueeze(-1)
+    res_mix = mixes[:, 2 * hc_mult :].view(-1, hc_mult, hc_mult)
+
+    res_mix = sinkhorn_normalize_ref(
+        res_mix, repeat=sinkhorn_repeat, eps=hc_sinkhorn_eps
+    )
+
+    layer_input = (residual * pre_mix).sum(-2).bfloat16()
+
+    return post_mix, res_mix, layer_input
+
+
+def mhc_post_ref(
+    x: torch.Tensor,
+    residual: torch.Tensor,
+    post_layer_mix: torch.Tensor,
+    comb_res_mix: torch.Tensor,
+) -> torch.Tensor:
+    """mHC post reference kernel from tilelang repo: https://github.com/tile-ai/tilelang/blob/d135bd1cd2d2eee74fbb41dd0a0831a427194c86/examples/deepseek_mhc/example_mhc_post.py#L68"""
+    term2 = torch.bmm(comb_res_mix.mT, residual.float())
+    return (x.float().unsqueeze(-2) * post_layer_mix + term2).bfloat16()
+
+
+@pytest.mark.skipif(
+    not current_platform.is_cuda(),
+    reason="CUDA 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_mhc_fused_post_pre(num_tokens, hidden_size, hc_mult):
+    torch.set_default_device(DEVICE)
+    set_random_seed(0)
+
+    x = torch.randn((num_tokens, hidden_size), dtype=torch.bfloat16)
+    residual = torch.randn((num_tokens, hc_mult, hidden_size), dtype=torch.bfloat16)
+    post_layer_mix = torch.randn((num_tokens, hc_mult, 1), dtype=torch.float32)
+    comb_res_mix = torch.randn((num_tokens, hc_mult, hc_mult), dtype=torch.float32)
+
+    hc_mult2 = hc_mult * hc_mult
+    hc_mult3 = hc_mult * 2 + hc_mult2
+    fn = (
+        torch.randn((hc_mult3, hc_mult, hidden_size), dtype=torch.float)
+        * 1e-4
+        * (1 + torch.arange(hc_mult).mul(0.01).view(1, -1, 1))
+    ).flatten(1, 2)
+    hc_scale = torch.randn((3,), dtype=torch.float) * 0.1
+    hc_base = torch.randn((hc_mult3,), dtype=torch.float) * 0.1
+
+    hc_sinkhorn_eps = hc_pre_eps = rms_eps = 1e-6
+    sinkhorn_repeat = 20
+    hc_post_alpha = 1.0
+
+    def run_ref():
+        residual_ref = mhc_post_ref(x, residual, post_layer_mix, comb_res_mix)
+        post_mix_ref, res_mix_ref, layer_input_ref = mhc_pre_ref(
+            residual_ref,
+            fn,
+            hc_scale,
+            hc_base,
+            rms_eps,
+            hc_pre_eps,
+            hc_sinkhorn_eps,
+            hc_post_alpha,
+            sinkhorn_repeat,
+        )
+        return residual_ref, post_mix_ref, res_mix_ref, layer_input_ref
+
+    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(
+        x,
+        residual,
+        post_layer_mix,
+        comb_res_mix,
+        fn,
+        hc_scale,
+        hc_base,
+        rms_eps,
+        hc_pre_eps,
+        hc_sinkhorn_eps,
+        hc_post_alpha,
+        sinkhorn_repeat,
+    )
+
+    torch.testing.assert_close(residual, residual_ref, atol=1e-2, rtol=1e-2)
+    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)
@@ -1,9 +1,12 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project

+import pytest
+
 import vllm
 import vllm.config
 from vllm.lora.request import LoRARequest
+from vllm.platforms import current_platform

 from ..utils import create_new_process_for_each_test, multi_gpu_test

@@ -50,6 +53,9 @@ def do_sample(llm: vllm.LLM, lora_path: str, lora_id: int) -> list[str]:
    return generated_texts


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@create_new_process_for_each_test()
 def test_chatglm3_lora(chatglm3_lora_files):
    llm = vllm.LLM(
@@ -70,6 +76,9 @@ def test_chatglm3_lora(chatglm3_lora_files):
        assert output2[i] == EXPECTED_LORA_OUTPUT[i]


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@multi_gpu_test(num_gpus=4)
 def test_chatglm3_lora_tp4(chatglm3_lora_files):
    llm = vllm.LLM(
@@ -11,6 +11,7 @@ import pytest
 from huggingface_hub import snapshot_download

 from vllm.lora.request import LoRARequest
+from vllm.platforms import current_platform

 from ..conftest import AudioTestAssets, VllmRunner
 from ..utils import create_new_process_for_each_test
@@ -76,6 +77,9 @@ def test_active_default_mm_lora(
    )


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@create_new_process_for_each_test()
 def test_inactive_default_mm_lora(
    vllm_runner: type[VllmRunner],
@@ -92,6 +96,9 @@ def test_inactive_default_mm_lora(
    )


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@create_new_process_for_each_test()
 def test_default_mm_lora_succeeds_with_redundant_lora_request(
    vllm_runner: type[VllmRunner],
@@ -107,6 +114,9 @@ def test_default_mm_lora_succeeds_with_redundant_lora_request(
    )


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@create_new_process_for_each_test()
 def test_default_mm_lora_fails_with_overridden_lora_request(
    vllm_runner: type[VllmRunner],
@@ -10,6 +10,7 @@ import vllm.config
 from vllm import LLM
 from vllm.lora.request import LoRARequest
 from vllm.model_executor.model_loader.tensorizer import TensorizerConfig
+from vllm.platforms import current_platform

 from ..utils import VLLM_PATH, create_new_process_for_each_test, multi_gpu_test

@@ -139,6 +140,9 @@ def test_llama_lora(llama32_lora_files, cudagraph_specialize_lora: bool):
    generate_and_test(llm, llama32_lora_files)


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@multi_gpu_test(num_gpus=4)
 def test_llama_lora_tp4(llama32_lora_files):
    llm = vllm.LLM(
@@ -68,6 +68,9 @@ def do_sample(llm: vllm.LLM, lora_path: str, lora_id: int) -> list[str]:
    return generated_texts


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
 def test_minicpmv_lora(minicpmv_lora_files):
    llm = vllm.LLM(
        MODEL_PATH,
@@ -11,6 +11,7 @@ from safetensors.torch import load_file, save_file

 import vllm
 from vllm.lora.request import LoRARequest
+from vllm.platforms import current_platform

 from ..utils import multi_gpu_test

@@ -110,6 +111,9 @@ def generate_and_test(
        )


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
 def test_olmoe_lora(olmoe_lora_files, maybe_enable_lora_dual_stream):
    # We enable enforce_eager=True here to reduce VRAM usage for lora-test CI,
    # Otherwise, the lora-test will fail due to CUDA OOM.
@@ -178,6 +182,9 @@ def test_olmoe_lora_mixed_random(
    assert outputs[0].outputs[0].text.strip().startswith(EXPECTED_LORA_OUTPUT[0])


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@pytest.mark.parametrize("fully_sharded_loras", [False, True])
@multi_gpu_test(num_gpus=2)
 def test_olmoe_lora_tp2(olmoe_lora_files, fully_sharded_loras):
@@ -1,12 +1,14 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project

+import pytest
 from transformers import AutoTokenizer

 import vllm
 import vllm.config
 from vllm.assets.image import ImageAsset
 from vllm.lora.request import LoRARequest
+from vllm.platforms import current_platform

 from ..utils import create_new_process_for_each_test, multi_gpu_test

@@ -311,6 +313,9 @@ def _assert_qwen35_text_vl_and_mixed_lora(
    )


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
@create_new_process_for_each_test()
 def test_qwen35_text_lora(
    qwen35_text_lora_files, qwen35_vl_lora_files, maybe_enable_lora_dual_stream
@@ -2,12 +2,14 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 from dataclasses import dataclass

+import pytest
 from packaging.version import Version
 from transformers import __version__ as TRANSFORMERS_VERSION

 import vllm
 from vllm.assets.image import ImageAsset
 from vllm.lora.request import LoRARequest
+from vllm.platforms import current_platform
 from vllm.sampling_params import BeamSearchParams


@@ -206,6 +208,9 @@ def test_qwen2vl_lora_beam_search(qwen2vl_lora_files):
        )


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
 def test_qwen25vl_lora(qwen25vl_lora_files):
    """Test Qwen 2.5 VL model with LoRA"""
    config = TestConfig(model_path=QWEN25VL_MODEL_PATH, lora_path=qwen25vl_lora_files)
@@ -216,6 +221,9 @@ def test_qwen25vl_lora(qwen25vl_lora_files):
        tester.run_test(TEST_IMAGES, expected_outputs=EXPECTED_OUTPUTS, lora_id=lora_id)


+@pytest.mark.skipif(
+    current_platform.is_cuda_alike(), reason="Skipping to avoid redundant model tests"
+)
 def test_qwen25vl_vision_lora(qwen25vl_vision_lora_files):
    config = TestConfig(
        model_path=QWEN25VL_MODEL_PATH,
@@ -124,30 +124,3 @@ def test_whisper_multi_lora(whisper_lora_files):
        f"Expected same outputs for same adapter with different IDs. "
        f"Got: {outputs_lora1} vs {outputs_lora2}"
    )
-
-
-@create_new_process_for_each_test()
-def test_whisper_with_and_without_lora(whisper_lora_files):
-    """Test that Whisper produces different outputs with and without LoRA.
-
-    This test verifies that the LoRA adapter actually affects the model output.
-    """
-    llm = create_whisper_llm(enable_lora=True)
-
-    # Run with LoRA
-    outputs_with_lora = run_whisper_inference(
-        llm, lora_path=whisper_lora_files, lora_id=1
-    )
-
-    # Run without LoRA (base model only)
-    outputs_without_lora = run_whisper_inference(llm, lora_path=None)
-
-    # Both should produce valid outputs
-    assert len(outputs_with_lora[0]) > 0
-    assert len(outputs_without_lora[0]) > 0
-
-    print(f"Output with LoRA: {outputs_with_lora[0]}")
-    print(f"Output without LoRA: {outputs_without_lora[0]}")
-
-    # Note: Outputs may or may not differ depending on the adapter
-    # The main verification is that both configurations work
@@ -4,9 +4,12 @@
 from collections.abc import Mapping

 import pytest
+import torch
 from PIL import Image as PILImage

+from vllm.model_executor.models.gemma4_mm import Gemma4ImagePixelInputs
 from vllm.multimodal import MULTIMODAL_REGISTRY
+from vllm.multimodal.inputs import MultiModalFieldConfig

 from ....conftest import ImageTestAssets
 from ...utils import build_model_context
@@ -15,6 +18,36 @@ from ...utils import build_model_context
 GEMMA4_MODEL_ID = "google/gemma-4-E2B-it"


+def test_gemma4_image_schema_accepts_variable_patch_counts():
+    Gemma4ImagePixelInputs(
+        pixel_values=[
+            torch.randn(10080, 768),
+            torch.randn(2520, 768),
+        ],
+        pixel_position_ids=[
+            torch.zeros(10080, 2, dtype=torch.long),
+            torch.zeros(2520, 2, dtype=torch.long),
+        ],
+    )
+
+
+def test_gemma4_image_batching_keeps_variable_patch_counts_unstacked():
+    field = MultiModalFieldConfig.batched("image").field
+    elems = field.build_elems(
+        "image",
+        "pixel_values",
+        [torch.randn(10080, 768), torch.randn(2520, 768)],
+    )
+
+    reduced = field.reduce_data(list(elems))
+
+    assert isinstance(reduced, list)
+    assert [tensor.shape for tensor in reduced] == [
+        torch.Size([10080, 768]),
+        torch.Size([2520, 768]),
+    ]
+
+
@pytest.mark.parametrize(
    "image_width,image_height,max_soft_tokens",
    [
@@ -0,0 +1,56 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+from types import SimpleNamespace
+
+import torch
+
+from vllm.model_executor.models.molmo2 import build_flat_image_bool_length
+
+
+def test_build_flat_image_bool_length_matches_molmoweb_processor_tokens():
+    hf_config = SimpleNamespace(
+        image_patch_id=151938,
+        low_res_image_start_token_id=151940,
+        image_start_token_id=151936,
+        image_col_id=151939,
+        image_end_token_id=151937,
+    )
+    image_grids = torch.tensor([[14, 14, 14, 23]], dtype=torch.long)
+
+    image_tokens, num_image_tokens = build_flat_image_bool_length(
+        image_grids,
+        hf_config,
+        image_use_col_tokens=True,
+        use_single_crop_col_tokens=None,
+        use_single_crop_start_token=False,
+    )
+
+    assert num_image_tokens.tolist() == [550]
+    assert len(image_tokens) == 550
+    assert image_tokens[0].item() == hf_config.image_start_token_id
+    assert (image_tokens == hf_config.image_col_id).sum().item() == 28
+
+
+def test_build_flat_image_bool_length_respects_disabled_col_tokens():
+    hf_config = SimpleNamespace(
+        image_patch_id=151938,
+        low_res_image_start_token_id=151940,
+        image_start_token_id=151936,
+        image_col_id=151939,
+        image_end_token_id=151937,
+    )
+    image_grids = torch.tensor([[2, 3, 5, 7]], dtype=torch.long)
+
+    image_tokens, num_image_tokens = build_flat_image_bool_length(
+        image_grids,
+        hf_config,
+        image_use_col_tokens=False,
+        use_single_crop_col_tokens=False,
+        use_single_crop_start_token=True,
+    )
+
+    assert num_image_tokens.tolist() == [45]
+    assert len(image_tokens) == 45
+    assert image_tokens[0].item() == hf_config.low_res_image_start_token_id
+    assert (image_tokens == hf_config.image_col_id).sum().item() == 0
@@ -53,6 +53,19 @@ class _VisionModel:
        self.loaded_weights = list(weights)


+class _FakeTensor:
+    """Sentinel stand-in for torch.Tensor in load_weights tests. Supports the
+    .detach().clone() chain used by load_weights for buffered mm weights;
+    both methods return self so identity (and the existing equality
+    assertions) are preserved through cloning."""
+
+    def detach(self):
+        return self
+
+    def clone(self):
+        return self
+
+
 def test_nano_nemotron_vl_skips_multimodal_weights_in_text_only_mode():
    model = object.__new__(NemotronH_Nano_VL_V2)
    language_model = _LanguageModel()
@@ -86,7 +99,7 @@ def test_nano_nemotron_vl_loads_vision_weights_without_sound_encoder():
    object.__setattr__(model, "sound_encoder", None)

    language_weight = object()
-    vision_weight = object()
+    vision_weight = _FakeTensor()
    model.load_weights(
        [
            ("language_model.layers.0.weight", language_weight),
@@ -28,6 +28,14 @@ def test_getattr_without_cache(monkeypatch: pytest.MonkeyPatch):
    assert not hasattr(envs.__getattr__, "cache_info")


+def test_nixl_side_channel_host_is_not_compile_factor(
+    monkeypatch: pytest.MonkeyPatch,
+):
+    monkeypatch.setenv("VLLM_NIXL_SIDE_CHANNEL_HOST", "10.0.0.15")
+
+    assert "VLLM_NIXL_SIDE_CHANNEL_HOST" not in envs.compile_factors()
+
+
 def test_getattr_with_cache(monkeypatch: pytest.MonkeyPatch):
    monkeypatch.setenv("VLLM_HOST_IP", "1.1.1.1")
    monkeypatch.setenv("VLLM_PORT", "1234")
@@ -6,6 +6,7 @@ import os
 from unittest.mock import patch

 from vllm.ray.ray_env import get_env_vars_to_copy
+from vllm.v1.executor.ray_utils import WORKER_SPECIFIC_ENV_VARS

 # ---------------------------------------------------------------------------
 # Default prefix matching
@@ -106,6 +107,19 @@ class TestExclusion:
        result = get_env_vars_to_copy(exclude_vars={"CUDA_VISIBLE_DEVICES"})
        assert "CUDA_VISIBLE_DEVICES" not in result

+    @patch.dict(
+        os.environ,
+        {
+            "VLLM_HOST_IP": "10.0.0.1",
+            "VLLM_NIXL_SIDE_CHANNEL_HOST": "10.0.0.1",
+        },
+        clear=False,
+    )
+    def test_worker_specific_host_vars_are_excluded(self):
+        result = get_env_vars_to_copy(exclude_vars=WORKER_SPECIFIC_ENV_VARS)
+        assert "VLLM_HOST_IP" not in result
+        assert "VLLM_NIXL_SIDE_CHANNEL_HOST" not in result
+
    @patch.dict(os.environ, {"LMCACHE_LOCAL_CPU": "True"}, clear=False)
    @patch(
        "vllm.ray.ray_env.RAY_NON_CARRY_OVER_ENV_VARS",
@@ -13,6 +13,7 @@ from vllm.model_executor.models.minimax_text_01 import MiniMaxText01LinearAttent
 from vllm.v1.attention.backends.linear_attn import LinearAttentionBackend
 from vllm.v1.attention.backends.mamba1_attn import Mamba1AttentionBackend
 from vllm.v1.attention.backends.mamba2_attn import Mamba2AttentionBackend
+from vllm.v1.attention.backends.registry import MambaAttentionBackendEnum
 from vllm.v1.attention.backends.short_conv_attn import ShortConvAttentionBackend


@@ -32,7 +33,7 @@ from vllm.v1.attention.backends.short_conv_attn import ShortConvAttentionBackend
                use_rms_norm=True,
            ),
            Mamba1AttentionBackend,
-            "mamba1",
+            MambaAttentionBackendEnum.MAMBA1,
        ),
        (
            MambaMixer2,
@@ -48,7 +49,7 @@ from vllm.v1.attention.backends.short_conv_attn import ShortConvAttentionBackend
                head_dim=32,
            ),
            Mamba2AttentionBackend,
-            "mamba2",
+            MambaAttentionBackendEnum.MAMBA2,
        ),
        (
            MiniMaxText01LinearAttention,
@@ -64,7 +65,7 @@ from vllm.v1.attention.backends.short_conv_attn import ShortConvAttentionBackend
                linear_layer_idx=0,
            ),
            LinearAttentionBackend,
-            "linear_attention",
+            MambaAttentionBackendEnum.LINEAR,
        ),
        (
            ShortConv,
@@ -74,7 +75,7 @@ from vllm.v1.attention.backends.short_conv_attn import ShortConvAttentionBackend
                layer_idx=0,
            ),
            ShortConvAttentionBackend,
-            "short_conv",
+            MambaAttentionBackendEnum.SHORT_CONV,
        ),
    ],
 )
@@ -97,10 +98,14 @@ def test_mamba_layers_get_attn_backend(
@pytest.mark.parametrize(
    "layer_class,expected_backend,expected_mamba_type",
    [
-        (MambaMixer, Mamba1AttentionBackend, "mamba1"),
-        (MambaMixer2, Mamba2AttentionBackend, "mamba2"),
-        (MiniMaxText01LinearAttention, LinearAttentionBackend, "linear_attention"),
-        (ShortConv, ShortConvAttentionBackend, "short_conv"),
+        (MambaMixer, Mamba1AttentionBackend, MambaAttentionBackendEnum.MAMBA1),
+        (MambaMixer2, Mamba2AttentionBackend, MambaAttentionBackendEnum.MAMBA2),
+        (
+            MiniMaxText01LinearAttention,
+            LinearAttentionBackend,
+            MambaAttentionBackendEnum.LINEAR,
+        ),
+        (ShortConv, ShortConvAttentionBackend, MambaAttentionBackendEnum.SHORT_CONV),
    ],
 )
 def test_mamba_layers_have_unified_interface(
@@ -0,0 +1,68 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import pytest
+import torch
+
+import vllm.envs as envs
+from tests.utils import TestFP8Layer, requires_fp8
+from vllm.model_executor.kernels.linear.scaled_mm.cutlass import (
+    CutlassFP8ScaledMMLinearKernel,
+)
+from vllm.model_executor.layers.quantization.utils.quant_utils import (
+    kFp8DynamicTokenSym,
+    kFp8StaticTensorSym,
+)
+from vllm.platforms import current_platform
+
+pytest.importorskip("torch.cuda")
+
+
+@pytest.fixture(autouse=True)
+def setup_cuda():
+    if not current_platform.is_cuda():
+        pytest.skip("CUTLASS FP8 kernels require CUDA.")
+    torch.set_default_device("cuda")
+
+
+@requires_fp8
+@pytest.mark.parametrize("weight_shape", [(1024, 2048), (4608, 4096)])
+@pytest.mark.parametrize("batch_size", [1, 16, 17, 32, 64, 65, 256, 257])
+@torch.inference_mode()
+def test_cutlass_fp8_batch_invariant_fixed_config(
+    weight_shape: tuple[int, int],
+    batch_size: int,
+    default_vllm_config,
+    monkeypatch: pytest.MonkeyPatch,
+):
+    monkeypatch.setenv("VLLM_BATCH_INVARIANT", "1")
+    monkeypatch.setattr(envs, "VLLM_BATCH_INVARIANT", True)
+
+    torch.manual_seed(0)
+    layer = TestFP8Layer(
+        weight_shape=weight_shape,
+        activation_quant_key=kFp8DynamicTokenSym,
+        weight_quant_key=kFp8StaticTensorSym,
+        input_dtype=torch.bfloat16,
+        out_dtype=torch.bfloat16,
+        device=torch.device("cuda"),
+        force_kernel=CutlassFP8ScaledMMLinearKernel,
+    )
+    assert isinstance(layer.kernel, CutlassFP8ScaledMMLinearKernel)
+
+    in_features = weight_shape[1]
+    needle = torch.randn((1, in_features), device="cuda", dtype=torch.bfloat16)
+    baseline = layer(needle)[0]
+
+    filler = torch.randn(
+        (max(batch_size - 1, 0), in_features), device="cuda", dtype=torch.bfloat16
+    )
+
+    front_batch = torch.cat([needle, filler], dim=0)
+    back_batch = torch.cat([filler, needle], dim=0)
+
+    front_output = layer(front_batch)[0]
+    back_output = layer(back_batch)[-1]
+
+    torch.testing.assert_close(front_output, baseline, rtol=0, atol=0)
+    torch.testing.assert_close(back_output, baseline, rtol=0, atol=0)
@@ -488,6 +488,10 @@ def _run_eagle_correctness(


@single_gpu_only
+@pytest.mark.skipif(
+    current_platform.is_device_capability_family(100),
+    reason="DeepSeek head_dim=192 not supported on SM100/SM110 (Blackwell)",
+)
@pytest.mark.parametrize(
    [
        "model_setup",
@@ -718,7 +722,15 @@ def test_eagle_correctness_heavy(
    ["model_setup", "mm_enabled", "expected_accuracy_threshold"],
    [
        (("mtp", "XiaomiMiMo/MiMo-7B-Base", 1), False, 0.5),  # ref: 65%-70%
-        (("mtp", "ZixiQi/DeepSeek-V3-4layers-MTP-FP8", 1), False, 0.0),  # dummy model
+        pytest.param(
+            ("mtp", "ZixiQi/DeepSeek-V3-4layers-MTP-FP8", 1),
+            False,
+            0.0,
+            marks=pytest.mark.skipif(
+                current_platform.is_device_capability_family(100),
+                reason="DeepSeek MTP: TRTLLM MoE top_k check fails on Blackwell",
+            ),
+        ),  # dummy model
        (
            ("mtp", "Qwen/Qwen3.5-0.8B-Base", 1),
            False,
@@ -0,0 +1,136 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import os
+import uuid
+from pathlib import Path
+from types import SimpleNamespace
+from typing import Any
+
+import pytest
+import ray
+
+from vllm.v1.engine.core import EngineCoreActorMixin
+from vllm.v1.engine.utils import CoreEngineActorManager, EngineZmqAddresses
+
+
+class _StubEngineCoreActor(EngineCoreActorMixin):
+    def __init__(
+        self,
+        vllm_config: Any,
+        local_client: bool,
+        addresses: EngineZmqAddresses,
+        executor_class: type[Any],
+        log_stats: bool,
+        dp_rank: int = 0,
+        local_dp_rank: int = 0,
+    ):
+        # Exercise the production Ray actor mixin without loading a model.
+        EngineCoreActorMixin.__init__(
+            self, vllm_config, addresses, dp_rank, local_dp_rank
+        )
+
+    def _set_visible_devices(self, vllm_config: Any, local_dp_rank: int) -> None:
+        pass
+
+    def wait_for_init(self) -> None:
+        pass
+
+    def run(self) -> None:
+        pass
+
+    def get_nixl_side_channel_host(self) -> str | None:
+        return os.environ.get("VLLM_NIXL_SIDE_CHANNEL_HOST")
+
+
+class _DummyExecutor:
+    pass
+
+
+def _make_vllm_config() -> SimpleNamespace:
+    return SimpleNamespace(
+        parallel_config=SimpleNamespace(
+            data_parallel_size=1,
+            data_parallel_size_local=1,
+            enable_elastic_ep=False,
+            world_size=1,
+        ),
+        model_config=SimpleNamespace(is_moe=False),
+        kv_transfer_config=None,
+    )
+
+
+def _make_addresses() -> EngineZmqAddresses:
+    return EngineZmqAddresses(
+        inputs=["tcp://127.0.0.1:12345"],
+        outputs=["tcp://127.0.0.1:12346"],
+    )
+
+
+def _make_cpu_placement_group():
+    pg = ray.util.placement_group(
+        [{"CPU": 0.001}, {"CPU": 1.0}],
+        strategy="PACK",
+    )
+    ray.get(pg.ready())
+    return pg
+
+
+@pytest.fixture
+def ray_context():
+    started_ray = False
+    if not ray.is_initialized():
+        project_root = str(Path(__file__).resolve().parents[3])
+        ray.init(
+            num_cpus=2,
+            runtime_env={"env_vars": {"PYTHONPATH": project_root}},
+            log_to_driver=False,
+        )
+        started_ray = True
+
+    yield
+
+    if started_ray:
+        ray.shutdown()
+
+
+@pytest.mark.usefixtures("ray_context")
+def test_driver_nixl_side_channel_host_does_not_leak_to_engine_core_actor(
+    monkeypatch: pytest.MonkeyPatch,
+) -> None:
+    driver_marker = f"driver-only-nixl-host-{uuid.uuid4()}"
+    created_placement_groups: list[Any] = []
+    manager: CoreEngineActorManager | None = None
+
+    def create_dp_placement_groups(vllm_config: Any):
+        pg = _make_cpu_placement_group()
+        created_placement_groups.append(pg)
+        return [pg], [0]
+
+    monkeypatch.setenv("VLLM_NIXL_SIDE_CHANNEL_HOST", driver_marker)
+    monkeypatch.setattr("vllm.v1.engine.core.EngineCoreActor", _StubEngineCoreActor)
+    monkeypatch.setattr(
+        CoreEngineActorManager,
+        "create_dp_placement_groups",
+        staticmethod(create_dp_placement_groups),
+    )
+
+    try:
+        manager = CoreEngineActorManager(
+            vllm_config=_make_vllm_config(),
+            addresses=_make_addresses(),
+            executor_class=_DummyExecutor,
+            log_stats=False,
+        )
+        actor = manager.local_engine_actors[0]
+        actor_host = ray.get(actor.get_nixl_side_channel_host.remote())
+        node_host = ray.util.get_node_ip_address()
+
+        assert actor_host != driver_marker
+        assert actor_host == node_host
+    finally:
+        if manager is not None:
+            manager.shutdown()
+        else:
+            for pg in created_placement_groups:
+                ray.util.remove_placement_group(pg)
@@ -0,0 +1,66 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Unit tests for LogprobsProcessor.
+
+These tests exercise the truncation invariant that the MRV2 sampler relies
+on: when the sampler returns a row wider than a request's own
+`num_logprobs + 1` (because another request in the batch needed a wider
+row), the trailing positions are populated with sentinel values
+(`token_id=0`, `logprob=-inf`). LogprobsProcessor must read only the first
+`num_logprobs + 1` entries so those sentinels never reach the user.
+"""
+
+import numpy as np
+
+from vllm.logprobs import create_sample_logprobs
+from vllm.v1.engine.logprobs import LogprobsProcessor
+from vllm.v1.outputs import LogprobsLists
+
+
+def _make_processor(num_logprobs: int) -> LogprobsProcessor:
+    return LogprobsProcessor(
+        tokenizer=None,
+        logprobs=create_sample_logprobs(flat_logprobs=False),
+        prompt_logprobs=None,
+        cumulative_logprob=0.0,
+        num_logprobs=num_logprobs,
+        num_prompt_logprobs=None,
+    )
+
+
+def test_drops_trailing_sentinel_columns():
+    """A request that asked for 3 custom token logprobs but ended up in a
+    batch padded to width 5 must not surface the trailing -inf entries."""
+    processor = _make_processor(num_logprobs=3)
+
+    sampled = 42
+    # Layout: [sampled, custom_1, custom_2, custom_3, SENTINEL, SENTINEL]
+    # Use float32-exact values so cumulative_logprob compares cleanly.
+    token_ids = np.array([[sampled, 100, 200, 300, 0, 0]], dtype=np.int32)
+    logprobs = np.array([[-0.5, -1.0, -2.0, -3.0, -np.inf, -np.inf]], dtype=np.float32)
+    ranks = np.array([1], dtype=np.int32)
+
+    processor._update_sample_logprobs(LogprobsLists(token_ids, logprobs, ranks))
+
+    assert len(processor.logprobs) == 1
+    pos = processor.logprobs[0]
+    # Exactly sampled + 3 requested tokens; trailing sentinels dropped.
+    assert set(pos.keys()) == {sampled, 100, 200, 300}
+    assert 0 not in pos
+    assert all(np.isfinite(lp.logprob) for lp in pos.values())
+    # cumulative_logprob comes from the sampled token's logprob only.
+    assert processor.cumulative_logprob == -0.5
+
+
+def test_accepts_exactly_sized_row():
+    """When the row is exactly num_logprobs+1, no truncation needed."""
+    processor = _make_processor(num_logprobs=2)
+
+    token_ids = np.array([[7, 11, 13]], dtype=np.int32)
+    logprobs = np.array([[-0.5, -1.5, -2.5]], dtype=np.float32)
+    ranks = np.array([1], dtype=np.int32)
+
+    processor._update_sample_logprobs(LogprobsLists(token_ids, logprobs, ranks))
+
+    pos = processor.logprobs[0]
+    assert set(pos.keys()) == {7, 11, 13}
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
khluuandClaude Opus 4.6	a8b0f65f44	Remove broken allocator pre-warming (initializes CUDA before fork) The pre-warming called torch.cuda.mem_get_info() in the parent process before the EngineCore subprocess forks, which initialized CUDA. This caused "Cannot re-initialize CUDA in forked subprocess" in the child process, breaking tests that previously passed. The pre-warming also wouldn't have helped — PyTorch's allocator pool is per-process, so the child's pool starts empty regardless. Signed-off-by: khluu <khluu000@gmail.com> Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-05-13 02:26:43 -07:00
khluuandClaude Opus 4.6	8a93cda9a9	[CI] Migrate all remaining gpu_1_queue jobs to h200_18gb MIG Move all test steps that were falling through to gpu_1_queue (L4) to explicitly use device: h200_18gb (H200 18GB MIG partition). This covers 52 steps across 15 test area files. Requires ci-infra PR #355 (PYTORCH_CUDA_ALLOC_CONF=expandable_segments:False) to be merged first — without it, PyTorch's CUDACachingAllocator hits an NVML assertion on MIG partitions when allocating large tensors. Signed-off-by: khluu <khluu000@gmail.com> Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-05-12 18:50:08 -07:00
khluuandClaude Opus 4.6	0ab46391b7	[Core] Add MIG detection and workaround for PyTorch NVML assertion On MIG (Multi-Instance GPU) partitions, NVML access is restricted. PyTorch's CUDACachingAllocator uses NVML internally when growing the memory pool, causing: NVML_SUCCESS == r INTERNAL ASSERT FAILED at "/pytorch/c10/cuda/CUDACachingAllocator.cpp":1165 This change: 1. Detects MIG via NVIDIA_VISIBLE_DEVICES="MIG-..." prefix 2. On MIG, sets PYTORCH_NVML_BASED_CUDA_CHECK=0 (use CUDA runtime instead of NVML for device checks) 3. On MIG, sets PYTORCH_CUDA_ALLOC_CONF=expandable_segments:False (avoids one NVML code path in the allocator) 4. Pre-warms the PyTorch allocator pool by allocating 90% of GPU memory upfront, so later KV cache allocations reuse the pool without triggering NVML-based pool growth Signed-off-by: khluu <khluu000@gmail.com> Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-05-12 15:57:19 -07:00
Kevin H. LuuandGitHub	1ff9d33535	[CI] Migrate remaining B200 jobs to b200-k8s with test fixes (#42387 ) Signed-off-by: khluu <khluu000@gmail.com>	2026-05-12 02:00:37 -07:00
chazen GitHub Chauncey	7f65f84428	[Bugfix] Fix empty channel/recipient in harmony for /v1/responses (#35540 ) Signed-off-by: kg6-sleipnir <christopherhazen42@gmail.com> Signed-off-by: chazen <45186108+kg6-sleipnir@users.noreply.github.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com>	2026-05-12 08:45:51 +00:00
amitz-nvandGitHub	ef34592a1a	[Bugfix] Fix double reduce in flashinfer_nvlink_two_sided and flashinfer_nvlink_one_sided backends (#41382 ) Signed-off-by: amitz-nv <203509407+amitz-nv@users.noreply.github.com>	2026-05-12 07:47:47 +00:00
Kevin H. LuuandGitHub	f69644caf8	[CI] Migrate more B200 jobs to b200-k8s queue (#42356 ) Signed-off-by: khluu <khluu000@gmail.com>	2026-05-12 00:38:31 -07:00
wang.yuqi GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	d37e25ffbe	[Frontend] Consolidate Speech to Text entrypoints. (#42370 ) Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io> Signed-off-by: wang.yuqi <noooop@126.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-05-12 07:06:57 +00:00
liuzhenwei GitHub Kunshang Ji	8517cdaf90	[XPU] update dp rank w/o env-var isolation (#39856 ) Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-05-12 14:49:54 +08:00
Lucas KabelaandGitHub	4e498b5e5c	[Bugfix][Performance Improvement] Improve penalties triton kernel performance (#40657 ) Signed-off-by: Lucas Kabela <lucaskabela@meta.com>	2026-05-12 05:47:20 +00:00
Yasmin Moslem GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	28ee78af54	Implement custom dataset class for ASR benchmarking (#41576 ) Signed-off-by: Yasmin Moslem <48152713+ymoslem@users.noreply.github.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-05-12 12:17:58 +08:00
ZiTian ZhaoandGitHub	630492da30	[Fix] Gemma4 Mixed-Resolution Image Co-Batching Crash (#42217 ) Signed-off-by: zitian.zhao <zitian.zhao@tencentmusic.com>	2026-05-12 03:13:03 +00:00
ChaunceyandGitHub	920bf3ec84	[Bugifx] [Qwen3CoderTool] Restore supports_required_and_named for required tool_choice (#42292 ) Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>	2026-05-12 02:09:56 +00:00
pschlan-amdandGitHub	39dff5ff39	Add VLLM_USE_SPINLOOP_EXT to use more efficient busy polling (#36517 ) Signed-off-by: Patrick Schlangen <pschlan@amd.com>	2026-05-11 16:11:49 -07:00
Wentao Ye GitHub Nick Hillmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	d7af6b34d8	[Model Runner V2] Bug fix: logprob dtype int64/int32 issue (#41761 ) Signed-off-by: yewentao256 <zhyanwentao@126.com> Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-11 21:55:43 +00:00
Nick Hill GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	bbee532988	[Perf][1/n] Eliminate various GPU<->CPU syncs (#41429 ) Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-11 20:36:03 +00:00
Jee Jee Li GitHub Yongye Zhu	53181384e0	[Bugfix] Fix DSV4 swiglu_limit on marlin backend (#42287 ) Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai> Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>	2026-05-11 13:03:56 -07:00
wang.yuqiandGitHub	a0dc7a0f36	[CI] Consolidate Speech to Text tests (#42274 ) Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>	2026-05-11 19:50:17 +00:00
Jeffrey Wang GitHub Tyler Michael Smith	56e5810ff1	[BugFix] Prevent orphaned process on NCCL destroy (#39846 ) Signed-off-by: Jeffrey Wang <jeffreywang@anyscale.com> Co-authored-by: Tyler Michael Smith <tyler@neuralmagic.com>	2026-05-11 15:25:26 -04:00
Flora FengandGitHub	639cbfd274	[CI] Add tests/parser to CI coverage (#41877 ) Signed-off-by: sfeng33 <4florafeng@gmail.com>	2026-05-11 19:08:54 +00:00
frida-andersson GitHub Markus Hartikainen Cursor Chuan Li Luka Govedič Frida Andersson TJian	a721315488	[ROCm][Perf] Fix RMSNorm+Quant fusion for gfx950 (non-fnuz) (#41825 ) Signed-off-by: Frida Andersson <fanderss@amd.com> Signed-off-by: Chuan Li <chuali@amd.com> Co-authored-by: Markus Hartikainen <markus.hartikainen@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Chuan Li <chuali@amd.com> Co-authored-by: Luka Govedič <ProExpertProg@users.noreply.github.com> Co-authored-by: Frida Andersson <frida-andersson@users.noreply.github.com> Co-authored-by: TJian <tunjian.tan@embeddedllm.com>	2026-05-11 15:00:51 -04:00
Vensen GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	6fdb49392e	[Bugfix] Fix int32 overflow in DeepGEMM SiLU/mul FP8 Triton kernel (#42201 ) Signed-off-by: vensen <vensenmu@gmail.com> Signed-off-by: Vensen <vensenmu@gmail.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	2026-05-11 14:52:31 -04:00
Alexa Griffith GitHub Claude Sonnet 4.5	cf0d279142	[Docs] Add Apple Silicon documentation for vLLM-Metal GPU support (#41987 ) Signed-off-by: alexagriffith <agriffith96@gmail.com> Co-authored-by: Claude Sonnet 4.5 <noreply@anthropic.com>	2026-05-11 11:34:25 -07:00
Vinay R Damodaran GitHub Russell Bryant	5497ffbf7c	Add documentation about vLLM FIPS compliance (#42190 ) Signed-off-by: Vinay Damodaran <vrdn@hey.com> Signed-off-by: Vinay R Damodaran <vrdn@hey.com> Co-authored-by: Russell Bryant <russell.bryant@gmail.com>	2026-05-11 18:17:02 +00:00
Nick HillandGitHub	9af6a5ed75	[Model Runner V2] Fix `seq_lens_cpu_upper_bound` (#42202 ) Signed-off-by: Nick Hill <nickhill123@gmail.com>	2026-05-11 10:37:50 -07:00
Hexiang WangandGitHub	7863fff6e5	[ROCm][DSv4] implement flash sparse mla with triton kernels (#41812 ) Signed-off-by: whx-sjtu <xiaowang990929@gmail.com>	2026-05-11 09:27:11 -07:00
Wentao YeandGitHub	0d453e2336	[Perf] Batch invariance with Cutlass fp8 support, 28.9% E2E latency improvement (#40408 ) Signed-off-by: yewentao256 <zhyanwentao@126.com> Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>	2026-05-11 12:20:58 -04:00
Wentao YeandGitHub	3f9c0c25b3	[Bug] Fix kimi dtype issue with `mm_projector_forward` (#42081 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-11 11:45:24 -04:00
Vadim GimpelsonandGitHub	a2e776d716	[Bugfix] Accept canonicalized `modelopt_*` quant_method in `_extract_modelopt_quant_algo` (#42181 ) Signed-off-by: Vadim Gimpelson <vadim.gimpelson@gmail.com>	2026-05-11 11:10:57 -04:00
Martin Hickey GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	4955990f1b	[kv_offload] Move `FilterReusedOffloadingManager` logic to `CPUOffloadingManager` (#41727 ) 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>	2026-05-11 18:09:29 +03:00
Wentao YeandGitHub	4b64fc2cbf	[Refactor] Cleanup batch invariant dead code (#41993 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-11 10:48:39 -04:00
pschlan-amdandGitHub	5f1b313900	[ROCm] Clean up a bit the AITER FA backend (#41942 ) Signed-off-by: Patrick Schlangen <pschlan@amd.com>	2026-05-11 22:45:18 +08:00
Thien Tran GitHub Yongye Zhu	724ed2fc35	[DSv4] Improved dequant gather K cache kernel (#42236 ) Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg> Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>	2026-05-11 10:41:12 -04:00
Rohan Potdar GitHub ElizaWszola	a51376b3f0	[Performance][DSR1]: Fused RoPE+KVCache+q_concat for MLA (#40392 ) Signed-off-by: Rohan138 <rohanpotdar138@gmail.com> Signed-off-by: Rohan Potdar <66227218+Rohan138@users.noreply.github.com> Co-authored-by: ElizaWszola <ewszola@redhat.com>	2026-05-11 14:10:50 +00:00
Martin HickeyandGitHub	8415bf2cdb	[kv_offload] Set offloading connector to prefer HND layout (#41928 ) Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com>	2026-05-11 15:05:41 +03:00
Noa NeriaandGitHub	ac062147fa	Avoid silent weights corruption when loading Nemotron Nano VL with reusable-buffer loaders like runai distributed streaming (#42244 ) Signed-off-by: Noa Neria <nneria@nvidia.com>	2026-05-11 12:03:14 +00:00
ChaunceyandGitHub	617239b70c	[Frontend]Responses API supports chat_template_kwargs (#42272 ) Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>	2026-05-11 11:59:39 +00:00
Kyungmin LeeandGitHub	27ae676364	Fix EXAONE-4.5 to align with Transformers update (#42246 ) Signed-off-by: lkm2835 <lkm2835@gmail.com>	2026-05-11 10:25:31 +00:00
haosdentandGitHub	17ed5e61f5	[CI] Make Python-only Installation optional (#42293 ) Signed-off-by: haosdent <haosdent@gmail.com>	2026-05-11 09:47:16 +00:00
Nicolò LucchesiandGitHub	5672d100ed	[KV Connector][NIXL][Bugfix] Fix NIXL handshake failures not honoring kv_load_failure_policy (#40364 ) When NIXL handshake fails (e.g., due to compatibility hash mismatch between prefill and decode instances), requests fail with "engine dead" error instead of gracefully falling back to local recomputation as configured by kv_load_failure_policy='recompute'. Signed-off-by: NickLucche <nlucches@redhat.com>	2026-05-11 09:37:21 +00:00
Nicolò LucchesiandGitHub	770e9bd6b3	[Nixl][PD] Lease renewal TTL KV blocks on P (#41383 ) Signed-off-by: NickLucche <nlucches@redhat.com>	2026-05-11 09:27:30 +00:00
Cyrus LeungandGitHub	9efdddca28	[Model] Fix missing `maybe_prefix` (#42280 ) Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk>	2026-05-11 09:04:06 +00:00
QiuandGitHub	b1b59720b2	bugfix(flashinfer,dcp): remove kv_cache_layout for BatchDCPPrefillWrapper._new_tokens. (#38895 ) Signed-off-by: QiuChunshuo <qiuchunshuo@huawei.com>	2026-05-11 08:11:49 +00:00
shaharmor98 GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	f9f770ca0b	fix nixl side-channel host selection (#41806 ) Signed-off-by: Shahar Mor <smor@nvidia.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-11 07:40:37 +00:00
Haoqing WangandGitHub	5cba6839e6	Document MolmoWeb hf_overrides (#42163 ) Signed-off-by: Haoqi Wang <78337154+hqhq1025@users.noreply.github.com>	2026-05-10 23:58:22 -07:00
Jee Jee LiandGitHub	05d610e5cd	[CI/Build] Reduce LoRA model tests. (#42266 ) Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>	2026-05-11 14:49:08 +08:00
汪志鹏 GitHubWang, Zhipeng \| RASIA <zhipeng.wang@rakuten.com>Cursor Chauncey	581b5e9afc	[Frontend] Return rendered prompt text in chat completion response (#42052 ) Signed-off-by: Wang, Zhipeng \| RASIA <zhipeng.wang@rakuten.com> Co-authored-by: Wang, Zhipeng \| RASIA <zhipeng.wang@rakuten.com> Co-authored-by: Cursor <cursor@cursor.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com>	2026-05-11 13:53:39 +08:00
wangxiyuanandGitHub	5536fc0c01	[Misc] Replace mamba_type string literals with MambaAttentionBackendEnum (#41188 ) Signed-off-by: wangxiyuan <wangxiyuan1007@gmail.com>	2026-05-11 03:59:36 +00:00
vllmellmandGitHub	7f95e66a11	[ROCm][Bugfix]: dynamically align BLOCK_DMODEL with Lv in MLA decode kernel (#41119 ) Signed-off-by: vllmellm <vllm.ellm@embeddedllm.com>	2026-05-11 11:14:19 +08:00
yzong-rhandGitHub	b1687527b8	[Bugfix] Gemma 4 chat template crash with missing tool name and tool id (#42188 ) Signed-off-by: Yifan <yzong@redhat.com>	2026-05-11 03:07:45 +00:00
gnovack GitHub george	171019ab19	add fused mhc_post_pre kernel (#41536 ) Signed-off-by: george <george@inferact.ai> Co-authored-by: george <george@inferact.ai>	2026-05-10 19:56:52 -07:00
Haoqing WangandGitHub	879a8c3180	Fix Molmo2 image token metadata (#42162 ) Signed-off-by: Haoqi Wang <78337154+hqhq1025@users.noreply.github.com>	2026-05-11 01:19:21 +00:00
bnellnm GitHub Jackmin801 Claude Robert Shaw Robert Shaw Jackmin801	1b57eb41f2	[MoE] Move various experts classes to fused_moe/experts/ (#41979 ) Signed-off-by: Jackmin801 <ongjackm@gmail.com> Signed-off-by: Robert Shaw <robertgshaw2@gmail.com> Signed-off-by: Jackmin801 <56836461+Jackmin801@users.noreply.github.com> Signed-off-by: Bill Nell <bnell@redhat.com> Co-authored-by: Jackmin801 <ongjackm@gmail.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Robert Shaw <robertgshaw2@gmail.com> Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com> Co-authored-by: Jackmin801 <56836461+Jackmin801@users.noreply.github.com>	2026-05-11 07:54:33 +08:00
Mohammad Miadh AngkadandGitHub	21943d4c25	[Performance] Make safetensors checkpoint prefetch settings configurable (#41499 ) Signed-off-by: Mohammad Miadh Angkad <MAngkad.BSDSBA2027@aim.edu>	2026-05-10 15:55:15 +00:00
Isotr0py GitHub qizixi	f396bee56f	[DSV4] Add PP support for deepseek-v4 (#41694 ) Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: qizixi <22851944+zixi-qi@users.noreply.github.com>	2026-05-10 15:47:26 +00:00
Vensen GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	215e2f7990	[Bugfix][Mamba] IMA in causal_conv1d kernel for long sequences (#41617 ) Signed-off-by: vensen <vensenmu@gmail.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-05-10 12:38:28 +00:00
Ronen SchafferandGitHub	e175192d33	[KV Offload] Pass ReqContext to touch(), complete_load(), and complete_store() (#41366 ) Signed-off-by: Ronen Schaffer <ronen.schaffer@ibm.com>	2026-05-10 15:09:25 +03:00
Jonathan Mamou GitHub Benjamin Chislett	a54f0d1049	[CPU] Fix spec decode kernel signatures for synthetic mode compatibility (#41932 ) Signed-off-by: jmamou <jonathan.mamou@intel.com> Signed-off-by: Jonathan Mamou <jonathan.mamou@intel.com> Co-authored-by: Benjamin Chislett <chislett.ben@gmail.com>	2026-05-10 12:07:15 +00:00
Isotr0pyandGitHub	48698b1b9b	[Bugfix] Fuse Qwen3.5 in_qkvz_proj forwarding with LoRA enabled (#37912 ) Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn> Signed-off-by: Isotr0py <Isotr0py@outlook.com>	2026-05-10 10:59:02 +00:00