[Bugfix] Fix DeepSeek V4 MTP HC state handling (#42320 )

Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com> (cherry picked from commit f1cc7aad3c)
[Bugfix] Add swiglu limits to deepgemm fp8 methods (#41986 )
2026-05-14 21:28:34 -07:00 · 2026-05-14 12:38:36 -07:00 · 2026-05-14 00:59:51 -07:00 · 2026-05-14 00:59:01 -07:00 · 2026-05-13 15:15:07 -07:00 · 2026-05-13 02:12:34 -07:00
414 changed files with 4783 additions and 14759 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_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
+  - 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

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

@@ -929,7 +929,6 @@ steps:
  - tests/tokenizers_
  - tests/reasoning
  - tests/tool_parsers
-  - tests/parser
  - tests/transformers_utils
  - tests/config
  commands:
@@ -943,7 +942,6 @@ 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

@@ -1322,6 +1320,7 @@ 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
@@ -1337,21 +1336,7 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
-
- 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
+  - 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

 - label: Entrypoints Integration (LLM) # TBD
  timeout_in_minutes: 180
@@ -1775,7 +1760,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_qwen3_with_multi_loras.py
+  - pytest -v -s -x lora/test_llm_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
@@ -2781,6 +2766,7 @@ 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
@@ -2796,21 +2782,7 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
-
- 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
+  - 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

 - label: Entrypoints Integration (Pooling) # TBD
  timeout_in_minutes: 180
@@ -11,7 +11,7 @@ steps:
  - tests/entrypoints/
  commands:
  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text
+  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling

 - label: Entrypoints Integration (LLM)
  key: entrypoints-integration-llm
@@ -44,6 +44,7 @@ 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
@@ -54,6 +55,7 @@ steps:
  - 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)
@@ -67,7 +69,7 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion  --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/speech_to_text/ --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
@@ -84,17 +86,6 @@ 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"
-  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
@@ -124,5 +115,5 @@ steps:
  - csrc/
  - vllm/entrypoints/openai/
  - vllm/model_executor/models/whisper.py
-  commands: # LMEval
+  commands: # LMEval+Transcription WER check
  - pytest -s entrypoints/openai/correctness/
@@ -91,7 +91,6 @@ 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
@@ -248,7 +247,7 @@ steps:
 - label: Kernels FusedMoE Layer Test (2 B200s)
  key: kernels-fusedmoe-layer-test-2-b200s
  timeout_in_minutes: 90
-  device: b200-k8s
+  device: b200
  num_devices: 2
  source_file_dependencies:
  - csrc/quantization/cutlass_w8a8/moe/
@@ -40,7 +40,7 @@ steps:
 - label: LM Eval Small Models (B200)
  key: lm-eval-small-models-b200
  timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
  optional: true
  source_file_dependencies:
  - csrc/
@@ -48,18 +48,6 @@ 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
@@ -104,7 +92,7 @@ steps:

 - label: MoE Refactor Integration Test (B200 DP - TEMPORARY)
  key: moe-refactor-integration-test-b200-dp-temporary
-  device: b200-k8s
+  device: b200
  optional: true
  num_devices: 2
  commands:
@@ -114,7 +102,6 @@ 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
@@ -140,7 +127,7 @@ steps:
 - label: GPQA Eval (GPT-OSS) (B200)
  key: gpqa-eval-gpt-oss-b200
  timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
  optional: true
  num_devices: 2
  source_file_dependencies:
@@ -9,7 +9,7 @@ steps:
  - 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_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 
+    - 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 
  parallelism: 4


@@ -31,7 +31,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_qwen3_with_multi_loras.py
+    - pytest -v -s -x lora/test_llm_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
@@ -79,7 +79,6 @@ steps:
    - tests/v1/core
    - tests/v1/executor
    - tests/v1/kv_offload
-    - tests/v1/simple_kv_offload
    - tests/v1/worker
    - tests/v1/kv_connector/unit
    - tests/v1/metrics
@@ -91,7 +90,6 @@ steps:
    - pytest -v -s -m 'not cpu_test' v1/core
    - pytest -v -s v1/executor
    - pytest -v -s v1/kv_offload
-    - pytest -v -s v1/simple_kv_offload
    - pytest -v -s v1/worker
    - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
    - pytest -v -s -m 'not cpu_test' v1/metrics
@@ -212,9 +210,7 @@ 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
@@ -286,7 +282,6 @@ steps:
  - tests/tokenizers_
  - tests/reasoning
  - tests/tool_parsers
-  - tests/parser
  - tests/transformers_utils
  - tests/config
  device: cpu-small
@@ -301,7 +296,6 @@ 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

@@ -106,12 +106,13 @@ steps:
  - vllm/v1/worker/gpu/
  - vllm/v1/worker/gpu_worker.py
  - tests/v1/spec_decode/test_max_len.py
-  - tests/v1/spec_decode/test_rejection_sampler_utils.py
+  - tests/v1/spec_decode/test_probabilistic_rejection_sampler_utils.py
+  - tests/v1/spec_decode/test_synthetic_rejection_sampler_utils.py
  - tests/v1/e2e/spec_decode/test_spec_decode.py
  commands:
  - set -x
  - export VLLM_USE_V2_MODEL_RUNNER=1
  - pytest -v -s v1/spec_decode/test_max_len.py -k "eagle or mtp"
-  - pytest -v -s v1/spec_decode/test_rejection_sampler_utils.py
+  - pytest -v -s v1/spec_decode/test_probabilistic_rejection_sampler_utils.py
  - pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py
  - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle or mtp"
@@ -5,7 +5,6 @@ steps:
 - label: Basic Models Tests (Initialization)
  key: basic-models-tests-initialization
  timeout_in_minutes: 45
-  device: h200_18gb
  torch_nightly: true
  source_file_dependencies:
  - vllm/
@@ -5,7 +5,6 @@ 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
@@ -16,7 +16,7 @@ steps:
 - label: Spec Decode Eagle Nightly B200
  key: spec-decode-eagle-nightly-b200
  timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
  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-k8s
+  device: b200
  optional: true
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -84,23 +84,23 @@ steps:
  commands:
    - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"

- label: Speculators Correctness
-  key: speculators-correctness
-  timeout_in_minutes: 60
+- label: DFlash Speculators Correctness
+  key: dflash-speculators-correctness
+  timeout_in_minutes: 30
  device: h100
  optional: true
  num_devices: 1
  source_file_dependencies:
    - vllm/v1/spec_decode/
    - vllm/model_executor/models/qwen3_dflash.py
-    - tests/v1/spec_decode/test_speculators_correctness.py
+    - tests/v1/spec_decode/test_speculators_dflash.py
  commands:
    - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
-    - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
+    - pytest -v -s v1/spec_decode/test_speculators_dflash.py -m slow_test

 - label: Spec Decode MTP hybrid (B200)
  timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
  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/speech_to_text/realtime @njhill
-/vllm/entrypoints/speech_to_text @NickLucche
+/vllm/entrypoints/openai/realtime @njhill
+/vllm/entrypoints/openai/speech_to_text @NickLucche
 /vllm/entrypoints/pooling @noooop
 /vllm/entrypoints/sagemaker @DarkLight1337
 /vllm/entrypoints/serve @njhill
@@ -109,24 +109,6 @@ 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.
 #
@@ -369,6 +351,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
  FetchContent_MakeAvailable(cutlass)

  list(APPEND VLLM_EXT_SRC
+    "csrc/quantization/awq/gemm_kernels.cu"
    "csrc/cutlass_extensions/common.cpp")

  set_gencode_flags_for_srcs(
@@ -500,6 +483,46 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
                   " in CUDA target architectures")
  endif()

+  # Only build AllSpark kernels if we are building for at least some compatible archs.
+  cuda_archs_loose_intersection(ALLSPARK_ARCHS "8.0;8.6;8.7;8.9" "${CUDA_ARCHS}")
+  if (ALLSPARK_ARCHS)
+    set(ALLSPARK_SRCS
+       "csrc/quantization/gptq_allspark/allspark_repack.cu"
+       "csrc/quantization/gptq_allspark/allspark_qgemm_w8a16.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${ALLSPARK_SRCS}"
+      CUDA_ARCHS "${ALLSPARK_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${ALLSPARK_SRCS}")
+    message(STATUS "Building AllSpark kernels for archs: ${ALLSPARK_ARCHS}")
+  else()
+    message(STATUS "Not building AllSpark kernels as no compatible archs found"
+                   " in CUDA target architectures")
+  endif()
+
+  # CUTLASS MLA Archs and flags
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
+  endif()
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
+    set(SRCS
+      "csrc/attention/mla/sm100_cutlass_mla_kernel.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${SRCS}"
+      CUDA_ARCHS "${MLA_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${SRCS}")
+    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
+    # Add MLA-specific include directories only to MLA source files
+    set_source_files_properties(${SRCS}
+      PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
+    message(STATUS "Building CUTLASS MLA for archs: ${MLA_ARCHS}")
+  else()
+    message(STATUS "Not building CUTLASS MLA as no compatible archs were found.")
+    # clear MLA_ARCHS
+    set(MLA_ARCHS)
+  endif()
+
  # Expert-specialization MXFP8 blockscaled grouped kernels (SM100+).
  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
    cuda_archs_loose_intersection(ES_MXFP8_GROUPED_MM_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
@@ -527,6 +550,24 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
    endif()
  endif()

+  # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
+  endif()
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
+    set(DSV3_FUSED_A_GEMM_SRC "csrc/dsv3_fused_a_gemm.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${DSV3_FUSED_A_GEMM_SRC}"
+      CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
+    list(APPEND VLLM_EXT_SRC ${DSV3_FUSED_A_GEMM_SRC})
+    message(STATUS "Building dsv3_fused_a_gemm for archs: ${DSV3_FUSED_A_GEMM_ARCHS}")
+  else()
+    message(STATUS "Not building dsv3_fused_a_gemm as no compatible archs found "
+                   "in CUDA target architectures.")
+  endif()
+
  #
  # Machete kernels

@@ -598,6 +639,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
  endif()


+  # Hadacore kernels
+  cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
+  if(HADACORE_ARCHS)
+    set(SRCS "csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${SRCS}"
+      CUDA_ARCHS "${HADACORE_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${SRCS}")
+    message(STATUS "Building hadacore")
+  endif()

 # if CUDA endif
 endif()
@@ -647,8 +698,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
    list(APPEND VLLM_STABLE_EXT_SRC
      "csrc/libtorch_stable/permute_cols.cu"
      "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
-      "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
-      "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu")
+      "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu")
  endif()

  if(VLLM_GPU_LANG STREQUAL "CUDA")
@@ -657,40 +707,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
      CUDA_ARCHS "${CUDA_ARCHS}")
  endif()

-  # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
-  endif()
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
-    set(SRCS "csrc/libtorch_stable/dsv3_fused_a_gemm.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building dsv3_fused_a_gemm for archs: ${DSV3_FUSED_A_GEMM_ARCHS}")
-  else()
-    message(STATUS "Not building dsv3_fused_a_gemm as no compatible archs found "
-                   "in CUDA target architectures.")
-  endif()
-
-  # Only build AllSpark kernels if we are building for at least some compatible archs.
-  cuda_archs_loose_intersection(ALLSPARK_ARCHS "8.0;8.6;8.7;8.9" "${CUDA_ARCHS}")
-  if (ALLSPARK_ARCHS)
-    set(SRCS
-       "csrc/libtorch_stable/quantization/gptq_allspark/allspark_repack.cu"
-       "csrc/libtorch_stable/quantization/gptq_allspark/allspark_qgemm_w8a16.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${ALLSPARK_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building AllSpark kernels for archs: ${ALLSPARK_ARCHS}")
-  else()
-    message(STATUS "Not building AllSpark kernels as no compatible archs found"
-                   " in CUDA target architectures")
-  endif()
-
  #
  # CUTLASS scaled_mm kernels (moved from _C to _C_stable_libtorch)
  #
@@ -1000,41 +1016,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
    endif()
  endif()

-  # CUTLASS MLA Archs and flags
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
-  endif()
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
-    set(SRCS
-      "csrc/libtorch_stable/attention/mla/sm100_cutlass_mla_kernel.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${MLA_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
-    # Add MLA-specific include directories only to MLA source files
-    set_source_files_properties(${SRCS}
-      PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
-    message(STATUS "Building CUTLASS MLA for archs: ${MLA_ARCHS}")
-  else()
-    message(STATUS "Not building CUTLASS MLA as no compatible archs were found.")
-    # clear MLA_ARCHS
-    set(MLA_ARCHS)
-  endif()
-
-  # Hadacore kernels
-  cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
-  if(HADACORE_ARCHS)
-    set(SRCS "csrc/libtorch_stable/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${HADACORE_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building hadacore")
-  endif()
-
  message(STATUS "Enabling C_stable extension.")
  define_extension_target(
    _C_stable_libtorch
@@ -53,6 +53,7 @@ backends:
  - FLASHINFER_MLA
  - FLASH_ATTN_MLA  # Hopper only
  - FLASHMLA        # Hopper only
+  - TOKENSPEED_MLA  # Blackwell + R1 dims + FP8 KV (use --kv-cache-dtype fp8)

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

 device: "cuda:0"
 repeats: 20
@@ -179,19 +179,27 @@ def create_minimal_vllm_config(

    if prefill_backend is not None:
        prefill_cfg = get_prefill_backend_config(prefill_backend)
-        if prefill_cfg["flash_attn_version"] is not None:
-            vllm_config.attention_config.flash_attn_version = prefill_cfg[
-                "flash_attn_version"
+        if prefill_cfg.get("mla_prefill_backend_enum") is not None:
+            # Registry-based backends bypass the deprecated boolean flags.
+            from vllm.v1.attention.backends.mla.prefill import MLAPrefillBackendEnum
+
+            vllm_config.attention_config.mla_prefill_backend = MLAPrefillBackendEnum[
+                prefill_cfg["mla_prefill_backend_enum"]
            ]
-        vllm_config.attention_config.disable_flashinfer_prefill = prefill_cfg[
-            "disable_flashinfer_prefill"
-        ]
-        vllm_config.attention_config.use_cudnn_prefill = prefill_cfg[
-            "use_cudnn_prefill"
-        ]
-        vllm_config.attention_config.use_trtllm_ragged_deepseek_prefill = prefill_cfg[
-            "use_trtllm_ragged_deepseek_prefill"
-        ]
+        else:
+            if prefill_cfg["flash_attn_version"] is not None:
+                vllm_config.attention_config.flash_attn_version = prefill_cfg[
+                    "flash_attn_version"
+                ]
+            vllm_config.attention_config.disable_flashinfer_prefill = prefill_cfg[
+                "disable_flashinfer_prefill"
+            ]
+            vllm_config.attention_config.use_cudnn_prefill = prefill_cfg[
+                "use_cudnn_prefill"
+            ]
+            vllm_config.attention_config.use_trtllm_ragged_deepseek_prefill = (
+                prefill_cfg["use_trtllm_ragged_deepseek_prefill"]
+            )

    return vllm_config

@@ -223,22 +231,17 @@ _PREFILL_BACKEND_CONFIG: dict[str, dict] = {
        "use_trtllm_ragged_deepseek_prefill": False,
    },
    "flashinfer": {
-        "flash_attn_version": None,
-        "disable_flashinfer_prefill": False,
-        "use_cudnn_prefill": False,
-        "use_trtllm_ragged_deepseek_prefill": False,
+        "mla_prefill_backend_enum": "FLASHINFER",
    },
    "cudnn": {
-        "flash_attn_version": None,
-        "disable_flashinfer_prefill": True,
-        "use_cudnn_prefill": True,
-        "use_trtllm_ragged_deepseek_prefill": False,
+        # cuDNN prefill backend was removed; AttentionConfig raises on use.
+        "mla_prefill_backend_enum": "FLASHINFER",
    },
    "trtllm": {
-        "flash_attn_version": None,
-        "disable_flashinfer_prefill": True,
-        "use_cudnn_prefill": False,
-        "use_trtllm_ragged_deepseek_prefill": True,
+        "mla_prefill_backend_enum": "TRTLLM_RAGGED",
+    },
+    "tokenspeed": {
+        "mla_prefill_backend_enum": "TOKENSPEED_MLA",
    },
 }

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

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

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

-        # Verify the actual prefill backend matches what was requested
+        # Verify the actual prefill backend matches what was requested. The
+        # selector + impl construction already raise on misuse; here we just
+        # check the resolved class against the requested name as a sanity guard.
        if prefill_backend is not None:
-            prefill_cfg = get_prefill_backend_config(prefill_backend)
-            fa_version = prefill_cfg["flash_attn_version"]
-
-            if fa_version is not None:
-                # FA backend: verify the impl's FA version
-                actual_fa_version = getattr(impl, "vllm_flash_attn_version", None)
+            expected_class = {
+                "fa2": "FlashAttnPrefillBackend",
+                "fa3": "FlashAttnPrefillBackend",
+                "fa4": "FlashAttnPrefillBackend",
+                "flashinfer": "FlashInferPrefillBackend",
+                "trtllm": "TrtllmRaggedPrefillBackend",
+                "tokenspeed": "TokenspeedMLAPrefillBackend",
+            }.get(prefill_backend)
+            actual_class = type(getattr(layer, "prefill_backend", None)).__name__
+            if expected_class and actual_class != expected_class:
+                raise RuntimeError(
+                    f"Prefill backend '{prefill_backend}' requested "
+                    f"{expected_class}, got {actual_class}. Check "
+                    f"attention_config plumbing or installed deps."
+                )
+            if prefill_backend in {"fa2", "fa3", "fa4"}:
+                fa_version = int(prefill_backend[2:])
+                actual_fa_version = getattr(
+                    layer.prefill_backend, "vllm_flash_attn_version", None
+                )
                if actual_fa_version != fa_version:
                    raise RuntimeError(
                        f"Prefill backend '{prefill_backend}' requested FA "
-                        f"version {fa_version}, but the impl is using FA "
-                        f"version {actual_fa_version}. Check "
-                        f"vllm/v1/attention/backends/fa_utils.py."
-                    )
-            else:
-                # Non-FA backend: verify the builder picked the right path
-                expected_flags = {
-                    "flashinfer": "_use_fi_prefill",
-                    "cudnn": "_use_cudnn_prefill",
-                    "trtllm": "_use_trtllm_ragged_prefill",
-                }
-                flag_name = expected_flags.get(prefill_backend)
-                if flag_name and not getattr(builder_instance, flag_name, False):
-                    raise RuntimeError(
-                        f"Prefill backend '{prefill_backend}' was requested "
-                        f"but the metadata builder did not enable it. This "
-                        f"usually means a dependency is missing (e.g., "
-                        f"flashinfer not installed) or the platform doesn't "
-                        f"support it."
+                        f"version {fa_version}, got "
+                        f"{actual_fa_version} on {actual_class}."
                    )

        # Run each benchmark with the shared impl
@@ -27,10 +27,10 @@ from vllm.model_executor.layers.fused_moe.config import (
    RoutingMethodType,
    _get_config_dtype_str,
 )
-from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
+from vllm.model_executor.layers.fused_moe.fused_moe import *
+from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
    TritonOrDeepGemmExperts,
 )
-from vllm.model_executor.layers.fused_moe.fused_moe import *
 from vllm.transformers_utils.config import get_config
 from vllm.triton_utils import triton
 from vllm.utils.argparse_utils import FlexibleArgumentParser
@@ -53,30 +53,14 @@ cuda_archs_loose_intersection(DEEPGEMM_ARCHS
 if(DEEPGEMM_ARCHS)
  message(STATUS "DeepGEMM CUDA architectures: ${DEEPGEMM_ARCHS}")

-  #
-  # DeepGEMM integration notes
-  # --------------------------
-  # We vendor DeepGEMM into vllm/third_party/deep_gemm/ and bundle a
-  # `_C.cpython-X.Y-*.so` for every CPython in `requires-python`. The
-  # per-Python build is delegated to tools/build_deepgemm_C.py.
-  #
-  # Why per-Python: DeepGEMM's binding uses PYBIND11_MODULE, which links
-  # private CPython symbols — a single `_C.abi3.so` is not viable today
-  # (see #41476 / #41512 for the failed attempt).
-  #
-  # TODOs (tracked in vllm-project/vllm#42431):
-  #   - Replace DeepGEMM's pybind11 binding with a TORCH_LIBRARY + shim
-  #     binding (cf. vllm-flash-attention/csrc/common/pytorch_shim.h) to
-  #     collapse to one `_C.abi3.so`. Needs either an upstream change or
-  #     a maintained binding fork in vLLM.
-  #   - AOT-compile DeepGEMM's CUDA kernels instead of runtime JIT to drop
-  #     the vendored CUTLASS/CCCL headers and the CUDA-toolkit-at-runtime
-  #     requirement.
-  #
-
-  # DEEPGEMM_PYTHON_INTERPRETERS: ":"-separated target Python paths.
-  # Empty/unset → fall back to the build interpreter (editable installs).
-  # (Empty-but-set env vars test as DEFINED in cmake — treat as unset.)
+  # Build _C once per interpreter in DEEPGEMM_PYTHON_INTERPRETERS (":"-
+  # separated paths) so the wheel imports cleanly on every supported Python.
+  # Unset → fall back to the build interpreter (editable / source builds).
+  # The compile is delegated to tools/build_deepgemm_C.py and always runs
+  # against the build interpreter's torch — target Pythons don't need torch.
+  # Note: empty-but-set env vars are still DEFINED in cmake; treat empty as
+  # unset so an empty interpreter list falls back to the build interpreter
+  # rather than silently skipping the per-Python build.
  if(NOT "$ENV{DEEPGEMM_PYTHON_INTERPRETERS}" STREQUAL "")
    string(REPLACE ":" ";" _dg_pythons "$ENV{DEEPGEMM_PYTHON_INTERPRETERS}")
  else()
@@ -84,8 +68,10 @@ if(DEEPGEMM_ARCHS)
  endif()
  message(STATUS "DeepGEMM _C will be built for: ${_dg_pythons}")

-  # add_custom_command does no implicit header scanning; glob explicitly so
-  # header-only edits in DeepGEMM/cutlass/fmt re-trigger the rebuild.
+  # Header set fed to add_custom_command's DEPENDS so a header-only edit
+  # (in upstream DeepGEMM or its vendored cutlass/fmt) re-triggers the
+  # rebuild. add_custom_command does no implicit header scanning, unlike
+  # add_library.
  file(GLOB_RECURSE _dg_headers
    "${deepgemm_SOURCE_DIR}/csrc/*.h"
    "${deepgemm_SOURCE_DIR}/csrc/*.hpp"
@@ -102,7 +88,8 @@ if(DEEPGEMM_ARCHS)
      OUTPUT_VARIABLE _dg_soabi
      OUTPUT_STRIP_TRAILING_WHITESPACE
      COMMAND_ERROR_IS_FATAL ANY)
-    # Dedup interpreters that resolve to the same CPython.
+    # Dedup so duplicate paths (or two paths resolving to the same CPython)
+    # don't register conflicting build rules.
    if(_dg_soabi IN_LIST _dg_seen_soabis)
      continue()
    endif()
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG bce29425653ec0fbc579d329883030e832d15ada
+          GIT_TAG f5bc33cfc02c744d24a2e9d50e6db656de40611c
          GIT_PROGRESS TRUE
          # Don't share the vllm-flash-attn build between build types
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
@@ -18,12 +18,13 @@ limitations under the License.
 * Taken from SGLANG PR https://github.com/sgl-project/sglang/pull/6929
 * by Alcanderian JieXin Liang
 */
-#include "libtorch_stable/torch_utils.h"
-
-#include <torch/csrc/stable/library.h>
+#include "core/registration.h"

+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
 #include <cutlass/cutlass.h>
 #include <cutlass/kernel_hardware_info.h>
+#include <torch/all.h>

 #include <cute/tensor.hpp>
 #include <iostream>
@@ -34,27 +35,27 @@ limitations under the License.
 // clang-format off
 #if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
 void sm100_cutlass_mla_decode(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    torch::Tensor const& out,
+    torch::Tensor const& lse,
+    torch::Tensor const& q_nope,
+    torch::Tensor const& q_pe,
+    torch::Tensor const& kv_c_and_k_pe_cache,
+    torch::Tensor const& seq_lens,
+    torch::Tensor const& page_table,
+    torch::Tensor const& workspace,
    double sm_scale,
    int64_t num_kv_splits) {
-  STD_TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
+  TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
 }
 int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_batches, int64_t sm_count, int64_t num_kv_splits) {
-  STD_TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
+  TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
 }
 #else

 #define CUTLASS_CHECK(status)                                                       \
  {                                                                                 \
    cutlass::Status error = status;                                                 \
-    STD_TORCH_CHECK(error == cutlass::Status::kSuccess, cutlassGetStatusString(error)); \
+    TORCH_CHECK(error == cutlass::Status::kSuccess, cutlassGetStatusString(error)); \
  }

 using namespace cute;
@@ -99,23 +100,23 @@ struct MlaSm100 {

 template <typename T>
 typename T::Fmha::Arguments args_from_options(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
+    at::Tensor const& out,
+    at::Tensor const& lse,
+    at::Tensor const& q_nope,
+    at::Tensor const& q_pe,
+    at::Tensor const& kv_c_and_k_pe_cache,
+    at::Tensor const& seq_lens,
+    at::Tensor const& page_table,
    double sm_scale,
    int64_t num_kv_splits) {
  cutlass::KernelHardwareInfo hw_info;
-  hw_info.device_id = q_nope.get_device_index();
+  hw_info.device_id = q_nope.device().index();
  hw_info.sm_count = cutlass::KernelHardwareInfo::query_device_multiprocessor_count(hw_info.device_id);

-  int batches = q_nope.size(0);
-  int page_count_per_seq = page_table.size(1);
-  int page_count_total = kv_c_and_k_pe_cache.size(0);
-  int page_size = kv_c_and_k_pe_cache.size(1);
+  int batches = q_nope.sizes()[0];
+  int page_count_per_seq = page_table.sizes()[1];
+  int page_count_total = kv_c_and_k_pe_cache.sizes()[0];
+  int page_size = kv_c_and_k_pe_cache.sizes()[1];
  int max_seq_len = page_size * page_count_per_seq;
  using TileShapeH = typename T::TileShapeH;
  using TileShapeD = typename T::TileShapeD;
@@ -185,14 +186,14 @@ typename T::Fmha::Arguments args_from_options(

 template <typename Element, typename ElementOut, bool IsPaged128, typename PersistenceOption>
 void runMla(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    at::Tensor const& out,
+    at::Tensor const& lse,
+    at::Tensor const& q_nope,
+    at::Tensor const& q_pe,
+    at::Tensor const& kv_c_and_k_pe_cache,
+    at::Tensor const& seq_lens,
+    at::Tensor const& page_table,
+    at::Tensor const& workspace,
    double sm_scale,
    int64_t num_kv_splits,
    cudaStream_t stream) {
@@ -219,37 +220,37 @@ void runMla(
  }()

 void sm100_cutlass_mla_decode(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    torch::Tensor const& out,
+    torch::Tensor const& lse,
+    torch::Tensor const& q_nope,
+    torch::Tensor const& q_pe,
+    torch::Tensor const& kv_c_and_k_pe_cache,
+    torch::Tensor const& seq_lens,
+    torch::Tensor const& page_table,
+    torch::Tensor const& workspace,
    double sm_scale,
    int64_t num_kv_splits) {
-  auto in_dtype = q_nope.scalar_type();
-  torch::stable::accelerator::DeviceGuard device_guard(q_nope.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream(q_nope.get_device_index());
-  const int page_size = kv_c_and_k_pe_cache.size(1);
-
+  auto in_dtype = q_nope.dtype();
+  at::cuda::CUDAGuard device_guard{(char)q_nope.get_device()};
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream(q_nope.get_device());
+  const int page_size = kv_c_and_k_pe_cache.sizes()[1];
+  
  // NOTE(alcanderian): IsPersistent has bug with manual split_kv.
  // Kernel will hang if batch is too large with large num_kv_splits. (for example bs=8, num_kv_splits=8)
  // Maybe per batch split kv will fix this.
  DISPATCH_BOOL(page_size == 128, IsPaged128, [&] {
    DISPATCH_BOOL(num_kv_splits <= 1, NotManualSplitKV, [&] {
-      if (in_dtype == torch::headeronly::ScalarType::Half) {
+      if (in_dtype == at::ScalarType::Half) {
        runMla<cutlass::half_t, cutlass::half_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
          out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
-      } else if (in_dtype == torch::headeronly::ScalarType::BFloat16) {
+      } else if (in_dtype == at::ScalarType::BFloat16) {
        runMla<cutlass::bfloat16_t, cutlass::bfloat16_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
          out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
-      } else if (in_dtype == torch::headeronly::ScalarType::Float8_e4m3fn) {
+      } else if (in_dtype == at::ScalarType::Float8_e4m3fn) {
        runMla<cutlass::float_e4m3_t, cutlass::bfloat16_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
          out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
      } else {
-        STD_TORCH_CHECK(false, "Unsupported input data type of MLA");
+        TORCH_CHECK(false, "Unsupported input data type of MLA");
      }
      return true;
    });
@@ -279,12 +280,12 @@ int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_ba

 #endif

-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("sm100_cutlass_mla_decode", TORCH_BOX(&sm100_cutlass_mla_decode));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+  m.impl("sm100_cutlass_mla_decode", &sm100_cutlass_mla_decode);
 }

-STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, m) {
-  m.impl("sm100_cutlass_mla_get_workspace_size", TORCH_BOX(&sm100_cutlass_mla_get_workspace_size));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CatchAll, m) {
+  m.impl("sm100_cutlass_mla_get_workspace_size", &sm100_cutlass_mla_get_workspace_size);
 }

 // clang-format on
@@ -21,33 +21,28 @@ 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, typename cos_sin_t, bool IS_NEOX,
-          typename raw_kv_scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt>
+template <typename qk_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 cos_sin_t* __restrict__ rope_cos_sin_cache,  // [max_position, 2,
-                                                       // rot_dim // 2]
+    const qk_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__ slot_mapping,  // [num_tokens]
+    const int64_t* __restrict__ kv_cache_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 cos_sin_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;
+  const qk_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;

  const int embed_dim = rot_dim / 2;

@@ -59,8 +54,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 = 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 cos = VLLM_LDG(cos_sin_ptr + pair_idx);
+    qk_t sin = 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;
@@ -86,15 +81,21 @@ __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 = 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 cos = VLLM_LDG(cos_sin_ptr + pair_idx);
+    qk_t sin = VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim);

    qk_t* k_pe_head_ptr = k_pe + token_idx * k_pe_stride;

@@ -164,43 +165,36 @@ __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;                                                  \
-      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>());                \
-            }                                                                 \
-          });                                                                 \
-    });                                                                       \
+#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>());                       \
+      }                                                                        \
+    });                                                                        \
  } while (false)

 // Executes RoPE on q_pe and k_pe, then writes k_pe and kv_c in the kv cache.
@@ -214,52 +208,43 @@ 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& slot_mapping,  // [num_tokens] or [num_actual_tokens]
+    torch::Tensor&
+        kv_cache_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) {
-  // 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 int64_t num_tokens = q_pe.size(0);

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

-  TORCH_CHECK_EQ(q_pe.dim(), 3);
-  TORCH_CHECK_EQ(q_pe.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(q_pe.size(0), num_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.dim(), 2);
-  TORCH_CHECK_EQ(k_pe.size(0), num_padded_tokens);
+  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.scalar_type(), q_pe.scalar_type());

-  TORCH_CHECK_EQ(kv_c.dim(), 2);
-  TORCH_CHECK_EQ(kv_c.size(0), num_padded_tokens);
+  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.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(slot_mapping.size(0), num_tokens);
-  TORCH_CHECK_EQ(slot_mapping.scalar_type(), c10::ScalarType::Long);
+  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(kv_cache.size(2), kv_lora_rank + rot_dim);
  TORCH_CHECK_EQ(kv_cache.dim(), 3);
@@ -1,6 +1,7 @@
 #pragma once
 #include <cstdlib>
 #include <string>
+#include <cctype>

 namespace vllm {

@@ -1,13 +1,7 @@
 #pragma once

-#include <cstdint>
-#include <string>
-#include <tuple>
-#include <utility>
-#include <variant>
-
-// For STD_TORCH_CHECK
-#include <torch/headeronly/util/Exception.h>
+// For TORCH_CHECK
+#include <torch/library.h>

 namespace vllm {

@@ -51,7 +45,7 @@ class ScalarType {
  // IEEE 754 compliant floating point type
  static constexpr ScalarType float_IEEE754(uint8_t exponent,
                                            uint8_t mantissa) {
-    STD_TORCH_CHECK(mantissa > 0 && exponent > 0);
+    TORCH_CHECK(mantissa > 0 && exponent > 0);
    return ScalarType(exponent, mantissa, true, 0, false, NAN_IEEE_754);
  }

@@ -59,12 +53,11 @@ class ScalarType {
  static constexpr ScalarType float_(uint8_t exponent, uint8_t mantissa,
                                     bool finite_values_only,
                                     NanRepr nan_repr) {
-    STD_TORCH_CHECK(nan_repr < NAN_REPR_ID_MAX, "Invalid NanRepr");
-    STD_TORCH_CHECK(mantissa > 0 && exponent > 0);
-    STD_TORCH_CHECK(
-        nan_repr != NAN_IEEE_754,
-        "use `float_IEEE754` constructor for floating point types that "
-        "follow IEEE 754 conventions");
+    TORCH_CHECK(nan_repr < NAN_REPR_ID_MAX, "Invalid NanRepr");
+    TORCH_CHECK(mantissa > 0 && exponent > 0);
+    TORCH_CHECK(nan_repr != NAN_IEEE_754,
+                "use `float_IEEE754` constructor for floating point types that "
+                "follow IEEE 754 conventions");
    return ScalarType(exponent, mantissa, true, 0, finite_values_only,
                      nan_repr);
  }
@@ -183,8 +176,8 @@ class ScalarType {

 private:
  double _floating_point_max() const {
-    STD_TORCH_CHECK(mantissa <= 52 && exponent <= 11,
-                    "Cannot represent max/min as a double for type ", str());
+    TORCH_CHECK(mantissa <= 52 && exponent <= 11,
+                "Cannot represent max/min as a double for type ", str());

    uint64_t max_mantissa = (uint64_t(1) << mantissa) - 1;
    if (nan_repr == NAN_EXTD_RANGE_MAX_MIN) {
@@ -193,8 +186,8 @@ class ScalarType {

    uint64_t max_exponent = (uint64_t(1) << exponent) - 2;
    if (nan_repr == NAN_EXTD_RANGE_MAX_MIN || nan_repr == NAN_NONE) {
-      STD_TORCH_CHECK(exponent < 11,
-                      "Cannot represent max/min as a double for type ", str());
+      TORCH_CHECK(exponent < 11,
+                  "Cannot represent max/min as a double for type ", str());
      max_exponent += 1;
    }

@@ -223,17 +216,16 @@ class ScalarType {
    if (is_floating_point()) {
      return {_floating_point_max()};
    } else {
-      STD_TORCH_CHECK(size_bits() < 64 || size_bits() == 64 && is_signed(),
-                      "Cannot represent max as a int64_t");
+      TORCH_CHECK(size_bits() < 64 || size_bits() == 64 && is_signed(),
+                  "Cannot represent max as a int64_t");
      return {(int64_t(1) << mantissa) - 1};
    }
  }

  constexpr std::variant<int64_t, double> _raw_min() const {
    if (is_floating_point()) {
-      STD_TORCH_CHECK(
-          is_signed(),
-          "We currently assume all floating point types are signed");
+      TORCH_CHECK(is_signed(),
+                  "We currently assume all floating point types are signed");
      constexpr uint64_t sign_bit_double = (uint64_t(1) << 63);

      double max = _floating_point_max();
@@ -241,8 +233,8 @@ class ScalarType {
      uint64_t min_raw = max_raw | sign_bit_double;
      return {*reinterpret_cast<double*>(&min_raw)};
    } else {
-      STD_TORCH_CHECK(!is_signed() || size_bits() <= 64,
-                      "Cannot represent min as a int64_t");
+      TORCH_CHECK(!is_signed() || size_bits() <= 64,
+                  "Cannot represent min as a int64_t");
      if (is_signed()) {
        // set the top bit to 1 (i.e. INT64_MIN) and the rest to 0
        // then perform an arithmetic shift right to set all the bits above
@@ -20,15 +20,13 @@
 * limitations under the License.
 */

-#include <torch/csrc/stable/library.h>
-#include <torch/csrc/stable/tensor.h>
-#include <torch/headeronly/core/ScalarType.h>
-
-#include "core/registration.h"
-#include "libtorch_stable/torch_utils.h"
-
+#include <ATen/ATen.h>
+#include <ATen/cuda/CUDAContext.h>
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
+#include <torch/all.h>
+
+#include "core/registration.h"

 #include <cstdlib>
 #include <mutex>
@@ -36,7 +34,7 @@
 namespace {

 inline int getSMVersion() {
-  auto* props = get_device_prop();
+  auto* props = at::cuda::getCurrentDeviceProperties();
  return props->major * 10 + props->minor;
 }

@@ -702,40 +700,37 @@ template void invokeFusedAGemm<__nv_bfloat16, 7168, 2112, 16>(
    __nv_bfloat16*, __nv_bfloat16 const*, __nv_bfloat16 const*, int num_tokens,
    cudaStream_t);

-void dsv3_fused_a_gemm(torch::stable::Tensor& output,
-                       torch::stable::Tensor const& mat_a,
-                       torch::stable::Tensor const& mat_b) {
-  STD_TORCH_CHECK(mat_a.dim() == 2 && mat_b.dim() == 2 && output.dim() == 2);
+void dsv3_fused_a_gemm(torch::Tensor& output, torch::Tensor const& mat_a,
+                       torch::Tensor const& mat_b) {
+  TORCH_CHECK(mat_a.dim() == 2 && mat_b.dim() == 2 && output.dim() == 2);
  int const num_tokens = mat_a.size(0);
  int const hd_in = mat_a.size(1);
  int const hd_out = mat_b.size(1);

  constexpr int kHdIn = 7168;
  constexpr int kHdOut = 2112;
-  STD_TORCH_CHECK(num_tokens >= 1 && num_tokens <= 16,
-                  "required 1 <= mat_a.shape[0] <= 16");
-  STD_TORCH_CHECK(hd_in == kHdIn, "required mat_a.shape[1] == 7168");
-  STD_TORCH_CHECK(hd_out == kHdOut, "required mat_b.shape[1] == 2112");
-  STD_TORCH_CHECK(output.size(0) == num_tokens,
-                  "required output.shape[0] == mat_a.shape[0]");
-  STD_TORCH_CHECK(output.size(1) == hd_out,
-                  "required output.shape[1] == mat_b.shape[1]");
+  TORCH_CHECK(num_tokens >= 1 && num_tokens <= 16,
+              "required 1 <= mat_a.shape[0] <= 16")
+  TORCH_CHECK(hd_in == kHdIn, "required mat_a.shape[1] == 7168")
+  TORCH_CHECK(hd_out == kHdOut, "required mat_b.shape[1] == 2112")
+  TORCH_CHECK(output.size(0) == num_tokens,
+              "required output.shape[0] == mat_a.shape[0]")
+  TORCH_CHECK(output.size(1) == hd_out,
+              "required output.shape[1] == mat_b.shape[1]")

-  STD_TORCH_CHECK(mat_a.stride(1) == 1, "mat_a must be a row major tensor");
-  STD_TORCH_CHECK(output.stride(1) == 1, "output must be a row major tensor");
-  STD_TORCH_CHECK(mat_b.stride(0) == 1, "mat_b must be a column major tensor");
+  TORCH_CHECK(mat_a.stride(1) == 1, "mat_a must be a row major tensor");
+  TORCH_CHECK(output.stride(1) == 1, "output must be a row major tensor");
+  TORCH_CHECK(mat_b.stride(0) == 1, "mat_b must be a column major tensor");

-  STD_TORCH_CHECK(
-      mat_a.scalar_type() == torch::headeronly::ScalarType::BFloat16 &&
-          mat_b.scalar_type() == torch::headeronly::ScalarType::BFloat16,
-      "Only BFloat16 input dtype is supported");
-  STD_TORCH_CHECK(
-      output.scalar_type() == torch::headeronly::ScalarType::BFloat16,
-      "Only BFloat16 output dtype is supported");
+  TORCH_CHECK(mat_a.scalar_type() == torch::kBFloat16 &&
+                  mat_b.scalar_type() == torch::kBFloat16,
+              "Only BFloat16 input dtype is supported")
+  TORCH_CHECK(output.scalar_type() == torch::kBFloat16,
+              "Only BFloat16 output dtype is supported")

-  STD_TORCH_CHECK(getSMVersion() >= 90, "required CUDA ARCH >= SM_90");
+  TORCH_CHECK(getSMVersion() >= 90, "required CUDA ARCH >= SM_90");

-  auto stream = get_current_cuda_stream(mat_a.get_device_index());
+  auto stream = at::cuda::getCurrentCUDAStream(mat_a.get_device());
  if (num_tokens <= 8) {
    invokeFusedAGemm<__nv_bfloat16, kHdIn, kHdOut, 8>(
        reinterpret_cast<__nv_bfloat16*>(output.mutable_data_ptr()),
@@ -751,6 +746,6 @@ void dsv3_fused_a_gemm(torch::stable::Tensor& output,
  }
 }

-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("dsv3_fused_a_gemm", TORCH_BOX(&dsv3_fused_a_gemm));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+  m.impl("dsv3_fused_a_gemm", &dsv3_fused_a_gemm);
 }
@@ -143,26 +143,4 @@ void cutlass_mxfp4_group_mm(torch::stable::Tensor& output,
                            const torch::stable::Tensor& expert_offsets,
                            const torch::stable::Tensor& sf_offsets);

-// AWQ ops
-torch::stable::Tensor awq_gemm(torch::stable::Tensor _in_feats,
-                               torch::stable::Tensor _kernel,
-                               torch::stable::Tensor _scaling_factors,
-                               torch::stable::Tensor _zeros,
-                               int64_t split_k_iters);
-
-torch::stable::Tensor awq_dequantize(torch::stable::Tensor _kernel,
-                                     torch::stable::Tensor _scaling_factors,
-                                     torch::stable::Tensor _zeros,
-                                     int64_t split_k_iters, int64_t thx,
-                                     int64_t thy);
-
-// DSV3 fused A GEMM: conditionally compiled so declaration and impl
-// registration are in the source file (dsv3_fused_a_gemm.cu)
-
-// AllSpark ops: declarations are in the source files
-// (allspark_repack.cu and allspark_qgemm_w8a16.cu)
-
 #endif
-
-torch::stable::Tensor hadacore_transform(torch::stable::Tensor& x,
-                                         bool inplace);
@@ -1,6 +1,5 @@
 #include "scaled_mm_kernels.hpp"
 #include "scaled_mm_sm100_fp8_dispatch.cuh"
-#include "core/batch_invariant.hpp"

 namespace vllm {

@@ -14,17 +13,9 @@ 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,34 +294,6 @@ 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,
@@ -348,28 +320,4 @@ 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,6 +1,5 @@
 #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 {
@@ -15,17 +14,9 @@ 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,26 +179,6 @@ 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,
@@ -222,26 +202,4 @@ 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,6 +1,5 @@
 #include "scaled_mm_kernels.hpp"
 #include "scaled_mm_sm90_fp8_dispatch.cuh"
-#include "core/batch_invariant.hpp"

 namespace vllm {

@@ -14,17 +13,9 @@ 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,35 +349,6 @@ 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,
@@ -404,28 +375,4 @@ 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,7 +9,6 @@
 #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;
@@ -192,17 +191,9 @@ 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,43 +381,4 @@ 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
@@ -218,54 +218,7 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
  ops.def(
      "cutlass_encode_and_reorder_int4b_grouped(Tensor b_tensors) -> (Tensor, "
      "Tensor)");
-
-  // SM100 CUTLASS MLA decode
-  // conditionally compiled so impl registrations are in source file
-  ops.def(
-      "sm100_cutlass_mla_decode(Tensor! out, Tensor! lse, Tensor q_nope,"
-      "                         Tensor q_pe, Tensor kv_c_and_k_pe_cache,"
-      "                         Tensor seq_lens, Tensor page_table,"
-      "                         Tensor workspace, float scale,"
-      "                         int num_kv_splits) -> ()");
-
-  ops.def(
-      "sm100_cutlass_mla_get_workspace_size(int max_seq_len, int num_batches,"
-      "                                     int sm_count, int num_kv_splits) "
-      "-> int");
-  // Quantized GEMM for AWQ.
-  ops.def(
-      "awq_gemm(Tensor _in_feats, Tensor _kernel, Tensor _scaling_factors, "
-      "Tensor _zeros, SymInt split_k_iters) -> Tensor");
-
-  // Dequantization for AWQ.
-  ops.def(
-      "awq_dequantize(Tensor _kernel, Tensor _scaling_factors, "
-      "Tensor _zeros, SymInt split_k_iters, int thx, int thy) -> Tensor");
-
-  // DeepSeek V3 fused A GEMM (SM 9.0+, bf16 only, 1-16 tokens).
-  // conditionally compiled so impl registration is in source file
-  ops.def(
-      "dsv3_fused_a_gemm(Tensor! output, Tensor mat_a, Tensor mat_b) -> ()");
-
-  // reorder weight for AllSpark Ampere W8A16 Fused Gemm kernel
-  ops.def(
-      "rearrange_kn_weight_as_n32k16_order(Tensor b_qweight, Tensor b_scales, "
-      "Tensor? b_zeros, "
-      "bool has_zp, Tensor! b_qweight_reorder, Tensor! b_scales_reorder, "
-      "Tensor!? b_zeros_reorder, "
-      "int K, int N, int N_32align) -> ()");
-
-  // AllSpark quantization ops
-  ops.def(
-      "allspark_w8a16_gemm(Tensor a, Tensor b_qweight, Tensor b_scales, "
-      "Tensor? b_qzeros, "
-      "SymInt n, SymInt group_size, SymInt sm_count, SymInt sm_version, SymInt "
-      "CUBLAS_M_THRESHOLD, bool has_zp, bool n32k16_reorder) -> Tensor");
 #endif
-
-  // Hadamard transforms
-  // conditionally compiled so impl registration is in source file
-  ops.def("hadacore_transform(Tensor! x, bool inplace) -> Tensor");
 }

 STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
@@ -301,16 +254,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
  ops.impl("silu_and_mul_nvfp4_quant", TORCH_BOX(&silu_and_mul_nvfp4_quant));
  // mxfp4_experts_quant: registered in mxfp4_experts_quant.cu (SM100 only).
  // W4A8 ops: registered in w4a8_mm_entry.cu / w4a8_grouped_mm_entry.cu.
-
-  // AWQ ops
-  ops.impl("awq_gemm", TORCH_BOX(&awq_gemm));
-  ops.impl("awq_dequantize", TORCH_BOX(&awq_dequantize));
-
-  // DSV3 fused A GEMM: conditionally compiled so impl registration is in
-  // source file (dsv3_fused_a_gemm.cu)
-
-  // AllSpark ops: conditionally compiled so impl registrations are in source
-  // files (allspark_repack.cu and allspark_qgemm_w8a16.cu)
 #endif
 }

@@ -6,71 +6,12 @@
 #include <torch/csrc/stable/tensor.h>
 #include <torch/headeronly/util/shim_utils.h>

-#include <cublas_v2.h>
 #include <cuda_runtime.h>

-#include <deque>
-#include <mutex>
-#include <string>
-#include <vector>
-
 // Stable ABI equivalent of TORCH_CHECK_NOT_IMPLEMENTED.
 #define STD_TORCH_CHECK_NOT_IMPLEMENTED(cond, ...) \
  STD_TORCH_CHECK(cond, "NotImplementedError: ", __VA_ARGS__)

-// Device properties cache for stable ABI compatibility.
-// Uses raw CUDA/HIP APIs instead of ATen functions.
-// Using inline ensures a single instance across all translation units.
-inline std::deque<std::once_flag> device_flags;
-inline std::vector<cudaDeviceProp> device_properties;
-inline std::once_flag vectors_init_flag;
-
-inline void do_init_device_vectors() {
-  int device_count;
-  cudaError_t err = cudaGetDeviceCount(&device_count);
-  if (err != cudaSuccess) {
-    STD_TORCH_CHECK(false, "cudaGetDeviceCount failed: " +
-                               std::string(cudaGetErrorString(err)));
-  }
-  device_flags.resize(device_count);
-  device_properties.resize(device_count);
-}
-
-inline void initDeviceVectors() {
-  std::call_once(vectors_init_flag, do_init_device_vectors);
-}
-
-inline void initDeviceProperty(int device_index) {
-  cudaDeviceProp device_prop{};
-  cudaError_t err = cudaGetDeviceProperties(&device_prop, device_index);
-  if (err != cudaSuccess) {
-    STD_TORCH_CHECK(false, "cudaGetDeviceProperties failed: " +
-                               std::string(cudaGetErrorString(err)));
-  }
-  device_properties[device_index] = device_prop;
-}
-
-// Get device properties using raw CUDA/HIP APIs (stable ABI compatible).
-// Caches results per device so cudaGetDeviceProperties is called at most once
-// per device.
-inline cudaDeviceProp* get_device_prop() {
-  initDeviceVectors();
-  int device_index;
-  cudaError_t err = cudaGetDevice(&device_index);
-  if (err != cudaSuccess) {
-    STD_TORCH_CHECK(
-        false, "cudaGetDevice failed: " + std::string(cudaGetErrorString(err)));
-  }
-  STD_TORCH_CHECK(device_index >= 0 && static_cast<size_t>(device_index) <
-                                           device_properties.size(),
-                  "CUDA device index " + std::to_string(device_index) +
-                      " out of range [0, " +
-                      std::to_string(device_properties.size()) + ")");
-
-  std::call_once(device_flags[device_index], initDeviceProperty, device_index);
-  return &device_properties[device_index];
-}
-
 // Utility to get the current CUDA stream for a given device using stable APIs.
 // Returns a cudaStream_t for use in kernel launches.
 inline cudaStream_t get_current_cuda_stream(int32_t device_index = -1) {
@@ -79,10 +20,3 @@ inline cudaStream_t get_current_cuda_stream(int32_t device_index = -1) {
      aoti_torch_get_current_cuda_stream(device_index, &stream_ptr));
  return reinterpret_cast<cudaStream_t>(stream_ptr);
 }
-
-// Utility to get the current cuBLAS handle using stable APIs.
-inline cublasHandle_t get_current_cuda_blas_handle() {
-  void* blas_handle_ptr = nullptr;
-  TORCH_ERROR_CODE_CHECK(torch_get_current_cuda_blas_handle(&blas_handle_ptr));
-  return reinterpret_cast<cublasHandle_t>(blas_handle_ptr);
-}
@@ -200,6 +200,19 @@ void cutlass_mla_decode(torch::Tensor const& out, torch::Tensor const& q_nope,

 torch::Tensor get_cuda_view_from_cpu_tensor(torch::Tensor& cpu_tensor);

+#ifndef USE_ROCM
+
+torch::Tensor awq_gemm(torch::Tensor _in_feats, torch::Tensor _kernel,
+                       torch::Tensor _scaling_factors, torch::Tensor _zeros,
+                       int64_t split_k_iters);
+
+torch::Tensor awq_dequantize(torch::Tensor _kernel,
+                             torch::Tensor _scaling_factors,
+                             torch::Tensor _zeros, int64_t split_k_iters,
+                             int64_t thx, int64_t thy);
+
+#endif
+
 torch::Tensor ggml_dequantize(torch::Tensor W, int64_t type, int64_t m,
                              int64_t n,
                              std::optional<at::ScalarType> const& dtype);
@@ -289,6 +302,8 @@ std::tuple<int64_t, torch::Tensor> allocate_shared_buffer_and_handle(
 int64_t open_mem_handle(torch::Tensor& mem_handle);
 void free_shared_buffer(int64_t buffer);

+torch::Tensor hadacore_transform(torch::Tensor& x, bool inplace);
+
 #ifdef USE_ROCM
 fptr_t init_custom_qr(int64_t rank, int64_t world_size,
                      std::optional<int64_t> qr_max_size = std::nullopt);
@@ -300,6 +315,11 @@ void qr_all_reduce(fptr_t _fa, torch::Tensor& inp, torch::Tensor& out,
 int64_t qr_max_size();
 #endif

+#ifndef USE_ROCM
+void dsv3_fused_a_gemm(torch::Tensor& output, torch::Tensor const& mat_a,
+                       torch::Tensor const& mat_b);
+#endif
+
 #ifndef USE_ROCM
 torch::Tensor minimax_allreduce_rms(torch::Tensor const& input,
                                    torch::Tensor const& norm_weight,
@@ -7,11 +7,10 @@ Shang and Dang, Xingyu and Han, Song}, journal={arXiv}, year={2023}
 }
 */

-#include <torch/csrc/stable/tensor.h>
-#include <torch/csrc/stable/ops.h>
-#include "libtorch_stable/torch_utils.h"
+#include <torch/all.h>
+#include <c10/cuda/CUDAGuard.h>

-#include "libtorch_stable/quantization/awq/dequantize.cuh"
+#include "dequantize.cuh"

 #include <cuda_fp16.h>

@@ -411,11 +410,10 @@ __global__ void __launch_bounds__(64)
 }  // namespace awq
 }  // namespace vllm

-torch::stable::Tensor awq_dequantize(torch::stable::Tensor _kernel,
-                                     torch::stable::Tensor _scaling_factors,
-                                     torch::stable::Tensor _zeros,
-                                     int64_t split_k_iters, int64_t thx,
-                                     int64_t thy) {
+torch::Tensor awq_dequantize(torch::Tensor _kernel,
+                             torch::Tensor _scaling_factors,
+                             torch::Tensor _zeros, int64_t split_k_iters,
+                             int64_t thx, int64_t thy) {
  int in_c = _kernel.size(0);
  int qout_c = _kernel.size(1);
  int out_c = qout_c * 8;
@@ -439,24 +437,23 @@ torch::stable::Tensor awq_dequantize(torch::stable::Tensor _kernel,
    y_blocks = (int)(in_c / 8);
  }

-  const torch::stable::accelerator::DeviceGuard device_guard(
-      _scaling_factors.get_device_index());
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(_scaling_factors));

-  auto _de_kernel =
-      torch::stable::empty({in_c, out_c}, _scaling_factors.scalar_type(),
-                           std::nullopt, _scaling_factors.device());
+  auto options = torch::TensorOptions()
+                     .dtype(_scaling_factors.dtype())
+                     .device(_scaling_factors.device());
+  at::Tensor _de_kernel = torch::empty({in_c, out_c}, options);

-  auto kernel = reinterpret_cast<int*>(_kernel.mutable_data_ptr<int>());
-  auto de_kernel = reinterpret_cast<half*>(
-      _de_kernel.mutable_data_ptr<torch::headeronly::Half>());
-  auto scaling_factors = reinterpret_cast<half*>(
-      _scaling_factors.mutable_data_ptr<torch::headeronly::Half>());
-  auto zeros = reinterpret_cast<int*>(_zeros.mutable_data_ptr<int>());
+  auto kernel = reinterpret_cast<int*>(_kernel.data_ptr<int>());
+  auto de_kernel = reinterpret_cast<half*>(_de_kernel.data_ptr<at::Half>());
+  auto scaling_factors =
+      reinterpret_cast<half*>(_scaling_factors.data_ptr<at::Half>());
+  auto zeros = reinterpret_cast<int*>(_zeros.data_ptr<int>());

  dim3 num_blocks(x_blocks, y_blocks);
  dim3 threads_per_block(x_thread, y_thread);

-  const cudaStream_t stream = get_current_cuda_stream();
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  vllm::awq::dequantize_weights<<<num_blocks, threads_per_block, 0, stream>>>(
      kernel, scaling_factors, zeros, de_kernel, G);

@@ -469,30 +466,27 @@ torch::stable::Tensor awq_dequantize(torch::stable::Tensor _kernel,
 // zeros: IC // G, OC // 8 [int32] -> cast to IC // G, OC [uint4b]
 // assume that batch_size < 16 for now

-torch::stable::Tensor awq_gemm(torch::stable::Tensor _in_feats,
-                               torch::stable::Tensor _kernel,
-                               torch::stable::Tensor _scaling_factors,
-                               torch::stable::Tensor _zeros,
-                               int64_t split_k_iters) {
+torch::Tensor awq_gemm(torch::Tensor _in_feats, torch::Tensor _kernel,
+                       torch::Tensor _scaling_factors, torch::Tensor _zeros,
+                       int64_t split_k_iters) {
  int num_in_feats = _in_feats.size(0);
  int num_in_channels = _in_feats.size(1);
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      _in_feats.get_device_index());
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(_in_feats));

-  auto _out_feats = torch::stable::empty(
-      {split_k_iters, num_in_feats, _kernel.size(1) * 8},
-      _in_feats.scalar_type(), std::nullopt, _in_feats.device());
+  auto options = torch::TensorOptions()
+                     .dtype(_in_feats.dtype())
+                     .device(_in_feats.device());
+  at::Tensor _out_feats =
+      torch::empty({split_k_iters, num_in_feats, _kernel.size(1) * 8}, options);
  int num_out_feats = _out_feats.size(-2);
  int num_out_channels = _out_feats.size(-1);

-  auto in_feats = reinterpret_cast<half*>(
-      _in_feats.mutable_data_ptr<torch::headeronly::Half>());
-  auto kernel = reinterpret_cast<int*>(_kernel.mutable_data_ptr<int>());
-  auto out_feats = reinterpret_cast<half*>(
-      _out_feats.mutable_data_ptr<torch::headeronly::Half>());
-  auto scaling_factors = reinterpret_cast<half*>(
-      _scaling_factors.mutable_data_ptr<torch::headeronly::Half>());
-  auto zeros = reinterpret_cast<int*>(_zeros.mutable_data_ptr<int>());
+  auto in_feats = reinterpret_cast<half*>(_in_feats.data_ptr<at::Half>());
+  auto kernel = reinterpret_cast<int*>(_kernel.data_ptr<int>());
+  auto out_feats = reinterpret_cast<half*>(_out_feats.data_ptr<at::Half>());
+  auto scaling_factors =
+      reinterpret_cast<half*>(_scaling_factors.data_ptr<at::Half>());
+  auto zeros = reinterpret_cast<int*>(_zeros.data_ptr<int>());
  int group_size = num_in_channels / _scaling_factors.size(0);

  if (num_out_channels % 64 != 0)
@@ -504,7 +498,7 @@ torch::stable::Tensor awq_gemm(torch::stable::Tensor _in_feats,
  if (num_out_channels % group_size != 0)
    throw std::invalid_argument("OC is not multiple of Group size");

-  const cudaStream_t stream = get_current_cuda_stream();
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  if (num_out_channels % 128 == 0) {
    int j_factors1 = num_out_channels / 128 / 1;
    dim3 num_blocks((num_out_feats + 16 - 1) / 16 * j_factors1 * split_k_iters);
@@ -528,5 +522,5 @@ torch::stable::Tensor awq_gemm(torch::stable::Tensor _in_feats,
            group_size, split_k_iters, in_feats, kernel, scaling_factors, zeros,
            num_in_feats, num_in_channels, num_out_channels, out_feats);
  }
-  return torch::stable::sum(_out_feats, 0);
+  return _out_feats.sum(0);
 }
@@ -1,28 +1,20 @@
 #include "allspark_utils.cuh"
-
-#include <torch/csrc/stable/c/shim.h>
-#include <torch/csrc/stable/library.h>
-#include <torch/csrc/stable/tensor.h>
-#include <torch/headeronly/core/ScalarType.h>
-
+#include <torch/all.h>
+#include "core/registration.h"
 #include <cublas_v2.h>

-#include "core/registration.h"
-#include "libtorch_stable/torch_utils.h"
-
-torch::stable::Tensor as_g_workspace;
+at::Tensor as_g_workspace;

 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 800

-torch::stable::Tensor allspark_w8a16_gemm(
-    torch::stable::Tensor const& a, torch::stable::Tensor const& b_qweight,
-    torch::stable::Tensor const& b_scales,
-    std::optional<torch::stable::Tensor> const& b_qzeros, int64_t n,
-    int64_t group_size, int64_t sm_count, int64_t sm_version,
+torch::Tensor allspark_w8a16_gemm(
+    torch::Tensor const& a, torch::Tensor const& b_qweight,
+    torch::Tensor const& b_scales, std::optional<torch::Tensor> const& b_qzeros,
+    int64_t n, int64_t group_size, int64_t sm_count, int64_t sm_version,
    int64_t CUBLAS_M_THRESHOLD, bool has_zp, bool n32k16_reorder) {
-  STD_TORCH_CHECK_NOT_IMPLEMENTED(
+  TORCH_CHECK_NOT_IMPLEMENTED(
      false, "allspark_w8a16_gemm(..) requires CUDA_ARCH >= 8.0");
-  return torch::stable::empty({1, 1});
+  return torch::empty({1, 1});
 }

 #else
@@ -856,8 +848,8 @@ void restore_N32_K16_dequantize_rhs_w8a16(const QT* qdata, const FT* scales,
                                          const int N_32align, const int N,
                                          const int K, const int GroupSize,
                                          cudaStream_t stream) {
-  STD_TORCH_CHECK(N % 8 == 0 && K % 16 == 0 && N_32align % 32 == 0,
-                  "Unsupported shape");
+  TORCH_CHECK(N % 8 == 0 && K % 16 == 0 && N_32align % 32 == 0,
+              "Unsupported shape");
  if (GroupSize == -1) {
    const int BLOCK = 128;
    dim3 grid(N_32align / 32, ((K / 16) + 3) / 4);
@@ -867,7 +859,7 @@ void restore_N32_K16_dequantize_rhs_w8a16(const QT* qdata, const FT* scales,
  }
  // TODO: Support SubChannel
  else {
-    STD_TORCH_CHECK(false, "Now only support PerChannel");
+    TORCH_CHECK(false, "Now only support PerChannel");
  }
 }

@@ -924,27 +916,24 @@ void allspark_qgemm_w8a16_perc_ampere(

 }  // namespace allspark

-torch::stable::Tensor allspark_w8a16_gemm(
-    torch::stable::Tensor const& a, torch::stable::Tensor const& b_qweight,
-    torch::stable::Tensor const& b_scales,
-    std::optional<torch::stable::Tensor> const& b_qzeros, int64_t n,
-    int64_t group_size, int64_t sm_count, int64_t sm_version,
+torch::Tensor allspark_w8a16_gemm(
+    torch::Tensor const& a, torch::Tensor const& b_qweight,
+    torch::Tensor const& b_scales, std::optional<torch::Tensor> const& b_qzeros,
+    int64_t n, int64_t group_size, int64_t sm_count, int64_t sm_version,
    int64_t CUBLAS_M_THRESHOLD, bool has_zp, bool n32k16_reorder) {
  // Verify device and strides
-  STD_TORCH_CHECK(a.device().is_cuda(), "A is not on GPU");
-  STD_TORCH_CHECK(a.is_contiguous(), "A is not contiguous");
+  TORCH_CHECK(a.device().is_cuda(), "A is not on GPU");
+  TORCH_CHECK(a.is_contiguous(), "A is not contiguous");

-  STD_TORCH_CHECK(b_qweight.device().is_cuda(), "b_qweight is not on GPU");
-  STD_TORCH_CHECK(b_qweight.is_contiguous(), "b_qweight is not contiguous");
+  TORCH_CHECK(b_qweight.device().is_cuda(), "b_qweight is not on GPU");
+  TORCH_CHECK(b_qweight.is_contiguous(), "b_qweight is not contiguous");

-  STD_TORCH_CHECK(b_scales.device().is_cuda(), "b_scales is not on GPU");
-  STD_TORCH_CHECK(b_scales.is_contiguous(), "b_scales is not contiguous");
+  TORCH_CHECK(b_scales.device().is_cuda(), "b_scales is not on GPU");
+  TORCH_CHECK(b_scales.is_contiguous(), "b_scales is not contiguous");

  if (has_zp) {
-    STD_TORCH_CHECK(b_qzeros.value().device().is_cuda(),
-                    "b_qzeros is not on GPU");
-    STD_TORCH_CHECK(b_qzeros.value().is_contiguous(),
-                    "b_qzeros is not contiguous");
+    TORCH_CHECK(b_qzeros.value().device().is_cuda(), "b_qzeros is not on GPU");
+    TORCH_CHECK(b_qzeros.value().is_contiguous(), "b_qzeros is not contiguous");
  }

  int m = a.size(0);
@@ -952,17 +941,16 @@ torch::stable::Tensor allspark_w8a16_gemm(
  int k = a.size(1);

  // Verify shape
-  STD_TORCH_CHECK(b_qweight.size(0) == n_32align,
-                  "Shape mismatch: b_qweight.size(0) = ", b_qweight.size(0),
-                  ", n_32align = ", n_32align);
-  STD_TORCH_CHECK(b_qweight.size(1) == k,
-                  "Shape mismatch: b_qweight.size(1) = ", b_qweight.size(1),
-                  ", k = ", k);
+  TORCH_CHECK(b_qweight.size(0) == n_32align,
+              "Shape mismatch: b_qweight.size(0) = ", b_qweight.size(0),
+              ", n_32align = ", n_32align);
+  TORCH_CHECK(b_qweight.size(1) == k,
+              "Shape mismatch: b_qweight.size(1) = ", b_qweight.size(1),
+              ", k = ", k);

-  STD_TORCH_CHECK(group_size == -1, "Currently only supports group_size = -1");
+  TORCH_CHECK(group_size == -1, "Currently only supports group_size = -1");

-  const torch::stable::accelerator::DeviceGuard device_guard(
-      a.get_device_index());
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(a));
  const void* a_ptr = reinterpret_cast<const void*>(a.data_ptr());
  const uint8_t* b_ptr = reinterpret_cast<const uint8_t*>(b_qweight.data_ptr());
  const void* b_scale_ptr = reinterpret_cast<const void*>(b_scales.data_ptr());
@@ -971,12 +959,12 @@ torch::stable::Tensor allspark_w8a16_gemm(
    b_zero_ptr = reinterpret_cast<const void*>(b_qzeros.value().data_ptr());
  }

-  auto c =
-      torch::stable::empty({m, n}, a.scalar_type(), std::nullopt, a.device());
-  void* c_ptr = reinterpret_cast<void*>(c.mutable_data_ptr());
+  auto c_options = torch::TensorOptions().dtype(a.dtype()).device(a.device());
+  torch::Tensor c = torch::empty({m, n}, c_options);
+  void* c_ptr = reinterpret_cast<void*>(c.data_ptr());

-  cudaStream_t stream = get_current_cuda_stream();
-  cublasHandle_t handle = get_current_cuda_blas_handle();
+  cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+  cublasHandle_t handle = at::cuda::getCurrentCUDABlasHandle();

  allspark::BlockTileSplitkParams fused_gemm_params;

@@ -988,15 +976,14 @@ torch::stable::Tensor allspark_w8a16_gemm(
        m, n, k, sm_count, fused_gemm_params);
  }

+  auto ws_options = torch::TensorOptions().dtype(at::kChar).device(a.device());
  if (as_g_workspace.numel() <
      ws_size) {  // ws_options: kChar, so numel() is bytes
-    as_g_workspace = torch::stable::empty({static_cast<int64_t>(ws_size)},
-                                          torch::headeronly::ScalarType::Char,
-                                          std::nullopt, a.device());
+    as_g_workspace = torch::empty({long(ws_size)}, ws_options);
  }
  void* ws = reinterpret_cast<void*>(as_g_workspace.data_ptr());

-  if (a.scalar_type() == torch::headeronly::ScalarType::Half) {
+  if (a.dtype() == at::ScalarType::Half) {
    allspark::allspark_qgemm_w8a16_perc_ampere<__half, uint8_t>(
        reinterpret_cast<const __half*>(a_ptr), b_ptr,
        reinterpret_cast<const __half*>(b_scale_ptr),
@@ -1004,7 +991,7 @@ torch::stable::Tensor allspark_w8a16_gemm(
        reinterpret_cast<__half*>(c_ptr), m, n_32align, n, k, ws,
        fused_gemm_params, group_size, CUBLAS_M_THRESHOLD, sm_version, stream,
        handle);
-  } else if (a.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
+  } else if (a.dtype() == at::ScalarType::BFloat16) {
    allspark::allspark_qgemm_w8a16_perc_ampere<__nv_bfloat16, uint8_t>(
        reinterpret_cast<const __nv_bfloat16*>(a_ptr), b_ptr,
        reinterpret_cast<const __nv_bfloat16*>(b_scale_ptr),
@@ -1019,6 +1006,6 @@ torch::stable::Tensor allspark_w8a16_gemm(

 #endif

-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("allspark_w8a16_gemm", TORCH_BOX(&allspark_w8a16_gemm));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+  m.impl("allspark_w8a16_gemm", &allspark_w8a16_gemm);
 }
@@ -1,11 +1,6 @@
 #include "allspark_utils.cuh"
-
-#include <torch/csrc/stable/library.h>
-#include <torch/csrc/stable/tensor.h>
-#include <torch/headeronly/core/ScalarType.h>
-
+#include <torch/all.h>
 #include "core/registration.h"
-#include "libtorch_stable/torch_utils.h"

 namespace allspark {

@@ -104,40 +99,36 @@ void rearrange_kn_weight_as_n32k16_order_ldg16(
 }  // namespace allspark

 void rearrange_kn_weight_as_n32k16_order(
-    torch::stable::Tensor const& b_qweight,
-    torch::stable::Tensor const& b_scales,
-    std::optional<torch::stable::Tensor> const& b_zeros, bool has_zp,
-    torch::stable::Tensor& b_qweight_reorder,
-    torch::stable::Tensor& b_scales_reorder,
-    std::optional<torch::stable::Tensor> const& b_zeros_reorder,
-    const int64_t K, const int64_t N, const int64_t N_32align) {
+    torch::Tensor const& b_qweight, torch::Tensor const& b_scales,
+    std::optional<torch::Tensor> const& b_zeros, bool has_zp,
+    torch::Tensor& b_qweight_reorder, torch::Tensor& b_scales_reorder,
+    std::optional<torch::Tensor> const& b_zeros_reorder, const int64_t K,
+    const int64_t N, const int64_t N_32align) {
  // Verify device and strides
-  STD_TORCH_CHECK(b_qweight.device().is_cuda(), "b_qweight is not on GPU");
-  STD_TORCH_CHECK(b_qweight.is_contiguous(), "b_qweight is not contiguous");
+  TORCH_CHECK(b_qweight.device().is_cuda(), "b_qweight is not on GPU");
+  TORCH_CHECK(b_qweight.is_contiguous(), "b_qweight is not contiguous");

-  STD_TORCH_CHECK(b_scales.device().is_cuda(), "b_scales is not on GPU");
-  STD_TORCH_CHECK(b_scales.is_contiguous(), "b_scales is not contiguous");
+  TORCH_CHECK(b_scales.device().is_cuda(), "b_scales is not on GPU");
+  TORCH_CHECK(b_scales.is_contiguous(), "b_scales is not contiguous");

-  STD_TORCH_CHECK(b_qweight_reorder.device().is_cuda(),
-                  "b_qweight_reorder is not on GPU");
-  STD_TORCH_CHECK(b_qweight_reorder.is_contiguous(),
-                  "b_qweight_reorder is not contiguous");
+  TORCH_CHECK(b_qweight_reorder.device().is_cuda(),
+              "b_qweight_reorder is not on GPU");
+  TORCH_CHECK(b_qweight_reorder.is_contiguous(),
+              "b_qweight_reorder is not contiguous");

-  STD_TORCH_CHECK(b_scales_reorder.device().is_cuda(),
-                  "b_scales_reorder is not on GPU");
-  STD_TORCH_CHECK(b_scales_reorder.is_contiguous(),
-                  "b_scales_reorder is not contiguous");
+  TORCH_CHECK(b_scales_reorder.device().is_cuda(),
+              "b_scales_reorder is not on GPU");
+  TORCH_CHECK(b_scales_reorder.is_contiguous(),
+              "b_scales_reorder is not contiguous");

  if (has_zp) {
-    STD_TORCH_CHECK(b_zeros.value().device().is_cuda(),
-                    "b_zeros is not on GPU");
-    STD_TORCH_CHECK(b_zeros.value().is_contiguous(),
-                    "b_zeros is not contiguous");
+    TORCH_CHECK(b_zeros.value().device().is_cuda(), "b_zeros is not on GPU");
+    TORCH_CHECK(b_zeros.value().is_contiguous(), "b_zeros is not contiguous");

-    STD_TORCH_CHECK(b_zeros_reorder.value().device().is_cuda(),
-                    "b_zeros_reorder is not on GPU");
-    STD_TORCH_CHECK(b_zeros_reorder.value().is_contiguous(),
-                    "b_zeros_reorder is not contiguous");
+    TORCH_CHECK(b_zeros_reorder.value().device().is_cuda(),
+                "b_zeros_reorder is not on GPU");
+    TORCH_CHECK(b_zeros_reorder.value().is_contiguous(),
+                "b_zeros_reorder is not contiguous");
  }

  const uint8_t* matB = reinterpret_cast<const uint8_t*>(b_qweight.data_ptr());
@@ -145,20 +136,18 @@ void rearrange_kn_weight_as_n32k16_order(
  const void* b_zero = has_zp ? b_zeros.value().data_ptr() : nullptr;

  uint8_t* matB_reorder =
-      reinterpret_cast<uint8_t*>(b_qweight_reorder.mutable_data_ptr());
-  void* b_scale_reorder = b_scales_reorder.mutable_data_ptr();
-  void* b_zero_reorder =
-      has_zp ? b_zeros_reorder.value().mutable_data_ptr() : nullptr;
+      reinterpret_cast<uint8_t*>(b_qweight_reorder.data_ptr());
+  void* b_scale_reorder = b_scales_reorder.data_ptr();
+  void* b_zero_reorder = has_zp ? b_zeros_reorder.value().data_ptr() : nullptr;

-  cudaStream_t stream = get_current_cuda_stream();
-  if (b_scales.scalar_type() == torch::headeronly::ScalarType::Half) {
+  cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+  if (b_scales.dtype() == at::ScalarType::Half) {
    allspark::rearrange_kn_weight_as_n32k16_order_ldg16<__half>(
        matB, reinterpret_cast<const __half*>(b_scale),
        reinterpret_cast<const __half*>(b_zero), matB_reorder,
        reinterpret_cast<__half*>(b_scale_reorder),
        reinterpret_cast<__half*>(b_zero_reorder), K, N, N_32align, stream);
-  } else if (b_scales.scalar_type() ==
-             torch::headeronly::ScalarType::BFloat16) {
+  } else if (b_scales.dtype() == at::ScalarType::BFloat16) {
    allspark::rearrange_kn_weight_as_n32k16_order_ldg16<__nv_bfloat16>(
        matB, reinterpret_cast<const __nv_bfloat16*>(b_scale),
        reinterpret_cast<const __nv_bfloat16*>(b_zero), matB_reorder,
@@ -168,7 +157,7 @@ void rearrange_kn_weight_as_n32k16_order(
  }
 }

-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
  m.impl("rearrange_kn_weight_as_n32k16_order",
-         TORCH_BOX(&rearrange_kn_weight_as_n32k16_order));
+         &rearrange_kn_weight_as_n32k16_order);
 }
@@ -1,12 +1,13 @@
 #pragma once

-#include <cuda_bf16.h>
-#include <cuda_fp16.h>
+#include <torch/all.h>
+#include <c10/cuda/CUDAGuard.h>
+#include <ATen/cuda/CUDAContext.h>
 #include <cuda_runtime.h>
-
+#include <cuda_fp16.h>
+#include <cuda_bf16.h>
 #include <iostream>
-
-#include "quantization/marlin/marlin_dtypes.cuh"
+#include "../marlin/marlin_dtypes.cuh"
 using marlin::MarlinScalarType2;

 namespace allspark {
@@ -11,16 +11,18 @@ Redistribution and use in source and binary forms, with or without modification,
 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS “AS IS” AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 ***********/

-#include "libtorch_stable/torch_utils.h"
-#include "libtorch_stable/dispatch_utils.h"
-
-#include <torch/csrc/stable/library.h>
-#include <torch/csrc/stable/macros.h>
-
+#include <torch/all.h>
 #include <stdint.h>
 #include <cuda_runtime.h>
 #include <mma.h>
 #include <cuda/annotated_ptr>
+#include <c10/cuda/CUDAException.h>
+
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
+
+#include "core/registration.h"
+#include "dispatch_utils.h"

 namespace hadacore {

@@ -63,12 +65,12 @@ constexpr int launch_configs_big[7][3] = {
 };

 // a 4x2, b 2x2, c 2x2
-template <torch::headeronly::ScalarType dtype>
+template <torch::ScalarType dtype>
 __device__ __forceinline__ void mma_m16_n8_k16_b16_b16_b16_noacc(b32 a0, b32 a1, b32 a2, b32 a3, b32 b0, b32 b1, b32& c0, b32& c1){
-    static_assert(dtype == torch::headeronly::ScalarType::Half || dtype == torch::headeronly::ScalarType::BFloat16);
+    static_assert(dtype == torch::ScalarType::Half || dtype == torch::ScalarType::BFloat16);
    // d, a, b, c
    b32 zero = 0;
-    if constexpr(dtype == torch::headeronly::ScalarType::Half) {
+    if constexpr(dtype == torch::ScalarType::Half) {
        asm (
            "mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 "
            "{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\n\t"
@@ -87,7 +89,7 @@ __device__ __forceinline__ void mma_m16_n8_k16_b16_b16_b16_noacc(b32 a0, b32 a1,
 }

 // a 4x2, b 4x2, c 4x2
-template <torch::headeronly::ScalarType dtype>
+template <torch::ScalarType dtype>
 __device__ __forceinline__ void mma_m16_n16_k16_b16_b16_b16_noacc(b32 a0, b32 a1, b32 a2, b32 a3, b32 b0, b32 b1, b32 b2, b32 b3, b32& c0, b32& c1, b32& c2, b32& c3){
    mma_m16_n8_k16_b16_b16_b16_noacc<dtype>(a0, a1, a2, a3, b0, b1, c0, c1);
    mma_m16_n8_k16_b16_b16_b16_noacc<dtype>(a0, a1, a2, a3, b2, b3, c2, c3);
@@ -106,11 +108,11 @@ __device__ __forceinline__ void matrix_transpose_m8_n8_b16_inplace(b32& a0) {
 #define n_p(i) ((val_1n[i] & 0x0000FFFF) | val_1p[i] << 16)
 #define n_n(i) ((val_1n[i] & 0x0000FFFF) | val_1n[i] << 16)

-template<int64_t num_chunks, int64_t warps_per_block, int64_t log_had_size, int64_t blocks_per_sm, bool enable_mask, torch::headeronly::ScalarType dtype>
+template<int64_t num_chunks, int64_t warps_per_block, int64_t log_had_size, int64_t blocks_per_sm, bool enable_mask, torch::ScalarType dtype>
 __global__ void __launch_bounds__(32 * warps_per_block, blocks_per_sm)
 // a is column major, b is row major
 hadamard_transform_kernel(b16* a, b16* out, int total_num_chunks) {
-    static_assert(dtype == torch::headeronly::ScalarType::Half || dtype == torch::headeronly::ScalarType::BFloat16, "Only fp16 and bf16 supported currently");
+    static_assert(dtype == torch::ScalarType::Half || dtype == torch::ScalarType::BFloat16, "Only fp16 and bf16 supported currently");

    b32 b_frag_all[num_chunks][4]; // for all chunks, holds matrix fragment (which takes 4 regs of b16x2 * 32 threads)

@@ -160,8 +162,8 @@ hadamard_transform_kernel(b16* a, b16* out, int total_num_chunks) {
    constexpr b16 bf16_1p[4] = {0b0011111100110101, 0b0011111100000000, 0b0011111010110101, 0b0011111010000000};
    constexpr b16 bf16_1n[4] = {0b1011111100110101, 0b1011111100000000, 0b1011111010110101, 0b1011111010000000};

-    #define val_type_1p(i) (((dtype) == torch::headeronly::ScalarType::Half) ? (fp16_1p[i]) : (bf16_1p[i]))
-    #define val_type_1n(i) (((dtype) == torch::headeronly::ScalarType::Half) ? (fp16_1n[i]) : (bf16_1n[i]))
+    #define val_type_1p(i) (((dtype) == torch::ScalarType::Half) ? (fp16_1p[i]) : (bf16_1p[i]))
+    #define val_type_1n(i) (((dtype) == torch::ScalarType::Half) ? (fp16_1n[i]) : (bf16_1n[i]))
    constexpr b16 val_1p[4] = {val_type_1p(0), val_type_1p(1), val_type_1p(2), val_type_1p(3)};
    constexpr b16 val_1n[4] = {val_type_1n(0), val_type_1n(1), val_type_1n(2), val_type_1n(3)};

@@ -682,14 +684,14 @@ constexpr int64_t ceil_div(int64_t a, int64_t b) {
    return (a + b - 1) / b;
 }

-template <torch::headeronly::ScalarType dtype, int64_t chunks_per_warp, int64_t warps_per_block, int64_t log_had_size, int64_t blocks_per_sm, bool check_masking = false>
+template <torch::ScalarType dtype, int64_t chunks_per_warp, int64_t warps_per_block, int64_t log_had_size, int64_t blocks_per_sm, bool check_masking = false>
 void __forceinline__ run_kernel(b16* a_mat, b16* out, int64_t num_chunks, cudaStream_t stream) {
    int64_t shared_size = chunks_per_warp * warps_per_block * 128 * 4;
    dim3 block_size = 32 * warps_per_block;

    #define CHECK_SHARED_LIM() {                                                                              \
        if (shared_size > 48 * 1024) {                                                                        \
-            STD_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 65536)); \
+            C10_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 65536)); \
        }                                                                                                     \
    }                                                                                                         \

@@ -712,10 +714,10 @@ void __forceinline__ run_kernel(b16* a_mat, b16* out, int64_t num_chunks, cudaSt
        kernel<<<dim3(grid_size), dim3(block_size), shared_size, stream>>>(a_mat, out, num_chunks);
    }
    
-    STD_CUDA_KERNEL_LAUNCH_CHECK();
+    C10_CUDA_KERNEL_LAUNCH_CHECK();
 }

-template <torch::headeronly::ScalarType dtype>
+template <torch::ScalarType dtype>
 void run_fht(void* a_mat_ptr, void* out_ptr, int64_t numel, int64_t had_size, cudaStream_t stream) {
    int64_t num_chunks = numel / 256; // caller required to ensure divisible by 256
    // for size 256, use (2, 1)
@@ -762,54 +764,54 @@ void run_fht(void* a_mat_ptr, void* out_ptr, int64_t numel, int64_t had_size, cu
    }
 }

-template void run_fht<torch::headeronly::ScalarType::Half>(void* a_mat_ptr, void* out_ptr, int64_t numel, int64_t had_size, cudaStream_t stream);
-template void run_fht<torch::headeronly::ScalarType::BFloat16>(void* a_mat_ptr, void* out_ptr, int64_t numel, int64_t had_size, cudaStream_t stream);
+template void run_fht<torch::ScalarType::Half>(void* a_mat_ptr, void* out_ptr, int64_t numel, int64_t had_size, cudaStream_t stream);
+template void run_fht<torch::ScalarType::BFloat16>(void* a_mat_ptr, void* out_ptr, int64_t numel, int64_t had_size, cudaStream_t stream);

 }  // namespace hadacore

 constexpr bool is_power_of_two(int x) { return x && !(x & (x - 1)); }

-torch::stable::Tensor hadacore_transform(torch::stable::Tensor& x, bool inplace) {
+torch::Tensor hadacore_transform(torch::Tensor& x, bool inplace) {
    auto dtype = x.scalar_type();
-    STD_TORCH_CHECK(dtype == torch::headeronly::ScalarType::Half || dtype == torch::headeronly::ScalarType::BFloat16, "Only fp16 and bf16 supported currently");
-    STD_TORCH_CHECK(x.is_cuda());
-
+    TORCH_CHECK(dtype == torch::ScalarType::Half || dtype == torch::ScalarType::BFloat16, "Only fp16 and bf16 supported currently");
+    TORCH_CHECK(x.is_cuda());
+    
    const int had_size = x.size(-1);
-    STD_TORCH_CHECK(is_power_of_two(had_size) && (had_size <= (1U << 15)),
+    TORCH_CHECK(is_power_of_two(had_size) && (had_size <= (1U << 15)),
        "Only power of two Hadamard sizes up to 2^15 are supported, got ", had_size);
-
+    
    const auto res_shape = x.sizes();
-    x = torch::stable::reshape(x, {-1, had_size});
-
+    x = x.reshape({-1, had_size});
+    
    auto numel = x.numel();
    if (numel % 256 != 0) {
-        x = torch::stable::pad(x, {0, 0, 0, (256 - numel % 256) / had_size});
+        x = torch::nn::functional::pad(x, torch::nn::functional::PadFuncOptions({0, 0, 0, (256 - numel % 256) / had_size}));
    }
-
+    
    if (x.stride(-1) != 1) {
-        x = torch::stable::contiguous(x);
+        x = x.contiguous();
    }
-    torch::stable::Tensor out = inplace ? x : torch::stable::empty_like(x);
+    torch::Tensor out = inplace ? x : torch::empty_like(x);

-    torch::stable::accelerator::DeviceGuard device_guard(x.get_device_index());
-    auto stream = get_current_cuda_stream();
+    at::cuda::CUDAGuard device_guard{(char)x.get_device()};
+    auto stream = at::cuda::getCurrentCUDAStream().stream();

-    VLLM_STABLE_DISPATCH_HALF_TYPES(x.scalar_type(), "hadacore_transform_runfht", [&] {
-      auto constexpr SCALAR_TYPE = torch::headeronly::CppTypeToScalarType<scalar_t>::value;
+    VLLM_DISPATCH_HALF_TYPES(x.scalar_type(), "hadacore_transform_runfht", [&] {
+      auto constexpr SCALAR_TYPE = c10::CppTypeToScalarType<scalar_t>::value;
      hadacore::run_fht<SCALAR_TYPE>(x.data_ptr(), x.data_ptr(), x.numel(), had_size, stream);
    });

    if (numel % 256 != 0) {
-        out = torch::stable::narrow(out, 0, 0, numel / had_size);
+        out = out.narrow(0, 0, numel / had_size);
    }

    if (inplace && out.data_ptr() != x.data_ptr()) {
-        torch::stable::copy_(x, torch::stable::view(out, res_shape));
+        x.copy_(out.view(res_shape));
        return x;
    }
-    return torch::stable::reshape(out, res_shape);
+    return out.reshape(res_shape);
 }

-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-    m.impl("hadacore_transform", TORCH_BOX(&hadacore_transform));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+    m.impl("hadacore_transform", &hadacore_transform);
 }
@@ -2,14 +2,10 @@

 #ifndef _marlin_cuh
  #define _marlin_cuh
-  // These torch headers are only needed by non-stable callers (e.g. ops.cu).
-  // Guard them so that stable ABI targets can still include marlin.cuh
-  // for Vec, constants, and cp_async helpers without pulling in torch/all.h.
-  #ifndef TORCH_TARGET_VERSION
-    #include <torch/all.h>
-    #include <ATen/cuda/CUDAContext.h>
-    #include <c10/cuda/CUDAGuard.h>
-  #endif
+  #include <torch/all.h>
+
+  #include <ATen/cuda/CUDAContext.h>
+  #include <c10/cuda/CUDAGuard.h>
  #include <cuda.h>
  #include <cuda_fp16.h>
  #include <cuda_runtime.h>
@@ -1,204 +0,0 @@
-#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"
@@ -263,6 +263,22 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {

  // Quantization ops
 #ifndef USE_ROCM
+  // DeepSeek V3 fused A GEMM (SM 9.0+, bf16 only, 1-16 tokens).
+  ops.def(
+      "dsv3_fused_a_gemm(Tensor! output, Tensor mat_a, Tensor mat_b) -> ()");
+  // conditionally compiled so impl registration is in source file
+
+  // Quantized GEMM for AWQ.
+  ops.def(
+      "awq_gemm(Tensor _in_feats, Tensor _kernel, Tensor _scaling_factors, "
+      "Tensor _zeros, SymInt split_k_iters) -> Tensor");
+  ops.impl("awq_gemm", torch::kCUDA, &awq_gemm);
+
+  // Dequantization for AWQ.
+  ops.def(
+      "awq_dequantize(Tensor _kernel, Tensor _scaling_factors, "
+      "Tensor _zeros, SymInt split_k_iters, int thx, int thy) -> Tensor");
+  ops.impl("awq_dequantize", torch::kCUDA, &awq_dequantize);

  // Note about marlin kernel 'workspace' arguments:
  // Technically these should be mutable since they are modified by the kernel.
@@ -392,6 +408,22 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      " -> ()");
  // conditionally compiled so impl registration is in source file

+  // SM100 CUTLASS MLA decode
+  ops.def(
+      "sm100_cutlass_mla_decode(Tensor! out, Tensor! lse, Tensor q_nope,"
+      "                         Tensor q_pe, Tensor kv_c_and_k_pe_cache,"
+      "                         Tensor seq_lens, Tensor page_table,"
+      "                         Tensor workspace, float scale,"
+      "                         int num_kv_splits) -> ()");
+  // conditionally compiled so impl in source file
+
+  // SM100 CUTLASS MLA workspace
+  ops.def(
+      "sm100_cutlass_mla_get_workspace_size(int max_seq_len, int num_batches,"
+      "                                     int sm_count, int num_kv_splits) "
+      "-> int");
+  // conditionally compiled so impl in source file
+
 #endif

  // Quantized GEMM for GPTQ.
@@ -464,7 +496,26 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      "Tensor? last_chunk_indices) -> ()");
  ops.impl("selective_scan_fwd", torch::kCUDA, &selective_scan_fwd);

+  // Hadamard transforms
+  ops.def("hadacore_transform(Tensor! x, bool inplace) -> Tensor");
+
 #ifndef USE_ROCM
+  // reorder weight for AllSpark Ampere W8A16 Fused Gemm kernel
+  ops.def(
+      "rearrange_kn_weight_as_n32k16_order(Tensor b_qweight, Tensor b_scales, "
+      "Tensor? b_zeros, "
+      "bool has_zp, Tensor! b_qweight_reorder, Tensor! b_scales_reorder, "
+      "Tensor!? b_zeros_reorder, "
+      "int K, int N, int N_32align) -> ()");
+  //  conditionally compiled so impl in source file
+
+  // AllSpark quantization ops
+  ops.def(
+      "allspark_w8a16_gemm(Tensor a, Tensor b_qweight, Tensor b_scales, "
+      "Tensor? b_qzeros, "
+      "SymInt n, SymInt group_size, SymInt sm_count, SymInt sm_version, SymInt "
+      "CUBLAS_M_THRESHOLD, bool has_zp, bool n32k16_reorder) -> Tensor");
+
  ops.def(
      "minimax_allreduce_rms("
      "Tensor input,"
@@ -304,9 +304,10 @@ RUN --mount=type=cache,target=/root/.cache/uv \
        python3 use_existing_torch.py --prefix; \
    fi

-# Provision one bare Python per `requires-python` entry; cmake reads
-# DEEPGEMM_PYTHON_INTERPRETERS to build DeepGEMM `_C` for each. See
-# cmake/external_projects/deepgemm.cmake for the full picture.
+# Provision a bare interpreter for each CPython covered by `requires-python`
+# so DeepGEMM `_C` is built once per Python and bundled side-by-side in the
+# wheel; cmake reads DEEPGEMM_PYTHON_INTERPRETERS in deepgemm.cmake's
+# foreach loop. The matrix is derived from pyproject.toml.
 COPY tools/setup_deepgemm_pythons.sh tools/build_deepgemm_C.py tools/
 ENV DEEPGEMM_VENV_PREFIX=/opt/dgenv
 RUN --mount=type=cache,target=/root/.cache/uv \
@@ -845,9 +846,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
                libcublas-dev-${CUDA_VERSION_DASH} \
                libcusolver-dev-${CUDA_VERSION_DASH}"; \
    if [ "$INSTALL_KV_CONNECTORS" = "true" ]; then \
-        if [ "$CUDA_MAJOR" -ge 13 ]; then \
-            uv pip install --system nixl-cu13; \
-        fi; \
        uv pip install --system -r /tmp/kv_connectors.txt --no-build || ( \
            # if the above fails, install from source
            apt-get update -y && \
@@ -857,6 +855,8 @@ RUN --mount=type=cache,target=/root/.cache/uv \
            # clean up -dev packages, keep runtime libraries
            rm -rf /var/lib/apt/lists/* \
        ); \
+        # Force-reinstall the matching CUDA wheel so the correct nixl_ep_cpp.so is installed.
+        uv pip install --system --force-reinstall --no-deps nixl-cu${CUDA_MAJOR}; \
    fi

 ENV VLLM_USAGE_SOURCE production-docker-image
@@ -39,8 +39,7 @@ 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 Audio | ✅ | ✅ | Local file: `audio_data.jsonl` |
-| Custom Image | ✅ | ✅ | Local file: `image_data.jsonl` |
+| Custom MM | ✅ | ✅ | Local file: `mm_data.jsonl` |

 Legend:

@@ -143,7 +142,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`. 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
+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

 ```json
 {"prompt": "What is the capital of India?"}
@@ -174,62 +173,9 @@ 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 Audio Dataset
+#### Custom multimodal dataset

-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`:
+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`:

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

@@ -34,7 +34,6 @@ you may contact the following individuals:

 - Simon Mo - <simon.mo@hey.com>
 - Russell Bryant - <rbryant@redhat.com>
- Juan Pérez de Algaba - <jperezde@redhat.com>
 - Huzaifa Sidhpurwala - <huzaifas@redhat.com>

 ## Slack Discussion
@@ -2,4 +2,4 @@

 vLLM can be deployed with [llm-d](https://github.com/llm-d/llm-d), a Kubernetes-native distributed inference serving stack providing well-lit paths for anyone to serve large generative AI models at scale. It helps achieve the fastest "time to state-of-the-art (SOTA) performance" for key OSS models across most hardware accelerators and infrastructure providers.

-You can use vLLM with llm-d directly by following [the official guides](https://llm-d.ai/docs/guides) or via [KServe's LLMInferenceService](https://kserve.github.io/website/docs/model-serving/generative-inference/llmisvc/llmisvc-overview).
+You can use vLLM with llm-d directly by following [this guide](https://llm-d.ai/docs/guide) or via [KServe's LLMInferenceService](https://kserve.github.io/website/docs/model-serving/generative-inference/llmisvc/llmisvc-overview).
@@ -125,12 +125,13 @@ Priority is **1 = highest** (tried first).
 | Priority | Backend |
 | -------- | ------- |
 | 1 | `FLASHINFER_MLA` |
-| 2 | `CUTLASS_MLA` |
-| 3 | `FLASH_ATTN_MLA` |
-| 4 | `FLASHMLA` |
-| 5 | `TRITON_MLA` |
-| 6 | `FLASHINFER_MLA_SPARSE`**\*** |
-| 7 | `FLASHMLA_SPARSE` |
+| 2 | `TOKENSPEED_MLA` |
+| 3 | `CUTLASS_MLA` |
+| 4 | `FLASH_ATTN_MLA` |
+| 5 | `FLASHMLA` |
+| 6 | `TRITON_MLA` |
+| 7 | `FLASHINFER_MLA_SPARSE`**\*** |
+| 8 | `FLASHMLA_SPARSE` |

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

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

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

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

 !!! note
    Encoder CUDA Graphs have currently been tested with `--mm-encoder-attn-backend=FLASH_ATTN` and `--mm-encoder-attn-backend=FLASHINFER` on Blackwell GPUs.
@@ -81,7 +81,7 @@ 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.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.experts.fused_batched_moe.BatchedTritonExperts] |
+| 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] |
@@ -91,7 +91,7 @@ To be used with a particular `FusedMoEPrepareAndFinalizeModular` subclass, MoE k
 | trtllm | standard | mxfp4,</br>nvfp4 | G(16),G(32) | <sup>5</sup> | N | Y | [`TrtLlmMxfp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsMonolithic],</br>[`TrtLlmMxfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsModular],</br>[`TrtLlmNvFp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsMonolithic],</br>[`TrtLlmNvfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsModular] |
 | rocm aiter moe | standard | mxfp4,</br>fp8 | G(32),G(128),A,T | silu, gelu,</br>swigluoai | Y | N | `rocm_aiter_fused_experts`,</br>`AiterExperts` |
 | cpu_fused_moe | standard | N/A | N/A | silu | N | N | [`CPUFusedMOE`][vllm.model_executor.layers.fused_moe.cpu_fused_moe.CPUFusedMOE] |
-| naive batched<sup>4</sup> | batched | int8,</br>fp8 | G,A,T | silu, gelu | <sup>6</sup> | Y | [`NaiveBatchedExperts`][vllm.model_executor.layers.fused_moe.experts.fused_batched_moe.NaiveBatchedExperts] |
+| naive batched<sup>4</sup> | batched | int8,</br>fp8 | G,A,T | silu, gelu | <sup>6</sup> | Y | [`NaiveBatchedExperts`][vllm.model_executor.layers.fused_moe.fused_batched_moe.NaiveBatchedExperts] |

 !!! info "Table key"
    1. All types: mxfp4, nvfp4, int4, int8, fp8
@@ -126,15 +126,9 @@ 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

- `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)
+- `VLLM_NIXL_ABORT_REQUEST_TIMEOUT`: Timeout (in seconds) for automatically releasing the prefiller’s KV cache for a particular request. (Optional)
    - Default: 480
-    - 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}}'`
+    - 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.

 ## Multi-Instance Setup

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

 - [GPU](gpu.md)
-    - [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))
+    - [NVIDIA CUDA](gpu.md#nvidia-cuda)
+    - [AMD ROCm](gpu.md#amd-rocm)
+    - [Intel XPU](gpu.md#intel-xpu)
 - [CPU](cpu.md)
    - [Intel/AMD x86](cpu.md#intelamd-x86)
    - [ARM AArch64](cpu.md#arm-aarch64)
@@ -1,125 +0,0 @@
-<!-- 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,10 +18,6 @@ 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
@@ -42,10 +38,6 @@ 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
@@ -64,10 +56,6 @@ 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"
@@ -82,10 +70,6 @@ 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"
@@ -100,10 +84,6 @@ 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
@@ -122,10 +102,6 @@ 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
@@ -144,10 +120,6 @@ 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
@@ -163,7 +135,3 @@ 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,9 +10,6 @@ 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"
@@ -76,18 +73,6 @@ 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.

@@ -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`, `allenai/MolmoWeb-4B`<sup>^</sup>, `allenai/MolmoWeb-8B`<sup>^</sup> | ✅︎ | ✅︎ |
+| `Molmo2ForConditionalGeneration` | Molmo2 | T + I<sup>+</sup> / V | `allenai/Molmo2-4B`, `allenai/Molmo2-8B`, `allenai/Molmo2-O-7B` | ✅︎ | ✅︎ |
 | `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,11 +666,6 @@ 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,43 +333,6 @@ 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).
@@ -2179,92 +2179,6 @@ def run_qwen3_vl_moe(questions: list[str], modality: str) -> ModelRequestData:
    )


-# Qwen3.5-Dense
-def run_qwen3_5(questions: list[str], modality: str) -> ModelRequestData:
-    model_name = "Qwen/Qwen3.5-4B"
-
-    mm_limit = {"image": 1, "video": 1} if modality == "image+video" else {modality: 1}
-    engine_args = EngineArgs(
-        model=model_name,
-        max_model_len=4096,
-        max_num_seqs=5,
-        mm_processor_kwargs={
-            "min_pixels": 28 * 28,
-            "max_pixels": 1280 * 28 * 28,
-            "fps": 1,
-        },
-        limit_mm_per_prompt=mm_limit,
-    )
-
-    image_placeholder = "<|vision_start|><|image_pad|><|vision_end|>"
-    video_placeholder = "<|vision_start|><|video_pad|><|vision_end|>"
-
-    if modality == "image":
-        placeholder = image_placeholder
-    elif modality == "video":
-        placeholder = video_placeholder
-    elif modality == "image+video":
-        placeholder = image_placeholder + video_placeholder
-
-    prompts = [
-        (
-            "<|im_start|>system\nYou are a helpful assistant.<|im_end|>\n"
-            f"<|im_start|>user\n{placeholder}"
-            f"{question}<|im_end|>\n"
-            "<|im_start|>assistant\n"
-        )
-        for question in questions
-    ]
-
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompts=prompts,
-    )
-
-
-# Qwen3.5-MoE
-def run_qwen3_5_moe(questions: list[str], modality: str) -> ModelRequestData:
-    model_name = "Qwen/Qwen3.5-35B-A3B"
-
-    mm_limit = {"image": 1, "video": 1} if modality == "image+video" else {modality: 1}
-    engine_args = EngineArgs(
-        model=model_name,
-        max_model_len=4096,
-        max_num_seqs=5,
-        mm_processor_kwargs={
-            "min_pixels": 28 * 28,
-            "max_pixels": 1280 * 28 * 28,
-            "fps": 1,
-        },
-        limit_mm_per_prompt=mm_limit,
-    )
-
-    image_placeholder = "<|vision_start|><|image_pad|><|vision_end|>"
-    video_placeholder = "<|vision_start|><|video_pad|><|vision_end|>"
-
-    if modality == "image":
-        placeholder = image_placeholder
-    elif modality == "video":
-        placeholder = video_placeholder
-    elif modality == "image+video":
-        placeholder = image_placeholder + video_placeholder
-
-    prompts = [
-        (
-            "<|im_start|>system\nYou are a helpful assistant.<|im_end|>\n"
-            f"<|im_start|>user\n{placeholder}"
-            f"{question}<|im_end|>\n"
-            "<|im_start|>assistant\n"
-        )
-        for question in questions
-    ]
-
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompts=prompts,
-    )
-
-
 # R-4B
 def run_r_vl(questions: list[str], modality: str) -> ModelRequestData:
    assert modality == "image"
@@ -2528,8 +2442,6 @@ model_example_map = {
    "qwen2_5_omni": run_qwen2_5_omni,
    "qwen3_vl": run_qwen3_vl,
    "qwen3_vl_moe": run_qwen3_vl_moe,
-    "qwen3_5": run_qwen3_5,
-    "qwen3_5_moe": run_qwen3_5_moe,
    "rvl": run_r_vl,
    "skywork_chat": run_skyworkr1v,
    "smolvlm": run_smolvlm,
@@ -2548,17 +2460,13 @@ MODELS_NEED_VIDEO_METADATA = [
    "molmo2",
    "qwen3_vl",
    "qwen3_vl_moe",
-    "qwen3_5",
-    "qwen3_5_moe",
 ]


 MODELS_SUPPORT_VIT_CUDA_GRAPH = [
-    "qwen2_5_vl",
    "qwen3_vl",
    "qwen3_vl_moe",
-    "qwen3_5",
-    "qwen3_5_moe",
+    "qwen2_5_vl",
 ]


@@ -123,8 +123,7 @@ 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/speech_to_text/transcription/test_transcription_validation.py",
+    "vllm/third_party/*", "vllm/entrypoints/serve/instrumentator/static/*", "tests/entrypoints/openai/speech_to_text/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
@@ -20,6 +20,9 @@ nvidia-cudnn-frontend>=1.13.0,<1.19.0
 # Required for faster safetensors model loading
 fastsafetensors >= 0.2.2

-# QuACK and Cutlass DSL for FA4 and DeepSeek-V4 kernels
-nvidia-cutlass-dsl[cu13]>=4.5.0
+# QuACK and Cutlass DSL for FA4 (cute-DSL implementation)
+nvidia-cutlass-dsl==4.4.2
 quack-kernels>=0.3.3
+
+# Tokenspeed_MLA for faster mla with spec decode
+tokenspeed-mla==0.1.2
@@ -15,4 +15,4 @@ torch==2.11.0+xpu
 torchaudio
 torchvision

-vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.8/vllm_xpu_kernels-0.1.8-cp38-abi3-manylinux_2_28_x86_64.whl
+vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.7/vllm_xpu_kernels-0.1.7-cp38-abi3-manylinux_2_28_x86_64.whl
@@ -686,7 +686,6 @@ 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",
                ]
@@ -997,8 +996,6 @@ 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"))

@@ -1,166 +0,0 @@
-# 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,
-            )
@@ -1,413 +0,0 @@
-# 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,6 +34,7 @@ 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
@@ -101,8 +102,13 @@ class QKRoPEKVCacheTestModel(torch.nn.Module):
        )

        # Initialize attn MetadataBuilder
-        self.builder = self.attn_backend.get_builder_cls()(
-            kv_cache_spec=self.attn.get_kv_cache_spec(vllm_config),
+        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,
+            ),
            layer_names=[self.attn.layer_name],
            vllm_config=vllm_config,
            device=device,
@@ -120,11 +126,12 @@ class QKRoPEKVCacheTestModel(torch.nn.Module):
        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(
+        attn_backend = self.attn.attn_backend
+        kv_cache_shape = 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()
+            kv_cache_stride_order = attn_backend.get_kv_cache_stride_order()
        except (AttributeError, NotImplementedError):
            kv_cache_stride_order = tuple(range(len(kv_cache_shape)))

@@ -224,19 +224,28 @@ def init_test_http_connection():
 def dist_init():
    from tests.utils import ensure_current_vllm_config

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

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


@pytest.fixture
@@ -1686,44 +1695,3 @@ def disable_log_dedup(monkeypatch):
    logger._print_warning_once = original_print_warning_once
    logger._print_info_once = original_print_info_once
    logger._print_debug_once = original_print_debug_once
-
-
-@pytest.fixture(scope="function")
-def fake_vllm_ir(monkeypatch):
-    """
-    Pytest fixture to allow isolated IR op registration in tests.
-
-    Replaces IrOp.registry with an empty dict and swaps ``vllm_ir_torch_lib`` for a
-    fresh ``Library`` with a unique namespace per test (see ``Library.ns``).
-
-    Torch keeps registrations for the process lifetime; reusing the fragment
-    name ``vllm_ir`` and defining the same op string again can segfault. A
-    random library name keeps each fixture run on a disjoint namespace.
-
-    The test Library is kept alive until after monkeypatch teardown so PyTorch's
-    C++ state is not freed while references may still exist.
-
-    Usage:
-        def test_my_ir_op(fake_vllm_ir):
-            @vllm.ir.register_op
-            def my_test_op(x: torch.Tensor) -> torch.Tensor:
-                return x * 2
-
-            result = my_test_op(torch.tensor([1, 2, 3]))
-            # Registry and library cleaned up automatically after the test
-    """
-    import secrets
-
-    from torch.library import Library
-    from vllm.ir.op import IrOp
-
-    monkeypatch.setattr(IrOp, "registry", {})
-
-    # Keep a local reference so the Library is not GC'd before monkeypatch
-    # teardown restores the original reference.
-    test_lib = Library(f"vllm_ir_{secrets.token_hex(8)}", "FRAGMENT")
-    monkeypatch.setattr("vllm.ir.op.vllm_ir_torch_lib", test_lib)
-
-    yield
-
-    del test_lib
@@ -10,7 +10,6 @@ import torch

 from tests.kernels.moe.utils import make_test_quant_config
 from vllm.config import VllmConfig, set_current_vllm_config
-from vllm.distributed.eplb.eplb_state import EplbLayerState
 from vllm.distributed.eplb.rebalance_execute import rearrange_expert_weights_inplace
 from vllm.distributed.parallel_state import (
    ensure_model_parallel_initialized,
@@ -202,7 +201,7 @@ def _test_eplb_fml(env, world_size: int, test_config: TestConfig):
                dtype=torch.int32,
                device=device,
            )
-            fml.eplb_state = EplbLayerState()
+            fml.enable_eplb = True
            fml.set_eplb_state(
                lidx,
                torch.zeros(
@@ -3,9 +3,7 @@

 import pytest

-from vllm.model_executor.layers.fused_moe.expert_map_manager import (
-    determine_expert_map,
-)
+from vllm.model_executor.layers.fused_moe.layer import determine_expert_map


 def verify_round_robin_pattern(expert_map, ep_rank, ep_size, global_num_experts):
@@ -9,10 +9,7 @@ from vllm.distributed.kv_events import BlockRemoved, BlockStored
 _FAKE_HASH: bytes = b"\xab" * 32


-def _make_block_stored(
-    group_idx: int | None = None,
-    kv_cache_spec_sliding_window: int | None = None,
-) -> BlockStored:
+def _make_block_stored(group_idx: int | None = None) -> BlockStored:
    return BlockStored(
        block_hashes=[_FAKE_HASH],
        parent_block_hash=None,
@@ -22,13 +19,10 @@ def _make_block_stored(
        medium="GPU",
        lora_name=None,
        group_idx=group_idx,
-        kv_cache_spec_sliding_window=kv_cache_spec_sliding_window,
    )


-def _make_block_removed(
-    group_idx: int | None = None,
-) -> BlockRemoved:
+def _make_block_removed(group_idx: int | None = None) -> BlockRemoved:
    return BlockRemoved(
        block_hashes=[_FAKE_HASH],
        medium="GPU",
@@ -78,9 +72,3 @@ def test_block_removed_hash_same_for_equal_group_idx():
    event_a = _make_block_removed(group_idx=1)
    event_b = _make_block_removed(group_idx=1)
    assert hash(event_a) == hash(event_b)
-
-
-def test_block_stored_hash_differs_by_sliding_window():
-    event_a = _make_block_stored(group_idx=1, kv_cache_spec_sliding_window=128)
-    event_b = _make_block_stored(group_idx=1, kv_cache_spec_sliding_window=256)
-    assert hash(event_a) != hash(event_b)
@@ -67,3 +67,58 @@ 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")
@@ -1,492 +0,0 @@
-# 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
-        )
@@ -10,7 +10,7 @@ import pybase64 as base64
 import pytest
 import websockets

-from tests.entrypoints.speech_to_text.conftest import add_attention_backend
+from tests.entrypoints.openai.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,11 +163,6 @@ 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."
            )


@@ -325,12 +325,8 @@ 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, enable_thinking: bool
+    client: openai.AsyncOpenAI, model_name: str, tool_choice
 ):
    # this links the "done" type with the "start" type
    # so every "done" type should have a corresponding "start" type
@@ -440,7 +436,6 @@ 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,
    )

@@ -1,274 +0,0 @@
-# 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=[],
-            )
@@ -4,6 +4,7 @@
 from unittest.mock import patch

 import pytest
+from openai.types.chat import ChatCompletionMessageParam
 from openai.types.responses.response_function_tool_call import ResponseFunctionToolCall
 from openai.types.responses.response_function_tool_call_output_item import (
    ResponseFunctionToolCallOutputItem,
@@ -16,8 +17,10 @@ from openai.types.responses.response_reasoning_item import (
    Summary,
 )

+from vllm.entrypoints.constants import MCP_PREFIX
 from vllm.entrypoints.openai.responses.utils import (
-    _construct_message_from_response_item,
+    _construct_single_message_from_response_item,
+    _maybe_combine_reasoning_and_tool_call,
    construct_chat_messages_with_tool_call,
    construct_input_messages,
    convert_tool_responses_to_completions_format,
@@ -25,96 +28,6 @@ from vllm.entrypoints.openai.responses.utils import (
 )


-def _single_chat_message(item):
-    message = _construct_message_from_response_item(item)
-    assert message is not None
-    return message
-
-
-def make_output_message(
-    text: str,
-    *,
-    id: str = "msg_1",
-    status: str = "completed",
-) -> ResponseOutputMessage:
-    return ResponseOutputMessage(
-        id=id,
-        content=[
-            ResponseOutputText(
-                annotations=[],
-                text=text,
-                type="output_text",
-                logprobs=None,
-            )
-        ],
-        role="assistant",
-        status=status,
-        type="message",
-    )
-
-
-def make_reasoning_item(
-    *,
-    content_text: str | None = None,
-    summary_text: str | None = None,
-    content: list[Content] | None = None,
-    summary: list[Summary] | None = None,
-    encrypted_content: str | None = None,
-    id: str = "reasoning_1",
-    status: str | None = None,
-) -> ResponseReasoningItem:
-    if content is None and content_text is not None:
-        content = [Content(text=content_text, type="reasoning_text")]
-    if summary is None and summary_text is not None:
-        summary = [Summary(text=summary_text, type="summary_text")]
-
-    return ResponseReasoningItem(
-        id=id,
-        summary=[] if summary is None else summary,
-        type="reasoning",
-        content=content,
-        encrypted_content=encrypted_content,
-        status=status,
-    )
-
-
-def make_function_call(
-    *,
-    call_id: str,
-    name: str = "test_function",
-    arguments: str = "{}",
-    id: str = "tool_id",
-    status: str | None = None,
-) -> ResponseFunctionToolCall:
-    kwargs = {
-        "type": "function_call",
-        "id": id,
-        "call_id": call_id,
-        "name": name,
-        "arguments": arguments,
-    }
-    if status is not None:
-        kwargs["status"] = status
-
-    return ResponseFunctionToolCall(**kwargs)
-
-
-def make_function_call_output(
-    *,
-    call_id: str,
-    output: str = "42",
-    id: str = "output_1",
-    status: str = "completed",
-) -> ResponseFunctionToolCallOutputItem:
-    return ResponseFunctionToolCallOutputItem(
-        id=id,
-        type="function_call_output",
-        call_id=call_id,
-        output=output,
-        status=status,
-    )
-
-
 class TestResponsesUtils:
    """Tests for convert_tool_responses_to_completions_format function."""

@@ -174,7 +87,7 @@ class TestResponsesUtils:
            message["tool_calls"][0]["function"]["arguments"] == '{"code": "123+456"}'
        )

-    def test_construct_chat_messages_preserves_single_item_conversions(self):
+    def test_construct_single_message_from_response_item(self):
        item = ResponseReasoningItem(
            id="lol",
            summary=[],
@@ -188,7 +101,7 @@ class TestResponsesUtils:
            encrypted_content=None,
            status=None,
        )
-        formatted_item = _single_chat_message(item)
+        formatted_item = _construct_single_message_from_response_item(item)
        assert formatted_item["role"] == "assistant"
        assert formatted_item["reasoning"] == "Leroy Jenkins"

@@ -206,7 +119,7 @@ class TestResponsesUtils:
            status=None,
        )

-        formatted_item = _single_chat_message(item)
+        formatted_item = _construct_single_message_from_response_item(item)
        assert formatted_item["role"] == "assistant"
        assert (
            formatted_item["reasoning"]
@@ -220,22 +133,11 @@ class TestResponsesUtils:
            output="1234",
            status="completed",
        )
-        formatted_item = _single_chat_message(tool_call_output)
+        formatted_item = _construct_single_message_from_response_item(tool_call_output)
        assert formatted_item["role"] == "tool"
        assert formatted_item["content"] == "1234"
        assert formatted_item["tool_call_id"] == "temp"

-        formatted_item = _single_chat_message(
-            {
-                "type": "function_call_output",
-                "call_id": "temp_dict",
-                "output": "5678",
-            }
-        )
-        assert formatted_item["role"] == "tool"
-        assert formatted_item["content"] == "5678"
-        assert formatted_item["tool_call_id"] == "temp_dict"
-
        item = ResponseReasoningItem(
            id="lol",
            summary=[],
@@ -245,7 +147,7 @@ class TestResponsesUtils:
            status=None,
        )
        with pytest.raises(ValueError):
-            construct_chat_messages_with_tool_call([item])
+            _construct_single_message_from_response_item(item)

        output_item = ResponseOutputMessage(
            id="msg_bf585bbbe3d500e0",
@@ -262,7 +164,7 @@ class TestResponsesUtils:
            type="message",
        )

-        formatted_item = _single_chat_message(output_item)
+        formatted_item = _construct_single_message_from_response_item(output_item)
        assert formatted_item["role"] == "assistant"
        assert formatted_item["content"] == "dongyi"

@@ -290,7 +192,7 @@ class TestReasoningItemContentPriority:
            encrypted_content=None,
            status=None,
        )
-        formatted = _single_chat_message(item)
+        formatted = _construct_single_message_from_response_item(item)
        assert formatted["reasoning"] == "This is the actual content"

    def test_content_only(self):
@@ -308,7 +210,7 @@ class TestReasoningItemContentPriority:
            encrypted_content=None,
            status=None,
        )
-        formatted = _single_chat_message(item)
+        formatted = _construct_single_message_from_response_item(item)
        assert formatted["reasoning"] == "Content without summary"

    @patch("vllm.entrypoints.openai.responses.utils.logger")
@@ -327,7 +229,7 @@ class TestReasoningItemContentPriority:
            encrypted_content=None,
            status=None,
        )
-        formatted = _single_chat_message(item)
+        formatted = _construct_single_message_from_response_item(item)
        assert formatted["reasoning"] == "Fallback summary text"
        mock_logger.warning.assert_called_once()
        assert (
@@ -350,7 +252,7 @@ class TestReasoningItemContentPriority:
            encrypted_content=None,
            status=None,
        )
-        formatted = _single_chat_message(item)
+        formatted = _construct_single_message_from_response_item(item)
        assert formatted["reasoning"] == "Summary when content empty"
        mock_logger.warning.assert_called_once()
        assert (
@@ -367,7 +269,7 @@ class TestReasoningItemContentPriority:
            encrypted_content=None,
            status=None,
        )
-        formatted = _single_chat_message(item)
+        formatted = _construct_single_message_from_response_item(item)
        assert formatted["reasoning"] == ""

    def test_encrypted_content_raises(self):
@@ -391,7 +293,7 @@ class TestReasoningItemContentPriority:
            status=None,
        )
        with pytest.raises(ValueError):
-            construct_chat_messages_with_tool_call([item])
+            _construct_single_message_from_response_item(item)

    @patch("vllm.entrypoints.openai.responses.utils.logger")
    def test_summary_with_multiple_entries_uses_first(self, mock_logger):
@@ -413,7 +315,7 @@ class TestReasoningItemContentPriority:
            encrypted_content=None,
            status=None,
        )
-        formatted = _single_chat_message(item)
+        formatted = _construct_single_message_from_response_item(item)
        assert formatted["reasoning"] == "First summary"
        mock_logger.warning.assert_called_once()
        assert (
@@ -441,7 +343,7 @@ class TestReasoningItemContentPriority:
            encrypted_content=None,
            status=None,
        )
-        construct_chat_messages_with_tool_call([item])
+        _construct_single_message_from_response_item(item)
        mock_logger.warning.assert_not_called()


@@ -724,125 +626,119 @@ class TestShouldContinueFinalMessage:
        assert should_continue_final_message([dict_item]) is False


-class TestConstructChatMessagesCombinePolicy:
-    """Tests for contiguous assistant-side merging."""
+class TestMaybeCombineReasoningAndToolCall:
+    """Tests for _maybe_combine_reasoning_and_tool_call function."""

-    @pytest.mark.parametrize(
-        ("items", "expected_content", "expected_reasoning", "expected_tool_call_ids"),
-        [
-            pytest.param(
-                [
-                    make_reasoning_item(content_text="Let me think"),
-                    make_output_message("Hello"),
-                ],
-                "Hello",
-                "Let me think",
-                None,
-                id="reasoning-output-messages",
-            ),
-            pytest.param(
-                [
-                    make_function_call(call_id="call_123"),
-                    make_function_call(call_id="call_456"),
-                ],
-                None,
-                None,
-                ["call_123", "call_456"],
-                id="consecutive-tool-calls",
-            ),
-            pytest.param(
-                [
-                    make_reasoning_item(content_text="Let me think"),
-                    make_function_call(call_id="call_123"),
-                ],
-                None,
-                "Let me think",
-                ["call_123"],
-                id="reasoning-tool-call",
-            ),
-            pytest.param(
-                [
-                    make_output_message("Hello"),
-                    make_function_call(call_id="call_123"),
-                ],
-                "Hello",
-                None,
-                ["call_123"],
-                id="output-tool-call",
-            ),
-            pytest.param(
-                [
-                    make_reasoning_item(content_text="Thinking"),
-                    make_output_message("Hello"),
-                    make_function_call(call_id="call_123"),
-                    make_function_call(call_id="call_456"),
-                ],
-                "Hello",
-                "Thinking",
-                ["call_123", "call_456"],
-                id="reasoning-output-tool-call",
-            ),
-        ],
-    )
-    def test_assistant_side_items_merge_until_tool_output(
-        self,
-        items,
-        expected_content,
-        expected_reasoning,
-        expected_tool_call_ids,
-    ):
-        messages = construct_chat_messages_with_tool_call(items)
+    def test_returns_none_when_item_id_is_none(self):
+        """
+        Test fix from PR #31999: when item.id is None, should return None
+        instead of raising TypeError on startswith().
+        """
+        item = ResponseFunctionToolCall(
+            type="function_call",
+            id=None,  # This was causing TypeError before the fix
+            call_id="call_123",
+            name="test_function",
+            arguments="{}",
+        )
+        messages: list[ChatCompletionMessageParam] = []

-        assert len(messages) == 1
-        assert messages[0]["role"] == "assistant"
-        if expected_content is None:
-            assert "content" not in messages[0]
-        else:
-            assert messages[0]["content"] == expected_content
-        if expected_reasoning is None:
-            assert "reasoning" not in messages[0]
-        else:
-            assert messages[0]["reasoning"] == expected_reasoning
-        if expected_tool_call_ids is None:
-            assert "tool_calls" not in messages[0]
-        else:
-            assert [tool_call["id"] for tool_call in messages[0]["tool_calls"]] == (
-                expected_tool_call_ids
-            )
+        result = _maybe_combine_reasoning_and_tool_call(item, messages)

-    @pytest.mark.parametrize(
-        ("items", "num_expected_messages"),
-        [
-            pytest.param(
-                [
-                    make_output_message("Hello"),
-                    make_output_message("World"),
-                ],
-                2,
-                id="consecutive-output-messages",
-            ),
-            pytest.param(
-                [
-                    make_reasoning_item(content_text="Let me think"),
-                    make_reasoning_item(content_text="Let me think more"),
-                ],
-                2,
-                id="consecutive-reasoning-messages",
-            ),
-            pytest.param(
-                [
-                    make_function_call(call_id="call_123"),
-                    make_function_call_output(call_id="call_123", output="42"),
-                    make_function_call(call_id="call_456"),
-                ],
-                3,
-                id="interrupted-by-non-assistant-item",
-            ),
-        ],
-    )
-    def test_merge_chain_breaks(self, items, num_expected_messages):
-        messages = construct_chat_messages_with_tool_call(items)
-        assert len(messages) == num_expected_messages
+        assert result is None
+
+    def test_returns_none_when_id_does_not_start_with_mcp_prefix(self):
+        """Test that non-MCP tool calls are not combined."""
+        item = ResponseFunctionToolCall(
+            type="function_call",
+            id="regular_id",  # Does not start with MCP_PREFIX
+            call_id="call_123",
+            name="test_function",
+            arguments="{}",
+        )
+        messages = [{"role": "assistant", "reasoning": "some reasoning"}]
+
+        result = _maybe_combine_reasoning_and_tool_call(item, messages)
+
+        assert result is None
+
+    def test_returns_none_when_last_message_is_not_assistant(self):
+        """Test that non-assistant last message returns None."""
+        item = ResponseFunctionToolCall(
+            type="function_call",
+            id=f"{MCP_PREFIX}tool_id",
+            call_id="call_123",
+            name="test_function",
+            arguments="{}",
+        )
+        messages = [{"role": "user", "content": "hello"}]
+
+        result = _maybe_combine_reasoning_and_tool_call(item, messages)
+
+        assert result is None
+
+    def test_returns_none_when_last_message_has_no_reasoning(self):
+        """Test that assistant message without reasoning returns None."""
+        item = ResponseFunctionToolCall(
+            type="function_call",
+            id=f"{MCP_PREFIX}tool_id",
+            call_id="call_123",
+            name="test_function",
+            arguments="{}",
+        )
+        messages = [{"role": "assistant", "content": "some content"}]
+
+        result = _maybe_combine_reasoning_and_tool_call(item, messages)
+
+        assert result is None
+
+    def test_combines_reasoning_and_mcp_tool_call(self):
+        """Test successful combination of reasoning message and MCP tool call."""
+        item = ResponseFunctionToolCall(
+            type="function_call",
+            id=f"{MCP_PREFIX}tool_id",
+            call_id="call_123",
+            name="test_function",
+            arguments='{"arg": "value"}',
+        )
+        messages = [{"role": "assistant", "reasoning": "I need to call this tool"}]
+
+        result = _maybe_combine_reasoning_and_tool_call(item, messages)
+
+        assert result is not None
+        assert result["role"] == "assistant"
+        assert result["reasoning"] == "I need to call this tool"
+        assert "tool_calls" in result
+        assert len(result["tool_calls"]) == 1
+        assert result["tool_calls"][0]["id"] == "call_123"
+        assert result["tool_calls"][0]["function"]["name"] == "test_function"
+        assert result["tool_calls"][0]["function"]["arguments"] == '{"arg": "value"}'
+        assert result["tool_calls"][0]["type"] == "function"
+
+    def test_returns_none_for_non_function_tool_call_type(self):
+        """Test that non-ResponseFunctionToolCall items return None."""
+        # Pass a dict instead of ResponseFunctionToolCall
+        item = {"type": "message", "content": "hello"}
+        messages = [{"role": "assistant", "reasoning": "some reasoning"}]
+
+        result = _maybe_combine_reasoning_and_tool_call(item, messages)
+
+        assert result is None
+
+    def test_returns_none_when_id_is_empty_string(self):
+        """Test that empty string id returns None (falsy check)."""
+        item = ResponseFunctionToolCall(
+            type="function_call",
+            id="",  # Empty string is falsy
+            call_id="call_123",
+            name="test_function",
+            arguments="{}",
+        )
+        messages = [{"role": "assistant", "reasoning": "some reasoning"}]
+
+        result = _maybe_combine_reasoning_and_tool_call(item, messages)
+
+        assert result is None


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

 import pytest

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


 async def _never_finishes():
@@ -24,14 +24,12 @@ from vllm.entrypoints.openai.engine.protocol import (
    RequestResponseMetadata,
 )
 from vllm.entrypoints.openai.models.serving import OpenAIServingModels
-from vllm.entrypoints.speech_to_text.base.serving import (
+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 (
    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.speech_to_text.conftest import add_attention_backend
+from tests.entrypoints.openai.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.speech_to_text.conftest import add_attention_backend
+from tests.entrypoints.openai.conftest import add_attention_backend
 from tests.utils import RemoteOpenAIServer
 from vllm.logger import init_logger
 from vllm.multimodal.media.audio import load_audio
@@ -145,7 +145,6 @@ def test_openapi_stateless(case: Case):
        ("POST", "/v1/chat/completions/batch"): LONG_TIMEOUT_SECONDS,
        ("POST", "/v1/completions"): LONG_TIMEOUT_SECONDS,
        ("POST", "/v1/messages"): LONG_TIMEOUT_SECONDS,
-        ("POST", "/inference/v1/generate"): LONG_TIMEOUT_SECONDS,
    }.get(key, DEFAULT_TIMEOUT_SECONDS)

    # No need to verify SSL certificate for localhost
@@ -1,70 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Unit tests for the disagg request/response protocol.
-
-These tests intentionally avoid spinning up a server — they exercise the
-pydantic validators on ``GenerateRequest`` directly so they run fast and
-fail loudly if the validator semantics ever drift.
-"""
-
-import json
-
-from vllm.entrypoints.serve.disagg.protocol import GenerateRequest
-from vllm.sampling_params import SamplingParams
-
-
-def _base_payload() -> dict:
-    return {"token_ids": [1, 2, 3], "sampling_params": {}}
-
-
-def test_omitted_max_tokens_is_not_provided():
-    """Body without ``max_tokens`` must surface as 'not provided' so the
-    server can apply its own default instead of the dataclass 16."""
-    req = GenerateRequest.model_validate(_base_payload())
-    # SamplingParams' dataclass default leaks through the parsed instance —
-    # this is exactly the bug the server-side defaulting works around.
-    assert req.sampling_params.max_tokens == 16
-    assert not req.is_sampling_param_provided("max_tokens")
-
-
-def test_explicit_max_tokens_is_provided():
-    """Even when the client picks the same value as the dataclass default,
-    it must register as explicitly set so the server won't override it."""
-    payload = _base_payload()
-    payload["sampling_params"] = {"max_tokens": 16}
-    req = GenerateRequest.model_validate(payload)
-    assert req.sampling_params.max_tokens == 16
-    assert req.is_sampling_param_provided("max_tokens")
-
-    payload["sampling_params"] = {"max_tokens": 256}
-    req = GenerateRequest.model_validate(payload)
-    assert req.sampling_params.max_tokens == 256
-    assert req.is_sampling_param_provided("max_tokens")
-
-
-def test_other_fields_tracked_independently():
-    payload = _base_payload()
-    payload["sampling_params"] = {"temperature": 0.0}
-    req = GenerateRequest.model_validate(payload)
-    assert not req.is_sampling_param_provided("max_tokens")
-    assert req.is_sampling_param_provided("temperature")
-
-
-def test_json_roundtrip_preserves_provided_keys():
-    payload = _base_payload()
-    payload["sampling_params"] = {"temperature": 0.5}
-    req = GenerateRequest.model_validate_json(json.dumps(payload))
-    assert not req.is_sampling_param_provided("max_tokens")
-    assert req.is_sampling_param_provided("temperature")
-
-
-def test_internal_instance_construction_treats_all_as_provided():
-    """When internal callers build ``GenerateRequest`` from a pre-resolved
-    ``SamplingParams`` instance, every field is considered explicitly set
-    so server-side defaulting can't clobber values resolved upstream."""
-    sp = SamplingParams(max_tokens=500, temperature=0.0)
-    req = GenerateRequest(token_ids=[1, 2, 3], sampling_params=sp)
-    assert req.is_sampling_param_provided("max_tokens")
-    assert req.is_sampling_param_provided("temperature")
-    # And keys we never touched should also count as provided in this path.
-    assert req.is_sampling_param_provided("top_p")
@@ -114,36 +114,6 @@ async def test_generate_endpoint(client):
    assert "choices" in data


-@pytest.mark.asyncio
-async def test_generate_defaults_max_tokens_when_omitted(client):
-    """Regression: omitting ``max_tokens`` must not silently cap at 16.
-
-    ``SamplingParams.max_tokens`` defaults to 16. Before the server-side
-    defaulting was wired up, every request that didn't set ``max_tokens``
-    truncated mid-generation. The server should now fill it in from
-    ``max_model_len - prompt_len`` (matching ``/v1/chat/completions``).
-    """
-    payload = {
-        "model": MODEL_NAME,
-        "token_ids": [1, 2, 3],
-        "sampling_params": {
-            "temperature": 0.0,
-            "ignore_eos": True,
-        },
-        "stream": False,
-    }
-    resp = await client.post(GEN_ENDPOINT, json=payload)
-    resp.raise_for_status()
-    data = resp.json()
-    completion_tokens = len(data["choices"][0]["token_ids"])
-    # max_model_len=1024 in the test fixture, prompt is 3 tokens, so we
-    # should get ~1021 tokens of output (capped at max_model_len boundary).
-    assert completion_tokens > 16, (
-        f"expected server-side default to exceed the legacy 16-token cap, "
-        f"got {completion_tokens}"
-    )
-
-
@pytest.mark.asyncio
 async def test_generate_stream(client):
    payload = {
@@ -1,61 +0,0 @@
-# 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")
@@ -1,3 +0,0 @@
-Qwen3-Next-80B-A3B-NVFP4-EP2.yaml
-Qwen3-Next-FP8-EP2.yaml
-Nemotron-3-Super-120B-A12B-NVFP4.yaml
@@ -3,3 +3,6 @@ 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
@@ -13,40 +13,20 @@ from torch.fx.experimental.proxy_tensor import make_fx
 import vllm.ir.op
 from vllm.ir.op import RESERVED_PROVIDERS, IrOp, IrOpImpl

+# This should not exist
+assert "_custom_add" not in IrOp.registry
+

 class CustomError(Exception):
    pass


-@pytest.fixture
-def custom_add_op(fake_vllm_ir):
-    """Register ``_custom_add`` plus impl_a, impl_b, impl_even for this test."""
-
-    @vllm.ir.register_op(allow_inplace=True)
-    def _custom_add(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
-        return x + y
-
-    @_custom_add.register_impl("impl_a")
-    def impl_a(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
-        return x + y + 10
-
-    @_custom_add.register_impl("impl_b", inplace=True)
-    def impl_b(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
-        """Computes x+y+20"""
-        x.add_(y)
-        x.add_(20)
-        return x
-
-    @_custom_add.register_impl(
-        "impl_even", supports_args=lambda x, y: x.size(1) % 2 == 0
-    )
-    def impl_even(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
-        return x + y + 50
-
-    return _custom_add
+@vllm.ir.register_op(allow_inplace=True)
+def _custom_add(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+    return x + y


-def test_registration_overloads(fake_vllm_ir):
+def test_registration_overloads():
    assert all(
        n not in IrOp.registry for n in ["_custom_sub", "_custom_mul", "_custom_div"]
    )
@@ -83,7 +63,7 @@ def test_registration_overloads(fake_vllm_ir):
            return x * y - 100


-def test_no_kw_only_args(fake_vllm_ir):
+def test_no_kw_only_args():
    # kw-only args not supported
    with pytest.raises(ValueError, match="keyword-only arguments"):

@@ -98,23 +78,18 @@ def test_no_kw_only_args(fake_vllm_ir):

 class TestIrOpCustomAdd:
    # Registration invariants
-    def test_decorated_object(self, custom_add_op):
+    def test_decorated_object(self):
        """Make sure that referring directly to an op is correct"""
-        _custom_add = custom_add_op
        assert isinstance(_custom_add, IrOp)
        assert "_custom_add" in IrOp.registry
        assert _custom_add is IrOp.registry["_custom_add"]

-    def test_torch_op_is_registered(self, custom_add_op):
-        _custom_add = custom_add_op
-        torch_ops = getattr(torch.ops, vllm.ir.op.vllm_ir_torch_lib.ns)
-        assert hasattr(torch_ops, "_custom_add")
-        assert callable(torch_ops._custom_add.default)
-        assert _custom_add.torch_op is torch_ops._custom_add.default
+    def test_torch_op_is_registered(self):
+        assert hasattr(torch.ops.vllm_ir, "_custom_add")
+        assert callable(torch.ops.vllm_ir._custom_add.default)

    # Semantic correctness
-    def test_semantics_match_native(self, custom_add_op):
-        _custom_add = custom_add_op
+    def test_semantics_match_native(self):
        x = torch.randn(4, 5)
        y = torch.randn(4, 5)

@@ -128,9 +103,7 @@ class TestIrOpCustomAdd:
    # Implementation registration
    # -------------------------

-    def test_register_impl_is_non_intrusive(self, custom_add_op):
-        _custom_add = custom_add_op
-
+    def test_register_impl_is_non_intrusive(self):
        @_custom_add.register_impl("dummy_provider")
        def dummy_impl(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
            return x + y + 123
@@ -144,8 +117,7 @@ class TestIrOpCustomAdd:
        # Native semantics must still hold
        torch.testing.assert_close(_custom_add(x, y), x + y)

-    def test_schema_contains_tensor_signature(self, custom_add_op):
-        _custom_add = custom_add_op
+    def test_schema_contains_tensor_signature(self):
        schema = _custom_add._schema_str

        assert "Tensor" in schema
@@ -159,19 +131,10 @@ class TestIrOpCustomAdd:
    @pytest.mark.parametrize("symbolic_trace", [True, False])
    @pytest.mark.parametrize("overload", ["default", "maybe_inplace"])
    def test_trace_sees_single_custom_op(
-        self,
-        custom_add_op,
-        symbolic_trace: bool,
-        enable_torch_wrap: bool,
-        overload: str,
+        self, symbolic_trace: bool, enable_torch_wrap: bool, overload: str
    ):
-        _custom_add = custom_add_op
        op_fn = _custom_add if overload == "default" else _custom_add.maybe_inplace
-        torch_op = (
-            _custom_add.torch_op
-            if overload == "default"
-            else _custom_add.maybe_inplace.torch_op
-        )
+        torch_op = getattr(torch.ops.vllm_ir._custom_add, overload)

        def fn(x, y):
            return op_fn(x, y)
@@ -212,13 +175,30 @@ class TestIrOpCustomAdd:
        assert len(ir_nodes) == 1, gm.code


+@_custom_add.register_impl("impl_a")
+def impl_a(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+    return x + y + 10
+
+
+@_custom_add.register_impl("impl_b", inplace=True)
+def impl_b(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+    """Computes x+y+20"""
+    x.add_(y)
+    x.add_(20)
+    return x
+
+
+@_custom_add.register_impl("impl_even", supports_args=lambda x, y: x.size(1) % 2 == 0)
+def impl_even(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+    return x + y + 50
+
+
 class TestIrOpImplDispatch:
-    def test_register_impl(self, custom_add_op):
-        _custom_add = custom_add_op
+    def test_register_impl(self):
        assert "impl_a" in _custom_add.impls
        impl = _custom_add.impls["impl_a"]

-        assert impl is _custom_add.impls["impl_a"]
+        assert impl is impl_a
        assert impl.op is _custom_add
        assert impl.provider == "impl_a"
        assert callable(impl.impl_fn)
@@ -231,15 +211,14 @@ class TestIrOpImplDispatch:
                return x + y + 30

        # Check the original impl is still intact
-        assert _custom_add.impls["impl_a"] is impl
+        assert _custom_add.impls["impl_a"] is impl_a

        # Check support all args
-        assert _custom_add.impls["impl_a"].supports_all_args
-        assert _custom_add.impls["impl_b"].supports_all_args
-        assert not _custom_add.impls["impl_even"].supports_all_args
+        assert impl_a.supports_all_args
+        assert impl_b.supports_all_args
+        assert not impl_even.supports_all_args

-    def test_reserved_provider_rejected(self, custom_add_op):
-        _custom_add = custom_add_op
+    def test_reserved_provider_rejected(self):
        for provider in RESERVED_PROVIDERS:
            with pytest.raises(AssertionError):

@@ -247,8 +226,7 @@ class TestIrOpImplDispatch:
                def bad_impl(x, y):
                    return x + y

-    def test_set_priority_scoped(self, custom_add_op):
-        _custom_add = custom_add_op
+    def test_set_priority_scoped(self):
        assert _custom_add.get_priority() == []

        with _custom_add.set_priority(["impl_even", "impl_b"]):
@@ -273,25 +251,20 @@ class TestIrOpImplDispatch:
        assert _custom_add.get_priority() == []

    @pytest.mark.parametrize("overload", ["default", "maybe_inplace"])
-    def test_dispatch_priority_order(self, custom_add_op, overload: str):
-        _custom_add = custom_add_op
+    def test_dispatch_priority_order(self, overload: str):
        op_fn = _custom_add if overload == "default" else _custom_add.maybe_inplace
-        torch_op = (
-            _custom_add.torch_op
-            if overload == "default"
-            else _custom_add.maybe_inplace.torch_op
-        )
+        torch_op = getattr(torch.ops.vllm_ir._custom_add, overload)

        x = torch.tensor(1, dtype=torch.int32)
        y = torch.tensor(2, dtype=torch.int32)

        with _custom_add.set_priority(["impl_b", "impl_a"]):
-            assert _custom_add.dispatch(x, y) is _custom_add.impls["impl_b"]
+            assert _custom_add.dispatch(x, y) is impl_b
            out1 = op_fn(x.clone(), y)
            out2 = torch_op(x.clone(), y)

            with _custom_add.set_priority(["impl_a"]):
-                assert _custom_add.dispatch(x, y) is _custom_add.impls["impl_a"]
+                assert _custom_add.dispatch(x, y) is impl_a
                out3 = op_fn(x.clone(), y)
                out4 = torch_op(x.clone(), y)

@@ -302,9 +275,7 @@ class TestIrOpImplDispatch:
        assert out3.item() == 1 + 2 + 10
        assert out4.item() == 1 + 2 + 10

-    def test_unsupported_impl_filtered(self, custom_add_op):
-        _custom_add = custom_add_op
-
+    def test_unsupported_impl_filtered(self):
        @_custom_add.register_impl("impl_unsupported", supported=False)
        def impl_unsupported(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
            return x + y + 999
@@ -320,9 +291,8 @@ class TestIrOpImplDispatch:
        assert out.item() == 1 + 2 + 10

    def test_supports_args_runtime_dispatch_and_warning(
-        self, custom_add_op, caplog_vllm: pytest.LogCaptureFixture
+        self, caplog_vllm: pytest.LogCaptureFixture
    ):
-        _custom_add = custom_add_op
        x1 = torch.ones((2, 2), dtype=torch.int32)
        y1 = torch.full((2, 2), 2, dtype=torch.int32)

@@ -342,7 +312,7 @@ class TestIrOpImplDispatch:

            # Check dispatching
            assert _custom_add.get_priority() == ["impl_even", "native"]
-            assert _custom_add.dispatch(x1, y1) is _custom_add.impls["impl_even"]
+            assert _custom_add.dispatch(x1, y1) is impl_even
            assert _custom_add.dispatch(x2, y2) is _custom_add.impls["native"]

            out1 = _custom_add(x1, y1)  # size(1) == 2 → impl_even
@@ -354,12 +324,8 @@ class TestIrOpImplDispatch:
        assert torch.all(out2 == 1 + 2)

    def test_default_priority(
-        self,
-        custom_add_op,
-        caplog_vllm: pytest.LogCaptureFixture,
-        disable_log_dedup,
+        self, caplog_vllm: pytest.LogCaptureFixture, disable_log_dedup
    ):
-        _custom_add = custom_add_op
        # Make sure logs are not deduplicated to properly test the warning
        x = torch.tensor([3], dtype=torch.int32)
        y = torch.tensor([4], dtype=torch.int32)
@@ -381,23 +347,15 @@ class TestIrOpImplDispatch:
        assert "priority not set" in message


-@pytest.fixture
-def custom_mm_op(fake_vllm_ir):
-    """Fixture that registers ``_custom_mm`` (isolated by ``fake_vllm_ir``)."""
-
-    @vllm.ir.register_op
-    def _custom_mm(
-        x: torch.Tensor, y: torch.Tensor, bias: torch.Tensor | None = None
-    ) -> torch.Tensor:
-        tmp = x @ y
-        return tmp if bias is None else tmp + bias
-
-    return _custom_mm
+@vllm.ir.register_op
+def _custom_mm(
+    x: torch.Tensor, y: torch.Tensor, bias: torch.Tensor | None = None
+) -> torch.Tensor:
+    tmp = x @ y
+    return tmp if bias is None else tmp + bias


-def test_default_args(custom_mm_op):
-    _custom_mm = custom_mm_op
-
+def test_default_args():
    # Test that default args are properly applied when dispatching and calling
    @_custom_mm.register_impl("impl_mm", supports_args=lambda x, y, bias=None: True)
    def impl_mm(
@@ -415,8 +373,7 @@ def test_default_args(custom_mm_op):
        assert _custom_mm.dispatch(x1, x2) is impl_mm


-def test_bad_impl_registrations(custom_mm_op):
-    _custom_mm = custom_mm_op
+def test_bad_impl_registrations():
    # Check bad schema
    with pytest.raises(ValueError, match="does not match native schema"):

@@ -489,8 +446,7 @@ def test_bad_impl_registrations(custom_mm_op):
        ) -> torch.Tensor:
            return x @ y + 40

-    # With fixture, each test gets a fresh op with only "native" impl
-    assert set(_custom_mm.impls.keys()) == {"native"}
+    assert set(_custom_mm.impls.keys()) == {"impl_mm", "native"}


 IMPL_OOT_SRC = """
@@ -504,14 +460,14 @@ def impl_mm_oot(
 """


-def load_custom_mm_module(file_path: Path, custom_mm_op):
+def load_custom_mm_module(file_path: Path):
    spec = importlib.util.spec_from_file_location("_custom_mm_oot", file_path)
    assert spec is not None
    module = importlib.util.module_from_spec(spec)

    # Inject the variable into the module's global namespace
    # This allows the @_custom_mm.register_impl decorator to work
-    module._custom_mm = custom_mm_op  # type: ignore[attr-defined]
+    module._custom_mm = _custom_mm  # type: ignore[attr-defined]

    # Execute the file; this triggers the decorator
    assert spec.loader is not None
@@ -519,13 +475,12 @@ def load_custom_mm_module(file_path: Path, custom_mm_op):
    return module


-def test_uuid_and_oot(custom_mm_op, tmp_path: Path):
-    _custom_mm = custom_mm_op
+def test_uuid_and_oot(tmp_path: Path):
    file_path = tmp_path / "_custom_mm_oot.py"
    file_path.write_text(IMPL_OOT_SRC)

    assert "impl_mm_oot" not in _custom_mm.impls
-    _ = load_custom_mm_module(file_path, _custom_mm)
+    _ = load_custom_mm_module(file_path)
    assert "impl_mm_oot" in _custom_mm.impls

    uuid = _custom_mm.impls["impl_mm_oot"].uuid()
@@ -534,7 +489,7 @@ def test_uuid_and_oot(custom_mm_op, tmp_path: Path):
    # Replace file source
    file_path.write_text(IMPL_OOT_SRC + " # added file source")
    assert "impl_mm_oot" not in _custom_mm.impls
-    _ = load_custom_mm_module(file_path, _custom_mm)
+    _ = load_custom_mm_module(file_path)
    assert "impl_mm_oot" in _custom_mm.impls

    uuid1 = _custom_mm.impls["impl_mm_oot"].uuid()
@@ -544,7 +499,7 @@ def test_uuid_and_oot(custom_mm_op, tmp_path: Path):
    # Back to original
    file_path.write_text(IMPL_OOT_SRC)
    assert "impl_mm_oot" not in _custom_mm.impls
-    _ = load_custom_mm_module(file_path, _custom_mm)
+    _ = load_custom_mm_module(file_path)
    assert "impl_mm_oot" in _custom_mm.impls

    uuid2 = _custom_mm.impls["impl_mm_oot"].uuid()
@@ -616,106 +571,3 @@ class TestTolerance:
        op = IrOp("_tol_test_unknown", _test_native)
        with pytest.raises(ValueError, match="No tolerance defined"):
            op.get_tolerance(torch.complex64)
-
-
-def test_naming_validation(fake_vllm_ir):
-    """Test that op and provider names are validated ([a-z_][a-z_0-9]*)."""
-
-    # Valid op and provider names
-    @vllm.ir.register_op
-    def _valid_name_123(x: torch.Tensor) -> torch.Tensor:
-        return x
-
-    @_valid_name_123.register_impl("valid_provider_123")
-    def valid_impl(x: torch.Tensor) -> torch.Tensor:
-        return x + 1
-
-    # Invalid op names should fail
-    with pytest.raises(ValueError, match="name.*invalid"):
-
-        @vllm.ir.register_op
-        def InvalidName(x: torch.Tensor) -> torch.Tensor:
-            return x
-
-    with pytest.raises(ValueError, match="name.*invalid"):
-
-        @vllm.ir.register_op(name="123invalid")
-        def some_func(x: torch.Tensor) -> torch.Tensor:
-            return x
-
-    # Invalid provider names should fail
-    with pytest.raises(ValueError, match="name.*invalid"):
-
-        @_valid_name_123.register_impl("Invalid-Provider")
-        def invalid_impl(x: torch.Tensor) -> torch.Tensor:
-            return x + 1
-
-
-def test_registration_stack_traces(fake_vllm_ir):
-    """Test that stack traces are captured for ops and impls."""
-
-    @vllm.ir.register_op
-    def _test_stack(x: torch.Tensor) -> torch.Tensor:
-        return x
-
-    @_test_stack.register_impl("test_provider")
-    def test_impl(x: torch.Tensor) -> torch.Tensor:
-        return x + 1
-
-    # Verify op stack trace
-    assert hasattr(_test_stack, "_registration_stack")
-    assert len(_test_stack._registration_stack) > 0
-    op_stack_str = "".join(_test_stack._registration_stack)
-    assert "test_op.py" in op_stack_str
-    # Last frame should be the decorator in user code, not internal decorator logic
-    assert "@vllm.ir.register_op" in _test_stack._registration_stack[-1]
-    assert "return decorator(f)" not in op_stack_str
-
-    # Verify impl stack trace
-    impl = _test_stack.impls["test_provider"]
-    assert hasattr(impl, "_registration_stack")
-    assert len(impl._registration_stack) > 0
-    impl_stack_str = "".join(impl._registration_stack)
-    assert "test_op.py" in impl_stack_str
-    # Last frame should be the decorator in user code
-    assert '@_test_stack.register_impl("test_provider")' in impl._registration_stack[-1]
-
-
-def test_op_repr_uses_docstring(fake_vllm_ir):
-    """Test that __str__ uses the function's docstring and __repr__ is simple."""
-
-    @vllm.ir.register_op
-    def _test_repr_with_doc(x: torch.Tensor) -> torch.Tensor:
-        """First line of docstring.
-
-        Additional details here.
-        """
-        return x
-
-    @vllm.ir.register_op
-    def _test_repr_no_doc(x: torch.Tensor) -> torch.Tensor:
-        return x
-
-    # __str__ with docstring: uses first line only
-    str_with = str(_test_repr_with_doc)
-    assert "IrOp('_test_repr_with_doc')" in str_with
-    assert "First line of docstring." in str_with
-    assert "Additional details" not in str_with
-
-    # __str__ without docstring: simple format
-    assert str(_test_repr_no_doc) == "IrOp('_test_repr_no_doc')"
-
-    # __repr__ should be simple for both
-    assert repr(_test_repr_with_doc) == "IrOp('_test_repr_with_doc')"
-    assert repr(_test_repr_no_doc) == "IrOp('_test_repr_no_doc')"
-
-
-def test_vllm_ir_fixture(fake_vllm_ir):
-    """Test that the fake_vllm_ir fixture provides test isolation."""
-
-    @vllm.ir.register_op
-    def _test_fixture(x: torch.Tensor) -> torch.Tensor:
-        return x
-
-    assert "_test_fixture" in IrOp.registry
-    # Fixture will automatically clean up after test
@@ -1,377 +0,0 @@
-# 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)
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Mohammad Miadh Angkadandkhluu	ad7125a431	[Bugfix] Fix DeepSeek V4 MTP HC state handling (#42320 ) Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com> (cherry picked from commit `f1cc7aad3c`)	2026-05-14 21:28:34 -07:00
Yongye Zhu khluu Claude	9da56fd18b	[Bugfix] Add swiglu limits to deepgemm fp8 methods (#41986 ) Cherry-picked from https://github.com/vllm-project/vllm/pull/41986 Plumb SwiGLU clamp limit through DeepGemm FP8/W4A8 MoE quant configs and experts. Extend silu_mul_per_token_group_quant_fp8_colmajor with clamp support and forward the limit on all FP8/MXFP8/MXFP4 paths. Co-authored-by: Yongye Zhu <zyy1102000@gmail.com> Co-authored-by: Claude <noreply@anthropic.com> Signed-off-by: khluu <khluu000@gmail.com>	2026-05-14 12:38:36 -07:00
Yongye Zhuandkhluu	800604bf53	[MLA Attention Backend] Add TOKENSPEED_MLA backend for DSR1/Kimi K25 prefill + decode on Blackwell (#41778 ) Signed-off-by: Yongye Zhu <zyy1102000@gmail.com> Signed-off-by: Roger Wang <hey@rogerw.io> Co-authored-by: Roger Wang <hey@rogerw.io> (cherry picked from commit `0d2732dd91`)	2026-05-14 00:59:51 -07:00
khluu	75a7914326	pin cutlass-dsl to 4.4.2 Signed-off-by: khluu <khluu000@gmail.com>	2026-05-14 00:59:01 -07:00
ovidiusmandkhluu	3b581add43	[PD] Fix broken NIXL EP installation (#42542 ) Signed-off-by: Ovidiu Mara <ovidium@nvidia.com> (cherry picked from commit `cca32d55a2`)	2026-05-13 15:15:07 -07:00
Kevin H. Luu	342cec8812	[CI] Use uv with Python 3.12 for PyPI wheel upload (#42470 ) Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> (cherry picked from commit `f6e868fbdf`)	2026-05-13 02:12:34 -07:00
Jiangyun Zhuandkhluu	135453b715	[Bugfix] Install nvidia-cutlass-dsl[cu13] extra on CUDA 13 platforms (#42438 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> (cherry picked from commit `140dc2ec30`)	2026-05-13 02:03:17 -07:00
sychen52andkhluu	a707288c1e	Patch SlidingWindowSpec.real_page_size_bytes for nvfp4 kv (#42464 ) Signed-off-by: Shiyang Chen <shiychen@nvidia.com> (cherry picked from commit `a8c13d2837`)	2026-05-13 02:03:07 -07:00
Alecandkhluu	638f8fa979	[PD] Bump NIXL connector dependency to 1.x (#42364 ) Signed-off-by: Alec Flowers <aflowers@nvidia.com> (cherry picked from commit `07534b8782`)	2026-05-13 02:02:55 -07:00
Chao Leiandkhluu	cbaa80fede	[KV Transfer] Add MooncakeStoreConnector for KV cache offloading via Mooncake distributed store (#40900 ) Signed-off-by: leichao.lc <leichao.lc@antgroup.com> Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai> Co-authored-by: leichao.lc <leichao.lc@antgroup.com> Co-authored-by: ivanium <yifanqiao@inferact.ai> Co-authored-by: aoshen524 <aoshen@inferact.ai> Co-authored-by: Dao007forever <daole@inferact.ai> Co-authored-by: Teng Ma <sima.mt@alibaba-inc.com> Co-authored-by: Pz1116 <zpbzpb123123@gmail.com> Co-authored-by: foraxe <1055696449@qq.com> Co-authored-by: Skywalker-EP <173423846@qq.com> Co-authored-by: fems14 <1804143737@qq.com> Co-authored-by: jianzs <zheng.shoujian@outlook.com> Co-authored-by: baxingpiaochong <771405853@qq.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> (cherry picked from commit `ebeb09d822`)	2026-05-13 02:02:44 -07:00
Kevin H. Luu	84a1066ccc	[CI] Inline build artifact annotations in release pipeline (#42357 ) Signed-off-by: khluu <khluu000@gmail.com> (cherry picked from commit `8c4fc4202a`)	2026-05-13 02:02:30 -07:00
Michael Goinandkhluu	d801ae8c26	[Build] Build bundled DeepGEMM `_C` per-Python so the wheel imports on every CPython (#41516 ) Signed-off-by: mgoin <mgoin64@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> (cherry picked from commit `d077622d60`)	2026-05-12 14:57:17 -07:00
Jiahan Chang (Cyrus)andkhluu	65df49eba3	[Perf] Use 2D-grid to eliminate divmod in W8W8 group quant (#42153 ) Signed-off-by: jiahanc <173873397+jiahanc@users.noreply.github.com> Co-authored-by: Yongye Zhu <zyy1102000@gmail.com> (cherry picked from commit `dd6b3a5ef5`)	2026-05-12 14:57:06 -07:00
Kevin H. Luu	2a2ac21d3d	[CI] Move DockerHub and PyPI publish steps to end of release pipeline (#42355 ) Signed-off-by: khluu <khluu000@gmail.com> (cherry picked from commit `e1c8776e90`)	2026-05-12 14:56:46 -07:00
Jee Jee Liandkhluu	c6fc95806b	[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> (cherry picked from commit `53181384e0`)	2026-05-12 14:56:29 -07:00