[CI] Filter import-only files using function-level coverage

Skip files where only module-level code ran (imports, class defs)
but no named functions were actually called. Uses the
functions_called field from stripped coverage JSON.

Reduces false-positive mappings by ~78% — e.g. ompmultiprocessing.py
drops from 73 steps to 0 (only used on CPU but imported everywhere).

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

.buildkite/ci_config_rocm.yaml

@@ -1,23 +0,0 @@
-name: vllm_rocm_ci
-job_dirs:
-  - ".buildkite/hardware_tests"
-run_all_patterns:
-  - "docker/Dockerfile.rocm"
-  - "docker/Dockerfile.rocm_base"
-  - "docker/ci-rocm.hcl"
-  - "docker/docker-bake-rocm.hcl"
-  - ".buildkite/hardware_tests/amd.yaml"
-  - ".buildkite/scripts/ci-bake-rocm.sh"
-  - ".buildkite/scripts/hardware_ci/run-amd-test.py"
-  - ".buildkite/scripts/hardware_ci/run-amd-test.sh"
-  - "CMakeLists.txt"
-  - "requirements/common.txt"
-  - "requirements/rocm.txt"
-  - "requirements/build/rocm.txt"
-  - "requirements/test/rocm.txt"
-  - "setup.py"
-  - "csrc/"
-  - "cmake/"
-run_all_exclude_patterns:
-  - "csrc/cpu/"
-  - "cmake/cpu_extension.cmake"

.buildkite/hardware_tests/amd.yaml

@@ -1,73 +1,22 @@
-group: Hardware - AMD Build
+group: Hardware - AMD Build 
 steps:
-  # Ensure ci_base is up-to-date before building the test image.
-  # Compares a content hash of ci_base-affecting files against the remote
-  # image label. If hashes match the build is skipped (< 30 s); if they
-  # differ ci_base is rebuilt and pushed automatically.
-  - label: "AMD: :docker: ensure ci_base"
-    key: ensure-ci-base-amd
+  - label: "AMD: :docker: build image"
+    key: image-build-amd
     depends_on: []
     device: amd_cpu
     no_plugin: true
     commands:
-      - bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
+    - >
+      docker build
+      --build-arg max_jobs=16
+      --build-arg REMOTE_VLLM=1
+      --build-arg ARG_PYTORCH_ROCM_ARCH='gfx90a;gfx942;gfx950'
+      --build-arg VLLM_BRANCH=$BUILDKITE_COMMIT
+      --tag "rocm/vllm-ci:${BUILDKITE_COMMIT}"
+      -f docker/Dockerfile.rocm
+      --target test
+      --no-cache
+      --progress plain .
+    - docker push "rocm/vllm-ci:${BUILDKITE_COMMIT}"
     env:
       DOCKER_BUILDKIT: "1"
-      VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
-      PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
-      REMOTE_VLLM: "1"
-      VLLM_BRANCH: "$BUILDKITE_COMMIT"
-    retry:
-      automatic:
-        - exit_status: -1  # Agent was lost
-          limit: 1
-        - exit_status: -10  # Agent was lost
-          limit: 1
-
-  - label: "AMD: :docker: build test image and artifacts"
-    key: image-build-amd
-    depends_on:
-      - ensure-ci-base-amd
-    device: amd_cpu
-    no_plugin: true
-    commands:
-      - |
-        if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" ]]; then
-          echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
-          IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
-        else
-          bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
-        fi
-      - |
-        docker run --rm --network=none --entrypoint /bin/bash "rocm/vllm-ci:${BUILDKITE_COMMIT}" -ec '
-          if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
-          if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
-          if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
-          if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
-          command -v python3
-          command -v uv
-          command -v pytest
-          if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
-            echo No ROCm CLI found in image >&2
-            exit 1
-          fi
-          python3 - <<PY
-        import torch, vllm
-        print(torch.__version__)
-        print(vllm.__version__)
-        PY
-          echo AMD image smoke OK
-        '
-    env:
-      DOCKER_BUILDKIT: "1"
-      VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
-      PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
-      IMAGE_TAG: "rocm/vllm-ci:$BUILDKITE_COMMIT"
-      REMOTE_VLLM: "1"
-      VLLM_BRANCH: "$BUILDKITE_COMMIT"
-    retry:
-      automatic:
-        - exit_status: -1  # Agent was lost
-          limit: 1
-        - exit_status: -10  # Agent was lost
-          limit: 1

.buildkite/hardware_tests/cpu.yaml

@@ -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 "
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
       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: []
@@ -81,10 +74,11 @@ steps:
   no_plugin: true
   source_file_dependencies:
   - csrc/cpu/
+  - vllm/model_executor/layers/quantization/cpu_wna16.py
   - vllm/model_executor/layers/quantization/auto_gptq.py
   - vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w8a8_int8.py
-  - vllm/model_executor/kernels/linear/mixed_precision/cpu.py
-  - vllm/model_executor/kernels/linear/scaled_mm/cpu.py
+  - vllm/model_executor/layers/quantization/kernels/scaled_mm/cpu.py
+  - vllm/model_executor/layers/quantization/kernels/mixed_precision/cpu.py
   - vllm/model_executor/layers/fused_moe/experts/cpu_moe.py
   - tests/quantization/test_compressed_tensors.py
   - tests/quantization/test_cpu_wna16.py

.buildkite/image_build/image_build.yaml

@@ -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
@@ -110,21 +98,3 @@ steps:
           limit: 2
         - exit_status: -10  # Agent was lost
           limit: 2
-
-  - label: ":docker: Build arm64 image"
-    key: arm64-image-build
-    depends_on: []
-    source_file_dependencies:
-      - ".buildkite/image_build/image_build.yaml"
-      - ".buildkite/image_build/image_build_arm64.sh"
-      - "docker/Dockerfile"
-    commands:
-    - .buildkite/image_build/image_build_arm64.sh $REGISTRY $REPO $BUILDKITE_COMMIT
-    env:
-      DOCKER_BUILDKIT: "1"
-    retry:
-      automatic:
-        - exit_status: -1  # Agent was lost
-          limit: 2
-        - exit_status: -10  # Agent was lost
-          limit: 2

.buildkite/image_build/image_build_arm64.sh

@@ -1,37 +0,0 @@
-#!/bin/bash
-set -e
-
-if [[ $# -lt 3 ]]; then
-  echo "Usage: $0 <registry> <repo> <commit>"
-  exit 1
-fi
-
-REGISTRY=$1
-REPO=$2
-BUILDKITE_COMMIT=$3
-
-# authenticate with AWS ECR
-aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
-
-# skip build if image already exists
-if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64) ]]; then
-  echo "Image not found, proceeding with build..."
-else
-  echo "Image found"
-  exit 0
-fi
-
-# build (Grace/GH200 is the arm64 GPU target; sm_90)
-docker build --file docker/Dockerfile \
-  --platform linux/arm64 \
-  --build-arg max_jobs=16 \
-  --build-arg nvcc_threads=4 \
-  --build-arg torch_cuda_arch_list="9.0" \
-  --build-arg USE_SCCACHE=1 \
-  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64 \
-  --target test \
-  --progress plain .
-
-# push
-docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64

.buildkite/performance-benchmarks/tests/latency-tests-arm64-cpu.json

@@ -2,6 +2,7 @@
     {
         "test_name": "latency_llama8B_tp1",
         "environment_variables": {
+            "VLLM_RPC_TIMEOUT": 100000,
             "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
             "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
             "VLLM_CPU_KVCACHE_SPACE": 40

.buildkite/performance-benchmarks/tests/latency-tests-cpu.json

@@ -2,6 +2,7 @@
     {
         "test_name": "latency_llama8B_tp2",
         "environment_variables": {
+            "VLLM_RPC_TIMEOUT": 100000,
 	    "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
 	    "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
 	    "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-arm64-cpu.json

@@ -13,6 +13,7 @@
       200
     ],
     "server_environment_variables": {
+      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-cpu-asr.json

@@ -5,6 +5,7 @@
     ],
     "max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
     "server_environment_variables": {
+      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120
     },
     "server_parameters": {

.buildkite/performance-benchmarks/tests/serving-tests-cpu-embed.json

@@ -9,6 +9,7 @@
       128
     ],
     "server_environment_variables": {
+      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-cpu-text.json

@@ -5,6 +5,7 @@
     ],
     "max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
     "server_environment_variables": {
+      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-cpu.json

@@ -5,6 +5,7 @@
     ],
     "max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
     "server_environment_variables": {
+      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/throughput-tests-arm64-cpu.json

@@ -2,6 +2,7 @@
     {
         "test_name": "throughput_llama8B_tp1",
         "environment_variables": {
+            "VLLM_RPC_TIMEOUT": 100000,
             "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
             "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
             "VLLM_CPU_KVCACHE_SPACE": 40

.buildkite/performance-benchmarks/tests/throughput-tests-cpu.json

@@ -2,6 +2,7 @@
     {
         "test_name": "throughput_llama8B_tp2",
         "environment_variables": {
+            "VLLM_RPC_TIMEOUT": 100000,
 	    "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
 	    "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
 	    "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/release-pipeline.yaml

@@ -737,7 +737,7 @@ steps:
       - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
     env:
       S3_BUCKET: "vllm-wheels"
-      VARIANT: "rocm723"
+      VARIANT: "rocm722"
 
   # ROCm Job 6: Build ROCm Release Docker Image
   - label: ":docker: Build release image - x86_64 - ROCm"

.buildkite/scripts/ci-fetch-log.sh

@@ -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
 

.buildkite/scripts/coverage/aggregate-coverage.py

@@ -0,0 +1,208 @@
+#!/usr/bin/env python3
+"""Aggregate per-step coverage JSON files into a test-selection mapping.
+
+Downloads all coverage_*.json artifacts from the current Buildkite build,
+then produces two output files:
+
+1. coverage_map.json — inverted index: {source_file: [step_keys]}
+   Used by the pipeline generator to determine which steps to trigger.
+
+2. step_coverage.json — forward index: {step_key: [source_files]}
+   Useful for debugging and understanding test coverage.
+
+Usage:
+    # Run as a Buildkite step at the end of nightly CI
+    python3 .buildkite/scripts/coverage/aggregate-coverage.py
+
+    # Or locally with downloaded artifacts
+    python3 .buildkite/scripts/coverage/aggregate-coverage.py --local-dir ./artifacts/
+"""
+
+import argparse
+import json
+import os
+import subprocess
+import sys
+import tempfile
+from collections import defaultdict
+from pathlib import Path
+
+
+def download_artifacts(dest_dir: str) -> list[str]:
+    """Download all coverage_*.json artifacts from the current build."""
+    try:
+        subprocess.run(
+            ["buildkite-agent", "artifact", "download", "coverage_*.json", dest_dir],
+            check=True,
+            capture_output=True,
+            text=True,
+        )
+    except FileNotFoundError:
+        print("buildkite-agent not found, skipping download", file=sys.stderr)
+        return []
+    except subprocess.CalledProcessError as e:
+        print(f"Artifact download failed: {e.stderr}", file=sys.stderr)
+        return []
+
+    return list(Path(dest_dir).glob("coverage_*.json"))
+
+
+def load_coverage_files(files: list[Path]) -> dict[str, list[str]]:
+    """Load coverage JSON files and extract source files per step.
+
+    Returns: {step_key: [source_files]}
+    """
+    step_coverage = {}
+
+    for filepath in files:
+        filename = filepath.name
+        # coverage_<step_key>.json -> step_key
+        step_key = filename.removeprefix("coverage_").removesuffix(".json")
+
+        try:
+            with open(filepath) as f:
+                data = json.load(f)
+        except (json.JSONDecodeError, OSError) as e:
+            print(f"Warning: skipping {filename}: {e}", file=sys.stderr)
+            continue
+
+        source_files = []
+        for fpath, fdata in data.get("files", {}).items():
+            # Skip files with zero executed lines — coverage.py reports
+            # all files in the source tree, not just those actually run.
+            # Supports both full format (summary.covered_lines) and
+            # stripped format (covered_lines directly).
+            covered = fdata.get("covered_lines") or fdata.get("summary", {}).get("covered_lines", 0)
+            if covered == 0:
+                continue
+            # If function-level data is available, skip import-only files
+            # (files where only module-level code ran but no named functions
+            # were actually called).
+            funcs_called = fdata.get("functions_called")
+            if funcs_called is not None and funcs_called == 0:
+                continue
+            # Normalize paths to be relative to the vllm package root.
+            # coverage.py may report absolute paths or paths relative to
+            # the installed package location. We only care about files
+            # under the vllm/ directory.
+            normalized = _normalize_path(fpath)
+            if normalized:
+                source_files.append(normalized)
+
+        if source_files:
+            step_coverage[step_key] = sorted(set(source_files))
+            print(f"  {step_key}: {len(source_files)} source files")
+
+    return step_coverage
+
+
+def _normalize_path(path: str) -> str | None:
+    """Normalize a coverage path to a vllm-relative path.
+
+    Returns None for paths outside the vllm package (tests, third-party, etc).
+    """
+    # Strip common prefixes from installed package paths
+    markers = ["/site-packages/", "/dist-packages/", "/vllm-workspace/src/"]
+    for marker in markers:
+        idx = path.find(marker)
+        if idx != -1:
+            path = path[idx + len(marker):]
+            break
+
+    # Also handle paths that are already relative
+    if path.startswith("vllm/"):
+        return path
+
+    # Handle absolute paths that contain /vllm/
+    idx = path.find("/vllm/")
+    if idx != -1:
+        return path[idx + 1:]
+
+    return None
+
+
+def build_inverted_index(
+    step_coverage: dict[str, list[str]],
+) -> dict[str, list[str]]:
+    """Build {source_file: [step_keys]} from {step_key: [source_files]}."""
+    inverted = defaultdict(list)
+    for step_key, source_files in step_coverage.items():
+        for src_file in source_files:
+            inverted[src_file].append(step_key)
+
+    # Sort step lists for deterministic output
+    return {k: sorted(v) for k, v in sorted(inverted.items())}
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument(
+        "--local-dir",
+        help="Directory containing coverage_*.json files (skip artifact download)",
+    )
+    parser.add_argument(
+        "--output-dir",
+        default=".",
+        help="Directory to write output files (default: cwd)",
+    )
+    args = parser.parse_args()
+
+    if args.local_dir:
+        artifact_dir = args.local_dir
+        files = list(Path(artifact_dir).glob("coverage_*.json"))
+    else:
+        artifact_dir = tempfile.mkdtemp(prefix="coverage_artifacts_")
+        files = download_artifacts(artifact_dir)
+
+    if not files:
+        print("No coverage files found. Nothing to aggregate.")
+        sys.exit(0)
+
+    print(f"Found {len(files)} coverage files:")
+
+    # Build the forward index: step -> source files
+    step_coverage = load_coverage_files(files)
+
+    if not step_coverage:
+        print("No valid coverage data found.")
+        sys.exit(0)
+
+    # Build the inverted index: source file -> steps
+    coverage_map = build_inverted_index(step_coverage)
+
+    # Write outputs
+    output_dir = Path(args.output_dir)
+    output_dir.mkdir(parents=True, exist_ok=True)
+
+    step_coverage_path = output_dir / "step_coverage.json"
+    with open(step_coverage_path, "w") as f:
+        json.dump(step_coverage, f, indent=2)
+    print(f"\nWrote {step_coverage_path} ({len(step_coverage)} steps)")
+
+    coverage_map_path = output_dir / "coverage_map.json"
+    with open(coverage_map_path, "w") as f:
+        json.dump(coverage_map, f, indent=2)
+    print(f"Wrote {coverage_map_path} ({len(coverage_map)} source files)")
+
+    # Summary stats
+    total_files = len(coverage_map)
+    total_mappings = sum(len(v) for v in coverage_map.values())
+    print(f"\nSummary: {total_files} source files mapped to "
+          f"{len(step_coverage)} steps ({total_mappings} total mappings)")
+
+    # Upload aggregated files as artifacts
+    for output_file in [step_coverage_path, coverage_map_path]:
+        try:
+            subprocess.run(
+                ["buildkite-agent", "artifact", "upload", str(output_file)],
+                check=True,
+                capture_output=True,
+                text=True,
+            )
+            print(f"Uploaded {output_file}")
+        except (FileNotFoundError, subprocess.CalledProcessError):
+            pass  # Not in Buildkite or upload failed — that's fine for local runs
+
+
+if __name__ == "__main__":
+    main()

.buildkite/scripts/coverage/upload-step-coverage.sh

@@ -0,0 +1,42 @@
+#!/bin/bash
+# Upload coverage data for the current Buildkite step.
+# Called automatically at the end of each step when COLLECT_COVERAGE=1.
+#
+# Expects:
+#   - .coverage.${BUILDKITE_STEP_KEY} data file from coverage run --append
+#   - BUILDKITE_STEP_KEY, BUILDKITE_BUILD_NUMBER env vars
+#
+# Produces:
+#   - coverage_${BUILDKITE_STEP_KEY}.json uploaded as a Buildkite artifact
+
+set -euo pipefail
+
+STEP_KEY="${BUILDKITE_STEP_KEY:-unknown}"
+DATA_FILE=".coverage.${STEP_KEY}"
+OUTPUT_JSON="coverage_${STEP_KEY}.json"
+
+if [ ! -f "$DATA_FILE" ]; then
+    echo "~~~ No coverage data file found ($DATA_FILE), skipping upload"
+    exit 0
+fi
+
+echo "~~~ :bar_chart: Exporting coverage data for step: ${STEP_KEY}"
+
+coverage json \
+    --data-file="$DATA_FILE" \
+    -o "$OUTPUT_JSON" \
+    --omit='*/tests/*,*/test_*,*/__pycache__/*' \
+    2>&1 || {
+        echo "Warning: coverage json export failed, skipping"
+        exit 0
+    }
+
+FILE_COUNT=$(python3 -c "import json; d=json.load(open('$OUTPUT_JSON')); print(len(d.get('files', {})))" 2>/dev/null || echo "?")
+echo "Coverage captured ${FILE_COUNT} source files for step ${STEP_KEY}"
+
+buildkite-agent artifact upload "$OUTPUT_JSON" 2>&1 || {
+    echo "Warning: artifact upload failed"
+    exit 0
+}
+
+echo "Uploaded $OUTPUT_JSON"

.buildkite/scripts/hardware_ci/run-amd-test.sh

@@ -35,9 +35,25 @@ export PYTHONPATH=".."
 # Helper Functions
 ###############################################################################
 
-report_docker_usage() {
-  echo "--- Docker usage"
-  docker system df || true
+cleanup_docker() {
+  # Get Docker's root directory
+  docker_root=$(docker info -f '{{.DockerRootDir}}')
+  if [ -z "$docker_root" ]; then
+    echo "Failed to determine Docker root directory."
+    exit 1
+  fi
+  echo "Docker root directory: $docker_root"
+
+  disk_usage=$(df "$docker_root" | tail -1 | awk '{print $5}' | sed 's/%//')
+  threshold=70
+  if [ "$disk_usage" -gt "$threshold" ]; then
+    echo "Disk usage is above $threshold%. Cleaning up Docker images and volumes..."
+    docker image prune -f
+    docker volume prune -f && docker system prune --force --filter "until=72h" --all
+    echo "Docker images and volumes cleanup completed."
+  else
+    echo "Disk usage is below $threshold%. No cleanup needed."
+  fi
 }
 
 cleanup_network() {
@@ -52,108 +68,6 @@ cleanup_network() {
   fi
 }
 
-prepare_artifact_image() {
-  if [[ "${VLLM_CI_USE_ARTIFACTS:-0}" != "1" ]]; then
-    return 1
-  fi
-  if ! command -v buildkite-agent >/dev/null 2>&1; then
-    echo "buildkite-agent not found; cannot download ROCm wheel artifact"
-    return 1
-  fi
-
-  local artifact_glob="${VLLM_CI_ARTIFACT_GLOB:-artifacts/vllm-rocm-install/vllm-rocm-install.tar.gz}"
-  local archive=""
-  local metadata_file=""
-  local base_image="${VLLM_CI_BASE_IMAGE:-rocm/vllm-dev:ci_base}"
-  local artifact_image=""
-  local artifact_key=""
-  local base_digest=""
-  local wheel_dir=""
-  local context_dir=""
-  local workspace_dir=""
-
-  artifact_work_dir=$(mktemp -d -t vllm-rocm-artifact.XXXXXX)
-  wheel_dir="${artifact_work_dir}/wheels"
-  context_dir="${artifact_work_dir}/context"
-  workspace_dir="${context_dir}/workspace"
-  mkdir -p "${wheel_dir}" "${context_dir}/wheels" "${workspace_dir}"
-
-  echo "--- Downloading ROCm wheel artifact"
-  if ! buildkite-agent artifact download "${artifact_glob}" "${artifact_work_dir}"; then
-    echo "Failed to download ${artifact_glob}"
-    return 1
-  fi
-  buildkite-agent artifact download \
-    "artifacts/vllm-rocm-install/ci-base-image.txt" \
-    "${artifact_work_dir}" >/dev/null 2>&1 || true
-
-  archive=$(find "${artifact_work_dir}" -name "vllm-rocm-install.tar.gz" -type f | head -1)
-  if [[ -z "${archive}" || ! -f "${archive}" ]]; then
-    echo "ROCm wheel artifact archive was not found"
-    return 1
-  fi
-
-  metadata_file=$(find "${artifact_work_dir}" -name "ci-base-image.txt" -type f | head -1)
-  if [[ -n "${metadata_file}" && -s "${metadata_file}" ]]; then
-    base_image=$(tr -d '[:space:]' < "${metadata_file}")
-  fi
-
-  echo "--- Preparing local ROCm test image"
-  echo "Base image: ${base_image}"
-  docker pull "${base_image}" || return 1
-  base_digest=$(
-    docker image inspect \
-      --format='{{if .RepoDigests}}{{index .RepoDigests 0}}{{else}}{{.Id}}{{end}}' \
-      "${base_image}" 2>/dev/null || printf '%s' "${base_image}"
-  )
-
-  artifact_key=$(
-    {
-      printf 'base-image:%s\n' "${base_digest}"
-      sha256sum "${archive}"
-    } | sha256sum | cut -c1-24
-  )
-  artifact_image="rocm/vllm-ci-artifact:${artifact_key}"
-
-  if docker image inspect "${artifact_image}" >/dev/null 2>&1; then
-    echo "Using existing local ROCm artifact image: ${artifact_image}"
-    image_name="${artifact_image}"
-    return 0
-  fi
-
-  tar -xzf "${archive}" -C "${wheel_dir}" || return 1
-  if ! ls "${wheel_dir}"/*.whl >/dev/null 2>&1; then
-    echo "ROCm wheel artifact did not contain a wheel"
-    return 1
-  fi
-  if [[ ! -d "${wheel_dir}/tests" ]]; then
-    echo "ROCm wheel artifact did not contain the test workspace"
-    return 1
-  fi
-
-  cp "${wheel_dir}"/*.whl "${context_dir}/wheels/" || return 1
-  tar -C "${wheel_dir}" --exclude='*.whl' -cf - . \
-    | tar -C "${workspace_dir}" -xf - || return 1
-  cat > "${context_dir}/Dockerfile" <<'EOF'
-ARG BASE_IMAGE
-FROM ${BASE_IMAGE}
-COPY wheels/ /tmp/vllm-wheels/
-COPY workspace/ /vllm-workspace/
-RUN python3 -m pip install --no-deps --force-reinstall /tmp/vllm-wheels/*.whl \
-    && rm -rf /tmp/vllm-wheels
-WORKDIR /vllm-workspace
-EOF
-
-  echo "--- Building local ROCm test image"
-  docker build \
-    --pull=false \
-    --build-arg "BASE_IMAGE=${base_image}" \
-    -t "${artifact_image}" \
-    "${context_dir}" || return 1
-  image_name="${artifact_image}"
-  return 0
-}
-
 is_multi_node() {
   local cmds="$1"
   # Primary signal: NUM_NODES environment variable set by the pipeline
@@ -340,36 +254,20 @@ re_quote_pytest_markers() {
 echo "--- ROCm info"
 rocminfo
 
-# --- Docker status ---
-report_docker_usage
+# --- Docker housekeeping ---
+cleanup_docker
 
 # --- Pull test image ---
 echo "--- Pulling container"
-image_name="${VLLM_CI_FALLBACK_IMAGE:-rocm/vllm-ci:${BUILDKITE_COMMIT:-local}}"
-artifact_work_dir=""
+image_name="rocm/vllm-ci:${BUILDKITE_COMMIT}"
 container_name="rocm_${BUILDKITE_COMMIT}_$(tr -dc A-Za-z0-9 < /dev/urandom | head -c 10; echo)"
+docker pull "${image_name}"
 
 remove_docker_container() {
-  if docker container inspect "${container_name}" >/dev/null 2>&1; then
-    docker rm -f "${container_name}" || true
-  fi
-  if [[ "${VLLM_CI_REMOVE_TEST_IMAGE:-0}" == "1" ]]; then
-    docker image rm -f "${image_name}" || true
-  else
-    # Keep images by default so later jobs on the same AMD node can reuse layers.
-    echo "Keeping ROCm test image locally: ${image_name}"
-  fi
-  if [[ -n "${artifact_work_dir}" ]]; then
-    rm -rf "${artifact_work_dir}"
-  fi
+  docker rm -f "${container_name}" || docker image rm -f "${image_name}" || true
 }
 trap remove_docker_container EXIT
 
-if ! prepare_artifact_image; then
-  echo "Using full ROCm CI image: ${image_name}"
-  docker pull "${image_name}" || exit 1
-fi
-
 # --- Prepare commands ---
 echo "--- Running container"
 

.buildkite/scripts/hardware_ci/run-cpu-test-arm.sh

@@ -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

.buildkite/scripts/install-kv-connectors.sh

@@ -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

.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh

@@ -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" \

.buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh

@@ -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 \

.buildkite/scripts/tool_call/run-bfcl-eval.sh

@@ -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

.buildkite/test-amd.yaml

@@ -449,6 +449,29 @@ steps:
   commands:
   - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_qwen3_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py
 
+#------------------------------------------------------  mi250 · model_executor  -------------------------------------------------------#
+
+- label: Model Executor # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
+  agent_pool: mi250_1
+  optional: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/engine/arg_utils.py
+  - vllm/config/model.py
+  - vllm/model_executor
+  - tests/model_executor
+  - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  commands:
+  - apt-get update && apt-get install -y curl libsodium23
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s model_executor -m '(not slow_test)'
+  - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
+
 #------------------------------------------------------  mi250 · models / basic  -------------------------------------------------------#
 
 - label: Basic Models Test (Other CPU) # TBD
@@ -1215,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
@@ -1235,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
@@ -1249,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
@@ -1346,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
@@ -1462,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
@@ -1504,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
@@ -1716,29 +1738,6 @@ steps:
   - pytest -v -s -x lora/test_gptoss_tp.py
   - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
 
-#------------------------------------------------------  mi300 · model_executor  -------------------------------------------------------#
-
-- label: Model Executor # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
-  agent_pool: mi300_1
-  optional: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/engine/arg_utils.py
-  - vllm/config/model.py
-  - vllm/model_executor
-  - tests/model_executor
-  - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - vllm/_aiter_ops.py
-  - vllm/platforms/rocm.py
-  commands:
-  - apt-get update && apt-get install -y curl libsodium23
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s model_executor -m '(not slow_test)'
-  - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
-
 #-----------------------------------------------------  mi300 · models / language  -----------------------------------------------------#
 
 - label: Language Models Test (Extended Pooling)  # TBD
@@ -2704,35 +2703,19 @@ steps:
   optional: true
   working_dir: "/vllm-workspace/"
   source_file_dependencies:
-  - csrc/custom_quickreduce.cu
-  - csrc/ops.h
-  - csrc/torch_bindings.cpp
   - vllm/distributed/
-  - vllm/model_executor/layers/
-  - vllm/entrypoints/llm.py
-  - vllm/config/parallel.py
-  - vllm/model_executor/layers/fused_moe/
-  - vllm/v1/engine/
-  - vllm/v1/executor/
-  - vllm/v1/worker/
   - vllm/v1/distributed/
+  - vllm/model_executor/layers/fused_moe/
   - vllm/v1/attention/backends/
   - vllm/v1/attention/selector.py
-  - vllm/_aiter_ops.py
-  - vllm/_custom_ops.py
-  - vllm/platforms/rocm.py
-  - vllm/envs.py
-  - examples/offline_inference/data_parallel.py
   - tests/distributed/test_context_parallel.py
-  - tests/distributed/test_rocm_quick_reduce.py
-  - tests/distributed/test_quick_all_reduce.py
   - tests/v1/distributed/test_dbo.py
-  - tests/utils.py
+  - examples/features/data_parallel/data_parallel_offline.py
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
   commands:
   - pytest -v -s tests/distributed/test_context_parallel.py
   - pytest -v -s tests/v1/distributed/test_dbo.py
-  - pytest -v -s tests/distributed/test_rocm_quick_reduce.py
-  - pytest -v -s tests/distributed/test_quick_all_reduce.py
 
 #--------------------------------------------------------  mi355 · entrypoints  --------------------------------------------------------#
 
@@ -2746,11 +2729,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 +2752,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 +2766,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 +2879,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

.buildkite/test_areas/disaggregated.yaml

@@ -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

.buildkite/test_areas/e2e_integration.yaml

@@ -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

.buildkite/test_areas/engine.yaml

@@ -38,7 +38,7 @@ steps:
     - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       timeout_in_minutes: 40
       depends_on:
       - image-build-amd

.buildkite/test_areas/entrypoints.yaml

@@ -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
@@ -28,8 +28,7 @@ steps:
   - pytest -v -s entrypoints/offline_mode # Needs to avoid interference with other tests
   mirror:
     amd:
-      device: mi325_1
-      soft_fail: true
+      device: mi300_1
       depends_on:
       - image-build-amd
 
@@ -43,11 +42,10 @@ 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
-      soft_fail: true
+      device: mi300_1
       timeout_in_minutes: 80
       depends_on:
       - image-build-amd
@@ -60,17 +58,12 @@ 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
-      soft_fail: true
+      device: mi300_1
       timeout_in_minutes: 60
       depends_on:
       - image-build-amd
@@ -89,8 +82,7 @@ steps:
   - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
   mirror:
     amd:
-      device: mi325_1
-      soft_fail: true
+      device: mi300_1
       timeout_in_minutes: 60
       depends_on:
       - image-build-amd
@@ -102,15 +94,17 @@ 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
-      soft_fail: true
+      device: mi300_1
       depends_on:
       - image-build-amd
 
@@ -154,5 +148,6 @@ steps:
   source_file_dependencies:
   - csrc/
   - vllm/entrypoints/openai/
+  - vllm/model_executor/models/whisper.py
   commands: # LMEval
   - pytest -s entrypoints/openai/correctness/

.buildkite/test_areas/kernels.yaml

@@ -38,28 +38,6 @@ steps:
   commands:
     - pytest -v -s kernels/core/test_minimax_reduce_rms.py
 
-- label: Deepseek V4 Kernel Test (H100)
-  key: deepseek-v4-kernel-test-h100
-  timeout_in_minutes: 15
-  device: h100
-  source_file_dependencies:
-  - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
-  - vllm/models/deepseek_v4/common/ops/
-  - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
-  commands:
-    - pytest -v -s kernels/test_fused_deepseek_v4_*.py
-
-- label: Deepseek V4 Kernel Test (B200)
-  key: deepseek-v4-kernel-test-b200
-  timeout_in_minutes: 15
-  device: b200-k8s
-  source_file_dependencies:
-  - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
-  - vllm/models/deepseek_v4/common/ops/
-  - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
-  commands:
-    - pytest -v -s kernels/test_fused_deepseek_v4_*.py
-
 - label: Kernels Attention Test %N
   key: kernels-attention-test
   timeout_in_minutes: 35
@@ -86,7 +64,7 @@ steps:
   parallelism: 2
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       source_file_dependencies:
       - csrc/quantization/
       - vllm/model_executor/layers/quantization

.buildkite/test_areas/misc.yaml

@@ -52,7 +52,7 @@ steps:
     - pytest -v -s v1/test_outputs.py
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       depends_on:
       - image-build-amd
 
@@ -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
@@ -281,7 +281,6 @@ steps:
   - vllm/model_executor/layers/quantization/quark/
   - vllm/multimodal/
   - vllm/outputs.py
-  - vllm/parser/
   - vllm/platforms/
   - vllm/pooling_params.py
   - vllm/ray/

.buildkite/test_areas/model_executor.yaml

@@ -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

.buildkite/test_areas/model_runner_v2.yaml

@@ -94,13 +94,11 @@ steps:
     - vllm/v1/worker/gpu_worker.py
     - tests/distributed/test_pipeline_parallel.py
     - tests/distributed/test_pp_cudagraph.py
-    - tests/v1/distributed/test_pp_dp_v2.py
   commands:
     - set -x
     - export VLLM_USE_V2_MODEL_RUNNER=1
     - pytest -v -s distributed/test_pipeline_parallel.py -k "not ray and not Jamba"
     - pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
-    - pytest -v -s v1/distributed/test_pp_dp_v2.py
 
 - label: Model Runner V2 Spec Decode
   device: h200_35gb

.buildkite/test_areas/models_basic.yaml

@@ -58,3 +58,37 @@ steps:
   device: cpu-small
   commands:
     - pytest -v -s models/test_utils.py models/test_vision.py
+
+- label: Transformers Nightly Models
+  device: h200_35gb
+  key: transformers-nightly-models
+  working_dir: "/vllm-workspace/"
+  optional: true
+  soft_fail: true
+  commands:
+    - pip install --upgrade git+https://github.com/huggingface/transformers
+    - pytest -v -s tests/models/test_initialization.py
+    - pytest -v -s tests/models/test_transformers.py
+    - pytest -v -s tests/models/multimodal/processing/
+    - pytest -v -s tests/models/multimodal/test_mapping.py
+    - python3 examples/basic/offline_inference/chat.py
+    - python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
+    # Whisper needs spawn method to avoid deadlock
+    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
+
+- label: Transformers Backward Compatibility Models Test
+  device: h200_35gb
+  key: transformers-backward-compatibility-models-test
+  working_dir: "/vllm-workspace/"
+  optional: true
+  soft_fail: true
+  commands:
+    - pip install transformers==4.57.5
+    - pytest -v -s tests/models/test_initialization.py
+    - pytest -v -s tests/models/test_transformers.py
+    - pytest -v -s tests/models/multimodal/processing/
+    - pytest -v -s tests/models/multimodal/test_mapping.py
+    - python3 examples/basic/offline_inference/chat.py
+    - python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
+    # Whisper needs spawn method to avoid deadlock
+    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper

.buildkite/test_areas/models_language.yaml

@@ -50,7 +50,7 @@ steps:
   mirror:
     torch_nightly: {}
     amd:
-      device: mi325_1
+      device: mi300_1
       depends_on:
       - image-build-amd
       commands:
@@ -96,7 +96,7 @@ steps:
     - pytest -v -s models/language/pooling -m 'not core_model'
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       timeout_in_minutes: 100
       depends_on:
       - image-build-amd

.buildkite/test_areas/models_multimodal.yaml

@@ -15,7 +15,7 @@ steps:
     - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       depends_on:
       - image-build-amd
 
@@ -33,7 +33,7 @@ steps:
     - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       depends_on:
       - image-build-amd
 
@@ -50,7 +50,7 @@ steps:
     - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       depends_on:
       - image-build-amd
 
@@ -118,7 +118,7 @@ steps:
     - pytest -v -s models/multimodal/test_mapping.py
   mirror:
     amd:
-      device: mi325_1
+      device: mi300_1
       depends_on:
       - image-build-amd
 

.buildkite/test_areas/rust_frontend.yaml

@@ -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

.buildkite/test_areas/spec_decode.yaml

@@ -32,7 +32,6 @@ steps:
   source_file_dependencies:
     - vllm/v1/spec_decode/
     - vllm/v1/worker/gpu/spec_decode/
-    - vllm/v1/attention/backends/
     - vllm/transformers_utils/configs/speculators/
     - tests/v1/e2e/spec_decode/
   commands:

.dockerignore

@@ -33,10 +33,3 @@ share/python-wheels/
 *.egg
 MANIFEST
 rust/target/
-# Not needed in Docker builds
-docs/
-.github/
-.pre-commit-config.yaml
-.clang-format
-.gitattributes
-format.sh

.github/CODEOWNERS

@@ -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,23 +72,21 @@
 /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
-/tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche @AndreasKaratzas
-/tests/evals @mgoin @vadiklyutiy @AndreasKaratzas
-/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye @AndreasKaratzas
+/tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche
+/tests/evals @mgoin @vadiklyutiy
+/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye
 /tests/kernels/ir @ProExpertProg @tjtanaa
-/tests/models @DarkLight1337 @ywang96 @AndreasKaratzas
+/tests/models @DarkLight1337 @ywang96
 /tests/multimodal @DarkLight1337 @ywang96 @NickLucche
-/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye @AndreasKaratzas
+/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye
 /tests/test_inputs.py @DarkLight1337 @ywang96
 /tests/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
 /tests/v1/structured_output @mgoin @russellb @aarnphm
@@ -179,20 +171,20 @@ mkdocs.yaml @hmellor
 
 # ROCm related: specify owner with write access to notify AMD folks for careful code review
 /vllm/**/*rocm* @tjtanaa @dllehr-amd
-/docker/Dockerfile.rocm* @tjtanaa @dllehr-amd @AndreasKaratzas
+/docker/Dockerfile.rocm* @tjtanaa @dllehr-amd
 /vllm/v1/attention/backends/rocm*.py @tjtanaa @dllehr-amd
 /vllm/v1/attention/backends/mla/rocm*.py @tjtanaa @dllehr-amd
 /vllm/v1/attention/ops/rocm*.py @tjtanaa @dllehr-amd
 /vllm/model_executor/layers/fused_moe/rocm*.py @tjtanaa @dllehr-amd
 /csrc/rocm @tjtanaa @dllehr-amd
-/requirements/*rocm* @tjtanaa @AndreasKaratzas
-/tests/**/*rocm* @tjtanaa @AndreasKaratzas
+/requirements/*rocm* @tjtanaa
+/tests/**/*rocm* @tjtanaa
 /docs/**/*rocm* @tjtanaa
 /vllm/**/*quark* @tjtanaa
-/tests/**/*quark* @tjtanaa @AndreasKaratzas
+/tests/**/*quark* @tjtanaa
 /docs/**/*quark* @tjtanaa
-/vllm/**/*aiter* @tjtanaa @AndreasKaratzas
-/tests/**/*aiter* @tjtanaa @AndreasKaratzas
+/vllm/**/*aiter* @tjtanaa
+/tests/**/*aiter* @tjtanaa
 
 # TPU
 /vllm/v1/worker/tpu* @NickLucche

.github/mergify.yml

@@ -103,19 +103,6 @@ pull_request_rules:
       add:
         - frontend
 
-- name: label-rust
-  description: Automatically apply rust label
-  conditions:
-    - label != stale
-    - or:
-      - files~=(?i)rust
-      - title~=(?i)rust
-      - title~=(?i)vllm-rs
-  actions:
-    label:
-      add:
-        - rust
-
 - name: label-llama
   description: Automatically apply llama label
   conditions:

.github/workflows/add_label_automerge.yml

@@ -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({

.github/workflows/issue_autolabel.yml

@@ -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();

.github/workflows/new_pr_bot.yml

@@ -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;

.github/workflows/pre-commit.yml

@@ -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({

.pre-commit-config.yaml

@@ -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

.readthedocs.yaml

@@ -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:

AGENTS.md

@@ -101,8 +101,6 @@ pre-commit run ruff-check --all-files
 pre-commit run mypy-3.10 --all-files --hook-stage manual
 ```
 
-The line length limit for Python code is 88 characters. If you are not sure, use pre-commit to check.
-
 ### Commit messages
 
 Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`). For example:

CMakeLists.txt

@@ -112,8 +112,6 @@ endif()
 #
 # spinloop extension (pure CXX; must stay above the non-CUDA device branch so
 # CPU builds define the target before the early return)
-# This extension requires SABI 3.11 since it relies on Py_buffer support. Loading
-# failure is handled gracefully on vLLM side for lower Python versions.
 #
 set(VLLM_SPINLOOP_EXT_SRC "csrc/spinloop.cpp")
 set(SPINLOOP_COMPILE_FLAGS "")
@@ -146,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
@@ -307,13 +305,26 @@ endif()
 #
 
 set(VLLM_EXT_SRC
+  "csrc/mamba/mamba_ssm/selective_scan_fwd.cu"
+  "csrc/cache_kernels.cu"
+  "csrc/cache_kernels_fused.cu"
+  "csrc/attention/paged_attention_v1.cu"
+  "csrc/attention/paged_attention_v2.cu"
+  "csrc/attention/merge_attn_states.cu"
+  "csrc/sampler.cu"
+  "csrc/topk.cu"
   "csrc/cuda_view.cu"
   "csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
   "csrc/quantization/activation_kernels.cu"
   "csrc/cuda_utils_kernels.cu"
-  "csrc/torch_bindings.cpp")
+  "csrc/custom_all_reduce.cu"
+  "csrc/torch_bindings.cpp"
+  "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
+  list(APPEND VLLM_EXT_SRC
+    "csrc/minimax_reduce_rms_kernel.cu")
+
   SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")
 
   # Set CUTLASS_REVISION. Used for FetchContent. Also fixes some bogus messages when building.
@@ -358,30 +369,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   # are not supported by Machete yet.
 
   # marlin arches for fp16 output
-  # Family-conditional 12.0f (one cubin for SM12x family) requires CUDA >= 13.0;
-  # fall back to architecture-specific 12.0a;12.1a on CUDA < 13.0 (e.g. 12.8).
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MARLIN_ARCHS "8.0+PTX;12.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(MARLIN_ARCHS "8.0+PTX;12.0a;12.1a" "${CUDA_ARCHS}")
-  endif()
+  cuda_archs_loose_intersection(MARLIN_ARCHS "8.0+PTX" "${CUDA_ARCHS}")
   # marlin has limited support for turing
   cuda_archs_loose_intersection(MARLIN_SM75_ARCHS "7.5" "${CUDA_ARCHS}")
   # marlin arches for bf16 output (we need 9.0 for bf16 atomicAdd PTX)
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MARLIN_BF16_ARCHS "8.0+PTX;9.0+PTX;12.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(MARLIN_BF16_ARCHS "8.0+PTX;9.0+PTX;12.0a;12.1a" "${CUDA_ARCHS}")
-  endif()
+  cuda_archs_loose_intersection(MARLIN_BF16_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
   # marlin arches for fp8 input
   # - sm80 doesn't support fp8 computation
   # - sm90 and sm100 don't support QMMA.16832.F32.E4M3.E4M3 SAAS instruction
   # so we only enable fp8 computation for SM89 (e.g. RTX 40x0)  and 12.0 (e.g. RTX 50x0)
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MARLIN_FP8_ARCHS "8.9;12.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(MARLIN_FP8_ARCHS "8.9;12.0a;12.1a" "${CUDA_ARCHS}")
-  endif()
+  cuda_archs_loose_intersection(MARLIN_FP8_ARCHS "8.9;12.0;12.1" "${CUDA_ARCHS}")
   # marlin arches for other files
   cuda_archs_loose_intersection(MARLIN_OTHER_ARCHS "7.5;8.0+PTX" "${CUDA_ARCHS}")
 
@@ -500,12 +497,12 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND ES_MXFP8_GROUPED_MM_ARCHS)
     set(SRCS
-      "csrc/libtorch_stable/moe/mxfp8_moe/cutlass_mxfp8_grouped_mm.cu"
-      "csrc/libtorch_stable/moe/mxfp8_moe/mxfp8_experts_quant.cu")
+      "csrc/moe/mxfp8_moe/cutlass_mxfp8_grouped_mm.cu"
+      "csrc/moe/mxfp8_moe/mxfp8_experts_quant.cu")
     set_gencode_flags_for_srcs(
       SRCS "${SRCS}"
       CUDA_ARCHS "${ES_MXFP8_GROUPED_MM_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
+    list(APPEND VLLM_EXT_SRC "${SRCS}")
     list(APPEND VLLM_GPU_FLAGS "-DENABLE_ES_MXFP8_GROUPED_MM_SM100=1")
     message(STATUS "Building ES MXFP8 grouped kernels for archs: ${ES_MXFP8_GROUPED_MM_ARCHS}")
   else()
@@ -595,7 +592,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
 endif()
 
 if (VLLM_GPU_LANG STREQUAL "HIP")
-  # Add QuickReduce kernels (ROCm-only; not part of stable ABI migration).
+  # Add QuickReduce kernels
   list(APPEND VLLM_EXT_SRC
     "csrc/custom_quickreduce.cu"
   )
@@ -630,26 +627,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     "csrc/libtorch_stable/activation_kernels.cu"
     "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
     "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
-    "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
-    "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
     "csrc/libtorch_stable/quantization/gptq/q_gemm.cu"
     "csrc/libtorch_stable/quantization/gguf/gguf_kernel.cu"
     "csrc/libtorch_stable/pos_encoding_kernels.cu"
     "csrc/libtorch_stable/fused_qknorm_rope_kernel.cu"
     "csrc/libtorch_stable/layernorm_kernels.cu"
     "csrc/libtorch_stable/layernorm_quant_kernels.cu"
-    "csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
-    "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.cu"
-    "csrc/libtorch_stable/cache_kernels_fused.cu"
-    "csrc/libtorch_stable/custom_all_reduce.cu"
-    "csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
+    "csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu")
 
   if(VLLM_GPU_LANG STREQUAL "CUDA")
     list(APPEND VLLM_STABLE_EXT_SRC
@@ -661,8 +645,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       "csrc/libtorch_stable/permute_cols.cu"
       "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
       "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
-      "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu"
-      "csrc/libtorch_stable/minimax_reduce_rms_kernel.cu")
+      "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu")
 
     set_gencode_flags_for_srcs(
       SRCS "${VLLM_STABLE_EXT_SRC}"
@@ -686,22 +669,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)
@@ -943,11 +910,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")
@@ -977,11 +946,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")
@@ -1152,11 +1121,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   # moe marlin arches
   # note that we always set `use_atomic_add=False` for moe marlin now,
   # so we don't need 9.0 for bf16 atomicAdd PTX
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MARLIN_MOE_ARCHS "8.0+PTX;12.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(MARLIN_MOE_ARCHS "8.0+PTX;12.0a;12.1a" "${CUDA_ARCHS}")
-  endif()
+  cuda_archs_loose_intersection(MARLIN_MOE_ARCHS "8.0+PTX" "${CUDA_ARCHS}")
   # moe marlin has limited support for turing
   cuda_archs_loose_intersection(MARLIN_MOE_SM75_ARCHS "7.5" "${CUDA_ARCHS}")
   # moe marlin arches for fp8 input
@@ -1259,22 +1224,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()
@@ -1301,14 +1268,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
@@ -1318,10 +1277,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

benchmarks/kernels/benchmark_moe_permute_unpermute.py

@@ -10,7 +10,6 @@ from transformers import AutoConfig
 
 from vllm.model_executor.layers.fused_moe import fused_topk
 from vllm.model_executor.layers.fused_moe.moe_permute_unpermute import (
-    MoEPermuteScratch,
     moe_permute,
     moe_unpermute,
 )
@@ -55,15 +54,6 @@ def benchmark_permute(
     topk_weights, topk_ids, token_expert_indices = fused_topk(
         qhidden_states, input_gating, topk, False
     )
-    scratch = MoEPermuteScratch(
-        max_num_tokens=num_tokens,
-        topk=topk,
-        num_experts=num_experts,
-        num_local_experts=num_experts,
-        device=qhidden_states.device,
-        hidden_size=hidden_size,
-        hidden_dtype=qhidden_states.dtype,
-    )
 
     def prepare(i: int):
         input_gating.copy_(gating_output[i])
@@ -75,7 +65,6 @@ def benchmark_permute(
             topk_ids=topk_ids,
             n_expert=num_experts,
             expert_map=None,
-            scratch=scratch,
         )
 
     # JIT compilation & warmup
@@ -134,15 +123,6 @@ def benchmark_unpermute(
     topk_weights, topk_ids, token_expert_indices = fused_topk(
         qhidden_states, input_gating, topk, False
     )
-    scratch = MoEPermuteScratch(
-        max_num_tokens=num_tokens,
-        topk=topk,
-        num_experts=num_experts,
-        num_local_experts=num_experts,
-        device=qhidden_states.device,
-        hidden_size=hidden_size,
-        hidden_dtype=qhidden_states.dtype,
-    )
 
     def prepare():
         (
@@ -157,7 +137,6 @@ def benchmark_unpermute(
             topk_ids=topk_ids,
             n_expert=num_experts,
             expert_map=None,
-            scratch=scratch,
         )
         # convert to fp16/bf16 as gemm output
         return (

benchmarks/kernels/benchmark_router_gemm.py

@@ -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)

cmake/cpu_extension.cmake

@@ -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.")

cmake/external_projects/vllm_flash_attn.cmake

@@ -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

cmake/hipify.py

@@ -14,18 +14,7 @@ import argparse
 import os
 import shutil
 
-from torch.utils.hipify.hipify_python import get_hip_file_path, hipify
-
-
-def _expected_hip_build_path(source_abs: str, output_directory: str) -> str:
-    """Match torch.utils.hipify.hipify_python.preprocessor fout_path naming."""
-    rel = os.path.relpath(source_abs, output_directory)
-    return os.path.abspath(
-        os.path.join(
-            output_directory, get_hip_file_path(rel, is_pytorch_extension=True)
-        )
-    )
-
+from torch.utils.hipify.hipify_python import hipify
 
 if __name__ == "__main__":
     parser = argparse.ArgumentParser()
@@ -64,11 +53,7 @@ if __name__ == "__main__":
     hipify_result = hipify(
         project_directory=args.project_dir,
         output_directory=args.output_dir,
-        # Hipify resolves quoted includes next to the including file first; vLLM
-        # uses paths relative to csrc/ (e.g. "libtorch_stable/torch_utils.h"
-        # from quantization/w8a8/fp8/*.cu). Without an include root here, those
-        # headers are never found and are not hipified or rewritten in dependents.
-        header_include_dirs=["."],
+        header_include_dirs=[],
         includes=includes,
         extra_files=extra_files,
         show_detailed=True,
@@ -79,20 +64,14 @@ if __name__ == "__main__":
     hipified_sources = []
     for source in args.sources:
         s_abs = os.path.abspath(source)
-        if s_abs in hipify_result and hipify_result[s_abs].hipified_path is not None:
-            path = hipify_result[s_abs].hipified_path
-            # PyTorch skips writing when is_pytorch_extension and text unchanged;
-            # hipified_path then stays *.cu. CMake expects *.hip under output_dir.
-            if s_abs.endswith(".cu") and path.endswith(".cu"):
-                dest = _expected_hip_build_path(s_abs, args.output_dir)
-                if os.path.normpath(path) != os.path.normpath(dest):
-                    os.makedirs(os.path.dirname(dest), exist_ok=True)
-                    shutil.copy2(path, dest)
-                hipified_s_abs = dest
-            else:
-                hipified_s_abs = path
-        else:
-            hipified_s_abs = s_abs
+        hipified_s_abs = (
+            hipify_result[s_abs].hipified_path
+            if (
+                s_abs in hipify_result
+                and hipify_result[s_abs].hipified_path is not None
+            )
+            else s_abs
+        )
         hipified_sources.append(hipified_s_abs)
 
     assert len(hipified_sources) == len(args.sources)

cmake/utils.cmake

@@ -81,14 +81,6 @@ function (hipify_sources_target OUT_SRCS NAME ORIG_SRCS)
   set_property(GLOBAL APPEND PROPERTY VLLM_HIPIFY_ALL_SRCS ${SRCS})
   set_property(GLOBAL APPEND PROPERTY VLLM_HIPIFY_ALL_BYPRODUCTS ${HIP_SRCS})
 
-  # Chain hipify targets so they run sequentially. Parallel hipify
-  # invocations race on shutil.copytree, overwriting .hip files
-  # produced by another target back to .cu originals.
-  if (DEFINED _VLLM_LAST_HIPIFY_TARGET)
-    add_dependencies(hipify${NAME} ${_VLLM_LAST_HIPIFY_TARGET})
-  endif()
-  set(_VLLM_LAST_HIPIFY_TARGET "hipify${NAME}" PARENT_SCOPE)
-
   # Swap out original extension sources with hipified sources.
   list(APPEND HIP_SRCS ${CXX_SRCS})
   set(${OUT_SRCS} ${HIP_SRCS} PARENT_SCOPE)
@@ -484,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

csrc/attention/attention_kernels.cuh

@@ -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))

csrc/attention/attention_utils.cuh

@@ -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>

csrc/attention/merge_attn_states.cu

@@ -1,14 +1,14 @@
 #include <optional>
+#include <torch/all.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
 #include <algorithm>
 #include <limits>
 
-#include "../torch_utils.h"
-#include "../dispatch_utils.h"
-#include <torch/headeronly/core/ScalarType.h>
-
-#include "../../attention/attention_dtypes.h"
+#include "attention_dtypes.h"
 #include "attention_utils.cuh"
-#include "../../quantization/w8a8/fp8/common.cuh"
+#include "../quantization/w8a8/fp8/common.cuh"
+#include "../dispatch_utils.h"
 
 namespace vllm {
 
@@ -196,17 +196,17 @@ __global__ void merge_attn_states_kernel(
 // The following macro is used to dispatch the conversion function based on
 // the output data type. The FN is a macro that calls a function with
 // template<typename scalar_t>.
-#define DISPATCH_BY_SCALAR_DTYPE(scalar_dtype, fn)                          \
-  {                                                                         \
-    if (scalar_dtype == torch::headeronly::ScalarType::Float) {             \
-      fn(float);                                                            \
-    } else if (scalar_dtype == torch::headeronly::ScalarType::Half) {       \
-      fn(uint16_t);                                                         \
-    } else if (scalar_dtype == torch::headeronly::ScalarType::BFloat16) {   \
-      fn(__nv_bfloat16);                                                    \
-    } else {                                                                \
-      STD_TORCH_CHECK(false, "Unsupported data type of O: ", scalar_dtype); \
-    }                                                                       \
+#define DISPATCH_BY_SCALAR_DTYPE(scalar_dtype, fn)                      \
+  {                                                                     \
+    if (scalar_dtype == at::ScalarType::Float) {                        \
+      fn(float);                                                        \
+    } else if (scalar_dtype == at::ScalarType::Half) {                  \
+      fn(uint16_t);                                                     \
+    } else if (scalar_dtype == at::ScalarType::BFloat16) {              \
+      fn(__nv_bfloat16);                                                \
+    } else {                                                            \
+      TORCH_CHECK(false, "Unsupported data type of O: ", scalar_dtype); \
+    }                                                                   \
   }
 
 #define LAUNCH_MERGE_ATTN_STATES(scalar_t, output_t, NUM_THREADS,           \
@@ -245,14 +245,11 @@ __global__ void merge_attn_states_kernel(
  */
 template <typename scalar_t>
 void merge_attn_states_launcher(
-    torch::stable::Tensor& output,
-    std::optional<torch::stable::Tensor> output_lse,
-    const torch::stable::Tensor& prefix_output,
-    const torch::stable::Tensor& prefix_lse,
-    const torch::stable::Tensor& suffix_output,
-    const torch::stable::Tensor& suffix_lse,
+    torch::Tensor& output, std::optional<torch::Tensor> output_lse,
+    const torch::Tensor& prefix_output, const torch::Tensor& prefix_lse,
+    const torch::Tensor& suffix_output, const torch::Tensor& suffix_lse,
     const std::optional<int64_t> prefill_tokens_with_context,
-    const std::optional<torch::stable::Tensor>& output_scale) {
+    const std::optional<torch::Tensor>& output_scale) {
   constexpr uint NUM_THREADS = 128;
   const uint num_tokens = output.size(0);
   const uint num_heads = output.size(1);
@@ -261,23 +258,23 @@ void merge_attn_states_launcher(
   const uint output_head_stride = output.stride(1);
   // Thread mapping is based on input BF16 pack_size
   const uint pack_size = 16 / sizeof(scalar_t);
-  STD_TORCH_CHECK(head_size % pack_size == 0,
-                  "headsize must be multiple of pack_size:", pack_size);
+  TORCH_CHECK(head_size % pack_size == 0,
+              "headsize must be multiple of pack_size:", pack_size);
 
   const uint prefix_num_tokens =
       prefill_tokens_with_context.has_value()
           ? static_cast<uint>(prefill_tokens_with_context.value())
           : num_tokens;
-  STD_TORCH_CHECK(prefix_num_tokens <= num_tokens,
-                  "prefix_num_tokens must be <= num_tokens");
+  TORCH_CHECK(prefix_num_tokens <= num_tokens,
+              "prefix_num_tokens must be <= num_tokens");
 
   float* output_lse_ptr = nullptr;
   if (output_lse.has_value()) {
-    output_lse_ptr = output_lse.value().mutable_data_ptr<float>();
+    output_lse_ptr = output_lse.value().data_ptr<float>();
   }
   float* output_scale_ptr = nullptr;
   if (output_scale.has_value()) {
-    output_scale_ptr = output_scale.value().mutable_data_ptr<float>();
+    output_scale_ptr = output_scale.value().data_ptr<float>();
   }
   // Process one pack elements per thread. for float, the
   // pack_size is 4 for half/bf16, the pack_size is 8.
@@ -287,15 +284,14 @@ void merge_attn_states_launcher(
   dim3 block(NUM_THREADS);
   dim3 grid((total_threads + NUM_THREADS - 1) / NUM_THREADS);
 
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      prefix_output.get_device_index());
-  auto stream = get_current_cuda_stream();
+  const c10::cuda::OptionalCUDAGuard device_guard(prefix_output.device());
+  auto stream = at::cuda::getCurrentCUDAStream();
 
   if (output_scale.has_value()) {
     // FP8 output path - dispatch on output FP8 type
-    VLLM_STABLE_DISPATCH_FP8_TYPES(
-        output.scalar_type(), "merge_attn_states_fp8",
-        [&] { LAUNCH_MERGE_ATTN_STATES(scalar_t, fp8_t, NUM_THREADS, true); });
+    VLLM_DISPATCH_FP8_TYPES(output.scalar_type(), "merge_attn_states_fp8", [&] {
+      LAUNCH_MERGE_ATTN_STATES(scalar_t, fp8_t, NUM_THREADS, true);
+    });
   } else {
     // Original BF16/FP16/FP32 output path
     LAUNCH_MERGE_ATTN_STATES(scalar_t, scalar_t, NUM_THREADS, false);
@@ -309,29 +305,26 @@ void merge_attn_states_launcher(
         suffix_lse, prefill_tokens_with_context, output_scale);       \
   }
 
-void merge_attn_states(
-    torch::stable::Tensor& output,
-    std::optional<torch::stable::Tensor> output_lse,
-    const torch::stable::Tensor& prefix_output,
-    const torch::stable::Tensor& prefix_lse,
-    const torch::stable::Tensor& suffix_output,
-    const torch::stable::Tensor& suffix_lse,
-    const std::optional<int64_t> prefill_tokens_with_context,
-    const std::optional<torch::stable::Tensor>& output_scale) {
+void merge_attn_states(torch::Tensor& output,
+                       std::optional<torch::Tensor> output_lse,
+                       const torch::Tensor& prefix_output,
+                       const torch::Tensor& prefix_lse,
+                       const torch::Tensor& suffix_output,
+                       const torch::Tensor& suffix_lse,
+                       std::optional<int64_t> prefill_tokens_with_context,
+                       const std::optional<torch::Tensor>& output_scale) {
   if (output_scale.has_value()) {
-    STD_TORCH_CHECK(
-        output.scalar_type() == torch::headeronly::ScalarType::Float8_e4m3fn ||
-            output.scalar_type() ==
-                torch::headeronly::ScalarType::Float8_e4m3fnuz,
-        "output must be FP8 when output_scale is provided, got: ",
-        output.scalar_type());
+    TORCH_CHECK(output.scalar_type() == at::ScalarType::Float8_e4m3fn ||
+                    output.scalar_type() == at::ScalarType::Float8_e4m3fnuz,
+                "output must be FP8 when output_scale is provided, got: ",
+                output.scalar_type());
   } else {
-    STD_TORCH_CHECK(
-        output.scalar_type() == prefix_output.scalar_type(), "output dtype (",
-        output.scalar_type(), ") must match prefix_output dtype (",
-        prefix_output.scalar_type(), ") when output_scale is not set");
+    TORCH_CHECK(output.scalar_type() == prefix_output.scalar_type(),
+                "output dtype (", output.scalar_type(),
+                ") must match prefix_output dtype (",
+                prefix_output.scalar_type(), ") when output_scale is not set");
   }
   // Always dispatch on prefix_output (input) dtype
-  DISPATCH_BY_SCALAR_DTYPE(prefix_output.scalar_type(),
+  DISPATCH_BY_SCALAR_DTYPE(prefix_output.dtype(),
                            CALL_MERGE_ATTN_STATES_LAUNCHER);
 }

csrc/attention/paged_attention_v1.cu

@@ -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)
 }
 

csrc/attention/paged_attention_v2.cu

@@ -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)
 }
 

csrc/cache_kernels_fused.cu

@@ -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);
 }

csrc/cpu/cpu_fused_moe.cpp

@@ -30,12 +30,7 @@
   }()
 
 namespace {
-enum class FusedMOEAct {
-  SiluAndMul,
-  SwigluOAIAndMul,
-  GeluAndMul,
-  GeluTanhAndMul,
-};
+enum class FusedMOEAct { SiluAndMul, SwigluOAIAndMul, GeluAndMul };
 
 FusedMOEAct get_act_type(const std::string& act) {
   if (act == "silu") {
@@ -44,8 +39,6 @@ FusedMOEAct get_act_type(const std::string& act) {
     return FusedMOEAct::SwigluOAIAndMul;
   } else if (act == "gelu") {
     return FusedMOEAct::GeluAndMul;
-  } else if (act == "gelu_tanh") {
-    return FusedMOEAct::GeluTanhAndMul;
   } else {
     TORCH_CHECK(false, "Invalid act type: " + act);
   }
@@ -150,44 +143,6 @@ void gelu_and_mul(float* __restrict__ input, scalar_t* __restrict__ output,
   }
 }
 
-template <typename scalar_t>
-void gelu_tanh_and_mul(float* __restrict__ input, scalar_t* __restrict__ output,
-                       const int32_t m_size, const int32_t n_size,
-                       const int32_t input_stride,
-                       const int32_t output_stride) {
-  using scalar_vec_t = typename cpu_utils::VecTypeTrait<scalar_t>::vec_t;
-  const int32_t dim = n_size / 2;
-  float* __restrict__ gate = input;
-  float* __restrict__ up = input + dim;
-  vec_op::FP32Vec16 one_vec(1.0);
-  vec_op::FP32Vec16 w1_vec(0.7978845608028654);
-  vec_op::FP32Vec16 w2_vec(0.5);
-  vec_op::FP32Vec16 w3_vec(0.044715);
-  alignas(64) float temp[16];
-
-  for (int32_t m = 0; m < m_size; ++m) {
-    for (int32_t n = 0; n < dim; n += 16) {
-      vec_op::FP32Vec16 gate_vec(gate + n);
-      vec_op::FP32Vec16 up_vec(up + n);
-      auto gate_pow3_vec = gate_vec * gate_vec * gate_vec;
-      auto inner_vec = w1_vec * (gate_vec + w3_vec * gate_pow3_vec);
-
-      inner_vec.save(temp);
-      for (int32_t i = 0; i < 16; ++i) {
-        temp[i] = std::tanh(temp[i]);
-      }
-      vec_op::FP32Vec16 tanh_vec(temp);
-      auto gelu_tanh = gate_vec * w2_vec * (one_vec + tanh_vec);
-      auto gated_output_fp32 = up_vec * gelu_tanh;
-      scalar_vec_t gated_output = scalar_vec_t(gated_output_fp32);
-      gated_output.save(output + n);
-    }
-    gate += input_stride;
-    up += input_stride;
-    output += output_stride;
-  }
-}
-
 template <typename scalar_t>
 FORCE_INLINE void apply_gated_act(const FusedMOEAct act,
                                   float* __restrict__ input,
@@ -205,9 +160,6 @@ FORCE_INLINE void apply_gated_act(const FusedMOEAct act,
     case FusedMOEAct::GeluAndMul:
       gelu_and_mul(input, output, m, n, input_stride, output_stride);
       return;
-    case FusedMOEAct::GeluTanhAndMul:
-      gelu_tanh_and_mul(input, output, m, n, input_stride, output_stride);
-      return;
     default:
       TORCH_CHECK(false, "Unsupported act type.");
   }

csrc/cpu/cpu_types_riscv_impl.hpp

@@ -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) {

csrc/cpu/cpu_types_vsx.hpp

@@ -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

csrc/cpu/sgl-kernels/blas_gemm.h

@@ -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

csrc/cpu/sgl-kernels/fla.cpp

@@ -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) {

csrc/cpu/sgl-kernels/gemm.h

@@ -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

csrc/cpu/sgl-kernels/vec.h

@@ -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;

csrc/cpu/shm.cpp

@@ -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__

csrc/cpu/torch_bindings.cpp

@@ -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 "

csrc/custom_all_reduce.cu

@@ -1,11 +1,7 @@
-#include "torch_utils.h"
-
-#include <torch/csrc/stable/macros.h>
-#include <torch/csrc/stable/accelerator.h>
-#include <torch/csrc/stable/tensor.h>
-#include <torch/csrc/stable/ops.h>
-#include <torch/headeronly/core/ScalarType.h>
-#include <torch/csrc/stable/device.h>
+#include <ATen/cuda/Exceptions.h>
+#include <c10/cuda/CUDAGuard.h>
+#include <c10/cuda/CUDAStream.h>
+#include <torch/all.h>
 
 #include "custom_all_reduce.cuh"
 
@@ -15,7 +11,7 @@ using fptr_t = int64_t;
 static_assert(sizeof(void*) == sizeof(fptr_t));
 
 fptr_t init_custom_ar(const std::vector<fptr_t>& fake_ipc_ptrs,
-                      torch::stable::Tensor& rank_data, int64_t rank,
+                      torch::Tensor& rank_data, int64_t rank,
                       bool fully_connected) {
   int world_size = fake_ipc_ptrs.size();
   if (world_size > 8)
@@ -29,9 +25,9 @@ fptr_t init_custom_ar(const std::vector<fptr_t>& fake_ipc_ptrs,
   for (int i = 0; i < world_size; i++) {
     ipc_ptrs[i] = reinterpret_cast<vllm::Signal*>(fake_ipc_ptrs[i]);
   }
-  return (fptr_t) new vllm::CustomAllreduce(
-      ipc_ptrs, rank_data.mutable_data_ptr(), rank_data.numel(), rank,
-      world_size, fully_connected);
+  return (fptr_t) new vllm::CustomAllreduce(ipc_ptrs, rank_data.data_ptr(),
+                                            rank_data.numel(), rank, world_size,
+                                            fully_connected);
 }
 
 /**
@@ -50,14 +46,10 @@ fptr_t init_custom_ar(const std::vector<fptr_t>& fake_ipc_ptrs,
  * 5. A[None].expand(2, -1, -1, -1): Not OK
  * 6. A[:, 1:, 1:]: Not OK
  */
-bool _is_weak_contiguous(torch::stable::Tensor& t) {
-  if (t.is_contiguous()) {
-    return true;
-  }
-  int64_t storage_nbytes = 0;
-  TORCH_ERROR_CODE_CHECK(aoti_torch_get_storage_size(t.get(), &storage_nbytes));
-  return storage_nbytes - t.storage_offset() * t.element_size() ==
-         static_cast<int64_t>(t.numel() * t.element_size());
+bool _is_weak_contiguous(torch::Tensor& t) {
+  return t.is_contiguous() ||
+         (t.storage().nbytes() - t.storage_offset() * t.element_size() ==
+          t.numel() * t.element_size());
 }
 
 /**
@@ -67,45 +59,42 @@ bool _is_weak_contiguous(torch::stable::Tensor& t) {
  * Otherwise, _reg_buffer is assumed to be IPC-registered and inp is first
  * copied into _reg_buffer.
  */
-void all_reduce(fptr_t _fa, torch::stable::Tensor& inp,
-                torch::stable::Tensor& out, fptr_t _reg_buffer,
-                int64_t reg_buffer_sz_bytes) {
+void all_reduce(fptr_t _fa, torch::Tensor& inp, torch::Tensor& out,
+                fptr_t _reg_buffer, int64_t reg_buffer_sz_bytes) {
   auto fa = reinterpret_cast<vllm::CustomAllreduce*>(_fa);
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      inp.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream(inp.get_device_index());
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(inp));
+  auto stream = c10::cuda::getCurrentCUDAStream().stream();
 
-  STD_TORCH_CHECK((inp.scalar_type()) == (out.scalar_type()));
-  STD_TORCH_CHECK((inp.numel()) == (out.numel()));
-  STD_TORCH_CHECK(_is_weak_contiguous(out));
-  STD_TORCH_CHECK(_is_weak_contiguous(inp));
+  TORCH_CHECK_EQ(inp.scalar_type(), out.scalar_type());
+  TORCH_CHECK_EQ(inp.numel(), out.numel());
+  TORCH_CHECK(_is_weak_contiguous(out));
+  TORCH_CHECK(_is_weak_contiguous(inp));
   auto input_size = inp.numel() * inp.element_size();
   auto reg_buffer = reinterpret_cast<void*>(_reg_buffer);
   if (reg_buffer) {
-    STD_TORCH_CHECK((input_size) <= (reg_buffer_sz_bytes));
-    STD_CUDA_CHECK(cudaMemcpyAsync(reg_buffer, inp.const_data_ptr(), input_size,
-                                   cudaMemcpyDeviceToDevice, stream));
+    TORCH_CHECK_LE(input_size, reg_buffer_sz_bytes);
+    AT_CUDA_CHECK(cudaMemcpyAsync(reg_buffer, inp.data_ptr(), input_size,
+                                  cudaMemcpyDeviceToDevice, stream));
   } else {
-    reg_buffer = inp.mutable_data_ptr();
+    reg_buffer = inp.data_ptr();
   }
   switch (out.scalar_type()) {
-    case torch::headeronly::ScalarType::Float: {
+    case at::ScalarType::Float: {
       fa->allreduce<float>(stream, reinterpret_cast<float*>(reg_buffer),
-                           reinterpret_cast<float*>(out.mutable_data_ptr()),
+                           reinterpret_cast<float*>(out.data_ptr()),
                            out.numel());
       break;
     }
-    case torch::headeronly::ScalarType::Half: {
+    case at::ScalarType::Half: {
       fa->allreduce<half>(stream, reinterpret_cast<half*>(reg_buffer),
-                          reinterpret_cast<half*>(out.mutable_data_ptr()),
-                          out.numel());
+                          reinterpret_cast<half*>(out.data_ptr()), out.numel());
       break;
     }
 #if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
-    case torch::headeronly::ScalarType::BFloat16: {
+    case at::ScalarType::BFloat16: {
       fa->allreduce<nv_bfloat16>(
           stream, reinterpret_cast<nv_bfloat16*>(reg_buffer),
-          reinterpret_cast<nv_bfloat16*>(out.mutable_data_ptr()), out.numel());
+          reinterpret_cast<nv_bfloat16*>(out.data_ptr()), out.numel());
       break;
     }
 #endif
@@ -123,7 +112,7 @@ int64_t meta_size() { return sizeof(vllm::Signal); }
 
 void register_buffer(fptr_t _fa, const std::vector<fptr_t>& fake_ipc_ptrs) {
   auto fa = reinterpret_cast<vllm::CustomAllreduce*>(_fa);
-  STD_TORCH_CHECK(fake_ipc_ptrs.size() == fa->world_size_);
+  TORCH_CHECK(fake_ipc_ptrs.size() == fa->world_size_);
   void* ipc_ptrs[8];
   for (int i = 0; i < fake_ipc_ptrs.size(); i++) {
     ipc_ptrs[i] = reinterpret_cast<void*>(fake_ipc_ptrs[i]);
@@ -154,49 +143,47 @@ void register_graph_buffers(fptr_t _fa,
   fa->register_graph_buffers(bytes, offsets);
 }
 
-std::tuple<fptr_t, torch::stable::Tensor> allocate_shared_buffer_and_handle(
+std::tuple<fptr_t, torch::Tensor> allocate_shared_buffer_and_handle(
     int64_t size) {
-  int device_index;
-  STD_CUDA_CHECK(cudaGetDevice(&device_index));
-  const torch::stable::accelerator::DeviceGuard device_guard(device_index);
+  auto device_index = c10::cuda::current_device();
+  at::DeviceGuard device_guard(at::Device(at::DeviceType::CUDA, device_index));
   void* buffer;
   cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
-  const cudaStream_t stream = get_current_cuda_stream(device_index);
-  STD_CUDA_CHECK(cudaThreadExchangeStreamCaptureMode(&mode));
+  auto stream = c10::cuda::getCurrentCUDAStream().stream();
+  AT_CUDA_CHECK(cudaThreadExchangeStreamCaptureMode(&mode));
 
   // Allocate buffer
 #if defined(USE_ROCM)
   // data buffers need to be "uncached" for signal on MI200
-  STD_CUDA_CHECK(
+  AT_CUDA_CHECK(
       hipExtMallocWithFlags((void**)&buffer, size, hipDeviceMallocUncached));
 #else
-  STD_CUDA_CHECK(cudaMalloc((void**)&buffer, size));
+  AT_CUDA_CHECK(cudaMalloc((void**)&buffer, size));
 #endif
-  STD_CUDA_CHECK(cudaMemsetAsync(buffer, 0, size, stream));
-  STD_CUDA_CHECK(cudaStreamSynchronize(stream));
-  STD_CUDA_CHECK(cudaThreadExchangeStreamCaptureMode(&mode));
+  AT_CUDA_CHECK(cudaMemsetAsync(buffer, 0, size, stream));
+  AT_CUDA_CHECK(cudaStreamSynchronize(stream));
+  AT_CUDA_CHECK(cudaThreadExchangeStreamCaptureMode(&mode));
 
   // Create IPC memhandle for the allocated buffer.
   // Will use it in open_mem_handle.
-  auto handle = torch::stable::empty(
-      {static_cast<int64_t>(sizeof(cudaIpcMemHandle_t))},
-      torch::headeronly::ScalarType::Byte, std::nullopt,
-      torch::stable::Device(torch::stable::DeviceType::CPU));
-  STD_CUDA_CHECK(cudaIpcGetMemHandle(
-      (cudaIpcMemHandle_t*)handle.mutable_data_ptr(), buffer));
+  auto options =
+      torch::TensorOptions().dtype(torch::kUInt8).device(torch::kCPU);
+  auto handle =
+      torch::empty({static_cast<int64_t>(sizeof(cudaIpcMemHandle_t))}, options);
+  AT_CUDA_CHECK(
+      cudaIpcGetMemHandle((cudaIpcMemHandle_t*)handle.data_ptr(), buffer));
 
   return std::make_tuple(reinterpret_cast<fptr_t>(buffer), handle);
 }
 
-fptr_t open_mem_handle(torch::stable::Tensor& mem_handle) {
+fptr_t open_mem_handle(torch::Tensor& mem_handle) {
   void* ipc_ptr;
-  STD_CUDA_CHECK(cudaIpcOpenMemHandle(
-      (void**)&ipc_ptr,
-      *((const cudaIpcMemHandle_t*)mem_handle.const_data_ptr()),
+  AT_CUDA_CHECK(cudaIpcOpenMemHandle(
+      (void**)&ipc_ptr, *((const cudaIpcMemHandle_t*)mem_handle.data_ptr()),
       cudaIpcMemLazyEnablePeerAccess));
   return reinterpret_cast<fptr_t>(ipc_ptr);
 }
 
 void free_shared_buffer(fptr_t buffer) {
-  STD_CUDA_CHECK(cudaFree(reinterpret_cast<void*>(buffer)));
+  AT_CUDA_CHECK(cudaFree(reinterpret_cast<void*>(buffer)));
 }

csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu

@@ -28,20 +28,7 @@
  *   [bs*576,       bs*576 + bs*8):   UE8M0 scales, 7 real + 1 pad per token
  */
 
-#include "torch_utils.h"
-
-#include <torch/csrc/stable/macros.h>
-#include <torch/csrc/stable/accelerator.h>
-#include <torch/csrc/stable/tensor.h>
-#include <torch/csrc/stable/ops.h>
-#include <torch/headeronly/core/ScalarType.h>
-#include <torch/csrc/stable/device.h>
-
 #include <cmath>
-#include "cuda_compat.h"
-#include "dispatch_utils.h"
-#include "type_convert.cuh"
-
 #ifndef USE_ROCM
   #include <cuda_fp8.h>
 #else
@@ -50,6 +37,14 @@
 #include <cuda_runtime.h>
 #include <type_traits>
 
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
+#include <torch/cuda.h>
+
+#include "cuda_compat.h"
+#include "dispatch_utils.h"
+#include "type_convert.cuh"
+
 #ifndef FINAL_MASK
   #ifdef USE_ROCM
     #define FINAL_MASK 0xffffffffffffffffULL
@@ -75,7 +70,7 @@ namespace deepseek_v4_fused_ops {
 
 namespace {
 inline int getSMVersion() {
-  auto* props = get_device_prop();
+  auto* props = at::cuda::getCurrentDeviceProperties();
   return props->major * 10 + props->minor;
 }
 }  // namespace
@@ -127,45 +122,18 @@ __device__ __forceinline__ float warpSum(float val) {
 // Per-slot inner pipeline
 // ────────────────────────────────────────────────────────────────────────────
 // Shared by both kernel variants: 1 CTA per (token, head) pair vs. 1 CTA per
-// token.  Templated on `kNumHeadsQPadded` so the KV-sentinel comparison and
-// q_out stride fold to compile-time constants.
-//
-// Slot layout (per token):
-//   slot < num_heads_q                          → live-Q   (RMSNorm + RoPE,
-//                                                           read q_in →
-//                                                           write q_out)
-//   num_heads_q <= slot < kNumHeadsQPadded      → pad-Q    (zero-fill q_out;
-//                                                           v0/v1 unused)
-//   slot == kNumHeadsQPadded                    → KV       (RoPE + UE8M0 quant
-//                                                           + paged-cache
-//                                                           insert)
-template <typename scalar_t_in, int kNumHeadsQPadded>
+// token
+template <typename scalar_t_in>
 __device__ __forceinline__ void processDeepseekV4Slot(
     uint4 v0, uint4 v1, int const tokenIdx, int const slotIdx,
     int const dim_base, int const laneId, int const num_heads_q,
-    float const eps, scalar_t_in* __restrict__ q_out,
+    float const eps, scalar_t_in* __restrict__ q_inout,
     uint8_t* __restrict__ k_cache, int64_t const* __restrict__ slot_mapping,
     int64_t const* __restrict__ position_ids,
     float const* __restrict__ cos_sin_cache, int const cache_block_size,
     int const kv_block_stride) {
   using Converter = vllm::_typeConvert<scalar_t_in>;
-  bool const isKV = (slotIdx == kNumHeadsQPadded);
-  bool const isPadQ = !isKV && (slotIdx >= num_heads_q);
-
-  // ── Pad-Q branch: write 32 B of zeros and exit. ─────────────────────────
-  // FlashMLA reads these slots; bf16 +0.0 is bit pattern 0x0000, so a uint4
-  // zero literal is correct.  Matches the live-Q branch's vectorized store.
-  if (isPadQ) {
-    scalar_t_in* dst =
-        q_out +
-        (static_cast<int64_t>(tokenIdx) * kNumHeadsQPadded + slotIdx) *
-            kHeadDim +
-        dim_base;
-    uint4 const zero4 = {0u, 0u, 0u, 0u};
-    *reinterpret_cast<uint4*>(dst) = zero4;
-    *reinterpret_cast<uint4*>(dst + 8) = zero4;
-    return;
-  }
+  bool const isKV = (slotIdx == num_heads_q);
 
   // ── Decode the bf16 → 16 fp32 registers ─────────────────────────────
   float elements[kElemsPerLane];
@@ -239,7 +207,7 @@ __device__ __forceinline__ void processDeepseekV4Slot(
   // triggering and per-iteration buffer rotation.
   // ═══════════════════════════════════════════════════════════════════
   if (!isKV) {
-    // ── Live-Q: cast back to bf16 and store into the padded q_out. ─────
+    // ── Q: cast back to bf16 and store. ────────────────────────────
     uint4 out0, out1;
     typename Converter::packed_hip_type* po0 =
         reinterpret_cast<typename Converter::packed_hip_type*>(&out0);
@@ -256,9 +224,8 @@ __device__ __forceinline__ void processDeepseekV4Slot(
           make_float2(elements[8 + 2 * i], elements[8 + 2 * i + 1]));
     }
     scalar_t_in* dst =
-        q_out +
-        (static_cast<int64_t>(tokenIdx) * kNumHeadsQPadded + slotIdx) *
-            kHeadDim +
+        q_inout +
+        (static_cast<int64_t>(tokenIdx) * num_heads_q + slotIdx) * kHeadDim +
         dim_base;
     *reinterpret_cast<uint4*>(dst) = out0;
     *reinterpret_cast<uint4*>(dst + 8) = out1;
@@ -346,32 +313,20 @@ __device__ __forceinline__ void processDeepseekV4Slot(
 // Kernel
 // ────────────────────────────────────────────────────────────────────────────
 //
-// Grid: 1D, gridDim.x = ceil(num_tokens_full * (kNumHeadsQPadded + 1) /
+// Grid: 1D, gridDim.x = ceil(num_tokens_full * (num_heads_q + 1) /
 // warps_per_block) Block: blockDim.x = 256 threads (8 warps per block) Each
-// warp handles one (token, head_slot) pair.
-//   slot < num_heads_q                              → live-Q branch
-//                                                     (RMSNorm + RoPE,
-//                                                      read q_in → write q_out)
-//   num_heads_q <= slot < kNumHeadsQPadded          → pad-Q branch
-//                                                     (zero-fill q_out)
-//   slot == kNumHeadsQPadded                        → KV branch
-//                                                     (RoPE + UE8M0 quant +
-//                                                      paged-cache insert)
-//
-// `kNumHeadsQPadded` is a template parameter (compile-time constant) so the
-// divisions in the grid math and the KV-sentinel comparison fold to fast
-// constant operations.  The launch wrapper dispatches the runtime value to
-// the matching instantiation.
+// warp handles one (token, head_slot) pair. head_slot < num_heads_q          →
+// Q branch (RMSNorm + RoPE, in place) head_slot == num_heads_q         → KV
+// branch (RoPE + UE8M0 quant + insert)
 //
 // With DP padding, q/kv/position_ids can have more rows than slot_mapping.
-// The live-Q and pad-Q branches cover all `num_tokens_full` rows (downstream
-// attention uses them).  The KV branch only inserts the first
-// `num_tokens_insert` tokens (= slot_mapping length) into the paged cache.
+// The Q branch covers all `num_tokens_full` rows (downstream attention uses
+// them).  The KV branch only inserts the first `num_tokens_insert` tokens
+// (= slot_mapping length) into the paged cache.
 //
-template <typename scalar_t_in, int kNumHeadsQPadded>
+template <typename scalar_t_in>
 __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
-    scalar_t_in const* __restrict__ q_in,      // [N, num_heads_q,      512]
-    scalar_t_in* __restrict__ q_out,           // [N, kNumHeadsQPadded, 512]
+    scalar_t_in* __restrict__ q_inout,         // [N, H, 512] bf16, in place
     scalar_t_in const* __restrict__ kv_in,     // [N, 512] bf16
     uint8_t* __restrict__ k_cache,             // [num_blocks, block_stride]
     int64_t const* __restrict__ slot_mapping,  // [num_tokens_insert] i64
@@ -380,7 +335,7 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
     float const eps,
     int const num_tokens_full,    // = q.size(0) = kv.size(0)
     int const num_tokens_insert,  // = slot_mapping.size(0), ≤ num_tokens_full
-    int const num_heads_q,        // live Q heads (input layout)
+    int const num_heads_q,        // H
     int const cache_block_size,   // tokens per paged-cache block
     int const kv_block_stride) {  // bytes per paged-cache block
 #if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
@@ -396,13 +351,12 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
     int const laneId = threadIdx.x % 32;
     int const globalWarpIdx = blockIdx.x * warpsPerBlock + warpId;
 
-    constexpr int kTotalSlotsPerToken = kNumHeadsQPadded + 1;
-    int const tokenIdx = globalWarpIdx / kTotalSlotsPerToken;
-    int const slotIdx = globalWarpIdx % kTotalSlotsPerToken;
+    int const total_slots_per_token = num_heads_q + 1;
+    int const tokenIdx = globalWarpIdx / total_slots_per_token;
+    int const slotIdx = globalWarpIdx % total_slots_per_token;
     if (tokenIdx >= num_tokens_full) return;
 
-    bool const isKV = (slotIdx == kNumHeadsQPadded);
-    bool const isPadQ = !isKV && (slotIdx >= num_heads_q);
+    bool const isKV = (slotIdx == num_heads_q);
     // KV branch: skip DP-padded tokens (no slot reserved for them).
     if (isKV && tokenIdx >= num_tokens_insert) return;
 
@@ -417,26 +371,22 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
     // Dim range this lane owns within the 512-wide head.
     int const dim_base = laneId * kElemsPerLane;  // in [0, 512) step 16
 
-    // Load only for live-Q and KV slots; pad-Q skips the read (q_in beyond
-    // num_heads_q is out of bounds) and the helper zero-fills its output.
-    uint4 v0, v1;
-    if (!isPadQ) {
-      scalar_t_in const* src_ptr;
-      if (isKV) {
-        src_ptr = kv_in + static_cast<int64_t>(tokenIdx) * kHeadDim + dim_base;
-      } else {
-        int64_t const q_row_offset =
-            (static_cast<int64_t>(tokenIdx) * num_heads_q + slotIdx) *
-                kHeadDim +
-            dim_base;
-        src_ptr = q_in + q_row_offset;
-      }
-      v0 = *reinterpret_cast<uint4 const*>(src_ptr);
-      v1 = *reinterpret_cast<uint4 const*>(src_ptr + 8);
+    // Two 16-byte loads per thread (8 bf16 each).  Use uint4 as the vector
+    // type; the shared per-slot helper bitcasts to scalar_t_in packed pairs.
+    scalar_t_in const* src_ptr;
+    if (isKV) {
+      src_ptr = kv_in + static_cast<int64_t>(tokenIdx) * kHeadDim + dim_base;
+    } else {
+      int64_t const q_row_offset =
+          (static_cast<int64_t>(tokenIdx) * num_heads_q + slotIdx) * kHeadDim +
+          dim_base;
+      src_ptr = q_inout + q_row_offset;
     }
+    uint4 const v0 = *reinterpret_cast<uint4 const*>(src_ptr);
+    uint4 const v1 = *reinterpret_cast<uint4 const*>(src_ptr + 8);
 
-    processDeepseekV4Slot<scalar_t_in, kNumHeadsQPadded>(
-        v0, v1, tokenIdx, slotIdx, dim_base, laneId, num_heads_q, eps, q_out,
+    processDeepseekV4Slot<scalar_t_in>(
+        v0, v1, tokenIdx, slotIdx, dim_base, laneId, num_heads_q, eps, q_inout,
         k_cache, slot_mapping, position_ids, cos_sin_cache, cache_block_size,
         kv_block_stride);
 
@@ -458,9 +408,9 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
 // Q branch (RMSNorm + RoPE, in place) head_slot == num_heads_q
 // KV branch (RoPE + UE8M0 quant + insert)
 //
-template <typename scalar_t_in, int kNumHeadsQPadded>
+template <typename scalar_t_in>
 __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid(
-    scalar_t_in const* __restrict__ q_in, scalar_t_in* __restrict__ q_out,
+    scalar_t_in* __restrict__ q_inout,  // [N, H, 512] bf16, in place
     scalar_t_in const* __restrict__ kv_in, uint8_t* __restrict__ k_cache,
     int64_t const* __restrict__ slot_mapping,
     int64_t const* __restrict__ position_ids,
@@ -485,32 +435,25 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid(
 #endif
 
     int const dim_base = laneId * kElemsPerLane;  // in [0, 512) step 16
-    // Slot enumeration: live-Q + pad-Q + (KV if this token has a slot).
-    int const slot_end = (tokenIdx >= num_tokens_insert)
-                             ? kNumHeadsQPadded
-                             : (kNumHeadsQPadded + 1);
+    int const slot_end =
+        (tokenIdx >= num_tokens_insert) ? num_heads_q : (num_heads_q + 1);
 
-    auto load_slot = [&](int s, uint4& va, uint4& vb) {
-      // pad-Q slots skip the load — q_in beyond num_heads_q is OOB.
-      if (s >= num_heads_q && s < kNumHeadsQPadded) return;
-      scalar_t_in const* src;
-      if (s == kNumHeadsQPadded) {
-        src = kv_in + static_cast<int64_t>(tokenIdx) * kHeadDim + dim_base;
-      } else {
-        src = q_in +
-              (static_cast<int64_t>(tokenIdx) * num_heads_q +
-               static_cast<int64_t>(s)) *
-                  kHeadDim +
-              dim_base;
+    auto src_for_slot = [&](int s) -> scalar_t_in const* {
+      if (s == num_heads_q) {
+        return kv_in + static_cast<int64_t>(tokenIdx) * kHeadDim + dim_base;
       }
-      va = *reinterpret_cast<uint4 const*>(src);
-      vb = *reinterpret_cast<uint4 const*>(src + 8);
+      return q_inout +
+             (static_cast<int64_t>(tokenIdx) * num_heads_q +
+              static_cast<int64_t>(s)) *
+                 kHeadDim +
+             dim_base;
     };
 
     if (warpId < slot_end) {
       int curr_slot = warpId;
-      uint4 v0_curr, v1_curr;
-      load_slot(curr_slot, v0_curr, v1_curr);
+      scalar_t_in const* src_curr = src_for_slot(curr_slot);
+      uint4 v0_curr = *reinterpret_cast<uint4 const*>(src_curr);
+      uint4 v1_curr = *reinterpret_cast<uint4 const*>(src_curr + 8);
 
       while (curr_slot < slot_end) {
         int const next_slot = curr_slot + warpsPerBlock;
@@ -519,12 +462,14 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid(
         // Prefetch src for the next slot
         uint4 v0_next, v1_next;
         if (has_next) {
-          load_slot(next_slot, v0_next, v1_next);
+          scalar_t_in const* src_next = src_for_slot(next_slot);
+          v0_next = *reinterpret_cast<uint4 const*>(src_next);
+          v1_next = *reinterpret_cast<uint4 const*>(src_next + 8);
         }
 
-        processDeepseekV4Slot<scalar_t_in, kNumHeadsQPadded>(
+        processDeepseekV4Slot<scalar_t_in>(
             v0_curr, v1_curr, tokenIdx, curr_slot, dim_base, laneId,
-            num_heads_q, eps, q_out, k_cache, slot_mapping, position_ids,
+            num_heads_q, eps, q_inout, k_cache, slot_mapping, position_ids,
             cos_sin_cache, cache_block_size, kv_block_stride);
 
         // ── Buffer rotation: hand the prefetched LDGs to the next iter.
@@ -545,10 +490,10 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid(
 // ────────────────────────────────────────────────────────────────────────────
 // Launch wrapper
 // ────────────────────────────────────────────────────────────────────────────
-template <typename scalar_t_in, int kNumHeadsQPadded>
-static void launchFusedDeepseekV4Templated(
-    scalar_t_in const* q_in, scalar_t_in* q_out, scalar_t_in const* kv_in,
-    uint8_t* k_cache, int64_t const* slot_mapping, int64_t const* position_ids,
+template <typename scalar_t_in>
+void launchFusedDeepseekV4QNormRopeKVRopeQuantInsert(
+    scalar_t_in* q_inout, scalar_t_in const* kv_in, uint8_t* k_cache,
+    int64_t const* slot_mapping, int64_t const* position_ids,
     float const* cos_sin_cache, float const eps, int const num_tokens_full,
     int const num_tokens_insert, int const num_heads_q,
     int const cache_block_size, int const kv_block_stride,
@@ -556,7 +501,7 @@ static void launchFusedDeepseekV4Templated(
   constexpr int kBlockSize = 256;
   constexpr int kWarpsPerBlock = kBlockSize / 32;
   int64_t const total_warps =
-      static_cast<int64_t>(num_tokens_full) * (kNumHeadsQPadded + 1);
+      static_cast<int64_t>(num_tokens_full) * (num_heads_q + 1);
   int const grid =
       static_cast<int>((total_warps + kWarpsPerBlock - 1) / kWarpsPerBlock);
 
@@ -569,7 +514,7 @@ static void launchFusedDeepseekV4Templated(
   // bf16 on pre-Ampere (sm_70/sm_75) because _typeConvert<BFloat16> is
   // unavailable there.  Refuse the launch loudly instead of silently
   // skipping the work.
-  STD_TORCH_CHECK(
+  TORCH_CHECK(
       sm_version >= 80,
       "fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert requires sm_80+ "
       "(Ampere or newer); got sm_",
@@ -587,151 +532,91 @@ static void launchFusedDeepseekV4Templated(
 
   if (num_tokens_full < NUM_TOKEN_CUTOFF) {
     cudaLaunchKernelEx(
-        &config,
-        fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel<scalar_t_in,
-                                                        kNumHeadsQPadded>,
-        q_in, q_out, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache,
-        eps, num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
+        &config, fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel<scalar_t_in>,
+        q_inout, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache, eps,
+        num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
         kv_block_stride);
   } else {
     config.gridDim = dim3(num_tokens_full);
     cudaLaunchKernelEx(
         &config,
-        fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid<
-            scalar_t_in, kNumHeadsQPadded>,
-        q_in, q_out, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache,
-        eps, num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
+        fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid<scalar_t_in>,
+        q_inout, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache, eps,
+        num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
         kv_block_stride);
   }
+
 #else
   // ROCm: use standard kernel launch syntax (no PDL/stream serialization)
   // clang-format off
-  fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel<scalar_t_in, kNumHeadsQPadded>
+  fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel<scalar_t_in>
       <<<grid, kBlockSize, 0, stream>>>(
-          q_in, q_out, kv_in, k_cache, slot_mapping, position_ids,
-          cos_sin_cache, eps, num_tokens_full, num_tokens_insert, num_heads_q,
+          q_inout, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache,
+          eps, num_tokens_full, num_tokens_insert, num_heads_q,
           cache_block_size, kv_block_stride);
 #endif
 }
 
-// Runtime dispatch into one of the precompiled `kNumHeadsQPadded`
-// instantiations.  Supported padded head counts: 8, 16, 32, 64, 128.
-template <typename scalar_t_in>
-void launchFusedDeepseekV4QNormRopeKVRopeQuantInsert(
-    scalar_t_in const* q_in, scalar_t_in* q_out, scalar_t_in const* kv_in,
-    uint8_t* k_cache, int64_t const* slot_mapping,
-    int64_t const* position_ids, float const* cos_sin_cache, float const eps,
-    int const num_tokens_full, int const num_tokens_insert,
-    int const num_heads_q, int const num_heads_q_padded,
-    int const cache_block_size, int const kv_block_stride,
-    cudaStream_t stream) {
-#define DISPATCH(N)                                                         \
-  case N:                                                                   \
-    launchFusedDeepseekV4Templated<scalar_t_in, N>(                         \
-        q_in, q_out, kv_in, k_cache, slot_mapping, position_ids,            \
-        cos_sin_cache, eps, num_tokens_full, num_tokens_insert, num_heads_q, \
-        cache_block_size, kv_block_stride, stream);                         \
-    return;
-
-  switch (num_heads_q_padded) {
-    DISPATCH(8)
-    DISPATCH(16)
-    DISPATCH(32)
-    DISPATCH(64)
-    DISPATCH(128)
-    default:
-      STD_TORCH_CHECK(false,
-                  "fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert: "
-                  "unsupported num_heads_q_padded=",
-                  num_heads_q_padded,
-                  " (compiled instantiations: 8, 16, 32, 64, 128).");
-  }
-#undef DISPATCH
-}
-
 }  // namespace deepseek_v4_fused_ops
 }  // namespace vllm
 
 // ────────────────────────────────────────────────────────────────────────────
 // Torch op wrapper
 // ────────────────────────────────────────────────────────────────────────────
-torch::stable::Tensor fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert(
-    torch::stable::Tensor const& q_in,           // [N, num_heads_q, 512] bf16
-    torch::stable::Tensor const& kv,             // [N, 512] bf16 (read-only)
-    torch::stable::Tensor& k_cache,              // [num_blocks, block_bytes] uint8
-    torch::stable::Tensor const& slot_mapping,   // [N] int64
-    torch::stable::Tensor const& position_ids,   // [N] int64
-    torch::stable::Tensor const& cos_sin_cache,  // [max_pos, rope_dim] bf16
-    int64_t q_head_padded,                       // padded Q head count for output
+void fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert(
+    torch::Tensor& q,                    // [N, H, 512] bf16, in place
+    torch::Tensor const& kv,             // [N, 512] bf16 (read-only)
+    torch::Tensor& k_cache,              // [num_blocks, block_bytes] uint8
+    torch::Tensor const& slot_mapping,   // [N] int64
+    torch::Tensor const& position_ids,   // [N] int64
+    torch::Tensor const& cos_sin_cache,  // [max_pos, rope_dim] bf16
     double eps, int64_t cache_block_size) {
-  STD_TORCH_CHECK(q_in.device().is_cuda() && q_in.is_contiguous(),
-                  "q_in must be contiguous CUDA");
-  STD_TORCH_CHECK(kv.device().is_cuda() && kv.is_contiguous(),
-                  "kv must be contiguous CUDA");
-  STD_TORCH_CHECK(k_cache.device().is_cuda(), "k_cache must be CUDA");
-  STD_TORCH_CHECK(slot_mapping.device().is_cuda() &&
-                      slot_mapping.scalar_type() ==
-                          torch::headeronly::ScalarType::Long,
-                  "slot_mapping must be int64 CUDA");
-  STD_TORCH_CHECK(position_ids.device().is_cuda() &&
-                      position_ids.scalar_type() ==
-                          torch::headeronly::ScalarType::Long,
-                  "position_ids must be int64 CUDA");
-  STD_TORCH_CHECK(cos_sin_cache.device().is_cuda(), "cos_sin_cache must be CUDA");
-  STD_TORCH_CHECK(q_in.dim() == 3 && q_in.size(2) == 512,
-                  "q_in shape [N, num_heads_q, 512]");
-  STD_TORCH_CHECK(kv.dim() == 2 && kv.size(1) == 512, "kv shape [N, 512]");
-  STD_TORCH_CHECK(q_in.scalar_type() == kv.scalar_type(),
-                  "q_in and kv dtype must match");
-  STD_TORCH_CHECK(q_head_padded >= q_in.size(1),
-                  "q_head_padded must be >= q_in.size(1) (num_heads_q)");
-  STD_TORCH_CHECK(k_cache.scalar_type() == torch::headeronly::ScalarType::Byte,
-                  "k_cache must be uint8");
-  STD_TORCH_CHECK(cos_sin_cache.dim() == 2 && cos_sin_cache.size(1) == 64,
-                  "cos_sin_cache shape [max_pos, 64]");
-  STD_TORCH_CHECK(cos_sin_cache.scalar_type() ==
-                      torch::headeronly::ScalarType::Float,
-                  "cos_sin_cache must be float32");
+  TORCH_CHECK(q.is_cuda() && q.is_contiguous(), "q must be contiguous CUDA");
+  TORCH_CHECK(kv.is_cuda() && kv.is_contiguous(), "kv must be contiguous CUDA");
+  TORCH_CHECK(k_cache.is_cuda(), "k_cache must be CUDA");
+  TORCH_CHECK(slot_mapping.is_cuda() && slot_mapping.dtype() == torch::kInt64,
+              "slot_mapping must be int64 CUDA");
+  TORCH_CHECK(position_ids.is_cuda() && position_ids.dtype() == torch::kInt64,
+              "position_ids must be int64 CUDA");
+  TORCH_CHECK(cos_sin_cache.is_cuda(), "cos_sin_cache must be CUDA");
+  TORCH_CHECK(q.dim() == 3 && q.size(2) == 512, "q shape [N, H, 512]");
+  TORCH_CHECK(kv.dim() == 2 && kv.size(1) == 512, "kv shape [N, 512]");
+  TORCH_CHECK(q.dtype() == kv.dtype(), "q and kv dtype must match");
+  TORCH_CHECK(k_cache.dtype() == torch::kUInt8, "k_cache must be uint8");
+  TORCH_CHECK(cos_sin_cache.dim() == 2 && cos_sin_cache.size(1) == 64,
+              "cos_sin_cache shape [max_pos, 64]");
+  TORCH_CHECK(cos_sin_cache.dtype() == torch::kFloat32,
+              "cos_sin_cache must be float32");
 
   // With DP padding, slot_mapping can be shorter than q/kv/positions.
   // Q-norm+RoPE runs on all q.size(0) rows (downstream attention uses them);
   // KV quant+insert runs only on the first slot_mapping.size(0) rows.
-  int const num_tokens_full = static_cast<int>(q_in.size(0));
+  int const num_tokens_full = static_cast<int>(q.size(0));
   int const num_tokens_insert = static_cast<int>(slot_mapping.size(0));
-  STD_TORCH_CHECK(static_cast<int>(kv.size(0)) == num_tokens_full &&
-                      static_cast<int>(position_ids.size(0)) == num_tokens_full,
-                  "q/kv/position_ids row counts must match");
-  STD_TORCH_CHECK(num_tokens_insert <= num_tokens_full,
-                  "slot_mapping must not exceed q row count");
-  int const num_heads_q = static_cast<int>(q_in.size(1));
-  int const num_heads_q_padded = static_cast<int>(q_head_padded);
+  TORCH_CHECK(static_cast<int>(kv.size(0)) == num_tokens_full &&
+                  static_cast<int>(position_ids.size(0)) == num_tokens_full,
+              "q/kv/position_ids row counts must match");
+  TORCH_CHECK(num_tokens_insert <= num_tokens_full,
+              "slot_mapping must not exceed q row count");
+  int const num_heads_q = static_cast<int>(q.size(1));
   int const cache_block_size_i = static_cast<int>(cache_block_size);
   int const kv_block_stride = static_cast<int>(k_cache.stride(0));
 
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      q_in.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream(q_in.get_device_index());
+  at::cuda::OptionalCUDAGuard device_guard(device_of(q));
+  auto stream = at::cuda::getCurrentCUDAStream();
 
-  // Allocate the padded q output.  The kernel writes every element (live
-  // region gets RMSNorm+RoPE; pad region gets zeros), so `empty` is safe.
-  auto q_out = torch::stable::new_empty(
-      q_in, {q_in.size(0), q_head_padded, q_in.size(2)}, q_in.scalar_type());
-
-  VLLM_STABLE_DISPATCH_HALF_TYPES(
-      q_in.scalar_type(), "fused_deepseek_v4_qnorm_rope_kv_insert", [&] {
+  VLLM_DISPATCH_HALF_TYPES(
+      q.scalar_type(), "fused_deepseek_v4_qnorm_rope_kv_insert", [&] {
         using qkv_scalar_t = scalar_t;
         vllm::deepseek_v4_fused_ops::
             launchFusedDeepseekV4QNormRopeKVRopeQuantInsert<qkv_scalar_t>(
-                reinterpret_cast<qkv_scalar_t const*>(q_in.const_data_ptr()),
-                reinterpret_cast<qkv_scalar_t*>(q_out.mutable_data_ptr()),
-                reinterpret_cast<qkv_scalar_t const*>(kv.const_data_ptr()),
-                reinterpret_cast<uint8_t*>(k_cache.mutable_data_ptr()),
-                slot_mapping.const_data_ptr<int64_t>(),
-                position_ids.const_data_ptr<int64_t>(),
-                cos_sin_cache.const_data_ptr<float>(), static_cast<float>(eps),
+                reinterpret_cast<qkv_scalar_t*>(q.data_ptr()),
+                reinterpret_cast<qkv_scalar_t const*>(kv.data_ptr()),
+                reinterpret_cast<uint8_t*>(k_cache.data_ptr()),
+                reinterpret_cast<int64_t const*>(slot_mapping.data_ptr()),
+                reinterpret_cast<int64_t const*>(position_ids.data_ptr()),
+                cos_sin_cache.data_ptr<float>(), static_cast<float>(eps),
                 num_tokens_full, num_tokens_insert, num_heads_q,
-                num_heads_q_padded, cache_block_size_i, kv_block_stride,
-                stream);
+                cache_block_size_i, kv_block_stride, stream);
       });
-  return q_out;
-}
+}

csrc/libtorch_stable/activation_kernels.cu

@@ -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"
 

csrc/libtorch_stable/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp

@@ -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,

csrc/libtorch_stable/fp32_router_gemm.cu

@@ -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

csrc/libtorch_stable/fp32_router_gemm_entry.cu

@@ -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));
-}

csrc/libtorch_stable/fused_qknorm_rope_kernel.cu

@@ -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"

csrc/libtorch_stable/layernorm_kernels.cu

@@ -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);
   }
 }
 

csrc/libtorch_stable/layernorm_quant_kernels.cu

@@ -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);
   }

csrc/libtorch_stable/moe/mxfp8_moe/cutlass_mxfp8_grouped_mm.cu

@@ -1,69 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-// Adapted from SGLang:
-// https://github.com/sgl-project/sglang/blob/ded068a76e00878881d52d5bfb791e0f60d7311b/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled.cu
-
-#include <torch/csrc/stable/library.h>
-#include <torch/csrc/stable/tensor.h>
-#include "libtorch_stable/torch_utils.h"
-
-#include "cutlass_mxfp8_grouped_mm_launcher.cuh"
-
-void cutlass_mxfp8_grouped_mm(const torch::stable::Tensor& a,
-                              const torch::stable::Tensor& b,
-                              const torch::stable::Tensor& sfa,
-                              const torch::stable::Tensor& sfb,
-                              torch::stable::Tensor& d,
-                              const torch::stable::Tensor& problem_sizes,
-                              const torch::stable::Tensor& expert_offsets,
-                              const torch::stable::Tensor& blockscale_offsets) {
-#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
-  STD_TORCH_CHECK(problem_sizes.dim() == 2, "problem_sizes must be 2D tensor");
-  STD_TORCH_CHECK(problem_sizes.size(1) == 3,
-                  "problem_sizes must have shape (num_experts, 3)");
-  STD_TORCH_CHECK(
-      problem_sizes.size(0) == expert_offsets.size(0),
-      "Number of experts in problem_sizes must match expert_offsets");
-  STD_TORCH_CHECK(
-      problem_sizes.scalar_type() == torch::headeronly::ScalarType::Int,
-      "problem_sizes must be int32");
-  STD_TORCH_CHECK(
-      expert_offsets.scalar_type() == torch::headeronly::ScalarType::Int,
-      "expert_offsets must be int32");
-  STD_TORCH_CHECK(
-      blockscale_offsets.scalar_type() == torch::headeronly::ScalarType::Int,
-      "blockscale_offsets must be int32");
-  STD_TORCH_CHECK(a.dim() == 2,
-                  "a must be a 2D tensor of shape (num_tokens, k)");
-  STD_TORCH_CHECK(b.dim() == 3,
-                  "b must be a 3D tensor of shape (num_experts, k, n)");
-  STD_TORCH_CHECK(a.size(1) == b.size(1) && a.size(1) % 128 == 0,
-                  "k should align 128");
-  STD_TORCH_CHECK(b.size(2) % 128 == 0, "n should align 128");
-  STD_TORCH_CHECK(a.stride(1) == 1, "a must be row major");
-  STD_TORCH_CHECK(b.stride(1) == 1, "b must be column major");
-
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      a.get_device_index());
-  auto stream = get_current_cuda_stream(a.get_device_index());
-  if (d.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
-    expert_specialization::cutlass_mxfp8_grouped_mm_dispatch_out_dtype<
-        cutlass::bfloat16_t>(a, b, sfa, sfb, d, problem_sizes, expert_offsets,
-                             blockscale_offsets, stream);
-  } else if (d.scalar_type() == torch::headeronly::ScalarType::Half) {
-    expert_specialization::cutlass_mxfp8_grouped_mm_dispatch_out_dtype<
-        cutlass::half_t>(a, b, sfa, sfb, d, problem_sizes, expert_offsets,
-                         blockscale_offsets, stream);
-  } else {
-    STD_TORCH_CHECK(false, "dtype must be kFloat16 or kBFloat16");
-  }
-#else
-  STD_TORCH_CHECK(false,
-                  "No implemented cutlass_mxfp8_grouped_mm for "
-                  "current device");
-#endif
-}
-
-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("cutlass_mxfp8_grouped_mm", TORCH_BOX(&cutlass_mxfp8_grouped_mm));
-}

csrc/libtorch_stable/moe/mxfp8_moe/mxfp8_experts_quant.cu

@@ -1,66 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-// Adapted from SGLang:
-// https://github.com/sgl-project/sglang/blob/ded068a76e00878881d52d5bfb791e0f60d7311b/sgl-kernel/csrc/expert_specialization/es_sm100_mxfp8_blockscaled_group_quant.cu
-
-#include <torch/csrc/stable/library.h>
-#include <torch/csrc/stable/tensor.h>
-#include "libtorch_stable/torch_utils.h"
-
-#include "mxfp8_experts_quant.cuh"
-
-void mxfp8_experts_quant(const torch::stable::Tensor& input,
-                         const torch::stable::Tensor& problem_sizes,
-                         const torch::stable::Tensor& expert_offsets,
-                         const torch::stable::Tensor& blockscale_offsets,
-                         torch::stable::Tensor& quant_output,
-                         torch::stable::Tensor& scale_factor) {
-#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
-  STD_TORCH_CHECK(input.dim() == 2, "input must be 2D tensor");
-  STD_TORCH_CHECK(input.size(1) % 128 == 0, "k must align to 128");
-  STD_TORCH_CHECK(input.stride(1) == 1, "input must be row major");
-  STD_TORCH_CHECK(problem_sizes.dim() == 2, "problem_sizes must be 2D tensor");
-  STD_TORCH_CHECK(
-      problem_sizes.scalar_type() == torch::headeronly::ScalarType::Int,
-      "problem_sizes must be int32");
-  STD_TORCH_CHECK(
-      expert_offsets.scalar_type() == torch::headeronly::ScalarType::Int,
-      "expert_offsets must be int32");
-  STD_TORCH_CHECK(
-      blockscale_offsets.scalar_type() == torch::headeronly::ScalarType::Int,
-      "blockscale_offsets must be int32");
-
-  auto groups = problem_sizes.size(0);
-  STD_TORCH_CHECK(
-      expert_offsets.dim() == 1 && expert_offsets.size(0) == groups,
-      "expert_offsets must be 1D and have size equal to the number of groups");
-  STD_TORCH_CHECK(
-      blockscale_offsets.dim() == 1 && blockscale_offsets.size(0) == groups,
-      "blockscale_offsets must be 1D and have size equal to the number of "
-      "groups");
-
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      input.get_device_index());
-  if (input.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
-    expert_specialization::launch_mxfp8_experts_quant<__nv_bfloat16>(
-        input, problem_sizes, expert_offsets, blockscale_offsets, quant_output,
-        scale_factor);
-  } else if (input.scalar_type() == torch::headeronly::ScalarType::Half) {
-    expert_specialization::launch_mxfp8_experts_quant<__half>(
-        input, problem_sizes, expert_offsets, blockscale_offsets, quant_output,
-        scale_factor);
-  } else {
-    STD_TORCH_CHECK(false, "dtype must be kFloat16 or kBFloat16");
-  }
-#else
-  STD_TORCH_CHECK(false,
-                  "No implemented mxfp8_experts_quant for "
-                  "current device");
-#endif
-}
-
-// Registered here (not torch_bindings.cpp) because ENABLE_ES_MXFP8_GROUPED_MM
-// is applied only under COMPILE_LANGUAGE:CUDA.
-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("mxfp8_experts_quant", TORCH_BOX(&mxfp8_experts_quant));
-}

csrc/libtorch_stable/ops.h

@@ -3,6 +3,10 @@
 #include <torch/csrc/stable/library.h>
 #include <torch/csrc/stable/tensor.h>
 
+#ifndef USE_ROCM
+torch::stable::Tensor permute_cols(torch::stable::Tensor const& A,
+                                   torch::stable::Tensor const& perm);
+
 void per_token_group_quant_fp8(const torch::stable::Tensor& input,
                                torch::stable::Tensor& output_q,
                                torch::stable::Tensor& output_s,
@@ -24,10 +28,6 @@ void per_token_group_quant_int8(const torch::stable::Tensor& input,
                                 int64_t group_size, double eps, double int8_min,
                                 double int8_max);
 
-#ifndef USE_ROCM
-torch::stable::Tensor permute_cols(torch::stable::Tensor const& A,
-                                   torch::stable::Tensor const& perm);
-
 bool cutlass_scaled_mm_supports_fp8(int64_t cuda_device_capability);
 bool cutlass_scaled_mm_supports_block_fp8(int64_t cuda_device_capability);
 bool cutlass_group_gemm_supported(int64_t cuda_device_capability);
@@ -164,17 +164,6 @@ torch::stable::Tensor awq_dequantize(torch::stable::Tensor _kernel,
 
 #endif
 
-// Attention kernels (shared CUDA/ROCm)
-void merge_attn_states(
-    torch::stable::Tensor& output,
-    std::optional<torch::stable::Tensor> output_lse,
-    const torch::stable::Tensor& prefix_output,
-    const torch::stable::Tensor& prefix_lse,
-    const torch::stable::Tensor& suffix_output,
-    const torch::stable::Tensor& suffix_lse,
-    const std::optional<int64_t> prefill_tokens_with_context,
-    const std::optional<torch::stable::Tensor>& output_scale = std::nullopt);
-
 torch::stable::Tensor hadacore_transform(torch::stable::Tensor& x,
                                          bool inplace);
 
@@ -231,89 +220,6 @@ void fused_qk_norm_rope(torch::stable::Tensor& qkv, int64_t num_heads_q,
                         torch::stable::Tensor& position_ids,
                         int64_t forced_token_heads_per_warp);
 
-torch::stable::Tensor fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert(
-    torch::stable::Tensor const& q_in, torch::stable::Tensor const& kv,
-    torch::stable::Tensor& k_cache, torch::stable::Tensor const& slot_mapping,
-    torch::stable::Tensor const& position_ids,
-    torch::stable::Tensor const& cos_sin_cache, int64_t q_head_padded,
-    double eps, int64_t cache_block_size);
-
-#ifndef USE_ROCM
-torch::stable::Tensor minimax_allreduce_rms(
-    torch::stable::Tensor const& input,
-    torch::stable::Tensor const& norm_weight, torch::stable::Tensor workspace,
-    int64_t const rank, int64_t const nranks, double const eps);
-std::tuple<torch::stable::Tensor, torch::stable::Tensor>
-minimax_allreduce_rms_qk(torch::stable::Tensor qkv,
-                         torch::stable::Tensor const& norm_weight_q,
-                         torch::stable::Tensor const& norm_weight_k,
-                         torch::stable::Tensor workspace, int64_t const q_size,
-                         int64_t const kv_size, int64_t const rank,
-                         int64_t const nranks, double const eps);
-#endif
-
-// Sampler kernels (shared CUDA/ROCm)
-void apply_repetition_penalties_(
-    torch::stable::Tensor& logits, const torch::stable::Tensor& prompt_mask,
-    const torch::stable::Tensor& output_mask,
-    const torch::stable::Tensor& repetition_penalties);
-
-void top_k_per_row_prefill(const torch::stable::Tensor& logits,
-                           const torch::stable::Tensor& rowStarts,
-                           const torch::stable::Tensor& rowEnds,
-                           torch::stable::Tensor& indices, int64_t numRows,
-                           int64_t stride0, int64_t stride1, int64_t topK);
-
-void top_k_per_row_decode(const torch::stable::Tensor& logits, int64_t next_n,
-                          const torch::stable::Tensor& seqLens,
-                          torch::stable::Tensor& indices, int64_t numRows,
-                          int64_t stride0, int64_t stride1, int64_t topK);
-
-void persistent_topk(const torch::stable::Tensor& logits,
-                     const torch::stable::Tensor& lengths,
-                     torch::stable::Tensor& output,
-                     torch::stable::Tensor& workspace, int64_t k,
-                     int64_t max_seq_len);
-
-void selective_scan_fwd(
-    const torch::stable::Tensor& u, const torch::stable::Tensor& delta,
-    const torch::stable::Tensor& A, const torch::stable::Tensor& B,
-    const torch::stable::Tensor& C,
-    const std::optional<torch::stable::Tensor>& D_,
-    const std::optional<torch::stable::Tensor>& z_,
-    const std::optional<torch::stable::Tensor>& delta_bias_,
-    bool delta_softplus,
-    const std::optional<torch::stable::Tensor>& query_start_loc,
-    const std::optional<torch::stable::Tensor>& cache_indices,
-    const std::optional<torch::stable::Tensor>& has_initial_state,
-    const torch::stable::Tensor& ssm_states, int64_t null_block_id,
-    int64_t block_size,
-    const std::optional<torch::stable::Tensor>& block_idx_first_scheduled_token,
-    const std::optional<torch::stable::Tensor>& block_idx_last_scheduled_token,
-    const std::optional<torch::stable::Tensor>& initial_state_idx,
-    const std::optional<torch::stable::Tensor>& cu_chunk_seqlen,
-    const std::optional<torch::stable::Tensor>& last_chunk_indices);
-
-using fptr_t = int64_t;
-fptr_t init_custom_ar(const std::vector<int64_t>& fake_ipc_ptrs,
-                      torch::stable::Tensor& rank_data, int64_t rank,
-                      bool fully_connected);
-void all_reduce(fptr_t _fa, torch::stable::Tensor& inp,
-                torch::stable::Tensor& out, fptr_t reg_buffer,
-                int64_t reg_buffer_sz_bytes);
-void dispose(fptr_t _fa);
-int64_t meta_size();
-void register_buffer(fptr_t _fa, const std::vector<int64_t>& fake_ipc_ptrs);
-std::tuple<std::vector<int64_t>, std::vector<int64_t>>
-get_graph_buffer_ipc_meta(fptr_t _fa);
-void register_graph_buffers(fptr_t _fa,
-                            const std::vector<std::vector<int64_t>>& handles,
-                            const std::vector<std::vector<int64_t>>& offsets);
-std::tuple<int64_t, torch::stable::Tensor> allocate_shared_buffer_and_handle(
-    int64_t size);
-int64_t open_mem_handle(torch::stable::Tensor& mem_handle);
-void free_shared_buffer(int64_t buffer);
-
 // Activation kernels (shared CUDA/ROCm)
 void silu_and_mul(torch::stable::Tensor& out, torch::stable::Tensor& input);
 void silu_and_mul_clamp(torch::stable::Tensor& out,
@@ -396,132 +302,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]

csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu

@@ -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.

csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu

@@ -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 {
 

csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu

@@ -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 {
 

csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu

@@ -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.

csrc/libtorch_stable/quantization/fp4/nvfp4_utils.cuh

@@ -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

csrc/libtorch_stable/quantization/gguf/gguf_kernel.cu

@@ -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"
 

csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu

@@ -7,11 +7,7 @@
 
 #include <cmath>
 
-#ifdef USE_ROCM
-  #include <hip/hip_fp8.h>
-#else
-  #include <cuda_fp8.h>
-#endif
+#include <cuda_fp8.h>
 
 #include "libtorch_stable/quantization/vectorization.cuh"
 #include "libtorch_stable/quantization/vectorization_utils.cuh"
@@ -19,23 +15,12 @@
 #include "libtorch_stable/torch_utils.h"
 
 __device__ __forceinline__ float GroupReduceMax(float val) {
-#ifdef USE_ROCM
-  // 16-thread logical groups may pack up to four per 64-lane wavefront; use a
-  // 64-bit mask and explicit width so shuffles stay within each group.
-  const int lane_in_wave = threadIdx.x % warpSize;
-  const unsigned long long mask = 0xFFFFull << ((lane_in_wave / 16) * 16);
-  val = fmaxf(val, __shfl_xor_sync(mask, val, 8, 16));
-  val = fmaxf(val, __shfl_xor_sync(mask, val, 4, 16));
-  val = fmaxf(val, __shfl_xor_sync(mask, val, 2, 16));
-  val = fmaxf(val, __shfl_xor_sync(mask, val, 1, 16));
-#else
   unsigned mask = threadIdx.x % 32 >= 16 ? 0xffff0000 : 0x0000ffff;
 
   val = fmaxf(val, __shfl_xor_sync(mask, val, 8));
   val = fmaxf(val, __shfl_xor_sync(mask, val, 4));
   val = fmaxf(val, __shfl_xor_sync(mask, val, 2));
   val = fmaxf(val, __shfl_xor_sync(mask, val, 1));
-#endif
   return val;
 }
 
@@ -252,12 +237,10 @@ void per_token_group_quant_8bit(const torch::stable::Tensor& input,
 
   VLLM_STABLE_DISPATCH_FLOATING_TYPES(
       input.scalar_type(), "per_token_group_quant_8bit", ([&] {
-        if (dst_type == torch::headeronly::ScalarType::Char) {
+        if (dst_type == torch::headeronly::ScalarType::Float8_e4m3fn) {
+          LAUNCH_KERNEL(scalar_t, __nv_fp8_e4m3);
+        } else if (dst_type == torch::headeronly::ScalarType::Char) {
           LAUNCH_KERNEL(scalar_t, int8_t);
-        } else {
-          VLLM_STABLE_DISPATCH_FP8_TYPES(
-              dst_type, "per_token_group_quant_8bit_fp8",
-              ([&] { LAUNCH_KERNEL(scalar_t, fp8_t); }));
         }
       }));
 
@@ -334,18 +317,10 @@ __global__ void per_token_group_quant_8bit_packed_register_kernel(
 
   // 8-lane subgroup shuffle reduce (octet of the warp). The mask selects the
   // 8 lanes within the warp that share a group.
-#ifdef USE_ROCM
-  const int lane_in_wave = threadIdx.x % warpSize;
-  const unsigned long long mask = 0xFFull << (lane_in_wave & ~7);
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 4, 8));
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 2, 8));
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 1, 8));
-#else
   unsigned mask = 0xffu << (threadIdx.x & 24u);
   local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 4));
   local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 2));
   local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 1));
-#endif
 
   float y_s = local_absmax / max_8bit;
   y_s = fmaxf(y_s, 1e-10f);
@@ -528,12 +503,15 @@ void per_token_group_quant_8bit_packed(const torch::stable::Tensor& input,
 
   VLLM_STABLE_DISPATCH_HALF_TYPES(
       input.scalar_type(), "per_token_group_quant_8bit_packed_register", ([&] {
-        if (dst_type == torch::headeronly::ScalarType::Char) {
+        if (dst_type == torch::headeronly::ScalarType::Float8_e4m3fn) {
+          LAUNCH_REG_KERNEL(scalar_t, __nv_fp8_e4m3);
+        } else if (dst_type == torch::headeronly::ScalarType::Char) {
           LAUNCH_REG_KERNEL(scalar_t, int8_t);
         } else {
-          VLLM_STABLE_DISPATCH_FP8_TYPES(
-              dst_type, "per_token_group_quant_8bit_packed_fp8",
-              ([&] { LAUNCH_REG_KERNEL(scalar_t, fp8_t); }));
+          STD_TORCH_CHECK(
+              false,
+              "per_token_group_quant_8bit_packed only supports FP8/INT8 "
+              "outputs.");
         }
       }));
 

csrc/libtorch_stable/torch_bindings.cpp

@@ -7,6 +7,11 @@
 // Note: We register under namespace "_C" so ops are accessible as
 // torch.ops._C.<op_name> for compatibility with existing code.
 STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
+#ifndef USE_ROCM
+  ops.def("permute_cols(Tensor A, Tensor perm) -> Tensor");
+#endif
+
+#ifndef USE_ROCM
   // Compute per-token-group FP8 quantized tensor and scaling factor.
   // The dummy arguments are here so we can correctly fuse with RMSNorm.
   ops.def(
@@ -27,11 +32,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "output_s, int group_size, float eps, float int8_min, float int8_max) -> "
       "()");
 
-#ifndef USE_ROCM
-  ops.def("permute_cols(Tensor A, Tensor perm) -> Tensor");
-#endif
-
-#ifndef USE_ROCM
   // CUTLASS w8a8 GEMM, supporting symmetric per-tensor or per-row/column
   // quantization, as well as bias
   ops.def(
@@ -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, "
@@ -267,20 +263,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "CUBLAS_M_THRESHOLD, bool has_zp, bool n32k16_reorder) -> Tensor");
 #endif
 
-  // Merge attn states
-  // Implements section 2.2 of https://www.arxiv.org/pdf/2501.01005
-  // can be used to combine partial attention results (in the split-KV case)
-  ops.def(
-      "merge_attn_states("
-      "    Tensor! output,"
-      "    Tensor!? output_lse,"
-      "    Tensor prefix_output,"
-      "    Tensor prefix_lse,"
-      "    Tensor suffix_output,"
-      "    Tensor suffix_lse,"
-      "    int!? prefill_tokens_with_context,"
-      "    Tensor? output_scale=None) -> ()");
-
   // Hadamard transforms
   // conditionally compiled so impl registration is in source file
   ops.def("hadacore_transform(Tensor! x, bool inplace) -> Tensor");
@@ -337,44 +319,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "bool is_neox, Tensor position_ids, "
       "int forced_token_heads_per_warp=-1) -> ()");
 
-  ops.def(
-      "fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert("
-      "Tensor q_in, Tensor kv, Tensor! k_cache, "
-      "Tensor slot_mapping, Tensor position_ids, Tensor cos_sin_cache, "
-      "int q_head_padded, float eps, int cache_block_size) -> Tensor");
-
-#ifndef USE_ROCM
-  ops.def(
-      "minimax_allreduce_rms("
-      "Tensor input, Tensor norm_weight, Tensor workspace, "
-      "int rank, int nranks, float eps) -> Tensor");
-  ops.def(
-      "minimax_allreduce_rms_qk("
-      "Tensor qkv, Tensor norm_weight_q, Tensor norm_weight_k, "
-      "Tensor workspace, int q_size, int kv_size, int rank, int nranks, "
-      "float eps) -> (Tensor, Tensor)");
-#endif
-
-  // Apply repetition penalties to logits in-place.
-  ops.def(
-      "apply_repetition_penalties_(Tensor! logits, Tensor prompt_mask, "
-      "Tensor output_mask, Tensor repetition_penalties) -> ()");
-
-  // Optimized top-k per row operations.
-  ops.def(
-      "top_k_per_row_prefill(Tensor logits, Tensor rowStarts, Tensor rowEnds, "
-      "Tensor! indices, int numRows, int stride0, "
-      "int stride1, int topK) -> ()");
-
-  ops.def(
-      "top_k_per_row_decode(Tensor logits, int next_n, "
-      "Tensor seq_lens, Tensor! indices, "
-      "int numRows, int stride0, int stride1, int topK) -> ()");
-
-  ops.def(
-      "persistent_topk(Tensor logits, Tensor lengths, Tensor! output, "
-      "Tensor workspace, int k, int max_seq_len) -> ()");
-
   // Activation ops
   // Activation function used in SwiGLU.
   ops.def("silu_and_mul(Tensor! result, Tensor input) -> ()");
@@ -478,54 +422,14 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "int type, SymInt row, SymInt tokens) -> Tensor");
 
   ops.def("ggml_moe_get_block_size(int type) -> int");
-
-  // Mamba selective scan kernel
-  ops.def(
-      "selective_scan_fwd(Tensor! u, Tensor! delta,"
-      "Tensor! A, Tensor! B, Tensor! C,"
-      "Tensor? D_, Tensor!? z_, Tensor? delta_bias_,"
-      "bool delta_softplus,"
-      "Tensor? query_start_loc,"
-      "Tensor? cache_indices,"
-      "Tensor? has_initial_state,"
-      "Tensor! ssm_states,"
-      "int null_block_id,"
-      "int block_size,"
-      "Tensor? block_idx_first_scheduled_token,"
-      "Tensor? block_idx_last_scheduled_token,"
-      "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) {
+#ifndef USE_ROCM
+  ops.impl("permute_cols", TORCH_BOX(&permute_cols));
+#endif
+
+#ifndef USE_ROCM
   // Per-token group quantization
   ops.impl("per_token_group_fp8_quant", TORCH_BOX(&per_token_group_quant_fp8));
   ops.impl("per_token_group_fp8_quant_packed",
@@ -533,11 +437,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
   ops.impl("per_token_group_quant_int8",
            TORCH_BOX(&per_token_group_quant_int8));
 
-#ifndef USE_ROCM
-  ops.impl("permute_cols", TORCH_BOX(&permute_cols));
-#endif
-
-#ifndef USE_ROCM
   // CUTLASS scaled_mm ops
   ops.impl("cutlass_scaled_mm", TORCH_BOX(&cutlass_scaled_mm));
   ops.impl("cutlass_scaled_mm_azp", TORCH_BOX(&cutlass_scaled_mm_azp));
@@ -570,8 +469,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
   // files (allspark_repack.cu and allspark_qgemm_w8a16.cu)
 #endif
 
-  ops.impl("merge_attn_states", TORCH_BOX(&merge_attn_states));
-
   // Layernorm kernels (shared CUDA/ROCm)
   ops.impl("rms_norm", TORCH_BOX(&rms_norm));
   ops.impl("fused_add_rms_norm", TORCH_BOX(&fused_add_rms_norm));
@@ -589,19 +486,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
   // Positional encoding kernels (shared CUDA/ROCm)
   ops.impl("rotary_embedding", TORCH_BOX(&rotary_embedding));
   ops.impl("fused_qk_norm_rope", TORCH_BOX(&fused_qk_norm_rope));
-  ops.impl("fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert",
-           TORCH_BOX(&fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert));
-#ifndef USE_ROCM
-  ops.impl("minimax_allreduce_rms", TORCH_BOX(&minimax_allreduce_rms));
-  ops.impl("minimax_allreduce_rms_qk", TORCH_BOX(&minimax_allreduce_rms_qk));
-#endif
-
-  // Sampler kernels (shared CUDA/ROCm)
-  ops.impl("apply_repetition_penalties_",
-           TORCH_BOX(&apply_repetition_penalties_));
-  ops.impl("top_k_per_row_prefill", TORCH_BOX(&top_k_per_row_prefill));
-  ops.impl("top_k_per_row_decode", TORCH_BOX(&top_k_per_row_decode));
-  ops.impl("persistent_topk", TORCH_BOX(&persistent_topk));
 
   // Activation kernels (shared CUDA/ROCm)
   ops.impl("silu_and_mul", TORCH_BOX(&silu_and_mul));
@@ -635,10 +519,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
   ops.impl("ggml_mul_mat_a8", TORCH_BOX(&ggml_mul_mat_a8));
   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
@@ -661,154 +541,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_FRAGMENT(_C_custom_ar, custom_ar) {
-  custom_ar.def(
-      "init_custom_ar(int[] ipc_tensors, Tensor rank_data, "
-      "int rank, bool fully_connected) -> int");
-  custom_ar.def(
-      "all_reduce(int fa, Tensor inp, Tensor! out, int reg_buffer, "
-      "int reg_buffer_sz_bytes) -> ()");
-  custom_ar.def("dispose(int fa) -> ()");
-  custom_ar.def("meta_size() -> int");
-  custom_ar.def("register_buffer(int fa, int[] ipc_tensors) -> ()");
-  custom_ar.def("get_graph_buffer_ipc_meta(int fa) -> (int[], int[])");
-  custom_ar.def(
-      "register_graph_buffers(int fa, int[][] handles, int[][] offsets) -> ()");
-  custom_ar.def("allocate_shared_buffer_and_handle(int size) -> (int, Tensor)");
-  custom_ar.def("open_mem_handle(Tensor mem_handle) -> int");
-  custom_ar.def("free_shared_buffer(int ptr) -> ()");
-}
-
-STABLE_TORCH_LIBRARY_IMPL(_C_custom_ar, CUDA, custom_ar) {
-  custom_ar.impl("init_custom_ar", TORCH_BOX(&init_custom_ar));
-  custom_ar.impl("all_reduce", TORCH_BOX(&all_reduce));
-}
-
-STABLE_TORCH_LIBRARY_IMPL(_C_custom_ar, CPU, custom_ar) {
-  custom_ar.impl("open_mem_handle", TORCH_BOX(&open_mem_handle));
-}
-
-STABLE_TORCH_LIBRARY_IMPL(_C_custom_ar, CompositeExplicitAutograd, custom_ar) {
-  custom_ar.impl("dispose", TORCH_BOX(&dispose));
-  custom_ar.impl("meta_size", TORCH_BOX(&meta_size));
-  custom_ar.impl("register_buffer", TORCH_BOX(&register_buffer));
-  custom_ar.impl("get_graph_buffer_ipc_meta",
-                 TORCH_BOX(&get_graph_buffer_ipc_meta));
-  custom_ar.impl("register_graph_buffers", TORCH_BOX(&register_graph_buffers));
-  custom_ar.impl("allocate_shared_buffer_and_handle",
-                 TORCH_BOX(&allocate_shared_buffer_and_handle));
-  custom_ar.impl("free_shared_buffer", TORCH_BOX(&free_shared_buffer));
-}
-
-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)

csrc/libtorch_stable/torch_utils.h

@@ -6,7 +6,11 @@
 #include <torch/csrc/stable/tensor.h>
 #include <torch/headeronly/util/shim_utils.h>
 
-#include <cuda_runtime.h>
+#ifndef USE_ROCM
+  #include <cuda_runtime.h>
+#else
+  #include <hip/hip_runtime.h>
+#endif
 #include <cublas_v2.h>
 
 #include <deque>

csrc/mamba/mamba_ssm/selective_scan.h

@@ -12,9 +12,6 @@
     #include <hip/hip_bf16.h>
 #endif
 #include <cuda_fp16.h>
-
-#include <torch/headeronly/util/Half.h>
-#include <torch/headeronly/util/BFloat16.h>
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 
 struct SSMParamsBase {
@@ -162,8 +159,8 @@ struct Converter{
 };
 
 template<int N>
-struct Converter<torch::headeronly::Half, N>{
-    static inline __device__ void to_float(const torch::headeronly::Half (&src)[N], float (&dst)[N]) {
+struct Converter<at::Half, N>{
+    static inline __device__ void to_float(const at::Half (&src)[N], float (&dst)[N]) {
         static_assert(N % 2 == 0);
         auto &src2 = reinterpret_cast<const half2 (&)[N / 2]>(src);
         auto &dst2 = reinterpret_cast<float2 (&)[N / 2]>(dst);
@@ -174,8 +171,8 @@ struct Converter<torch::headeronly::Half, N>{
 
 #if __CUDA_ARCH__ >= 800
 template<int N>
-struct Converter<torch::headeronly::BFloat16, N>{
-    static inline __device__ void to_float(const torch::headeronly::BFloat16 (&src)[N], float (&dst)[N]) {
+struct Converter<at::BFloat16, N>{
+    static inline __device__ void to_float(const at::BFloat16 (&src)[N], float (&dst)[N]) {
         static_assert(N % 2 == 0);
         auto &src2 = reinterpret_cast<const nv_bfloat162 (&)[N / 2]>(src);
         auto &dst2 = reinterpret_cast<float2 (&)[N / 2]>(dst);

csrc/mamba/mamba_ssm/selective_scan_fwd.cu

@@ -1,9 +1,18 @@
 // clang-format off
 // adapted from https://github.com/state-spaces/mamba/blob/main/csrc/selective_scan/selective_scan_fwd_kernel.cuh
-#include "../torch_utils.h"
-#include <torch/csrc/stable/macros.h>
+#include <torch/all.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
 #include "selective_scan.h"
 
+#include <c10/util/BFloat16.h>
+#include <c10/util/Half.h>
+#ifdef USE_ROCM
+    #include <c10/hip/HIPException.h>  // For C10_HIP_CHECK and C10_HIP_KERNEL_LAUNCH_CHECK
+#else
+    #include <c10/cuda/CUDAException.h>  // For C10_CUDA_CHECK and C10_CUDA_KERNEL_LAUNCH_CHECK
+#endif
+
 #ifndef USE_ROCM
     #include <cub/block/block_load.cuh>
     #include <cub/block/block_store.cuh>
@@ -407,15 +416,15 @@ void selective_scan_fwd_launch(SSMParamsBase &params, cudaStream_t stream) {
                 auto kernel = &selective_scan_fwd_kernel<Ktraits>;
                 if (kSmemSize >= 48 * 1024) {
 #ifdef USE_ROCM
-                    STD_CUDA_CHECK(hipFuncSetAttribute(
+                    C10_HIP_CHECK(hipFuncSetAttribute(
                         reinterpret_cast<const void*>(kernel), hipFuncAttributeMaxDynamicSharedMemorySize, kSmemSize));
 #else
-                    STD_CUDA_CHECK(cudaFuncSetAttribute(
+                    C10_CUDA_CHECK(cudaFuncSetAttribute(
                         kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, kSmemSize));
 #endif
                 }
                 kernel<<<grid, Ktraits::kNThreads, kSmemSize, stream>>>(params);
-                STD_CUDA_KERNEL_LAUNCH_CHECK();
+                C10_CUDA_KERNEL_LAUNCH_CHECK();
             });
         });
     });
@@ -453,46 +462,46 @@ void selective_scan_fwd_cuda(SSMParamsBase &params, cudaStream_t stream) {
     #endif
 }
 
-template void selective_scan_fwd_cuda<torch::headeronly::BFloat16, float, torch::headeronly::BFloat16>(SSMParamsBase &params, cudaStream_t stream);
-template void selective_scan_fwd_cuda<torch::headeronly::BFloat16, float, float>(SSMParamsBase &params, cudaStream_t stream);
-template void selective_scan_fwd_cuda<torch::headeronly::Half, float, torch::headeronly::Half>(SSMParamsBase &params, cudaStream_t stream);
-template void selective_scan_fwd_cuda<torch::headeronly::Half, float, float>(SSMParamsBase &params, cudaStream_t stream);
+template void selective_scan_fwd_cuda<at::BFloat16, float, at::BFloat16>(SSMParamsBase &params, cudaStream_t stream);
+template void selective_scan_fwd_cuda<at::BFloat16, float, float>(SSMParamsBase &params, cudaStream_t stream);
+template void selective_scan_fwd_cuda<at::Half, float, at::Half>(SSMParamsBase &params, cudaStream_t stream);
+template void selective_scan_fwd_cuda<at::Half, float, float>(SSMParamsBase &params, cudaStream_t stream);
 template void selective_scan_fwd_cuda<float, float, float>(SSMParamsBase &params, cudaStream_t stream);
 
-#define CHECK_SHAPE(x, ...) STD_TORCH_CHECK(x.sizes().equals(torch::headeronly::IntHeaderOnlyArrayRef({__VA_ARGS__})), #x " must have shape (" #__VA_ARGS__ ")")
+#define CHECK_SHAPE(x, ...) TORCH_CHECK(x.sizes() == torch::IntArrayRef({__VA_ARGS__}), #x " must have shape (" #__VA_ARGS__ ")")
 
 #define DISPATCH_WTYPE_ITYPE_FLOAT_AND_HALF_AND_BF16(ITYPE, STYPE, NAME, ...)       \
-    if (ITYPE == torch::headeronly::ScalarType::Half) {                             \
-        using input_t = torch::headeronly::Half;                                                  \
+    if (ITYPE == at::ScalarType::Half) {                                            \
+        using input_t = at::Half;                                                   \
         using weight_t = float;                                                     \
-        if (STYPE == torch::headeronly::ScalarType::Half) {                         \
-            using state_t = torch::headeronly::Half;                                              \
+        if (STYPE == at::ScalarType::Half) {                                        \
+            using state_t = at::Half;                                               \
             __VA_ARGS__();                                                          \
-        } else if (STYPE == torch::headeronly::ScalarType::Float) {                 \
+        } else if (STYPE == at::ScalarType::Float) {                                \
             using state_t = float;                                                  \
             __VA_ARGS__();                                                          \
         } else {                                                                    \
-            STD_TORCH_CHECK(false, #NAME " not implemented for state type '", STYPE, "'"); \
+            AT_ERROR(#NAME, " not implemented for state type '", toString(STYPE), "'"); \
         }                                                                           \
-    } else if (ITYPE == torch::headeronly::ScalarType::BFloat16) {                  \
-        using input_t = torch::headeronly::BFloat16;                                              \
+    } else if (ITYPE == at::ScalarType::BFloat16) {                                 \
+        using input_t = at::BFloat16;                                               \
         using weight_t = float;                                                     \
-        if (STYPE == torch::headeronly::ScalarType::BFloat16) {                     \
-            using state_t = torch::headeronly::BFloat16;                                          \
+        if (STYPE == at::ScalarType::BFloat16) {                                    \
+            using state_t = at::BFloat16;                                           \
             __VA_ARGS__();                                                          \
-        } else if (STYPE == torch::headeronly::ScalarType::Float) {                 \
+        } else if (STYPE == at::ScalarType::Float) {                                \
             using state_t = float;                                                  \
             __VA_ARGS__();                                                          \
         } else {                                                                    \
-            STD_TORCH_CHECK(false, #NAME " not implemented for state type '", STYPE, "'"); \
+            AT_ERROR(#NAME, " not implemented for state type '", toString(STYPE), "'"); \
         }                                                                           \
-    } else if (ITYPE == torch::headeronly::ScalarType::Float)  {                    \
+    } else if (ITYPE == at::ScalarType::Float)  {                                   \
         using input_t = float;                                                      \
         using weight_t = float;                                                     \
         using state_t = float;                                                      \
         __VA_ARGS__();                                                              \
     } else {                                                                        \
-        STD_TORCH_CHECK(false, #NAME " not implemented for input type '", ITYPE, "'"); \
+        AT_ERROR(#NAME, " not implemented for input type '", toString(ITYPE), "'"); \
     }
 
 
@@ -509,30 +518,30 @@ void set_ssm_params_fwd(SSMParamsBase &params,
                         const bool is_variable_B,
                         const bool is_variable_C,
                         // device pointers
-                        const torch::stable::Tensor u,
-                        const torch::stable::Tensor delta,
-                        const torch::stable::Tensor A,
-                        const torch::stable::Tensor B,
-                        const torch::stable::Tensor C,
-                        const torch::stable::Tensor out,
-                        const torch::stable::Tensor z,
-                        const torch::stable::Tensor out_z,
-                        const std::optional<torch::stable::Tensor>& D,
-                        const std::optional<torch::stable::Tensor>& delta_bias,
-                        const torch::stable::Tensor ssm_states,
+                        const torch::Tensor u,
+                        const torch::Tensor delta,
+                        const torch::Tensor A,
+                        const torch::Tensor B,
+                        const torch::Tensor C,
+                        const torch::Tensor out,
+                        const torch::Tensor z,
+                        const torch::Tensor out_z,
+                        const std::optional<at::Tensor>& D,
+                        const std::optional<at::Tensor>& delta_bias,
+                        const torch::Tensor ssm_states,
                         bool has_z,
                         bool delta_softplus,
-                        const std::optional<torch::stable::Tensor>& query_start_loc,
-                        const std::optional<torch::stable::Tensor>& cache_indices,
-                        const std::optional<torch::stable::Tensor>& has_initial_state,
+                        const std::optional<at::Tensor>& query_start_loc,
+                        const std::optional<at::Tensor>& cache_indices,
+                        const std::optional<at::Tensor>& has_initial_state,
                         bool varlen,
                         int64_t null_block_id,
                         int64_t block_size,
-                        const std::optional<torch::stable::Tensor> &block_idx_first_scheduled_token,
-                        const std::optional<torch::stable::Tensor> &block_idx_last_scheduled_token,
-                        const std::optional<torch::stable::Tensor> &initial_state_idx,
-                        const std::optional<torch::stable::Tensor> &cu_chunk_seqlen,
-                        const std::optional<torch::stable::Tensor> &last_chunk_indices) {
+                        const std::optional<torch::Tensor> &block_idx_first_scheduled_token,
+                        const std::optional<torch::Tensor> &block_idx_last_scheduled_token,
+                        const std::optional<torch::Tensor> &initial_state_idx,
+                        const std::optional<torch::Tensor> &cu_chunk_seqlen,
+                        const std::optional<torch::Tensor> &last_chunk_indices) {
 
     // Reset the parameters
     memset(&params, 0, sizeof(params));
@@ -645,45 +654,45 @@ void set_ssm_params_fwd(SSMParamsBase &params,
     }
 }
 
-void selective_scan_fwd(const torch::stable::Tensor &u, const torch::stable::Tensor &delta,
-                  const torch::stable::Tensor &A, const torch::stable::Tensor &B, const torch::stable::Tensor &C,
-                  const std::optional<torch::stable::Tensor> &D_,
-                  const std::optional<torch::stable::Tensor> &z_,
-                  const std::optional<torch::stable::Tensor> &delta_bias_,
+void selective_scan_fwd(const torch::Tensor &u, const torch::Tensor &delta,
+                  const torch::Tensor &A, const torch::Tensor &B, const torch::Tensor &C,
+                  const std::optional<torch::Tensor> &D_,
+                  const std::optional<torch::Tensor> &z_,
+                  const std::optional<torch::Tensor> &delta_bias_,
                   bool delta_softplus,
-                  const std::optional<torch::stable::Tensor> &query_start_loc,
-                  const std::optional<torch::stable::Tensor> &cache_indices,
-                  const std::optional<torch::stable::Tensor> &has_initial_state,
-                  const torch::stable::Tensor &ssm_states,
+                  const std::optional<torch::Tensor> &query_start_loc,
+                  const std::optional<torch::Tensor> &cache_indices,
+                  const std::optional<torch::Tensor> &has_initial_state,
+                  const torch::Tensor &ssm_states,
                   // used to identify padding entries if cache_indices provided
                   // in case of padding, the kernel will return early
                   int64_t null_block_id,
                   int64_t block_size,
-                  const std::optional<torch::stable::Tensor> &block_idx_first_scheduled_token,
-                  const std::optional<torch::stable::Tensor> &block_idx_last_scheduled_token,
-                  const std::optional<torch::stable::Tensor> &initial_state_idx,
-                  const std::optional<torch::stable::Tensor> &cu_chunk_seqlen,
-                  const std::optional<torch::stable::Tensor> &last_chunk_indices) {
+                  const std::optional<torch::Tensor> &block_idx_first_scheduled_token,
+                  const std::optional<torch::Tensor> &block_idx_last_scheduled_token,
+                  const std::optional<torch::Tensor> &initial_state_idx,
+                  const std::optional<torch::Tensor> &cu_chunk_seqlen,
+                  const std::optional<torch::Tensor> &last_chunk_indices) {
     auto input_type = u.scalar_type();
     auto weight_type = A.scalar_type();
-    STD_TORCH_CHECK(input_type == torch::headeronly::ScalarType::Float || input_type == torch::headeronly::ScalarType::Half || input_type == torch::headeronly::ScalarType::BFloat16);
-    STD_TORCH_CHECK(weight_type == torch::headeronly::ScalarType::Float);
+    TORCH_CHECK(input_type == at::ScalarType::Float || input_type == at::ScalarType::Half || input_type == at::ScalarType::BFloat16);
+    TORCH_CHECK(weight_type == at::ScalarType::Float);
 
     const bool is_variable_B = B.dim() >= 3;
     const bool is_variable_C = C.dim() >= 3;
 
-    STD_TORCH_CHECK(delta.scalar_type() == input_type);
-    STD_TORCH_CHECK(B.scalar_type() == (!is_variable_B ? weight_type : input_type));
-    STD_TORCH_CHECK(C.scalar_type() == (!is_variable_C ? weight_type : input_type));
+    TORCH_CHECK(delta.scalar_type() == input_type);
+    TORCH_CHECK(B.scalar_type() == (!is_variable_B ? weight_type : input_type));
+    TORCH_CHECK(C.scalar_type() == (!is_variable_C ? weight_type : input_type));
 
-    STD_TORCH_CHECK(u.is_cuda());
-    STD_TORCH_CHECK(delta.is_cuda());
-    STD_TORCH_CHECK(A.is_cuda());
-    STD_TORCH_CHECK(B.is_cuda());
-    STD_TORCH_CHECK(C.is_cuda());
+    TORCH_CHECK(u.is_cuda());
+    TORCH_CHECK(delta.is_cuda());
+    TORCH_CHECK(A.is_cuda());
+    TORCH_CHECK(B.is_cuda());
+    TORCH_CHECK(C.is_cuda());
 
-    STD_TORCH_CHECK(u.stride(-1) == 1 || u.size(-1) == 1);
-    STD_TORCH_CHECK(delta.stride(-1) == 1 || delta.size(-1) == 1);
+    TORCH_CHECK(u.stride(-1) == 1 || u.size(-1) == 1);
+    TORCH_CHECK(delta.stride(-1) == 1 || delta.size(-1) == 1);
 
     const auto sizes = u.sizes();
     const bool varlen = query_start_loc.has_value();
@@ -693,7 +702,7 @@ void selective_scan_fwd(const torch::stable::Tensor &u, const torch::stable::Ten
     const int dstate = A.size(1);
     const int n_groups = varlen ? B.size(0) : B.size(1);
 
-    STD_TORCH_CHECK(dstate <= 256, "selective_scan only supports state dimension <= 256");
+    TORCH_CHECK(dstate <= 256, "selective_scan only supports state dimension <= 256");
 
     if (varlen) {
         CHECK_SHAPE(u, dim, seqlen);
@@ -703,94 +712,94 @@ void selective_scan_fwd(const torch::stable::Tensor &u, const torch::stable::Ten
         CHECK_SHAPE(delta, batch_size, dim, seqlen);
     }
     CHECK_SHAPE(A, dim, dstate);
-    STD_TORCH_CHECK(is_variable_B, "is_variable_B = False is disabled in favor of reduced binary size");
+    TORCH_CHECK(is_variable_B, "is_variable_B = False is disabled in favor of reduced binary size")
     if (varlen) {
         CHECK_SHAPE(B, n_groups, dstate, seqlen);
     } else {
-        CHECK_SHAPE(B, batch_size, n_groups, dstate, seqlen);
+        CHECK_SHAPE(B, batch_size, n_groups, dstate, seqlen); 
     }
-    STD_TORCH_CHECK(B.stride(-1) == 1 || B.size(-1) == 1);
+    TORCH_CHECK(B.stride(-1) == 1 || B.size(-1) == 1);
 
-    STD_TORCH_CHECK(is_variable_C, "is_variable_C = False is disabled in favor of reduced binary size");
+    TORCH_CHECK(is_variable_C, "is_variable_C = False is disabled in favor of reduced binary size")
     if (varlen) {
         CHECK_SHAPE(C, n_groups, dstate, seqlen);
     } else {
-        CHECK_SHAPE(C, batch_size, n_groups, dstate, seqlen);
+        CHECK_SHAPE(C, batch_size, n_groups, dstate, seqlen); 
     }
-    STD_TORCH_CHECK(C.stride(-1) == 1 || C.size(-1) == 1);
+    TORCH_CHECK(C.stride(-1) == 1 || C.size(-1) == 1);
 
     if (D_.has_value()) {
         auto D = D_.value();
-        STD_TORCH_CHECK(D.scalar_type() == torch::headeronly::ScalarType::Float);
-        STD_TORCH_CHECK(D.is_cuda());
-        STD_TORCH_CHECK(D.stride(-1) == 1 || D.size(-1) == 1);
+        TORCH_CHECK(D.scalar_type() == at::ScalarType::Float);
+        TORCH_CHECK(D.is_cuda());
+        TORCH_CHECK(D.stride(-1) == 1 || D.size(-1) == 1);
         CHECK_SHAPE(D, dim);
     }
 
     if (delta_bias_.has_value()) {
         auto delta_bias = delta_bias_.value();
-        STD_TORCH_CHECK(delta_bias.scalar_type() == torch::headeronly::ScalarType::Float);
-        STD_TORCH_CHECK(delta_bias.is_cuda());
-        STD_TORCH_CHECK(delta_bias.stride(-1) == 1 || delta_bias.size(-1) == 1);
+        TORCH_CHECK(delta_bias.scalar_type() == at::ScalarType::Float);
+        TORCH_CHECK(delta_bias.is_cuda());
+        TORCH_CHECK(delta_bias.stride(-1) == 1 || delta_bias.size(-1) == 1);
         CHECK_SHAPE(delta_bias, dim);
     }
 
 
     if (has_initial_state.has_value()) {
         auto has_initial_state_ = has_initial_state.value();
-        STD_TORCH_CHECK(has_initial_state_.scalar_type() == torch::headeronly::ScalarType::Bool);
-        STD_TORCH_CHECK(has_initial_state_.is_cuda());
+        TORCH_CHECK(has_initial_state_.scalar_type() == at::ScalarType::Bool);
+        TORCH_CHECK(has_initial_state_.is_cuda());
         CHECK_SHAPE(has_initial_state_, batch_size);
     }
 
 
     if (query_start_loc.has_value()) {
         auto query_start_loc_ = query_start_loc.value();
-        STD_TORCH_CHECK(query_start_loc_.scalar_type() == torch::headeronly::ScalarType::Int);
-        STD_TORCH_CHECK(query_start_loc_.is_cuda());
+        TORCH_CHECK(query_start_loc_.scalar_type() == at::ScalarType::Int);
+        TORCH_CHECK(query_start_loc_.is_cuda());
     }
 
 
     if (cache_indices.has_value()) {
         auto cache_indices_ = cache_indices.value();
-        STD_TORCH_CHECK(cache_indices_.scalar_type() == torch::headeronly::ScalarType::Int);
-        STD_TORCH_CHECK(cache_indices_.is_cuda());
+        TORCH_CHECK(cache_indices_.scalar_type() == at::ScalarType::Int);
+        TORCH_CHECK(cache_indices_.is_cuda());
 
         // cache_indices can be either 1D (batch_size,) for non-APC mode
         // or 2D (batch_size, max_positions) for APC mode
         const bool is_apc_mode = block_idx_first_scheduled_token.has_value();
         if (is_apc_mode) {
-            STD_TORCH_CHECK(cache_indices_.dim() == 2, "cache_indices must be 2D for APC mode");
-            STD_TORCH_CHECK(cache_indices_.size(0) == batch_size, "cache_indices first dimension must match batch_size");
+            TORCH_CHECK(cache_indices_.dim() == 2, "cache_indices must be 2D for APC mode");
+            TORCH_CHECK(cache_indices_.size(0) == batch_size, "cache_indices first dimension must match batch_size");
         } else {
             CHECK_SHAPE(cache_indices_, batch_size);
         }
     }
+   
 
-
-    torch::stable::Tensor z, out_z;
+    at::Tensor z, out_z;
     const bool has_z = z_.has_value();
     if (has_z) {
         z = z_.value();
-        STD_TORCH_CHECK(z.scalar_type() == input_type);
-        STD_TORCH_CHECK(z.is_cuda());
-        STD_TORCH_CHECK(z.stride(-1) == 1 || z.size(-1) == 1);
+        TORCH_CHECK(z.scalar_type() == input_type);
+        TORCH_CHECK(z.is_cuda());
+        TORCH_CHECK(z.stride(-1) == 1 || z.size(-1) == 1);
         if (varlen){
             CHECK_SHAPE(z, dim, seqlen);
         } else {
             CHECK_SHAPE(z, batch_size, dim, seqlen);
         }
-
+        
         out_z = z;
     }
 
     // Right now u has BHL layout and delta has HBL layout, and we want out to have HBL layout
-    torch::stable::Tensor out = delta;
+    at::Tensor out = delta;
     // ssm_states can now be either the same as input_type or float32
     auto state_type = ssm_states.scalar_type();
-    STD_TORCH_CHECK(state_type == input_type || state_type == torch::headeronly::ScalarType::Float);
-    STD_TORCH_CHECK(ssm_states.is_cuda());
-    STD_TORCH_CHECK(ssm_states.stride(-1) == 1);
+    TORCH_CHECK(state_type == input_type || state_type == at::ScalarType::Float);
+    TORCH_CHECK(ssm_states.is_cuda());
+    TORCH_CHECK(ssm_states.stride(-1) == 1);
 
     SSMParamsBase params;
     set_ssm_params_fwd(params, batch_size, dim, seqlen, dstate, n_groups, is_variable_B, is_variable_C,
@@ -814,8 +823,8 @@ void selective_scan_fwd(const torch::stable::Tensor &u, const torch::stable::Ten
                        );
 
     
-    const torch::stable::accelerator::DeviceGuard device_guard(u.get_device_index());
-    auto stream = get_current_cuda_stream();
+    const at::cuda::OptionalCUDAGuard device_guard(device_of(u));
+    auto stream = at::cuda::getCurrentCUDAStream().stream();
     DISPATCH_WTYPE_ITYPE_FLOAT_AND_HALF_AND_BF16(u.scalar_type(), ssm_states.scalar_type(), "selective_scan_fwd", [&] {
         selective_scan_fwd_cuda<input_t, weight_t, state_t>(params, stream);
     });