test: touch tilelang.py for coverage comparison

2026-05-28 13:33:39 -07:00
451 changed files with 5587 additions and 21004 deletions
@@ -16,7 +16,6 @@ steps:
  - tests/kernels/test_onednn.py
  - tests/kernels/test_awq_int4_to_int8.py
  - tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
-  - tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
  commands:
    - |
      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
@@ -25,8 +24,7 @@ steps:
      pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
      pytest -x -v -s tests/kernels/test_onednn.py
      pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
-      pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
-      pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
+      pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py"

 - label: CPU-Compatibility Tests
  depends_on: []
@@ -64,16 +62,11 @@ steps:
  source_file_dependencies:
  - vllm/v1/worker/cpu/
  - vllm/v1/worker/gpu/
-  - vllm/v1/sample/ops/topk_topp_triton.py
-  - vllm/v1/sample/ops/topk_topp_sampler.py
-  - tests/v1/sample/test_topk_topp_sampler.py
  commands:
    - |
      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
      uv pip install git+https://github.com/triton-lang/triton-cpu.git@270e696d
-      VLLM_USE_V2_MODEL_RUNNER=1 pytest -x -v -s tests/models/language/generation/test_granite.py -m cpu_model
-      # TODO: move to CPU-Kernel Tests once triton-cpu has a pre-built wheel
-      pytest -x -v -s tests/v1/sample/test_topk_topp_sampler.py::TestTritonTopkTopp"
+      VLLM_USE_V2_MODEL_RUNNER=1 pytest -x -v -s tests/models/language/generation/test_granite.py -m cpu_model"

 - label: CPU-Quantization Model Tests
  depends_on: []
@@ -6,26 +6,14 @@ steps:
    timeout_in_minutes: 600
    commands:
    - if [[ "$BUILDKITE_BRANCH" == "main" ]]; then .buildkite/image_build/image_build.sh $REGISTRY $REPO $BUILDKITE_COMMIT $BRANCH $IMAGE_TAG $IMAGE_TAG_LATEST; else .buildkite/image_build/image_build.sh $REGISTRY $REPO $BUILDKITE_COMMIT $BRANCH $IMAGE_TAG; fi
-    retry:
-      automatic:
-        - exit_status: -1  # Agent was lost
-          limit: 2
-        - exit_status: -10  # Agent was lost
-          limit: 2
-
-  - label: ":docker: :smoking: Non-root smoke tests"
-    key: image-build-smoke-test
-    depends_on:
-      - image-build
-    commands:
-    # Smoke 1: the default (root) image must still be importable
+    # Non-root smoke 1: the default (root) image must still be importable
    # under a non-root UID via `--user 2000:0`. Validates the `vllm` passwd
    # entry + group-0-writable /home/vllm + uv path cleanup from #31959.
    # Uses `import vllm` rather than `vllm serve --help` because the latter
    # instantiates `VllmConfig` which requires a GPU attached to the
    # container.
    - docker run --rm --user 2000:0 --entrypoint python3 "$IMAGE_TAG" -c "import vllm; print(vllm.__version__)"
-    # Smoke 2: assert the non-root enabling invariants are baked
+    # Non-root smoke 2: assert the non-root enabling invariants are baked
    # into the image. Runs as UID 2000:0 via a shell so we can verify
    # filesystem perms + passwd/group file state + wrapper presence without
    # triggering vLLM's GPU-requiring config-init path. The opt-in
@@ -9,13 +9,6 @@
 # Find <build_number> and <job_uuid> via:
 #   gh pr checks <PR> --repo vllm-project/vllm
 # Each failing row's URL is .../builds/<build_number>#<job_uuid>.
-#
-# Default output path: ci-<build>-<uuid_first_13_chars>.log (e.g.
-# ci-68478-019e6b07-daae.log). Jobs in the same build share the UUID's
-# first 8 chars, so the second segment is needed for uniqueness when
-# fetching multiple jobs in parallel. The script refuses to overwrite an
-# existing output file; pass an explicit path or set CI_FETCH_LOG_FORCE=1
-# to override.

 set -euo pipefail

@@ -33,12 +26,12 @@ if [ $# -lt 1 ]; then usage; fi
 if [[ "$1" == https://* ]]; then
    BUILD=$(echo "$1" | sed -nE 's#.*/builds/([0-9]+).*#\1#p')
    JOB=$(echo "$1" | grep -oE '[0-9a-f]{8}-[0-9a-f-]+' | head -n 1)
-    OUT="${2:-}"
+    OUT="${2:-ci-${BUILD}-${JOB:0:8}.log}"
 else
    if [ $# -lt 2 ]; then usage; fi
    BUILD="$1"
    JOB="$2"
-    OUT="${3:-}"
+    OUT="${3:-ci-${BUILD}-${JOB:0:8}.log}"
 fi

 if [ -z "$BUILD" ] || [ -z "$JOB" ]; then
@@ -46,18 +39,6 @@ if [ -z "$BUILD" ] || [ -z "$JOB" ]; then
    usage
 fi

-# Jobs in the same build share the UUID's first segment, so include the
-# second segment (chars 9-13, e.g. "019e6b07-daae") to keep default filenames
-# unique when fetching multiple jobs from one build in parallel.
-if [ -z "$OUT" ]; then
-    OUT="ci-${BUILD}-${JOB:0:13}.log"
-fi
-
-if [ -e "$OUT" ] && [ -z "${CI_FETCH_LOG_FORCE:-}" ]; then
-    echo "Refusing to overwrite existing $OUT (set CI_FETCH_LOG_FORCE=1 or pass an explicit output path)." >&2
-    exit 1
-fi
-
 COOKIES=$(mktemp)
 trap 'rm -f "$COOKIES"' EXIT

@@ -37,8 +37,7 @@ function cpu_tests() {
    pytest -x -v -s tests/kernels/test_onednn.py
    pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
    pytest -x -v -s tests/kernels/core/test_cpu_activation.py
-    pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic
-    pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
+    pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic"

  # skip tests requiring model downloads if HF_TOKEN is not set
  # due to rate-limits
@@ -1,39 +0,0 @@
-#!/bin/bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-set -euo pipefail
-
-REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-/vllm-workspace/requirements/kv_connectors.txt}"
-
-uv pip install --system -r "${REQUIREMENTS_FILE}"
-
-NIXL_METADATA=$(python3 - <<'PY'
-import importlib.metadata as metadata
-
-import torch
-
-cuda_version = torch.version.cuda
-if cuda_version is None:
-    raise SystemExit("torch.version.cuda is not set")
-
-print(cuda_version.split(".", 1)[0], metadata.version("nixl"))
-PY
-)
-read -r CUDA_MAJOR NIXL_VERSION <<<"${NIXL_METADATA}"
-
-# nixl>=1.1.0 can install multiple CUDA wheel variants. Keep only the variant
-# matching this CI image so nixl_ep_cpp links against the available libcudart.
-uv pip uninstall --system nixl-cu12 nixl-cu13 2>/dev/null || true
-uv pip install --system --no-deps "nixl-cu${CUDA_MAJOR}==${NIXL_VERSION}"
-
-python3 - <<'PY'
-import importlib.metadata as metadata
-
-for package_name in ("nixl", "nixl-cu12", "nixl-cu13"):
-    try:
-        version = metadata.version(package_name)
-    except metadata.PackageNotFoundError:
-        version = "not installed"
-    print(f"{package_name}: {version}")
-PY
@@ -49,7 +49,6 @@ for BACK in "${BACKENDS[@]}"; do
    --data-parallel-size 2 \
    --enable-expert-parallel \
    --enable-eplb \
-    --eplb-config '{"use_async": false}' \
    --trust-remote-code \
    --max-model-len 2048 \
    --all2all-backend "$BACK" \
@@ -48,7 +48,7 @@ for BACK in "${BACKENDS[@]}"; do
    --enforce-eager \
    --enable-eplb \
    --all2all-backend "$BACK" \
-    --eplb-config '{"window_size":10, "step_interval":100, "num_redundant_experts":0, "log_balancedness":true, "use_async":false}' \
+    --eplb-config '{"window_size":10, "step_interval":100, "num_redundant_experts":0, "log_balancedness":true}' \
    --tensor-parallel-size "${TENSOR_PARALLEL_SIZE}" \
    --data-parallel-size "${DATA_PARALLEL_SIZE}" \
    --enable-expert-parallel \
@@ -70,7 +70,7 @@ echo "============================================"
 # ---- Install bfcl-eval if missing ----
 if ! python3 -c "import bfcl_eval" 2>/dev/null; then
    echo "Installing bfcl-eval..."
-    uv pip install "bfcl-eval>=2025.10.20.1,<2026"
+    pip install "bfcl-eval>=2025.10.20.1,<2026"
 fi

 # ---- Cleanup handler ----
@@ -100,7 +100,7 @@ SERVE_ARGS=(
    --tensor-parallel-size "$TP_SIZE"
    --max-model-len "$MAX_MODEL_LEN"
    --enforce-eager
-    --enable-prefix-caching
+    --no-enable-prefix-caching
 )

 # Append reasoning parser if specified
@@ -1238,11 +1238,14 @@ steps:
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
-  - tests/entrypoints/serve
+  - tests/entrypoints/rpc
+  - tests/entrypoints/serve/instrumentator
+  - tests/tool_use
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
-  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
+  - pytest -v -s entrypoints/serve/instrumentator
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
+  - pytest -v -s tool_use

 - label: Entrypoints Integration (API Server openai - Part 1) # TBD
  timeout_in_minutes: 180
@@ -1258,7 +1261,7 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
+  - 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) # TBD
  timeout_in_minutes: 180
@@ -1272,14 +1275,10 @@ steps:
  - vllm/
  - tests/entrypoints/openai
  - tests/entrypoints/test_chat_utils
-  - tests/entrypoints/generate
-  - tests/tool_use
  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/test_chat_utils.py
-  - pytest -v -s entrypoints/generate
-  - pytest -v -s tool_use

 - label: Entrypoints Integration (API Server openai - Part 3) # TBD
  timeout_in_minutes: 180
@@ -1369,7 +1368,7 @@ steps:
  - vllm/platforms/rocm.py
  commands:
  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/offline_mode --ignore=entrypoints/openai --ignore=entrypoints/serve --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text --ignore=tests/entrypoints/generate
+  - 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: OpenAI API correctness # TBD
  timeout_in_minutes: 180
@@ -1485,7 +1484,7 @@ steps:
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt

- label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100-4xMI300) # TBD
+- label: DeepSeek V2-Lite Accuracy (4xH100-4xMI300) # TBD
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_4
@@ -1527,7 +1526,7 @@ steps:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
  - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4

- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100-4xMI300) # TBD
+- label: Qwen3-30B-A3B-FP8-block Accuracy (4xH100-4xMI300) # TBD
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
  agent_pool: mi300_4
@@ -2746,11 +2745,14 @@ steps:
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
-  - tests/entrypoints/serve
+  - tests/entrypoints/rpc
+  - tests/entrypoints/serve/instrumentator
+  - tests/tool_use
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
-  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
+  - pytest -v -s entrypoints/serve/instrumentator
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
+  - pytest -v -s tool_use

 - label: Entrypoints Integration (API Server openai - Part 1) # TBD
  timeout_in_minutes: 180
@@ -2766,7 +2768,7 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
+  - 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) # TBD
  timeout_in_minutes: 180
@@ -2780,14 +2782,10 @@ steps:
  - vllm/
  - tests/entrypoints/openai
  - tests/entrypoints/test_chat_utils
-  - tests/entrypoints/generate
-  - tests/tool_use
  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/test_chat_utils.py
-  - pytest -v -s entrypoints/generate
-  - pytest -v -s tool_use

 - label: Entrypoints Integration (API Server openai - Part 3) # TBD
  timeout_in_minutes: 180
@@ -2897,7 +2895,7 @@ steps:
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt

- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (B200-MI355) # TBD
+- label: Qwen3-30B-A3B-FP8-block Accuracy (B200-MI355) # TBD
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
  agent_pool: mi355_2
@@ -11,7 +11,7 @@ steps:
    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
  key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
@@ -22,7 +22,7 @@ steps:
    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh

 - label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
@@ -34,7 +34,7 @@ steps:
    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - DP_EP=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh

 - label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
@@ -46,7 +46,7 @@ steps:
    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh

 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
@@ -58,7 +58,7 @@ steps:
    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh

 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
@@ -73,7 +73,7 @@ steps:
    - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh

 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
@@ -87,7 +87,7 @@ steps:
    - vllm/v1/worker/kv_connector_model_runner_mixin.py
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh

 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
@@ -102,5 +102,5 @@ steps:
    - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
    - tests/v1/kv_connector/nixl_integration/
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
-    - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
@@ -2,8 +2,8 @@ group: E2E Integration
 depends_on: 
  - image-build
 steps:
- label: DeepSeek V2-Lite Sync EPLB Accuracy
-  key: deepseek-v2-lite-sync-eplb-accuracy
+- label: DeepSeek V2-Lite Accuracy
+  key: deepseek-v2-lite-accuracy
  timeout_in_minutes: 60
  device: h100
  optional: true
@@ -12,8 +12,8 @@ steps:
  commands:
  - bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh 0.25 200 8010

- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy
-  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy
+- label: Qwen3-30B-A3B-FP8-block Accuracy
+  key: qwen3-30b-a3b-fp8-block-accuracy
  timeout_in_minutes: 60
  device: h100
  optional: true
@@ -22,8 +22,8 @@ steps:
  commands:
  - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020

- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (B200)
-  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-b200
+- label: Qwen3-30B-A3B-FP8-block Accuracy (B200)
+  key: qwen3-30b-a3b-fp8-block-accuracy-b200
  timeout_in_minutes: 60
  device: b200-k8s
  optional: true
@@ -11,7 +11,7 @@ steps:
  - tests/entrypoints/
  commands:
  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/offline_mode --ignore=entrypoints/openai --ignore=entrypoints/serve --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text --ignore=tests/entrypoints/generate
+  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text

 - label: Entrypoints Integration (LLM)
  key: entrypoints-integration-llm
@@ -43,7 +43,7 @@ steps:
  - tests/entrypoints/test_chat_utils
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
+  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
  mirror:
    amd:
      device: mi325_1
@@ -60,13 +60,9 @@ steps:
  - vllm/
  - tests/entrypoints/openai
  - tests/entrypoints/test_chat_utils
-  - tests/entrypoints/generate
-  - tests/tool_use
  commands:
  - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
  - pytest -v -s entrypoints/test_chat_utils.py
-  - pytest -v -s entrypoints/generate
-  - pytest -v -s tool_use
  mirror:
    amd:
      device: mi325_1
@@ -102,11 +98,14 @@ steps:
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
-  - tests/entrypoints/serve
+  - tests/entrypoints/rpc
+  - tests/entrypoints/serve/instrumentator
+  - tests/tool_use
  commands:
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
-  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
+  - pytest -v -s entrypoints/serve/instrumentator
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
+  - pytest -v -s tool_use
  mirror:
    amd:
      device: mi325_1
@@ -154,5 +153,6 @@ steps:
  source_file_dependencies:
  - csrc/
  - vllm/entrypoints/openai/
+  - vllm/model_executor/models/whisper.py
  commands: # LMEval
  - pytest -s entrypoints/openai/correctness/
@@ -86,7 +86,7 @@ steps:
    - tests/v1/metrics
    - tests/entrypoints/openai/correctness/test_lmeval.py
  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
    # split the test to avoid interference
    - pytest -v -s -m 'not cpu_test' v1/core
@@ -14,12 +14,5 @@ steps:
  commands:
    - apt-get update && apt-get install -y curl libsodium23
    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-    # Dump tracebacks of all threads if a test hangs, so a wedged GPU/CUDA
-    # init surfaces a stack instead of silently stalling.
-    - export PYTHONFAULTHANDLER=1
-    # Per-test watchdog: a single hung test (e.g. stuck during engine/CUDA
-    # init) fails fast with a traceback instead of running until the global
-    # build timeout. The `thread` method also handles hangs inside C/CUDA
-    # calls that the signal method cannot interrupt.
-    - pytest -v -s model_executor -m '(not slow_test)' --timeout=900 --timeout-method=thread
-    - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
+    - pytest -v -s model_executor -m '(not slow_test)'
+    - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
@@ -16,7 +16,7 @@ steps:
  - tests/benchmarks/test_serve_cli.py
  - tests/entrypoints/openai/chat_completion/test_chat_completion.py
  # - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
-
+  # - tests/entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py
  # - tests/entrypoints/openai/completion/test_prompt_validation.py
  - tests/entrypoints/openai/completion/test_shutdown.py
  # - tests/entrypoints/openai/test_return_token_ids.py
@@ -28,7 +28,7 @@ steps:
  - pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
  - pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py
  # - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
-
+  # - pytest -v -s entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py
  # - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
  - pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
  # - pytest -v -s entrypoints/openai/test_return_token_ids.py
@@ -45,19 +45,19 @@ steps:
  - vllm/entrypoints/serve/
  - vllm/v1/engine/
  - tests/utils.py
-  # - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
+  # - tests/entrypoints/rpc/test_collective_rpc.py
  - tests/entrypoints/serve/disagg/test_serving_tokens.py
  - tests/entrypoints/serve/instrumentator/test_basic.py
  - tests/entrypoints/serve/instrumentator/test_metrics.py
-  # - tests/entrypoints/serve/dev/test_sleep.py
+  # - tests/entrypoints/serve/instrumentator/test_sleep.py
  commands:
  - export VLLM_USE_RUST_FRONTEND=1
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  # - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
+  # - pytest -v -s entrypoints/rpc/test_collective_rpc.py
  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
  - pytest -v -s entrypoints/serve/disagg/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
  - pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
-  # - pytest -v -s entrypoints/serve/dev/test_sleep.py
+  # - pytest -v -s entrypoints/serve/instrumentator/test_sleep.py

 - label: Rust Frontend Core Correctness
  timeout_in_minutes: 30
@@ -40,12 +40,6 @@
 /vllm/entrypoints/chat_utils.py @DarkLight1337
 /vllm/entrypoints/llm.py @DarkLight1337

-# Rust Frontend
-/rust/ @BugenZhao @njhill
-/build_rust.sh @BugenZhao @njhill
-/rust-toolchain.toml @BugenZhao @njhill
-/.buildkite/test_areas/rust* @BugenZhao @njhill
-
 # Input/Output Processing
 /vllm/sampling_params.py @njhill @NickLucche
 /vllm/pooling_params.py @noooop @DarkLight1337
@@ -78,13 +72,11 @@
 /vllm/v1/worker/gpu/kv_connector.py @orozery

 # CI & building
-/.buildkite @Harry-Chen @khluu
-/docker/Dockerfile @Harry-Chen @khluu
-/pyproject.toml @khluu
-/setup.py @khluu
+/.buildkite @Harry-Chen
+/docker/Dockerfile @Harry-Chen

 # Test ownership
-/.buildkite/lm-eval-harness @mgoin
+/.buildkite/lm-eval-harness @mgoin 
 /tests/distributed/test_multi_node_assignment.py @youkaichao
 /tests/distributed/test_pipeline_parallel.py @youkaichao
 /tests/distributed/test_same_node.py @youkaichao
@@ -10,7 +10,7 @@ jobs:
        runs-on: ubuntu-latest
        steps:
            -   name: Add label
-                uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
+                uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
                with:
                    script: |
                        github.rest.issues.addLabels({
@@ -14,7 +14,7 @@ jobs:
    steps:
      - name: Label issues based on keywords
        id: label-step
-        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3  # v9.0.0
+        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd  # v8.0.0
        with:
          script: |
            // Configuration: Add new labels and keywords here
@@ -315,7 +315,7 @@ jobs:

      - name: CC users for labeled issues
        if: steps.label-step.outputs.labels_added != '[]'
-        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3  # v9.0.0
+        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd  # v8.0.0
        with:
          script: |
            // Configuration: Map labels to GitHub users to CC
@@ -392,7 +392,7 @@ jobs:

      - name: Request missing ROCm info from issue author
        if: contains(steps.label-step.outputs.labels_added, 'rocm') && contains(toJSON(github.event.issue.labels.*.name), 'bug')
-        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3  # v9.0.0
+        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd  # v8.0.0
        with:
          script: |
            const body = (context.payload.issue.body || '').toLowerCase();
@@ -12,7 +12,7 @@ jobs:
    runs-on: ubuntu-latest
    steps:
      - name: Update PR description
-        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
+        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
        with:
          script: |
            const { owner, repo } = context.repo;
@@ -55,7 +55,7 @@ jobs:
    runs-on: ubuntu-latest
    steps:
      - name: Post welcome comment for first-time contributors
-        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
+        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
        with:
          script: |
            const { owner, repo } = context.repo;
@@ -20,7 +20,7 @@ jobs:
    runs-on: ubuntu-latest
    steps:
    - name: Check PR label and author merge count
-      uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
+      uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
      with:
        script: |
          const { data: pr } = await github.rest.pulls.get({
@@ -49,7 +49,7 @@ jobs:
    runs-on: ubuntu-latest
    steps:
    - uses: actions/checkout@8e8c483db84b4bee98b60c0593521ed34d9990e8 # v6.0.1
-    - uses: actions/setup-python@a309ff8b426b58ec0e2a45f0f869d46889d02405 # v6.2.0
+    - uses: actions/setup-python@83679a892e2d95755f2dac6acb0bfd1e9ac5d548 # v6.1.0
      with:
        python-version: "3.12"
    - run: echo "::add-matcher::.github/workflows/matchers/actionlint.json"
@@ -21,7 +21,7 @@ repos:
  rev: v21.1.2
  hooks:
  - id: clang-format
-    exclude: 'csrc/(moe/topk_softmax_kernels.cu|libtorch_stable/quantization/gguf/(ggml-common.h|dequantize.cuh|vecdotq.cuh|mmq.cuh|mmvq.cuh))|vllm/third_party/.*'
+    exclude: 'csrc/(moe/topk_softmax_kernels.cu|quantization/gguf/(ggml-common.h|dequantize.cuh|vecdotq.cuh|mmq.cuh|mmvq.cuh))|vllm/third_party/.*'
    types_or: [c++, cuda]
    args: [--style=file, --verbose]
 - repo: https://github.com/DavidAnson/markdownlint-cli2
@@ -9,8 +9,8 @@ build:
    python: "3.12"
  jobs:
    post_checkout:
-      - git fetch origin main --unshallow --no-tags --filter=blob:none || true
      - bash docs/pre_run_check.sh
+      - git fetch origin main --unshallow --no-tags --filter=blob:none || true
    pre_create_environment:
      - pip install uv
    create_environment:
@@ -144,14 +144,14 @@ endif()
 # Set up GPU language and check the torch version and warn if it isn't
 # what is expected.
 #
-if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
+if (NOT HIP_FOUND AND CUDA_FOUND)
  set(VLLM_GPU_LANG "CUDA")

  if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
    message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
      "expected for CUDA build, saw ${Torch_VERSION} instead.")
  endif()
-elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
+elseif(HIP_FOUND)
  set(VLLM_GPU_LANG "HIP")

  # Importing torch recognizes and sets up some HIP/ROCm configuration but does
@@ -305,6 +305,10 @@ endif()
 #

 set(VLLM_EXT_SRC
+  "csrc/cache_kernels.cu"
+  "csrc/cache_kernels_fused.cu"
+  "csrc/attention/paged_attention_v1.cu"
+  "csrc/attention/paged_attention_v2.cu"
  "csrc/cuda_view.cu"
  "csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
  "csrc/quantization/activation_kernels.cu"
@@ -643,11 +647,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    "csrc/libtorch_stable/attention/merge_attn_states.cu"
    "csrc/libtorch_stable/sampler.cu"
    "csrc/libtorch_stable/topk.cu"
-    "csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
-    "csrc/libtorch_stable/attention/paged_attention_v1.cu"
-    "csrc/libtorch_stable/attention/paged_attention_v2.cu"
-    "csrc/libtorch_stable/cache_kernels.cu"
-    "csrc/libtorch_stable/cache_kernels_fused.cu")
+    "csrc/libtorch_stable/mamba/selective_scan_fwd.cu")

  if(VLLM_GPU_LANG STREQUAL "CUDA")
    list(APPEND VLLM_STABLE_EXT_SRC
@@ -683,22 +683,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
                   "in CUDA target architectures.")
  endif()

-  # FP32 router GEMM (H=3072, E=256, M<=32). Requires SM90+ and CUDA >= 12.0.
-  cuda_archs_sm90plus(FP32_ROUTER_GEMM_ARCHS "${CUDA_ARCHS}")
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND FP32_ROUTER_GEMM_ARCHS)
-    set(SRCS
-      "csrc/libtorch_stable/fp32_router_gemm_entry.cu"
-      "csrc/libtorch_stable/fp32_router_gemm.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${FP32_ROUTER_GEMM_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building fp32_router_gemm for archs: ${FP32_ROUTER_GEMM_ARCHS}")
-  else()
-    message(STATUS "Not building fp32_router_gemm as no compatible archs found "
-                   "(requires SM90+ and CUDA >= 12.0).")
-  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)
@@ -940,11 +924,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
      SRCS "${SRCS}"
      CUDA_ARCHS "${FP4_ARCHS}")
    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    set(NVFP4_KV_SRC "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
+    # nvfp4_kv_cache_kernels uses non-stable torch API and is called directly
+    # from cache_kernels.cu, so it belongs in _C rather than _C_stable.
+    set(NVFP4_KV_SRC "csrc/nvfp4_kv_cache_kernels.cu")
    set_gencode_flags_for_srcs(
      SRCS "${NVFP4_KV_SRC}"
      CUDA_ARCHS "${FP4_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${NVFP4_KV_SRC}")
+    target_sources(_C PRIVATE ${NVFP4_KV_SRC})
    target_compile_definitions(_C PRIVATE ENABLE_NVFP4_SM120=1)
    list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM120=1")
    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM120=1")
@@ -974,11 +960,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
      SRCS "${SRCS}"
      CUDA_ARCHS "${FP4_ARCHS}")
    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    set(NVFP4_KV_SRC "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
+    set(NVFP4_KV_SRC "csrc/nvfp4_kv_cache_kernels.cu")
    set_gencode_flags_for_srcs(
      SRCS "${NVFP4_KV_SRC}"
      CUDA_ARCHS "${FP4_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${NVFP4_KV_SRC}")
+    target_sources(_C PRIVATE ${NVFP4_KV_SRC})
    target_compile_definitions(_C PRIVATE ENABLE_NVFP4_SM100=1)
    list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM100=1")
    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM100=1")
@@ -1256,22 +1242,24 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
                   " in CUDA target architectures")
  endif()

-  # DeepSeek V3 router GEMM kernel requires SM90+ and CUDA >= 12.0.
-  # (fp32_router_gemm has been migrated to _C_stable_libtorch above.)
-  cuda_archs_sm90plus(SM90PLUS_ROUTER_GEMM_ARCHS "${CUDA_ARCHS}")
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND SM90PLUS_ROUTER_GEMM_ARCHS)
+  # DeepSeek V3 router GEMM kernel - requires SM90+
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(DSV3_ROUTER_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(DSV3_ROUTER_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_ROUTER_GEMM_ARCHS)
    set(DSV3_ROUTER_GEMM_SRC
      "csrc/moe/dsv3_router_gemm_entry.cu"
      "csrc/moe/dsv3_router_gemm_float_out.cu"
      "csrc/moe/dsv3_router_gemm_bf16_out.cu")
    set_gencode_flags_for_srcs(
      SRCS "${DSV3_ROUTER_GEMM_SRC}"
-      CUDA_ARCHS "${SM90PLUS_ROUTER_GEMM_ARCHS}")
+      CUDA_ARCHS "${DSV3_ROUTER_GEMM_ARCHS}")
    list(APPEND VLLM_MOE_EXT_SRC "${DSV3_ROUTER_GEMM_SRC}")
-
-    message(STATUS "Building DSV3 router GEMM kernels for archs: ${SM90PLUS_ROUTER_GEMM_ARCHS}")
+    message(STATUS "Building DSV3 router GEMM kernel for archs: ${DSV3_ROUTER_GEMM_ARCHS}")
  else()
-    message(STATUS "Not building DSV3 router GEMM kernels as no compatible archs found"
+    message(STATUS "Not building DSV3 router GEMM kernel as no compatible archs found"
                   " (requires SM90+ and CUDA >= 12.0)")
  endif()
 endif()
@@ -1298,14 +1286,6 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
    "csrc/rocm/skinny_gemms.cu"
    "csrc/rocm/attention.cu")

-  set(VLLM_ROCM_HAS_GFX1100 OFF)
-  if(VLLM_GPU_ARCHES MATCHES "gfx1100")
-    set(VLLM_ROCM_HAS_GFX1100 ON)
-    list(APPEND VLLM_ROCM_EXT_SRC
-      "csrc/rocm/q_gemm_rdna3.cu"
-      "csrc/rocm/q_gemm_rdna3_wmma.cu")
-  endif()
-
  define_extension_target(
    _rocm_C
    DESTINATION vllm
@@ -1315,10 +1295,6 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
    ARCHITECTURES ${VLLM_GPU_ARCHES}
    USE_SABI 3
    WITH_SOABI)
-
-  if(VLLM_ROCM_HAS_GFX1100)
-    target_compile_definitions(_rocm_C PRIVATE VLLM_ROCM_GFX1100)
-  endif()
 endif()

 # Must run after the last HIP `define_extension_target` so every extension
@@ -1,154 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import torch
-import torch.nn.functional as F
-
-from vllm import _custom_ops as ops
-from vllm.platforms import current_platform
-from vllm.transformers_utils.config import get_config
-from vllm.triton_utils import triton
-from vllm.utils.argparse_utils import FlexibleArgumentParser
-
-# Dimensions supported by the DSV3 specialized kernel
-DSV3_SUPPORTED_NUM_EXPERTS = [256, 384]
-DSV3_SUPPORTED_HIDDEN_SIZES = [7168]
-
-# Dimensions supported by the gpt-oss specialized kernel
-GPT_OSS_SUPPORTED_NUM_EXPERTS = [32, 128]
-GPT_OSS_SUPPORTED_HIDDEN_SIZES = [2880]
-
-# Dimensions supported by the fp32 specialized kernel (MiniMax-M2)
-FP32_SUPPORTED_NUM_EXPERTS = [256]
-FP32_SUPPORTED_HIDDEN_SIZES = [3072]
-FP32_MAX_TOKENS = 32
-
-
-def get_batch_size_range(max_batch_size):
-    return [2**x for x in range(14) if 2**x <= max_batch_size]
-
-
-def get_model_params(config):
-    if config.architectures[0] in (
-        "DeepseekV2ForCausalLM",
-        "DeepseekV3ForCausalLM",
-        "DeepseekV32ForCausalLM",
-    ):
-        num_experts = config.n_routed_experts
-        hidden_size = config.hidden_size
-    elif config.architectures[0] in ("GptOssForCausalLM",) or config.architectures[
-        0
-    ] in ("MiniMaxM2ForCausalLM",):
-        num_experts = config.num_local_experts
-        hidden_size = config.hidden_size
-    else:
-        raise ValueError(f"Unsupported architecture: {config.architectures}")
-    return num_experts, hidden_size
-
-
-def get_benchmark(model, max_batch_size, trust_remote_code):
-    @triton.testing.perf_report(
-        triton.testing.Benchmark(
-            x_names=["batch_size"],
-            x_vals=get_batch_size_range(max_batch_size),
-            x_log=False,
-            line_arg="provider",
-            line_vals=[
-                "torch",
-                "vllm",
-            ],
-            line_names=["PyTorch", "vLLM"],
-            styles=([("blue", "-"), ("red", "-")]),
-            ylabel="TFLOPs",
-            plot_name=f"{model} router gemm throughput",
-            args={},
-        )
-    )
-    def benchmark(batch_size, provider):
-        config = get_config(model=model, trust_remote_code=trust_remote_code)
-        num_experts, hidden_size = get_model_params(config)
-
-        is_hopper_or_blackwell = current_platform.is_device_capability(
-            90
-        ) or current_platform.is_device_capability_family(100)
-        allow_dsv3_router_gemm = (
-            is_hopper_or_blackwell
-            and num_experts in DSV3_SUPPORTED_NUM_EXPERTS
-            and hidden_size in DSV3_SUPPORTED_HIDDEN_SIZES
-        )
-        allow_gpt_oss_router_gemm = (
-            is_hopper_or_blackwell
-            and num_experts in GPT_OSS_SUPPORTED_NUM_EXPERTS
-            and hidden_size in GPT_OSS_SUPPORTED_HIDDEN_SIZES
-        )
-        is_fp32_router_model = (
-            is_hopper_or_blackwell
-            and num_experts in FP32_SUPPORTED_NUM_EXPERTS
-            and hidden_size in FP32_SUPPORTED_HIDDEN_SIZES
-        )
-        allow_fp32_router_gemm = is_fp32_router_model and batch_size <= FP32_MAX_TOKENS
-
-        # Weight dtype: fp32 kernel requires fp32 weights; others use bf16.
-        weight_dtype = torch.float32 if is_fp32_router_model else torch.bfloat16
-        mat_a = torch.randn(
-            (batch_size, hidden_size), dtype=torch.bfloat16, device="cuda"
-        ).contiguous()
-        mat_b = torch.randn(
-            (num_experts, hidden_size), dtype=weight_dtype, device="cuda"
-        ).contiguous()
-        bias = torch.randn(
-            num_experts, dtype=torch.bfloat16, device="cuda"
-        ).contiguous()
-
-        has_bias = allow_gpt_oss_router_gemm
-
-        quantiles = [0.5, 0.2, 0.8]
-
-        if provider == "torch":
-
-            def runner():
-                if allow_fp32_router_gemm:
-                    F.linear(mat_a.float(), mat_b)
-                elif has_bias:
-                    F.linear(mat_a, mat_b, bias)
-                else:
-                    F.linear(mat_a, mat_b)
-        elif provider == "vllm":
-
-            def runner():
-                if allow_dsv3_router_gemm:
-                    ops.dsv3_router_gemm(mat_a, mat_b, torch.bfloat16)
-                elif allow_fp32_router_gemm:
-                    ops.fp32_router_gemm(mat_a, mat_b)
-                elif allow_gpt_oss_router_gemm:
-                    ops.gpt_oss_router_gemm(mat_a, mat_b, bias)
-                elif is_fp32_router_model:
-                    # batch_size > FP32_MAX_TOKENS: fall back to F.linear
-                    F.linear(mat_a.float(), mat_b)
-                else:
-                    F.linear(mat_a, mat_b)
-
-        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
-            runner, quantiles=quantiles
-        )
-
-        def tflops(t_ms):
-            flops = 2 * batch_size * hidden_size * num_experts
-            return flops / (t_ms * 1e-3) / 1e12
-
-        return tflops(ms), tflops(max_ms), tflops(min_ms)
-
-    return benchmark
-
-
-if __name__ == "__main__":
-    parser = FlexibleArgumentParser()
-    parser.add_argument("--model", type=str, default="openai/gpt-oss-20b")
-    parser.add_argument("--max-batch-size", default=16, type=int)
-    parser.add_argument("--trust-remote-code", action="store_true")
-    args = parser.parse_args()
-
-    # Get the benchmark function
-    benchmark = get_benchmark(args.model, args.max_batch_size, args.trust_remote_code)
-    # Run performance benchmark
-    benchmark.run(print_data=True)
@@ -369,18 +369,6 @@ else()
    add_compile_definitions(-DVLLM_NUMA_DISABLED)
 endif()

-# check if the pytorch wheel ships libopenblas.so.
-set(VLLM_OPENBLAS_LIB "")
-if (NOT ENABLE_X86_ISA)
-    file(GLOB _VLLM_TORCH_OPENBLAS_LIBS
-        "${TORCH_INSTALL_PREFIX}/lib/libopenblas*.so*")
-    # Note: we don't link openblas directly to _C extension, as it's available through libtorch.so 
-    if (_VLLM_TORCH_OPENBLAS_LIBS)
-        list(GET _VLLM_TORCH_OPENBLAS_LIBS 0 VLLM_OPENBLAS_LIB)
-        message(STATUS "CPU OpenBLAS library: ${VLLM_OPENBLAS_LIB}")
-    endif()
-endif()
-
 #
 # Generate CPU attention dispatch header
 #
@@ -399,7 +387,6 @@ endif()
 #
 set(VLLM_EXT_SRC
    "csrc/cpu/activation.cpp"
-    "csrc/cpu/sgl-kernels/fla.cpp"
    "csrc/cpu/utils.cpp"
    "csrc/cpu/spec_decode_utils.cpp"
    "csrc/cpu/layernorm.cpp"
@@ -409,13 +396,6 @@ set(VLLM_EXT_SRC
    "csrc/cpu/cpu_attn.cpp"
    "csrc/cpu/torch_bindings.cpp")

-if (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64" AND VLLM_RVV_VLEN AND
-        VLLM_RVV_VLEN GREATER 0 AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
-    set(VLLM_EXT_SRC
-        "csrc/cpu/cpu_wna16.cpp"
-        ${VLLM_EXT_SRC})
-endif()
-
 if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
    set(VLLM_EXT_SRC
        "csrc/cpu/shm.cpp"
@@ -423,12 +403,6 @@ if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
        ${VLLM_EXT_SRC})
 endif()

-if (POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND)	
-    set(VLLM_EXT_SRC
-        "csrc/cpu/shm.cpp"
-        ${VLLM_EXT_SRC})
-endif()
-
 if(USE_ONEDNN)
    set(VLLM_EXT_SRC
        "csrc/cpu/dnnl_kernels.cpp"
@@ -437,6 +411,7 @@ endif()

 if (ENABLE_X86_ISA)
    set(VLLM_EXT_SRC_SGL
+        "csrc/cpu/sgl-kernels/fla.cpp"
        "csrc/cpu/sgl-kernels/conv.cpp"
        "csrc/cpu/sgl-kernels/gemm.cpp"
        "csrc/cpu/sgl-kernels/gemm_int8.cpp"
@@ -448,7 +423,6 @@ if (ENABLE_X86_ISA)
        "csrc/cpu/sgl-kernels/moe_fp8.cpp")

    set(VLLM_EXT_SRC_AVX512
-        "csrc/cpu/sgl-kernels/fla.cpp"
        "csrc/cpu/shm.cpp"
        "csrc/cpu/cpu_wna16.cpp"
        "csrc/cpu/cpu_fused_moe.cpp"
@@ -465,7 +439,6 @@ if (ENABLE_X86_ISA)
        "csrc/moe/dynamic_4bit_int_moe_cpu.cpp") 

    set(VLLM_EXT_SRC_AVX2
-        "csrc/cpu/sgl-kernels/fla.cpp"
        "csrc/cpu/utils.cpp"
        "csrc/cpu/spec_decode_utils.cpp"
        "csrc/cpu/cpu_attn.cpp"
@@ -539,9 +512,6 @@ else()
        USE_SABI 3
        WITH_SOABI
    )
-    if (VLLM_OPENBLAS_LIB)
-        target_compile_definitions(_C PRIVATE VLLM_HAS_OPENBLAS)
-    endif()
 endif()

 message(STATUS "Enabling C extension.")
@@ -31,7 +31,7 @@ endif()

 if(VLLM_FLASH_ATTN_SRC_DIR)
  FetchContent_Declare(
-          vllm-flash-attn SOURCE_DIR
+          vllm-flash-attn SOURCE_DIR 
          ${VLLM_FLASH_ATTN_SRC_DIR}
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
  )
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG dd62dac706b1cf7895bd99b18c6cb7e7e117ee25
+          GIT_TAG bce29425653ec0fbc579d329883030e832d15ada
          GIT_PROGRESS TRUE
          # Don't share the vllm-flash-attn build between build types
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
@@ -476,16 +476,6 @@ function(cuda_archs_loose_intersection OUT_CUDA_ARCHS SRC_CUDA_ARCHS TGT_CUDA_AR
  set(${OUT_CUDA_ARCHS} ${_CUDA_ARCHS} PARENT_SCOPE)
 endfunction()

-
-function(cuda_archs_sm90plus OUT_CUDA_ARCHS TGT_CUDA_ARCHS)
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(_archs "9.0a;10.0f;11.0f" "${TGT_CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(_archs "9.0a;10.0a;10.1a;10.3a" "${TGT_CUDA_ARCHS}")
-  endif()
-  set(${OUT_CUDA_ARCHS} ${_archs} PARENT_SCOPE)
-endfunction()
-
 #
 # Override the GPU architectures detected by cmake/torch and filter them by
 # `GPU_SUPPORTED_ARCHES`. Sets the final set of architectures in
@@ -17,18 +17,21 @@
 * limitations under the License.
 */

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

-#include "../../attention/attention_dtypes.h"
+#include "attention_dtypes.h"
 #include "attention_utils.cuh"
-#include "../../cuda_compat.h"
+#include "../cuda_compat.h"

 #ifdef USE_ROCM
  #include <hip/hip_bf16.h>
-  #include "../../quantization/w8a8/fp8/amd/quant_utils.cuh"
+  #include "../quantization/w8a8/fp8/amd/quant_utils.cuh"
 typedef __hip_bfloat16 __nv_bfloat16;
 #else
-  #include "../../quantization/w8a8/fp8/nvidia/quant_utils.cuh"
+  #include "../quantization/w8a8/fp8/nvidia/quant_utils.cuh"
 #endif

 #define MAX(a, b) ((a) > (b) ? (a) : (b))
@@ -18,8 +18,8 @@
 */
 #pragma once

-#include "../../cuda_compat.h"
-#include "../../attention/attention_dtypes.h"
+#include "../cuda_compat.h"
+#include "attention_dtypes.h"

 #include <float.h>
 #include <type_traits>
@@ -16,9 +16,8 @@
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */
-#include "../torch_utils.h"
 #include "attention_kernels.cuh"
-#include "../../cuda_compat.h"
+#include "../cuda_compat.h"

 #define MAX(a, b) ((a) > (b) ? (a) : (b))
 #define MIN(a, b) ((a) < (b) ? (a) : (b))
@@ -45,15 +44,13 @@ template <typename T, typename CACHE_T, int BLOCK_SIZE,
          vllm::Fp8KVCacheDataType KV_DTYPE, bool IS_BLOCK_SPARSE,
          int NUM_THREADS = 128>
 void paged_attention_v1_launcher(
-    torch::stable::Tensor& out, torch::stable::Tensor& query,
-    torch::stable::Tensor& key_cache, torch::stable::Tensor& value_cache,
-    int num_kv_heads, float scale, torch::stable::Tensor& block_tables,
-    torch::stable::Tensor& seq_lens, int max_seq_len,
-    const std::optional<torch::stable::Tensor>& alibi_slopes,
-    torch::stable::Tensor& k_scale, torch::stable::Tensor& v_scale,
-    const int tp_rank, const int blocksparse_local_blocks,
-    const int blocksparse_vert_stride, const int blocksparse_block_size,
-    const int blocksparse_head_sliding_step) {
+    torch::Tensor& out, torch::Tensor& query, torch::Tensor& key_cache,
+    torch::Tensor& value_cache, int num_kv_heads, float scale,
+    torch::Tensor& block_tables, torch::Tensor& seq_lens, int max_seq_len,
+    const std::optional<torch::Tensor>& alibi_slopes, torch::Tensor& k_scale,
+    torch::Tensor& v_scale, const int tp_rank,
+    const int blocksparse_local_blocks, const int blocksparse_vert_stride,
+    const int blocksparse_block_size, const int blocksparse_head_sliding_step) {
  int num_seqs = query.size(0);
  int num_heads = query.size(1);
  int head_size = query.size(2);
@@ -72,8 +69,8 @@ void paged_attention_v1_launcher(
  T* query_ptr = reinterpret_cast<T*>(query.data_ptr());
  CACHE_T* key_cache_ptr = reinterpret_cast<CACHE_T*>(key_cache.data_ptr());
  CACHE_T* value_cache_ptr = reinterpret_cast<CACHE_T*>(value_cache.data_ptr());
-  int* block_tables_ptr = block_tables.mutable_data_ptr<int>();
-  int* seq_lens_ptr = seq_lens.mutable_data_ptr<int>();
+  int* block_tables_ptr = block_tables.data_ptr<int>();
+  int* seq_lens_ptr = seq_lens.data_ptr<int>();
  const float* k_scale_ptr = reinterpret_cast<const float*>(k_scale.data_ptr());
  const float* v_scale_ptr = reinterpret_cast<const float*>(v_scale.data_ptr());

@@ -88,9 +85,8 @@ void paged_attention_v1_launcher(

  dim3 grid(num_heads, num_seqs, 1);
  dim3 block(NUM_THREADS);
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      query.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(query));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  switch (head_size) {
    // NOTE(woosuk): To reduce the compilation time, we only compile for the
    // head sizes that we use in the model. However, we can easily extend this
@@ -123,7 +119,7 @@ void paged_attention_v1_launcher(
      LAUNCH_PAGED_ATTENTION_V1(256);
      break;
    default:
-      STD_TORCH_CHECK(false, "Unsupported head size: ", head_size);
+      TORCH_CHECK(false, "Unsupported head size: ", head_size);
      break;
  }
 }
@@ -145,43 +141,43 @@ void paged_attention_v1_launcher(

 // NOTE(woosuk): To reduce the compilation time, we omitted block sizes
 // 1, 2, 4, 64, 128, 256.
-#define CALL_V1_LAUNCHER_BLOCK_SIZE(T, CACHE_T, KV_DTYPE)             \
-  switch (block_size) {                                               \
-    case 8:                                                           \
-      CALL_V1_LAUNCHER_SPARSITY(T, CACHE_T, 8, KV_DTYPE);             \
-      break;                                                          \
-    case 16:                                                          \
-      CALL_V1_LAUNCHER_SPARSITY(T, CACHE_T, 16, KV_DTYPE);            \
-      break;                                                          \
-    case 32:                                                          \
-      CALL_V1_LAUNCHER_SPARSITY(T, CACHE_T, 32, KV_DTYPE);            \
-      break;                                                          \
-    default:                                                          \
-      STD_TORCH_CHECK(false, "Unsupported block size: ", block_size); \
-      break;                                                          \
+#define CALL_V1_LAUNCHER_BLOCK_SIZE(T, CACHE_T, KV_DTYPE)         \
+  switch (block_size) {                                           \
+    case 8:                                                       \
+      CALL_V1_LAUNCHER_SPARSITY(T, CACHE_T, 8, KV_DTYPE);         \
+      break;                                                      \
+    case 16:                                                      \
+      CALL_V1_LAUNCHER_SPARSITY(T, CACHE_T, 16, KV_DTYPE);        \
+      break;                                                      \
+    case 32:                                                      \
+      CALL_V1_LAUNCHER_SPARSITY(T, CACHE_T, 32, KV_DTYPE);        \
+      break;                                                      \
+    default:                                                      \
+      TORCH_CHECK(false, "Unsupported block size: ", block_size); \
+      break;                                                      \
  }

 void paged_attention_v1(
-    torch::stable::Tensor& out,    // [num_seqs, num_heads, head_size]
-    torch::stable::Tensor& query,  // [num_seqs, num_heads, head_size]
-    torch::stable::Tensor&
+    torch::Tensor& out,    // [num_seqs, num_heads, head_size]
+    torch::Tensor& query,  // [num_seqs, num_heads, head_size]
+    torch::Tensor&
        key_cache,  // [num_blocks, num_heads, head_size/x, block_size, x]
-    torch::stable::Tensor&
+    torch::Tensor&
        value_cache,       // [num_blocks, num_heads, head_size, block_size]
    int64_t num_kv_heads,  // [num_heads]
    double scale,
-    torch::stable::Tensor& block_tables,  // [num_seqs, max_num_blocks_per_seq]
-    torch::stable::Tensor& seq_lens,      // [num_seqs]
+    torch::Tensor& block_tables,  // [num_seqs, max_num_blocks_per_seq]
+    torch::Tensor& seq_lens,      // [num_seqs]
    int64_t block_size, int64_t max_seq_len,
-    const std::optional<torch::stable::Tensor>& alibi_slopes,
-    const std::string& kv_cache_dtype, torch::stable::Tensor& k_scale,
-    torch::stable::Tensor& v_scale, const int64_t tp_rank,
+    const std::optional<torch::Tensor>& alibi_slopes,
+    const std::string& kv_cache_dtype, torch::Tensor& k_scale,
+    torch::Tensor& v_scale, const int64_t tp_rank,
    const int64_t blocksparse_local_blocks,
    const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
    const int64_t blocksparse_head_sliding_step) {
  const bool is_block_sparse = (blocksparse_vert_stride > 1);

-  DISPATCH_BY_KV_CACHE_DTYPE(query.scalar_type(), kv_cache_dtype,
+  DISPATCH_BY_KV_CACHE_DTYPE(query.dtype(), kv_cache_dtype,
                             CALL_V1_LAUNCHER_BLOCK_SIZE)
 }

@@ -16,9 +16,8 @@
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */
-#include "../torch_utils.h"
 #include "attention_kernels.cuh"
-#include "../../cuda_compat.h"
+#include "../cuda_compat.h"

 #define MAX(a, b) ((a) > (b) ? (a) : (b))
 #define MIN(a, b) ((a) < (b) ? (a) : (b))
@@ -45,16 +44,14 @@ template <typename T, typename CACHE_T, int BLOCK_SIZE,
          vllm::Fp8KVCacheDataType KV_DTYPE, bool IS_BLOCK_SPARSE,
          int NUM_THREADS = 128, int PARTITION_SIZE = 512>
 void paged_attention_v2_launcher(
-    torch::stable::Tensor& out, torch::stable::Tensor& exp_sums,
-    torch::stable::Tensor& max_logits, torch::stable::Tensor& tmp_out,
-    torch::stable::Tensor& query, torch::stable::Tensor& key_cache,
-    torch::stable::Tensor& value_cache, int num_kv_heads, float scale,
-    torch::stable::Tensor& block_tables, torch::stable::Tensor& seq_lens,
-    int max_seq_len, const std::optional<torch::stable::Tensor>& alibi_slopes,
-    torch::stable::Tensor& k_scale, torch::stable::Tensor& v_scale,
-    const int tp_rank, const int blocksparse_local_blocks,
-    const int blocksparse_vert_stride, const int blocksparse_block_size,
-    const int blocksparse_head_sliding_step) {
+    torch::Tensor& out, torch::Tensor& exp_sums, torch::Tensor& max_logits,
+    torch::Tensor& tmp_out, torch::Tensor& query, torch::Tensor& key_cache,
+    torch::Tensor& value_cache, int num_kv_heads, float scale,
+    torch::Tensor& block_tables, torch::Tensor& seq_lens, int max_seq_len,
+    const std::optional<torch::Tensor>& alibi_slopes, torch::Tensor& k_scale,
+    torch::Tensor& v_scale, const int tp_rank,
+    const int blocksparse_local_blocks, const int blocksparse_vert_stride,
+    const int blocksparse_block_size, const int blocksparse_head_sliding_step) {
  int num_seqs = query.size(0);
  int num_heads = query.size(1);
  int head_size = query.size(2);
@@ -76,8 +73,8 @@ void paged_attention_v2_launcher(
  T* query_ptr = reinterpret_cast<T*>(query.data_ptr());
  CACHE_T* key_cache_ptr = reinterpret_cast<CACHE_T*>(key_cache.data_ptr());
  CACHE_T* value_cache_ptr = reinterpret_cast<CACHE_T*>(value_cache.data_ptr());
-  int* block_tables_ptr = block_tables.mutable_data_ptr<int>();
-  int* seq_lens_ptr = seq_lens.mutable_data_ptr<int>();
+  int* block_tables_ptr = block_tables.data_ptr<int>();
+  int* seq_lens_ptr = seq_lens.data_ptr<int>();
  const float* k_scale_ptr = reinterpret_cast<const float*>(k_scale.data_ptr());
  const float* v_scale_ptr = reinterpret_cast<const float*>(v_scale.data_ptr());

@@ -94,9 +91,8 @@ void paged_attention_v2_launcher(
  int reduce_shared_mem_size = 2 * max_num_partitions * sizeof(float);

  dim3 block(NUM_THREADS);
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      query.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(query));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  switch (head_size) {
    // NOTE(woosuk): To reduce the compilation time, we only compile for the
    // head sizes that we use in the model. However, we can easily extend this
@@ -129,7 +125,7 @@ void paged_attention_v2_launcher(
      LAUNCH_PAGED_ATTENTION_V2(256);
      break;
    default:
-      STD_TORCH_CHECK(false, "Unsupported head size: ", head_size);
+      TORCH_CHECK(false, "Unsupported head size: ", head_size);
      break;
  }
 }
@@ -152,48 +148,46 @@ void paged_attention_v2_launcher(

 // NOTE(woosuk): To reduce the compilation time, we omitted block sizes
 // 1, 2, 4, 64, 128, 256.
-#define CALL_V2_LAUNCHER_BLOCK_SIZE(T, CACHE_T, KV_DTYPE)             \
-  switch (block_size) {                                               \
-    case 8:                                                           \
-      CALL_V2_LAUNCHER_SPARSITY(T, CACHE_T, 8, KV_DTYPE);             \
-      break;                                                          \
-    case 16:                                                          \
-      CALL_V2_LAUNCHER_SPARSITY(T, CACHE_T, 16, KV_DTYPE);            \
-      break;                                                          \
-    case 32:                                                          \
-      CALL_V2_LAUNCHER_SPARSITY(T, CACHE_T, 32, KV_DTYPE);            \
-      break;                                                          \
-    default:                                                          \
-      STD_TORCH_CHECK(false, "Unsupported block size: ", block_size); \
-      break;                                                          \
+#define CALL_V2_LAUNCHER_BLOCK_SIZE(T, CACHE_T, KV_DTYPE)         \
+  switch (block_size) {                                           \
+    case 8:                                                       \
+      CALL_V2_LAUNCHER_SPARSITY(T, CACHE_T, 8, KV_DTYPE);         \
+      break;                                                      \
+    case 16:                                                      \
+      CALL_V2_LAUNCHER_SPARSITY(T, CACHE_T, 16, KV_DTYPE);        \
+      break;                                                      \
+    case 32:                                                      \
+      CALL_V2_LAUNCHER_SPARSITY(T, CACHE_T, 32, KV_DTYPE);        \
+      break;                                                      \
+    default:                                                      \
+      TORCH_CHECK(false, "Unsupported block size: ", block_size); \
+      break;                                                      \
  }

 void paged_attention_v2(
-    torch::stable::Tensor& out,  // [num_seqs, num_heads, head_size]
-    torch::stable::Tensor&
-        exp_sums,  // [num_seqs, num_heads, max_num_partitions]
-    torch::stable::Tensor&
-        max_logits,  // [num_seqs, num_heads, max_num_partitions]
-    torch::stable::Tensor&
+    torch::Tensor& out,         // [num_seqs, num_heads, head_size]
+    torch::Tensor& exp_sums,    // [num_seqs, num_heads, max_num_partitions]
+    torch::Tensor& max_logits,  // [num_seqs, num_heads, max_num_partitions]
+    torch::Tensor&
        tmp_out,  // [num_seqs, num_heads, max_num_partitions, head_size]
-    torch::stable::Tensor& query,  // [num_seqs, num_heads, head_size]
-    torch::stable::Tensor&
+    torch::Tensor& query,  // [num_seqs, num_heads, head_size]
+    torch::Tensor&
        key_cache,  // [num_blocks, num_heads, head_size/x, block_size, x]
-    torch::stable::Tensor&
+    torch::Tensor&
        value_cache,       // [num_blocks, num_heads, head_size, block_size]
    int64_t num_kv_heads,  // [num_heads]
    double scale,
-    torch::stable::Tensor& block_tables,  // [num_seqs, max_num_blocks_per_seq]
-    torch::stable::Tensor& seq_lens,      // [num_seqs]
+    torch::Tensor& block_tables,  // [num_seqs, max_num_blocks_per_seq]
+    torch::Tensor& seq_lens,      // [num_seqs]
    int64_t block_size, int64_t max_seq_len,
-    const std::optional<torch::stable::Tensor>& alibi_slopes,
-    const std::string& kv_cache_dtype, torch::stable::Tensor& k_scale,
-    torch::stable::Tensor& v_scale, const int64_t tp_rank,
+    const std::optional<torch::Tensor>& alibi_slopes,
+    const std::string& kv_cache_dtype, torch::Tensor& k_scale,
+    torch::Tensor& v_scale, const int64_t tp_rank,
    const int64_t blocksparse_local_blocks,
    const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
    const int64_t blocksparse_head_sliding_step) {
  const bool is_block_sparse = (blocksparse_vert_stride > 1);
-  DISPATCH_BY_KV_CACHE_DTYPE(query.scalar_type(), kv_cache_dtype,
+  DISPATCH_BY_KV_CACHE_DTYPE(query.dtype(), kv_cache_dtype,
                             CALL_V2_LAUNCHER_BLOCK_SIZE)
 }

@@ -1,13 +1,15 @@
-#include "torch_utils.h"
+#include <torch/all.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
+
+#include "cuda_compat.h"
 #include "dispatch_utils.h"

-#include "../cuda_compat.h"
-
-#include "../quantization/w8a8/fp8/common.cuh"
+#include "quantization/w8a8/fp8/common.cuh"
 #ifdef USE_ROCM
-  #include "../quantization/w8a8/fp8/amd/quant_utils.cuh"
+  #include "quantization/w8a8/fp8/amd/quant_utils.cuh"
 #else
-  #include "../quantization/w8a8/fp8/nvidia/quant_utils.cuh"
+  #include "quantization/w8a8/fp8/nvidia/quant_utils.cuh"
 #endif

 #ifdef USE_ROCM
@@ -162,52 +164,43 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(

 }  // namespace vllm

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

 // Executes RoPE on q_pe and k_pe, then writes k_pe and kv_c in the kv cache.
@@ -215,69 +208,64 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
 // Replaces DeepseekScalingRotaryEmbedding.self.rotary_emb and
 // concat_and_cache_mla.
 void concat_and_cache_mla_rope_fused(
-    torch::stable::Tensor& positions,  // [num_tokens]
-    torch::stable::Tensor& q_pe,       // [num_tokens, num_q_heads, rot_dim]
-    torch::stable::Tensor& k_pe,       // [num_tokens, rot_dim]
-    torch::stable::Tensor& kv_c,       // [num_tokens, kv_lora_rank]
-    torch::stable::Tensor& rope_cos_sin_cache,  // [max_position, rot_dim]
+    torch::Tensor& positions,           // [num_tokens]
+    torch::Tensor& q_pe,                // [num_tokens, num_q_heads, rot_dim]
+    torch::Tensor& k_pe,                // [num_tokens, rot_dim]
+    torch::Tensor& kv_c,                // [num_tokens, kv_lora_rank]
+    torch::Tensor& rope_cos_sin_cache,  // [max_position, rot_dim]
    bool rope_is_neox,
-    torch::stable::Tensor& slot_mapping,  // [num_tokens] or [num_actual_tokens]
-    torch::stable::Tensor&
+    torch::Tensor& slot_mapping,  // [num_tokens] or [num_actual_tokens]
+    torch::Tensor&
        kv_cache,  // [num_blocks, block_size, (kv_lora_rank + rot_dim)]
-    const std::string& kv_cache_dtype,
-    torch::stable::Tensor& kv_cache_quant_scale) {
+    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
+  // 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.
-  const int64_t num_tokens = slot_mapping.size(0);
-  const int64_t num_padded_tokens = q_pe.size(0);
-  STD_TORCH_CHECK(num_padded_tokens >= num_tokens);
+  // For compatibility with both cases, we use slot_mapping.size(0) as the
+  // number of tokens.
+  int num_tokens = slot_mapping.size(0);
+  int num_padded_tokens = q_pe.size(0);
+  TORCH_CHECK_GE(num_padded_tokens, num_tokens);

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

-  STD_TORCH_CHECK(positions.size(0) == num_padded_tokens);
-  STD_TORCH_CHECK(positions.dim() == 1);
-  STD_TORCH_CHECK(positions.scalar_type() ==
-                  torch::headeronly::ScalarType::Long);
+  TORCH_CHECK_EQ(positions.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(positions.dim(), 1);
+  TORCH_CHECK_EQ(positions.scalar_type(), c10::ScalarType::Long);

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

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

-  STD_TORCH_CHECK(kv_c.dim() == 2);
-  STD_TORCH_CHECK(kv_c.size(0) == num_padded_tokens);
-  STD_TORCH_CHECK(kv_c.size(1) == kv_lora_rank);
-  STD_TORCH_CHECK(kv_c.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(1), kv_lora_rank);
+  TORCH_CHECK_EQ(kv_c.scalar_type(), q_pe.scalar_type());
+  TORCH_CHECK_EQ(kv_c.dtype(), q_pe.dtype());

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

-  STD_TORCH_CHECK(slot_mapping.size(0) == num_tokens);
-  STD_TORCH_CHECK(slot_mapping.scalar_type() ==
-                  torch::headeronly::ScalarType::Long);
+  TORCH_CHECK_EQ(slot_mapping.size(0), num_tokens);
+  TORCH_CHECK_EQ(slot_mapping.scalar_type(), c10::ScalarType::Long);

-  STD_TORCH_CHECK(kv_cache.size(2) == kv_lora_rank + rot_dim);
-  STD_TORCH_CHECK(kv_cache.dim() == 3);
+  TORCH_CHECK_EQ(kv_cache.size(2), kv_lora_rank + rot_dim);
+  TORCH_CHECK_EQ(kv_cache.dim(), 3);

-  STD_TORCH_CHECK(kv_cache_quant_scale.numel() == 1);
-  STD_TORCH_CHECK(kv_cache_quant_scale.scalar_type() ==
-                  torch::headeronly::ScalarType::Float);
+  TORCH_CHECK_EQ(kv_cache_quant_scale.numel(), 1);
+  TORCH_CHECK_EQ(kv_cache_quant_scale.scalar_type(), c10::ScalarType::Float);

  int64_t q_pe_stride_token = q_pe.stride(0);
  int64_t q_pe_stride_head = q_pe.stride(1);
@@ -298,10 +286,9 @@ void concat_and_cache_mla_rope_fused(
  dim3 grid(num_tokens, 1, 1);
  dim3 block(thread_block_size, 1, 1);

-  const torch::stable::accelerator::DeviceGuard device_guard(
-      positions.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(positions));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

-  DISPATCH_BY_KV_CACHE_DTYPE(kv_c.scalar_type(), kv_cache_dtype,
+  DISPATCH_BY_KV_CACHE_DTYPE(kv_c.dtype(), kv_cache_dtype,
                             CALL_CONCAT_AND_CACHE_MLA_ROPE_FUSED);
 }
@@ -94,10 +94,6 @@ struct FP16Vec16 : public Vec<FP16Vec16> {
      : reg(RVVI(__riscv_vle16_v_f16, LMUL_256)(
            static_cast<const _Float16*>(ptr), VEC_ELEM_NUM)) {};

-  explicit FP16Vec16(const c10::Half v)
-      : reg(RVVI4(__riscv_vreinterpret_v_u16, LMUL_256, _f16, LMUL_256)(
-            RVVI(__riscv_vmv_v_x_u16, LMUL_256)(v.x, VEC_ELEM_NUM))) {};
-
  explicit FP16Vec16(const FP32Vec16& vec);

  void save(void* ptr) const {
@@ -169,9 +165,6 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
            reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};

  explicit BF16Vec16(fixed_bf16x16_t data) : reg(data) {};
-  explicit BF16Vec16(const c10::BFloat16 v)
-      : reg(RVVI4(__riscv_vreinterpret_v_u16, LMUL_256, _bf16, LMUL_256)(
-            RVVI(__riscv_vmv_v_x_u16, LMUL_256)(v.x, VEC_ELEM_NUM))) {};
  explicit BF16Vec16(const FP32Vec16&);

  void save(void* ptr) const {
@@ -297,9 +290,6 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
    }
    reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_512)(tmp, 16);
  }
-  explicit BF16Vec16(const c10::BFloat16 v)
-      : reg_fp32(RVVI(__riscv_vfmv_v_f_f32, LMUL_512)(static_cast<float>(v),
-                                                      VEC_ELEM_NUM)) {}
  explicit BF16Vec16(const FP32Vec16&);
  void save(void* ptr) const {
    float tmp[16];
@@ -639,19 +629,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
      : reg(RVVI4(__riscv_vcreate_v_f32, LMUL_256, _f32, LMUL_512)(
            data.reg, data.reg)) {};
  explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
-  explicit FP32Vec16(int64_t value, const FP32Vec16& lut) {
-    const uint64_t q_values = static_cast<uint64_t>(value);
-    auto packed = RVVI(__riscv_vmv_v_x_u64, LMUL_1024)(q_values, VEC_ELEM_NUM);
-    auto lane_ids = RVVI(__riscv_vid_v_u64, LMUL_1024)(VEC_ELEM_NUM);
-    auto shifts =
-        RVVI(__riscv_vsll_vx_u64, LMUL_1024)(lane_ids, 2, VEC_ELEM_NUM);
-    auto shifted =
-        RVVI(__riscv_vsrl_vv_u64, LMUL_1024)(packed, shifts, VEC_ELEM_NUM);
-    auto idx64 =
-        RVVI(__riscv_vand_vx_u64, LMUL_1024)(shifted, 0xF, VEC_ELEM_NUM);
-    auto idx32 = RVVI(__riscv_vnsrl_wx_u32, LMUL_512)(idx64, 0, VEC_ELEM_NUM);
-    reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx32, VEC_ELEM_NUM);
-  }
  explicit FP32Vec16(const FP16Vec16& v);

 #ifdef __riscv_zvfbfmin
@@ -664,10 +641,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
  explicit FP32Vec16(const BF16Vec16& v) : reg(v.reg_fp32) {};
 #endif

-  // FP8 stub: dead code on RISC-V (fp8 KV cache is x86-only), needed for
-  // load_b_pair_vec template to compile on all platforms.
-  explicit FP32Vec16(const BF16Vec32&, int) : FP32Vec16() {}
-
  FP32Vec16 operator+(const FP32Vec16& b) const {
    return FP32Vec16(
        RVVI(__riscv_vfadd_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
@@ -918,30 +891,6 @@ inline void fma(FP32Vec16& acc, const FP32Vec16& a, const FP32Vec16& b) {
  acc = acc.fma(a, b);
 }

-template <typename VecT>
-static void interleave_save_16b(const VecT& vec0, const VecT& vec1, void* ptr) {
-  alignas(64) uint16_t values0[VecT::VEC_ELEM_NUM];
-  alignas(64) uint16_t values1[VecT::VEC_ELEM_NUM];
-  vec0.save(values0);
-  vec1.save(values1);
-
-  auto* packed = reinterpret_cast<uint32_t*>(ptr);
-  for (int32_t i = 0; i < VecT::VEC_ELEM_NUM; ++i) {
-    packed[i] = static_cast<uint32_t>(values0[i]) |
-                (static_cast<uint32_t>(values1[i]) << 16);
-  }
-}
-
-static void interleave_save(const FP16Vec16& vec0, const FP16Vec16& vec1,
-                            void* ptr) {
-  interleave_save_16b(vec0, vec1, ptr);
-}
-
-static void interleave_save(const BF16Vec16& vec0, const BF16Vec16& vec1,
-                            void* ptr) {
-  interleave_save_16b(vec0, vec1, ptr);
-}
-
 #ifdef __riscv_zvfbfmin
 template <>
 inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
@@ -89,35 +89,6 @@ struct BF16Vec8 : public Vec<BF16Vec8> {
  }
 };

-struct FP16Vec16 : public Vec<FP16Vec16> {
-  constexpr static int VEC_ELEM_NUM = 16;
-  ss16x8x2_t reg;
-
-  explicit FP16Vec16(const void* ptr) {
-    reg.val[0] = (__vector signed short)vec_xl(0, (signed short*)ptr);
-    reg.val[1] = (__vector signed short)vec_xl(16, (signed short*)ptr);
-  }
-
-  explicit FP16Vec16(bool, const void* ptr) : FP16Vec16(ptr) {}
-
-  explicit FP16Vec16(const FP32Vec16&);
-
-  void save(void* ptr) const {
-    vec_xst(reg.val[0], 0, (signed short*)ptr);
-    vec_xst(reg.val[1], 16, (signed short*)ptr);
-  }
-
-  void save(void* ptr, int elem_num) const {
-    int num = std::max(0, std::min(elem_num, VEC_ELEM_NUM));
-    if (num <= 8) {
-      vec_xst_len(reg.val[0], (signed short*)ptr, num * 2);
-    } else {
-      vec_xst(reg.val[0], 0, (signed short*)ptr);
-      vec_xst_len(reg.val[1], (signed short*)ptr + 8, (num - 8) * 2);
-    }
-  }
-};
-
 struct BF16Vec16 : public Vec<BF16Vec16> {
  constexpr static int VEC_ELEM_NUM = 16;

@@ -129,8 +100,6 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
    reg.val[1] = (__vector signed short)vec_xl(16, (signed short*)ptr);
  }

-  explicit BF16Vec16(bool, const void* ptr) : BF16Vec16(ptr) {}
-
  explicit BF16Vec16(const FP32Vec16&);

  void save(void* ptr) const {
@@ -410,8 +379,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
    reg.val[3] = vec_xl(48, ptr);
  }

-  explicit FP32Vec16(bool, const float* ptr) : FP32Vec16(ptr) {}
-
  explicit FP32Vec16(f32x4x4_t data) : reg(data) {}

  explicit FP32Vec16(const FP32Vec16& data) {
@@ -435,7 +402,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
    reg.val[3] = data.reg.val[1];
  }

-  explicit FP32Vec16(const FP16Vec16& v);
  explicit FP32Vec16(const BF16Vec16& v) {
    reg.val[0] = (__vector float)vec_mergeh(zero, v.reg.val[0]);
    reg.val[1] = (__vector float)vec_mergel(zero, v.reg.val[0]);
@@ -769,40 +735,6 @@ inline BF16Vec8::BF16Vec8(const FP32Vec8& v) {
 #endif
 }

-inline FP16Vec16::FP16Vec16(const FP32Vec16& v) {
-  alignas(16) float temp_fp32[16];
-  alignas(16) c10::Half temp_fp16[16];
-
-  vec_xst(v.reg.val[0], 0, temp_fp32);
-  vec_xst(v.reg.val[1], 16, temp_fp32);
-  vec_xst(v.reg.val[2], 32, temp_fp32);
-  vec_xst(v.reg.val[3], 48, temp_fp32);
-
-  for (int i = 0; i < 16; i++) {
-    temp_fp16[i] = c10::Half(temp_fp32[i]);
-  }
-
-  reg.val[0] = (__vector signed short)vec_xl(0, (signed short*)temp_fp16);
-  reg.val[1] = (__vector signed short)vec_xl(16, (signed short*)temp_fp16);
-}
-
-inline FP32Vec16::FP32Vec16(const FP16Vec16& v) {
-  alignas(16) c10::Half temp_fp16[16];
-  alignas(16) float temp_fp32[16];
-
-  vec_xst(v.reg.val[0], 0, (signed short*)temp_fp16);
-  vec_xst(v.reg.val[1], 16, (signed short*)temp_fp16);
-
-  for (int i = 0; i < 16; i++) {
-    temp_fp32[i] = float(temp_fp16[i]);
-  }
-
-  reg.val[0] = vec_xl(0, temp_fp32);
-  reg.val[1] = vec_xl(16, temp_fp32);
-  reg.val[2] = vec_xl(32, temp_fp32);
-  reg.val[3] = vec_xl(48, temp_fp32);
-}
-
 inline BF16Vec16::BF16Vec16(const FP32Vec16& v) {
 #ifdef _ARCH_PWR10
  __vector signed short ret[4];
@@ -862,43 +794,6 @@ inline void prefetch(const void* addr) {
  __asm__ __volatile__("dcbt 0, %0" : : "r"(addr) : "memory");
 }

-struct INT8Vec64 {
-  __vector signed char data[4];
-
-  INT8Vec64() = default;
-
-  explicit INT8Vec64(const int8_t* ptr) {
-    data[0] = vec_xl(0, ptr);
-    data[1] = vec_xl(16, ptr);
-    data[2] = vec_xl(32, ptr);
-    data[3] = vec_xl(48, ptr);
-  }
-
-  explicit INT8Vec64(bool, const int8_t* ptr) : INT8Vec64(ptr) {}
-
-  void save(int8_t* ptr) const {
-    vec_xst(data[0], 0, ptr);
-    vec_xst(data[1], 16, ptr);
-    vec_xst(data[2], 32, ptr);
-    vec_xst(data[3], 48, ptr);
-  }
-
-  void save(int8_t* ptr, int elem_num) const {
-    if (elem_num <= 0) return;
-
-    int full_vecs = elem_num / 16;
-    for (int i = 0; i < full_vecs && i < 4; i++) {
-      vec_xst(data[i], i * 16, ptr);
-    }
-
-    int remaining = elem_num % 16;
-    if (remaining > 0 && full_vecs < 4) {
-      vec_xst_len(data[full_vecs], ptr + full_vecs * 16, remaining);
-    }
-  }
-
-  void nt_save(int8_t* ptr) const { save(ptr); }
-};
-}  // namespace vec_op
+};  // namespace vec_op

 #endif
@@ -1,82 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-#include <ATen/native/CPUBlas.h>
-
-// Unlike brgemm, PyTorch does not publicly expose at::native::cpublas::gemm
-// If OpenBLS is available in the PyTorch wheel, we rely on it for fast
-// bf16:bf16->fp32 GEMMs Otherwise, we fall back to PyTorch reference BLAS path.
-#if defined(VLLM_HAS_OPENBLAS)
-extern "C" void sbgemm_(char* transa, char* transb, int* m, int* n, int* k,
-                        float* alpha, const at::BFloat16* a, int* lda,
-                        const at::BFloat16* b, int* ldb, float* beta, float* c,
-                        int* ldc);
-
-extern "C" void sgemm_(char* transa, char* transb, int* m, int* n, int* k,
-                       float* alpha, const float* a, int* lda, const float* b,
-                       int* ldb, float* beta, float* c, int* ldc);
-
-inline char blas_transpose(at::native::TransposeType trans) {
-  switch (trans) {
-    case at::native::TransposeType::NoTranspose:
-      return 'n';
-    case at::native::TransposeType::Transpose:
-      return 't';
-    case at::native::TransposeType::ConjTranspose:
-      return 'c';
-  }
-  return 'n';
-}
-
-inline void blas_gemm(at::native::TransposeType transa,
-                      at::native::TransposeType transb, int64_t m, int64_t n,
-                      int64_t k, float alpha, const at::BFloat16* a,
-                      int64_t lda, const at::BFloat16* b, int64_t ldb,
-                      float beta, float* c, int64_t ldc) {
-  char transa_ = blas_transpose(transa);
-  char transb_ = blas_transpose(transb);
-  int m_ = static_cast<int>(m);
-  int n_ = static_cast<int>(n);
-  int k_ = static_cast<int>(k);
-  int lda_ = static_cast<int>(lda);
-  int ldb_ = static_cast<int>(ldb);
-  int ldc_ = static_cast<int>(ldc);
-  sbgemm_(&transa_, &transb_, &m_, &n_, &k_, &alpha, a, &lda_, b, &ldb_, &beta,
-          c, &ldc_);
-}
-
-inline void blas_gemm(at::native::TransposeType transa,
-                      at::native::TransposeType transb, int64_t m, int64_t n,
-                      int64_t k, float alpha, const float* a, int64_t lda,
-                      const float* b, int64_t ldb, float beta, float* c,
-                      int64_t ldc) {
-  char transa_ = blas_transpose(transa);
-  char transb_ = blas_transpose(transb);
-  int m_ = static_cast<int>(m);
-  int n_ = static_cast<int>(n);
-  int k_ = static_cast<int>(k);
-  int lda_ = static_cast<int>(lda);
-  int ldb_ = static_cast<int>(ldb);
-  int ldc_ = static_cast<int>(ldc);
-  sgemm_(&transa_, &transb_, &m_, &n_, &k_, &alpha, a, &lda_, b, &ldb_, &beta,
-         c, &ldc_);
-}
-
-inline void blas_gemm(at::native::TransposeType, at::native::TransposeType,
-                      int64_t, int64_t, int64_t, float, const at::Half*,
-                      int64_t, const at::Half*, int64_t, float, float*,
-                      int64_t) {
-  TORCH_CHECK(false, "CPU OpenBLAS hgemm is not available.");
-}
-#else
-template <typename scalar_t>
-inline void blas_gemm(at::native::TransposeType transa,
-                      at::native::TransposeType transb, int64_t m, int64_t n,
-                      int64_t k, float alpha, const scalar_t* a, int64_t lda,
-                      const scalar_t* b, int64_t ldb, float beta, float* c,
-                      int64_t ldc) {
-  auto gemm = at::native::cpublas::gemm_no_downcast_stub.DEFAULT;
-  gemm(c10::CppTypeToScalarType<scalar_t>::value, transa, transb, m, n, k,
-       at::Scalar(alpha), a, lda, b, ldb, at::Scalar(beta), c, ldc);
-}
-#endif
@@ -301,42 +301,25 @@ void chunk_gated_delta_rule_kernel_impl(
      // attn = k_beta @ key.transpose(-1, -2)
      // attn: [B, HV, num_chunk, chunk_size, chunk_size]
      // transpose and pack for key
-      if constexpr (brgemm_supported()) {
-        pack_vnni<scalar_t>(
-            /*    dst */ k_transpose,
-            /*    src */ curr_k_pad,
-            /*     N  */ chunk_size,
-            /*     K  */ qk_head_size,
-            /* ld_src */ qk_head_size,
-            /* ld_dst */ chunk_size);
-        // k_beta @ key.transpose(-1, -2)
-        at::native::cpublas::brgemm(
-            /*     M */ chunk_size,
-            /*     N */ chunk_size,
-            /*     K */ qk_head_size,
-            /*   lda */ qk_head_size,
-            /*   ldb */ chunk_size,
-            /*   ldc */ chunk_size,
-            /* add_C */ false,
-            /*     A */ curr_k_beta,
-            /*     B */ k_transpose,
-            /*     C */ curr_attn);
-      } else {
-        blas_gemm(
-            at::native::TransposeType::Transpose,
-            at::native::TransposeType::NoTranspose,
-            chunk_size,
-            chunk_size,
-            qk_head_size,
-            1.0f,
-            curr_k_pad,
-            qk_head_size,
-            curr_k_beta,
-            qk_head_size,
-            0.0f,
-            curr_attn,
-            chunk_size);
-      }
+      pack_vnni<scalar_t>(
+          /*    dst */ k_transpose,
+          /*    src */ curr_k_pad,
+          /*     N  */ chunk_size,
+          /*     K  */ qk_head_size,
+          /* ld_src */ qk_head_size,
+          /* ld_dst */ chunk_size);
+      // k_beta @ key.transpose(-1, -2)
+      at::native::cpublas::brgemm(
+          /*     M */ chunk_size,
+          /*     N */ chunk_size,
+          /*     K */ qk_head_size,
+          /*   lda */ qk_head_size,
+          /*   ldb */ chunk_size,
+          /*   ldc */ chunk_size,
+          /* add_C */ false,
+          /*     A */ curr_k_beta,
+          /*     B */ k_transpose,
+          /*     C */ curr_attn);
      // attn = attn * decay_mask
      for (int64_t m = 0; m < chunk_size; m++) {
        at::vec::map2<float>(
@@ -430,42 +413,25 @@ void chunk_gated_delta_rule_kernel_impl(
      // k_beta_g = k_beta * g: [B, HV, num_chunk, chunk_size, EK]
      // k_cumdecay: [B, HV, num_chunk, chunk_size, EK]
      // pack for value
-      if constexpr (brgemm_supported()) {
-        pack_vnni2<scalar_t>(
-            /*    dst */ v_pack,
-            /*    src */ curr_v_beta,
-            /*     N  */ chunk_size,
-            /*     K  */ v_head_size,
-            /* ld_src */ v_head_size,
-            /* ld_dst */ v_head_size);
-        // value = attn @ v_beta
-        at::native::cpublas::brgemm(
-            /*     M */ chunk_size,
-            /*     N */ v_head_size,
-            /*     K */ chunk_size,
-            /*   lda */ chunk_size,
-            /*   ldb */ v_head_size,
-            /*   ldc */ v_head_size,
-            /* add_C */ false,
-            /*     A */ curr_attn_reduced,
-            /*     B */ v_pack,
-            /*     C */ curr_value);
-      } else {
-        blas_gemm(
-            at::native::TransposeType::NoTranspose,
-            at::native::TransposeType::NoTranspose,
-            v_head_size,
-            chunk_size,
-            chunk_size,
-            1.0f,
-            curr_v_beta,
-            v_head_size,
-            curr_attn_reduced,
-            chunk_size,
-            0.0f,
-            curr_value,
-            v_head_size);
-      }
+      pack_vnni2<scalar_t>(
+          /*    dst */ v_pack,
+          /*    src */ curr_v_beta,
+          /*     N  */ chunk_size,
+          /*     K  */ v_head_size,
+          /* ld_src */ v_head_size,
+          /* ld_dst */ v_head_size);
+      // value = attn @ v_beta
+      at::native::cpublas::brgemm(
+          /*     M */ chunk_size,
+          /*     N */ v_head_size,
+          /*     K */ chunk_size,
+          /*   lda */ chunk_size,
+          /*   ldb */ v_head_size,
+          /*   ldc */ v_head_size,
+          /* add_C */ false,
+          /*     A */ curr_attn_reduced,
+          /*     B */ v_pack,
+          /*     C */ curr_value);
      // k_beta_g = k_beta * g.exp().unsqueeze(-1)
      for (int64_t j = 0; j < chunk_size; j++) {
        int64_t i = 0;
@@ -479,42 +445,25 @@ void chunk_gated_delta_rule_kernel_impl(
        }
      }
      // pack for k_beta_g
-      if constexpr (brgemm_supported()) {
-        pack_vnni2<scalar_t>(
-            /*    dst */ k_beta_g_pack,
-            /*    src */ k_beta_g,
-            /*     N  */ chunk_size,
-            /*     K  */ qk_head_size,
-            /* ld_src */ qk_head_size,
-            /* ld_dst */ qk_head_size);
-        // k_cumdecay = attn @ k_beta_g
-        at::native::cpublas::brgemm(
-            /*     M */ chunk_size,
-            /*     N */ qk_head_size,
-            /*     K */ chunk_size,
-            /*   lda */ chunk_size,
-            /*   ldb */ qk_head_size,
-            /*   ldc */ qk_head_size,
-            /* add_C */ false,
-            /*     A */ curr_attn_reduced,
-            /*     B */ k_beta_g_pack,
-            /*     C */ k_cumdecay);
-      } else {
-        blas_gemm(
-            at::native::TransposeType::NoTranspose,
-            at::native::TransposeType::NoTranspose,
-            qk_head_size,
-            chunk_size,
-            chunk_size,
-            1.0f,
-            k_beta_g,
-            qk_head_size,
-            curr_attn_reduced,
-            chunk_size,
-            0.0f,
-            k_cumdecay,
-            qk_head_size);
-      }
+      pack_vnni2<scalar_t>(
+          /*    dst */ k_beta_g_pack,
+          /*    src */ k_beta_g,
+          /*     N  */ chunk_size,
+          /*     K  */ qk_head_size,
+          /* ld_src */ qk_head_size,
+          /* ld_dst */ qk_head_size);
+      // k_cumdecay = attn @ k_beta_g
+      at::native::cpublas::brgemm(
+          /*     M */ chunk_size,
+          /*     N */ qk_head_size,
+          /*     K */ chunk_size,
+          /*   lda */ chunk_size,
+          /*   ldb */ qk_head_size,
+          /*   ldc */ qk_head_size,
+          /* add_C */ false,
+          /*     A */ curr_attn_reduced,
+          /*     B */ k_beta_g_pack,
+          /*     C */ k_cumdecay);
      for (int i = 0; i < chunk_size; i++) {
        at::vec::map<scalar_t>(
            [](fVec x) { return x; },
@@ -602,42 +551,25 @@ void chunk_gated_delta_rule_kernel_impl(

        // attn_i = (q_i @ k_i.transpose(-1, -2) * decay_mask[:, :, i]).masked_fill_(mask, 0)
        // k_transpose_i = k_i.transpose(-1, -2)
-        if constexpr (brgemm_supported()) {
-          pack_vnni<scalar_t>(
-              /*    dst */ k_transpose_i,
-              /*    src */ k_i,
-              /*     N  */ chunk_size,
-              /*     K  */ qk_head_size,
-              /* ld_src */ qk_head_size,
-              /* ld_dst */ chunk_size);
-          // attn_i = q_i @ k_transpose_i
-          at::native::cpublas::brgemm(
-              /* M */ chunk_size,
-              /* N */ chunk_size,
-              /* K */ qk_head_size,
-              /* lda */ qk_head_size,
-              /* ldb */ chunk_size,
-              /* ldc */ chunk_size,
-              /* add_C */ false,
-              /* A */ q_i,
-              /* B */ k_transpose_i,
-              /* C */ attn_i);
-        } else {
-          blas_gemm(
-              at::native::TransposeType::Transpose,
-              at::native::TransposeType::NoTranspose,
-              chunk_size,
-              chunk_size,
-              qk_head_size,
-              1.0f,
-              k_i,
-              qk_head_size,
-              q_i,
-              qk_head_size,
-              0.0f,
-              attn_i,
-              chunk_size);
-        }
+        pack_vnni<scalar_t>(
+            /*    dst */ k_transpose_i,
+            /*    src */ k_i,
+            /*     N  */ chunk_size,
+            /*     K  */ qk_head_size,
+            /* ld_src */ qk_head_size,
+            /* ld_dst */ chunk_size);
+        // attn_i = q_i @ k_transpose_i
+        at::native::cpublas::brgemm(
+            /* M */ chunk_size,
+            /* N */ chunk_size,
+            /* K */ qk_head_size,
+            /* lda */ qk_head_size,
+            /* ldb */ chunk_size,
+            /* ldc */ chunk_size,
+            /* add_C */ false,
+            /* A */ q_i,
+            /* B */ k_transpose_i,
+            /* C */ attn_i);
        // attn_i = attn_i * decay_mask_i
        for (int64_t m = 0; m < chunk_size; m++) {
          auto attn_i_m = attn_i + m * chunk_size;
@@ -677,45 +609,28 @@ void chunk_gated_delta_rule_kernel_impl(
        }

        // pack for curr_last_recurrent_state
-        if constexpr (brgemm_supported()) {
-          pack_vnni2<scalar_t>(
-              /*    dst */ curr_last_recurrent_state_pack_reduced,
-              /*    src */ curr_last_recurrent_state_reduced,
-              /*     N  */ qk_head_size,
-              /*     K  */ v_head_size,
-              /* ld_src */ v_head_size,
-              /* ld_dst */ v_head_size);
+        pack_vnni2<scalar_t>(
+            /*    dst */ curr_last_recurrent_state_pack_reduced,
+            /*    src */ curr_last_recurrent_state_reduced,
+            /*     N  */ qk_head_size,
+            /*     K  */ v_head_size,
+            /* ld_src */ v_head_size,
+            /* ld_dst */ v_head_size);

-          // v_prime = k_cumdecay_i @ curr_last_recurrent_state: [chunk_size, EV]
-          // k_cumdecay_i: [chunk_size, EK]
-          // curr_last_recurrent_state: [EK, EV]
-          at::native::cpublas::brgemm(
-              /*     M */ chunk_size,
-              /*     N */ v_head_size,
-              /*     K */ qk_head_size,
-              /*   lda */ qk_head_size,
-              /*   ldb */ v_head_size,
-              /*   ldc */ v_head_size,
-              /* add_C */ false,
-              /*     A */ k_cumdecay_i_reduced,
-              /*     B */ curr_last_recurrent_state_pack_reduced,
-              /*     C */ v_prime);
-        } else {
-          blas_gemm(
-              at::native::TransposeType::NoTranspose,
-              at::native::TransposeType::NoTranspose,
-              v_head_size,
-              chunk_size,
-              qk_head_size,
-              1.0f,
-              curr_last_recurrent_state_reduced,
-              v_head_size,
-              k_cumdecay_i_reduced,
-              qk_head_size,
-              0.0f,
-              v_prime,
-              v_head_size);
-        }
+        // v_prime = k_cumdecay_i @ curr_last_recurrent_state: [chunk_size, EV]
+        // k_cumdecay_i: [chunk_size, EK]
+        // curr_last_recurrent_state: [EK, EV]
+        at::native::cpublas::brgemm(
+            /*     M */ chunk_size,
+            /*     N */ v_head_size,
+            /*     K */ qk_head_size,
+            /*   lda */ qk_head_size,
+            /*   ldb */ v_head_size,
+            /*   ldc */ v_head_size,
+            /* add_C */ false,
+            /*     A */ k_cumdecay_i_reduced,
+            /*     B */ curr_last_recurrent_state_pack_reduced,
+            /*     C */ v_prime);

        // v_new = v_prime = v_i - v_prime
        // v_i: [chunk_size, EV]
@@ -748,75 +663,41 @@ void chunk_gated_delta_rule_kernel_impl(
        }
        // attn_inter = qg @ curr_last_recurrent_state: [chunk_size, EV]
        // curr_last_recurrent_state: [EK, EV]
-        if constexpr (brgemm_supported()) {
-          at::native::cpublas::brgemm(
-              /* M */ chunk_size,
-              /* N */ v_head_size,
-              /* K */ qk_head_size,
-              /* lda */ qk_head_size,
-              /* ldb */ v_head_size,
-              /* ldc */ v_head_size,
-              /* add_C */ false,
-              /* A */ qg,
-              /* B */ curr_last_recurrent_state_pack_reduced,
-              /* C */ attn_inter);
-        } else {
-          blas_gemm(
-              at::native::TransposeType::NoTranspose,
-              at::native::TransposeType::NoTranspose,
-              v_head_size,
-              chunk_size,
-              qk_head_size,
-              1.0f,
-              curr_last_recurrent_state_reduced,
-              v_head_size,
-              qg,
-              qk_head_size,
-              0.0f,
-              attn_inter,
-              v_head_size);
-        }
+        at::native::cpublas::brgemm(
+            /* M */ chunk_size,
+            /* N */ v_head_size,
+            /* K */ qk_head_size,
+            /* lda */ qk_head_size,
+            /* ldb */ v_head_size,
+            /* ldc */ v_head_size,
+            /* add_C */ false,
+            /* A */ qg,
+            /* B */ curr_last_recurrent_state_pack_reduced,
+            /* C */ attn_inter);

        // core_attn_out[:, :, i] = attn_inter + attn_i @ v_new
        // pack for v_prime
-        if constexpr (brgemm_supported()) {
-          pack_vnni2<scalar_t>(
-              /*    dst */ v_prime_pack_reduced,
-              /*    src */ v_prime_reduced,
-              /*     N  */ chunk_size,
-              /*     K  */ v_head_size,
-              /* ld_src */ v_head_size,
-              /* ld_dst */ v_head_size);
-          // attn_inter = attn_inter + attn_i @ v_new: [chunk_size, EV]
-          // attn_i: [chunk_size, chunk_size]
-          // v_new: [chunk_size, EV]
-          at::native::cpublas::brgemm(
-              /* M */ chunk_size,
-              /* N */ v_head_size,
-              /* K */ chunk_size,
-              /* lda */ chunk_size,
-              /* ldb */ v_head_size,
-              /* ldc */ v_head_size,
-              /* add_C */ true,
-              /* A */ attn_i_reduced,
-              /* B */ v_prime_pack_reduced,
-              /* C */ attn_inter);
-        } else {
-          blas_gemm(
-              at::native::TransposeType::NoTranspose,
-              at::native::TransposeType::NoTranspose,
-              v_head_size,
-              chunk_size,
-              chunk_size,
-              1.0f,
-              v_prime_reduced,
-              v_head_size,
-              attn_i_reduced,
-              chunk_size,
-              1.0f,
-              attn_inter,
-              v_head_size);
-        }
+        pack_vnni2<scalar_t>(
+            /*    dst */ v_prime_pack_reduced,
+            /*    src */ v_prime_reduced,
+            /*     N  */ chunk_size,
+            /*     K  */ v_head_size,
+            /* ld_src */ v_head_size,
+            /* ld_dst */ v_head_size);
+        // attn_inter = attn_inter + attn_i @ v_new: [chunk_size, EV]
+        // attn_i: [chunk_size, chunk_size]
+        // v_new: [chunk_size, EV]
+        at::native::cpublas::brgemm(
+            /* M */ chunk_size,
+            /* N */ v_head_size,
+            /* K */ chunk_size,
+            /* lda */ chunk_size,
+            /* ldb */ v_head_size,
+            /* ldc */ v_head_size,
+            /* add_C */ true,
+            /* A */ attn_i_reduced,
+            /* B */ v_prime_pack_reduced,
+            /* C */ attn_inter);

        // core_attn_out[:, :, i] = attn_inter
        for (int64_t m = 0; m < chunk_size; m++) {
@@ -881,34 +762,17 @@ void chunk_gated_delta_rule_kernel_impl(
            /* ld_dst */ chunk_size);
        // kgv = kg.transpose(-1, -2) @ v_new
        // v_new: [chunk_size, EV]
-        if constexpr (brgemm_supported()) {
-          at::native::cpublas::brgemm(
-              /* M */ qk_head_size,
-              /* N */ v_head_size,
-              /* K */ chunk_size,
-              /* lda */ chunk_size,
-              /* ldb */ v_head_size,
-              /* ldc */ v_head_size,
-              /* add_C */ false,
-              /* A */ kg_transpose,
-              /* B */ v_prime_pack_reduced,
-              /* C */ kgv);
-        } else {
-          blas_gemm(
-              at::native::TransposeType::NoTranspose,
-              at::native::TransposeType::NoTranspose,
-              v_head_size,
-              qk_head_size,
-              chunk_size,
-              1.0f,
-              v_prime_reduced,
-              v_head_size,
-              kg_transpose,
-              chunk_size,
-              0.0f,
-              kgv,
-              v_head_size);
-        }
+        at::native::cpublas::brgemm(
+            /* M */ qk_head_size,
+            /* N */ v_head_size,
+            /* K */ chunk_size,
+            /* lda */ chunk_size,
+            /* ldb */ v_head_size,
+            /* ldc */ v_head_size,
+            /* add_C */ false,
+            /* A */ kg_transpose,
+            /* B */ v_prime_pack_reduced,
+            /* C */ kgv);
        // last_recurrent_state = 1) + 2)
        for (int64_t m = 0; m < qk_head_size; m++) {
          at::vec::map2<float>(
@@ -1057,8 +921,7 @@ void fused_sigmoid_gating_delta_rule_update_kernel_impl(
      float k_scale = use_qk_l2norm_in_kernel ? qk_scale_buf[k_scale_offset] : 1.0f;
      int64_t v_offset = si * v_strideS + bi * v_strideB + ni * v_strideH;
      int64_t o_offset = ((bi * seq_len + si) * v_num_heads + ni) * v_head_dim;
-      // See: https://github.com/sgl-project/sglang/pull/26634
-      float beta_val = 1 / (1 + std::exp(-b_ptr[bi * v_num_heads + ni]));
+      float beta_val = 1 / (1 + std::exp(-b_ptr[ni]));
      fVec beta_vec = fVec(beta_val);
      int64_t dvi = 0;
      for (; dvi <= v_head_dim - VecSize; dvi += VecSize) {
@@ -4,12 +4,9 @@
 // clang-format off

 #pragma once
-#include "common.h"
-#include "blas_gemm.h"
+#include <ATen/native/CPUBlas.h>

-#if defined(__AVX512F__) && defined(__AVX512BF16__) && defined(__AMX_BF16__)
-#define CPU_CAPABILITY_AVX512
-#endif
+#include "common.h"

 // amx-bf16
 #define TILE_M 16
@@ -24,39 +21,31 @@ constexpr int block_size_n() {
  return 2 * TILE_N;
 }

-constexpr bool brgemm_supported() {
-#if defined(CPU_CAPABILITY_AVX512)
-  return true;
-#else
-  return false;
-#endif
-}
-
 // define threshold using brgemm (intel AMX)
 template <typename T>
 inline bool can_use_brgemm(int M);
 template <>
 inline bool can_use_brgemm<at::BFloat16>(int M) {
-  return brgemm_supported() && M > 4;
+  return M > 4;
 }
 template <>
 inline bool can_use_brgemm<at::Half>(int M) {
-  return brgemm_supported();
+  return true;
 }
 // this requires PyTorch 2.7 or above
 template <>
 inline bool can_use_brgemm<int8_t>(int M) {
-  return brgemm_supported() && M > 4;
+  return M > 4;
 }

 template <>
 inline bool can_use_brgemm<uint8_t>(int M) {
-  return brgemm_supported() && M > 4;
+  return M > 4;
 }

 template <>
 inline bool can_use_brgemm<at::Float8_e4m3fn>(int M) {
-  return brgemm_supported() && M > 4;
+  return M > 4;
 }

 // work around compiler internal error
@@ -11,9 +11,7 @@

 #include <ATen/cpu/vec/functional.h>
 #include <ATen/cpu/vec/vec.h>
-#if defined(CPU_CAPABILITY_AVX512)
 #include <immintrin.h>
-#endif
 namespace {

 using namespace at::vec;
@@ -5,7 +5,7 @@
 #include <sys/stat.h>
 #include <unistd.h>

-#if defined(__aarch64__) || defined(__powerpc64__)
+#ifdef __aarch64__
  #include <atomic>
 #endif

@@ -38,7 +38,7 @@ struct KernelVecType<c10::Half> {
 };

 struct ThreadSHMContext {
-#if defined(__aarch64__) || defined(__powerpc64__)
+#ifdef __aarch64__
  // memory model is weaker on AArch64, so we use atomic variables for
  // consumer (load-acquire) and producer (store-release) to make sure
  // that a stamp cannot be ready before the corresponding data is ready.
@@ -75,7 +75,7 @@ struct ThreadSHMContext {
    TORCH_CHECK(group_size <= MAX_SHM_RANK_NUM);
    TORCH_CHECK((size_t)this % 64 == 0);
    TORCH_CHECK((size_t)thread_shm_ptr % 64 == 0);
-#if defined(__aarch64__) || defined(__powerpc64__)
+#ifdef __aarch64__
    _curr_thread_stamp[0].store(1, std::memory_order_relaxed);
    _curr_thread_stamp[1].store(1, std::memory_order_relaxed);
    _ready_thread_stamp[0].store(0, std::memory_order_relaxed);
@@ -124,7 +124,7 @@ struct ThreadSHMContext {
  }

  char get_curr_stamp(int idx) const {
-#if defined(__aarch64__) || defined(__powerpc64__)
+#ifdef __aarch64__
    return _curr_thread_stamp[idx].load(std::memory_order_acquire);
 #else
    return _curr_thread_stamp[idx];
@@ -132,7 +132,7 @@ struct ThreadSHMContext {
  }

  char get_ready_stamp(int idx) const {
-#if defined(__aarch64__) || defined(__powerpc64__)
+#ifdef __aarch64__
    return _ready_thread_stamp[idx].load(std::memory_order_acquire);
 #else
    return _ready_thread_stamp[idx];
@@ -140,7 +140,7 @@ struct ThreadSHMContext {
  }

  void next_stamp() {
-#if defined(__aarch64__) || defined(__powerpc64__)
+#ifdef __aarch64__
    _curr_thread_stamp[local_stamp_buffer_idx].fetch_add(
        1, std::memory_order_release);
 #else
@@ -150,7 +150,7 @@ struct ThreadSHMContext {
  }

  void commit_ready_stamp() {
-#if defined(__aarch64__) || defined(__powerpc64__)
+#ifdef __aarch64__
    _ready_thread_stamp[local_stamp_buffer_idx].store(
        _curr_thread_stamp[local_stamp_buffer_idx].load(
            std::memory_order_relaxed),
@@ -186,10 +186,8 @@ struct ThreadSHMContext {
        break;
      }
      ++_spinning_count;
-#if defined(__aarch64__)
+#ifdef __aarch64__
      __asm__ __volatile__("yield");
-#elif defined(__powerpc64__)
-      __asm__ __volatile__("or 1,1,1");
 #else
      _mm_pause();
 #endif  // __aarch64__
@@ -378,8 +378,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
 #endif

 // SHM CCL
-#if defined(__AVX512F__) || (defined(__aarch64__) && !defined(__APPLE__)) || \
-    defined(__powerpc64__)
+#if defined(__AVX512F__) || (defined(__aarch64__) && !defined(__APPLE__))
  ops.def(
      "init_shm_manager(str name, int group_size, int rank, int thread_num) -> "
      "int",
@@ -448,25 +447,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      "bool is_vnni) -> Tensor");
  ops.impl("fp8_scaled_mm_cpu", torch::kCPU, &fp8_scaled_mm_cpu);

-  // Adapted from sglang: casual_conv1d kernels
-  ops.def("causal_conv1d_weight_pack(Tensor weight) -> Tensor");
-  ops.impl("causal_conv1d_weight_pack", torch::kCPU,
-           &causal_conv1d_weight_pack);
-  ops.def(
-      "causal_conv1d_fwd_cpu(Tensor x, Tensor weight, Tensor? bias, Tensor? "
-      "conv_states, Tensor? query_start_loc,"
-      "Tensor? cache_indices, Tensor? has_initial_state, bool silu_activation, "
-      "int pad_slot_id, bool is_vnni) -> "
-      "Tensor");
-  ops.impl("causal_conv1d_fwd_cpu", torch::kCPU, &causal_conv1d_fwd_cpu);
-  ops.def(
-      "causal_conv1d_update_cpu(Tensor x, Tensor(a!) conv_states, Tensor "
-      "weight, Tensor? bias, bool silu_activation,"
-      "Tensor? cache_seqlens, Tensor? conv_state_indices, int pad_slot_id, "
-      "bool is_vnni) -> Tensor");
-  ops.impl("causal_conv1d_update_cpu", torch::kCPU, &causal_conv1d_update_cpu);
-#endif
-
  // Adapted from sglang: GDN kernels
  ops.def(
      "chunk_gated_delta_rule_cpu(Tensor query, Tensor key, Tensor value, "
@@ -490,6 +470,25 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      "-> (Tensor, Tensor)");
  ops.impl("fused_gdn_gating_cpu", torch::kCPU, &fused_gdn_gating_cpu);

+  // Adapted from sglang: casual_conv1d kernels
+  ops.def("causal_conv1d_weight_pack(Tensor weight) -> Tensor");
+  ops.impl("causal_conv1d_weight_pack", torch::kCPU,
+           &causal_conv1d_weight_pack);
+  ops.def(
+      "causal_conv1d_fwd_cpu(Tensor x, Tensor weight, Tensor? bias, Tensor? "
+      "conv_states, Tensor? query_start_loc,"
+      "Tensor? cache_indices, Tensor? has_initial_state, bool silu_activation, "
+      "int pad_slot_id, bool is_vnni) -> "
+      "Tensor");
+  ops.impl("causal_conv1d_fwd_cpu", torch::kCPU, &causal_conv1d_fwd_cpu);
+  ops.def(
+      "causal_conv1d_update_cpu(Tensor x, Tensor(a!) conv_states, Tensor "
+      "weight, Tensor? bias, bool silu_activation,"
+      "Tensor? cache_seqlens, Tensor? conv_state_indices, int pad_slot_id, "
+      "bool is_vnni) -> Tensor");
+  ops.impl("causal_conv1d_update_cpu", torch::kCPU, &causal_conv1d_update_cpu);
+#endif
+
  // CPU attention kernels
  ops.def(
      "get_scheduler_metadata(int num_req, int num_heads_q, int num_heads_kv, "
@@ -519,7 +518,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
  ops.def("dynamic_per_token_scaled_fp8_quant() -> ()", placeholder_op);

  // WNA16
-#if defined(__AVX512F__) || defined(__riscv_v)
+#if defined(__AVX512F__)
  ops.def(
      "cpu_gemm_wna16(Tensor input, Tensor q_weight, Tensor(a2!) output, "
      "Tensor scales, Tensor? zeros, Tensor? g_idx, Tensor? bias, SymInt "
@@ -4,7 +4,7 @@
 #include <cmath>

 #include "../cuda_compat.h"
-#include "cuda_vec_utils.cuh"
+#include "../cuda_vec_utils.cuh"
 #include "dispatch_utils.h"
 #include "torch_utils.h"

@@ -7,7 +7,7 @@
 #include <torch/headeronly/core/ScalarType.h>

 #include "../../attention/attention_dtypes.h"
-#include "attention_utils.cuh"
+#include "../../attention/attention_utils.cuh"
 #include "../../quantization/w8a8/fp8/common.cuh"

 namespace vllm {
@@ -2,7 +2,7 @@

 #include <torch/csrc/stable/tensor.h>

-#include "broadcast_load_epilogue_c2x.hpp"
+#include "cutlass_extensions/epilogue/broadcast_load_epilogue_c2x.hpp"

 /*
   This file defines custom epilogues for fusing channel scales, token scales,
@@ -1,223 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-//
-// Router GEMM: activation(T) x weight(fp32) -> fp32, H=3072, E=256, M<=32.
-// Supports bf16 or fp32 activation; weight is always fp32.
-// Adapted from dsv3_router_gemm_float_out.cu.
-
-#include <cuda_bf16.h>
-#include <cuda_runtime.h>
-
-// ---------------------------------------------------------------------------
-// Load helpers
-// ---------------------------------------------------------------------------
-
-// Load VPT fp32 values from the weight matrix (always fp32).
-//   VPT=4 when activation is fp32 (one float4 load)
-//   VPT=8 when activation is bf16 (two float4 loads)
-template <int VPT>
-__device__ __forceinline__ void load_weight(float const* ptr, float* dst);
-
-template <>
-__device__ __forceinline__ void load_weight<4>(float const* ptr, float* dst) {
-  float4 v = *reinterpret_cast<float4 const*>(ptr);
-  dst[0] = v.x;
-  dst[1] = v.y;
-  dst[2] = v.z;
-  dst[3] = v.w;
-}
-
-template <>
-__device__ __forceinline__ void load_weight<8>(float const* ptr, float* dst) {
-  float4 v0 = *reinterpret_cast<float4 const*>(ptr);
-  float4 v1 = *reinterpret_cast<float4 const*>(ptr + 4);
-  dst[0] = v0.x;
-  dst[1] = v0.y;
-  dst[2] = v0.z;
-  dst[3] = v0.w;
-  dst[4] = v1.x;
-  dst[5] = v1.y;
-  dst[6] = v1.z;
-  dst[7] = v1.w;
-}
-
-// Load VPT activation values and convert to fp32.
-template <typename T, int VPT>
-__device__ __forceinline__ void load_activation(T const* ptr, float* dst);
-
-// fp32 activation: one float4 load, no conversion needed.
-template <>
-__device__ __forceinline__ void load_activation<float, 4>(float const* ptr,
-                                                          float* dst) {
-  float4 v = *reinterpret_cast<float4 const*>(ptr);
-  dst[0] = v.x;
-  dst[1] = v.y;
-  dst[2] = v.z;
-  dst[3] = v.w;
-}
-
-// bf16 activation: one uint4 load (8 × bf16) + element-wise conversion.
-template <>
-__device__ __forceinline__ void load_activation<__nv_bfloat16, 8>(
-    __nv_bfloat16 const* ptr, float* dst) {
-  uint4 v = *reinterpret_cast<uint4 const*>(ptr);
-  __nv_bfloat16 const* bf16_ptr = reinterpret_cast<__nv_bfloat16 const*>(&v);
-#pragma unroll
-  for (int i = 0; i < 8; i++) dst[i] = __bfloat162float(bf16_ptr[i]);
-}
-
-// ---------------------------------------------------------------------------
-// Kernel
-// ---------------------------------------------------------------------------
-
-// InputT : type of activation (float or __nv_bfloat16)
-// Weight is always fp32; output is always fp32.
-// VPT = 16 / sizeof(InputT):  4 for fp32, 8 for bf16
-template <typename InputT, int kBlockSize, int kNumTokens, int kNumExperts,
-          int kHiddenDim>
-__global__ __launch_bounds__(128, 1) void fp32_router_gemm_kernel(
-    float* out, InputT const* mat_a, float const* mat_b) {
-  constexpr int VPT = 16 / sizeof(InputT);
-  constexpr int k_elems_per_k_iteration = VPT * kBlockSize;
-  constexpr int k_iterations = kHiddenDim / k_elems_per_k_iteration;
-  constexpr int kWarpSize = 32;
-  constexpr int kNumWarps = kBlockSize / kWarpSize;
-
-  int const n_idx = blockIdx.x;
-  int const tid = threadIdx.x;
-  int const warpId = tid / kWarpSize;
-  int const laneId = tid % kWarpSize;
-
-  float acc[kNumTokens] = {};
-  __shared__ float sm_reduction[kNumTokens][kNumWarps];
-
-  float const* b_col = mat_b + n_idx * kHiddenDim;
-
-  int k_bases[k_iterations];
-#pragma unroll
-  for (int ki = 0; ki < k_iterations; ki++) {
-    k_bases[ki] = ki * k_elems_per_k_iteration + tid * VPT;
-  }
-
-#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-  asm volatile("griddepcontrol.wait;");
-#endif
-
-  for (int ki = 0; ki < k_iterations; ki++) {
-    int const k_base = k_bases[ki];
-
-    float b_float[VPT];
-    load_weight<VPT>(b_col + k_base, b_float);
-
-#pragma unroll
-    for (int m_idx = 0; m_idx < kNumTokens; m_idx++) {
-      float a_float[VPT];
-      load_activation<InputT, VPT>(mat_a + m_idx * kHiddenDim + k_base,
-                                   a_float);
-#pragma unroll
-      for (int k = 0; k < VPT; k++) {
-        acc[m_idx] += a_float[k] * b_float[k];
-      }
-    }
-  }
-
-  // Warp-level butterfly reduction
-#pragma unroll
-  for (int m = 0; m < kNumTokens; m++) {
-    float sum = acc[m];
-    sum += __shfl_xor_sync(0xffffffff, sum, 16);
-    sum += __shfl_xor_sync(0xffffffff, sum, 8);
-    sum += __shfl_xor_sync(0xffffffff, sum, 4);
-    sum += __shfl_xor_sync(0xffffffff, sum, 2);
-    sum += __shfl_xor_sync(0xffffffff, sum, 1);
-    if (laneId == 0) sm_reduction[m][warpId] = sum;
-  }
-
-  __syncthreads();
-
-  if (tid == 0) {
-#pragma unroll
-    for (int m = 0; m < kNumTokens; m++) {
-      float final_sum = 0.0f;
-#pragma unroll
-      for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][w];
-      out[m * kNumExperts + n_idx] = final_sum;
-    }
-  }
-
-#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-  asm volatile("griddepcontrol.launch_dependents;");
-#endif
-}
-
-// ---------------------------------------------------------------------------
-// Launcher
-// ---------------------------------------------------------------------------
-
-template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
-void invokeFp32RouterGemm(float* output, InputT const* mat_a,
-                          float const* mat_b, cudaStream_t stream) {
-  constexpr int kBlockSize = 128;
-  cudaLaunchConfig_t config;
-  config.gridDim = kNumExperts;
-  config.blockDim = kBlockSize;
-  config.dynamicSmemBytes = 0;
-  config.stream = stream;
-  cudaLaunchAttribute attrs[1];
-  attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-  attrs[0].val.programmaticStreamSerializationAllowed = 1;
-  config.numAttrs = 1;
-  config.attrs = attrs;
-  cudaLaunchKernelEx(&config,
-                     fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens,
-                                             kNumExperts, kHiddenDim>,
-                     output, mat_a, mat_b);
-}
-
-// ---------------------------------------------------------------------------
-// Explicit instantiations: M=1..32, E=256, H=3072, for both input types
-// ---------------------------------------------------------------------------
-
-#define INSTANTIATE(T, M)                              \
-  template void invokeFp32RouterGemm<T, M, 256, 3072>( \
-      float*, T const*, float const*, cudaStream_t);
-
-#define INSTANTIATE_ALL(T) \
-  INSTANTIATE(T, 1)        \
-  INSTANTIATE(T, 2)        \
-  INSTANTIATE(T, 3)        \
-  INSTANTIATE(T, 4)        \
-  INSTANTIATE(T, 5)        \
-  INSTANTIATE(T, 6)        \
-  INSTANTIATE(T, 7)        \
-  INSTANTIATE(T, 8)        \
-  INSTANTIATE(T, 9)        \
-  INSTANTIATE(T, 10)       \
-  INSTANTIATE(T, 11)       \
-  INSTANTIATE(T, 12)       \
-  INSTANTIATE(T, 13)       \
-  INSTANTIATE(T, 14)       \
-  INSTANTIATE(T, 15)       \
-  INSTANTIATE(T, 16)       \
-  INSTANTIATE(T, 17)       \
-  INSTANTIATE(T, 18)       \
-  INSTANTIATE(T, 19)       \
-  INSTANTIATE(T, 20)       \
-  INSTANTIATE(T, 21)       \
-  INSTANTIATE(T, 22)       \
-  INSTANTIATE(T, 23)       \
-  INSTANTIATE(T, 24)       \
-  INSTANTIATE(T, 25)       \
-  INSTANTIATE(T, 26)       \
-  INSTANTIATE(T, 27)       \
-  INSTANTIATE(T, 28)       \
-  INSTANTIATE(T, 29)       \
-  INSTANTIATE(T, 30)       \
-  INSTANTIATE(T, 31)       \
-  INSTANTIATE(T, 32)
-
-INSTANTIATE_ALL(float)
-INSTANTIATE_ALL(__nv_bfloat16)
-
-#undef INSTANTIATE_ALL
-#undef INSTANTIATE
@@ -1,127 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-#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 <cuda_bf16.h>
-#include <cuda_runtime.h>
-
-#include <stdexcept>
-
-namespace {
-
-inline int getSMVersion() {
-  auto* props = get_device_prop();
-  return props->major * 10 + props->minor;
-}
-
-}  // namespace
-
-static constexpr int FP32_NUM_EXPERTS = 256;
-static constexpr int FP32_HIDDEN_DIM = 3072;
-static constexpr int FP32_MAX_TOKENS = 32;
-
-// Forward declarations — 4 template params must match fp32_router_gemm.cu
-template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
-void invokeFp32RouterGemm(float* output, InputT const* mat_a,
-                          float const* mat_b, cudaStream_t stream);
-
-// LoopUnroller templated on InputT
-template <typename InputT, int kBegin, int kEnd>
-struct Fp32LoopUnroller {
-  static void unroll(int num_tokens, float* output, InputT const* mat_a,
-                     float const* mat_b, cudaStream_t stream) {
-    if (num_tokens == kBegin) {
-      invokeFp32RouterGemm<InputT, kBegin, FP32_NUM_EXPERTS, FP32_HIDDEN_DIM>(
-          output, mat_a, mat_b, stream);
-    } else {
-      Fp32LoopUnroller<InputT, kBegin + 1, kEnd>::unroll(num_tokens, output,
-                                                         mat_a, mat_b, stream);
-    }
-  }
-};
-
-template <typename InputT, int kEnd>
-struct Fp32LoopUnroller<InputT, kEnd, kEnd> {
-  static void unroll(int num_tokens, float* output, InputT const* mat_a,
-                     float const* mat_b, cudaStream_t stream) {
-    if (num_tokens == kEnd) {
-      invokeFp32RouterGemm<InputT, kEnd, FP32_NUM_EXPERTS, FP32_HIDDEN_DIM>(
-          output, mat_a, mat_b, stream);
-    } else {
-      throw std::invalid_argument(
-          "fp32_router_gemm: num_tokens must be in [1, 32]");
-    }
-  }
-};
-
-void fp32_router_gemm(
-    torch::stable::Tensor& output,       // [num_tokens, num_experts]
-    torch::stable::Tensor const& mat_a,  // [num_tokens, hidden_dim]
-    torch::stable::Tensor const& mat_b   // [num_experts, hidden_dim]
-) {
-  STD_TORCH_CHECK(output.dim() == 2 && mat_a.dim() == 2 && mat_b.dim() == 2);
-  STD_TORCH_CHECK(output.is_cuda() && mat_a.is_cuda() && mat_b.is_cuda(),
-                  "fp32_router_gemm: all tensors must be CUDA tensors");
-  STD_TORCH_CHECK(output.get_device_index() == mat_a.get_device_index() &&
-                      output.get_device_index() == mat_b.get_device_index(),
-                  "fp32_router_gemm: all tensors must be on the same device");
-  STD_TORCH_CHECK(
-      output.is_contiguous() && mat_a.is_contiguous() && mat_b.is_contiguous(),
-      "fp32_router_gemm: all tensors must be contiguous");
-
-  const int num_tokens = mat_a.size(0);
-  const int num_experts = mat_b.size(0);
-  const int hidden_dim = mat_a.size(1);
-
-  STD_TORCH_CHECK(output.size(0) == num_tokens && output.size(1) == num_experts,
-                  "fp32_router_gemm: output must have shape [num_tokens, "
-                  "num_experts]");
-  STD_TORCH_CHECK(
-      mat_a.size(1) == mat_b.size(1),
-      "fp32_router_gemm: mat_a and mat_b must have the same hidden_dim");
-  STD_TORCH_CHECK(hidden_dim == FP32_HIDDEN_DIM,
-                  "fp32_router_gemm: expected hidden_dim=3072");
-  STD_TORCH_CHECK(num_experts == FP32_NUM_EXPERTS,
-                  "fp32_router_gemm: expected num_experts=256");
-  STD_TORCH_CHECK(num_tokens <= FP32_MAX_TOKENS,
-                  "fp32_router_gemm: num_tokens must be in [0, 32]");
-  STD_TORCH_CHECK(
-      mat_a.scalar_type() == torch::headeronly::ScalarType::Float ||
-          mat_a.scalar_type() == torch::headeronly::ScalarType::BFloat16,
-      "fp32_router_gemm: mat_a must be float32 or bfloat16");
-  STD_TORCH_CHECK(mat_b.scalar_type() == torch::headeronly::ScalarType::Float,
-                  "fp32_router_gemm: mat_b (weight) must be float32");
-  STD_TORCH_CHECK(output.scalar_type() == torch::headeronly::ScalarType::Float,
-                  "fp32_router_gemm: output must be float32");
-
-  if (num_tokens == 0) {
-    return;
-  }
-
-  STD_TORCH_CHECK(getSMVersion() >= 90, "fp32_router_gemm: requires SM90+");
-
-  auto stream = get_current_cuda_stream(mat_a.get_device_index());
-  float* out_ptr = reinterpret_cast<float*>(output.mutable_data_ptr());
-  float const* mat_b_ptr = reinterpret_cast<float const*>(mat_b.data_ptr());
-
-  if (mat_a.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
-    auto const* mat_a_ptr =
-        reinterpret_cast<__nv_bfloat16 const*>(mat_a.data_ptr());
-    Fp32LoopUnroller<__nv_bfloat16, 1, FP32_MAX_TOKENS>::unroll(
-        num_tokens, out_ptr, mat_a_ptr, mat_b_ptr, stream);
-  } else {
-    auto const* mat_a_ptr = reinterpret_cast<float const*>(mat_a.data_ptr());
-    Fp32LoopUnroller<float, 1, FP32_MAX_TOKENS>::unroll(
-        num_tokens, out_ptr, mat_a_ptr, mat_b_ptr, stream);
-  }
-}
-
-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("fp32_router_gemm", TORCH_BOX(&fp32_router_gemm));
-}
@@ -20,7 +20,7 @@

 #include "torch_utils.h"

-#include "async_util.cuh"
+#include "../async_util.cuh"
 #include "../cuda_compat.h"
 #include "../type_convert.cuh"
 #include "dispatch_utils.h"
@@ -78,7 +78,8 @@ __global__ void rms_norm_kernel(
 #pragma unroll
    for (int j = 0; j < VEC_SIZE; j++) {
      float x = static_cast<float>(src1.val[j]);
-      dst.val[j] = static_cast<scalar_t>(x * s_variance) * src2.val[j];
+      float w = static_cast<float>(src2.val[j]);
+      dst.val[j] = static_cast<scalar_t>(x * s_variance * w);
    }
    v_out[i] = dst;
  }
@@ -142,7 +143,8 @@ fused_add_rms_norm_kernel(
 #pragma unroll
    for (int j = 0; j < width; ++j) {
      float x = Converter::convert(res.data[j]);
-      out.data[j] = Converter::convert(x * s_variance) * w.data[j];
+      float wf = Converter::convert(w.data[j]);
+      out.data[j] = Converter::convert(x * s_variance * wf);
    }
    input_v[strided_id] = out;
  }
@@ -181,8 +183,8 @@ fused_add_rms_norm_kernel(

  for (int idx = threadIdx.x; idx < hidden_size; idx += blockDim.x) {
    float x = (float)residual[blockIdx.x * hidden_size + idx];
-    input[blockIdx.x * input_stride + idx] =
-        (scalar_t)(x * s_variance) * weight[idx];
+    float w = (float)weight[idx];
+    input[blockIdx.x * input_stride + idx] = (scalar_t)(x * s_variance * w);
  }
 }

@@ -66,8 +66,13 @@ __global__ void rms_norm_static_fp8_quant_kernel(
 #pragma unroll
    for (int j = 0; j < VEC_SIZE; j++) {
      float x = static_cast<float>(src1.val[j]);
-      // Multiply in weight's native dtype to match rms_norm_kernel.
-      scalar_t out_norm = static_cast<scalar_t>(x * s_variance) * src2.val[j];
+      float w = static_cast<float>(src2.val[j]);
+      // Round normalized result through scalar_t to match the precision of the
+      // unfused composite (rms_norm writes scalar_t, then
+      // static_scaled_fp8_quant re-loads it as float before FP8 conversion).
+      // Without this round, the fused path is strictly more accurate and
+      // disagrees with the composite at exact E4M3 quantization tie boundaries.
+      scalar_t out_norm = static_cast<scalar_t>(x * s_variance * w);
      out[blockIdx.x * hidden_size + idx * VEC_SIZE + j] =
          scaled_fp8_conversion<true, fp8_type>(static_cast<float>(out_norm),
                                                scale_inv);
@@ -137,8 +142,12 @@ fused_add_rms_norm_static_fp8_quant_kernel(
 #pragma unroll
    for (int i = 0; i < width; ++i) {
      float x = Converter::convert(res.data[i]);
-      // Multiply in weight's native dtype to match fused_add_rms_norm_kernel.
-      HipT out_norm_h = Converter::convert(x * s_variance) * w.data[i];
+      float wf = Converter::convert(w.data[i]);
+      // See note in rms_norm_static_fp8_quant_kernel: round through scalar_t
+      // to match the unfused composite path at FP8 boundaries. We use the
+      // backend's hip_type for the intermediate since c10::Half/BFloat16 has
+      // ambiguous conversions on CUDA and no implicit conversion on ROCm.
+      HipT out_norm_h = Converter::convert(x * s_variance * wf);
      out[id * width + i] = scaled_fp8_conversion<true, fp8_type>(
          Converter::convert(out_norm_h), scale_inv);
    }
@@ -183,8 +192,10 @@ fused_add_rms_norm_static_fp8_quant_kernel(

  for (int idx = threadIdx.x; idx < hidden_size; idx += blockDim.x) {
    float x = (float)residual[blockIdx.x * hidden_size + idx];
-    // Multiply in weight's native dtype to match fused_add_rms_norm_kernel.
-    scalar_t out_norm = static_cast<scalar_t>(x * s_variance) * weight[idx];
+    float w = (float)weight[idx];
+    // See note in rms_norm_static_fp8_quant_kernel: round through scalar_t
+    // to match the unfused composite path at FP8 boundaries.
+    scalar_t out_norm = static_cast<scalar_t>(x * s_variance * w);
    out[blockIdx.x * hidden_size + idx] = scaled_fp8_conversion<true, fp8_type>(
        static_cast<float>(out_norm), scale_inv);
  }
@@ -355,132 +355,3 @@ torch::stable::Tensor ggml_moe_a8_vec(torch::stable::Tensor X,
                                      int64_t tokens);

 int64_t ggml_moe_get_block_size(int64_t type);
-
-void paged_attention_v1(
-    torch::stable::Tensor& out, torch::stable::Tensor& query,
-    torch::stable::Tensor& key_cache, torch::stable::Tensor& value_cache,
-    int64_t num_kv_heads, double scale, torch::stable::Tensor& block_tables,
-    torch::stable::Tensor& seq_lens, int64_t block_size, int64_t max_seq_len,
-    const std::optional<torch::stable::Tensor>& alibi_slopes,
-    const std::string& kv_cache_dtype, torch::stable::Tensor& k_scale,
-    torch::stable::Tensor& v_scale, const int64_t tp_rank,
-    const int64_t blocksparse_local_blocks,
-    const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
-    const int64_t blocksparse_head_sliding_step);
-
-void paged_attention_v2(
-    torch::stable::Tensor& out, torch::stable::Tensor& exp_sums,
-    torch::stable::Tensor& max_logits, torch::stable::Tensor& tmp_out,
-    torch::stable::Tensor& query, torch::stable::Tensor& key_cache,
-    torch::stable::Tensor& value_cache, int64_t num_kv_heads, double scale,
-    torch::stable::Tensor& block_tables, torch::stable::Tensor& seq_lens,
-    int64_t block_size, int64_t max_seq_len,
-    const std::optional<torch::stable::Tensor>& alibi_slopes,
-    const std::string& kv_cache_dtype, torch::stable::Tensor& k_scale,
-    torch::stable::Tensor& v_scale, const int64_t tp_rank,
-    const int64_t blocksparse_local_blocks,
-    const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
-    const int64_t blocksparse_head_sliding_step);
-
-// Cache ops (shared CUDA/ROCm)
-void swap_blocks(torch::stable::Tensor& src, torch::stable::Tensor& dst,
-                 int64_t block_size_in_bytes,
-                 const torch::stable::Tensor& block_mapping);
-
-// Batch swap: submit all block copies in a single driver call.
-void swap_blocks_batch(const torch::stable::Tensor& src_ptrs,
-                       const torch::stable::Tensor& dst_ptrs,
-                       const torch::stable::Tensor& sizes,
-                       bool is_src_access_order_any);
-
-void reshape_and_cache(torch::stable::Tensor& key, torch::stable::Tensor& value,
-                       torch::stable::Tensor& key_cache,
-                       torch::stable::Tensor& value_cache,
-                       torch::stable::Tensor& slot_mapping,
-                       const std::string& kv_cache_dtype,
-                       torch::stable::Tensor& k_scale,
-                       torch::stable::Tensor& v_scale);
-
-void reshape_and_cache_flash(
-    torch::stable::Tensor& key, torch::stable::Tensor& value,
-    torch::stable::Tensor& key_cache, torch::stable::Tensor& value_cache,
-    torch::stable::Tensor& slot_mapping, const std::string& kv_cache_dtype,
-    torch::stable::Tensor& k_scale, torch::stable::Tensor& v_scale);
-
-void concat_and_cache_mla(torch::stable::Tensor& kv_c,
-                          torch::stable::Tensor& k_pe,
-                          torch::stable::Tensor& kv_cache,
-                          torch::stable::Tensor& slot_mapping,
-                          const std::string& kv_cache_dtype,
-                          torch::stable::Tensor& scale);
-
-// NOTE: k_pe and kv_c order is flipped compared to concat_and_cache_mla
-void concat_and_cache_mla_rope_fused(
-    torch::stable::Tensor& positions, torch::stable::Tensor& q_pe,
-    torch::stable::Tensor& k_pe, torch::stable::Tensor& kv_c,
-    torch::stable::Tensor& rope_cos_sin_cache, bool rope_is_neox,
-    torch::stable::Tensor& slot_mapping, torch::stable::Tensor& kv_cache,
-    const std::string& kv_cache_dtype,
-    torch::stable::Tensor& kv_cache_quant_scale);
-
-// Just for unittest
-void convert_fp8(torch::stable::Tensor& dst_cache,
-                 torch::stable::Tensor& src_cache, const double scale,
-                 const std::string& kv_cache_dtype);
-
-void gather_and_maybe_dequant_cache(
-    torch::stable::Tensor const& src_cache,     // [NUM_BLOCKS, BLOCK_SIZE,
-                                                // ENTRIES...]
-    torch::stable::Tensor const& dst,           // [TOT_TOKENS, ENTRIES...]
-    torch::stable::Tensor const& block_table,   // [BATCH, BLOCK_INDICES]
-    torch::stable::Tensor const& cu_seq_lens,   // [BATCH+1]
-    torch::stable::Tensor const& token_to_seq,  // [MAX_TOKEN_ACROSS_CHUNKS]
-    int64_t num_tokens, const std::string& kv_cache_dtype,
-    torch::stable::Tensor const& scale,
-    std::optional<torch::stable::Tensor> seq_starts = std::nullopt);
-
-// TODO(hc): cp_gather_cache need support scaled kvcahe in the future.
-void cp_gather_cache(
-    torch::stable::Tensor const& src_cache,    // [NUM_BLOCKS, BLOCK_SIZE,
-                                               // ENTRIES...]
-    torch::stable::Tensor const& dst,          // [TOT_TOKENS, ENTRIES...]
-    torch::stable::Tensor const& block_table,  // [BATCH, BLOCK_INDICES]
-    torch::stable::Tensor const& cu_seq_lens,  // [BATCH+1]
-    int64_t batch_size,
-    std::optional<torch::stable::Tensor> seq_starts = std::nullopt);
-
-// Gather and upconvert FP8 KV cache to BF16 workspace
-void cp_gather_and_upconvert_fp8_kv_cache(
-    torch::stable::Tensor const& src_cache,         // [NUM_BLOCKS, BLOCK_SIZE,
-                                                    // 656]
-    torch::stable::Tensor const& dst,               // [TOT_TOKENS, 576]
-    torch::stable::Tensor const& block_table,       // [BATCH, BLOCK_INDICES]
-    torch::stable::Tensor const& seq_lens,          // [BATCH]
-    torch::stable::Tensor const& workspace_starts,  // [BATCH]
-    int64_t batch_size);
-
-// Indexer K quantization and cache function
-void indexer_k_quant_and_cache(
-    torch::stable::Tensor& k,             // [num_tokens, head_dim]
-    torch::stable::Tensor& kv_cache,      // [num_blocks, block_size,
-                                          // cache_stride]
-    torch::stable::Tensor& slot_mapping,  // [num_tokens]
-    int64_t quant_block_size,             // quantization block size
-    const std::string& scale_fmt);
-
-// Concatenate query nope and rope for MLA/DSA attention
-void concat_mla_q(
-    torch::stable::Tensor& ql_nope,  // [num_tokens, num_heads, nope_dim]
-    torch::stable::Tensor& q_pe,     // [num_tokens, num_heads, rope_dim]
-    torch::stable::Tensor& q_out);   // [num_tokens, num_heads, nope_dim +
-                                     // rope_dim]
-
-// Extract function to gather quantized K cache
-void cp_gather_indexer_k_quant_cache(
-    const torch::stable::Tensor& kv_cache,      // [num_blocks, block_size,
-                                                // cache_stride]
-    torch::stable::Tensor& dst_k,               // [num_tokens, head_dim]
-    torch::stable::Tensor& dst_scale,           // [num_tokens, head_dim /
-                                                // quant_block_size * 4]
-    const torch::stable::Tensor& block_table,   // [batch_size, num_blocks]
-    const torch::stable::Tensor& cu_seq_lens);  // [batch_size + 1]
@@ -17,7 +17,7 @@
 #include <torch/csrc/stable/tensor.h>
 #include "libtorch_stable/torch_utils.h"
 #include "libtorch_stable/dispatch_utils.h"
-#include "../../cuda_vec_utils.cuh"
+#include "cuda_vec_utils.cuh"

 #include <cuda_runtime_api.h>
 #include <cuda_runtime.h>
@@ -25,7 +25,7 @@
 #include <cuda_fp8.h>

 #include "cuda_utils.h"
-#include "libtorch_stable/launch_bounds_utils.h"
+#include "launch_bounds_utils.h"

 // Define before including nvfp4_utils.cuh so the header
 // can use this macro during compilation.
@@ -27,14 +27,14 @@
 #include <torch/csrc/stable/tensor.h>
 #include "libtorch_stable/torch_utils.h"
 #include "libtorch_stable/dispatch_utils.h"
-#include "../../cuda_vec_utils.cuh"
+#include "cuda_vec_utils.cuh"
 #include "cuda_utils.h"

 #include "nvfp4_utils.cuh"
 static_assert(CVT_FP4_ELTS_PER_THREAD == 16,
              "MXFP4 experts quant requires PACK16 mode (CUDA >= 12.9)");

-#include "libtorch_stable/launch_bounds_utils.h"
+#include "launch_bounds_utils.h"

 namespace vllm {

@@ -17,7 +17,7 @@
 #include <torch/csrc/stable/tensor.h>
 #include "libtorch_stable/torch_utils.h"
 #include "libtorch_stable/dispatch_utils.h"
-#include "../../cuda_vec_utils.cuh"
+#include "cuda_vec_utils.cuh"

 #include <cuda_runtime_api.h>
 #include <cuda_runtime.h>
@@ -26,7 +26,7 @@

 #include "cuda_utils.h"
 #include "nvfp4_utils.cuh"
-#include "libtorch_stable/launch_bounds_utils.h"
+#include "launch_bounds_utils.h"

 namespace vllm {

@@ -23,10 +23,10 @@

 #include "libtorch_stable/torch_utils.h"
 #include "libtorch_stable/dispatch_utils.h"
-#include "../../cuda_vec_utils.cuh"
+#include "cuda_vec_utils.cuh"

 #include "cuda_utils.h"
-#include "libtorch_stable/launch_bounds_utils.h"
+#include "launch_bounds_utils.h"

 // Define before including nvfp4_utils.cuh so the header
 // can use this macro during compilation.
@@ -20,7 +20,7 @@
 #include <cuda_fp8.h>
 #include <utility>

-#include "../../cuda_vec_utils.cuh"
+#include "cuda_vec_utils.cuh"

 #if defined(NVFP4_ENABLE_ELTS16) && defined(CUDA_VERSION) && \
    CUDA_VERSION >= 12090
@@ -7,11 +7,14 @@

 #include <torch/csrc/stable/ops.h>

-#include "ggml-common.h"
-#include "vecdotq.cuh"
-#include "dequantize.cuh"
-#include "mmvq.cuh"
-#include "mmq.cuh"
+// NOTE: These headers are intentionally kept in csrc/quantization/gguf/ (not
+// moved to libtorch_stable) to avoid unnecessary reformatting that would break
+// git rename detection and pollute blame history.
+#include "../../../quantization/gguf/ggml-common.h"
+#include "../../../quantization/gguf/vecdotq.cuh"
+#include "../../../quantization/gguf/dequantize.cuh"
+#include "../../../quantization/gguf/mmvq.cuh"
+#include "../../../quantization/gguf/mmq.cuh"
 #include "moe.cuh"
 #include "moe_vec.cuh"

@@ -7,7 +7,7 @@
 #include "torch_utils.h"

 #ifndef USE_ROCM
-  #include "persistent_topk.cuh"
+  #include "../persistent_topk.cuh"
 #endif

 namespace {
@@ -247,10 +247,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
  ops.def(
      "dsv3_fused_a_gemm(Tensor! output, Tensor mat_a, Tensor mat_b) -> ()");

-  // BF16/FP32 x FP32 -> FP32 router GEMM for H=3072, E=256, M<=32 (SM90+).
-  // conditionally compiled so impl registration is in source file
-  ops.def("fp32_router_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, "
@@ -478,33 +474,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
      "Tensor? initial_state_idx,"
      "Tensor? cu_chunk_seqlen,"
      "Tensor? last_chunk_indices) -> ()");
-
-  // Attention ops
-  // Compute the attention between an input query and the cached
-  // keys/values using PagedAttention.
-  ops.def(
-      "paged_attention_v1("
-      "    Tensor! out, Tensor query, Tensor key_cache,"
-      "    Tensor value_cache, int num_kv_heads, float scale,"
-      "    Tensor block_tables, Tensor seq_lens, int block_size,"
-      "    int max_seq_len, Tensor? alibi_slopes,"
-      "    str kv_cache_dtype, Tensor k_scale, Tensor v_scale,"
-      "    int tp_rank, int blocksparse_local_blocks,"
-      "    int blocksparse_vert_stride, int blocksparse_block_size,"
-      "    int blocksparse_head_sliding_step) -> ()");
-
-  // PagedAttention V2.
-  ops.def(
-      "paged_attention_v2("
-      "    Tensor! out, Tensor! exp_sums, Tensor! max_logits,"
-      "    Tensor! tmp_out, Tensor query, Tensor key_cache,"
-      "    Tensor value_cache, int num_kv_heads, float scale,"
-      "    Tensor block_tables, Tensor seq_lens, int block_size,"
-      "    int max_seq_len, Tensor? alibi_slopes,"
-      "    str kv_cache_dtype, Tensor k_scale, Tensor v_scale,"
-      "    int tp_rank, int blocksparse_local_blocks,"
-      "    int blocksparse_vert_stride, int blocksparse_block_size,"
-      "    int blocksparse_head_sliding_step) -> ()");
 }

 STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
@@ -612,9 +581,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
  ops.impl("ggml_moe_a8", TORCH_BOX(&ggml_moe_a8));
  ops.impl("ggml_moe_a8_vec", TORCH_BOX(&ggml_moe_a8_vec));
  ops.impl("selective_scan_fwd", TORCH_BOX(&selective_scan_fwd));
-
-  ops.impl("paged_attention_v1", TORCH_BOX(&paged_attention_v1));
-  ops.impl("paged_attention_v2", TORCH_BOX(&paged_attention_v2));
 }

 // These capability-check functions take only primitive args (no tensors), so
@@ -637,115 +603,4 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, ops) {
  ops.impl("ggml_moe_get_block_size", TORCH_BOX(&ggml_moe_get_block_size));
 }

-// Cache ops
-STABLE_TORCH_LIBRARY_FRAGMENT(_C_cache_ops, ops) {
-  // Swap in (out) the cache blocks from src to dst.
-  ops.def(
-      "swap_blocks(Tensor src, Tensor! dst,"
-      "            int block_size_in_bytes, Tensor block_mapping) -> ()");
-
-  // Batch swap: submit all block copies in a single driver call.
-  ops.def(
-      "swap_blocks_batch(Tensor src_ptrs, Tensor dst_ptrs,"
-      "                  Tensor sizes,"
-      "                  bool is_src_access_order_any=False) -> ()");
-
-  // Reshape the key and value tensors and cache them.
-  ops.def(
-      "reshape_and_cache(Tensor key, Tensor value,"
-      "                  Tensor! key_cache, Tensor! value_cache,"
-      "                  Tensor slot_mapping,"
-      "                  str kv_cache_dtype,"
-      "                  Tensor k_scale, Tensor v_scale) -> ()");
-
-  // Reshape the key and value tensors and cache them.
-  ops.def(
-      "reshape_and_cache_flash(Tensor key, Tensor value,"
-      "                        Tensor! key_cache,"
-      "                        Tensor! value_cache,"
-      "                        Tensor slot_mapping,"
-      "                        str kv_cache_dtype,"
-      "                        Tensor k_scale, Tensor v_scale) -> ()");
-
-  // Concat kv_c and k_pe and cache them.
-  ops.def(
-      "concat_and_cache_mla(Tensor kv_c, Tensor k_pe,"
-      "                     Tensor! kv_cache,"
-      "                     Tensor slot_mapping,"
-      "                     str kv_cache_dtype,"
-      "                     Tensor scale) -> ()");
-
-  // Rotate Q and K, then write to kv cache for MLA
-  ops.def(
-      "concat_and_cache_mla_rope_fused("
-      "                     Tensor positions,"
-      "                     Tensor! q_pe,"
-      "                     Tensor! k_pe,"
-      "                     Tensor kv_c,"
-      "                     Tensor cos_sin_cache,"
-      "                     bool is_neox,"
-      "                     Tensor slot_mapping,"
-      "                     Tensor! kv_cache,"
-      "                     str kv_cache_dtype,"
-      "                     Tensor kv_cache_scale) -> ()");
-
-  // Convert the key and value cache to fp8 data type.
-  ops.def(
-      "convert_fp8(Tensor! dst_cache, Tensor src_cache, float scale, "
-      "str kv_cache_dtype) -> ()");
-
-  // Gather cache blocks from src_cache to dst, dequantizing from
-  // src_cache's dtype to dst's dtype if necessary.
-  ops.def(
-      "gather_and_maybe_dequant_cache(Tensor src_cache, Tensor! dst, "
-      "                               Tensor block_table, Tensor cu_seq_lens, "
-      "                               Tensor token_to_seq, "
-      "                               int num_tokens, "
-      "                               str kv_cache_dtype, "
-      "                               Tensor scale, Tensor? seq_starts) -> ()");
-
-  ops.def(
-      "cp_gather_cache(Tensor src_cache, Tensor! dst, Tensor block_table, "
-      "Tensor cu_seq_lens, int batch_size, Tensor? seq_starts) -> ()");
-
-  ops.def(
-      "cp_gather_and_upconvert_fp8_kv_cache(Tensor src_cache, Tensor! dst, "
-      "Tensor block_table, Tensor seq_lens, Tensor workspace_starts, int "
-      "batch_size) -> ()");
-
-  ops.def(
-      "indexer_k_quant_and_cache(Tensor k, Tensor! kv_cache, Tensor "
-      "slot_mapping, "
-      "int quant_block_size, str kv_cache_dtype) -> ()");
-
-  ops.def("concat_mla_q(Tensor ql_nope, Tensor q_pe, Tensor! q_out) -> ()");
-
-  ops.def(
-      "cp_gather_indexer_k_quant_cache(Tensor kv_cache, Tensor! dst_k, Tensor! "
-      "dst_scale, Tensor block_table, Tensor cu_seq_lens) -> ()");
-}
-
-STABLE_TORCH_LIBRARY_IMPL(_C_cache_ops, CPU, ops) {
-  ops.impl("swap_blocks_batch", TORCH_BOX(&swap_blocks_batch));
-}
-
-STABLE_TORCH_LIBRARY_IMPL(_C_cache_ops, CUDA, ops) {
-  ops.impl("swap_blocks", TORCH_BOX(&swap_blocks));
-  ops.impl("reshape_and_cache", TORCH_BOX(&reshape_and_cache));
-  ops.impl("reshape_and_cache_flash", TORCH_BOX(&reshape_and_cache_flash));
-  ops.impl("concat_and_cache_mla", TORCH_BOX(&concat_and_cache_mla));
-  ops.impl("concat_and_cache_mla_rope_fused",
-           TORCH_BOX(&concat_and_cache_mla_rope_fused));
-  ops.impl("convert_fp8", TORCH_BOX(&convert_fp8));
-  ops.impl("gather_and_maybe_dequant_cache",
-           TORCH_BOX(&gather_and_maybe_dequant_cache));
-  ops.impl("cp_gather_cache", TORCH_BOX(&cp_gather_cache));
-  ops.impl("cp_gather_and_upconvert_fp8_kv_cache",
-           TORCH_BOX(&cp_gather_and_upconvert_fp8_kv_cache));
-  ops.impl("indexer_k_quant_and_cache", TORCH_BOX(&indexer_k_quant_and_cache));
-  ops.impl("concat_mla_q", TORCH_BOX(&concat_mla_q));
-  ops.impl("cp_gather_indexer_k_quant_cache",
-           TORCH_BOX(&cp_gather_indexer_k_quant_cache));
-}
-
 REGISTER_EXTENSION(_C_stable_libtorch)
@@ -17,8 +17,11 @@
 #define NVFP4_ENABLE_ELTS16 1
 #include "libtorch_stable/quantization/fp4/nvfp4_utils.cuh"

-#include "libtorch_stable/dispatch_utils.h"
-#include "libtorch_stable/torch_utils.h"
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
+#include <torch/all.h>
+
+#include "dispatch_utils.h"

 namespace vllm {

@@ -181,13 +184,12 @@ __global__ void reshape_and_cache_nvfp4_kernel(
 // Receives key_cache/value_cache as kv_cache[:, 0] and kv_cache[:, 1].
 // Each KV side contains both data and scale:
 //   page = [K_data | K_scale | V_data | V_scale]
-void reshape_and_cache_nvfp4_dispatch(torch::stable::Tensor& key,
-                                      torch::stable::Tensor& value,
-                                      torch::stable::Tensor& key_cache,
-                                      torch::stable::Tensor& value_cache,
-                                      torch::stable::Tensor& slot_mapping,
-                                      torch::stable::Tensor& k_scale,
-                                      torch::stable::Tensor& v_scale) {
+void reshape_and_cache_nvfp4_dispatch(torch::Tensor& key, torch::Tensor& value,
+                                      torch::Tensor& key_cache,
+                                      torch::Tensor& value_cache,
+                                      torch::Tensor& slot_mapping,
+                                      torch::Tensor& k_scale,
+                                      torch::Tensor& v_scale) {
  int num_tokens = slot_mapping.size(0);
  int num_heads = key.size(1);
  int head_size = key.size(2);
@@ -198,18 +200,17 @@ void reshape_and_cache_nvfp4_dispatch(torch::stable::Tensor& key,
  // key_cache is kv_cache[:, 0] with shape
  // [num_blocks, block_size, num_heads, full_dim] in logical order.
  // Strides encode the physical layout (HND or NHD).
-  STD_TORCH_CHECK(key_cache.dim() == 4, "key_cache must be 4D");
-  STD_TORCH_CHECK(key_cache.size(3) == full_dim,
-                  "key_cache last dim must be data_dim + scale_dim, got ",
-                  key_cache.size(3), " expected ", full_dim);
+  TORCH_CHECK(key_cache.dim() == 4, "key_cache must be 4D");
+  TORCH_CHECK(key_cache.size(3) == full_dim,
+              "key_cache last dim must be data_dim + scale_dim, got ",
+              key_cache.size(3), " expected ", full_dim);

  int block_size = key_cache.size(1);

-  STD_TORCH_CHECK(head_size % 16 == 0,
-                  "head_size must be divisible by 16 for NVFP4 KV cache");
-  STD_TORCH_CHECK(block_size % 4 == 0,
-                  "block_size must be divisible by 4 for NVFP4 KV cache "
-                  "swizzle");
+  TORCH_CHECK(head_size % 16 == 0,
+              "head_size must be divisible by 16 for NVFP4 KV cache");
+  TORCH_CHECK(block_size % 4 == 0,
+              "block_size must be divisible by 4 for NVFP4 KV cache swizzle");

  // Detect physical layout from strides (based on full_dim).
  // HND: head stride > block_offset stride.
@@ -229,9 +230,8 @@ void reshape_and_cache_nvfp4_dispatch(torch::stable::Tensor& key,
  // Scale follows data within each KV side.
  int64_t data_per_kv = (int64_t)num_heads * block_size * data_dim;

-  uint8_t* key_scale_ptr = key_cache.mutable_data_ptr<uint8_t>() + data_per_kv;
-  uint8_t* value_scale_ptr =
-      value_cache.mutable_data_ptr<uint8_t>() + data_per_kv;
+  uint8_t* key_scale_ptr = key_cache.data_ptr<uint8_t>() + data_per_kv;
+  uint8_t* value_scale_ptr = value_cache.data_ptr<uint8_t>() + data_per_kv;

  // Scale strides: same page stride, inner strides from layout.
  int64_t scale_block_stride = data_block_stride;
@@ -244,8 +244,8 @@ void reshape_and_cache_nvfp4_dispatch(torch::stable::Tensor& key,
    scale_block_offset_stride = (int64_t)num_heads * scale_dim;
  }

-  const float* k_scale_ptr = k_scale.const_data_ptr<float>();
-  const float* v_scale_ptr = v_scale.const_data_ptr<float>();
+  const float* k_scale_ptr = k_scale.data_ptr<float>();
+  const float* v_scale_ptr = v_scale.data_ptr<float>();

  int groups_per_head = head_size / CVT_FP4_SF_VEC_SIZE;
  int total_groups = num_heads * groups_per_head;
@@ -256,22 +256,20 @@ void reshape_and_cache_nvfp4_dispatch(torch::stable::Tensor& key,
  dim3 grid(num_tokens);
  dim3 block(num_threads);

-  const torch::stable::accelerator::DeviceGuard device_guard(
-      key.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(key));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

-  VLLM_STABLE_DISPATCH_HALF_TYPES(
+  AT_DISPATCH_REDUCED_FLOATING_TYPES(
      key.scalar_type(), "reshape_and_cache_nvfp4", [&] {
        vllm::reshape_and_cache_nvfp4_kernel<scalar_t>
            <<<grid, block, 0, stream>>>(
-                key.const_data_ptr<scalar_t>(),
-                value.const_data_ptr<scalar_t>(),
-                key_cache.mutable_data_ptr<uint8_t>(),
-                value_cache.mutable_data_ptr<uint8_t>(), key_scale_ptr,
-                value_scale_ptr, slot_mapping.const_data_ptr<int64_t>(),
-                k_scale_ptr, v_scale_ptr, key.stride(0), value.stride(0),
-                num_heads, head_size, block_size, data_block_stride,
-                data_head_stride, data_block_offset_stride, scale_block_stride,
-                scale_head_stride, scale_block_offset_stride);
+                key.data_ptr<scalar_t>(), value.data_ptr<scalar_t>(),
+                key_cache.data_ptr<uint8_t>(), value_cache.data_ptr<uint8_t>(),
+                key_scale_ptr, value_scale_ptr,
+                slot_mapping.data_ptr<int64_t>(), k_scale_ptr, v_scale_ptr,
+                key.stride(0), value.stride(0), num_heads, head_size,
+                block_size, data_block_stride, data_head_stride,
+                data_block_offset_stride, scale_block_stride, scale_head_stride,
+                scale_block_offset_stride);
      });
 }
@@ -31,6 +31,29 @@ torch::Tensor weak_ref_tensor(torch::Tensor& tensor) {
  return new_tensor;
 }

+void paged_attention_v1(
+    torch::Tensor& out, torch::Tensor& query, torch::Tensor& key_cache,
+    torch::Tensor& value_cache, int64_t num_kv_heads, double scale,
+    torch::Tensor& block_tables, torch::Tensor& seq_lens, int64_t block_size,
+    int64_t max_seq_len, const std::optional<torch::Tensor>& alibi_slopes,
+    const std::string& kv_cache_dtype, torch::Tensor& k_scale,
+    torch::Tensor& v_scale, const int64_t tp_rank,
+    const int64_t blocksparse_local_blocks,
+    const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
+    const int64_t blocksparse_head_sliding_step);
+
+void paged_attention_v2(
+    torch::Tensor& out, torch::Tensor& exp_sums, torch::Tensor& max_logits,
+    torch::Tensor& tmp_out, torch::Tensor& query, torch::Tensor& key_cache,
+    torch::Tensor& value_cache, int64_t num_kv_heads, double scale,
+    torch::Tensor& block_tables, torch::Tensor& seq_lens, int64_t block_size,
+    int64_t max_seq_len, const std::optional<torch::Tensor>& alibi_slopes,
+    const std::string& kv_cache_dtype, torch::Tensor& k_scale,
+    torch::Tensor& v_scale, const int64_t tp_rank,
+    const int64_t blocksparse_local_blocks,
+    const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
+    const int64_t blocksparse_head_sliding_step);
+
 // rms_norm and fused_add_rms_norm declarations also exist in
 // csrc/libtorch_stable/ops.h (torch::stable ABI for CUDA). They remain here
 // because the CPU build still uses these torch::Tensor declarations.
@@ -6,7 +6,6 @@
 #include <hip/hip_bfloat16.h>

 #include "../../../../attention/attention_dtypes.h"
-#include <torch/headeronly/core/ScalarType.h>

 namespace vllm {
 #ifdef USE_ROCM
@@ -643,29 +642,27 @@ __inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) {
    vllm::Fp8KVCacheDataType KV_CACHE_DTYPE =                                  \
        vllm::get_fp8_kv_cache_data_type(KV_DTYPE);                            \
    if (KV_CACHE_DTYPE == vllm::Fp8KVCacheDataType::kAuto) {                   \
-      if (SRC_DTYPE == torch::headeronly::ScalarType::Float) {                 \
+      if (SRC_DTYPE == at::ScalarType::Float) {                                \
        FN(float, float, vllm::Fp8KVCacheDataType::kAuto);                     \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::Half) {           \
+      } else if (SRC_DTYPE == at::ScalarType::Half) {                          \
        FN(uint16_t, uint16_t, vllm::Fp8KVCacheDataType::kAuto);               \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::BFloat16) {       \
+      } else if (SRC_DTYPE == at::ScalarType::BFloat16) {                      \
        FN(__nv_bfloat16, __nv_bfloat16, vllm::Fp8KVCacheDataType::kAuto);     \
      } else {                                                                 \
-        STD_TORCH_CHECK(false,                                                 \
-                        "Unsupported input type of kv cache: ", SRC_DTYPE);    \
+        TORCH_CHECK(false, "Unsupported input type of kv cache: ", SRC_DTYPE); \
      }                                                                        \
    } else if (KV_CACHE_DTYPE == vllm::Fp8KVCacheDataType::kFp8E4M3) {         \
-      if (SRC_DTYPE == torch::headeronly::ScalarType::Float) {                 \
+      if (SRC_DTYPE == at::ScalarType::Float) {                                \
        FN(float, uint8_t, vllm::Fp8KVCacheDataType::kFp8E4M3);                \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::Half) {           \
+      } else if (SRC_DTYPE == at::ScalarType::Half) {                          \
        FN(uint16_t, uint8_t, vllm::Fp8KVCacheDataType::kFp8E4M3);             \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::BFloat16) {       \
+      } else if (SRC_DTYPE == at::ScalarType::BFloat16) {                      \
        FN(__nv_bfloat16, uint8_t, vllm::Fp8KVCacheDataType::kFp8E4M3);        \
      } else {                                                                 \
-        STD_TORCH_CHECK(false,                                                 \
-                        "Unsupported input type of kv cache: ", SRC_DTYPE);    \
+        TORCH_CHECK(false, "Unsupported input type of kv cache: ", SRC_DTYPE); \
      }                                                                        \
    } else {                                                                   \
-      STD_TORCH_CHECK(false, "Unsupported data type of kv cache: ", KV_DTYPE); \
+      TORCH_CHECK(false, "Unsupported data type of kv cache: ", KV_DTYPE);     \
    }

 }  // namespace fp8
@@ -1,7 +1,6 @@
 #pragma once

 #include "../../../../attention/attention_dtypes.h"
-#include <torch/headeronly/core/ScalarType.h>
 #include <assert.h>
 #include <float.h>
 #include <stdint.h>
@@ -547,40 +546,37 @@ __inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) {
    vllm::Fp8KVCacheDataType KV_CACHE_DTYPE =                                  \
        vllm::get_fp8_kv_cache_data_type(KV_DTYPE);                            \
    if (KV_CACHE_DTYPE == vllm::Fp8KVCacheDataType::kAuto) {                   \
-      if (SRC_DTYPE == torch::headeronly::ScalarType::Float) {                 \
+      if (SRC_DTYPE == at::ScalarType::Float) {                                \
        FN(float, float, vllm::Fp8KVCacheDataType::kAuto);                     \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::Half) {           \
+      } else if (SRC_DTYPE == at::ScalarType::Half) {                          \
        FN(uint16_t, uint16_t, vllm::Fp8KVCacheDataType::kAuto);               \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::BFloat16) {       \
+      } else if (SRC_DTYPE == at::ScalarType::BFloat16) {                      \
        FN(__nv_bfloat16, __nv_bfloat16, vllm::Fp8KVCacheDataType::kAuto);     \
      } else {                                                                 \
-        STD_TORCH_CHECK(false,                                                 \
-                        "Unsupported input type of kv cache: ", SRC_DTYPE);    \
+        TORCH_CHECK(false, "Unsupported input type of kv cache: ", SRC_DTYPE); \
      }                                                                        \
    } else if (KV_CACHE_DTYPE == vllm::Fp8KVCacheDataType::kFp8E4M3) {         \
-      if (SRC_DTYPE == torch::headeronly::ScalarType::Float) {                 \
+      if (SRC_DTYPE == at::ScalarType::Float) {                                \
        FN(float, uint8_t, vllm::Fp8KVCacheDataType::kFp8E4M3);                \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::Half) {           \
+      } else if (SRC_DTYPE == at::ScalarType::Half) {                          \
        FN(uint16_t, uint8_t, vllm::Fp8KVCacheDataType::kFp8E4M3);             \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::BFloat16) {       \
+      } else if (SRC_DTYPE == at::ScalarType::BFloat16) {                      \
        FN(__nv_bfloat16, uint8_t, vllm::Fp8KVCacheDataType::kFp8E4M3);        \
      } else {                                                                 \
-        STD_TORCH_CHECK(false,                                                 \
-                        "Unsupported input type of kv cache: ", SRC_DTYPE);    \
+        TORCH_CHECK(false, "Unsupported input type of kv cache: ", SRC_DTYPE); \
      }                                                                        \
    } else if (KV_CACHE_DTYPE == vllm::Fp8KVCacheDataType::kFp8E5M2) {         \
-      if (SRC_DTYPE == torch::headeronly::ScalarType::Float) {                 \
+      if (SRC_DTYPE == at::ScalarType::Float) {                                \
        FN(float, uint8_t, vllm::Fp8KVCacheDataType::kFp8E5M2);                \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::Half) {           \
+      } else if (SRC_DTYPE == at::ScalarType::Half) {                          \
        FN(uint16_t, uint8_t, vllm::Fp8KVCacheDataType::kFp8E5M2);             \
-      } else if (SRC_DTYPE == torch::headeronly::ScalarType::BFloat16) {       \
+      } else if (SRC_DTYPE == at::ScalarType::BFloat16) {                      \
        FN(__nv_bfloat16, uint8_t, vllm::Fp8KVCacheDataType::kFp8E5M2);        \
      } else {                                                                 \
-        STD_TORCH_CHECK(false,                                                 \
-                        "Unsupported input type of kv cache: ", SRC_DTYPE);    \
+        TORCH_CHECK(false, "Unsupported input type of kv cache: ", SRC_DTYPE); \
      }                                                                        \
    } else {                                                                   \
-      STD_TORCH_CHECK(false, "Unsupported data type of kv cache: ", KV_DTYPE); \
+      TORCH_CHECK(false, "Unsupported data type of kv cache: ", KV_DTYPE);     \
    }

 }  // namespace fp8
@@ -18,15 +18,6 @@ void wvSplitKQ(const at::Tensor& in_a, const at::Tensor& in_b,
               const at::Tensor& scale_a, const at::Tensor& scale_b,
               const int64_t CuCount);

-torch::Tensor gptq_gemm_rdna3(torch::Tensor a, torch::Tensor b_q_weight,
-                              torch::Tensor b_qzeros, torch::Tensor b_scales,
-                              torch::Tensor b_g_idx, bool use_v2_format);
-
-torch::Tensor gptq_gemm_rdna3_wmma(torch::Tensor a, torch::Tensor b_q_weight,
-                                   torch::Tensor b_qzeros,
-                                   torch::Tensor b_scales,
-                                   torch::Tensor b_g_idx, bool use_v2_format);
-
 void paged_attention(
    torch::Tensor& out, torch::Tensor& exp_sums, torch::Tensor& max_logits,
    torch::Tensor& tmp_out, torch::Tensor& query, torch::Tensor& key_cache,
@@ -1,780 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-//
-// W4A16 GPTQ kernel for RDNA3 (gfx1100 / RX 7900 XTX class), templated on the
-// activation dtype (half or __hip_bfloat16). Adapted from exllamav2's 4-bit
-// kernel (csrc/quantization/gptq/q_gemm.cu) with the following changes:
-//
-//   1. Direct write to the T-typed output via packed CAS-loop on a 64-bit
-//      word (atomic_add_pk4_{f16,bf16}). gfx11 has no native
-//      v_global_atomic_pk_add_{f16,bf16}, so the kernel emulates one with
-//      global_atomic_cmpswap_b64. This avoids the M*N*4-byte FP32 scratch
-//      buffer + memset + cast-pass that an fp32-accumulator design would
-//      need; the caller passes a zero-initialised T-typed output tensor
-//      and every block atomically adds its partial sum into it.
-//
-//   2. The bf16 path uses a dedicated bit-trick that avoids the fp16-only
-//      "upper nibble * 16" trick, which would overflow the 7-bit bf16
-//      mantissa. See qdq_4_rdna3.cuh for details.
-//
-//   3. Wave32 geometry sized for high CU saturation: THREADS_X=256
-//      (8 waves per block) and BLOCK_KN_SIZE=256, with each thread
-//      computing 4 N output columns. gridDim.z = K / BLOCK_KN_SIZE
-//      splits K and the output is atomically accumulated. fp16 uses
-//      v_dot2_f32_f16 (__builtin_amdgcn_fdot2) for the inner dot;
-//      bf16 widens to fp32 (no v_pk_fma_bf16 on gfx11) and accumulates
-//      with v_fma_f32. M_COUNT ∈ {1,2,4,8} is selected at launch
-//      based on size_m.
-//
-//   4. The bf16 dispatch with M >= 16 forwards to the WMMA kernel in
-//      q_gemm_rdna3_wmma.cu (separate translation unit) where
-//      v_wmma_f32_16x16x16_bf16_w32 wins. The fp16 path always stays
-//      scalar (the bit-trick dequant beats WMMA below M=64).
-
-#include <cstdint>
-#include <cstdio>
-
-#include <torch/all.h>
-#include <c10/cuda/CUDAGuard.h>
-#include <ATen/cuda/CUDAContext.h>
-
-#include <hip/hip_runtime.h>
-#include <hip/hip_bf16.h>
-#include <hip/hip_fp16.h>
-
-#include "qdq_4_rdna3.cuh"
-
-#if defined(__HIPCC__) && defined(__gfx1100__)
-  #define __HIP__RDNA3__
-#endif
-
-namespace vllm {
-namespace gptq_rdna3 {
-
-// BLOCK_KN_SIZE = 256 (was 128 in exllama). Each block covers 256 K
-// elements and THREADS_X*4 = 1024 N columns. For Qwen-class K=4096 this
-// halves gridDim.z (32 → 16) and therefore halves the atomic count per
-// output position vs the exllama default. THREADS_X=256 = 8 waves on RDNA3
-// wave32; with ~32 wave slots per CU we still fit 4 blocks per CU at peak.
-//
-// We tried BLOCK_KN_SIZE=512 (microbench on Qwen3.6-27B): bf16 improved
-// 5-10% at large M (atomic CAS halved), but fp16 decode regressed up to
-// +40% on qkv-square (32 → 45 μs at M=1). Cause: 16 waves/block × 16
-// total blocks for [M=1, K=N=4096] only saturates ~8 of the 96 CUs,
-// breaking memory-latency hiding for the fp16 path which is already
-// memory-bound. Reverted to 256; bf16 keeps most of its gains from the
-// fp32 dequant rewrite alone.
-#define BLOCK_KN_SIZE 256
-#define THREADS_X 256
-
-// Device code below is RDNA3-only; non-RDNA3 device passes fall through to
-// the empty __global__ stub at the #else below for symbol parity.
-#if defined(__HIP__RDNA3__) || !defined(__HIP_DEVICE_COMPILE__)
-
-// ---------------------------------------------------------------------------
-// Per-dtype helpers. We avoid heavy template metaprogramming and just provide
-// overloaded inline functions; the kernel below selects via `if constexpr`.
-// ---------------------------------------------------------------------------
-
-// Type-generic zero — both half and bf16_t in HIP/ROCm have a converting
-// constructor from float, but going through __float2half_rn / __float2bfloat16
-// is the unambiguously correct path on every ROCm version.
-template <typename T>
-__forceinline__ __device__ T tzero();
-
-template <>
-__forceinline__ __device__ half tzero<half>() {
-  return __float2half_rn(0.0f);
-}
-
-template <>
-__forceinline__ __device__ bf16_t tzero<bf16_t>() {
-  return __float2bfloat16(0.0f);
-}
-
-__forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) {
-  // RDNA3 has v_dot2_f32_f16 (`__builtin_amdgcn_fdot2`) which computes
-  // fp32 += a.x*b.x + a.y*b.y in a single instruction with the accumulator
-  // staying in fp32 throughout. hipcc 7.2 does NOT peephole the obvious
-  // `__hfma2 + cast + add` pattern into v_dot2 (verified by ISA
-  // disassembly: 0 v_dot2_f32_f16 vs 256 v_cvt_f32_f16 + 218 v_add_f32 in
-  // the M_COUNT=8 kernel before this change), so we issue the builtin
-  // explicitly. Saves the trailing 2× v_cvt_f32_f16 + v_add_f32 (3 ops)
-  // per dot22_8_f call vs the half2-accumulator form. With 128 calls per
-  // K=32 step that's ~384 ops/K-step less issue pressure on the VALU.
-  //
-  // Numerical bonus: accumulator stays fp32 throughout the dot. The old
-  // form accumulated 8 muladds in fp16 (10-bit mantissa) before casting,
-  // which could lose ~3 bits of precision on borderline magnitudes.
-  float result = 0.0f;
-  const half2* a2_ptr = (const half2*)a_ptr;
-  #pragma unroll
-  for (int i = 0; i < 4; i++) {
-    result = __builtin_amdgcn_fdot2(dq[i], *a2_ptr++, result, /*clamp=*/false);
-  }
-  return result;
-}
-
-__forceinline__ __device__ float dot22_8_f(bf162_t (&dq)[4],
-                                           const bf16_t* a_ptr) {
-  // RDNA3 (gfx1100) lacks a packed bf16 FMA: there is no v_pk_fma_bf16 in
-  // the gfx11 ISA (it only landed on CDNA3+ / gfx94x and later). hipcc
-  // therefore lowers __hfma2(bf162_t, bf162_t, bf162_t) to a serialised
-  // fallback (single-element FMAs or fp32 round-trips), which empirically
-  // runs ~2× the cycle count of v_pk_fma_f16 on the same VALU. The bf16
-  // decode path was paying that tax in full, scaling linearly with M (the
-  // fp16 path scales sub-linearly because its v_pk_fma_f16 is full rate
-  // and the kernel becomes memory-bound).
-  //
-  // Fix: widen bf16 → fp32 explicitly (a left-shift by 16, free in VGPRs)
-  // and accumulate with v_fma_f32, which IS full rate on RDNA3. Same FMA
-  // count, but each FMA is fast. Bonus: the accumulator is now fp32
-  // throughout instead of bf16, which is also numerically more accurate
-  // (no compounding bf16-rounding inside the dot loop).
-  float result = 0.0f;
-  #pragma unroll
-  for (int i = 0; i < 4; i++) {
-    uint32_t aw, dw;
-    __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
-    __builtin_memcpy(&dw, &dq[i], sizeof(uint32_t));
-    // bf16 in low 16 bits  → fp32 by left-shifting into the upper half.
-    // bf16 in high 16 bits → already aligned with fp32's upper half.
-    float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
-    float a_y = __uint_as_float(aw & 0xFFFF0000u);
-    float d_x = __uint_as_float((dw & 0xFFFFu) << 16);
-    float d_y = __uint_as_float(dw & 0xFFFF0000u);
-    result = __fmaf_rn(d_x, a_x, result);
-    result = __fmaf_rn(d_y, a_y, result);
-  }
-  return result;
-}
-
-// fp32-input dot product: paired with dequant_4bit_8_bf16_f32 which already
-// produces fp32 dq[8]. Saves the bf16→fp32 widening that the bf162_t
-// overload above does for dq (still need to widen A from bf16). Wins more
-// at high N: the bf162_t version's per-call widening cost scales with the
-// number of dequants × M_COUNT × 4 dot calls; the fp32 version pays only
-// for A widening (M_COUNT × 4 × 4 widens, half as many).
-__forceinline__ __device__ float dot22_8_f(float (&dq)[8],
-                                           const bf16_t* a_ptr) {
-  float result = 0.0f;
-  #pragma unroll
-  for (int i = 0; i < 4; i++) {
-    uint32_t aw;
-    __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
-    float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
-    float a_y = __uint_as_float(aw & 0xFFFF0000u);
-    result = __fmaf_rn(dq[2 * i + 0], a_x, result);
-    result = __fmaf_rn(dq[2 * i + 1], a_y, result);
-  }
-  return result;
-}
-
-// ---------------------------------------------------------------------------
-// Packed atomic-add via CAS-loop on a 64-bit word (4 fp16/bf16 lanes per CAS).
-// RDNA3 (gfx11) does NOT have native v_global_atomic_pk_add_f16 / _bf16 (those
-// landed on gfx940 / gfx1250 respectively), so this lowers to
-// global_atomic_cmpswap_b64 plus retry. We use this in the kernel epilogue to
-// write 4 output columns per row in a single atomic operation — half the
-// atomic instruction count and half the contention vs two 32-bit CAS calls.
-//
-// Writing directly to fp16/bf16 (instead of through an FP32 scratch buffer +
-// cast pass) saves M*N*4 bytes of allocation, the memset, and the epilogue
-// cast pass that an fp32-accumulator design would need.
-//
-// 64-bit alignment: the kernel writes at `out + n` where n = offset_n + t*4
-// (always multiple of 4), and partition_weight_shape[1] is required to be a
-// multiple of 8 by can_implement(), so every (m, n) write target is 8-byte
-// aligned. Required by global_atomic_cmpswap_b64.
-// ---------------------------------------------------------------------------
-
-__forceinline__ __device__ void atomic_add_pk4_f16(half* addr, half2 v01,
-                                                   half2 v23) {
-  unsigned long long* addr_u = reinterpret_cast<unsigned long long*>(addr);
-  unsigned long long old = *addr_u;
-  while (true) {
-    union {
-      unsigned long long u;
-      half2 h2[2];
-    } cur, sum;
-    cur.u = old;
-    sum.h2[0] = __hadd2(cur.h2[0], v01);
-    sum.h2[1] = __hadd2(cur.h2[1], v23);
-    unsigned long long prev = atomicCAS(addr_u, old, sum.u);
-    if (prev == old) break;
-    old = prev;
-  }
-}
-
-__forceinline__ __device__ void atomic_add_pk4_bf16(bf16_t* addr, bf162_t v01,
-                                                    bf162_t v23) {
-  unsigned long long* addr_u = reinterpret_cast<unsigned long long*>(addr);
-  unsigned long long old = *addr_u;
-  while (true) {
-    union {
-      unsigned long long u;
-      bf162_t b2[2];
-    } cur, sum;
-    cur.u = old;
-    sum.b2[0] = __hadd2(cur.b2[0], v01);
-    sum.b2[1] = __hadd2(cur.b2[1], v23);
-    unsigned long long prev = atomicCAS(addr_u, old, sum.u);
-    if (prev == old) break;
-    old = prev;
-  }
-}
-
-// Load one row's worth of 4 packed zeros (column n..n+3) from a [groups, N/8]
-// uint32 tensor. n is a multiple of 4 by construction (n = offset_n + t*4 with
-// offset_n = blockIdx.x * 512), so the 4 nibbles always live within one or two
-// uint32 words; in practice within one because n & 7 is 0 or 4.
-__forceinline__ __device__ void load4_zeros(const uint32_t* qzeros_row, int n,
-                                            int (&zeros)[4]) {
-  int qcol = n / 8;
-  int shift = (n & 0x07) * 4;
-  uint32_t d = qzeros_row[qcol] >> shift;
-  zeros[0] = (int)(d & 0xF);
-  zeros[1] = (int)((d >> 4) & 0xF);
-  zeros[2] = (int)((d >> 8) & 0xF);
-  zeros[3] = (int)((d >> 12) & 0xF);
-}
-
-template <typename T>
-__forceinline__ __device__ void load4_scales(const T* scales_row, int n,
-                                             T (&scales)[4]) {
-  scales[0] = scales_row[n + 0];
-  scales[1] = scales_row[n + 1];
-  scales[2] = scales_row[n + 2];
-  scales[3] = scales_row[n + 3];
-}
-
-// ---------------------------------------------------------------------------
-// Main kernel.
-// ---------------------------------------------------------------------------
-
-template <typename T, int M_COUNT>
-__global__ void gemm_q4_kernel_rdna3(
-    const T* __restrict__ a, const uint32_t* __restrict__ b_q_weight,
-    const uint32_t* __restrict__ b_qzeros, const T* __restrict__ b_scales,
-    T* __restrict__ c, const int size_m, const int size_n, const int size_k,
-    const int groups, const int zero_offset, const int* __restrict__ b_q_perm) {
-  const int t = threadIdx.x;
-  const int offset_n = blockIdx.x * BLOCK_KN_SIZE * 4;
-  const int offset_m = blockIdx.y * M_COUNT;
-  const int offset_k = blockIdx.z * BLOCK_KN_SIZE;
-  const int end_k = min(offset_k + BLOCK_KN_SIZE, size_k);
-  const int n = offset_n + t * 4;
-
-  // LDS layout: [M_COUNT][BLOCK_KN_SIZE + LDS_PAD]. The PAD=8 elements per M
-  // row break the natural 256-element/512-byte alignment that would otherwise
-  // collide on the same LDS bank when a thread reads block_a[0..M_COUNT-1][k]
-  // (same k, different m). Row stride becomes 264 elements * 2B = 528B = 132
-  // 4-byte banks, so m-stride hits banks (m*132)%32 = (m*4)%32 — distinct for
-  // all M_COUNT ≤ 8. Cost: 16B LDS per block, irrelevant.
-  constexpr int LDS_PAD = 8;
-  __shared__ T block_a[M_COUNT][BLOCK_KN_SIZE + LDS_PAD];
-
-  // Stage A: each thread loads 1 K element per M row into LDS (with optional
-  // act-order permutation). THREADS_X == BLOCK_KN_SIZE so this is a 1:1 map.
-  // For M_COUNT > 1 with size_m not a multiple of M_COUNT, slots past size_m
-  // are zero-padded so the dot product contribution is 0 (we then skip the
-  // atomic write for those rows below).
-  //
-  // M=1 fast path: skip LDS staging + __syncthreads entirely. All 256 threads
-  // read the SAME 8-element A window per inner step (a_off is uniform across
-  // the block), so the cache-line broadcast through L1 makes global reads as
-  // cheap as LDS reads. Measured: ~1% on 4B b=1, ~6% on 27B b=1 in=128.
-  static_assert(BLOCK_KN_SIZE == THREADS_X,
-                "BLOCK_KN_SIZE must equal THREADS_X (1 K element per thread)");
-  // The M=1 fast path (skip LDS) only has a global-read code path for bf16
-  // (the v_dot2_f32_bf16 branch).  The fp16 inner loop still indexes
-  // block_a[m][a_off] unconditionally, so for fp16 we MUST stage A through
-  // LDS even at M=1 to avoid reading uninitialized shared memory.
-  constexpr bool USE_LDS_A = (M_COUNT > 1) || std::is_same<T, half>::value;
-  if constexpr (USE_LDS_A) {
-    if (offset_k + t < end_k) {
-  #pragma unroll
-      for (int m = 0; m < M_COUNT; ++m) {
-        T av;
-        if (offset_m + m < size_m) {
-          const T* a_row = a + (offset_m + m) * size_k;
-          if (b_q_perm)
-            av = a_row[b_q_perm[offset_k + t]];
-          else
-            av = a_row[offset_k + t];
-        } else {
-          av = tzero<T>();  // zero-pad invalid M rows
-        }
-        block_a[m][t] = av;
-      }
-    }
-
-    // Threads beyond the right edge of N have nothing to do. Note: we must NOT
-    // return before __syncthreads() if any thread in the block participates in
-    // the LDS load above — but here all THREADS_X (=256) threads always do,
-    // regardless of whether their `n` is in bounds.
-    __syncthreads();
-  } else if (b_q_perm) {
-    // bf16 M=1 fast path skips LDS, but its global read below is sequential
-    // and cannot apply act-order. When a permutation is present, stage the
-    // single A row through LDS (as fp16 / M>1 do) so the read picks it up.
-    // b_q_perm is block-uniform, so the __syncthreads is non-divergent.
-    if (offset_k + t < end_k)
-      block_a[0][t] = a[offset_m * size_k + b_q_perm[offset_k + t]];
-    __syncthreads();
-  }
-  if (n >= size_n) return;
-
-  // Group bookkeeping. We require size_k % groups == 0 (groupsize divides K).
-  const int groupsize = size_k / groups;
-  int group = offset_k / groupsize;
-  int nextgroup = (group + 1) * groupsize;
-
-  // qweight stride: weights are [K/8, N] uint32 with K packed at dim 0.
-  int qk = offset_k / 8;
-  const uint32_t* b_ptr = b_q_weight + qk * size_n + n;
-
-  // Per-column dequant constants. We hold one set of (z, y) pairs per column.
-  // fp16 uses the exllama (z1z16, y1y16) double-pair to enable the upper-
-  // nibble-*16 trick. bf16 uses fp32 scalars (z, y) because the dequant
-  // produces fp32 directly — see prep_zero_scale_bf16_f32 / the FMA
-  // bypass for the missing v_pk_fma_bf16 on gfx11.
-  half2 z1z16_h[4][2], y1y16_h[4][2];
-  float z_b_f[4], y_b_f[4];
-
-  auto refresh_group = [&](int g) {
-    const uint32_t* qz_row = b_qzeros + g * (size_n / 8);
-    const T* sc_row = b_scales + g * size_n;
-    int zeros[4];
-    T scales[4];
-    load4_zeros(qz_row, n, zeros);
-    load4_scales<T>(sc_row, n, scales);
-    if constexpr (std::is_same<T, half>::value) {
-  #pragma unroll
-      for (int i = 0; i < 4; ++i) {
-        prep_zero_scale_fp16((uint32_t)(zeros[i] + zero_offset), scales[i],
-                             z1z16_h[i], y1y16_h[i]);
-      }
-    } else {
-  #pragma unroll
-      for (int i = 0; i < 4; ++i) {
-        prep_zero_scale_bf16_f32((uint32_t)(zeros[i] + zero_offset), scales[i],
-                                 z_b_f[i], y_b_f[i]);
-      }
-    }
-  };
-
-  refresh_group(group);
-
-  float block_c[M_COUNT][4];
-  #pragma unroll
-  for (int m = 0; m < M_COUNT; ++m) {
-  #pragma unroll
-    for (int j = 0; j < 4; ++j) block_c[m][j] = 0.0f;
-  }
-
-  // Note on group-transition granularity: we check `k == nextgroup` at the
-  // start of each outer iteration (which advances K by 32). This is correct
-  // when group_size >= 32 OR group_size divides 32 evenly (groupsize is one
-  // of {1,2,4,8,16,32,64,128,...}). For group_size in {16, 8, 4, ...} the
-  // inner loop would cross a group boundary between j-iterations; we require
-  // group_size >= 32 here, mirroring exllama's assumption.
-  //
-  // Software pipelining: we issue all 4 vectorized weight loads up front
-  // before any dequant/FMA depends on them. This gives the AMDGPU backend
-  // freedom to schedule the global_loads early and overlap their latency
-  // with dequant + v_pk_fma_f16 of earlier iterations. Cost: 4×int4 = 16
-  // VGPRs in flight per thread, plenty of headroom on RDNA3.
-  int k = offset_k;
-  while (k < end_k) {
-    if (k == nextgroup) {
-      group++;
-      nextgroup += groupsize;
-      refresh_group(group);
-    }
-
-    // Prefetch all four j-iterations' weight words. The compiler emits 4
-    // global_load_b128 instructions back-to-back; the dependent dequant +
-    // FMA work below hides their latency.
-    int4 b_w[4];
-  #pragma unroll
-    for (int j = 0; j < 4; ++j) {
-      b_w[j] = *(const int4*)(b_ptr + j * size_n);
-    }
-    b_ptr += 4 * size_n;
-
-  #pragma unroll
-    for (int j = 0; j < 4; ++j) {
-      const int a_off = (k - offset_k) + 8 * j;
-
-      if constexpr (std::is_same<T, half>::value) {
-        half2 dq[4][4];
-        dequant_4bit_8_fp16((uint32_t)b_w[j].x, dq[0], z1z16_h[0], y1y16_h[0]);
-        dequant_4bit_8_fp16((uint32_t)b_w[j].y, dq[1], z1z16_h[1], y1y16_h[1]);
-        dequant_4bit_8_fp16((uint32_t)b_w[j].z, dq[2], z1z16_h[2], y1y16_h[2]);
-        dequant_4bit_8_fp16((uint32_t)b_w[j].w, dq[3], z1z16_h[3], y1y16_h[3]);
-
-  #pragma unroll
-        for (int m = 0; m < M_COUNT; ++m) {
-          const half* a_ptr = reinterpret_cast<const half*>(&block_a[m][a_off]);
-          block_c[m][0] += dot22_8_f(dq[0], a_ptr);
-          block_c[m][1] += dot22_8_f(dq[1], a_ptr);
-          block_c[m][2] += dot22_8_f(dq[2], a_ptr);
-          block_c[m][3] += dot22_8_f(dq[3], a_ptr);
-        }
-      } else if constexpr (M_COUNT == 1) {
-        // bf16 decode (M=1), v_dot2_f32_bf16 path. Mirrors the data-flow of
-        // Hybrid PR #40977's wvSplitK_int4 kernel exactly so clang's
-        // InstCombine cannot fold the bf16→fp32 widening (LLVM #76000):
-        //   * activations and magic-value weights share a fp32-aliased
-        //     union (bytes written as uint32, read as bf16x2_t for the
-        //     dot — pointer-cast opacity defeats the fold)
-        //   * sum_a computed via a *second* v_dot2 with bf162(1,1) as the
-        //     second operand, avoiding any explicit bf16→fp32 widen of A
-        //   * bias correction y_b_f * partial + z_b_f * sum_a, identical
-        //     to the previous fp32-FMA-chain path
-        //
-        // Net: 20 v_dot2_f32_bf16 + 8 fp32 FMA per int32 weight vs the
-        // previous 40 fp32 FMA. v_dot2 runs at full rate on gfx1100, so
-        // the substitution is ~2× cheaper for the inner accumulator.
-        typedef short __attribute__((ext_vector_type(2))) bf16x2_t;
-        constexpr uint32_t BF16_MAGIC = 0x43004300u;  // bf162(128, 128)
-        constexpr uint32_t BF16_ONES = 0x3F803F80u;   // bf162(1.0, 1.0)
-        union pack4 {
-          float f[4];
-          uint32_t u[4];
-        };
-
-        uint32_t w[4];
-        __builtin_memcpy(w, &b_w[j], sizeof(int4));
-
-        // Load 8 bf16 activations as 4 uint32s (= 4 bf16x2 pairs) into a
-        // fp32-aliased union. Storing as uint32 keeps the IR-level type
-        // opaque so the inner v_dot2 cannot be folded to fp32 widening.
-        //
-        // A is read direct from global (no LDS staging — see USE_LDS_A above),
-        // except under act-order, where it comes from the permuted LDS copy.
-        pack4 a_pack;
-        {
-          const uint32_t* a_words =
-              b_q_perm
-                  ? reinterpret_cast<const uint32_t*>(&block_a[0][a_off])
-                  : reinterpret_cast<const uint32_t*>(a + offset_k + a_off);
-          a_pack.u[0] = a_words[0];
-          a_pack.u[1] = a_words[1];
-          a_pack.u[2] = a_words[2];
-          a_pack.u[3] = a_words[3];
-        }
-
-        // sum_a = Σ a[i]. Computed via 4× v_dot2_f32_bf16 with bf162(1,1) as
-        // the second operand — every bf16 pair contributes 1·a_lo + 1·a_hi.
-        // No fp32 widening of activations: the bytes go straight from LDS
-        // through v_dot2 into the fp32 accumulator.
-        float sum_a = 0.0f;
-  #pragma unroll
-        for (int b = 0; b < 4; ++b) {
-          sum_a = __builtin_amdgcn_fdot2_f32_bf16(
-              *((bf16x2_t*)(&a_pack.f[b])), *((const bf16x2_t*)&BF16_ONES),
-              sum_a, /*clamp=*/false);
-        }
-
-        // unroll 1 keeps q_pack alive only one col at a time (8 fp32 VGPRs
-        // recycled across cols), avoiding straight-line expansion that
-        // would inflate live-range to 32 VGPRs.
-  #pragma unroll 1
-        for (int col = 0; col < 4; ++col) {
-          // Build dequant magic values bf16(128 + nibble) directly into a
-          // fp32-aliased union via uint32 stores. No fp32 in the data flow
-          // until v_dot2 consumes the bytes.
-          pack4 q_pack;
-          const uint32_t qa = w[col];
-          q_pack.u[0] = ((qa >> 0) & 0x000F000Fu) | BF16_MAGIC;
-          q_pack.u[1] = ((qa >> 4) & 0x000F000Fu) | BF16_MAGIC;
-          q_pack.u[2] = ((qa >> 8) & 0x000F000Fu) | BF16_MAGIC;
-          q_pack.u[3] = ((qa >> 12) & 0x000F000Fu) | BF16_MAGIC;
-
-          // partial = Σ (128 + nibble[i]) · a[i], via 4× v_dot2_f32_bf16.
-          float partial = 0.0f;
-  #pragma unroll
-          for (int b = 0; b < 4; ++b) {
-            partial = __builtin_amdgcn_fdot2_f32_bf16(
-                *((bf16x2_t*)(&a_pack.f[b])), *((bf16x2_t*)(&q_pack.f[b])),
-                partial, /*clamp=*/false);
-          }
-
-          // block_c += y_b_f * partial + z_b_f * sum_a
-          //   y_b_f = scale, z_b_f = -(128+zero)*scale
-          // partial holds (128 + nibble) · a; subtracting (128+zero)·sum_a
-          // and scaling yields scale · (nibble - zero) · a as required.
-          block_c[0][col] =
-              __fmaf_rn(y_b_f[col], partial,
-                        __fmaf_rn(z_b_f[col], sum_a, block_c[0][col]));
-        }
-      } else {
-        // bf16 M_COUNT > 1 path with v_dot2_f32_bf16. Same opacity trick as
-        // the M=1 branch: activations + magic-value weights stored in
-        // fp32-aliased unions, dot via __builtin_amdgcn_fdot2_f32_bf16 with
-        // pointer-cast to bf16x2_t. sum_a[m] computed via second v_dot2
-        // with BF16_ONES; bias correction (y_b_f * partial + z_b_f * sum_a)
-        // applied after the dot. Magic values built once per col and reused
-        // across all M rows — amortizes dequant cost across M_COUNT.
-        typedef short __attribute__((ext_vector_type(2))) bf16x2_t;
-        constexpr uint32_t BF16_MAGIC = 0x43004300u;  // bf162(128, 128)
-        constexpr uint32_t BF16_ONES = 0x3F803F80u;   // bf162(1.0, 1.0)
-        union pack4 {
-          float f[4];
-          uint32_t u[4];
-        };
-
-        uint32_t w[4];
-        __builtin_memcpy(w, &b_w[j], sizeof(int4));
-
-        // Load M_COUNT × 8 bf16 activations as 4 uint32s each into pack4
-        // unions. Stored as uint32 to keep IR-level types opaque (defeats
-        // InstCombine fold). At M_COUNT=8 this is 32 fp32 VGPRs — within RDNA3
-        // budget.
-        pack4 a_pack[M_COUNT];
-  #pragma unroll
-        for (int m = 0; m < M_COUNT; ++m) {
-          const uint32_t* a_words =
-              reinterpret_cast<const uint32_t*>(&block_a[m][a_off]);
-          a_pack[m].u[0] = a_words[0];
-          a_pack[m].u[1] = a_words[1];
-          a_pack[m].u[2] = a_words[2];
-          a_pack[m].u[3] = a_words[3];
-        }
-
-        // sum_a[m] = Σ a[m][i] via 4× v_dot2 with bf162(1,1) — no fp32 widen.
-        float sum_a[M_COUNT];
-  #pragma unroll
-        for (int m = 0; m < M_COUNT; ++m) {
-          float s = 0.0f;
-  #pragma unroll
-          for (int b = 0; b < 4; ++b) {
-            s = __builtin_amdgcn_fdot2_f32_bf16(*((bf16x2_t*)(&a_pack[m].f[b])),
-                                                *((const bf16x2_t*)&BF16_ONES),
-                                                s, /*clamp=*/false);
-          }
-          sum_a[m] = s;
-        }
-
-        // Per col: build magic-value pack, dot against all M activations.
-        // unroll 1 keeps q_pack live one col at a time (8 fp32 VGPRs recycled)
-        // — same register-pressure trick as the previous fp32 path.
-  #pragma unroll 1
-        for (int col = 0; col < 4; ++col) {
-          pack4 q_pack;
-          const uint32_t qa = w[col];
-          q_pack.u[0] = ((qa >> 0) & 0x000F000Fu) | BF16_MAGIC;
-          q_pack.u[1] = ((qa >> 4) & 0x000F000Fu) | BF16_MAGIC;
-          q_pack.u[2] = ((qa >> 8) & 0x000F000Fu) | BF16_MAGIC;
-          q_pack.u[3] = ((qa >> 12) & 0x000F000Fu) | BF16_MAGIC;
-
-  #pragma unroll
-          for (int m = 0; m < M_COUNT; ++m) {
-            float partial = 0.0f;
-  #pragma unroll
-            for (int b = 0; b < 4; ++b) {
-              partial = __builtin_amdgcn_fdot2_f32_bf16(
-                  *((bf16x2_t*)(&a_pack[m].f[b])), *((bf16x2_t*)(&q_pack.f[b])),
-                  partial, /*clamp=*/false);
-            }
-            // block_c += y_b_f * partial + z_b_f * sum_a (same correction as
-            // M=1)
-            block_c[m][col] =
-                __fmaf_rn(y_b_f[col], partial,
-                          __fmaf_rn(z_b_f[col], sum_a[m], block_c[m][col]));
-          }
-        }
-      }
-    }
-    k += 32;  // 4 weight words * 8 nibbles = 32 K elements
-  }
-
-  // Pack the 4 FP32 partial sums into 2 packed pairs and atomically add all
-  // four lanes in a single 64-bit CAS write directly to the T-typed output
-  // (caller pre-zeros it). On gfx11 the packed atomic is a CAS-loop, but with
-  // a single b64 op we halve the atomic instruction count vs two b32 CAS
-  // calls, AND save the FP32 buffer + memset + cast pass entirely.
-  #pragma unroll
-  for (int m = 0; m < M_COUNT; ++m) {
-    if (offset_m + m >= size_m) continue;  // skip padding rows past size_m
-    T* out = c + (offset_m + m) * size_n + n;
-    if constexpr (std::is_same<T, half>::value) {
-      half2 r01 = __halves2half2(__float2half_rn(block_c[m][0]),
-                                 __float2half_rn(block_c[m][1]));
-      half2 r23 = __halves2half2(__float2half_rn(block_c[m][2]),
-                                 __float2half_rn(block_c[m][3]));
-      atomic_add_pk4_f16(out, r01, r23);
-    } else {
-      bf162_t r01;
-      r01.x = __float2bfloat16(block_c[m][0]);
-      r01.y = __float2bfloat16(block_c[m][1]);
-      bf162_t r23;
-      r23.x = __float2bfloat16(block_c[m][2]);
-      r23.y = __float2bfloat16(block_c[m][3]);
-      atomic_add_pk4_bf16(out, r01, r23);
-    }
-  }
-}
-
-#else  // non-RDNA3 device pass: empty __global__ for symbol parity.
-
-template <typename T, int M_COUNT>
-__global__ void gemm_q4_kernel_rdna3(const T*, const uint32_t*, const uint32_t*,
-                                     const T*, T*, const int, const int,
-                                     const int, const int, const int,
-                                     const int*) {}
-
-#endif  // __HIP__RDNA3__ || !__HIP_DEVICE_COMPILE__
-
-// ---------------------------------------------------------------------------
-// Launcher.
-// ---------------------------------------------------------------------------
-
-template <typename T, int M_COUNT>
-void launch_gemm_q4_for_mcount(const T* a, const uint32_t* b_q_weight,
-                               const uint32_t* b_qzeros, const T* b_scales,
-                               const int* b_q_perm, T* c, int size_m,
-                               int size_n, int size_k, int groups,
-                               int zero_offset, cudaStream_t stream) {
-  dim3 block(THREADS_X);
-  dim3 grid((size_n + BLOCK_KN_SIZE * 4 - 1) / (BLOCK_KN_SIZE * 4),
-            (size_m + M_COUNT - 1) / M_COUNT,
-            (size_k + BLOCK_KN_SIZE - 1) / BLOCK_KN_SIZE);
-
-  gemm_q4_kernel_rdna3<T, M_COUNT><<<grid, block, 0, stream>>>(
-      a, b_q_weight, b_qzeros, b_scales, c, size_m, size_n, size_k, groups,
-      zero_offset, b_q_perm);
-}
-
-// Dispatch to the largest M_COUNT template that doesn't waste more than
-// half a tile. Caps at 8: above that, the WMMA-prefill kernel (M >= 16) is
-// the right tool, not bigger M_COUNT in the scalar dot-product path.
-//
-// Tile-waste table:
-//   M=1   -> M_COUNT=1   (no waste)
-//   M=2,3 -> M_COUNT=2   (M=3 wastes 1/2 of last tile)
-//   M=4-7 -> M_COUNT=4   (worst case M=5: wastes 3/4 of last tile)
-//   M=8-15-> M_COUNT=8   (worst case M=9: wastes 7/8 of last tile)
-// "Wasted" rows are zero-padded in LDS and skip the atomic write, so they
-// only burn instructions on the last block, never affect correctness.
-template <typename T>
-void launch_gemm_q4(const T* a, const uint32_t* b_q_weight,
-                    const uint32_t* b_qzeros, const T* b_scales,
-                    const int* b_q_perm, T* c, int size_m, int size_n,
-                    int size_k, int groups, bool use_v2_format,
-                    cudaStream_t stream) {
-  const int zero_offset = use_v2_format ? 0 : 1;
-
-  if (size_m == 1) {
-    launch_gemm_q4_for_mcount<T, 1>(a, b_q_weight, b_qzeros, b_scales, b_q_perm,
-                                    c, size_m, size_n, size_k, groups,
-                                    zero_offset, stream);
-  } else if (size_m <= 3) {
-    launch_gemm_q4_for_mcount<T, 2>(a, b_q_weight, b_qzeros, b_scales, b_q_perm,
-                                    c, size_m, size_n, size_k, groups,
-                                    zero_offset, stream);
-  } else if (size_m <= 7) {
-    launch_gemm_q4_for_mcount<T, 4>(a, b_q_weight, b_qzeros, b_scales, b_q_perm,
-                                    c, size_m, size_n, size_k, groups,
-                                    zero_offset, stream);
-  } else {
-    // M_COUNT=8 covers M up to 15 here; M >= 16 should ideally take the
-    // WMMA path, but if it falls through we still produce correct output —
-    // just leaving 3-5× of throughput on the table for prefill workloads.
-    launch_gemm_q4_for_mcount<T, 8>(a, b_q_weight, b_qzeros, b_scales, b_q_perm,
-                                    c, size_m, size_n, size_k, groups,
-                                    zero_offset, stream);
-  }
-}
-
-}  // namespace gptq_rdna3
-}  // namespace vllm
-
-// ---------------------------------------------------------------------------
-// Public entry point.
-// ---------------------------------------------------------------------------
-//
-// Inputs:
-//   a         [M, K]            half or bfloat16
-//   b_q_weight[K/8, N]          uint32 (already shuffled via gptq_shuffle)
-//   b_qzeros  [groups, N/8]     uint32 (packed 4-bit zeros)
-//   b_scales  [groups, N]       half or bfloat16
-//   b_g_idx   [K] or empty      int32 (act-order permutation; empty=identity)
-//   use_v2_format                bool   (true = GPTQv2, no +1 zero offset)
-//
-// Output:
-//   c         [M, N]            same dtype as a
-
-torch::Tensor gptq_gemm_rdna3_wmma(torch::Tensor a, torch::Tensor b_q_weight,
-                                   torch::Tensor b_qzeros,
-                                   torch::Tensor b_scales,
-                                   torch::Tensor b_g_idx, bool use_v2_format);
-
-torch::Tensor gptq_gemm_rdna3(torch::Tensor a, torch::Tensor b_q_weight,
-                              torch::Tensor b_qzeros, torch::Tensor b_scales,
-                              torch::Tensor b_g_idx, bool use_v2_format) {
-  if (a.dim() == 2 && b_q_weight.dim() == 2 && a.size(1) % 16 == 0 &&
-      b_q_weight.size(1) % 16 == 0 &&
-      ((a.scalar_type() == torch::kBFloat16 && a.size(0) >= 16) ||
-       (a.scalar_type() == torch::kHalf && a.size(0) >= 64))) {
-    return gptq_gemm_rdna3_wmma(a, b_q_weight, b_qzeros, b_scales, b_g_idx,
-                                use_v2_format);
-  }
-
-  TORCH_CHECK(a.is_cuda(), "a must be a CUDA/HIP tensor");
-  TORCH_CHECK(b_q_weight.is_cuda(), "b_q_weight must be a CUDA/HIP tensor");
-  TORCH_CHECK(b_qzeros.is_cuda(), "b_qzeros must be a CUDA/HIP tensor");
-  TORCH_CHECK(b_scales.is_cuda(), "b_scales must be a CUDA/HIP tensor");
-  TORCH_CHECK(a.dim() == 2, "a must be 2D [M, K]");
-  TORCH_CHECK(b_q_weight.dim() == 2, "b_q_weight must be 2D [K/8, N]");
-  TORCH_CHECK(
-      a.scalar_type() == torch::kHalf || a.scalar_type() == torch::kBFloat16,
-      "a must be half or bfloat16");
-  TORCH_CHECK(a.scalar_type() == b_scales.scalar_type(),
-              "b_scales dtype must match a");
-
-  const at::cuda::OptionalCUDAGuard device_guard(device_of(a));
-  auto stream = at::cuda::getCurrentCUDAStream();
-
-  int size_m = (int)a.size(0);
-  int size_k = (int)a.size(1);
-  int size_n = (int)b_q_weight.size(1);
-  int groups = (int)b_qzeros.size(0);
-
-  TORCH_CHECK(b_q_weight.size(0) * 8 == size_k,
-              "b_q_weight first dim must be K/8");
-  TORCH_CHECK(b_scales.size(0) == groups,
-              "b_scales must have same group count as qzeros");
-  TORCH_CHECK(b_scales.size(1) == size_n, "b_scales last dim must be N");
-  TORCH_CHECK(size_n % 8 == 0, "N must be a multiple of 8 (64-bit atomic CAS)");
-
-  auto opts = torch::TensorOptions().dtype(a.dtype()).device(a.device());
-  at::Tensor c = torch::zeros({size_m, size_n}, opts);
-
-  const int* g_idx_ptr = nullptr;
-  if (!b_g_idx.device().is_meta() && b_g_idx.numel() > 0) {
-    TORCH_CHECK(b_g_idx.scalar_type() == torch::kInt32,
-                "b_g_idx must be int32");
-    g_idx_ptr = (const int*)b_g_idx.data_ptr();
-  }
-
-  if (a.scalar_type() == torch::kHalf) {
-    vllm::gptq_rdna3::launch_gemm_q4<half>(
-        (const half*)a.data_ptr(), (const uint32_t*)b_q_weight.data_ptr(),
-        (const uint32_t*)b_qzeros.data_ptr(), (const half*)b_scales.data_ptr(),
-        g_idx_ptr, (half*)c.data_ptr(), size_m, size_n, size_k, groups,
-        use_v2_format, stream);
-  } else {
-    vllm::gptq_rdna3::launch_gemm_q4<vllm::gptq_rdna3::bf16_t>(
-        (const vllm::gptq_rdna3::bf16_t*)a.data_ptr(),
-        (const uint32_t*)b_q_weight.data_ptr(),
-        (const uint32_t*)b_qzeros.data_ptr(),
-        (const vllm::gptq_rdna3::bf16_t*)b_scales.data_ptr(), g_idx_ptr,
-        (vllm::gptq_rdna3::bf16_t*)c.data_ptr(), size_m, size_n, size_k, groups,
-        use_v2_format, stream);
-  }
-
-  return c;
-}
@@ -1,239 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-//
-// W4A16 dequant primitives for RDNA3 (gfx1100/gfx1101/gfx1102), templated on
-// the activation/scale dtype (half or __hip_bfloat16). The fp16 path reuses
-// the classic exllamav2 bit-trick:
-//
-//   (qa & 0x000F000F) | 0x64006400  ->  half2(1024+q_lo, 1024+q_hi)
-//   (qa & 0x00F000F0) | 0x64006400  ->  half2(1024+q_lo*16, 1024+q_hi*16)
-//
-// The "*16 then divide by 16 in the FMA" trick for the upper-nibble pairs
-// works in fp16 because the mantissa (10 bits) is wide enough to hold a value
-// shifted by 4 bits. In bf16 the mantissa is only 7 bits, so shifting an upper
-// nibble into bits [7:4] would spill into the exponent. To avoid that, the
-// bf16 path shifts each pair of nibbles down to bits [3:0]/[19:16] with a
-// single right-shift before the OR with 0x43004300 (= bf162(128, 128)).
-
-#ifndef _qdq_4_rdna3_cuh
-#define _qdq_4_rdna3_cuh
-
-#include <cstdint>
-
-#include <hip/hip_bf16.h>
-#include <hip/hip_fp16.h>
-
-namespace vllm {
-namespace gptq_rdna3 {
-
-using bf16_t = __hip_bfloat16;
-using bf162_t = __hip_bfloat162;
-
-// Bit-shuffle for an int32 holding 8 sequential 4-bit weights q[0..7]:
-//   in:  q[7] q[6] q[5] q[4] q[3] q[2] q[1] q[0]   (LSB first)
-//   out: q[7] q[5] q[3] q[1] q[6] q[4] q[2] q[0]   (even/odd interleaved)
-//
-// After shuffle, q[2k]   sits at bits [4k   : 4k+3]   (lower 16)
-//                q[2k+1] sits at bits [16+4k: 16+4k+3] (upper 16)
-// so a single mask 0x000F000F selects the matching even/odd pair, ready to
-// bitcast to half2 / bfloat162 after OR-ing with the magic constant.
-__forceinline__ __device__ void shuffle_4bit_8(uint32_t* q) {
-  uint32_t qa = q[0];
-  uint32_t qb = 0;
-#pragma unroll
-  for (int i = 0; i < 4; i++) {
-    uint32_t qa0 = qa & 0x0F;
-    uint32_t qa1 = (qa & 0xF0) >> 4;
-    qa >>= 8;
-    qb |= (qa1 << (i * 4 + 16));
-    qb |= (qa0 << (i * 4));
-  }
-  q[0] = qb;
-}
-
-// ---------------------------------------------------------------------------
-// fp16 path
-// ---------------------------------------------------------------------------
-
-// Precompute scale-baked constants for a single zero/scale pair.
-//   z1z16[0] = scale * (-1024 - zero)            (used for "low" pairs)
-//   z1z16[1] = scale * (-64   - zero)            (used for "high" pairs)
-//   y1y16[0] = scale * 1                          (low pairs are q + 1024)
-//   y1y16[1] = scale * (1/16)                     (high pairs are q*16 + 1024)
-__forceinline__ __device__ void prep_zero_scale_fp16(uint32_t zero, half scale,
-                                                     half2 (&z1z16)[2],
-                                                     half2 (&y1y16)[2]) {
-  // half(-1024 - zero) via the exllamav2 bit-trick:
-  //   half bits 0xE400 == -1024.0 ; ORing the zero into mantissa subtracts it.
-  union {
-    uint16_t u;
-    half h;
-  } z1u;
-  z1u.u = (uint16_t)(0xE400 | zero);
-  half z1 = z1u.h;
-  half z16 = __hsub(__int2half_rn(-64), __int2half_rn((int)zero));
-
-  half2 scale2 = __half2half2(scale);
-  z1z16[0] = __hmul2(scale2, __half2half2(z1));
-  z1z16[1] = __hmul2(scale2, __half2half2(z16));
-
-  half y1 = __float2half_rn(1.0f);
-  half y16 = __float2half_rn(1.0f / 16.0f);
-  y1y16[0] = __hmul2(scale2, __half2half2(y1));
-  y1y16[1] = __hmul2(scale2, __half2half2(y16));
-}
-
-// Dequantize one int32 (8 shuffled 4-bit weights) into 4 half2 pairs:
-//   dq[0] = (q[0], q[1]) * scale - zero*scale
-//   dq[1] = (q[2], q[3]) * scale - zero*scale
-//   dq[2] = (q[4], q[5]) * scale - zero*scale
-//   dq[3] = (q[6], q[7]) * scale - zero*scale
-__forceinline__ __device__ void dequant_4bit_8_fp16(uint32_t qa, half2 (&dq)[4],
-                                                    half2 (&z1z16)[2],
-                                                    half2 (&y1y16)[2]) {
-  const uint32_t c0 = 0x64006400;
-
-  union {
-    uint32_t u;
-    half2 h2;
-  } q0, q1, q2, q3;
-  q0.u = (qa & 0x000F000F) | c0;  // half2(q[0]+1024, q[1]+1024)
-  q1.u = (qa & 0x00F000F0) | c0;  // half2(q[2]*16+1024, q[3]*16+1024)
-  uint32_t qa_hi = qa >> 8;
-  q2.u = (qa_hi & 0x000F000F) | c0;  // half2(q[4]+1024, q[5]+1024)
-  q3.u = (qa_hi & 0x00F000F0) | c0;  // half2(q[6]*16+1024, q[7]*16+1024)
-
-  dq[0] = __hfma2(q0.h2, y1y16[0], z1z16[0]);
-  dq[1] = __hfma2(q1.h2, y1y16[1], z1z16[1]);
-  dq[2] = __hfma2(q2.h2, y1y16[0], z1z16[0]);
-  dq[3] = __hfma2(q3.h2, y1y16[1], z1z16[1]);
-}
-
-// ---------------------------------------------------------------------------
-// bf16 path
-// ---------------------------------------------------------------------------
-
-// Bit-trick magic for bf16:
-//   bf16(128) == 0x4300 (sign 0, exp 134, mantissa 0).
-//   For nibble n in [0..15], bits [3:0] of mantissa hold n exactly because
-//   bf16's ULP at 128 is 1 (mantissa step = 2^(7-7) = 1). So
-//   ((qa & 0x000F000F) | 0x43004300) bitcasts to bfloat162(128+n_lo, 128+n_hi).
-//
-// Because bf16's mantissa is only 7 bits, we cannot use the fp16 "upper nibble
-// * 16" trick. Instead each pair of nibbles is shifted down to [3:0]/[19:16]
-// via a single 4/8/12-bit right-shift before the OR. That costs one extra
-// shift per pair vs fp16, but keeps the FMA structure identical.
-__forceinline__ __device__ void prep_zero_scale_bf16(uint32_t zero,
-                                                     bf16_t scale,
-                                                     bf162_t& z_prep,
-                                                     bf162_t& y_prep) {
-  // z = scale * -(128 + zero); y = scale.
-  float scale_f = __bfloat162float(scale);
-  float zf = -(128.0f + (float)zero) * scale_f;
-  bf16_t zb = __float2bfloat16(zf);
-  z_prep = __bfloat162bfloat162(zb);
-  y_prep = __bfloat162bfloat162(scale);
-}
-
-__forceinline__ __device__ void dequant_4bit_8_bf16(uint32_t qa,
-                                                    bf162_t (&dq)[4],
-                                                    bf162_t z_prep,
-                                                    bf162_t y_prep) {
-  const uint32_t c0 = 0x43004300;
-
-  union {
-    uint32_t u;
-    bf162_t b2;
-  } q0, q1, q2, q3;
-  q0.u = ((qa >> 0) & 0x000F000F) | c0;   // bf162(128+q[0], 128+q[1])
-  q1.u = ((qa >> 4) & 0x000F000F) | c0;   // bf162(128+q[2], 128+q[3])
-  q2.u = ((qa >> 8) & 0x000F000F) | c0;   // bf162(128+q[4], 128+q[5])
-  q3.u = ((qa >> 12) & 0x000F000F) | c0;  // bf162(128+q[6], 128+q[7])
-
-  // dq = q_b * scale + (-(128+zero)*scale) = (q - zero) * scale
-  dq[0] = __hfma2(q0.b2, y_prep, z_prep);
-  dq[1] = __hfma2(q1.b2, y_prep, z_prep);
-  dq[2] = __hfma2(q2.b2, y_prep, z_prep);
-  dq[3] = __hfma2(q3.b2, y_prep, z_prep);
-}
-
-// ---------------------------------------------------------------------------
-// bf16-input → fp32-output dequant (RDNA3 scalar path).
-//
-// RDNA3 (gfx1100) has no v_pk_fma_bf16; packed bf16 FMA lowers to a slow
-// fallback. Rather than computing dq in bf16 and widening at FMA time in
-// the dot product, we widen to fp32 here once (a free left-shift by 16) and
-// emit the (q - zero) * scale FMA directly in fp32. This:
-//   * Replaces 4× slow bf16 packed FMA with 8× fast fp32 FMA per int32.
-//   * Eliminates 4× bf16→fp32 widens that the dot product would do.
-//   * Keeps the dot product accumulator in fp32 without a roundtrip.
-//
-// Output: fp32 dq[8], one element per K position (consumed by the
-// fp32-overload of dot22_8_f in q_gemm_rdna3.cu).
-__forceinline__ __device__ void prep_zero_scale_bf16_f32(uint32_t zero,
-                                                         bf16_t scale,
-                                                         float& z_prep,
-                                                         float& y_prep) {
-  float scale_f = __bfloat162float(scale);
-  z_prep = -(128.0f + (float)zero) * scale_f;
-  y_prep = scale_f;
-}
-
-// Pure-q dequant for the M_COUNT=1 factored path: outputs the unscaled fp32
-// values 128+nibble, without folding scale/zero. The caller folds scale/zb
-// into the accumulator outside the inner loop using a precomputed sum_a,
-// which saves ~27% of the FMA count vs the per-col-dequant approach above
-// (only beneficial at M_COUNT=1; break-even at M_COUNT=2).
-//
-// Cost: 0 FMAs (pure bit-trick + as_float reinterprets).
-__forceinline__ __device__ void dequant_4bit_8_bf16_q_only(uint32_t qa,
-                                                           float (&q_f32)[8]) {
-  const uint32_t c0 = 0x43004300;
-  const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0;
-  const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0;
-  const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0;
-  const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0;
-  q_f32[0] = __uint_as_float((q0 & 0xFFFFu) << 16);
-  q_f32[1] = __uint_as_float(q0 & 0xFFFF0000u);
-  q_f32[2] = __uint_as_float((q1 & 0xFFFFu) << 16);
-  q_f32[3] = __uint_as_float(q1 & 0xFFFF0000u);
-  q_f32[4] = __uint_as_float((q2 & 0xFFFFu) << 16);
-  q_f32[5] = __uint_as_float(q2 & 0xFFFF0000u);
-  q_f32[6] = __uint_as_float((q3 & 0xFFFFu) << 16);
-  q_f32[7] = __uint_as_float(q3 & 0xFFFF0000u);
-}
-
-__forceinline__ __device__ void dequant_4bit_8_bf16_f32(uint32_t qa,
-                                                        float (&dq)[8],
-                                                        float z_prep,
-                                                        float y_prep) {
-  const uint32_t c0 = 0x43004300;
-  const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0;
-  const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0;
-  const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0;
-  const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0;
-  // bf16(128+nibble) bits → fp32(128+nibble) bits via left-shift by 16
-  // (just zero-extends the mantissa from 7 to 23 bits; exponent preserved).
-  const float q0x = __uint_as_float((q0 & 0xFFFFu) << 16);
-  const float q0y = __uint_as_float(q0 & 0xFFFF0000u);
-  const float q1x = __uint_as_float((q1 & 0xFFFFu) << 16);
-  const float q1y = __uint_as_float(q1 & 0xFFFF0000u);
-  const float q2x = __uint_as_float((q2 & 0xFFFFu) << 16);
-  const float q2y = __uint_as_float(q2 & 0xFFFF0000u);
-  const float q3x = __uint_as_float((q3 & 0xFFFFu) << 16);
-  const float q3y = __uint_as_float(q3 & 0xFFFF0000u);
-  // dq[i] = q_f32 * scale + (-(128+zero)*scale) = (nibble - zero) * scale
-  dq[0] = __fmaf_rn(q0x, y_prep, z_prep);
-  dq[1] = __fmaf_rn(q0y, y_prep, z_prep);
-  dq[2] = __fmaf_rn(q1x, y_prep, z_prep);
-  dq[3] = __fmaf_rn(q1y, y_prep, z_prep);
-  dq[4] = __fmaf_rn(q2x, y_prep, z_prep);
-  dq[5] = __fmaf_rn(q2y, y_prep, z_prep);
-  dq[6] = __fmaf_rn(q3x, y_prep, z_prep);
-  dq[7] = __fmaf_rn(q3y, y_prep, z_prep);
-}
-
-}  // namespace gptq_rdna3
-}  // namespace vllm
-
-#endif  // _qdq_4_rdna3_cuh
@@ -39,19 +39,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, rocm_ops) {
      "          Tensor scale_b, int CuCount) -> ()");
  rocm_ops.impl("wvSplitKQ", torch::kCUDA, &wvSplitKQ);

-#ifdef VLLM_ROCM_GFX1100
-  // W4A16 GPTQ kernels for AMD RDNA3 (gfx1100).
-  rocm_ops.def(
-      "gptq_gemm_rdna3(Tensor a, Tensor b_q_weight, Tensor b_qzeros, "
-      "Tensor b_scales, Tensor b_g_idx, bool use_v2_format) -> Tensor");
-  rocm_ops.impl("gptq_gemm_rdna3", torch::kCUDA, &gptq_gemm_rdna3);
-
-  rocm_ops.def(
-      "gptq_gemm_rdna3_wmma(Tensor a, Tensor b_q_weight, Tensor b_qzeros, "
-      "Tensor b_scales, Tensor b_g_idx, bool use_v2_format) -> Tensor");
-  rocm_ops.impl("gptq_gemm_rdna3_wmma", torch::kCUDA, &gptq_gemm_rdna3_wmma);
-#endif
-
  // Custom attention op
  // Compute the attention between an input query and the cached
  // keys/values using PagedAttention.
@@ -1,7 +1,4 @@
-// Provides torch::Tensor for ops.h (previously included transitively via
-// cache.h, which is no longer included here after cache ops moved to
-// _C_stable_libtorch).
-#include <torch/all.h>
+#include "cache.h"
 #include "cuda_utils.h"
 #include "ops.h"
 #include "core/registration.h"
@@ -36,6 +33,35 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
  ops.impl("get_cuda_view_from_cpu_tensor", torch::kCPU,
           &get_cuda_view_from_cpu_tensor);

+  // Attention ops
+  // Compute the attention between an input query and the cached
+  // keys/values using PagedAttention.
+  ops.def(
+      "paged_attention_v1("
+      "    Tensor! out, Tensor query, Tensor key_cache,"
+      "    Tensor value_cache, int num_kv_heads, float scale,"
+      "    Tensor block_tables, Tensor seq_lens, int block_size,"
+      "    int max_seq_len, Tensor? alibi_slopes,"
+      "    str kv_cache_dtype, Tensor k_scale, Tensor v_scale,"
+      "    int tp_rank, int blocksparse_local_blocks,"
+      "    int blocksparse_vert_stride, int blocksparse_block_size,"
+      "    int blocksparse_head_sliding_step) -> ()");
+  ops.impl("paged_attention_v1", torch::kCUDA, &paged_attention_v1);
+
+  // PagedAttention V2.
+  ops.def(
+      "paged_attention_v2("
+      "    Tensor! out, Tensor! exp_sums, Tensor! max_logits,"
+      "    Tensor! tmp_out, Tensor query, Tensor key_cache,"
+      "    Tensor value_cache, int num_kv_heads, float scale,"
+      "    Tensor block_tables, Tensor seq_lens, int block_size,"
+      "    int max_seq_len, Tensor? alibi_slopes,"
+      "    str kv_cache_dtype, Tensor k_scale, Tensor v_scale,"
+      "    int tp_rank, int blocksparse_local_blocks,"
+      "    int blocksparse_vert_stride, int blocksparse_block_size,"
+      "    int blocksparse_head_sliding_step) -> ()");
+  ops.impl("paged_attention_v2", torch::kCUDA, &paged_attention_v2);
+
  // Activation ops (quantized only — basic ops moved to _C_stable_libtorch)
  ops.def(
      "silu_and_mul_quant(Tensor! result, Tensor input, Tensor scale) -> ()");
@@ -191,6 +217,114 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
 #endif
 }

+TORCH_LIBRARY_EXPAND(CONCAT(TORCH_EXTENSION_NAME, _cache_ops), cache_ops) {
+  // Cache ops
+  // Swap in (out) the cache blocks from src to dst.
+  cache_ops.def(
+      "swap_blocks(Tensor src, Tensor! dst,"
+      "            int block_size_in_bytes, Tensor block_mapping) -> ()");
+  cache_ops.impl("swap_blocks", torch::kCUDA, &swap_blocks);
+
+  // Batch swap: submit all block copies in a single driver call.
+  cache_ops.def(
+      "swap_blocks_batch(Tensor src_ptrs, Tensor dst_ptrs,"
+      "                  Tensor sizes,"
+      "                  bool is_src_access_order_any=False) -> ()");
+  cache_ops.impl("swap_blocks_batch", torch::kCPU, &swap_blocks_batch);
+
+  // Reshape the key and value tensors and cache them.
+  cache_ops.def(
+      "reshape_and_cache(Tensor key, Tensor value,"
+      "                  Tensor! key_cache, Tensor! value_cache,"
+      "                  Tensor slot_mapping,"
+      "                  str kv_cache_dtype,"
+      "                  Tensor k_scale, Tensor v_scale) -> ()");
+  cache_ops.impl("reshape_and_cache", torch::kCUDA, &reshape_and_cache);
+
+  // Reshape the key and value tensors and cache them.
+  cache_ops.def(
+      "reshape_and_cache_flash(Tensor key, Tensor value,"
+      "                        Tensor! key_cache,"
+      "                        Tensor! value_cache,"
+      "                        Tensor slot_mapping,"
+      "                        str kv_cache_dtype,"
+      "                        Tensor k_scale, Tensor v_scale) -> ()");
+  cache_ops.impl("reshape_and_cache_flash", torch::kCUDA,
+                 &reshape_and_cache_flash);
+
+  // Concat kv_c and k_pe and cache them.
+  cache_ops.def(
+      "concat_and_cache_mla(Tensor kv_c, Tensor k_pe,"
+      "                     Tensor! kv_cache,"
+      "                     Tensor slot_mapping,"
+      "                     str kv_cache_dtype,"
+      "                     Tensor scale) -> ()");
+  cache_ops.impl("concat_and_cache_mla", torch::kCUDA, &concat_and_cache_mla);
+
+  // Rotate Q and K, then write to kv cache for MLA
+  cache_ops.def(
+      "concat_and_cache_mla_rope_fused("
+      "                     Tensor positions,"
+      "                     Tensor! q_pe,"
+      "                     Tensor! k_pe,"
+      "                     Tensor kv_c,"
+      "                     Tensor cos_sin_cache,"
+      "                     bool is_neox,"
+      "                     Tensor slot_mapping,"
+      "                     Tensor! kv_cache,"
+      "                     str kv_cache_dtype,"
+      "                     Tensor kv_cache_scale) -> ()");
+  cache_ops.impl("concat_and_cache_mla_rope_fused", torch::kCUDA,
+                 &concat_and_cache_mla_rope_fused);
+
+  // Convert the key and value cache to fp8 data type.
+  cache_ops.def(
+      "convert_fp8(Tensor! dst_cache, Tensor src_cache, float scale, "
+      "str kv_cache_dtype) -> ()");
+  cache_ops.impl("convert_fp8", torch::kCUDA, &convert_fp8);
+
+  // Gather cache blocks from src_cache to dst, dequantizing from
+  // src_cache's dtype to dst's dtype if necessary.
+  cache_ops.def(
+      "gather_and_maybe_dequant_cache(Tensor src_cache, Tensor! dst, "
+      "                               Tensor block_table, Tensor cu_seq_lens, "
+      "                               Tensor token_to_seq, "
+      "                               int num_tokens, "
+      "                               str kv_cache_dtype, "
+      "                               Tensor scale, Tensor? seq_starts) -> ()");
+  cache_ops.impl("gather_and_maybe_dequant_cache", torch::kCUDA,
+                 &gather_and_maybe_dequant_cache);
+
+  cache_ops.def(
+      "cp_gather_cache(Tensor src_cache, Tensor! dst, Tensor block_table, "
+      "Tensor cu_seq_lens, int batch_size, Tensor? seq_starts) -> ()");
+  cache_ops.impl("cp_gather_cache", torch::kCUDA, &cp_gather_cache);
+
+  cache_ops.def(
+      "cp_gather_and_upconvert_fp8_kv_cache(Tensor src_cache, Tensor! dst, "
+      "Tensor block_table, Tensor seq_lens, Tensor workspace_starts, int "
+      "batch_size) -> ()");
+  cache_ops.impl("cp_gather_and_upconvert_fp8_kv_cache", torch::kCUDA,
+                 &cp_gather_and_upconvert_fp8_kv_cache);
+
+  cache_ops.def(
+      "indexer_k_quant_and_cache(Tensor k, Tensor! kv_cache, Tensor "
+      "slot_mapping, "
+      "int quant_block_size, str kv_cache_dtype) -> ()");
+  cache_ops.impl("indexer_k_quant_and_cache", torch::kCUDA,
+                 &indexer_k_quant_and_cache);
+
+  cache_ops.def(
+      "concat_mla_q(Tensor ql_nope, Tensor q_pe, Tensor! q_out) -> ()");
+  cache_ops.impl("concat_mla_q", torch::kCUDA, &concat_mla_q);
+
+  cache_ops.def(
+      "cp_gather_indexer_k_quant_cache(Tensor kv_cache, Tensor! dst_k, Tensor! "
+      "dst_scale, Tensor block_table, Tensor cu_seq_lens) -> ()");
+  cache_ops.impl("cp_gather_indexer_k_quant_cache", torch::kCUDA,
+                 &cp_gather_indexer_k_quant_cache);
+}
+
 TORCH_LIBRARY_EXPAND(CONCAT(TORCH_EXTENSION_NAME, _cuda_utils), cuda_utils) {
  // Cuda utils

@@ -165,23 +165,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
    --mount=type=cache,target=/vllm-workspace/.deps,sharing=locked \
    VLLM_TARGET_DEVICE=cpu python3 setup.py bdist_wheel --dist-dir=dist --py-limited-api=cp38

-######################### TRITON-CPU BUILD IMAGE #########################
-FROM base AS vllm-triton-cpu-build
-
-WORKDIR /vllm-workspace
-
-RUN mkdir dist
-
-RUN --mount=type=cache,target=/root/.cache/uv \
-    --mount=type=cache,target=/root/.cache/ccache \
-    --mount=type=cache,target=/vllm-workspace/.deps,sharing=locked \
-    if [ "$TARGETARCH" = "amd64" ] || [ "$VLLM_CPU_X86" != "0" ]; then \
-        git clone --recurse-submodules "https://github.com/triton-lang/triton-cpu.git"; \
-        cd triton-cpu; \
-        git checkout "270e696d"; \
-        uv build --wheel --out-dir=../dist; \
-    fi
-
 ######################### TEST DEPS #########################
 FROM base AS vllm-test-deps

@@ -274,14 +257,7 @@ WORKDIR /vllm-workspace
 RUN --mount=type=cache,target=/root/.cache/uv \
    --mount=type=cache,target=/root/.cache/ccache \
    --mount=type=bind,from=vllm-build,src=/vllm-workspace/dist,target=dist \
-    uv pip install "$(realpath dist/*.whl)[audio]"
-
-RUN --mount=type=cache,target=/root/.cache/uv \
-    --mount=type=cache,target=/root/.cache/ccache \
-    --mount=type=bind,from=vllm-triton-cpu-build,src=/vllm-workspace/dist,target=dist \
-    if [ "$TARGETARCH" = "amd64" ] || [ "$VLLM_CPU_X86" != "0" ]; then \
-        uv pip install "$(realpath dist/*.whl)"; \
-    fi
+    uv pip install "$(realpath dist/*.whl)[audio,triton-cpu]"

 # Add labels to document build configuration
 LABEL org.opencontainers.image.title="vLLM CPU"
@@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.9.0"
 ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git"
 ARG FA_BRANCH="0e60e394"
 ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git"
-ARG AITER_BRANCH="v0.1.13.post1"
+ARG AITER_BRANCH="v0.1.13"
 ARG AITER_REPO="https://github.com/ROCm/aiter.git"
 ARG MORI_BRANCH="v1.1.0"
 ARG MORI_REPO="https://github.com/ROCm/mori.git"
@@ -246,12 +246,6 @@ Every image listed in "image_files" is added to the request in the listed order

 The "image" shorthand accepts the same values as "image_files". The "image_url" field accepts either an OpenAI-style object with a "url" field or a URL string.

-By default, image references are sent to the serving endpoint as provided, with local image paths converted to `file://` URLs.
-
-If the benchmark client should load local and HTTP(S) images before sending requests, pass `--custom-ensure-client-side-data` to encode them as base64 data URLs on the client side.
-
-Existing `data:image/...` URLs are already self-contained and are kept unchanged.
-
 ```bash
 # need a model with vision capability here
 vllm serve Qwen/Qwen2-VL-7B-Instruct
@@ -259,13 +253,13 @@ vllm serve Qwen/Qwen2-VL-7B-Instruct

 ```bash
 # run benchmarking script
-vllm bench serve --save-result --save-detailed \
+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> \
-  --custom-ensure-client-side-data
+  --allowed-local-media-path /path/to/image/folder
 ```

 Note that we need to use the `openai-chat` backend and `/v1/chat/completions` endpoint for multimodal inputs.
@@ -46,14 +46,14 @@ In V1, **chunked prefill is enabled by default whenever possible**. With chunked

 This policy has two benefits:

- It improves inter-token latency (ITL) and generation decode because decode requests are prioritized.
+- It improves ITL and generation decode because decode requests are prioritized.
 - It helps achieve better GPU utilization by locating compute-bound (prefill) and memory-bound (decode) requests to the same batch.

 ### Performance Tuning with Chunked Prefill

 You can tune the performance by adjusting `max_num_batched_tokens`:

- Smaller values (e.g., 2048) achieve better ITL because there are fewer prefills slowing down decodes.
+- Smaller values (e.g., 2048) achieve better inter-token latency (ITL) because there are fewer prefills slowing down decodes.
 - Higher values achieve better time to first token (TTFT) as you can process more prefill tokens in a batch.
 - For optimal throughput, we recommend setting `max_num_batched_tokens > 8192` especially for smaller models on large GPUs.
 - If `max_num_batched_tokens` is the same as `max_model_len`, that's almost the equivalent to the V0 default scheduling policy (except that it still prioritizes decodes).
@@ -177,7 +177,7 @@ Priority is **1 = highest** (tried first).
 | `FLASH_ATTN` | FA4* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ✅ | ✅ | ❌ | ✅ | All | ≥10.0 |
 | `FLASH_ATTN_DIFFKV` | | fp16, bf16 | `auto` | Any | Any | ❌ | ❌ | ❌ | ✅ | Decoder | Any |
 | `FLEX_ATTENTION` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ✅ | ✅ | ❌ | Decoder, Encoder Only | Any |
-| `ROCM_AITER_FA` | | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32 | 64, 128, 256 | ✅ | ✅ | ❌ | ❌ | Decoder | N/A |
+| `ROCM_AITER_FA` | | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32 | 64, 128, 256 | ❌ | ✅ | ❌ | ❌ | Decoder | N/A |
 | `ROCM_AITER_UNIFIED_ATTN` | | fp16, bf16 | `auto` | %16 | Any | ✅ | ❌ | ✅ | ❌ | All | N/A |
 | `ROCM_ATTN` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | 32, 64, 80, 96, 128, 160, 192, 224, 256 | ❌ | ✅ | ✅ | ❌ | Decoder, Encoder, Encoder Only | N/A |
 | `TRITON_ATTN` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2`, `int8_per_token_head`, `fp8_per_token_head` | %16 | Any | ✅ | ❌ | ✅ | ❌ | All | Any |
@@ -778,7 +778,7 @@ Then, you can use the OpenAI client as follows:
        base_url=openai_api_base,
    )

-    video_url = "https://huggingface.co/datasets/raushan-testing-hf/videos-test/resolve/main/sample_demo_1.mp4"
+    video_url = "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/ForBiggerFun.mp4"

    ## Use video url in the payload
    chat_completion_from_url = client.chat.completions.create(
@@ -41,7 +41,7 @@ pip install -v -r requirements/xpu.txt

    ```bash
    pip uninstall -y triton triton-xpu
-    pip install triton-xpu==3.7.0 --extra-index-url https://download.pytorch.org/whl/xpu
+    pip install triton-xpu==3.6.0 --extra-index-url https://download.pytorch.org/whl/xpu
    ```

    !!! note
@@ -17,7 +17,6 @@ Sorted alphabetically by GitHub handle:
 - [@bbrowning](https://github.com/bbrowning): Tool use and reasoning parser
 - [@benchislett](https://github.com/benchislett): Engine core and spec decode
 - [@bigPYJ1151](https://github.com/bigPYJ1151): Intel CPU/XPU integration
- [@BugenZhao](https://github.com/BugenZhao): Rust frontend
 - [@chaunceyjiang](https://github.com/chaunceyjiang): Tool use and reasoning parser
 - [@DarkLight1337](https://github.com/DarkLight1337): Multimodality, API server
 - [@esmeetu](https://github.com/esmeetu): developer marketing, community
@@ -131,8 +130,6 @@ If you have PRs touching the area, please feel free to ping the area owner for r
    - @DarkLight1337
 - API Server: The OpenAI-compatible API server
    - @DarkLight1337, @njhill, @aarnphm, @simon-mo, @heheda12345 (Responses API)
- Rust Frontend: The experimental API server in Rust
-    - @BugenZhao, @njhill
 - Batch Runner: The OpenAI-compatible batch runner
    - @simon-mo

@@ -437,7 +437,6 @@ th {
 | `LongcatFlashForCausalLM` | LongCat-Flash | `meituan-longcat/LongCat-Flash-Chat`, `meituan-longcat/LongCat-Flash-Chat-FP8` | ✅︎ | ✅︎ |
 | `MambaForCausalLM` | Mamba | `state-spaces/mamba-130m-hf`, `state-spaces/mamba-790m-hf`, `state-spaces/mamba-2.8b-hf`, etc. | | ✅︎ |
 | `Mamba2ForCausalLM` | Mamba2 | `mistralai/Mamba-Codestral-7B-v0.1`, etc. | | ✅︎ |
-| `MellumForCausalLM` | Mellum 2 | `JetBrains/Mellum2-12B-A2.5B-Base`, etc. | | ✅︎ |
 | `MiMoForCausalLM` | MiMo | `XiaomiMiMo/MiMo-7B-RL`, etc. | ✅︎ | ✅︎ |
 | `MiMoV2FlashForCausalLM` | MiMoV2Flash | `XiaomiMiMo/MiMo-V2-Flash`, etc. | | ✅︎ |
 | `MiMoV2ForCausalLM` | MiMoV2Pro | `XiaomiMiMo/MiMo-V2.5-Pro`, etc. | | ✅︎ |
@@ -634,7 +633,6 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `SmolVLMForConditionalGeneration` | SmolVLM2 | T + I | `SmolVLM2-2.2B-Instruct` | ✅︎ | |
 | `Step3VLForConditionalGeneration` | Step3-VL | T + I<sup>+</sup> | `stepfun-ai/step3` | | ✅︎ |
 | `StepVLForConditionalGeneration` | Step3-VL-10B | T + I<sup>+</sup> | `stepfun-ai/Step3-VL-10B` | | ✅︎ |
-| `Step3p7ForConditionalGeneration` | Step-3.7-Flash | T + I<sup>+</sup> | `stepfun-ai/Step-3.7-Flash` | | ✅︎ |
 | `TarsierForConditionalGeneration` | Tarsier | T + I<sup>E+</sup> | `omni-search/Tarsier-7b`, `omni-search/Tarsier-34b` | | ✅︎ |
 | `Tarsier2ForConditionalGeneration`<sup>^</sup> | Tarsier2 | T + I<sup>E+</sup> + V<sup>E+</sup> | `omni-research/Tarsier2-Recap-7b`, `omni-research/Tarsier2-7b-0115` | | ✅︎ |
 | `UltravoxModel` | Ultravox | T + A<sup>E+</sup> | `fixie-ai/ultravox-v0_5-llama-3_2-1b` | ✅︎ | ✅︎ |
@@ -1,60 +1,41 @@
-if [ "$READTHEDOCS_VERSION_TYPE" != "external" ]; then
-  echo "Not a PR build (version type=$READTHEDOCS_VERSION_TYPE); skipping pre-run-check gate."
-  exit 0
-fi
-
-echo "Checking for changes to docs-affecting files vs origin/main..."
-DOCS_PATHS=(
-  docs/                       # Actual docs content
-  examples/                   # Examples are rendered in docs
-  vllm/                       # API & CLI reference
-  requirements/test/cuda.txt  # CLI reference (see docs/mkdocs/hooks/generate_argparse.py)
-  mkdocs.yaml                 # Affects build process
-  .readthedocs.yaml           # Affects build process
-  requirements/docs.txt       # Affects build process
-  requirements/docs.in        # Affects build process
-)
-if git diff --quiet origin/main -- "${DOCS_PATHS[@]}"; then
-  echo "No docs-affecting files changed vs origin/main; cancelling build."
-  # See https://docs.readthedocs.com/platform/latest/guides/build/skip-build.html for info on exit code
-  exit 183
-fi
-echo "Docs-affecting files changed; continuing pre-run-check."
-echo "Checking pre-commit/pre-run-check status..."
-MAX_WAIT=300
-INTERVAL=60
-ELAPSED=0
-while :; do
-  RAW=$(curl -sS -w "\n%{http_code}" "https://api.github.com/repos/vllm-project/vllm/commits/${READTHEDOCS_GIT_COMMIT_HASH}/check-runs?check_name=pre-run-check&filter=latest")
-  HTTP_CODE=$(printf %s "$RAW" | tail -n1)
-  BODY=$(printf %s "$RAW" | sed '$d')
-  if [ "$HTTP_CODE" != "200" ]; then
-    echo "GitHub API returned HTTP $HTTP_CODE (likely rate-limited); skipping pre-commit/pre-run-check gate."
-    break
-  fi
-  STATUS=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"status\") or \"\") if r else \"none\")")
-  CONCLUSION=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"conclusion\") or \"\") if r else \"\")")
-  CHECK_URL=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"html_url\") or \"\") if r else \"\")")
-  if [ "$STATUS" = "none" ]; then
-    echo "no pre-commit/pre-run-check found for this commit; skipping gate."
-    break
-  fi
-  if [ -n "$CONCLUSION" ]; then
-    echo "pre-commit/pre-run-check conclusion: $CONCLUSION"
-    if [ "$CONCLUSION" = "failure" ] || [ "$CONCLUSION" = "cancelled" ] || [ "$CONCLUSION" = "timed_out" ]; then
-      echo "pre-commit/pre-run-check did not pass; skipping docs build."
-      if [ -n "$CHECK_URL" ]; then
-        echo "pre-commit/pre-run-check failure reason: $CHECK_URL"
-      fi
-      exit 1
+if [ "$READTHEDOCS_VERSION_TYPE" = "external" ]; then
+  MAX_WAIT=300
+  INTERVAL=60
+  ELAPSED=0
+  while :; do
+    RAW=$(curl -sS -w "\n%{http_code}" "https://api.github.com/repos/vllm-project/vllm/commits/${READTHEDOCS_GIT_COMMIT_HASH}/check-runs?check_name=pre-run-check&filter=latest")
+    HTTP_CODE=$(printf %s "$RAW" | tail -n1)
+    BODY=$(printf %s "$RAW" | sed '$d')
+    if [ "$HTTP_CODE" != "200" ]; then
+      echo "GitHub API returned HTTP $HTTP_CODE (likely rate-limited); skipping pre-run-check gate."
+      break
    fi
-    break
-  fi
-  if [ "$ELAPSED" -ge "$MAX_WAIT" ]; then
-    echo "pre-commit/pre-run-check status=$STATUS after ${MAX_WAIT}s; skipping gate."
-    break
-  fi
-  echo "pre-commit/pre-run-check status=$STATUS; waiting ${INTERVAL}s..."
-  sleep "$INTERVAL"
-  ELAPSED=$((ELAPSED + INTERVAL))
-done
+    STATUS=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"status\") or \"\") if r else \"none\")")
+    CONCLUSION=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"conclusion\") or \"\") if r else \"\")")
+    CHECK_URL=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"html_url\") or \"\") if r else \"\")")
+    if [ "$STATUS" = "none" ]; then
+      echo "no pre-run-check found for this commit; skipping gate."
+      break
+    fi
+    if [ -n "$CONCLUSION" ]; then
+      echo "pre-run-check conclusion: $CONCLUSION"
+      if [ "$CONCLUSION" = "failure" ] || [ "$CONCLUSION" = "cancelled" ] || [ "$CONCLUSION" = "timed_out" ]; then
+        echo "pre-run-check did not pass; skipping docs build."
+        if [ -n "$CHECK_URL" ]; then
+          echo "pre-run-check failure reason: $CHECK_URL"
+        fi
+        exit 1
+      fi
+      break
+    fi
+    if [ "$ELAPSED" -ge "$MAX_WAIT" ]; then
+      echo "pre-run-check status=$STATUS after ${MAX_WAIT}s; skipping gate."
+      break
+    fi
+    echo "pre-run-check status=$STATUS; waiting ${INTERVAL}s..."
+    sleep "$INTERVAL"
+    ELAPSED=$((ELAPSED + INTERVAL))
+  done
+else
+  echo "Not a PR build (version type=$READTHEDOCS_VERSION_TYPE); skipping pre-run-check gate."
+fi
@@ -151,7 +151,7 @@ Configure EPLB with the `--eplb-config` argument, which accepts a JSON string. T
 | `step_interval` | Frequency of rebalancing (every N engine steps) | 3000 |
 | `log_balancedness` | Log balancedness metrics (avg tokens per expert ÷ max tokens per expert) | `false` |
 | `num_redundant_experts` | Additional global experts per EP rank beyond equal distribution | `0` |
-| `use_async` | Use non-blocking EPLB for reduced latency overhead | `true` |
+| `use_async` | Use non-blocking EPLB for reduced latency overhead | `false` |
 | `policy` | The policy type for expert parallel load balancing | `"default"` |
 | `communicator` | Backend for expert weight transfers: `"torch_nccl"`, `"torch_gloo"`, `"pynccl"`, `"nixl"`,  or `null` (auto) | `null` |

@@ -100,44 +100,14 @@ For further details on renderer APIs, please refer to [this page](renderer.md).
 - `/version` - Version information
 - `/load` - Server load metrics

-## Server in development mode
-
-When using the flag VLLM_SERVER_DEV_MODE=1, you enable development endpoints.
-
-**SECURITY WARNING: These endpoints should NOT be used in production!**
-
-### Cache Management APIs
-
- `/reset_prefix_cache` - Reset prefix cache (can disrupt service)
- `/reset_mm_cache` - Reset multimodal cache (can disrupt service)
- `/reset_encoder_cache` - Reset encoder cache (can disrupt service)
-
-### Weight Transfer APIs (RL Training)
-
-For further details on Weight Transfer, please refer to [this page](../../training/weight_transfer/README.md).
-
- `/pause` - Pause generation (causes denial of service)
- `/resume` - Resume generation
- `/is_paused` - Check if generation is paused
- `/init_weight_transfer_engine` - Initialize weight transfer engine for RLHF
- `/update_weights` - Update model weights (can alter model behavior)
- `/get_world_size` - Get distributed world size
-
-### Collective RPC
-
- `/collective_rpc` - Execute arbitrary RPC methods on the engine (extremely dangerous)
-
-### Server info
-
- `/server_info` - Get detailed server configuration
-
-### Sleep Mode APIs
+## Sleep Mode APIs

 For further details on sleep mode, please refer to [this page](../../features/sleep_mode.md).

 - `/sleep` - Put engine to sleep (causes denial of service)
 - `/wake_up` - Wake engine from sleep
 - `/is_sleeping` - Check if engine is sleeping
+- `/collective_rpc` - Execute arbitrary RPC methods on the engine (extremely dangerous)

 ## Chat Template

@@ -84,10 +84,7 @@ Both the trainer (`NCCLTrainerSendWeightsArgs`) and inference side (`NCCLWeightT

 ## Receiving Weights (Inference Side)

-The inference side triggers weight reception using the four-phase protocol:
-`init_weight_transfer_engine`, `start_weight_update`, `update_weights`,
-`finish_weight_update`. The init phase is shown [above](#initialization). The
-remaining three steps are:
+The inference side triggers weight reception using the four-phase protocol — `init_weight_transfer_engine`, `start_weight_update`, `update_weights`, `finish_weight_update`. The init phase is shown [above](#initialization). The remaining three steps are:

 ```python
 from vllm.distributed.weight_transfer.base import WeightTransferUpdateRequest
@@ -111,24 +108,12 @@ llm.update_weights(
 llm.finish_weight_update()
 ```

-The `names`, `dtype_names`, and `shapes` lists describe each parameter. These
-must match the order in which the trainer iterates over its parameters.
+The `names`, `dtype_names`, and `shapes` lists describe each parameter. These must match the order in which the trainer iterates over its parameters.

-`start_weight_update` must be called before `update_weights`, and
-`finish_weight_update` must be called after all weight chunks have been
-transferred. The `is_checkpoint_format` flag controls whether layerwise reload
-processing is applied (`True` for checkpoint-format weights, `False` for
-pre-processed kernel-format weights).
-
-Sparse NCCL patches still use `update_kind="sparse_flat"` inside
-`update_info`, but they should be wrapped in
-`start_weight_update(is_checkpoint_format=False)` because sparse patches apply
-directly to runtime/kernel-format parameters. The current sparse MVP requires
-`TP=1` and `PP=1`.
+`start_weight_update` must be called before `update_weights`, and `finish_weight_update` must be called after all weight chunks have been transferred. The `is_checkpoint_format` flag controls whether layerwise reload processing is applied (`True` for checkpoint-format weights, `False` for pre-processed kernel-format weights).

 ## Examples

 - [RLHF with NCCL weight syncing (offline, Ray)](../../../examples/rl/rlhf_nccl.py) - Trainer on one GPU, 2x tensor-parallel vLLM engine on two others, with packed NCCL weight broadcast
- [RLHF with sparse NCCL weight syncing (offline, Ray)](../../../examples/rl/rlhf_sparse_nccl.py) - Dense-vs-sparse equivalence demo with a real model on a 2-GPU trainer/inference setup; sparse patches use `start_weight_update(is_checkpoint_format=False)` and currently require `TP=1` and `PP=1`
 - [RLHF with async weight syncing (offline, Ray)](../../../examples/rl/rlhf_async_new_apis.py) - Async generation with mid-flight pause, weight sync, resume, and validation against a fresh model
 - [RLHF with NCCL weight syncing (online serving, HTTP)](../../../examples/rl/rlhf_http_nccl.py) - Weight transfer with a running vLLM HTTP server using HTTP control plane and NCCL data plane
@@ -203,7 +203,7 @@ def run_multi_image(model: str, max_completion_tokens: int) -> None:

 # Video input inference
 def run_video(model: str, max_completion_tokens: int) -> None:
-    video_url = "https://huggingface.co/datasets/raushan-testing-hf/videos-test/resolve/main/sample_demo_1.mp4"
+    video_url = "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/ForBiggerFun.mp4"
    video_base64 = encode_base64_content_from_url(video_url)

    ## Use video url in the payload
@@ -1,526 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Demonstrates dense-vs-sparse NCCL weight syncing with a real model.
-
-This example mirrors the validation story used for the sparse NCCL MVP:
-both the dense update path and the sparse patch path start from the same real
-checkpoint and apply the same deterministic trainer-side patch. The script then
-checks that greedy 1-token outputs match between the dense and sparse vLLM
-engines after the update.
-
-The example performs the following steps:
-* Load a training model on one GPU via a Ray actor.
-* Launch a vLLM engine with the same real model on a second GPU.
-* Verify trainer vs vLLM baseline agreement before any update.
-* Apply a deterministic patch to ``model.embed_tokens.weight`` on the trainer.
-* Run a dense NCCL update into a fresh vLLM engine and collect post-update
-  outputs.
-* Reset the trainer back to the baseline checkpoint.
-* Apply the same deterministic patch again.
-* Run a sparse NCCL update into another fresh vLLM engine and collect
-  post-update outputs.
-* Compare dense vs sparse baseline outputs, dense vs sparse post-update
-  outputs, estimated payload sizes, and trainer-side send times.
-
-Current sparse weight transfer MVP limitations:
-* ``TP=1`` and ``PP=1`` only
-* sparse updates use runtime/kernel-format parameter names
-* sparse updates are not composable with checkpoint-format or packed updates
-
-This example assumes a single-node cluster with two GPUs.
-"""
-
-import hashlib
-import os
-import time
-from collections.abc import Sequence
-
-import ray
-import torch
-from ray.util.placement_group import placement_group
-from ray.util.scheduling_strategies import PlacementGroupSchedulingStrategy
-from transformers import AutoModelForCausalLM, AutoTokenizer
-
-from vllm import LLM, SamplingParams
-from vllm.config import WeightTransferConfig
-from vllm.distributed.weight_transfer.base import SparseWeightPatch
-from vllm.distributed.weight_transfer.nccl_engine import (
-    NCCLTrainerSendWeightsArgs,
-    NCCLWeightTransferEngine,
-)
-from vllm.utils.network_utils import get_ip, get_open_port
-
-MODEL_NAME = "Qwen/Qwen2.5-0.5B-Instruct"
-PATCHED_PARAM_NAME = "model.embed_tokens.weight"
-MAX_PATCH_ROWS = 32
-PROMPTS = [
-    "Hello, my name is",
-    "The president of the United States is",
-    "The capital of France is",
-    "The future of AI is",
-]
-SAMPLING_PARAMS = SamplingParams(temperature=0.0, max_tokens=1)
-
-
-class MyLLM(LLM):
-    """Configure the vLLM worker for Ray placement group execution."""
-
-    def __init__(self, *args, **kwargs):
-        os.environ["VLLM_RAY_BUNDLE_INDICES"] = "0"
-        super().__init__(*args, **kwargs)
-
-
-@ray.remote(num_gpus=1)
-class TrainModel:
-    """Ray actor that owns the trainer-side model and deterministic patch state."""
-
-    def __init__(self, model_name: str):
-        self.model_name = model_name
-        self.tokenizer = AutoTokenizer.from_pretrained(model_name)
-        if self.tokenizer.pad_token_id is None:
-            self.tokenizer.pad_token = self.tokenizer.eos_token
-
-        self.model = None
-        self.patched_param = None
-        self.pending_sparse_patches: list[SparseWeightPatch] | None = None
-        self.model_update_group = None
-        self.master_address = get_ip()
-        self.port = get_open_port()
-        self.reset_model()
-
-    def reset_model(self) -> None:
-        self.model = AutoModelForCausalLM.from_pretrained(
-            self.model_name,
-            torch_dtype=torch.bfloat16,
-        ).to("cuda:0")
-        self.model.eval()
-
-        try:
-            self.patched_param = self.model.get_parameter(PATCHED_PARAM_NAME)
-        except AttributeError as exc:
-            raise RuntimeError(
-                f"Expected trainer model to expose `{PATCHED_PARAM_NAME}`"
-            ) from exc
-
-        self.pending_sparse_patches = None
-
-    def create_rendezvous(self) -> tuple[str, int]:
-        self.port = get_open_port()
-        return self.master_address, self.port
-
-    def init_weight_transfer_group(self, world_size: int) -> None:
-        self.model_update_group = NCCLWeightTransferEngine.trainer_init(
-            dict(
-                master_address=self.master_address,
-                master_port=self.port,
-                world_size=world_size,
-            )
-        )
-
-    def get_dense_update_info(self, packed: bool = False) -> tuple[dict, int]:
-        names = []
-        dtype_names = []
-        shapes = []
-        payload_bytes = 0
-        for name, param in self.model.named_parameters():
-            names.append(name)
-            dtype_names.append(str(param.dtype).split(".")[-1])
-            shapes.append(list(param.shape))
-            payload_bytes += param.numel() * param.element_size()
-
-        return (
-            dict(
-                names=names,
-                dtype_names=dtype_names,
-                shapes=shapes,
-                packed=packed,
-            ),
-            payload_bytes,
-        )
-
-    @torch.inference_mode()
-    def generate(
-        self,
-        prompts: Sequence[str],
-        max_new_tokens: int = 1,
-    ) -> list[dict[str, object]]:
-        generations = []
-        for prompt in prompts:
-            model_inputs = self.tokenizer(prompt, return_tensors="pt").to("cuda:0")
-            output = self.model.generate(
-                **model_inputs,
-                max_new_tokens=max_new_tokens,
-                do_sample=False,
-                pad_token_id=self.tokenizer.pad_token_id,
-            )
-            new_token_ids = output[0, model_inputs["input_ids"].shape[1] :].tolist()
-            generations.append(
-                {
-                    "token_ids": new_token_ids,
-                    "text": self.tokenizer.decode(
-                        new_token_ids,
-                        skip_special_tokens=False,
-                    ),
-                }
-            )
-        return generations
-
-    def prepare_sparse_patch(
-        self,
-        prompts: Sequence[str],
-        max_patch_rows: int = MAX_PATCH_ROWS,
-    ) -> tuple[dict[str, object], list[int], str, int]:
-        selected_token_ids: list[int] = []
-        special_ids = set(self.tokenizer.all_special_ids)
-        for prompt in prompts:
-            token_ids = self.tokenizer(prompt, add_special_tokens=False)["input_ids"]
-            for token_id in token_ids:
-                if token_id in special_ids or token_id in selected_token_ids:
-                    continue
-                selected_token_ids.append(token_id)
-                if len(selected_token_ids) == max_patch_rows:
-                    break
-            if len(selected_token_ids) == max_patch_rows:
-                break
-
-        if not selected_token_ids:
-            raise ValueError("Could not derive any non-special token IDs to patch")
-
-        vocab_size = self.patched_param.shape[0]
-        next_token_id = selected_token_ids[-1]
-        while len(selected_token_ids) < max_patch_rows:
-            next_token_id = (next_token_id + 1) % vocab_size
-            if next_token_id in special_ids or next_token_id in selected_token_ids:
-                continue
-            selected_token_ids.append(next_token_id)
-
-        row_ids = torch.tensor(
-            selected_token_ids,
-            device=self.patched_param.device,
-            dtype=torch.long,
-        )
-        hidden_size = self.patched_param.shape[1]
-        column_offsets = torch.arange(
-            hidden_size,
-            device=self.patched_param.device,
-            dtype=torch.long,
-        )
-
-        with torch.no_grad():
-            # Rotate the selected embedding rows instead of zeroing them so the
-            # patch remains deterministic while avoiding a degenerate collapse
-            # to the same special token after the update.
-            replacement_rows = self.patched_param[row_ids].roll(shifts=1, dims=0)
-            self.patched_param[row_ids] = replacement_rows
-
-        flat_indices = (
-            row_ids.unsqueeze(1).mul(hidden_size).add(column_offsets).reshape(-1)
-        )
-        flat_values = self.patched_param[row_ids].reshape(-1).contiguous()
-        self.pending_sparse_patches = [
-            SparseWeightPatch(
-                name=PATCHED_PARAM_NAME,
-                indices=flat_indices.to(torch.int32),
-                values=flat_values,
-            )
-        ]
-        patch_digest = hashlib.sha256(
-            self.pending_sparse_patches[0].indices.cpu().numpy().tobytes()
-            + self.pending_sparse_patches[0]
-            .values.detach()
-            .float()
-            .cpu()
-            .numpy()
-            .tobytes()
-        ).hexdigest()
-
-        sparse_payload_bytes = (
-            flat_indices.numel() * torch.tensor([], dtype=torch.int32).element_size()
-            + flat_values.numel() * flat_values.element_size()
-        )
-        update_info = dict(
-            names=[PATCHED_PARAM_NAME],
-            dtype_names=[str(self.patched_param.dtype).split(".")[-1]],
-            shapes=[list(self.patched_param.shape)],
-            num_updates_list=[flat_indices.numel()],
-            update_kind="sparse_flat",
-        )
-        return update_info, selected_token_ids, patch_digest, sparse_payload_bytes
-
-    def broadcast_weights(self, packed: bool = False) -> float:
-        if self.model_update_group is None:
-            raise RuntimeError("Weight transfer group is not initialized")
-
-        trainer_args = NCCLTrainerSendWeightsArgs(
-            group=self.model_update_group,
-            packed=packed,
-        )
-        start = time.perf_counter()
-        NCCLWeightTransferEngine.trainer_send_weights(
-            iterator=self.model.named_parameters(),
-            trainer_args=trainer_args,
-        )
-        torch.accelerator.synchronize()
-        return (time.perf_counter() - start) * 1000.0
-
-    def broadcast_pending_sparse_patch(self) -> float:
-        if self.model_update_group is None:
-            raise RuntimeError("Weight transfer group is not initialized")
-        if self.pending_sparse_patches is None:
-            raise RuntimeError("Sparse patch has not been prepared")
-
-        start = time.perf_counter()
-        NCCLWeightTransferEngine.trainer_send_sparse_weights(
-            iter(self.pending_sparse_patches),
-            NCCLTrainerSendWeightsArgs(group=self.model_update_group),
-        )
-        torch.accelerator.synchronize()
-        self.pending_sparse_patches = None
-        return (time.perf_counter() - start) * 1000.0
-
-
-def launch_llm(
-    scheduling_inference: PlacementGroupSchedulingStrategy,
-):
-    return ray.remote(
-        num_cpus=0,
-        num_gpus=0,
-        scheduling_strategy=scheduling_inference,
-    )(MyLLM).remote(
-        model=MODEL_NAME,
-        enforce_eager=True,
-        tensor_parallel_size=1,
-        distributed_executor_backend="ray",
-        gpu_memory_utilization=0.7,
-        weight_transfer_config=WeightTransferConfig(backend="nccl"),
-    )
-
-
-def collect_vllm_generations(llm_handle) -> list[dict[str, object]]:
-    outputs = ray.get(llm_handle.generate.remote(PROMPTS, SAMPLING_PARAMS))
-    generations = []
-    for output in outputs:
-        generations.append(
-            {
-                "token_ids": output.outputs[0].token_ids,
-                "text": output.outputs[0].text,
-            }
-        )
-    return generations
-
-
-def token_sequences_match(
-    left: Sequence[dict[str, object]],
-    right: Sequence[dict[str, object]],
-) -> bool:
-    return [item["token_ids"] for item in left] == [item["token_ids"] for item in right]
-
-
-def print_generations(label: str, prompts: Sequence[str], generations) -> None:
-    print(f"\n{label}")
-    print("-" * 50)
-    for prompt, generation in zip(prompts, generations):
-        print(f"Prompt: {prompt!r}")
-        print(f"Token IDs: {generation['token_ids']}")
-        print(f"Text: {generation['text']!r}")
-        print("-" * 50)
-
-
-def run_dense_phase(
-    train_model,
-    scheduling_inference: PlacementGroupSchedulingStrategy,
-) -> dict[str, object]:
-    ray.get(train_model.reset_model.remote())
-    llm = launch_llm(scheduling_inference)
-    try:
-        dense_before = collect_vllm_generations(llm)
-
-        ray.get(llm.sleep.remote(level=0))
-        master_address, master_port = ray.get(train_model.create_rendezvous.remote())
-        world_size = ray.get(llm.get_world_size.remote()) + 1
-        inference_init = llm.init_weight_transfer_engine.remote(
-            dict(
-                init_info=dict(
-                    master_address=master_address,
-                    master_port=master_port,
-                    rank_offset=1,
-                    world_size=world_size,
-                )
-            )
-        )
-        trainer_init = train_model.init_weight_transfer_group.remote(world_size)
-        ray.get([trainer_init, inference_init])
-        ray.get(llm.start_weight_update.remote(is_checkpoint_format=True))
-
-        dense_update_info, dense_payload_bytes = ray.get(
-            train_model.get_dense_update_info.remote()
-        )
-        _, selected_token_ids, patch_digest, _ = ray.get(
-            train_model.prepare_sparse_patch.remote(PROMPTS)
-        )
-
-        inference_update = llm.update_weights.remote(
-            dict(update_info=dense_update_info)
-        )
-        dense_send_ms, _ = ray.get(
-            [
-                train_model.broadcast_weights.remote(packed=False),
-                inference_update,
-            ]
-        )
-        ray.get(llm.finish_weight_update.remote())
-        ray.get(llm.wake_up.remote(tags=["scheduling"]))
-
-        dense_after = collect_vllm_generations(llm)
-
-        return {
-            "dense_before": dense_before,
-            "dense_after": dense_after,
-            "selected_token_ids": selected_token_ids,
-            "patch_digest": patch_digest,
-            "dense_payload_bytes": dense_payload_bytes,
-            "dense_send_ms": dense_send_ms,
-        }
-    finally:
-        ray.kill(llm)
-
-
-def run_sparse_phase(
-    train_model,
-    scheduling_inference: PlacementGroupSchedulingStrategy,
-) -> dict[str, object]:
-    ray.get(train_model.reset_model.remote())
-    llm = launch_llm(scheduling_inference)
-    try:
-        sparse_before = collect_vllm_generations(llm)
-
-        ray.get(llm.sleep.remote(level=0))
-        master_address, master_port = ray.get(train_model.create_rendezvous.remote())
-        world_size = ray.get(llm.get_world_size.remote()) + 1
-        inference_init = llm.init_weight_transfer_engine.remote(
-            dict(
-                init_info=dict(
-                    master_address=master_address,
-                    master_port=master_port,
-                    rank_offset=1,
-                    world_size=world_size,
-                )
-            )
-        )
-        trainer_init = train_model.init_weight_transfer_group.remote(world_size)
-        ray.get([trainer_init, inference_init])
-        ray.get(llm.start_weight_update.remote(is_checkpoint_format=False))
-
-        sparse_update_info, selected_token_ids, patch_digest, sparse_payload_bytes = (
-            ray.get(train_model.prepare_sparse_patch.remote(PROMPTS))
-        )
-
-        inference_update = llm.update_weights.remote(
-            dict(update_info=sparse_update_info)
-        )
-        sparse_send_ms, _ = ray.get(
-            [
-                train_model.broadcast_pending_sparse_patch.remote(),
-                inference_update,
-            ]
-        )
-        ray.get(llm.finish_weight_update.remote())
-        ray.get(llm.wake_up.remote(tags=["scheduling"]))
-
-        sparse_after = collect_vllm_generations(llm)
-
-        return {
-            "sparse_before": sparse_before,
-            "sparse_after": sparse_after,
-            "selected_token_ids": selected_token_ids,
-            "patch_digest": patch_digest,
-            "sparse_payload_bytes": sparse_payload_bytes,
-            "sparse_send_ms": sparse_send_ms,
-        }
-    finally:
-        ray.kill(llm)
-
-
-ray.init()
-
-try:
-    train_model = TrainModel.remote(MODEL_NAME)
-
-    pg_inference = placement_group([{"GPU": 1, "CPU": 0}])
-    ray.get(pg_inference.ready())
-    scheduling_inference = PlacementGroupSchedulingStrategy(
-        placement_group=pg_inference,
-        placement_group_capture_child_tasks=True,
-        placement_group_bundle_index=0,
-    )
-
-    dense_results = run_dense_phase(train_model, scheduling_inference)
-    sparse_results = run_sparse_phase(train_model, scheduling_inference)
-
-    baseline_equal = token_sequences_match(
-        dense_results["dense_before"],
-        sparse_results["sparse_before"],
-    )
-    patch_selection_equal = (
-        dense_results["selected_token_ids"] == sparse_results["selected_token_ids"]
-    )
-    patch_digest_equal = dense_results["patch_digest"] == sparse_results["patch_digest"]
-    after_equal = token_sequences_match(
-        dense_results["dense_after"],
-        sparse_results["sparse_after"],
-    )
-    any_output_changed = any(
-        before["token_ids"] != after["token_ids"]
-        for before, after in zip(
-            dense_results["dense_before"],
-            dense_results["dense_after"],
-        )
-    )
-    dense_payload_mb = dense_results["dense_payload_bytes"] / (1024 * 1024)
-    sparse_payload_mb = sparse_results["sparse_payload_bytes"] / (1024 * 1024)
-
-    print_generations(
-        "Dense baseline outputs",
-        PROMPTS,
-        dense_results["dense_before"],
-    )
-    print_generations(
-        "Sparse baseline outputs", PROMPTS, sparse_results["sparse_before"]
-    )
-    print_generations(
-        "Dense outputs after update", PROMPTS, dense_results["dense_after"]
-    )
-    print_generations(
-        "Sparse outputs after update",
-        PROMPTS,
-        sparse_results["sparse_after"],
-    )
-
-    print(f"patched_token_ids = {dense_results['selected_token_ids']}")
-    print(f"patch_selection_equal = {patch_selection_equal}")
-    print(f"dense_patch_digest = {dense_results['patch_digest']}")
-    print(f"sparse_patch_digest = {sparse_results['patch_digest']}")
-    print(f"patch_digest_equal = {patch_digest_equal}")
-    print(f"baseline_equal = {baseline_equal}")
-    print(f"after_equal = {after_equal}")
-    print(f"any_output_changed = {any_output_changed}")
-    print(f"dense_payload_mb = {dense_payload_mb:.2f}")
-    print(f"sparse_payload_mb = {sparse_payload_mb:.2f}")
-    print(f"dense_send_ms = {dense_results['dense_send_ms']:.2f}")
-    print(f"sparse_send_ms = {sparse_results['sparse_send_ms']:.2f}")
-
-    if not baseline_equal:
-        raise RuntimeError(
-            "Dense and sparse phases did not start from the same baseline"
-        )
-    if not patch_selection_equal:
-        raise RuntimeError("Dense and sparse phases used different sparse patches")
-    if not patch_digest_equal:
-        raise RuntimeError("Dense and sparse phases produced different patch values")
-    if not after_equal:
-        raise RuntimeError("Dense and sparse updates produced different outputs")
-    if not any_output_changed:
-        raise RuntimeError("Patch did not change the observed outputs")
-finally:
-    ray.shutdown()
@@ -295,15 +295,6 @@
                                {%- endif -%}
                            {%- endfor -%}
                            {{- format_tool_response_block(ns_tname.name, ns_txt.s) -}}
-                            {%- for part in tool_body -%}
-                                {%- if part.get('type') == 'image' -%}
-                                    {{- '<|image|>' -}}
-                                {%- elif part.get('type') == 'audio' -%}
-                                    {{- '<|audio|>' -}}
-                                {%- elif part.get('type') == 'video' -%}
-                                    {{- '<|video|>' -}}
-                                {%- endif -%}
-                            {%- endfor -%}
                        {%- else -%}
                            {{- format_tool_response_block(ns_tname.name, tool_body) -}}
                        {%- endif -%}
@@ -17,4 +17,4 @@ torchaudio
 torchvision

 auto_round_lib>=0.13.0
-vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.9/vllm_xpu_kernels-0.1.9-cp38-abi3-manylinux_2_28_x86_64.whl
+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
@@ -93,7 +93,7 @@ pub struct ChatLlm {
    text: TextLlm,
    backend: DynChatBackend,
    /// Effective model dtype reported by the engine.
-    model_dtype: ModelDtype,
+    model_dtype: Option<ModelDtype>,
    /// Tool-call parser selection.
    tool_call_parser: ParserSelection,
    /// Reasoning parser selection.
@@ -135,7 +135,7 @@ impl ChatLlm {
    }

    /// Override the effective model dtype used for multimodal tensor encoding.
-    pub fn with_model_dtype(mut self, model_dtype: ModelDtype) -> Self {
+    pub fn with_model_dtype(mut self, model_dtype: Option<ModelDtype>) -> Self {
        self.model_dtype = model_dtype;
        self
    }
@@ -233,7 +233,7 @@ mod tests {
        )
        .unwrap_err();

-        expect_test::expect!["tool parser `definitely_missing_tool_parser` is not registered (choose from: deepseek_v3, deepseek_v31, deepseek_v32, deepseek_v4, gemma4, glm45, glm47, hermes, hy_v3, internlm, kimi_k2, llama3_json, llama4_json, minimax_m2, mistral, qwen3_coder, qwen3_xml)"].assert_eq(&error.to_report_string());
+        expect_test::expect!["tool parser `definitely_missing_tool_parser` is not registered (choose from: deepseek_v3, deepseek_v31, deepseek_v32, deepseek_v4, gemma4, glm45, glm47, hermes, hy_v3, kimi_k2, llama3_json, llama4_json, minimax_m2, mistral, qwen3_coder, qwen3_xml)"].assert_eq(&error.to_report_string());
    }

    #[test]
@@ -12,7 +12,7 @@
 use std::collections::{HashMap, HashSet};
 use std::fs;
 use std::path::Path;
-use std::sync::{Arc, LazyLock};
+use std::sync::{Arc, LazyLock, Once};

 use itertools::izip;
 use llm_multimodal::{
@@ -239,7 +239,7 @@ pub(crate) async fn finalize_rendered_prompt(
    request: &ChatRequest,
    rendered: RenderedPrompt,
    info: Option<&MultimodalModelInfo>,
-    model_dtype: ModelDtype,
+    model_dtype: Option<ModelDtype>,
 ) -> Result<(Prompt, Option<MmFeatures>)> {
    if !request.has_multimodal() {
        return Ok((rendered.prompt, None));
@@ -249,6 +249,16 @@ pub(crate) async fn finalize_rendered_prompt(
        bail_multimodal!("multimodal chat renderer must return a text prompt before expansion");
    };
    let media_parts = extract_media_parts(request)?;
+    let model_dtype = model_dtype.unwrap_or_else(|| {
+        static WARN_ONCE: Once = Once::new();
+        WARN_ONCE.call_once(|| {
+            warn!(
+                "engine handshake did not report model dtype; \
+                 falling back to float32 for multimodal tensor encoding"
+            );
+        });
+        ModelDtype::Float32
+    });

    let mut prompt_token_ids = info
        .context
@@ -5,9 +5,9 @@ use std::sync::LazyLock;
 pub use vllm_tool_parser::{
    DeepSeekV3ToolParser, DeepSeekV4ToolParser, DeepSeekV31ToolParser, DeepSeekV32ToolParser,
    Gemma4ToolParser, Glm45MoeToolParser, Glm47MoeToolParser, HermesToolParser, HyV3ToolParser,
-    Internlm2ToolParser, KimiK2ToolParser, Llama3JsonToolParser, MinimaxM2ToolParser,
-    MistralToolParser, Qwen3CoderToolParser, Qwen3XmlToolParser, ToolCallDelta, ToolParser,
-    ToolParserError, ToolParserOutput,
+    KimiK2ToolParser, Llama3JsonToolParser, MinimaxM2ToolParser, MistralToolParser,
+    Qwen3CoderToolParser, Qwen3XmlToolParser, ToolCallDelta, ToolParser, ToolParserError,
+    ToolParserOutput,
 };

 use crate::parser::ParserFactory;
@@ -24,9 +24,6 @@ pub mod names {
    pub const GEMMA4: &str = "gemma4";
    pub const HERMES: &str = "hermes";
    pub const HY_V3: &str = "hy_v3";
-    // Matches the Python CLI name `--tool-call-parser internlm`, which Python
-    // also routes to `Internlm2ToolParser` despite the version-agnostic name.
-    pub const INTERNLM: &str = "internlm";
    pub const KIMI_K2: &str = "kimi_k2";
    pub const LLAMA3_JSON: &str = "llama3_json";
    pub const LLAMA4_JSON: &str = "llama4_json";
@@ -65,7 +62,6 @@ impl ToolParserFactory {
            .register_parser::<Gemma4ToolParser>(names::GEMMA4)
            .register_parser::<HermesToolParser>(names::HERMES)
            .register_parser::<HyV3ToolParser>(names::HY_V3)
-            .register_parser::<Internlm2ToolParser>(names::INTERNLM)
            .register_parser::<KimiK2ToolParser>(names::KIMI_K2)
            .register_parser::<Llama3JsonToolParser>(names::LLAMA3_JSON)
            .register_parser::<Llama3JsonToolParser>(names::LLAMA4_JSON)
@@ -84,12 +80,6 @@ impl ToolParserFactory {
            .register_pattern("hermes", names::HERMES)
            .register_pattern("hy3", names::HY_V3)
            .register_pattern("hy_v3", names::HY_V3)
-            // Narrow to `internlm2` substring so it matches `internlm2-chat-7b`
-            // and `internlm2_5-7b-chat` but NOT `internlm-chat-7b` (InternLM v1,
-            // routes to Llama), `internlm3-*` (also Llama-architecture per
-            // vllm/model_executor/models/registry.py:146), or `Intern-S1` /
-            // `Intern-S1-Pro` (separate intern-s1 parser, see PR #40115).
-            .register_pattern("internlm2", names::INTERNLM)
            .register_pattern("llama-4", names::LLAMA4_JSON)
            .register_pattern("llama-3.2", names::LLAMA3_JSON)
            .register_pattern("llama-3.1", names::LLAMA3_JSON)
--- a/Show More
+++ b/Show More