[Bugfix] Add MigCudaPlatform to avoid CUDA init before fork

NonNvmlCudaPlatform uses torch.cuda for get_device_capability(), which initializes the CUDA context. When this runs at import time (e.g. via requires_fp8 in test utils), it prevents fork-based tests from working ("Cannot re-initialize CUDA in forked subprocess"). MigCudaPlatform inherits from NvmlCudaPlatform (NVML-based capability queries don't init CUDA) and only overrides the memory methods that fail on MIG partitions. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Signed-off-by: khluu <khluu000@gmail.com>
lint
2026-04-16 11:49:08 -07:00 · 2026-04-16 02:35:54 -07:00 · 2026-04-15 16:09:29 -07:00 · 2026-04-14 23:44:59 -07:00
1161 changed files with 23938 additions and 85055 deletions
@@ -46,7 +46,7 @@ steps:
  - tests/models/language/pooling/
  commands:
    - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
      pytest -x -v -s tests/models/language/generation -m cpu_model
      pytest -x -v -s tests/models/language/pooling -m cpu_model"

@@ -69,11 +69,11 @@ steps:
      pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs
      pytest -x -v -s tests/quantization/test_cpu_wna16.py"
      
- label: CPU-Distributed Tests (PP+TP)
+- label: CPU-Distributed Tests
  depends_on: []
  device: intel_cpu
  no_plugin: true
-  source_file_dependencies: &cpu_distributed_deps
+  source_file_dependencies:
  - csrc/cpu/shm.cpp
  - vllm/v1/worker/cpu_worker.py
  - vllm/v1/worker/gpu_worker.py
@@ -82,21 +82,10 @@ steps:
  - vllm/platforms/cpu.py
  - vllm/distributed/parallel_state.py
  - vllm/distributed/device_communicators/cpu_communicator.py
-  - .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh
  commands:
    - |
      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 10m "
-      bash .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh tp_pp"
-
- label: CPU-Distributed Tests (DP+TP)
-  depends_on: []
-  device: intel_cpu
-  no_plugin: true
-  source_file_dependencies: *cpu_distributed_deps
-  commands:
-    - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 10m "
-      bash .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh dp_tp"
+      bash .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh"

 - label: CPU-Multi-Modal Model Tests %N
  depends_on: []
@@ -110,7 +99,7 @@ steps:
    - |
      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
-  parallelism: 3
+  parallelism: 2

 - label: "Arm CPU Test"
  depends_on: []
@@ -92,8 +92,8 @@ check_and_skip_if_image_exists() {
 }

 ecr_login() {
-    aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
-    aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com || true
+    aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
+    aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com
 }

 prepare_cache_tags() {
@@ -192,7 +192,6 @@ export BUILDKITE_COMMIT
 export PARENT_COMMIT
 export IMAGE_TAG
 export IMAGE_TAG_LATEST
-export COMMIT="${COMMIT:-${BUILDKITE_COMMIT}}"
 export CACHE_FROM
 export CACHE_FROM_BASE_BRANCH
 export CACHE_FROM_MAIN
@@ -11,7 +11,7 @@ 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
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"

 # skip build if image already exists
 if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu) ]]; then
@@ -11,7 +11,7 @@ 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
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"

 # skip build if image already exists
 if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu) ]]; then
@@ -46,7 +46,7 @@ echo "Image not found, proceeding with build..."

 # --- CUDA 13.0 for nightly builds ---
 # Nightly CI uses CUDA 13.0 while regular CI stays on CUDA 12.9
-NIGHTLY_CUDA_VERSION="13.0.2"
+NIGHTLY_CUDA_VERSION="13.0.0"
 NIGHTLY_BUILD_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-devel-ubuntu22.04"
 NIGHTLY_FINAL_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-base-ubuntu22.04"

@@ -1,21 +0,0 @@
-group: Engine Intel
-depends_on:
-  - image-build-xpu
-steps:
- label: Engine (1 GPU)
-  timeout_in_minutes: 30
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/v1/engine/
-    - tests/v1/engine/
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py'
@@ -1,21 +0,0 @@
-group: Kernels Intel
-depends_on: 
-  - image-build-xpu
-steps:
- label: vLLM IR Tests
-  timeout_in_minutes: 30
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/ir
-    - vllm/kernels
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      pytest -v -s kernels/ir'
@@ -1,129 +0,0 @@
-group: LoRA Intel
-depends_on:
-  - image-build-xpu
-steps:
- label: LoRA Runtime + Utils
-  timeout_in_minutes: 45
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/lora
-    - tests/lora
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      pytest -v -s lora/test_layers.py &&
-      pytest -v -s lora/test_lora_checkpoints.py &&
-      (pytest -v -s lora/test_lora_functions.py --deselect="tests/lora/test_lora_functions.py::test_lora_functions_sync" --deselect="tests/lora/test_lora_functions.py::test_lora_functions_async" || true) &&
-      pytest -v -s lora/test_lora_huggingface.py &&
-      pytest -v -s lora/test_lora_manager.py &&
-      pytest -v -s lora/test_lora_utils.py &&
-      pytest -v -s lora/test_peft_helper.py &&
-      pytest -v -s lora/test_resolver.py &&
-      pytest -v -s lora/test_utils.py &&
-      (pytest -v -s lora/test_add_lora.py --deselect="tests/lora/test_add_lora.py::test_add_lora" || true) &&
-      (pytest -v -s lora/test_worker.py --deselect="tests/lora/test_worker.py::test_worker_apply_lora" || true)'
-
- label: LoRA Fused/MoE Kernels
-  timeout_in_minutes: 45
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/lora
-    - tests/lora
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      pytest -v -s lora/test_fused_moe_lora_kernel.py && 
-      pytest -v -s lora/test_moe_lora_align_sum.py'
-
- label: LoRA Punica Kernels
-  timeout_in_minutes: 45
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/lora
-    - tests/lora
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      set -o pipefail &&
-      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-2-2049-64-32-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype1-2-64000-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype1-1-102656-32-4-4]"'
-
- label: LoRA Punica FP8/XPU Ops
-  timeout_in_minutes: 45
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/lora
-    - tests/lora
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      pytest -v -s lora/test_punica_ops_fp8.py &&
-      pytest -v -s lora/test_punica_xpu_ops.py'
-
- label: LoRA Models
-  timeout_in_minutes: 45
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/lora
-    - tests/lora
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      (pytest -v -s lora/test_mixtral.py --deselect="tests/lora/test_mixtral.py::test_mixtral_lora[4]" || true) &&
-      pytest -v -s lora/test_quant_model.py --deselect="tests/lora/test_quant_model.py::test_quant_model_lora[model0]" --deselect="tests/lora/test_quant_model.py::test_quant_model_lora[model1]" --deselect="tests/lora/test_quant_model.py::test_quant_model_tp_equality[model0]" &&
-      pytest -v -s lora/test_qwen35_densemodel_lora.py &&
-      pytest -v -s lora/test_transformers_model.py'
-
- label: LoRA Multimodal
-  timeout_in_minutes: 45
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/lora
-    - tests/lora
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      pytest -v -s lora/test_default_mm_loras.py && 
-      (pytest -v -s lora/test_qwen3_unembed.py || true) &&
-      pytest -v -s lora/test_whisper.py'
@@ -1,55 +0,0 @@
-group: Miscellaneous Intel
-depends_on:
-  - image-build-xpu
-steps:
- label: V1 Core + KV + Metrics
-  timeout_in_minutes: 30
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/
-    - tests/v1/core
-    - tests/v1/executor
-    - tests/v1/kv_offload
-    - tests/v1/worker
-    - tests/v1/kv_connector/unit
-    - tests/v1/metrics
-    - tests/entrypoints/openai/correctness/test_lmeval.py
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install -r requirements/kv_connectors.txt &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
-      cd tests &&
-      pytest -v -s v1/executor'
-
- label: V1 Sample + Logits
-  timeout_in_minutes: 30
-  device: intel_gpu
-  no_plugin: true
-  working_dir: "."
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  source_file_dependencies:
-    - vllm/
-    - tests/v1/sample
-    - tests/v1/logits_processors
-    - tests/v1/test_oracle.py
-    - tests/v1/test_request.py
-    - tests/v1/test_outputs.py
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
-      cd tests &&
-      pytest -v -s v1/logits_processors &&
-      pytest -v -s v1/test_oracle.py &&
-      pytest -v -s v1/test_request.py &&
-      pytest -v -s v1/test_outputs.py'
@@ -1,13 +1,3 @@
-# CUDA architecture lists — following PyTorch RELEASE.md
-# (https://github.com/pytorch/pytorch/blob/main/RELEASE.md)
-# SM86 included for broader Ampere coverage; SM89 for marlin fp8 support
-env:
-  CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
-  # aarch64 only architectures: 8.7 for Orin, 11.0 for Thor (since CUDA 13)
-  CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0+PTX"
-  CUDA_ARCH_X86_CU129: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
-  CUDA_ARCH_AARCH64_CU129: "8.0 8.7 8.9 9.0 10.0 12.0"
-
 steps:
  - input: "Provide Release version here"
    id: input-release-version
@@ -24,7 +14,9 @@ steps:
        agents:
          queue: arm64_cpu_queue_release
        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64_CU129}\" --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          # #NOTE: torch_cuda_arch_list is derived from upstream PyTorch build files here:
+          # https://github.com/pytorch/pytorch/blob/main/.ci/aarch64_linux/aarch64_ci_build.sh#L7
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
@@ -37,10 +29,12 @@ steps:
        agents:
          queue: arm64_cpu_queue_release
        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04  --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          # #NOTE: torch_cuda_arch_list is derived from upstream PyTorch build files here:
+          # https://github.com/pytorch/pytorch/blob/main/.ci/aarch64_linux/aarch64_ci_build.sh#L7
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04  --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
        env:
          DOCKER_BUILDKIT: "1"

@@ -53,7 +47,7 @@ steps:
          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --build-arg VLLM_BUILD_ACL=ON --tag vllm-ci:build-image --target vllm-build --progress plain -f docker/Dockerfile.cpu ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
        env:
          DOCKER_BUILDKIT: "1"

@@ -63,10 +57,10 @@ steps:
        agents:
          queue: cpu_queue_release
        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86_CU129}\" --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_31"
        env:
          DOCKER_BUILDKIT: "1"

@@ -76,10 +70,10 @@ steps:
        agents:
          queue: cpu_queue_release
        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
        env:
          DOCKER_BUILDKIT: "1"

@@ -92,7 +86,7 @@ steps:
          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --build-arg VLLM_CPU_X86=true --tag vllm-ci:build-image --target vllm-build --progress plain -f docker/Dockerfile.cpu ."
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
        env:
          DOCKER_BUILDKIT: "1"

@@ -104,205 +98,99 @@ steps:
    commands:
      - "bash .buildkite/scripts/generate-and-upload-nightly-index.sh"

-  - block: "Unblock to build release Docker images"
-    depends_on: ~
-    key: block-build-release-images
-    if: build.env("NIGHTLY") != "1"
-
  - group: "Build release Docker images"
    key: "build-release-images"
-    depends_on: block-build-release-images
-    allow_dependency_failure: true
    steps:
-      - label: "Build release image - x86_64 - CUDA 13.0"
+      - label: "Build release image - x86_64 - CUDA 12.9"
        depends_on: ~
        id: build-release-image-x86
        agents:
          queue: cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"
          # re-tag to default image tag and push, just in case arm64 build fails
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"

-      - label: "Build release image - aarch64 - CUDA 13.0"
+      - label: "Build release image - aarch64 - CUDA 12.9"
        depends_on: ~
        id: build-release-image-arm64
        agents:
          queue: arm64_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"

-      - label: "Build release image - x86_64 - CUDA 12.9"
+      - label: "Build release image - x86_64 - CUDA 13.0"
        depends_on: ~
-        id: build-release-image-x86-cuda-12-9
+        id: build-release-image-x86-cuda-13-0
        agents:
          queue: cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
-          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130 --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130"
          # re-tag to default image tag and push, just in case arm64 build fails
-          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
-          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
+          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130"
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130"

-      - label: "Build release image - aarch64 - CUDA 12.9"
+      - label: "Build release image - aarch64 - CUDA 13.0"
        depends_on: ~
-        id: build-release-image-arm64-cuda-12-9
+        id: build-release-image-arm64-cuda-13-0
        agents:
          queue: arm64_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
-          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"
+          # compute capability 12.0 for RTX-50 series / RTX PRO 6000 Blackwell, 12.1 for DGX Spark
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0 12.1' --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130 --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130"

-      - label: "Build release image - x86_64 - CUDA 13.0 - Ubuntu 24.04"
+      - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
        depends_on: ~
        id: build-release-image-x86-ubuntu2404
        agents:
          queue: cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"

-      - label: "Build release image - aarch64 - CUDA 13.0 - Ubuntu 24.04"
+      - label: "Build release image - aarch64 - CUDA 12.9 - Ubuntu 24.04"
        depends_on: ~
        id: build-release-image-arm64-ubuntu2404
        agents:
          queue: arm64_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"

-      - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
+      - label: "Build release image - x86_64 - CUDA 13.0 - Ubuntu 24.04"
        depends_on: ~
-        id: build-release-image-x86-cuda-12-9-ubuntu2404
+        id: build-release-image-x86-cuda-13-0-ubuntu2404
        agents:
          queue: cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129-ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
-          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
-          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
-          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0 12.1' --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404"
+          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404"
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404"

-      - label: "Build release image - aarch64 - CUDA 12.9 - Ubuntu 24.04"
+      - label: "Build release image - aarch64 - CUDA 13.0 - Ubuntu 24.04"
        depends_on: ~
-        id: build-release-image-arm64-cuda-12-9-ubuntu2404
+        id: build-release-image-arm64-cuda-13-0-ubuntu2404
        agents:
          queue: arm64_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129-ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
-          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0 12.1' --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404"

      - block: "Build release image for x86_64 CPU"
        key: block-cpu-release-image-build
@@ -343,7 +231,7 @@ steps:
  - group: "Publish release images"
    key: "publish-release-images"
    steps:
-      - label: "Create multi-arch manifest - CUDA 13.0"
+      - label: "Create multi-arch manifest - CUDA 12.9"
        depends_on:
          - build-release-image-x86
          - build-release-image-arm64
@@ -355,7 +243,7 @@ steps:
          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64 --amend"
          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"

-      - label: "Annotate release workflow - CUDA 13.0"
+      - label: "Annotate release workflow - CUDA 12.9"
        depends_on:
          - create-multi-arch-manifest
        id: annotate-release-workflow
@@ -364,19 +252,19 @@ steps:
        commands:
          - "bash .buildkite/scripts/annotate-release.sh"

-      - label: "Create multi-arch manifest - CUDA 12.9"
+      - label: "Create multi-arch manifest - CUDA 13.0"
        depends_on:
-          - build-release-image-x86-cuda-12-9
-          - build-release-image-arm64-cuda-12-9
-        id: create-multi-arch-manifest-cuda-12-9
+          - build-release-image-x86-cuda-13-0
+          - build-release-image-arm64-cuda-13-0
+        id: create-multi-arch-manifest-cuda-13-0
        agents:
          queue: small_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129 --amend"
-          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
+          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu130 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu130 --amend"
+          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130"

-      - label: "Create multi-arch manifest - CUDA 13.0 - Ubuntu 24.04"
+      - label: "Create multi-arch manifest - CUDA 12.9 - Ubuntu 24.04"
        depends_on:
          - build-release-image-x86-ubuntu2404
          - build-release-image-arm64-ubuntu2404
@@ -388,17 +276,17 @@ steps:
          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-ubuntu2404 --amend"
          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"

-      - label: "Create multi-arch manifest - CUDA 12.9 - Ubuntu 24.04"
+      - label: "Create multi-arch manifest - CUDA 13.0 - Ubuntu 24.04"
        depends_on:
-          - build-release-image-x86-cuda-12-9-ubuntu2404
-          - build-release-image-arm64-cuda-12-9-ubuntu2404
-        id: create-multi-arch-manifest-cuda-12-9-ubuntu2404
+          - build-release-image-x86-cuda-13-0-ubuntu2404
+          - build-release-image-arm64-cuda-13-0-ubuntu2404
+        id: create-multi-arch-manifest-cuda-13-0-ubuntu2404
        agents:
          queue: small_cpu_queue_release
        commands:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129-ubuntu2404 --amend"
-          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
+          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu130-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu130-ubuntu2404 --amend"
+          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404"

      - label: "Publish nightly multi-arch image to DockerHub"
        depends_on:
@@ -418,16 +306,16 @@ steps:
          DOCKER_BUILDKIT: "1"
          DOCKERHUB_USERNAME: "vllmbot"

-      - label: "Publish nightly multi-arch image to DockerHub - CUDA 12.9"
+      - label: "Publish nightly multi-arch image to DockerHub - CUDA 13.0"
        depends_on:
-          - create-multi-arch-manifest-cuda-12-9
+          - create-multi-arch-manifest-cuda-13-0
        if: build.env("NIGHTLY") == "1"
        agents:
          queue: small_cpu_queue_release
        commands:
-          - "bash .buildkite/scripts/push-nightly-builds.sh cu129"
+          - "bash .buildkite/scripts/push-nightly-builds.sh cu130"
          # Clean up old nightly builds (keep only last 14)
-          - "bash .buildkite/scripts/cleanup-nightly-builds.sh cu129-nightly-"
+          - "bash .buildkite/scripts/cleanup-nightly-builds.sh cu130-nightly-"
        plugins:
          - docker-login#v3.0.0:
              username: vllmbot
@@ -729,8 +617,6 @@ steps:
  - label: ":docker: Build release image - x86_64 - ROCm"
    id: build-rocm-release-image
    depends_on:
-      - step: block-build-release-images
-        allow_failure: true
      - step: build-rocm-base-wheels
        allow_failure: false
    agents:
@@ -13,12 +13,12 @@ ROCM_BASE_CACHE_KEY=$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
 buildkite-agent annotate --style 'info' --context 'release-workflow' << EOF
 To download the wheel (by commit):
 \`\`\`
-aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_35_x86_64.whl .
-aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_35_aarch64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_31_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_31_aarch64.whl .

-(Optional) For CUDA 12.9:
-aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu129-cp38-abi3-manylinux_2_31_x86_64.whl .
-aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu129-cp38-abi3-manylinux_2_31_aarch64.whl .
+(Optional) For CUDA 13.0:
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu130-cp38-abi3-manylinux_2_35_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu130-cp38-abi3-manylinux_2_35_aarch64.whl .

 (Optional) For CPU:
 aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cpu-cp38-abi3-manylinux_2_35_x86_64.whl .
@@ -33,8 +33,8 @@ To download and upload the image:

 docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64
 docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu129
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu129
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu130
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu130
 docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base
 docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm
 docker pull public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:v${RELEASE_VERSION}
@@ -50,11 +50,11 @@ docker tag vllm/vllm-openai:x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
 docker push vllm/vllm-openai:latest-x86_64
 docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64

-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu129 vllm/vllm-openai:x86_64-cu129
-docker tag vllm/vllm-openai:x86_64-cu129 vllm/vllm-openai:latest-x86_64-cu129
-docker tag vllm/vllm-openai:x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
-docker push vllm/vllm-openai:latest-x86_64-cu129
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu130 vllm/vllm-openai:x86_64-cu130
+docker tag vllm/vllm-openai:x86_64-cu130 vllm/vllm-openai:latest-x86_64-cu130
+docker tag vllm/vllm-openai:x86_64-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu130
+docker push vllm/vllm-openai:latest-x86_64-cu130
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu130

 docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64 vllm/vllm-openai:aarch64
 docker tag vllm/vllm-openai:aarch64 vllm/vllm-openai:latest-aarch64
@@ -62,11 +62,11 @@ docker tag vllm/vllm-openai:aarch64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
 docker push vllm/vllm-openai:latest-aarch64
 docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64

-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu129 vllm/vllm-openai:aarch64-cu129
-docker tag vllm/vllm-openai:aarch64-cu129 vllm/vllm-openai:latest-aarch64-cu129
-docker tag vllm/vllm-openai:aarch64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-docker push vllm/vllm-openai:latest-aarch64-cu129
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu130 vllm/vllm-openai:aarch64-cu130
+docker tag vllm/vllm-openai:aarch64-cu130 vllm/vllm-openai:latest-aarch64-cu130
+docker tag vllm/vllm-openai:aarch64-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu130
+docker push vllm/vllm-openai:latest-aarch64-cu130
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu130

 ## ROCm

@@ -104,11 +104,11 @@ docker manifest create vllm/vllm-openai:v${RELEASE_VERSION} vllm/vllm-openai:v${
 docker manifest push vllm/vllm-openai:latest
 docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}

-docker manifest rm vllm/vllm-openai:latest-cu129
-docker manifest create vllm/vllm-openai:latest-cu129 vllm/vllm-openai:latest-x86_64-cu129 vllm/vllm-openai:latest-aarch64-cu129
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-docker manifest push vllm/vllm-openai:latest-cu129
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu129
+docker manifest rm vllm/vllm-openai:latest-cu130
+docker manifest create vllm/vllm-openai:latest-cu130 vllm/vllm-openai:latest-x86_64-cu130 vllm/vllm-openai:latest-aarch64-cu130
+docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu130
+docker manifest push vllm/vllm-openai:latest-cu130
+docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu130

 docker manifest rm vllm/vllm-openai-cpu:latest || true
 docker manifest create vllm/vllm-openai-cpu:latest vllm/vllm-openai-cpu:latest-x86_64 vllm/vllm-openai-cpu:latest-arm64
@@ -29,7 +29,7 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
        TORCH_INDEX_URL=""
    fi
 else
-    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu129"
 fi
 echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"

@@ -1,142 +0,0 @@
-#!/usr/bin/env python3
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Detect the manylinux platform tag for a wheel and rename it in place.
-
-vLLM's build images produce wheels with the generic ``linux_<arch>`` platform
-tag, which installers like ``pip`` won't accept off PyPI/our index. We need to
-rewrite the platform tag to the appropriate ``manylinux_<major>_<minor>_<arch>``
-before uploading.
-
-Historically the tag was hard-coded per build (``manylinux_2_31`` for the
-Ubuntu 20.04-based image, ``manylinux_2_35`` for the Ubuntu 22.04-based
-images). That is brittle: bumping the base image silently produces wheels
-labelled with the wrong glibc requirement. This script asks ``auditwheel``
-to derive the tag from the symbol versions actually referenced by the
-binaries inside the wheel, so the label tracks reality.
-
-We can't simply call ``auditwheel repair`` -- it tries to graft external
-shared libraries into the wheel and fails on vLLM's CUDA/cuBLAS dependencies.
-Instead we use ``auditwheel.wheel_abi.analyze_wheel_abi`` directly, which is
-the same call that powers ``auditwheel show``, and read off
-``winfo.sym_policy.name``.
-
-Usage:
-    detect-manylinux-tag.py <wheel_path>
-
-The wheel is renamed in place; the new path is printed on stdout. All
-diagnostics go to stderr so callers can capture stdout safely.
-"""
-
-from __future__ import annotations
-
-import argparse
-import sys
-from pathlib import Path
-
-from auditwheel.error import (
-    AuditwheelError,
-    NonPlatformWheelError,
-    WheelToolsError,
-)
-from auditwheel.wheel_abi import analyze_wheel_abi
-from auditwheel.wheeltools import get_wheel_architecture, get_wheel_libc
-
-
-def detect_platform_tag(wheel_path: Path) -> str:
-    """Return the most precise platform tag the wheel is consistent with.
-
-    Mirrors ``auditwheel show`` but returns ``sym_policy`` rather than
-    ``overall_policy``: we only care about the glibc symbol versions used,
-    not about other policy axes (ISA extensions, blacklist, etc.) that
-    ``overall_policy`` folds in.
-    """
-    fn = wheel_path.name
-
-    try:
-        arch = get_wheel_architecture(fn)
-    except (WheelToolsError, NonPlatformWheelError):
-        # Architecture isn't deducible from the filename; let auditwheel
-        # infer it from the ELF binaries inside the wheel.
-        arch = None
-
-    try:
-        libc = get_wheel_libc(fn)
-    except WheelToolsError:
-        # An unrepaired wheel uses ``linux_<arch>``, which doesn't encode
-        # libc. Let auditwheel infer it from the ELF binaries.
-        libc = None
-
-    winfo = analyze_wheel_abi(
-        libc,
-        arch,
-        wheel_path,
-        frozenset(),
-        disable_isa_ext_check=False,
-        allow_graft=False,
-    )
-    return winfo.sym_policy.name
-
-
-def rename_wheel(wheel_path: Path, new_platform_tag: str) -> Path:
-    """Rename the wheel in place, replacing only its platform tag."""
-    # Wheel filename per PEP 427:
-    #   {distribution}-{version}(-{build})?-{python}-{abi}-{platform}.whl
-    # The platform tag is always the last ``-``-separated token before
-    # ``.whl``. Compound tags like ``manylinux_2_31_x86_64`` use ``_`` as the
-    # internal separator, so ``-``-splitting is unambiguous.
-    parts = wheel_path.stem.split("-")
-    if len(parts) < 5:
-        raise ValueError(f"Unrecognised wheel filename: {wheel_path.name}")
-    parts[-1] = new_platform_tag
-    new_path = wheel_path.with_name("-".join(parts) + ".whl")
-    if new_path != wheel_path:
-        wheel_path.rename(new_path)
-    return new_path
-
-
-def main() -> int:
-    parser = argparse.ArgumentParser(
-        description="Detect a wheel's manylinux platform tag with "
-        "auditwheel and rename the wheel in place."
-    )
-    parser.add_argument(
-        "wheel",
-        type=Path,
-        help="Path to the wheel to inspect and rename.",
-    )
-    args = parser.parse_args()
-
-    wheel_path: Path = args.wheel
-    if not wheel_path.is_file():
-        print(f"error: {wheel_path} is not a file", file=sys.stderr)
-        return 1
-
-    # Catch the things that ``analyze_wheel_abi`` and ``rename_wheel`` can
-    # raise: any subclass of ``AuditwheelError`` (pure-Python wheels,
-    # invalid libc, malformed wheels), filesystem errors, or our own
-    # ``ValueError`` for an unrecognised wheel filename. Print a single
-    # ``ERROR_TYPE: message`` line to stderr instead of a Python
-    # traceback, which is much friendlier in CI logs.
-    try:
-        new_tag = detect_platform_tag(wheel_path)
-        print(f"detected platform tag: {new_tag}", file=sys.stderr)
-        new_path = rename_wheel(wheel_path, new_tag)
-    except (AuditwheelError, ValueError, OSError) as e:
-        print(
-            f"error: failed to retag {wheel_path.name}: {type(e).__name__}: {e}",
-            file=sys.stderr,
-        )
-        return 2
-
-    if new_path != wheel_path:
-        print(f"renamed {wheel_path.name} -> {new_path.name}", file=sys.stderr)
-    else:
-        print(f"wheel already tagged {new_tag}", file=sys.stderr)
-
-    print(new_path)
-    return 0
-
-
-if __name__ == "__main__":
-    sys.exit(main())
@@ -1,54 +0,0 @@
-#!/bin/bash
-# Emit docker build flags for release image provenance metadata.
-# Keep this helper best-effort: missing Buildkite metadata should fall back to
-# local/default values instead of blocking the Docker build.
-
-# Variant examples: "", "cu129", "ubuntu2404", "cu129-ubuntu2404".
-variant="${1:-}"
-variant_suffix="${variant:+-${variant}}"
-
-image_name="${VLLM_DOCKER_IMAGE_NAME:-vllm/vllm-openai}"
-staging_repo="${VLLM_STAGING_IMAGE_REPO:-public.ecr.aws/q9t5s3a7/vllm-release-repo}"
-build_commit="${VLLM_BUILD_COMMIT:-${BUILDKITE_COMMIT:-unknown}}"
-build_pipeline="${VLLM_BUILD_PIPELINE:-${BUILDKITE_PIPELINE_ID:-${BUILDKITE_PIPELINE_SLUG:-local}}}"
-build_url="${VLLM_BUILD_URL:-${BUILDKITE_BUILD_URL:-}}"
-tag_commit="${BUILDKITE_COMMIT:-${build_commit}}"
-
-if [[ -n "${BUILDKITE:-}" || -n "${BUILDKITE_COMMIT:-}" ]]; then
-  release_version="${RELEASE_VERSION:-}"
-  if command -v buildkite-agent >/dev/null 2>&1; then
-    release_version="${release_version:-$(buildkite-agent meta-data get release-version 2>/dev/null)}"
-  fi
-  release_version="${release_version#v}"
-  release_version="${release_version:-${tag_commit}}"
-
-  staging_image_ref="${staging_repo}:${tag_commit}-$(uname -m)${variant_suffix}"
-
-  if [[ "${NIGHTLY:-}" == "1" ]]; then
-    if [[ -z "${variant}" ]]; then
-      image_tag="${image_name}:nightly-${tag_commit}"
-    elif [[ "${variant}" == cu* ]]; then
-      cuda_variant="${variant%%-*}"
-      remaining_variant="${variant#${cuda_variant}}"
-      image_tag="${image_name}:${cuda_variant}-nightly-${tag_commit}${remaining_variant}"
-    else
-      image_tag="${image_name}:nightly-${tag_commit}${variant_suffix}"
-    fi
-  else
-    image_tag="${image_name}:v${release_version}${variant_suffix}"
-  fi
-else
-  image_tag="${VLLM_IMAGE_TAG:-local/vllm-openai:dev}"
-  staging_image_ref="${image_tag}"
-fi
-
-emit_arg() {
-  printf -- "--build-arg %s=%s " "$1" "$2"
-}
-
-emit_arg VLLM_BUILD_COMMIT "${build_commit}"
-emit_arg VLLM_BUILD_PIPELINE "${build_pipeline}"
-emit_arg VLLM_BUILD_URL "${build_url}"
-# This is the intended public tag. The final digest is only known after push.
-emit_arg VLLM_IMAGE_TAG "${image_tag}"
-printf -- "--tag %s " "${staging_image_ref}"
@@ -9,14 +9,21 @@ set -ex

 BUCKET="vllm-wheels"
 INDICES_OUTPUT_DIR="indices"
-DEFAULT_VARIANT_ALIAS="cu130" # align with vLLM_MAIN_CUDA_VERSION in vllm/envs.py
+DEFAULT_VARIANT_ALIAS="cu129" # align with vLLM_MAIN_CUDA_VERSION in vllm/envs.py
+PYTHON="${PYTHON_PROG:-python3}" # try to read from env var, otherwise use python3
 SUBPATH=$BUILDKITE_COMMIT
 S3_COMMIT_PREFIX="s3://$BUCKET/$SUBPATH/"

-# Select python3 (>= 3.12) -- local if available, else a docker fallback.
-# shellcheck source=lib/select-python.sh
-source .buildkite/scripts/lib/select-python.sh
-select_python
+# detect if python3.12+ is available
+has_new_python=$($PYTHON -c "print(1 if __import__('sys').version_info >= (3,12) else 0)")
+if [[ "$has_new_python" -eq 0 ]]; then
+    # use new python from docker
+    docker pull python:3-slim
+    PYTHON="docker run --rm -u $(id -u):$(id -g) -v $(pwd):/app -w /app python:3-slim python3"
+fi
+
+echo "Using python interpreter: $PYTHON"
+echo "Python version: $($PYTHON --version)"

 # ======== generate and upload indices ========

@@ -3,37 +3,42 @@ set -euox pipefail
 export VLLM_CPU_CI_ENV=0
 export VLLM_CPU_KVCACHE_SPACE=1 # avoid OOM

-MODE=${1:-all}
+echo "--- PP+TP"
+vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -pp=2 --max-model-len=4096 &
+server_pid=$!
+timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
+vllm bench serve \
+    --backend vllm \
+    --dataset-name random \
+    --model meta-llama/Llama-3.2-3B-Instruct \
+    --num-prompts 20 \
+    --result-dir ./test_results \
+    --result-filename tp_pp.json \
+    --save-result \
+    --endpoint /v1/completions
+kill -s SIGTERM $server_pid; wait $server_pid || true
+failed_req=$(jq '.failed' ./test_results/tp_pp.json)
+if [ "$failed_req" -ne 0 ]; then
+  echo "Some requests were failed!"
+  exit 1
+fi

-run_scenario() {
-    local label="$1" result_file="$2"
-    shift 2
-    echo "--- $label"
-    vllm serve meta-llama/Llama-3.2-3B-Instruct "$@" --max-model-len=4096 &
-    local server_pid=$!
-    timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
-    vllm bench serve \
-        --backend vllm \
-        --dataset-name random \
-        --model meta-llama/Llama-3.2-3B-Instruct \
-        --num-prompts 20 \
-        --result-dir ./test_results \
-        --result-filename "$result_file" \
-        --save-result \
-        --endpoint /v1/completions
-    kill -s SIGTERM "$server_pid"; wait "$server_pid" || true
-    if [ "$(jq '.failed' "./test_results/$result_file")" -ne 0 ]; then
-        echo "Some requests were failed in $label!"
-        exit 1
-    fi
-}
-
-case "$MODE" in
-    tp_pp) run_scenario "PP+TP" tp_pp.json -tp=2 -pp=2 ;;
-    dp_tp) run_scenario "DP+TP" dp_tp.json -tp=2 -dp=2 ;;
-    all)
-        run_scenario "PP+TP" tp_pp.json -tp=2 -pp=2
-        run_scenario "DP+TP" dp_tp.json -tp=2 -dp=2
-        ;;
-    *) echo "ERROR: unknown mode '$MODE' (expected: tp_pp | dp_tp | all)" >&2; exit 1 ;;
-esac
+#echo "--- DP+TP"
+#vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -dp=2 --max-model-len=4096 &
+#server_pid=$!
+#timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
+#vllm bench serve \
+#    --backend vllm \
+#    --dataset-name random \
+#    --model meta-llama/Llama-3.2-3B-Instruct \
+#    --num-prompts 20 \
+#    --result-dir ./test_results \
+#    --result-filename dp_pp.json \
+#    --save-result \
+#    --endpoint /v1/completions
+#kill -s SIGTERM $server_pid; wait $server_pid || true
+#failed_req=$(jq '.failed' ./test_results/dp_pp.json)
+#if [ "$failed_req" -ne 0 ]; then
+#  echo "Some requests were failed!"
+#  exit 1
+#fi
@@ -51,7 +51,6 @@ function cpu_tests() {
    set -e
    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"

  # basic online serving
@@ -16,5 +16,5 @@ echo "--- :docker: Building Docker image"
 docker build --progress plain --tag "$IMAGE_NAME" --target vllm-test -f docker/Dockerfile.cpu .

 # Run the image, setting --shm-size=4g for tensor parallel.
-docker run --rm --cpuset-cpus="$CORE_RANGE" --cpuset-mems="$NUMA_NODE" -v ~/.cache/huggingface:/root/.cache/huggingface --privileged=true -e HF_TOKEN -e VLLM_CPU_KVCACHE_SPACE=16 -e VLLM_CPU_CI_ENV=1 -e VLLM_CPU_SIM_MULTI_NUMA=1 -e VLLM_CPU_ATTN_SPLIT_KV=0 --shm-size=4g "$IMAGE_NAME" \
+docker run --rm --cpuset-cpus="$CORE_RANGE" --cpuset-mems="$NUMA_NODE" -v ~/.cache/huggingface:/root/.cache/huggingface --privileged=true -e HF_TOKEN -e VLLM_CPU_KVCACHE_SPACE=16 -e VLLM_CPU_CI_ENV=1 -e VLLM_CPU_SIM_MULTI_NUMA=1 --shm-size=4g "$IMAGE_NAME" \
        timeout "$TIMEOUT_VAL" bash -c "set -euox pipefail; echo \"--- Print packages\"; pip list; echo \"--- Running tests\"; ${TEST_COMMAND}"
@@ -25,100 +25,22 @@ export PYTHONPATH=".."
 ###############################################################################

 cleanup_docker() {
-  # Share the same lock with image pull to avoid cleanup/pull races on one node.
-  local docker_lock="/tmp/docker-pull.lock"
-  exec 9>"$docker_lock"
-  flock 9
-
  docker_root=$(docker info -f '{{.DockerRootDir}}')
  if [ -z "$docker_root" ]; then
    echo "Failed to determine Docker root directory." >&2
-    flock -u 9
-    return 1
+    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%. Running aggressive CI image cleanup..."
-    cleanup_old_ci_images "${REGISTRY}/${REPO}" "${image_name}" "${DOCKER_IMAGE_CLEANUP_HOURS:-72}" 1
+    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%. Checking old CI images anyway."
-    cleanup_old_ci_images "${REGISTRY}/${REPO}" "${image_name}" "${DOCKER_IMAGE_CLEANUP_HOURS:-72}" 0
-  fi
-  echo "Old CI image cleanup completed."
-
-  flock -u 9
-}
-
-cleanup_old_ci_images() {
-  local repo_prefix="$1"
-  local current_image_ref="$2"
-  local ttl_hours="$3"
-  local aggressive_cleanup="$4"
-
-  if [[ -z "$repo_prefix" || "$repo_prefix" == "/" ]]; then
-    echo "Skip old-image cleanup: invalid repo prefix '${repo_prefix}'"
-    return 0
-  fi
-
-  if ! [[ "$ttl_hours" =~ ^[0-9]+$ ]]; then
-    echo "Invalid DOCKER_IMAGE_CLEANUP_HOURS='${ttl_hours}', fallback to 72"
-    ttl_hours=72
-  fi
-
-  local now_epoch cutoff_epoch
-  now_epoch=$(date +%s)
-  cutoff_epoch=$((now_epoch - ttl_hours * 3600))
-
-  local -a used_image_ids
-  mapfile -t used_image_ids < <(docker ps -aq | xargs -r docker inspect --format '{{.Image}}' | sort -u)
-
-  local removed_count=0
-  local examined_count=0
-  declare -A seen_ids=()
-
-  while read -r image_ref image_id; do
-    [[ -z "$image_ref" || -z "$image_id" ]] && continue
-    ((examined_count++))
-
-    # Keep the image this job is going to use.
-    if [[ "$image_ref" == "$current_image_ref" ]]; then
-      continue
-    fi
-
-    # Avoid duplicate deletes when multiple tags point to same image id.
-    if [[ -n "${seen_ids[$image_id]:-}" ]]; then
-      continue
-    fi
-    seen_ids[$image_id]=1
-
-    # Never delete images that are used by any container on this node.
-    if printf '%s\n' "${used_image_ids[@]}" | grep -qx "$image_id"; then
-      continue
-    fi
-
-    local created created_epoch
-    created=$(docker image inspect -f '{{.Created}}' "$image_id" 2>/dev/null || true)
-    [[ -z "$created" ]] && continue
-    created_epoch=$(date -d "$created" +%s 2>/dev/null || true)
-    [[ -z "$created_epoch" ]] && continue
-
-    if (( created_epoch < cutoff_epoch )) || [[ "$aggressive_cleanup" == "1" ]]; then
-      if docker image rm -f "$image_id" >/dev/null 2>&1; then
-        ((removed_count++))
-      fi
-    fi
-  done < <(docker image ls --no-trunc "$repo_prefix" --format '{{.Repository}}:{{.Tag}} {{.ID}}')
-
-  # Also trim old dangling layers; this is safe and does not remove referenced images.
-  docker image prune -f --filter "until=${ttl_hours}h" >/dev/null 2>&1 || true
-
-  if [[ "$aggressive_cleanup" == "1" ]]; then
-    echo "Examined ${examined_count} images under ${repo_prefix}, removed ${removed_count} unused images under disk pressure."
-  else
-    echo "Examined ${examined_count} images under ${repo_prefix}, removed ${removed_count} old images (>${ttl_hours}h)."
+    echo "Disk usage is below $threshold%. No cleanup needed."
  fi
 }

@@ -318,6 +240,7 @@ fi
 cleanup_docker

 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
+aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com

 # --- Build or pull test image ---
 IMAGE="${IMAGE_TAG_XPU:-${image_name}}"
@@ -343,6 +266,8 @@ fi

 remove_docker_container() {
  docker rm -f "${container_name}" || true
+  docker image rm -f "${image_name}" || true
+  docker system prune -f || true
 }
 trap remove_docker_container EXIT

@@ -358,7 +283,6 @@ docker run \
    --ipc=host \
    --privileged \
    -v /dev/dri/by-path:/dev/dri/by-path \
-    -v ${HOME}/.cache/huggingface:/root/.cache/huggingface \
    --entrypoint="" \
    -e "HF_TOKEN=${HF_TOKEN:-}" \
    -e "ZE_AFFINITY_MASK=${ZE_AFFINITY_MASK:-}" \
@@ -12,7 +12,9 @@ docker build -t "${image_name}" -f docker/Dockerfile.xpu .

 # Setup cleanup
 remove_docker_container() {
-  docker rm -f "${container_name}" || true
+  docker rm -f "${container_name}" || true;
+  docker image rm -f "${image_name}" || true;
+  docker system prune -f || true;
 }
 trap remove_docker_container EXIT

@@ -1,127 +0,0 @@
-#!/usr/bin/env bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Shared helper for rewriting a wheel's platform tag from the generic
-# ``linux_<arch>`` to the correct ``manylinux_<major>_<minor>_<arch>``.
-# After sourcing, call ``apply_manylinux_tag <wheel>`` on each wheel
-# that still carries the generic tag; the renamed path is printed on
-# stdout (logs go to stderr).
-#
-# Why a pinned Docker container instead of using whatever Python
-# happens to be on the agent:
-#   - vLLM's release agents are heterogeneous -- they don't agree on
-#     a Python minor version, and we can't rely on a particular
-#     ``auditwheel`` being installed.
-#   - ``detect-manylinux-tag.py`` reads ``auditwheel.wheel_abi`` and
-#     ``Policy.sym_policy``, which are *internal* APIs without a
-#     stability promise. Pinning both Python and auditwheel makes the
-#     detected tag a function of the inputs alone, and shifts version
-#     bumps from "implicit drift" to "deliberate, retested change".
-#   - Other release scripts (``generate-and-upload-nightly-index.sh``,
-#     ``upload-rocm-wheels.sh``) already use the python:3-slim image
-#     when the agent's interpreter is too old; this is the same idea
-#     made stricter.
-#
-# To keep the per-wheel cost down (the ROCm upload retags ~10 wheels
-# each run), we install auditwheel into a long-lived helper container
-# once on source, then ``docker exec`` into it for each call.
-#
-# Trap behaviour:
-# - Sourcing installs an EXIT trap that calls ``manylinux_cleanup`` to
-#   tear down the helper container. Any EXIT trap that was already in
-#   place when this file was sourced is captured and run AFTER our
-#   cleanup, so we don't silently clobber it.
-# - If a caller sets a new EXIT trap *after* sourcing, that trap will
-#   replace ours; in that case the caller should call
-#   ``manylinux_cleanup`` from their own handler.
-
-if [[ -n "${_MANYLINUX_LIB_SOURCED:-}" ]]; then
-    return 0
-fi
-_MANYLINUX_LIB_SOURCED=1
-
-# Pin both sides. Bump these deliberately and re-run a representative
-# wheel from each build target through the detection.
-_MANYLINUX_PYTHON_IMAGE="python:3.12-slim"
-_MANYLINUX_AUDITWHEEL_VERSION="6.6.0"
-
-# Resolve our own directory (and the sibling detect script) using the
-# canonical, symlink-resolved path. The container mounts cwd at the
-# same absolute path on both sides, so all paths we hand to it -- the
-# script, the wheel -- must canonicalise to a location under cwd.
-_MANYLINUX_LIB_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd -P)"
-_MANYLINUX_DETECT_SCRIPT="$(cd "${_MANYLINUX_LIB_DIR}/.." && pwd -P)/detect-manylinux-tag.py"
-_MANYLINUX_CWD="$(pwd -P)"
-
-docker pull --quiet "$_MANYLINUX_PYTHON_IMAGE" >/dev/null
-
-# Spin up a long-lived helper container so we install auditwheel once
-# and then ``docker exec`` into it for each wheel.
-#
-# The container runs as root so ``pip install`` can write into the
-# system site-packages; individual ``docker exec`` calls below pin
-# themselves to the host UID so any file rename happens with host
-# ownership, not root.
-_MANYLINUX_CONTAINER="$(docker run -d --rm \
-    -v "$_MANYLINUX_CWD:$_MANYLINUX_CWD" \
-    -w "$_MANYLINUX_CWD" \
-    "$_MANYLINUX_PYTHON_IMAGE" \
-    sleep infinity)"
-docker exec "$_MANYLINUX_CONTAINER" \
-    pip install --quiet --disable-pip-version-check \
-    --root-user-action=ignore \
-    "auditwheel==${_MANYLINUX_AUDITWHEEL_VERSION}"
-
-# Public cleanup -- safe to call multiple times.
-manylinux_cleanup() {
-    if [[ -n "${_MANYLINUX_CONTAINER:-}" ]]; then
-        docker rm -f "$_MANYLINUX_CONTAINER" >/dev/null 2>&1 || true
-        _MANYLINUX_CONTAINER=""
-    fi
-}
-
-# Capture any EXIT trap that was already in place so we can chain to
-# it rather than overwrite it. ``trap -p EXIT`` prints the handler in
-# eval-able form (``trap -- 'CMD' EXIT``) or nothing if unset; we
-# strip the wrapper to recover ``CMD``. Handles the common case --
-# CMDs without embedded single quotes -- and degrades gracefully (we
-# still run our own cleanup) for the pathological case.
-_manylinux_prev_exit_trap_cmd=""
-_manylinux_existing_exit_trap="$(trap -p EXIT)"
-if [[ -n "$_manylinux_existing_exit_trap" ]]; then
-    _tmp="${_manylinux_existing_exit_trap#trap -- \'}"
-    _manylinux_prev_exit_trap_cmd="${_tmp%\' EXIT}"
-    unset _tmp
-fi
-unset _manylinux_existing_exit_trap
-
-_manylinux_run_exit_chain() {
-    manylinux_cleanup
-    if [[ -n "$_manylinux_prev_exit_trap_cmd" ]]; then
-        eval "$_manylinux_prev_exit_trap_cmd"
-    fi
-}
-trap _manylinux_run_exit_chain EXIT
-
-# Detect the manylinux platform tag for a single wheel and rename it
-# in place, printing the renamed wheel path on stdout. Returns
-# non-zero on failure (which under ``set -e`` propagates to caller).
-#
-# The wheel must be reachable via a path under the host cwd so it's
-# visible inside the helper container; in CI the wheels always live
-# under ``artifacts/`` so this is fine.
-apply_manylinux_tag() {
-    local wheel="$1"
-    local abs_wheel
-    abs_wheel="$(realpath "$wheel")"
-    local new_wheel
-    new_wheel="$(docker exec -u "$(id -u):$(id -g)" \
-        "$_MANYLINUX_CONTAINER" \
-        python "$_MANYLINUX_DETECT_SCRIPT" "$abs_wheel")"
-    if [[ -z "$new_wheel" || ! -f "$new_wheel" ]]; then
-        echo "apply_manylinux_tag: detect-manylinux-tag.py did not produce a valid wheel path for $wheel" >&2
-        return 1
-    fi
-    printf '%s\n' "$new_wheel"
-}
@@ -1,41 +0,0 @@
-#!/usr/bin/env bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Pick a Python interpreter for buildkite scripts: prefer a local
-# ``python3`` if it is recent enough (>= 3.12), otherwise fall back to
-# a one-shot Docker container running ``python:3-slim``. After
-# ``select_python`` returns, ``$PYTHON`` is set in the caller's shell
-# and is safe to use as a command (e.g. ``$PYTHON some_script.py``).
-#
-# The 3.12 threshold matches what the existing nightly-index work
-# expects -- typing features used by ``generate-nightly-index.py``.
-# This helper does not pin the *minor* version; if you need stricter
-# reproducibility (e.g. relying on auditwheel internals), invoke
-# Docker yourself with a pinned tag rather than calling this.
-
-if [[ -n "${_SELECT_PYTHON_LIB_SOURCED:-}" ]]; then
-    return 0
-fi
-_SELECT_PYTHON_LIB_SOURCED=1
-
-# Sets ``PYTHON`` in the caller's shell and exports it. Idempotent --
-# calling twice is safe and the second call simply re-runs the probe.
-select_python() {
-    local py="${PYTHON_PROG:-python3}"
-    local has_new_python
-    has_new_python=$("$py" -c \
-        "print(1 if __import__('sys').version_info >= (3,12) else 0)" \
-        2>/dev/null || echo 0)
-    if [[ "$has_new_python" -eq 0 ]]; then
-        # ``-u $(id -u):$(id -g)`` so files created via the container
-        # end up owned by the host user, not root.
-        docker pull python:3-slim
-        PYTHON="docker run --rm -u $(id -u):$(id -g) -v $(pwd):/app -w /app python:3-slim python3"
-    else
-        PYTHON="$py"
-    fi
-    export PYTHON
-    echo "Using python interpreter: $PYTHON"
-    echo "Python version: $($PYTHON --version)"
-}
@@ -1,55 +0,0 @@
-#!/usr/bin/env bash
-set -euxo pipefail
-
-# args: [THRESHOLD] [NUM_QUESTIONS] [START_PORT]
-THRESHOLD=${1:-0.8}
-NUM_Q=${2:-1319}
-PORT=${3:-8050}
-OUT_DIR=${OUT_DIR:-/tmp/vllm-scheduled}
-mkdir -p "${OUT_DIR}"
-
-wait_for_server() {
-  local port=$1
-  timeout 600 bash -c '
-    until curl -sf "http://127.0.0.1:'"$port"'/health" > /dev/null; do
-      sleep 1
-    done'
-}
-
-MODEL="Qwen/Qwen3-30B-A3B-FP8"
-BACK="allgather_reducescatter"
-
-cleanup() {
-  if [[ -n "${SERVER_PID:-}" ]] && kill -0 "${SERVER_PID}" 2>/dev/null; then
-    kill "${SERVER_PID}" 2>/dev/null || true
-    for _ in {1..20}; do
-      kill -0 "${SERVER_PID}" 2>/dev/null || break
-      sleep 0.5
-    done
-    kill -9 "${SERVER_PID}" 2>/dev/null || true
-  fi
-}
-trap cleanup EXIT
-
-VLLM_DEEP_GEMM_WARMUP=skip \
-vllm serve "$MODEL" \
--enforce-eager \
--data-parallel-size 4 \
--enable-expert-parallel \
--enable-eplb \
--all2all-backend "$BACK" \
--eplb-config '{"window_size":20, "step_interval":100, "use_async":true}' \
--trust-remote-code \
--max-model-len 2048 \
--port "$PORT" &
-SERVER_PID=$!
-wait_for_server "$PORT"
-
-TAG=$(echo "$MODEL" | tr '/: \\n' '_____')
-OUT="${OUT_DIR}/${TAG}_${BACK}.json"
-python3 tests/evals/gsm8k/gsm8k_eval.py --host http://127.0.0.1 --port "$PORT" --num-questions "${NUM_Q}" --save-results "${OUT}"
-python3 - <<PY
-import json; acc=json.load(open('${OUT}'))['accuracy']
-print(f"${MODEL} ${BACK}: accuracy {acc:.3f}")
-assert acc >= ${THRESHOLD}, f"${MODEL} ${BACK} accuracy {acc}"
-PY
@@ -28,7 +28,6 @@
 #   BFCL_MAX_MODEL_LEN  - Max model length (default: 4096)
 #   BFCL_PORT           - Server port (default: 8000)
 #   BFCL_REASONING_PARSER - Reasoning parser name (default: disabled)
-#   BFCL_TEMPERATURE    - Temperature (default: 0.0)
 #   BFCL_EXTRA_ARGS     - Additional vLLM server args

 set -euo pipefail
@@ -44,7 +43,6 @@ TP_SIZE="${BFCL_TP_SIZE:-1}"
 MAX_MODEL_LEN="${BFCL_MAX_MODEL_LEN:-4096}"
 PORT="${BFCL_PORT:-8000}"
 REASONING_PARSER="${BFCL_REASONING_PARSER:-}"
-TEMPERATURE="${BFCL_TEMPERATURE:-0.0}"
 EXTRA_ARGS="${BFCL_EXTRA_ARGS:-}"

 # Set up output directory
@@ -141,7 +139,7 @@ echo "vLLM server is ready. (started in ${SECONDS_WAITED}s)"
 # be patched in-process so BFCL knows to use the OpenAI-compatible handler
 # against our local vLLM server.
 bfcl_exit_code=0
-python3 - "$MODEL" "$TEST_CATEGORY" "$NUM_THREADS" "$PORT" "$API_TYPE" "$TEMPERATURE" "$OUTPUT_DIR" << 'PYEOF' || bfcl_exit_code=$?
+python3 - "$MODEL" "$TEST_CATEGORY" "$NUM_THREADS" "$PORT" "$API_TYPE" "$OUTPUT_DIR" << 'PYEOF' || bfcl_exit_code=$?
 import os
 import sys

@@ -150,8 +148,7 @@ test_category = sys.argv[2]
 num_threads = int(sys.argv[3])
 port = sys.argv[4]
 api_type = sys.argv[5]
-temperature = float(sys.argv[6])
-output_dir = sys.argv[7] if len(sys.argv) > 7 and sys.argv[7] else os.getcwd()
+output_dir = sys.argv[6] if len(sys.argv) > 6 and sys.argv[6] else os.getcwd()

 os.environ["OPENAI_BASE_URL"] = f"http://localhost:{port}/v1"
 os.environ["OPENAI_API_KEY"] = "dummy"
@@ -207,7 +204,6 @@ gen_kwargs["model"] = [model]
 gen_kwargs["test_category"] = [c.strip() for c in test_category.split(",")]
 gen_kwargs["skip_server_setup"] = True
 gen_kwargs["num_threads"] = num_threads
-gen_kwargs["temperature"] = temperature
 generate(**gen_kwargs)

 # ---- evaluate ----
@@ -2,18 +2,14 @@

 set -ex

-# Upload a single wheel to S3, after detecting and applying the appropriate
-# manylinux platform tag with auditwheel.
+# Upload a single wheel to S3 (rename linux -> manylinux).
 # Index generation is handled separately by generate-and-upload-nightly-index.sh.

-# shellcheck source=lib/manylinux.sh
-source .buildkite/scripts/lib/manylinux.sh
-
 BUCKET="vllm-wheels"
 SUBPATH=$BUILDKITE_COMMIT
 S3_COMMIT_PREFIX="s3://$BUCKET/$SUBPATH/"

-# ========= locate the wheel ==========
+# ========= collect, rename & upload the wheel ==========

 # Assume wheels are in artifacts/dist/*.whl
 wheel_files=(artifacts/dist/*.whl)
@@ -25,9 +21,19 @@ if [[ ${#wheel_files[@]} -ne 1 ]]; then
 fi
 wheel="${wheel_files[0]}"

-# ========= detect manylinux tag and rename ==========
+# default build image uses ubuntu 20.04, which corresponds to manylinux_2_31
+# we also accept params as manylinux tag
+# refer to https://github.com/mayeut/pep600_compliance?tab=readme-ov-file#acceptable-distros-to-build-wheels
+manylinux_version="${1:-manylinux_2_31}"

-wheel="$(apply_manylinux_tag "$wheel")"
+# Rename 'linux' to the appropriate manylinux version in the wheel filename
+if [[ "$wheel" != *"linux"* ]]; then
+  echo "Error: Wheel filename does not contain 'linux': $wheel"
+  exit 1
+fi
+new_wheel="${wheel/linux/$manylinux_version}"
+mv -- "$wheel" "$new_wheel"
+wheel="$new_wheel"
 echo "Renamed wheel to: $wheel"

 # Extract the version from the wheel
@@ -20,6 +20,10 @@ BUCKET="${S3_BUCKET:-vllm-wheels}"
 ROCM_SUBPATH="rocm/${BUILDKITE_COMMIT}"
 S3_COMMIT_PREFIX="s3://$BUCKET/$ROCM_SUBPATH/"
 INDICES_OUTPUT_DIR="rocm-indices"
+PYTHON="${PYTHON_PROG:-python3}"
+
+# ROCm uses manylinux_2_35 (Ubuntu 22.04 based)
+MANYLINUX_VERSION="manylinux_2_35"

 echo "========================================"
 echo "ROCm Wheel Upload Configuration"
@@ -30,21 +34,19 @@ echo "Commit: $BUILDKITE_COMMIT"
 echo "Branch: $BUILDKITE_BRANCH"
 echo "========================================"

-# ======== Part 0: Setup Python and helpers ========
+# ======== Part 0: Setup Python ========

-# Pick a Python interpreter for index generation -- local if recent
-# enough, else a one-shot docker fallback.
-# shellcheck source=lib/select-python.sh
-source .buildkite/scripts/lib/select-python.sh
-select_python
+# Detect if python3.12+ is available
+has_new_python=$($PYTHON -c "print(1 if __import__('sys').version_info >= (3,12) else 0)" 2>/dev/null || echo 0)
+if [[ "$has_new_python" -eq 0 ]]; then
+    # Use new python from docker
+    # Use --user to ensure files are created with correct ownership (not root)
+    docker pull python:3-slim
+    PYTHON="docker run --rm --user $(id -u):$(id -g) -v $(pwd):/app -w /app python:3-slim python3"
+fi

-# Set up auditwheel-in-a-container for the manylinux retagging step.
-# Distinct from select_python: ``manylinux.sh`` deliberately pins both
-# the Python and auditwheel versions (the script reads auditwheel
-# internals) and so always runs in a known-good container regardless
-# of what's on the agent.
-# shellcheck source=lib/manylinux.sh
-source .buildkite/scripts/lib/manylinux.sh
+echo "Using python interpreter: $PYTHON"
+echo "Python version: $($PYTHON --version)"

 # ======== Part 1: Collect and prepare wheels ========

@@ -61,18 +63,11 @@ if [ "$WHEEL_COUNT" -eq 0 ]; then
    exit 1
 fi

-# Detect the appropriate manylinux platform tag for any wheel that still
-# carries the generic ``linux_<arch>`` tag, and rename it in place. We use
-# auditwheel via ``apply_manylinux_tag`` (see lib/manylinux.sh) rather than
-# a hard-coded ``manylinux_2_35`` string so that the label tracks the actual
-# glibc symbol versions used by the binaries (and stays correct if the
-# rocm_base image is rebased).
-#
-# The ``linux``/``manylinux`` filter below skips both pre-tagged wheels
-# (e.g. upstream torch) and pure-Python ``-any.whl`` wheels.
+# Rename linux to manylinux in wheel filenames
 for wheel in all-rocm-wheels/*.whl; do
    if [[ "$wheel" == *"linux"* ]] && [[ "$wheel" != *"manylinux"* ]]; then
-        new_wheel="$(apply_manylinux_tag "$wheel")"
+        new_wheel="${wheel/linux/$MANYLINUX_VERSION}"
+        mv -- "$wheel" "$new_wheel"
        echo "Renamed: $(basename "$wheel") -> $(basename "$new_wheel")"
    fi
 done
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: V1 attention (H100)
-  key: v1-attention-h100
  timeout_in_minutes: 30
  device: h100
  source_file_dependencies:
@@ -15,7 +14,6 @@ steps:
    - pytest -v -s v1/attention

 - label: V1 attention (B200)
-  key: v1-attention-b200
  timeout_in_minutes: 30
  device: b200
  source_file_dependencies:
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Basic Correctness
-  key: basic-correctness
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Benchmarks CLI Test
-  key: benchmarks-cli-test
  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
@@ -13,7 +12,6 @@ steps:
  - pytest -v -s benchmarks/

 - label: Attention Benchmarks Smoke Test (B200)
-  key: attention-benchmarks-smoke-test-b200
  device: b200
  num_gpus: 2
  optional: true
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Sequence Parallel Correctness Tests (2 GPUs)
-  key: sequence-parallel-correctness-tests-2-gpus
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/"
  num_devices: 2
@@ -18,7 +17,6 @@ steps:
  - pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py

 - label: Sequence Parallel Correctness Tests (2xH100)
-  key: sequence-parallel-correctness-tests-2xh100
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/"
  device: h100
@@ -29,7 +27,6 @@ steps:
  - pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py

 - label: AsyncTP Correctness Tests (2xH100)
-  key: asynctp-correctness-tests-2xh100
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/"
  device: h100
@@ -40,7 +37,6 @@ steps:
  - pytest -v -s tests/compile/correctness_e2e/test_async_tp.py

 - label: AsyncTP Correctness Tests (B200)
-  key: asynctp-correctness-tests-b200
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/"
  device: b200
@@ -51,7 +47,6 @@ steps:
  - pytest -v -s tests/compile/correctness_e2e/test_async_tp.py

 - label: Distributed Compile Unit Tests (2xH100)
-  key: distributed-compile-unit-tests-2xh100
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/"
  device: h100
@@ -65,7 +60,6 @@ steps:
  - pytest -s -v tests/compile/passes/distributed

 - label: Fusion and Compile Unit Tests (2xB200)
-  key: fusion-and-compile-unit-tests-2xb200
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/"
  device: b200
@@ -95,7 +89,6 @@ steps:
    - pytest -v -s tests/compile/fullgraph/test_full_graph.py::test_fp8_kv_scale_compile

 - label: Fusion E2E Quick (H100)
-  key: fusion-e2e-quick-h100
  timeout_in_minutes: 15
  working_dir: "/vllm-workspace/"
  device: h100
@@ -114,7 +107,6 @@ steps:
    - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "inductor_partition and not +rms_norm and +quant_fp8 and (qwen3 or deepseek)"

 - label: Fusion E2E Config Sweep (H100)
-  key: fusion-e2e-config-sweep-h100
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: h100
@@ -134,7 +126,6 @@ steps:
    - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "llama-3"

 - label: Fusion E2E Config Sweep (B200)
-  key: fusion-e2e-config-sweep-b200
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: b200
@@ -148,7 +139,6 @@ steps:
    - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "inductor_partition and (FLASHINFER and not +rms_norm and (not +quant_fp8 or +quant_fp8 and (qwen3 or deepseek)) or llama-3)"

 - label: Fusion E2E TP2 Quick (H100)
-  key: fusion-e2e-tp2-quick-h100
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/"
  device: h100
@@ -166,7 +156,6 @@ steps:
    - pytest -v -s tests/compile/fusions_e2e/test_tp2_async_tp.py -k "inductor_partition and not +rms_norm and (not +quant_fp8 or +quant_fp8 and (qwen3 or deepseek))"

 - label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
-  key: fusion-e2e-tp2-ar-rms-config-sweep-h100
  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/"
  device: h100
@@ -186,7 +175,6 @@ steps:
    - pytest -v -s tests/compile/fusions_e2e/test_tp2_ar_rms.py -k "llama-3"

 - label: Fusion E2E TP2 AsyncTP Config Sweep (H100)
-  key: fusion-e2e-tp2-asynctp-config-sweep-h100
  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/"
  device: h100
@@ -206,7 +194,6 @@ steps:
    - pytest -v -s tests/compile/fusions_e2e/test_tp2_async_tp.py -k "llama-3"

 - label: Fusion E2E TP2 (B200)
-  key: fusion-e2e-tp2-b200
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/"
  device: b200
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Platform Tests (CUDA)
-  key: platform-tests-cuda
  timeout_in_minutes: 15
  device: h200_18gb
  source_file_dependencies:
@@ -14,7 +13,6 @@ steps:
    - pytest -v -s cuda/test_platform_no_cuda_init.py

 - label: Cudagraph
-  key: cudagraph
  timeout_in_minutes: 20
  source_file_dependencies:
  - tests/v1/cudagraph
@@ -1,106 +0,0 @@
-group: Disaggregated
-depends_on: 
-  - image-build
-steps:
- label: Distributed NixlConnector PD accuracy (4 GPUs)
-  key: distributed-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 4
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - 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
-  timeout_in_minutes: 30
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 4
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - 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)
-  key: dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 30
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 4
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - 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)
-  key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 30
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 4
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - 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)
-  key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 20
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 4
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - 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)
-  key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
-  timeout_in_minutes: 30
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 2
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
-    - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
-    - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - 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)
-  key: nixlconnector-pd-spec-decode-acceptance-2-gpus
-  timeout_in_minutes: 30
-  device: a100
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 2
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - vllm/v1/worker/kv_connector_model_runner_mixin.py
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
-    - bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
-
- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
-  key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
-  timeout_in_minutes: 30
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 2
-  source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
-    - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
-    - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
-    - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
-    - tests/v1/kv_connector/nixl_integration/
-  commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
-    - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Distributed Comm Ops
-  key: distributed-comm-ops
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -17,7 +16,6 @@ steps:
  - pytest -v -s distributed/test_shm_storage.py

 - label: Distributed DP Tests (2 GPUs)
-  key: distributed-dp-tests-2-gpus
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -39,7 +37,6 @@ steps:
  - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py

 - label: Distributed Compile + RPC Tests (2 GPUs)
-  key: distributed-compile-rpc-tests-2-gpus
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -62,7 +59,6 @@ steps:
  - pytest -v -s ./compile/test_wrapper.py

 - label: Distributed Torchrun + Shutdown Tests (2 GPUs)
-  key: distributed-torchrun-shutdown-tests-2-gpus
  timeout_in_minutes: 20
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -85,7 +81,6 @@ steps:
  - pytest -v -s v1/worker/test_worker_memory_snapshot.py

 - label: Distributed Torchrun + Examples (4 GPUs)
-  key: distributed-torchrun-examples-4-gpus
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace"
  num_devices: 4
@@ -93,8 +88,9 @@ steps:
  - vllm/distributed/
  - tests/distributed/test_torchrun_example.py
  - tests/distributed/test_torchrun_example_moe.py
+  - examples/offline_inference/rlhf_colocate.py
  - examples/rl/
-  - tests/examples/features/data_parallel/data_parallel_offline.py
+  - tests/examples/offline_inference/data_parallel.py
  commands:
  # https://github.com/NVIDIA/nccl/issues/1838
  - export NCCL_CUMEM_HOST_ENABLE=0
@@ -111,13 +107,12 @@ steps:
  # test with torchrun tp=2 and dp=2 with ep
  - TP_SIZE=2 DP_SIZE=2 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
  # test with internal dp
-  - python3 examples/features/data_parallel/data_parallel_offline.py --enforce-eager
+  - python3 examples/offline_inference/data_parallel.py --enforce-eager
  # rlhf examples
  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_nccl.py
  - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_ipc.py

 - label: Distributed DP Tests (4 GPUs)
-  key: distributed-dp-tests-4-gpus
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -138,7 +133,6 @@ steps:
  - pytest -v -s distributed/test_utils.py

 - label: Distributed Compile + Comm (4 GPUs)
-  key: distributed-compile-comm-4-gpus
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -160,28 +154,24 @@ steps:
  - pytest -v -s distributed/test_multiproc_executor.py::test_multiproc_executor_multi_node

 - label: Distributed Tests (8 GPUs)(H100)
-  key: distributed-tests-8-gpus-h100
  timeout_in_minutes: 10
  device: h100
  num_devices: 8
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
-  - examples/features/torchrun/torchrun_dp_example_offline.py
+  - examples/offline_inference/torchrun_dp_example.py
  - vllm/config/parallel.py
  - vllm/distributed/
  - vllm/v1/engine/llm_engine.py
  - vllm/v1/executor/uniproc_executor.py
  - vllm/v1/worker/gpu_worker.py
-  - tests/distributed/test_mnnvl_alltoall.py
-
  commands:
  # https://github.com/NVIDIA/nccl/issues/1838
  - export NCCL_CUMEM_HOST_ENABLE=0
  # test with torchrun tp=2 and dp=4 with ep
-  - torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
+  - torchrun --nproc-per-node=8 ../examples/offline_inference/torchrun_dp_example.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep

 - label: Distributed Tests (4 GPUs)(A100)
-  key: distributed-tests-4-gpus-a100
  device: a100
  optional: true
  num_devices: 4
@@ -196,7 +186,6 @@ steps:
  - pytest -v -s -x lora/test_mixtral.py

 - label: Distributed Tests (2 GPUs)(H100)
-  key: distributed-tests-2-gpus-h100
  timeout_in_minutes: 15
  device: h100
  optional: true
@@ -205,13 +194,11 @@ steps:
  commands:
    - pytest -v -s tests/distributed/test_context_parallel.py
    - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
-    - VLLM_USE_DEEP_GEMM=1 VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
+    - VLLM_USE_DEEP_GEMM=1 VLLM_LOGGING_LEVEL=DEBUG python3 examples/offline_inference/data_parallel.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
    - pytest -v -s tests/v1/distributed/test_dbo.py
-    - VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s tests/distributed/test_weight_transfer.py
-    - pytest -v -s tests/distributed/test_packed_tensor.py
+    - TP_SIZE=1 DP_SIZE=2 pytest -v -s tests/v1/distributed/test_eagle_dp.py

 - label: Distributed Tests (2 GPUs)(B200)
-  key: distributed-tests-2-gpus-b200
  device: b200
  optional: true
  working_dir: "/vllm-workspace/"
@@ -220,12 +207,8 @@ steps:
    - pytest -v -s tests/distributed/test_context_parallel.py
    - pytest -v -s tests/distributed/test_nccl_symm_mem_allreduce.py
    - pytest -v -s tests/v1/distributed/test_dbo.py
-    - pytest -v -s tests/distributed/test_mnnvl_alltoall.py
-
-    

 - label: 2 Node Test (4 GPUs)
-  key: 2-node-test-4-gpus
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -238,12 +221,96 @@ steps:
  - vllm/executor/
  - vllm/model_executor/models/
  - tests/distributed/
-  - tests/examples/features/data_parallel/data_parallel_offline.py
+  - tests/examples/offline_inference/data_parallel.py
  commands:
-    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_multi_node_assignment.py && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"
+    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_multi_node_assignment.py && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"
+
+- label: Distributed NixlConnector PD accuracy (4 GPUs)
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 4
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+
+- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 4
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - 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)
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 4
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - 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: Hyrbid SSM NixlConnector PD accuracy tests (4 GPUs)
+  timeout_in_minutes: 20
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 4
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - 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)
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - 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)
+  timeout_in_minutes: 30
+  device: a100
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+    - vllm/v1/worker/kv_connector_model_runner_mixin.py
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
+
+- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh

 - label: Pipeline + Context Parallelism (4 GPUs)
-  key: pipeline-context-parallelism-4-gpus
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -258,7 +325,6 @@ steps:
  - pytest -v -s distributed/test_pipeline_parallel.py

 - label: RayExecutorV2 (4 GPUs)
-  key: rayexecutorv2-4-gpus
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -1,16 +0,0 @@
-group: Docker
-depends_on:
-  - image-build-cpu
-steps:
- label: Docker Build Metadata
-  timeout_in_minutes: 10
-  device: cpu-small
-  source_file_dependencies:
-    - .buildkite/release-pipeline.yaml
-    - .buildkite/scripts/docker-build-metadata-args.sh
-    - docker/Dockerfile
-    - docker/Dockerfile.cpu
-    - docker/docker-bake.hcl
-    - tests/tools/test_docker_build_metadata_args.py
-  commands:
-    - pytest -v -s tools/test_docker_build_metadata_args.py
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: DeepSeek V2-Lite Accuracy
-  key: deepseek-v2-lite-accuracy
  timeout_in_minutes: 60
  device: h100
  optional: true
@@ -13,7 +12,6 @@ steps:
  - bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh 0.25 200 8010

 - label: Qwen3-30B-A3B-FP8-block Accuracy
-  key: qwen3-30b-a3b-fp8-block-accuracy
  timeout_in_minutes: 60
  device: h100
  optional: true
@@ -23,7 +21,6 @@ steps:
  - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020

 - label: Qwen3-30B-A3B-FP8-block Accuracy (B200)
-  key: qwen3-30b-a3b-fp8-block-accuracy-b200
  timeout_in_minutes: 60
  device: b200
  optional: true
@@ -32,18 +29,7 @@ steps:
  commands:
  - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020 2 1

- label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
-  key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
-  timeout_in_minutes: 60
-  device: h100
-  optional: true
-  num_devices: 4
-  working_dir: "/vllm-workspace"
-  commands:
-  - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_dp4_async_eplb.sh 0.8 200 8050
-
 - label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
-  key: deepseek-v2-lite-prefetch-offload-accuracy-h100
  timeout_in_minutes: 60
  device: h100
  optional: true
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Engine
-  key: engine
  timeout_in_minutes: 15
  device: h200_18gb
  source_file_dependencies:
@@ -17,7 +16,6 @@ steps:
  - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py

 - label: Engine (1 GPU)
-  key: engine-1-gpu
  timeout_in_minutes: 30
  source_file_dependencies:
    - vllm/v1/engine/
@@ -27,7 +25,6 @@ steps:
    - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py

 - label: e2e Scheduling (1 GPU)
-  key: e2e-scheduling-1-gpu
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -37,7 +34,6 @@ steps:
    - pytest -v -s v1/e2e/general/test_async_scheduling.py

 - label: e2e Core (1 GPU)
-  key: e2e-core-1-gpu
  timeout_in_minutes: 30
  source_file_dependencies:
    - vllm/v1/
@@ -46,7 +42,6 @@ steps:
    - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py

 - label: V1 e2e (2 GPUs)
-  key: v1-e2e-2-gpus
  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
  optional: true
  num_devices: 2
@@ -63,7 +58,6 @@ steps:
      - image-build-amd

 - label: V1 e2e (4 GPUs)
-  key: v1-e2e-4-gpus
  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
  optional: true
  num_devices: 4
@@ -80,7 +74,6 @@ steps:
      - image-build-amd

 - label: V1 e2e (4xH100)
-  key: v1-e2e-4xh100
  timeout_in_minutes: 60
  device: h100
  num_devices: 4
@@ -2,8 +2,7 @@ group: Entrypoints
 depends_on: 
  - image-build
 steps:
- label: Entrypoints Unit Tests
-  key: entrypoints-unit-tests
+- label: Entrypoints Unit Tests  
  timeout_in_minutes: 10
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -14,7 +13,6 @@ steps:
  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling

 - label: Entrypoints Integration (LLM)
-  key: entrypoints-integration-llm
  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -28,7 +26,6 @@ steps:
  - pytest -v -s entrypoints/offline_mode # Needs to avoid interference with other tests

 - label: Entrypoints Integration (API Server openai - Part 1)
-  key: entrypoints-integration-api-server-openai-part-1
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -46,7 +43,6 @@ steps:


 - label: Entrypoints Integration (API Server openai - Part 2)
-  key: entrypoints-integration-api-server-openai-part-2
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -64,7 +60,6 @@ steps:
      - image-build-amd

 - label: Entrypoints Integration (API Server openai - Part 3)
-  key: entrypoints-integration-api-server-openai-part-3
  timeout_in_minutes: 50
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
@@ -77,7 +72,6 @@ steps:
  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py

 - label: Entrypoints Integration (API Server 2)
-  key: entrypoints-integration-api-server-2
  timeout_in_minutes: 130
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -92,7 +86,6 @@ steps:
  - pytest -v -s tool_use

 - label: Entrypoints Integration (Pooling)
-  key: entrypoints-integration-pooling
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -103,7 +96,6 @@ steps:
  - pytest -v -s entrypoints/pooling

 - label: Entrypoints Integration (Responses API)
-  key: entrypoints-integration-responses-api
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -113,7 +105,6 @@ steps:
  - pytest -v -s entrypoints/openai/responses

 - label: OpenAI API Correctness
-  key: openai-api-correctness
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: EPLB Algorithm
-  key: eplb-algorithm
  timeout_in_minutes: 15
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
@@ -16,7 +15,6 @@ steps:
  - pytest -v -s distributed/test_eplb_utils.py

 - label: EPLB Execution # 17min
-  key: eplb-execution
  timeout_in_minutes: 27
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -28,7 +26,6 @@ steps:
  - pytest -v -s distributed/test_eplb_spec_decode.py

 - label: Elastic EP Scaling Test
-  key: elastic-ep-scaling-test
  timeout_in_minutes: 20
  device: h100
  working_dir: "/vllm-workspace/tests"
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: vLLM IR Tests
-  key: vllm-ir-tests
  timeout_in_minutes: 10
  device: h200_18gb
  working_dir: "/vllm-workspace/"
@@ -15,7 +14,6 @@ steps:
    - pytest -v -s tests/kernels/ir

 - label: Kernels Core Operation Test
-  key: kernels-core-operation-test
  timeout_in_minutes: 75
  source_file_dependencies:
  - csrc/
@@ -25,7 +23,6 @@ steps:
    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py  kernels/test_concat_mla_q.py

 - label: Kernels MiniMax Reduce RMS Test (2 GPUs)
-  key: kernels-minimax-reduce-rms-test-2-gpus
  timeout_in_minutes: 15
  num_devices: 2
  device: h100
@@ -39,7 +36,6 @@ steps:
    - pytest -v -s kernels/core/test_minimax_reduce_rms.py

 - label: Kernels Attention Test %N
-  key: kernels-attention-test
  timeout_in_minutes: 35
  source_file_dependencies:
  - csrc/attention/
@@ -53,7 +49,6 @@ steps:
  parallelism: 2

 - label: Kernels Quantization Test %N
-  key: kernels-quantization-test
  timeout_in_minutes: 90
  source_file_dependencies:
  - csrc/quantization/
@@ -64,7 +59,6 @@ steps:
  parallelism: 2

 - label: Kernels MoE Test %N
-  key: kernels-moe-test
  timeout_in_minutes: 25
  source_file_dependencies:
  - csrc/quantization/cutlass_w8a8/moe/
@@ -80,7 +74,6 @@ steps:
  parallelism: 5

 - label: Kernels Mamba Test
-  key: kernels-mamba-test
  timeout_in_minutes: 45
  source_file_dependencies:
  - csrc/mamba/
@@ -89,18 +82,7 @@ steps:
  commands:
    - pytest -v -s kernels/mamba

- label: Kernels KDA Test
-  timeout_in_minutes: 20
-  source_file_dependencies:
-  - vllm/model_executor/layers/fla/ops/kda.py
-  - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
-  - vllm/model_executor/layers/fla/ops/l2norm.py
-  - tests/kernels/test_kda.py
-  commands:
-    - pytest -v -s kernels/test_kda.py
-
 - label: Kernels DeepGEMM Test (H100)
-  key: kernels-deepgemm-test-h100
  timeout_in_minutes: 45
  device: h100
  num_devices: 1
@@ -113,16 +95,13 @@ steps:
  - tests/kernels/moe/test_deepgemm.py
  - tests/kernels/moe/test_batched_deepgemm.py
  - tests/kernels/attention/test_deepgemm_attention.py
-  - tests/quantization/test_cutlass_w4a16.py
  commands:
    - pytest -v -s kernels/quantization/test_block_fp8.py
    - pytest -v -s kernels/moe/test_deepgemm.py
    - pytest -v -s kernels/moe/test_batched_deepgemm.py
    - pytest -v -s kernels/attention/test_deepgemm_attention.py
-    - pytest -v -s quantization/test_cutlass_w4a16.py

 - label: Kernels (B200)
-  key: kernels-b200
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: b200
@@ -162,7 +141,6 @@ steps:
    - pytest -v -s tests/kernels/quantization/test_nvfp4_qutlass.py
    - pytest -v -s tests/kernels/quantization/test_mxfp4_qutlass.py
    - pytest -v -s tests/kernels/moe/test_nvfp4_moe.py
-    - pytest -v -s tests/kernels/moe/test_mxfp4_moe.py
    - pytest -v -s tests/kernels/moe/test_ocp_mx_moe.py
    - pytest -v -s tests/kernels/moe/test_flashinfer.py
    - pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
@@ -171,19 +149,17 @@ steps:
    - pytest -v -s tests/models/quantization/test_nvfp4.py

 - label: Kernels Helion Test
-  key: kernels-helion-test
  timeout_in_minutes: 30
  device: h100
  source_file_dependencies:
  - vllm/utils/import_utils.py
  - tests/kernels/helion/
  commands:
-    - pip install helion==1.0.0
+    - pip install helion==0.3.3
    - pytest -v -s kernels/helion/

 
 - label: Kernels FP8 MoE Test (1 H100)
-  key: kernels-fp8-moe-test-1-h100
  timeout_in_minutes: 90
  device: h100
  num_devices: 1
@@ -200,7 +176,6 @@ steps:
    - pytest -v -s kernels/moe/test_triton_moe_ptpc_fp8.py

 - label: Kernels FP8 MoE Test (2 H100s)
-  key: kernels-fp8-moe-test-2-h100s
  timeout_in_minutes: 90
  device: h100
  num_devices: 2
@@ -210,7 +185,6 @@ steps:
    - pytest -v -s kernels/moe/test_deepep_moe.py

 - label: Kernels Fp4 MoE Test (B200)
-  key: kernels-fp4-moe-test-b200
  timeout_in_minutes: 60
  device: b200
  num_devices: 1
@@ -223,7 +197,6 @@ steps:


 - label: Kernels FusedMoE Layer Test (2 H100s)
-  key: kernels-fusedmoe-layer-test-2-h100s
  timeout_in_minutes: 90
  device: h100
  num_devices: 2
@@ -240,7 +213,6 @@ steps:


 - label: Kernels FusedMoE Layer Test (2 B200s)
-  key: kernels-fusedmoe-layer-test-2-b200s
  timeout_in_minutes: 90
  device: b200
  num_devices: 2
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: LM Eval Small Models
-  key: lm-eval-small-models
  timeout_in_minutes: 75
  source_file_dependencies:
  - csrc/
@@ -25,7 +24,6 @@ steps:
 #   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4

 - label: LM Eval Large Models (4 GPUs)(H100)
-  key: lm-eval-large-models-4-gpus-h100
  device: h100
  optional: true
  num_devices: 4
@@ -38,7 +36,6 @@ steps:
    - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-hopper.txt --tp-size=4

 - label: LM Eval Small Models (B200)
-  key: lm-eval-small-models-b200
  timeout_in_minutes: 120
  device: b200
  optional: true
@@ -49,7 +46,6 @@ steps:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt

 - label: LM Eval Qwen3.5 Models (B200)
-  key: lm-eval-qwen3-5-models-b200
  timeout_in_minutes: 120
  device: b200
  optional: true
@@ -66,7 +62,6 @@ steps:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt

 - label: LM Eval Large Models (H200)
-  key: lm-eval-large-models-h200
  timeout_in_minutes: 60
  device: h200
  optional: true
@@ -75,7 +70,6 @@ steps:
    - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200.txt

 - label: MoE Refactor Integration Test (H100 - TEMPORARY)
-  key: moe-refactor-integration-test-h100-temporary
  device: h100
  optional: true
  num_devices: 2
@@ -83,7 +77,6 @@ steps:
    - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor/config-h100.txt
  
 - label: MoE Refactor Integration Test (B200 - TEMPORARY)
-  key: moe-refactor-integration-test-b200-temporary
  device: b200
  optional: true
  num_devices: 2
@@ -91,7 +84,6 @@ steps:
    - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor/config-b200.txt

 - label: MoE Refactor Integration Test (B200 DP - TEMPORARY)
-  key: moe-refactor-integration-test-b200-dp-temporary
  device: b200
  optional: true
  num_devices: 2
@@ -100,7 +92,6 @@ steps:


 - label: LM Eval TurboQuant KV Cache
-  key: lm-eval-turboquant-kv-cache
  timeout_in_minutes: 75
  source_file_dependencies:
  - vllm/model_executor/layers/quantization/turboquant/
@@ -111,7 +102,6 @@ steps:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/models-turboquant.txt

 - label: GPQA Eval (GPT-OSS) (H100)
-  key: gpqa-eval-gpt-oss-h100
  timeout_in_minutes: 120
  device: h100
  optional: true
@@ -125,7 +115,6 @@ steps:
    - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-h100.txt

 - label: GPQA Eval (GPT-OSS) (B200)
-  key: gpqa-eval-gpt-oss-b200
  timeout_in_minutes: 120
  device: b200
  optional: true
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: LoRA %N
-  key: lora
  timeout_in_minutes: 30
  source_file_dependencies:
  - vllm/lora
@@ -14,7 +13,6 @@ steps:


 - label: LoRA TP (Distributed)
-  key: lora-tp-distributed
  timeout_in_minutes: 30
  num_devices: 4
  source_file_dependencies:
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: V1 Spec Decode
-  key: v1-spec-decode
  timeout_in_minutes: 30
  source_file_dependencies:
    - vllm/
@@ -19,7 +18,6 @@ steps:
      - image-build-amd

 - label: V1 Sample + Logits
-  key: v1-sample-logits
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -43,7 +41,6 @@ steps:
      - image-build-amd

 - label: V1 Core + KV + Metrics
-  key: v1-core-kv-metrics
  timeout_in_minutes: 30
  source_file_dependencies:
    - vllm/
@@ -74,7 +71,6 @@ steps:
      - image-build-amd

 - label: V1 Others (CPU)
-  key: v1-others-cpu
  depends_on:
    - image-build-cpu
  source_file_dependencies:
@@ -90,7 +86,6 @@ steps:
    - pytest -v -s -m 'cpu_test' v1/metrics

 - label: Regression
-  key: regression
  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
@@ -102,7 +97,6 @@ steps:
  working_dir: "/vllm-workspace/tests" # optional

 - label: Examples
-  key: examples
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/examples"
  source_file_dependencies:
@@ -119,22 +113,21 @@ steps:
    - python3 basic/offline_inference/embed.py
    - python3 basic/offline_inference/score.py
    # for multi-modal models
-    - python3 generate/multimodal/audio_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
-    - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+    - python3 offline_inference/audio_language.py --seed 0
+    - python3 offline_inference/vision_language.py --seed 0
+    - python3 offline_inference/vision_language_multi_image.py --seed 0
+    - python3 offline_inference/encoder_decoder_multimodal.py --model-type whisper --seed 0
     # for pooling models
    - python3 pooling/embed/vision_embedding_offline.py --seed 0
    # for features demo
-    - python3 features/automatic_prefix_caching/prefix_caching_offline.py
+    - python3 offline_inference/prefix_caching.py
    - python3 offline_inference/llm_engine_example.py
    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+    - python3 offline_inference/spec_decode.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
    # https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+    - python3 offline_inference/spec_decode.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536

 - label: Metrics, Tracing (2 GPUs)
-  key: metrics-tracing-2-gpus
  timeout_in_minutes: 20
  num_devices: 2
  source_file_dependencies:
@@ -149,7 +142,6 @@ steps:
  - pytest -v -s v1/tracing

 - label: Python-only Installation
-  key: python-only-installation
  depends_on: ~
  timeout_in_minutes: 20
  source_file_dependencies:
@@ -159,7 +151,6 @@ steps:
  - bash standalone_tests/python_only_compile.sh

 - label: Async Engine, Inputs, Utils, Worker
-  key: async-engine-inputs-utils-worker
  timeout_in_minutes: 50
  source_file_dependencies:
  - vllm/
@@ -172,8 +163,7 @@ steps:
  - pytest -v -s utils_

 - label: Async Engine, Inputs, Utils, Worker, Config (CPU)
-  key: async-engine-inputs-utils-worker-config-cpu
-  depends_on:
+  depends_on: 
  - image-build-cpu
  timeout_in_minutes: 30
  source_file_dependencies:
@@ -206,7 +196,6 @@ steps:
  - pytest -v -s config

 - label: Batch Invariance (H100)
-  key: batch-invariance-h100
  timeout_in_minutes: 30
  device: h100
  source_file_dependencies:
@@ -222,7 +211,6 @@ steps:
    - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]

 - label: Batch Invariance (B200)
-  key: batch-invariance-b200
  timeout_in_minutes: 30
  device: b200
  source_file_dependencies:
@@ -239,7 +227,6 @@ steps:
    - pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
  
 - label: Acceptance Length Test (Large Models) # optional
-  key: acceptance-length-test-large-models
  timeout_in_minutes: 25
  gpu: h100
  optional: true
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Model Executor
-  key: model-executor
  timeout_in_minutes: 35
  source_file_dependencies:
  - vllm/engine/arg_utils.py
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Model Runner V2 Core Tests
-  key: model-runner-v2-core-tests
  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/v1/worker/gpu/
@@ -26,16 +25,14 @@ steps:
  - pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"

 - label: Model Runner V2 Examples
-  key: model-runner-v2-examples
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/examples"
  source_file_dependencies:
    - vllm/v1/worker/gpu/
    - vllm/v1/core/sched/
    - vllm/v1/worker/gpu_worker.py
+    - examples/offline_inference/
    - examples/basic/offline_inference/
-    - examples/generate/multimodal/
-    - examples/features/
    - examples/pooling/embed/vision_embedding_offline.py
    - examples/others/tensorize_vllm_model.py
  commands:
@@ -47,22 +44,21 @@ steps:
    #- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10  # TODO
    #- python3 basic/offline_inference/embed.py   # TODO
    # for multi-modal models
-    - python3 generate/multimodal/audio_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
-    - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+    - python3 offline_inference/audio_language.py --seed 0
+    - python3 offline_inference/vision_language.py --seed 0
+    - python3 offline_inference/vision_language_multi_image.py --seed 0
+    - python3 offline_inference/encoder_decoder_multimodal.py --model-type whisper --seed 0
    # for pooling models
    - python3 pooling/embed/vision_embedding_offline.py --seed 0
    # for features demo
-    - python3 features/automatic_prefix_caching/prefix_caching_offline.py
+    - python3 offline_inference/prefix_caching.py
    - python3 offline_inference/llm_engine_example.py
    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+    - python3 offline_inference/spec_decode.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
    # https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+    - python3 offline_inference/spec_decode.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536

 - label: Model Runner V2 Distributed (2 GPUs)
-  key: model-runner-v2-distributed-2-gpus
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -83,7 +79,6 @@ steps:
    - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py

 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
-  key: model-runner-v2-pipeline-parallelism-4-gpus
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -99,7 +94,6 @@ steps:
    - pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"

 - label: Model Runner V2 Spec Decode
-  key: model-runner-v2-spec-decode
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -3,8 +3,8 @@ depends_on:
  - image-build
 steps:
 - label: Basic Models Tests (Initialization)
-  key: basic-models-tests-initialization
  timeout_in_minutes: 45
+  device: h200_18gb
  torch_nightly: true
  source_file_dependencies:
  - vllm/
@@ -17,7 +17,6 @@ steps:
    torch_nightly: {}

 - label: Basic Models Tests (Extra Initialization) %N
-  key: basic-models-tests-extra-initialization
  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/model_executor/models/
@@ -33,7 +32,6 @@ steps:
    torch_nightly: {}

 - label: Basic Models Tests (Other)
-  key: basic-models-tests-other
  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/
@@ -50,7 +48,6 @@ steps:


 - label: Basic Models Test (Other CPU) # 5min
-  key: basic-models-test-other-cpu
  depends_on:
  - image-build-cpu
  timeout_in_minutes: 10
@@ -63,7 +60,6 @@ steps:
    - pytest -v -s models/test_utils.py models/test_vision.py

 - label: Transformers Nightly Models
-  key: transformers-nightly-models
  working_dir: "/vllm-workspace/"
  optional: true
  soft_fail: true
@@ -74,22 +70,6 @@ steps:
    - 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
+    - python3 examples/offline_inference/vision_language.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
-  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
+    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/offline_inference/audio_language.py --model-type whisper
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Distributed Model Tests (2 GPUs)
-  key: distributed-model-tests-2-gpus
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Language Models Tests (Standard)
-  key: language-models-tests-standard
  timeout_in_minutes: 25
  source_file_dependencies:
  - vllm/
@@ -16,7 +15,6 @@ steps:
    torch_nightly: {}

 - label: Language Models Tests (Extra Standard) %N
-  key: language-models-tests-extra-standard
  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/model_executor/models/
@@ -33,7 +31,6 @@ steps:
    torch_nightly: {}

 - label: Language Models Tests (Hybrid) %N
-  key: language-models-tests-hybrid
  timeout_in_minutes: 75
  source_file_dependencies:
  - vllm/
@@ -50,7 +47,6 @@ steps:
    torch_nightly: {}

 - label: Language Models Test (Extended Generation) # 80min
-  key: language-models-test-extended-generation
  timeout_in_minutes: 110
  optional: true
  source_file_dependencies:
@@ -73,7 +69,6 @@ steps:
      - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'

 - label: Language Models Test (PPL)
-  key: language-models-test-ppl
  timeout_in_minutes: 110
  device: h200_18gb
  optional: true
@@ -84,7 +79,6 @@ steps:
    - pytest -v -s models/language/generation_ppl_test

 - label: Language Models Test (Extended Pooling)  # 36min
-  key: language-models-test-extended-pooling
  timeout_in_minutes: 50
  optional: true
  source_file_dependencies:
@@ -99,7 +93,6 @@ steps:
      - image-build-amd

 - label: Language Models Test (MTEB)
-  key: language-models-test-mteb
  timeout_in_minutes: 110
  device: h200_18gb
  optional: true
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
-  key: multi-modal-models-standard-1-qwen2
  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
@@ -20,7 +19,6 @@ steps:
      - image-build-amd

 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
-  key: multi-modal-models-standard-2-qwen3-gemma
  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
@@ -30,7 +28,6 @@ steps:
    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
    - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
    - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
-    - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
  mirror:
    amd:
      device: mi325_1
@@ -38,7 +35,6 @@ steps:
      - image-build-amd

 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
-  key: multi-modal-models-standard-3-llava-qwen2-vl
  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/
@@ -54,7 +50,6 @@ steps:
      - image-build-amd

 - label: "Multi-Modal Models (Standard) 4: other + whisper"
-  key: multi-modal-models-standard-4-other-whisper
  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/
@@ -71,8 +66,7 @@ steps:
      - image-build-amd

 - label: Multi-Modal Processor (CPU)
-  key: multi-modal-processor-cpu
-  depends_on:
+  depends_on: 
  - image-build-cpu
  timeout_in_minutes: 60
  source_file_dependencies:
@@ -85,7 +79,6 @@ steps:
    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py

 - label: Multi-Modal Processor # 44min
-  key: multi-modal-processor
  timeout_in_minutes: 60
  device: h200_18gb
  source_file_dependencies:
@@ -97,7 +90,6 @@ steps:
    - pytest -v -s models/multimodal/processing/test_tensor_schema.py

 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
-  key: multi-modal-accuracy-eval-small-models
  timeout_in_minutes: 70
  working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
  source_file_dependencies:
@@ -108,7 +100,6 @@ steps:
  - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-mm-small.txt --tp-size=1

 - label: Multi-Modal Models (Extended Generation 1)
-  key: multi-modal-models-extended-generation-1
  optional: true
  source_file_dependencies:
  - vllm/
@@ -125,7 +116,6 @@ steps:
      - image-build-amd

 - label: Multi-Modal Models (Extended Generation 2)
-  key: multi-modal-models-extended-generation-2
  optional: true
  source_file_dependencies:
  - vllm/
@@ -135,7 +125,6 @@ steps:
    - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'

 - label: Multi-Modal Models (Extended Generation 3)
-  key: multi-modal-models-extended-generation-3
  optional: true
  source_file_dependencies:
  - vllm/
@@ -145,7 +134,6 @@ steps:
    - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'

 - label: Multi-Modal Models (Extended Pooling)
-  key: multi-modal-models-extended-pooling
  optional: true
  device: h200_18gb
  source_file_dependencies:
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Plugin Tests (2 GPUs)
-  key: plugin-tests-2-gpus
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: PyTorch Compilation Unit Tests
-  key: pytorch-compilation-unit-tests
  timeout_in_minutes: 10
  source_file_dependencies:
    - vllm/
@@ -19,7 +18,6 @@ steps:
  - "find compile/ -maxdepth 1 -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"

 - label: PyTorch Compilation Unit Tests (H100)
-  key: pytorch-compilation-unit-tests-h100
  timeout_in_minutes: 30
  device: h100
  num_devices: 1
@@ -30,7 +28,6 @@ steps:
  - "find compile/h100/ -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"

 - label: PyTorch Compilation Passes Unit Tests
-  key: pytorch-compilation-passes-unit-tests
  timeout_in_minutes: 20
  source_file_dependencies:
    - vllm/
@@ -39,7 +36,6 @@ steps:
  - pytest -s -v compile/passes --ignore compile/passes/distributed

 - label: PyTorch Fullgraph Smoke Test
-  key: pytorch-fullgraph-smoke-test
  timeout_in_minutes: 35
  source_file_dependencies:
  - vllm/
@@ -52,7 +48,6 @@ steps:
  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"

 - label: PyTorch Fullgraph
-  key: pytorch-fullgraph
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -63,7 +58,6 @@ steps:
  - pytest -v -s compile/fullgraph/test_full_graph.py -k 'not test_fp8_kv_scale_compile'

 - label: Pytorch Nightly Dependency Override Check # 2min
-  key: pytorch-nightly-dependency-override-check
  # if this test fails, it means the nightly torch version is not compatible with some
  # of the dependencies. Please check the error message and add the package to whitelist
  # in /vllm/tools/pre_commit/generate_nightly_torch_test.py
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Quantization
-  key: quantization
  timeout_in_minutes: 90
  source_file_dependencies:
  - csrc/
@@ -22,7 +21,6 @@ steps:
  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py

 - label: Quantized MoE Test (B200)
-  key: quantized-moe-test-b200
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/"
  device: b200
@@ -40,7 +38,6 @@ steps:
    - pytest -s -v tests/quantization/test_blackwell_moe.py

 - label: Quantized Models Test
-  key: quantized-models-test
  timeout_in_minutes: 60
  source_file_dependencies:
  - vllm/model_executor/layers/quantization
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Ray Dependency Compatibility Check
-  key: ray-dependency-compatibility-check
  # Informational only — does not block the pipeline.
  # If this fails, it means the PR introduces a dependency that
  # conflicts with Ray's dependency constraints.
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Samplers Test
-  key: samplers-test
  timeout_in_minutes: 75
  source_file_dependencies:
  - vllm/model_executor/layers
@@ -11,9 +10,7 @@ steps:
  - tests/samplers
  - tests/conftest.py
  commands:
-    # VLLM_USE_FLASHINFER_SAMPLER defaults to 1 now, so we need to pin both
-    # values explicitly to still cover the PyTorch-native (Triton) path.
-    - VLLM_USE_FLASHINFER_SAMPLER=0 pytest -v -s samplers
+    - pytest -v -s samplers
    - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
  mirror:
    amd:
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Spec Decode Eagle
-  key: spec-decode-eagle
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -13,20 +12,7 @@ steps:
  commands:
    - pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"

- label: Spec Decode Eagle Nightly B200
-  key: spec-decode-eagle-nightly-b200
-  timeout_in_minutes: 30
-  device: b200
-  optional: true
-  source_file_dependencies:
-    - vllm/v1/spec_decode/
-    - vllm/v1/worker/gpu/spec_decode/
-    - tests/v1/e2e/spec_decode/
-  commands:
-    - pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
-
 - label: Spec Decode Speculators + MTP
-  key: spec-decode-speculators-mtp
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -37,21 +23,7 @@ steps:
  commands:
    - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"

- label: Spec Decode Speculators + MTP Nightly B200
-  key: spec-decode-speculators-mtp-nightly-b200
-  timeout_in_minutes: 30
-  device: b200
-  optional: true
-  source_file_dependencies:
-    - vllm/v1/spec_decode/
-    - vllm/v1/worker/gpu/spec_decode/
-    - vllm/transformers_utils/configs/speculators/
-    - tests/v1/e2e/spec_decode/
-  commands:
-    - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
-  
 - label: Spec Decode Ngram + Suffix
-  key: spec-decode-ngram-suffix
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -62,7 +34,6 @@ steps:
    - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"

 - label: Spec Decode Draft Model
-  key: spec-decode-draft-model
  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
@@ -71,42 +42,3 @@ steps:
    - tests/v1/e2e/spec_decode/
  commands:
    - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
-
- label: Spec Decode Draft Model Nightly B200
-  key: spec-decode-draft-model-nightly-b200
-  timeout_in_minutes: 30
-  device: b200
-  optional: true
-  source_file_dependencies:
-    - vllm/v1/spec_decode/
-    - vllm/v1/worker/gpu/spec_decode/
-    - tests/v1/e2e/spec_decode/
-  commands:
-    - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
-
- label: DFlash Speculators Correctness
-  key: dflash-speculators-correctness
-  timeout_in_minutes: 30
-  device: h100
-  optional: true
-  num_devices: 1
-  source_file_dependencies:
-    - vllm/v1/spec_decode/
-    - vllm/model_executor/models/qwen3_dflash.py
-    - tests/v1/spec_decode/test_speculators_dflash.py
-  commands:
-    - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
-    - pytest -v -s v1/spec_decode/test_speculators_dflash.py -m slow_test
-
- label: Spec Decode MTP hybrid (B200)
-  timeout_in_minutes: 30
-  device: b200
-  optional: true
-  source_file_dependencies:
-    - vllm/v1/spec_decode/
-    - vllm/v1/worker/gpu/spec_decode/
-    - vllm/model_executor/models/qwen3_5.py
-    - vllm/model_executor/models/qwen3_5_mtp.py
-    - tests/v1/e2e/spec_decode/
-  commands:
-    - pytest -v -s v1/e2e/spec_decode -k "qwen3_5-hybrid"
@@ -3,7 +3,6 @@ depends_on:
  - image-build
 steps:
 - label: Weight Loading Multiple GPU  # 33min
-  key: weight-loading-multiple-gpu
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -44,9 +44,8 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
 /vllm/pooling_params.py @noooop @DarkLight1337
 /vllm/tokenizers @DarkLight1337 @njhill
 /vllm/renderers @DarkLight1337 @njhill
-/vllm/reasoning @aarnphm @chaunceyjiang @sfeng33 @bbrowning
-/vllm/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
-/vllm/parser @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/vllm/reasoning @aarnphm @chaunceyjiang
+/vllm/tool_parsers @aarnphm @chaunceyjiang

 # vLLM V1
 /vllm/v1/attention @LucasWilkinson @MatthewBonanni
@@ -92,10 +91,7 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
 /tests/v1/kv_connector/nixl_integration @NickLucche
 /tests/v1/kv_connector @ApostaC @orozery
 /tests/v1/kv_offload @ApostaC @orozery
-/tests/v1/determinism @yewentao256
-/tests/reasoning @aarnphm @chaunceyjiang @sfeng33 @bbrowning
-/tests/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
-/tests/tool_use @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/tests/v1/determinism @yewentao256 

 # Transformers modeling backend
 /vllm/model_executor/models/transformers @hmellor
@@ -15,6 +15,7 @@ PLEASE FILL IN THE PR DESCRIPTION HERE ENSURING ALL CHECKLIST ITEMS (AT THE BOTT
 - [ ] The test plan, such as providing test command.
 - [ ] The test results, such as pasting the results comparison before and after, or e2e results
 - [ ] (Optional) The necessary documentation update, such as updating `supported_models.md` and `examples` for a new model.
+- [ ] (Optional) Release notes update. If your change is user facing, please update the release notes draft in the [Google Doc](https://docs.google.com/document/d/1YyVqrgX4gHTtrstbq8oWUImOyPCKSGnJ7xtTpmXzlRs/edit?tab=t.0).
 </details>

 **BEFORE SUBMITTING, PLEASE READ <https://docs.vllm.ai/en/latest/contributing>** (anything written below this line will be removed by GitHub Actions)
@@ -262,7 +262,7 @@ pull_request_rules:
      - files~=^docker/Dockerfile.xpu
      - files~=^\\.buildkite/intel_jobs/
      - files=\.buildkite/ci_config_intel.yaml
-      - files=vllm/model_executor/layers/fused_moe/experts/xpu_moe.py
+      - files=vllm/model_executor/layers/fused_moe/xpu_fused_moe.py
      - files=vllm/model_executor/kernels/linear/mixed_precision/xpu.py
      - files=vllm/model_executor/kernels/linear/mxfp8/xpu.py
      - files=vllm/model_executor/kernels/linear/scaled_mm/xpu.py
@@ -308,7 +308,8 @@ pull_request_rules:
      - files=benchmarks/benchmark_serving_structured_output.py
      - files=benchmarks/run_structured_output_benchmark.sh
      - files=docs/features/structured_outputs.md
-      - files=^examples/features/structured_outputs/
+      - files=examples/offline_inference/structured_outputs.py
+      - files=examples/online_serving/structured_outputs/structured_outputs.py
      - files~=^tests/v1/structured_output/
      - files=tests/entrypoints/llm/test_struct_output_generate.py
      - files~=^vllm/v1/structured_output/
@@ -324,7 +325,7 @@ pull_request_rules:
    - or:
      - files~=^vllm/v1/spec_decode/
      - files~=^tests/v1/spec_decode/
-      - files=^examples/features/speculative_decoding/
+      - files~=^examples/.*(spec_decode|mlpspeculator|eagle|speculation).*\.py
      - files~=^vllm/model_executor/models/.*eagle.*\.py
      - files=vllm/model_executor/models/mlp_speculator.py
      - files~=^vllm/transformers_utils/configs/(eagle|medusa|mlp_speculator)\.py
@@ -388,7 +389,11 @@ pull_request_rules:
      - files~=^tests/entrypoints/anthropic/.*tool.*
      - files~=^vllm/tool_parsers/
      - files=docs/features/tool_calling.md
-      - files~=^examples/tool_calling/
+      - files~=^examples/tool_chat_*
+      - files=examples/offline_inference/chat_with_tools.py
+      - files=examples/online_serving/openai_chat_completion_client_with_tools_required.py
+      - files=examples/online_serving/openai_chat_completion_tool_calls_with_reasoning.py
+      - files=examples/online_serving/openai_chat_completion_client_with_tools.py
  actions:
    label:
      add:
@@ -45,7 +45,6 @@ jobs:
      - name: Smoke test vllm serve
        run: |
          # Start server in background
-          VLLM_CPU_KVCACHE_SPACE=1 \
          vllm serve Qwen/Qwen3-0.6B \
            --max-model-len=2K \
            --load-format=dummy \
@@ -62,14 +62,14 @@ jobs:
            const prAuthor = context.payload.pull_request.user.login;

            const { data: searchResults } = await github.rest.search.issuesAndPullRequests({
-              q: `repo:${owner}/${repo} type:pr is:merged author:${prAuthor}`,
+              q: `repo:${owner}/${repo} type:pr author:${prAuthor}`,
              per_page: 1,
            });

-            const mergedPRCount = searchResults.total_count;
-            console.log(`Found ${mergedPRCount} merged PRs by ${prAuthor}`);
+            const authorPRCount = searchResults.total_count;
+            console.log(`Found ${authorPRCount} PRs by ${prAuthor}`);

-            if (mergedPRCount === 0) {
+            if (authorPRCount === 1) {
              console.log(`Posting welcome comment for first-time contributor: ${prAuthor}`);
              await github.rest.issues.createComment({
                owner,
@@ -98,5 +98,5 @@ jobs:
                ].join('\n'),
              });
            } else {
-              console.log(`Skipping comment for ${prAuthor} - not a first-time contributor (${mergedPRCount} merged PRs)`);
+              console.log(`Skipping comment for ${prAuthor} - not their first PR (${authorPRCount} PRs found)`);
            }
@@ -16,7 +16,11 @@ permissions:

 jobs:
  pre-run-check:
-    if: github.event_name == 'pull_request'
+    if: >-
+      github.event_name == 'pull_request' &&
+      (github.event.action != 'labeled' ||
+       github.event.label.name == 'ready' ||
+       github.event.label.name == 'verified')
    runs-on: ubuntu-latest
    steps:
    - name: Check PR label and author merge count
@@ -45,7 +49,12 @@ jobs:

  pre-commit:
    needs: pre-run-check
-    if: always() && (needs.pre-run-check.result == 'success' || needs.pre-run-check.result == 'skipped')
+    if: >-
+      always() &&
+      (github.event.action != 'labeled' ||
+       github.event.label.name == 'ready' ||
+       github.event.label.name == 'verified') &&
+      (needs.pre-run-check.result == 'success' || needs.pre-run-check.result == 'skipped')
    runs-on: ubuntu-latest
    steps:
    - uses: actions/checkout@8e8c483db84b4bee98b60c0593521ed34d9990e8 # v6.0.1
@@ -14,7 +14,7 @@ $python_executable -m pip install -r requirements/build/cuda.txt -r requirements
 # Limit the number of parallel jobs to avoid OOM
 export MAX_JOBS=1
 # Make sure release wheels are built for the following architectures
-export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
+export TORCH_CUDA_ARCH_LIST="7.0 7.5 8.0 8.6 8.9 9.0+PTX"

 bash tools/check_repo.sh

@@ -237,7 +237,6 @@ ep_kernels_workspace/

 # Allow tracked library source folders under submodules (e.g., benchmarks/lib)
 !vllm/benchmarks/lib/
-!.buildkite/scripts/lib/

 # Generated gRPC protobuf files (compiled at build time from vllm_engine.proto)
 vllm/grpc/vllm_engine_pb2.py
@@ -34,10 +34,10 @@ install(CODE "set(CMAKE_INSTALL_LOCAL_ONLY TRUE)" ALL_COMPONENTS)
 # Supported python versions.  These versions will be searched in order, the
 # first match will be selected.  These should be kept in sync with setup.py.
 #
-set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13" "3.14")
+set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13")

 # Supported AMD GPU architectures.
-set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
+set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")

 # ROCm installation prefix. Default to /opt/rocm but allow override via
 # -DROCM_PATH=/your/rocm/path when invoking cmake.
@@ -94,15 +94,12 @@ find_package(Torch REQUIRED)
 # This check must happen after find_package(Torch) because that's when CMAKE_CUDA_COMPILER_VERSION gets defined
 if(DEFINED CMAKE_CUDA_COMPILER_VERSION AND
   CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
-   # starting from CUDA 12.9 and Blackwell (10.0), we use family-specific targets (10.0f, 12.0f, etc)
-   # to support the whole generation without specifying all sub-architectures
-   # see: https://developer.nvidia.com/blog/nvidia-blackwell-and-nvidia-cuda-12-9-introduce-family-specific-architecture-features/
-  set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;11.0;12.0")
+  set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;11.0;12.0;12.1")
 elseif(DEFINED CMAKE_CUDA_COMPILER_VERSION AND
   CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.8)
-  set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;10.1;10.3;12.0;12.1")
+  set(CUDA_SUPPORTED_ARCHS "7.0;7.2;7.5;8.0;8.6;8.7;8.9;9.0;10.0;10.1;12.0;12.1")
 else()
-  set(CUDA_SUPPORTED_ARCHS "7.0;7.5;8.0;8.6;8.7;8.9;9.0")
+  set(CUDA_SUPPORTED_ARCHS "7.0;7.2;7.5;8.0;8.6;8.7;8.9;9.0")
 endif()

 #
@@ -310,9 +307,7 @@ set(VLLM_EXT_SRC
  "csrc/torch_bindings.cpp")

 if(VLLM_GPU_LANG STREQUAL "CUDA")
-  list(APPEND VLLM_EXT_SRC
-    "csrc/minimax_reduce_rms_kernel.cu"
-    "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
+  list(APPEND VLLM_EXT_SRC "csrc/minimax_reduce_rms_kernel.cu")

  SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")

@@ -928,14 +923,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
      SRCS "${SRCS}"
      CUDA_ARCHS "${FP4_ARCHS}")
    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    # 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}")
-    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")
    message(STATUS "Building NVFP4 for archs: ${FP4_ARCHS}")
@@ -957,19 +944,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
      "csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
      "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu")
+      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu")
    set_gencode_flags_for_srcs(
      SRCS "${SRCS}"
      CUDA_ARCHS "${FP4_ARCHS}")
    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    set(NVFP4_KV_SRC "csrc/nvfp4_kv_cache_kernels.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${NVFP4_KV_SRC}"
-      CUDA_ARCHS "${FP4_ARCHS}")
-    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")
    message(STATUS "Building NVFP4 for archs: ${FP4_ARCHS}")
@@ -1053,8 +1032,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
  list(APPEND VLLM_MOE_EXT_SRC
    "csrc/moe/moe_wna16.cu"
    "csrc/moe/grouped_topk_kernels.cu"
-    "csrc/moe/router_gemm.cu"
-    "csrc/moe/topk_softplus_sqrt_kernels.cu")
+    "csrc/moe/router_gemm.cu")
 endif()

 if(VLLM_GPU_LANG STREQUAL "CUDA")
@@ -14,7 +14,7 @@ Easy, fast, and cheap LLM serving for everyone
 | <a href="https://docs.vllm.ai"><b>Documentation</b></a> | <a href="https://blog.vllm.ai/"><b>Blog</b></a> | <a href="https://arxiv.org/abs/2309.06180"><b>Paper</b></a> | <a href="https://x.com/vllm_project"><b>Twitter/X</b></a> | <a href="https://discuss.vllm.ai"><b>User Forum</b></a> | <a href="https://slack.vllm.ai"><b>Developer Slack</b></a> |
 </p>

-🔥 We have built a vLLM website to help you get started with vLLM. Please visit [vllm.ai](https://vllm.ai) to learn more.
+🔥 We have built a vllm website to help you get started with vllm. Please visit [vllm.ai](https://vllm.ai) to learn more.
 For events, please visit [vllm.ai/events](https://vllm.ai/events) to join us.

 ---
@@ -50,7 +50,7 @@ vLLM is flexible and easy to use with:
 - Efficient multi-LoRA support for dense and MoE layers
 - Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.

-vLLM seamlessly supports 200+ model architectures on Hugging Face, including:
+vLLM seamlessly supports 200+ model architectures on HuggingFace, including:

 - Decoder-only LLMs (e.g., Llama, Qwen, Gemma)
 - Mixture-of-Expert LLMs (e.g., Mixtral, DeepSeek-V3, Qwen-MoE, GPT-OSS)
@@ -404,7 +404,6 @@ def _build_attention_metadata(
        query_start_loc=q_start_gpu,
        query_start_loc_cpu=q_start_cpu,
        seq_lens=seq_lens_gpu,
-        seq_lens_cpu_upper_bound=seq_lens_cpu,
        _seq_lens_cpu=seq_lens_cpu,
        _num_computed_tokens_cpu=num_computed_tokens_cpu,
        slot_mapping=slot_mapping,
@@ -16,7 +16,7 @@ from vllm.model_executor.layers.fused_moe.all2all_utils import (
    maybe_make_prepare_finalize,
 )
 from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_config
-from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import CutlassExpertsFp8
+from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp8
 from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts, fused_topk
 from vllm.platforms import current_platform
 from vllm.utils.argparse_utils import FlexibleArgumentParser
@@ -22,7 +22,7 @@ from vllm.model_executor.layers.fused_moe.config import (
    fp8_w8a8_moe_quant_config,
    nvfp4_moe_quant_config,
 )
-from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import (
+from vllm.model_executor.layers.fused_moe.cutlass_moe import (
    CutlassExpertsFp4,
 )
 from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts, fused_topk
@@ -13,7 +13,7 @@ from vllm.model_executor.layers.fused_moe.all2all_utils import (
    maybe_make_prepare_finalize,
 )
 from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_config
-from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import CutlassExpertsFp8
+from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp8
 from vllm.model_executor.layers.fused_moe.fused_moe import (
    fused_experts,
    fused_topk,
@@ -1,324 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-# Benchmarks FP8 vs BF16 ViT attention via FlashInfer cuDNN backend.
-#
-# == Usage Examples ==
-#
-# Benchmark mode (default, FlashInfer CUDAGraph Bench)
-#   python3 benchmark_vit_fp8_attn.py
-#
-# Profile mode (PyTorch profiler, saves TensorBoard traces):
-#   python3 benchmark_vit_fp8_attn.py --profile
-#   python3 benchmark_vit_fp8_attn.py --profile --profile-output-dir ./profile_traces
-#
-# Custom seq_lens:
-#   python3 benchmark_vit_fp8_attn.py --seq-lens 4096 8192 16384
-
-from functools import partial
-
-import numpy as np
-import torch
-from torch.profiler import ProfilerActivity, profile, record_function
-
-from vllm.utils.argparse_utils import FlexibleArgumentParser
-
-# Qwen3-VL defaults
-NUM_HEADS = 16
-HEAD_DIM = 72
-DEFAULT_SEQ_LENS = [2304, 4096, 8192, 16384]
-
-
-def _setup_fp8_attention(num_heads: int, head_dim: int) -> tuple:
-    """Create FP8 and BF16 attention modules + workspace."""
-    from types import SimpleNamespace
-    from unittest.mock import patch
-
-    from vllm.config import VllmConfig, set_current_vllm_config
-    from vllm.config.multimodal import MultiModalConfig
-    from vllm.model_executor.layers.attention.mm_encoder_attention import (
-        MMEncoderAttention,
-        _get_flashinfer_workspace_buffer,
-    )
-    from vllm.v1.attention.backends.registry import AttentionBackendEnum
-
-    old_dtype = torch.get_default_dtype()
-    torch.set_default_dtype(torch.bfloat16)
-
-    backend_patch = patch(
-        "vllm.model_executor.layers.attention.mm_encoder_attention"
-        ".get_vit_attn_backend",
-        return_value=AttentionBackendEnum.FLASHINFER,
-    )
-
-    # FP8 attention
-    mm_config_fp8 = MultiModalConfig(mm_encoder_attn_dtype="fp8")
-    vllm_config_fp8 = VllmConfig()
-    vllm_config_fp8.model_config = SimpleNamespace(multimodal_config=mm_config_fp8)
-    with set_current_vllm_config(vllm_config_fp8), backend_patch:
-        attn_fp8 = MMEncoderAttention(
-            num_heads=num_heads,
-            head_size=head_dim,
-            prefix="visual.blocks.0.attn",
-        ).to("cuda")
-
-    # BF16 attention (no FP8)
-    with set_current_vllm_config(VllmConfig()), backend_patch:
-        attn_bf16 = MMEncoderAttention(
-            num_heads=num_heads,
-            head_size=head_dim,
-            prefix="visual.blocks.0.attn",
-        ).to("cuda")
-
-    torch.set_default_dtype(old_dtype)
-
-    workspace = _get_flashinfer_workspace_buffer()
-    return attn_fp8, attn_bf16, workspace
-
-
-def _build_meta(
-    seq_len: int,
-    num_heads: int,
-    head_dim: int,
-    fp8: bool,
-):
-    """Build cu_seqlens, max_seqlen, sequence_lengths."""
-    from vllm.model_executor.layers.attention.mm_encoder_attention import (
-        MMEncoderAttention,
-    )
-    from vllm.utils.math_utils import round_up
-    from vllm.v1.attention.backends.registry import AttentionBackendEnum
-
-    cu_np = np.array([0, seq_len], dtype=np.int32)
-    fp8_padded = num_heads * round_up(head_dim, 16) if fp8 else None
-
-    seq_lengths = MMEncoderAttention.maybe_compute_seq_lens(
-        AttentionBackendEnum.FLASHINFER, cu_np, torch.device("cuda")
-    )
-    max_seqlen = torch.tensor(
-        MMEncoderAttention.compute_max_seqlen(AttentionBackendEnum.FLASHINFER, cu_np),
-        dtype=torch.int32,
-    )
-    cu_seqlens = MMEncoderAttention.maybe_recompute_cu_seqlens(
-        AttentionBackendEnum.FLASHINFER,
-        cu_np,
-        num_heads * head_dim,
-        1,
-        torch.device("cuda"),
-        fp8_padded_hidden_size=fp8_padded,
-    )
-    return cu_seqlens, max_seqlen, seq_lengths
-
-
-def run_benchmark(
-    seq_lens: list[int],
-    num_heads: int,
-    head_dim: int,
-    method: str,
-):
-    """Benchmark FP8 vs BF16 attention across seq_lens.
-
-    Uses FlashInfer GPU-level timing to measure pure kernel time,
-    excluding CPU launch overhead.
-    """
-    if method == "cupti":
-        from flashinfer.testing import bench_gpu_time_with_cupti as bench_fn
-
-        bench_fn = partial(bench_fn, use_cuda_graph=True, cold_l2_cache=False)
-    elif method == "cudagraph":
-        from flashinfer.testing import (
-            bench_gpu_time_with_cudagraph as bench_fn,
-        )
-
-        bench_fn = partial(bench_fn, cold_l2_cache=False)
-    else:
-        raise ValueError(f"Invalid method: {method}")
-
-    attn_fp8, attn_bf16, workspace = _setup_fp8_attention(num_heads, head_dim)
-
-    print(f"Timing method: {method}")
-    print(f"{'seq_len':>8} {'BF16 (us)':>12} {'FP8 (us)':>12} {'Speedup':>10}")
-    print("-" * 46)
-
-    for seq_len in seq_lens:
-        torch.manual_seed(42)
-
-        q = torch.randn(
-            seq_len,
-            num_heads,
-            head_dim,
-            device="cuda",
-            dtype=torch.bfloat16,
-        )
-        k = torch.randn_like(q)
-        v = torch.randn_like(q)
-
-        cu_fp8, max_s, seq_l = _build_meta(seq_len, num_heads, head_dim, fp8=True)
-        # we can reuse cu_fp8 for cu_bf16 since q, k, and v are contiguous
-        cu_bf16 = cu_fp8.clone()
-
-        def bf16_fn(q=q, k=k, v=v, cu=cu_bf16, ms=max_s, sl=seq_l):
-            attn_bf16._forward_flashinfer(q, k, v, cu, ms, sl)
-
-        def fp8_fn(q=q, k=k, v=v, cu=cu_fp8, ms=max_s, sl=seq_l):
-            attn_fp8._forward_flashinfer(q, k, v, cu, ms, sl)
-
-        # bench_fn returns List[float] of per-iteration times in ms
-        bf16_times = bench_fn(bf16_fn)
-        fp8_times = bench_fn(fp8_fn)
-
-        bf16_us = np.median(bf16_times) * 1e3  # ms -> us
-        fp8_us = np.median(fp8_times) * 1e3
-        speedup = bf16_us / fp8_us if fp8_us > 0 else float("inf")
-
-        print(f"{seq_len:>8} {bf16_us:>12.1f} {fp8_us:>12.1f} {speedup:>9.2f}x")
-
-
-def _make_trace_handler(output_dir: str, worker_name: str, label: str):
-    """Create a trace handler that saves to TensorBoard and prints summary."""
-
-    def handler(prof):
-        torch.profiler.tensorboard_trace_handler(output_dir, worker_name)(prof)
-        print(f"\n{'=' * 80}")
-        print(label)
-        print(f"{'=' * 80}")
-        print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))
-
-    return handler
-
-
-def run_profile(
-    seq_len: int,
-    num_heads: int,
-    head_dim: int,
-    warmup: int,
-    output_dir: str,
-):
-    """Profile FP8 vs BF16 attention with PyTorch profiler."""
-    attn_fp8, attn_bf16, workspace = _setup_fp8_attention(num_heads, head_dim)
-
-    torch.manual_seed(42)
-    q = torch.randn(
-        seq_len,
-        num_heads,
-        head_dim,
-        device="cuda",
-        dtype=torch.bfloat16,
-    )
-    k = torch.randn_like(q)
-    v = torch.randn_like(q)
-
-    cu_fp8, max_s, seq_l = _build_meta(seq_len, num_heads, head_dim, fp8=True)
-    # we can reuse cu_fp8 for cu_bf16 since q, k, and v are contiguous
-    cu_bf16 = cu_fp8.clone()
-
-    sched = torch.profiler.schedule(wait=0, warmup=warmup, active=1)
-
-    # Profile BF16 (warmup handled by profiler schedule)
-    with profile(
-        activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
-        schedule=sched,
-        on_trace_ready=_make_trace_handler(
-            output_dir,
-            f"bf16_h{head_dim}_s{seq_len}",
-            f"BF16 Attention (seq_len={seq_len}, heads={num_heads}, "
-            f"head_dim={head_dim})",
-        ),
-    ) as prof_bf16:
-        for _ in range(warmup + 1):
-            with record_function("bf16_attention"):
-                attn_bf16._forward_flashinfer(
-                    q.clone(), k.clone(), v.clone(), cu_bf16, max_s, seq_l
-                )
-                torch.accelerator.synchronize()
-            prof_bf16.step()
-
-    # Profile FP8 (warmup handled by profiler schedule)
-    with profile(
-        activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
-        schedule=sched,
-        on_trace_ready=_make_trace_handler(
-            output_dir,
-            f"fp8_h{head_dim}_s{seq_len}",
-            f"FP8 Attention (seq_len={seq_len}, heads={num_heads}, "
-            f"head_dim={head_dim})",
-        ),
-    ) as prof_fp8:
-        for _ in range(warmup + 1):
-            with record_function("fp8_attention"):
-                attn_fp8._forward_flashinfer(
-                    q.clone(), k.clone(), v.clone(), cu_fp8, max_s, seq_l
-                )
-                torch.accelerator.synchronize()
-            prof_fp8.step()
-
-    print(f"\nTensorBoard traces saved to: {output_dir}")
-    print(f"View with: tensorboard --logdir={output_dir}")
-
-
-if __name__ == "__main__":
-    parser = FlexibleArgumentParser(description="Benchmark FP8 vs BF16 ViT attention.")
-    parser.add_argument(
-        "--seq-lens",
-        type=int,
-        nargs="+",
-        default=DEFAULT_SEQ_LENS,
-        help="Sequence lengths to benchmark",
-    )
-    parser.add_argument(
-        "--num-heads",
-        type=int,
-        default=NUM_HEADS,
-    )
-    parser.add_argument(
-        "--head-dim",
-        type=int,
-        default=HEAD_DIM,
-    )
-    parser.add_argument(
-        "--method",
-        choices=["cupti", "cudagraph"],
-        default="cudagraph",
-        help="GPU timing method: cupti (CUPTI kernel timing) or "
-        "cudagraph (CUDA graph capture/replay). Default: cudagraph",
-    )
-    parser.add_argument(
-        "--warmup",
-        type=int,
-        default=10,
-        help="Warmup iterations (profile mode only)",
-    )
-    parser.add_argument(
-        "--profile",
-        action="store_true",
-        help="Run PyTorch profiler instead of benchmark",
-    )
-    parser.add_argument(
-        "--profile-seq-len",
-        type=int,
-        default=8192,
-        help="Sequence length for profiling (default: 8192)",
-    )
-    parser.add_argument(
-        "--profile-output-dir",
-        type=str,
-        default="./profile_traces",
-        help="Output directory for TensorBoard traces (default: ./profile_traces)",
-    )
-    args = parser.parse_args()
-
-    if args.profile:
-        run_profile(
-            args.profile_seq_len,
-            args.num_heads,
-            args.head_dim,
-            args.warmup,
-            args.profile_output_dir,
-        )
-    else:
-        run_benchmark(
-            args.seq_lens,
-            args.num_heads,
-            args.head_dim,
-            args.method,
-        )
@@ -1,378 +0,0 @@
-#!/usr/bin/env python3
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Generic benchmark harness for vLLM IR ops.
-
-Usage:
-    python benchmarks/kernels/ir/bench_ir_ops.py
-    python benchmarks/kernels/ir/bench_ir_ops.py --ops rms_norm
-    python benchmarks/kernels/ir/bench_ir_ops.py --ops rms_norm,silu_mul
-    python benchmarks/kernels/ir/bench_ir_ops.py --no-cuda-graph
-    python benchmarks/kernels/ir/bench_ir_ops.py --ops rms_norm --save-path ./results/
-"""
-
-import argparse
-import contextlib
-import csv
-import dataclasses
-import datetime
-import math
-import os
-import subprocess
-import sys
-import tempfile
-
-# Ensure repo root is on sys.path so `benchmarks` is importable as a package.
-_REPO_ROOT = os.path.abspath(os.path.join(os.path.dirname(__file__), "../../.."))
-if _REPO_ROOT not in sys.path:
-    sys.path.insert(0, _REPO_ROOT)
-
-# Suppress noisy C++ warnings from vllm kernel registration (written to fd 2
-# directly by the dynamic linker, so Python-level sys.stderr redirect won't
-# catch them).
-_saved_fd = os.dup(2)
-try:
-    with open(os.devnull, "w") as _devnull:
-        os.dup2(_devnull.fileno(), 2)
-        import torch
-
-        import vllm.kernels  # noqa: E402, F401
-finally:
-    os.dup2(_saved_fd, 2)
-    os.close(_saved_fd)
-
-from tqdm import tqdm  # noqa: E402
-
-from benchmarks.kernels.ir.shapes import SHAPE_CONFIGS  # noqa: E402  # isort: skip
-from vllm.ir.op import IrOp  # noqa: E402
-from vllm.platforms import current_platform  # noqa: E402
-from vllm.triton_utils import triton  # noqa: E402
-
-
-@dataclasses.dataclass(frozen=True)
-class BenchConfig:
-    use_cuda_graph: bool = True
-    warmup: int = 25
-    rep: int = 100
-
-
-def _pkg_version(name: str) -> str:
-    from importlib.metadata import PackageNotFoundError, version
-
-    with contextlib.suppress(PackageNotFoundError):
-        return version(name)
-    return "not installed"
-
-
-_METADATA_LABELS = {
-    "timestamp": "Timestamp",
-    "git_commit": "Git commit",
-    "vllm": "vLLM",
-    "pytorch": "PyTorch",
-    "cuda_runtime": "CUDA runtime",
-    "triton": "Triton",
-    "cutlass": "CUTLASS",
-    "helion": "Helion",
-    "device": "Device",
-    "bench_mode": "Bench mode",
-    "warmup": "Warmup",
-    "rep": "Repetitions",
-}
-
-
-def collect_env_metadata(cfg: BenchConfig) -> dict[str, str]:
-    from vllm.collect_env import get_env_info
-
-    env = get_env_info()
-
-    git_sha = "unknown"
-    with contextlib.suppress(subprocess.CalledProcessError, FileNotFoundError):
-        git_sha = (
-            subprocess.check_output(
-                ["git", "rev-parse", "--short", "HEAD"], stderr=subprocess.DEVNULL
-            )
-            .decode()
-            .strip()
-        )
-
-    device_name = current_platform.get_device_name()
-
-    warmup_note = " ms" if not cfg.use_cuda_graph else " ms (ignored)"
-    rep_note = " replays" if cfg.use_cuda_graph else " ms"
-
-    return {
-        "timestamp": datetime.datetime.now().strftime("%Y-%m-%d %H:%M:%S"),
-        "git_commit": git_sha,
-        "vllm": str(env.vllm_version),
-        "pytorch": str(env.torch_version),
-        "cuda_runtime": str(env.cuda_runtime_version),
-        "triton": triton.__version__,
-        "cutlass": _pkg_version("nvidia-cutlass-dsl"),
-        "helion": _pkg_version("helion"),
-        "device": device_name,
-        "bench_mode": "cuda_graph" if cfg.use_cuda_graph else "eager",
-        "warmup": f"{cfg.warmup}{warmup_note}",
-        "rep": f"{cfg.rep}{rep_note}",
-    }
-
-
-def print_metadata(metadata: dict[str, str]):
-    print("=" * 60)
-    for key, val in metadata.items():
-        print(f"{_METADATA_LABELS.get(key, key) + ':':<16}{val}")
-    print("=" * 60)
-
-
-def _clone_args(args: tuple) -> tuple:
-    return tuple(a.clone() if isinstance(a, torch.Tensor) else a for a in args)
-
-
-# TODO(gmagogsfm): When the `maybe_inplace` PR lands, ops marked as
-# inplace=True will mutate bench_args across iterations. Both CUDA graph
-# and eager modes will accumulate drift from repeated in-place mutation.
-# We need to re-clone inputs per iteration for inplace ops.
-def _bench_one(fn, args, cfg: BenchConfig) -> float:
-    bench_args = _clone_args(args)
-    bench_fn = lambda: fn(*bench_args)
-
-    if cfg.use_cuda_graph:
-        ms = triton.testing.do_bench_cudagraph(bench_fn, rep=cfg.rep, quantiles=[0.5])
-    else:
-        ms = triton.testing.do_bench(
-            bench_fn, warmup=cfg.warmup, rep=cfg.rep, quantiles=[0.5]
-        )
-    return ms * 1000
-
-
-# TODO(gmagogsfm): Once compiled native implementation lands (#38775),
-# the benchmark baseline should be the compiled native (what vLLM runs by
-# default) rather than the uncompiled native implementation.
-def collect_timings(
-    op: IrOp, shape_configs: list[dict], cfg: BenchConfig
-) -> tuple[list[str], list[str], dict[str, dict[str, float]]]:
-    def fmt(v) -> str:
-        return str(v).split(".")[-1] if isinstance(v, torch.dtype) else str(v)
-
-    case_names = [
-        "_".join(f"{k}={fmt(v)}" for k, v in kwargs.items()) for kwargs in shape_configs
-    ]
-    providers = [n for n, impl in op.impls.items() if impl.supported]
-
-    results: dict[str, dict[str, float]] = {c: {} for c in case_names}
-    for provider in providers:
-        impl = op.impls[provider]
-        desc = f"{op.name} / {provider}"
-        for case_name, kwargs in tqdm(
-            zip(case_names, shape_configs),
-            desc=desc,
-            total=len(case_names),
-            unit=" cases",
-        ):
-            args = op.generate_inputs(**kwargs)
-            if impl.supports_args(*args):
-                results[case_name][provider] = _bench_one(impl.impl_fn, args, cfg)
-            else:
-                results[case_name][provider] = float("nan")
-
-    return case_names, providers, results
-
-
-def analyze_results(
-    op_name: str,
-    case_names: list[str],
-    providers: list[str],
-    results: dict[str, dict[str, float]],
-) -> tuple[list[dict[str, str]], list[dict[str, str]], list[str]]:
-    native_col = "native"
-    non_native = [p for p in providers if p != native_col]
-
-    header_cols = ["case"]
-    for p in providers:
-        header_cols.append(f"{p} (us)")
-    for p in non_native:
-        header_cols.append(f"{p} speedup")
-
-    detail_rows: list[dict[str, str]] = []
-    speedup_data: dict[str, list[tuple[float, str]]] = {p: [] for p in non_native}
-
-    for case_name in case_names:
-        timings = results[case_name]
-        row: dict[str, str] = {"case": case_name}
-
-        for p in providers:
-            val = timings.get(p, float("nan"))
-            row[f"{p} (us)"] = f"{val:.2f}" if not math.isnan(val) else "n/a"
-
-        native_us = timings.get(native_col, float("nan"))
-        for p in non_native:
-            p_us = timings.get(p, float("nan"))
-            if not math.isnan(native_us) and not math.isnan(p_us) and p_us > 0:
-                speedup = native_us / p_us
-                row[f"{p} speedup"] = f"{speedup:.2f}x"
-                speedup_data[p].append((speedup, case_name))
-            else:
-                row[f"{p} speedup"] = "n/a"
-
-        detail_rows.append(row)
-
-    summary_rows: list[dict[str, str]] = []
-    for p in non_native:
-        entries = speedup_data[p]
-        if not entries:
-            continue
-        speedups = [s for s, _ in entries]
-        geomean = math.exp(sum(math.log(s) for s in speedups) / len(speedups))
-        best_val, best_case = max(entries)
-        worst_val, worst_case = min(entries)
-        wins = sum(1 for s in speedups if s > 1.0)
-        losses = sum(1 for s in speedups if s < 1.0)
-        total = len(speedups)
-
-        print(f"\n{p} vs native ({wins}/{total} faster, {losses}/{total} slower):")
-        print(f"  geomean speedup: {geomean:.2f}x")
-        print(f"  best:            {best_val:.2f}x  ({best_case})")
-        print(f"  worst:           {worst_val:.2f}x  ({worst_case})")
-
-        summary_rows.append(
-            {
-                "op": op_name,
-                "provider": p,
-                "geomean_speedup": f"{geomean:.2f}",
-                "best_speedup": f"{best_val:.2f}",
-                "best_case": best_case,
-                "worst_speedup": f"{worst_val:.2f}",
-                "worst_case": worst_case,
-                "wins": str(wins),
-                "losses": str(losses),
-                "total": str(total),
-            }
-        )
-
-    return detail_rows, summary_rows, header_cols
-
-
-def write_csv(path: str, rows: list[dict[str, str]], fieldnames: list[str]):
-    with open(path, "w", newline="") as f:
-        writer = csv.DictWriter(f, fieldnames=fieldnames)
-        writer.writeheader()
-        writer.writerows(rows)
-
-
-def save_results(
-    save_dir: str,
-    op_name: str,
-    detail_rows: list[dict[str, str]],
-    header_cols: list[str],
-    all_summary_rows: list[dict[str, str]],
-    metadata: dict[str, str],
-):
-    write_csv(
-        os.path.join(save_dir, f"{op_name}_detail.csv"),
-        detail_rows,
-        header_cols,
-    )
-    if all_summary_rows:
-        write_csv(
-            os.path.join(save_dir, "summary.csv"),
-            all_summary_rows,
-            list(all_summary_rows[0].keys()),
-        )
-    write_csv(
-        os.path.join(save_dir, "metadata.csv"),
-        [metadata],
-        list(metadata.keys()),
-    )
-
-
-def parse_args():
-    parser = argparse.ArgumentParser(description="Benchmark vLLM IR ops")
-    parser.add_argument(
-        "--ops",
-        type=str,
-        default=None,
-        help="Comma-separated list of op names to benchmark (substring match)",
-    )
-    parser.add_argument(
-        "--no-cuda-graph",
-        action="store_true",
-        help="Disable CUDA graph; use do_bench with L2 cache flushing instead",
-    )
-    parser.add_argument(
-        "--warmup",
-        type=int,
-        default=25,
-        help="Warmup time in ms (do_bench) or ignored with CUDA graph (default: 25)",
-    )
-    parser.add_argument(
-        "--rep",
-        type=int,
-        default=100,
-        help="Repetition time in ms (do_bench) or number of graph replays "
-        "(do_bench_cudagraph) (default: 100)",
-    )
-    parser.add_argument(
-        "--save-path",
-        type=str,
-        default=None,
-        help="Directory to save results (default: auto-created temp dir)",
-    )
-    return parser.parse_args()
-
-
-def main():
-    args = parse_args()
-    cfg = BenchConfig(
-        use_cuda_graph=not args.no_cuda_graph,
-        warmup=args.warmup,
-        rep=args.rep,
-    )
-
-    torch.set_default_device(current_platform.device_type)
-
-    metadata = collect_env_metadata(cfg)
-    print_metadata(metadata)
-
-    timestamp = datetime.datetime.now().strftime("%Y%m%d_%H%M%S")
-    save_dir = args.save_path or os.path.join(
-        tempfile.gettempdir(), f"vllm_ir_bench_{timestamp}"
-    )
-    os.makedirs(save_dir, exist_ok=True)
-
-    op_filters = [f.strip() for f in args.ops.split(",")] if args.ops else None
-    all_summary_rows: list[dict[str, str]] = []
-
-    for op in IrOp.registry.values():
-        if op_filters and not any(f in op.name for f in op_filters):
-            continue
-        if not op.has_input_generator:
-            print(f"Skipping op '{op.name}': no input generator registered")
-            continue
-        if op.name not in SHAPE_CONFIGS:
-            raise RuntimeError(
-                f"No benchmark shape config for op '{op.name}'. "
-                f"Add it to benchmarks/kernels/ir/shapes.py"
-            )
-
-        case_names, providers, results = collect_timings(
-            op, SHAPE_CONFIGS[op.name], cfg
-        )
-        detail_rows, summary_rows, header_cols = analyze_results(
-            op.name, case_names, providers, results
-        )
-        all_summary_rows.extend(summary_rows)
-
-        save_results(
-            save_dir,
-            op.name,
-            detail_rows,
-            header_cols,
-            all_summary_rows,
-            metadata,
-        )
-
-    print(f"\nResults saved to: {save_dir}")
-
-
-if __name__ == "__main__":
-    main()
@@ -1,29 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Shape configurations for IR op benchmarks.
-"""
-
-import torch
-
-NUM_TOKENS = [1, 2, 4, 16, 64, 256, 1024, 4096, 16384]
-COMMON_HIDDEN_SIZES = [
-    2048,  # Llama 3.2 1B, Qwen 3 MoE 30B-A3B, Gemma 3n
-    3072,  # Gemma 7B/9B
-    4096,  # Llama 3 8B, Qwen 3 8B, Mistral 7B
-    5120,  # Llama 4 Scout 17B-16E
-    7168,  # DeepSeek V3
-    8192,  # Llama 3 70B
-    16384,  # Llama 3 405B
-]
-
-# Each entry maps an op name to a list of kwarg dicts that will be passed
-# to that op's registered input generator via op.generate_inputs(**kwargs).
-SHAPE_CONFIGS: dict[str, list[dict]] = {
-    "rms_norm": [
-        {"num_tokens": n, "hidden_size": d, "dtype": dtype}
-        for dtype in [torch.float16, torch.bfloat16, torch.float32]
-        for d in COMMON_HIDDEN_SIZES
-        for n in NUM_TOKENS
-    ],
-}
@@ -217,7 +217,6 @@ async def send_request(
    min_tokens: int | None = None,
    max_tokens: int | None = None,
    timeout_sec: int = 120,
-    conversation_id: str | None = None,
 ) -> ServerResponse:
    payload = {
        "model": model,
@@ -226,9 +225,6 @@ async def send_request(
        "temperature": 0.0,
    }

-    if conversation_id is not None:
-        payload["conversation_id"] = conversation_id
-
    if stream:
        payload["stream"] = True
        payload["stream_options"] = {"include_usage": False}
@@ -423,7 +419,6 @@ async def send_turn(
        min_tokens,
        max_tokens,
        req_args.timeout_sec,
-        conversation_id=conv_id,
    )

    if response.valid is False:
@@ -30,21 +30,6 @@ else()
    list(APPEND CXX_COMPILE_FLAGS
        "-fopenmp"
        "-DVLLM_CPU_EXTENSION")
-
-    # locate PyTorch's libgomp (e.g. site-packages/torch.libs/libgomp-947d5fa1.so.1.0.0)
-    # and create a local shim dir with it
-    vllm_prepare_torch_gomp_shim(VLLM_TORCH_GOMP_SHIM_DIR)
-
-    find_library(OPEN_MP
-        NAMES gomp
-        PATHS ${VLLM_TORCH_GOMP_SHIM_DIR}
-        NO_DEFAULT_PATH
-        REQUIRED
-    )
-    # Set LD_LIBRARY_PATH to include the shim dir at build time to use the same libgomp as PyTorch
-    if (OPEN_MP)
-        set(ENV{LD_LIBRARY_PATH} "${VLLM_TORCH_GOMP_SHIM_DIR}:$ENV{LD_LIBRARY_PATH}")
-    endif()
 endif()

 if (NOT MACOSX_FOUND)
@@ -161,49 +146,16 @@ elseif (S390_FOUND)
        "-mtune=native")
 elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
    message(STATUS "RISC-V detected")
-    # VLLM_RVV_VLEN selects the target VLEN. Auto-detected from /proc/cpuinfo
-    # by default; override with -DVLLM_RVV_VLEN=128 or -DVLLM_RVV_VLEN=256.
-    if(NOT DEFINED VLLM_RVV_VLEN)
-        # Auto-detect: find the largest zvl<N>b in /proc/cpuinfo isa line.
-        if(EXISTS /proc/cpuinfo)
-            file(READ /proc/cpuinfo _cpuinfo)
-            set(_best 0)
-            foreach(_n IN ITEMS 128 256 512 1024)
-                if(_cpuinfo MATCHES "zvl${_n}b")
-                    set(_best ${_n})
-                endif()
-            endforeach()
-            if(_best GREATER 0)
-                set(VLLM_RVV_VLEN ${_best})
-            endif()
-        endif()
-        # If auto-detect failed (no /proc/cpuinfo or no zvl<N>b reported)
-        # but the compiler supports RVV, require explicit specification.
-        if(NOT DEFINED VLLM_RVV_VLEN AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
-            message(FATAL_ERROR
-                "RISC-V RVV is available but VLEN could not be auto-detected. "
-                "Please specify VLEN explicitly:\n"
-                "  -DVLLM_RVV_VLEN=128   (for VLEN=128 hardware)\n"
-                "  -DVLLM_RVV_VLEN=256   (for VLEN=256 hardware, e.g. Spacemit X100)\n"
-                "  -DVLLM_RVV_VLEN=0     (force scalar, no RVV)")
-        endif()
-    endif()
-    if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
-        message(STATUS "RISC-V RVV VLEN=${VLLM_RVV_VLEN}")
-        if(RVV_BF16_FOUND)
-            message(STATUS "BF16 extension detected")
-            set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
-            add_compile_definitions(RISCV_BF16_SUPPORT)
-        elseif(RVV_FP16_FOUND)
-            message(WARNING "BF16 functionality is not available")
-            set(MARCH_FLAGS -march=rv64gcv_zvfh_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
-        else()
-            message(STATUS "compile riscv with scalar (no FP16/BF16)")
-            set(MARCH_FLAGS -march=rv64gc)
-        endif()
+    if(RVV_BF16_FOUND)
+        message(STATUS "BF16 extension detected")
+        set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl128b -mrvv-vector-bits=zvl -mabi=lp64d)
+        add_compile_definitions(RISCV_BF16_SUPPORT)
+    elseif (RVV_FP16_FOUND)
+        message(WARNING "BF16 functionality is not available")
+        set(MARCH_FLAGS -march=rv64gcv_zvfh_zvl128b -mrvv-vector-bits=zvl -mabi=lp64d)
    else()
        message(STATUS "compile riscv with scalar")
-        set(MARCH_FLAGS -march=rv64gc)
+        list(APPEND CXX_COMPILE_FLAGS "-march=rv64gc")
    endif()
    list(APPEND CXX_COMPILE_FLAGS ${MARCH_FLAGS})
 else()
@@ -223,6 +175,20 @@ if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND
        if(NOT NPROC)
            set(NPROC 4)
        endif()
+        # locate PyTorch's libgomp (e.g. site-packages/torch.libs/libgomp-947d5fa1.so.1.0.0)
+        # and create a local shim dir with it
+        vllm_prepare_torch_gomp_shim(VLLM_TORCH_GOMP_SHIM_DIR)
+
+        find_library(OPEN_MP
+            NAMES gomp
+            PATHS ${VLLM_TORCH_GOMP_SHIM_DIR}
+            NO_DEFAULT_PATH
+            REQUIRED
+        )
+        # Set LD_LIBRARY_PATH to include the shim dir at build time to use the same libgomp as PyTorch
+        if (OPEN_MP)
+            set(ENV{LD_LIBRARY_PATH} "${VLLM_TORCH_GOMP_SHIM_DIR}:$ENV{LD_LIBRARY_PATH}")
+        endif()

        # Fetch and populate ACL
        if(DEFINED ENV{ACL_ROOT_DIR} AND IS_DIRECTORY "$ENV{ACL_ROOT_DIR}")
@@ -394,7 +360,6 @@ set(VLLM_EXT_SRC
 if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
    set(VLLM_EXT_SRC
        "csrc/cpu/shm.cpp"
-        "csrc/cpu/activation_lut_bf16.cpp"
        ${VLLM_EXT_SRC})
 endif()

@@ -20,7 +20,7 @@ else()
  FetchContent_Declare(
    deepgemm
    GIT_REPOSITORY https://github.com/deepseek-ai/DeepGEMM.git
-    GIT_TAG 891d57b4db1071624b5c8fa0d1e51cb317fa709f
+    GIT_TAG 477618cd51baffca09c4b0b87e97c03fe827ef03
    GIT_SUBMODULES "third-party/cutlass" "third-party/fmt"
    GIT_PROGRESS TRUE
    CONFIGURE_COMMAND ""
@@ -120,11 +120,6 @@ if(DEEPGEMM_ARCHS)
    COMPONENT _deep_gemm_C
    FILES_MATCHING PATTERN "*.py")

-  install(DIRECTORY "${deepgemm_SOURCE_DIR}/deep_gemm/mega/"
-    DESTINATION vllm/third_party/deep_gemm/mega
-    COMPONENT _deep_gemm_C
-    FILES_MATCHING PATTERN "*.py")
-
  # Generate envs.py (normally generated by DeepGEMM's setup.py build step)
  file(WRITE "${CMAKE_CURRENT_BINARY_DIR}/deep_gemm_envs.py"
    "# Pre-installed environment variables\npersistent_envs = dict()\n")
@@ -19,7 +19,7 @@ else()
  FetchContent_Declare(
        flashmla
        GIT_REPOSITORY https://github.com/vllm-project/FlashMLA
-        GIT_TAG a6ec2ba7bd0a7dff98b3f4d3e6b52b159c48d78b
+        GIT_TAG 692917b1cda61b93ac9ee2d846ec54e75afe87b1
        GIT_PROGRESS TRUE
        CONFIGURE_COMMAND ""
        BUILD_COMMAND ""
@@ -11,74 +11,29 @@
 namespace vllm {

 template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&),
-          bool act_first, bool HAS_CLAMP>
+          bool act_first>
 __device__ __forceinline__ scalar_t compute(const scalar_t& x,
-                                            const scalar_t& y,
-                                            const float limit) {
-  if constexpr (act_first) {
-    scalar_t gate = x;
-    scalar_t up = y;
-    if constexpr (HAS_CLAMP) {
-      gate = (scalar_t)fminf((float)gate, limit);
-      up = (scalar_t)fmaxf(fminf((float)up, limit), -limit);
-    }
-    return ACT_FN(gate) * up;
-  } else {
-    scalar_t gate = x;
-    scalar_t up = y;
-    if constexpr (HAS_CLAMP) {
-      gate = (scalar_t)fmaxf(fminf((float)gate, limit), -limit);
-      up = (scalar_t)fminf((float)up, limit);
-    }
-    return gate * ACT_FN(up);
-  }
+                                            const scalar_t& y) {
+  return act_first ? ACT_FN(x) * y : x * ACT_FN(y);
 }

 template <typename packed_t, packed_t (*PACKED_ACT_FN)(const packed_t&),
-          bool act_first, bool HAS_CLAMP>
+          bool act_first>
 __device__ __forceinline__ packed_t packed_compute(const packed_t& x,
-                                                   const packed_t& y,
-                                                   const float limit) {
-  if constexpr (act_first) {
-    packed_t gate = x;
-    packed_t up = y;
-    if constexpr (HAS_CLAMP) {
-      float2 g = cast_to_float2(gate);
-      float2 u = cast_to_float2(up);
-      g.x = fminf(g.x, limit);
-      g.y = fminf(g.y, limit);
-      u.x = fmaxf(fminf(u.x, limit), -limit);
-      u.y = fmaxf(fminf(u.y, limit), -limit);
-      gate = cast_to_packed<packed_t>(g);
-      up = cast_to_packed<packed_t>(u);
-    }
-    return packed_mul(PACKED_ACT_FN(gate), up);
-  } else {
-    packed_t gate = x;
-    packed_t up = y;
-    if constexpr (HAS_CLAMP) {
-      float2 g = cast_to_float2(gate);
-      float2 u = cast_to_float2(up);
-      g.x = fmaxf(fminf(g.x, limit), -limit);
-      g.y = fmaxf(fminf(g.y, limit), -limit);
-      u.x = fminf(u.x, limit);
-      u.y = fminf(u.y, limit);
-      gate = cast_to_packed<packed_t>(g);
-      up = cast_to_packed<packed_t>(u);
-    }
-    return packed_mul(gate, PACKED_ACT_FN(up));
-  }
+                                                   const packed_t& y) {
+  return act_first ? packed_mul(PACKED_ACT_FN(x), y)
+                   : packed_mul(x, PACKED_ACT_FN(y));
 }

 // Activation and gating kernel template.
 template <typename scalar_t, typename packed_t,
          scalar_t (*ACT_FN)(const scalar_t&),
          packed_t (*PACKED_ACT_FN)(const packed_t&), bool act_first,
-          bool use_vec, bool HAS_CLAMP, bool use_256b = false>
+          bool use_vec, bool use_256b = false>
 __global__ void act_and_mul_kernel(
    scalar_t* __restrict__ out,          // [..., d]
    const scalar_t* __restrict__ input,  // [..., 2, d]
-    const int d, const float limit) {
+    const int d) {
  const scalar_t* x_ptr = input + blockIdx.x * 2 * d;
  const scalar_t* y_ptr = x_ptr + d;
  scalar_t* out_ptr = out + blockIdx.x * d;
@@ -103,9 +58,8 @@ __global__ void act_and_mul_kernel(
      }
 #pragma unroll
      for (int j = 0; j < pvec_t::NUM_ELTS; j++) {
-        x.elts[j] =
-            packed_compute<packed_t, PACKED_ACT_FN, act_first, HAS_CLAMP>(
-                x.elts[j], y.elts[j], limit);
+        x.elts[j] = packed_compute<packed_t, PACKED_ACT_FN, act_first>(
+            x.elts[j], y.elts[j]);
      }
      if constexpr (use_256b) {
        st256(x, &out_vec[i]);
@@ -118,8 +72,7 @@ __global__ void act_and_mul_kernel(
    for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
      const scalar_t x = VLLM_LDG(&x_ptr[idx]);
      const scalar_t y = VLLM_LDG(&y_ptr[idx]);
-      out_ptr[idx] =
-          compute<scalar_t, ACT_FN, act_first, HAS_CLAMP>(x, y, limit);
+      out_ptr[idx] = compute<scalar_t, ACT_FN, act_first>(x, y);
    }
  }
 }
@@ -198,11 +151,8 @@ packed_gelu_tanh_kernel(const packed_t& val) {

 // Launch activation and gating kernel.
 // Use ACT_FIRST (bool) indicating whether to apply the activation function
-// first. HAS_CLAMP (bool) enables pre-activation clamping: gate input is
-// clamped (max only) and up input is clamped (both sides) before the
-// activation function is applied.
-#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, PACKED_KERNEL, ACT_FIRST,        \
-                                      HAS_CLAMP, LIMIT)                        \
+// first.
+#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, PACKED_KERNEL, ACT_FIRST)        \
  auto dtype = input.scalar_type();                                            \
  int d = input.size(-1) / 2;                                                  \
  int64_t num_tokens = input.numel() / input.size(-1);                         \
@@ -227,8 +177,8 @@ packed_gelu_tanh_kernel(const packed_t& val) {
            scalar_t, typename vllm::PackedTypeConverter<scalar_t>::Type,      \
            KERNEL<scalar_t>,                                                  \
            PACKED_KERNEL<typename vllm::PackedTypeConverter<scalar_t>::Type>, \
-            ACT_FIRST, true, HAS_CLAMP, true><<<grid, block, 0, stream>>>(     \
-            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, LIMIT);   \
+            ACT_FIRST, true, true><<<grid, block, 0, stream>>>(                \
+            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d);          \
      });                                                                      \
    } else {                                                                   \
      VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {          \
@@ -236,8 +186,8 @@ packed_gelu_tanh_kernel(const packed_t& val) {
            scalar_t, typename vllm::PackedTypeConverter<scalar_t>::Type,      \
            KERNEL<scalar_t>,                                                  \
            PACKED_KERNEL<typename vllm::PackedTypeConverter<scalar_t>::Type>, \
-            ACT_FIRST, true, HAS_CLAMP, false><<<grid, block, 0, stream>>>(    \
-            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, LIMIT);   \
+            ACT_FIRST, true, false><<<grid, block, 0, stream>>>(               \
+            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d);          \
      });                                                                      \
    }                                                                          \
  } else {                                                                     \
@@ -247,8 +197,8 @@ packed_gelu_tanh_kernel(const packed_t& val) {
          scalar_t, typename vllm::PackedTypeConverter<scalar_t>::Type,        \
          KERNEL<scalar_t>,                                                    \
          PACKED_KERNEL<typename vllm::PackedTypeConverter<scalar_t>::Type>,   \
-          ACT_FIRST, false, HAS_CLAMP><<<grid, block, 0, stream>>>(            \
-          out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, LIMIT);     \
+          ACT_FIRST, false><<<grid, block, 0, stream>>>(                       \
+          out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d);            \
    });                                                                        \
  }

@@ -256,14 +206,7 @@ void silu_and_mul(torch::Tensor& out,    // [..., d]
                  torch::Tensor& input)  // [..., 2 * d]
 {
  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
-                                true, false, 0.0f);
-}
-
-void silu_and_mul_clamp(torch::Tensor& out,    // [..., d]
-                        torch::Tensor& input,  // [..., 2 * d]
-                        double limit) {
-  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
-                                true, true, (float)limit);
+                                true);
 }

 void mul_and_silu(torch::Tensor& out,    // [..., d]
@@ -272,21 +215,21 @@ void mul_and_silu(torch::Tensor& out,    // [..., d]
  // The difference between mul_and_silu and silu_and_mul is that mul_and_silu
  // applies the silu to the latter half of the input.
  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
-                                false, false, 0.0f);
+                                false);
 }

 void gelu_and_mul(torch::Tensor& out,    // [..., d]
                  torch::Tensor& input)  // [..., 2 * d]
 {
  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_kernel, vllm::packed_gelu_kernel,
-                                true, false, 0.0f);
+                                true);
 }

 void gelu_tanh_and_mul(torch::Tensor& out,    // [..., d]
                       torch::Tensor& input)  // [..., 2 * d]
 {
-  LAUNCH_ACTIVATION_GATE_KERNEL(
-      vllm::gelu_tanh_kernel, vllm::packed_gelu_tanh_kernel, true, false, 0.0f);
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_tanh_kernel,
+                                vllm::packed_gelu_tanh_kernel, true);
 }

 namespace vllm {
@@ -599,11 +599,6 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
  const int head_idx = (blockIdx.y * blockDim.x + threadIdx.x) * VEC_SIZE;
  // Find batch index within a block
  __shared__ int batch_idx[BLOCK_Y_SIZE];
-  if (threadIdx.x == 0) {
-    batch_idx[threadIdx.y] = -1;
-  }
-  __syncthreads();
-
  for (int iter = 0; iter < cuda_utils::ceil_div(batch_size, int(blockDim.x));
       iter++) {
    int tid = iter * blockDim.x + threadIdx.x;
@@ -616,18 +611,16 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
    }
  }

-  __syncthreads();
+#ifndef USE_ROCM
+  __syncwarp();
+#endif

-  // num_tokens may be an allocation upper bound when Python avoids a D2H sync.
-  // Only tokens covered by the exact device-side cu_seq_lens are valid to
-  // gather.
-  const int batch = batch_idx[threadIdx.y];
-  if (head_idx >= head_dim || token_idx >= num_tokens || batch < 0) {
+  if (head_idx >= head_dim || token_idx >= num_tokens) {
    return;
  }
-  const int inbatch_seq_idx = token_idx - cu_seq_lens[batch];
-  const int block_idx =
-      block_table[batch * num_blocks + inbatch_seq_idx / cache_block_size];
+  const int inbatch_seq_idx = token_idx - cu_seq_lens[batch_idx[threadIdx.y]];
+  const int block_idx = block_table[batch_idx[threadIdx.y] * num_blocks +
+                                    inbatch_seq_idx / cache_block_size];
  const int64_t src_block_offset = block_idx * block_stride;
  const int64_t cache_inblock_offset =
      (inbatch_seq_idx % cache_block_size) * head_dim + head_idx;
@@ -731,28 +724,6 @@ void reshape_and_cache_flash(
  int num_tokens = slot_mapping.size(0);
  int num_heads = key.size(1);
  int head_size = key.size(2);
-
-  const at::cuda::OptionalCUDAGuard device_guard(device_of(key));
-  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
-
-  if (kv_cache_dtype == "nvfp4") {
-#if defined(ENABLE_NVFP4_SM100) || defined(ENABLE_NVFP4_SM120)
-    // NVFP4 dispatch is compiled separately for SM100+.
-    extern void reshape_and_cache_nvfp4_dispatch(
-        torch::Tensor & key, torch::Tensor & value, torch::Tensor & key_cache,
-        torch::Tensor & value_cache, torch::Tensor & slot_mapping,
-        torch::Tensor & k_scale, torch::Tensor & v_scale);
-    reshape_and_cache_nvfp4_dispatch(key, value, key_cache, value_cache,
-                                     slot_mapping, k_scale, v_scale);
-    return;
-#else
-    TORCH_CHECK(false,
-                "NVFP4 KV cache requires SM100+ (Blackwell). "
-                "Please rebuild vllm with a Blackwell-compatible CUDA target.");
-#endif
-  }
-
-  // Original FP8/auto path.
  int block_size = key_cache.size(1);

  int64_t key_stride = key.stride(0);
@@ -770,6 +741,8 @@ void reshape_and_cache_flash(

  dim3 grid(num_tokens);
  dim3 block(std::min(num_heads * head_size, 512));
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(key));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

  DISPATCH_BY_KV_CACHE_DTYPE(key.dtype(), kv_cache_dtype,
                             CALL_RESHAPE_AND_CACHE_FLASH);
@@ -1497,9 +1470,6 @@ void concat_mla_q(torch::Tensor& ql_nope,  // [num_tokens, num_heads, nope_dim]
  TORCH_CHECK(ql_nope.stride(2) == 1, "ql_nope must have stride 1 in dim 2");
  TORCH_CHECK(q_pe.stride(2) == 1, "q_pe must have stride 1 in dim 2");
  TORCH_CHECK(q_out.stride(2) == 1, "q_out must have stride 1 in dim 2");
-  TORCH_CHECK(ql_nope.scalar_type() == at::ScalarType::Half ||
-                  ql_nope.scalar_type() == at::ScalarType::BFloat16,
-              "ql_nope must be float16 or bfloat16 dtype");

  if (num_tokens == 0) return;

@@ -1511,7 +1481,7 @@ void concat_mla_q(torch::Tensor& ql_nope,  // [num_tokens, num_heads, nope_dim]
  const at::cuda::OptionalCUDAGuard device_guard(device_of(ql_nope));
  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

-  VLLM_DISPATCH_HALF_TYPES(ql_nope.scalar_type(), "concat_mla_q", [&] {
+  VLLM_DISPATCH_FLOATING_TYPES(ql_nope.scalar_type(), "concat_mla_q", [&] {
    vllm::ConcatMLAQKernel<scalar_t, 512><<<grid_size, block_size, 0, stream>>>(
        q_out.data_ptr<scalar_t>(), ql_nope.data_ptr<scalar_t>(),
        q_pe.data_ptr<scalar_t>(), num_tokens, num_heads, q_out.stride(0),
@@ -1,71 +0,0 @@
-#include "cpu_types.hpp"
-
-#include <array>
-#include <cstdint>
-#include <mutex>
-#include <string>
-
-#include <ATen/ops/empty.h>
-#include <ATen/ops/gelu.h>
-#include <c10/util/BFloat16.h>
-
-constexpr uint32_t ActivationLutSize = 1u << 16;
-
-at::Tensor gelu_reference(const at::Tensor& x) { return at::gelu(x, "none"); }
-
-void maybe_init_activation_lut_bf16(
-    uint16_t* lut, std::once_flag& once,
-    at::Tensor (*activation)(const at::Tensor&)) {
-  std::call_once(once, [&]() {
-    auto lut_input =
-        at::empty({static_cast<int64_t>(ActivationLutSize)},
-                  at::TensorOptions().device(at::kCPU).dtype(at::kFloat));
-    auto* lut_input_ptr = lut_input.data_ptr<float>();
-#pragma omp parallel for
-    for (uint32_t i = 0; i < ActivationLutSize; ++i) {
-      lut_input_ptr[i] = c10::detail::f32_from_bits(static_cast<uint16_t>(i));
-    }
-
-    auto lut_output = activation(lut_input);
-    const auto* lut_output_ptr = lut_output.data_ptr<float>();
-#pragma omp parallel for
-    for (uint32_t i = 0; i < ActivationLutSize; ++i) {
-      lut[i] = c10::detail::round_to_nearest_even(lut_output_ptr[i]);
-    }
-  });
-}
-
-void activation_lut_bf16(torch::Tensor& out, torch::Tensor& input,
-                         const uint16_t* lut, const char* op_name) {
-  TORCH_CHECK(input.scalar_type() == at::kBFloat16, op_name,
-              ": input must be bfloat16");
-  TORCH_CHECK(out.scalar_type() == at::kBFloat16, op_name,
-              ": out must be bfloat16");
-  TORCH_CHECK(input.is_contiguous(), op_name, ": input must be contiguous");
-  TORCH_CHECK(out.is_contiguous(), op_name, ": out must be contiguous");
-
-  const auto* src =
-      reinterpret_cast<const uint16_t*>(input.data_ptr<at::BFloat16>());
-  auto* dst = reinterpret_cast<uint16_t*>(out.data_ptr<at::BFloat16>());
-  const int64_t n = input.numel();
-
-  CPU_KERNEL_GUARD_IN(activation_lut_bf16_impl)
-#pragma omp parallel for
-  for (int64_t i = 0; i < n; ++i) {
-    dst[i] = lut[src[i]];
-  }
-  CPU_KERNEL_GUARD_OUT(activation_lut_bf16_impl)
-}
-
-void activation_lut_bf16(torch::Tensor& out, torch::Tensor& input,
-                         const std::string& activation) {
-  if (activation == "gelu") {
-    static std::array<uint16_t, ActivationLutSize> lut{};
-    static std::once_flag once;
-    maybe_init_activation_lut_bf16(lut.data(), once, gelu_reference);
-    activation_lut_bf16(out, input, lut.data(), "gelu_lut");
-    return;
-  }
-
-  TORCH_CHECK(false, "Unsupported activation: ", activation);
-}
@@ -61,23 +61,8 @@
 #endif

 #ifdef __aarch64__
-  // Implementation of neon_expf copied from Arm Optimized Routines (expf
-  // AdvSIMD)
+  // Implementation copied from Arm Optimized Routines (expf AdvSIMD)
  // https://github.com/ARM-software/optimized-routines/blob/master/math/aarch64/advsimd/expf.c
-  //
-  // Additional fast exponential intended for cases where outputs will be
-  // downcasted to FP16 / BF16 (e.g. attention softmax). Accurate within 1 ULP
-  // for FP16 Accurate within 1 ULP for BF16 for inputs in [-87.683, 88.376] &
-  // clamps inputs outside this range to 0 / inf. Implementation is similar to
-  // exp_u20, but:
-  // - uses a third degree polynomial approximation for exp(r) instead of a
-  // fifth degree one, with coefficients re-tuned.
-  // - does not split natural log (ln) into high / low parts
-  // - clamps exp(x) to 0 for x < -87.683113f and inf for x > 88.3762589f
-  // exp(x) = 2^n (exp(r))
-  // r = x - n*ln2, with n = round(x/ln2)
-  // exp(r) ~ poly(r) = 1 + r + r^2 * (c3 + c2 * r)
-  // n = round(x / ln2), r = x - n*ln2
  #include <limits>
  #define DEFINE_FAST_EXP                                                      \
    const float32x4_t inv_ln2 = vdupq_n_f32(0x1.715476p+0f);                   \
@@ -121,38 +106,7 @@
                          result.val[2] = neon_expf(vec.reg.val[2]);           \
                          result.val[3] = neon_expf(vec.reg.val[3]);           \
                          return vec_op::FP32Vec16(result);                    \
-                        };                                                     \
-    const float32x4_t lower_bound = vdupq_n_f32(-0x1.5ebb82p+6f);              \
-    const float32x4_t upper_bound = vdupq_n_f32(0x1.61814ap+6f);               \
-    constexpr float ln2 = 0x1.62e43p-1f;                                       \
-    constexpr float f_c2 = 0x1.5592ecp-3f;                                     \
-    const float32x4_t f_c3 = vdupq_n_f32(0x1.017d34p-1f);                      \
-    auto neon_expf_f16 = [&](float32x4_t values) __attribute__((               \
-                             always_inline)) {                                 \
-      const uint32x4_t lt_lower = vcltq_f32(values, lower_bound);              \
-      const uint32x4_t gt_upper = vcgtq_f32(values, upper_bound);              \
-      float32x4_t n = vrndaq_f32(vmulq_f32(values, inv_ln2));                  \
-      float32x4_t r = vfmsq_n_f32(values, n, ln2);                             \
-      uint32x4_t e = vshlq_n_u32(vreinterpretq_u32_s32(vcvtq_s32_f32(n)), 23); \
-      float32x4_t r2 = vmulq_f32(r, r);                                        \
-      float32x4_t q = vfmaq_n_f32(f_c3, r, f_c2);                              \
-      float32x4_t s = vaddq_f32(vdupq_n_f32(1.0f), r);                         \
-      float32x4_t p = vfmaq_f32(s, q, r2);                                     \
-      float32x4_t y =                                                          \
-          vreinterpretq_f32_u32(vaddq_u32(vreinterpretq_u32_f32(p), e));       \
-      y = vbslq_f32(lt_lower, vdupq_n_f32(0.0f), y);                           \
-      y = vbslq_f32(gt_upper, vdupq_n_f32(INFINITY), y);                       \
-      return y;                                                                \
-    };                                                                         \
-    auto fast_exp_f16 = [&](const vec_op::FP32Vec16& vec)                      \
-                            __attribute__((always_inline)) {                   \
-                              float32x4x4_t result;                            \
-                              result.val[0] = neon_expf_f16(vec.reg.val[0]);   \
-                              result.val[1] = neon_expf_f16(vec.reg.val[1]);   \
-                              result.val[2] = neon_expf_f16(vec.reg.val[2]);   \
-                              result.val[3] = neon_expf_f16(vec.reg.val[3]);   \
-                              return vec_op::FP32Vec16(result);                \
-                            };
+                        };

 #endif  // __aarch64__

@@ -1,16 +1,5 @@
 #include "cpu_attn_dispatch_generated.h"

-// Maps kv_cache_dtype string to Fp8KVCacheDataType enum.
-// "auto" -> kAuto(0); "fp8"/"fp8_e4m3" -> kFp8E4M3; "fp8_e5m2" -> kFp8E5M2.
-static inline cpu_attention::Fp8KVCacheDataType parse_fp8_kv_dtype(
-    const std::string& kv_cache_dtype) {
-  if (kv_cache_dtype == "fp8_e5m2")
-    return cpu_attention::Fp8KVCacheDataType::kFp8E5M2;
-  if (kv_cache_dtype == "fp8_e4m3" || kv_cache_dtype == "fp8")
-    return cpu_attention::Fp8KVCacheDataType::kFp8E4M3;
-  return cpu_attention::Fp8KVCacheDataType::kAuto;
-}
-
 torch::Tensor get_scheduler_metadata(
    const int64_t num_req, const int64_t num_heads_q,
    const int64_t num_heads_kv, const int64_t head_dim,
@@ -60,7 +49,7 @@ torch::Tensor get_scheduler_metadata(
  input.enable_kv_split = enable_kv_split;

  VLLM_DISPATCH_FLOATING_TYPES(dtype, "get_scheduler_metadata", [&]() {
-    CPU_ATTN_DISPATCH(head_dim, isa, 0, [&]() {
+    CPU_ATTN_DISPATCH(head_dim, isa, [&]() {
      input.elem_size = sizeof(scalar_t);
      input.q_buffer_elem_size = sizeof(attn_impl::q_buffer_t);
      input.logits_buffer_elem_size = sizeof(attn_impl::logits_buffer_t);
@@ -83,9 +72,7 @@ void cpu_attn_reshape_and_cache(
        key_cache,  // [num_blocks, num_kv_heads, block_size, head_size]
    torch::Tensor&
        value_cache,  // [num_blocks, num_kv_heads, block_size, head_size]
-    const torch::Tensor& slot_mapping, const std::string& isa,
-    const double k_scale = 1.0, const double v_scale = 1.0,
-    const std::string& kv_cache_dtype = "auto") {
+    const torch::Tensor& slot_mapping, const std::string& isa) {
  TORCH_CHECK_EQ(key.dim(), 3);
  TORCH_CHECK_EQ(value.dim(), 3);
  TORCH_CHECK_EQ(key_cache.dim(), 4);
@@ -93,30 +80,18 @@ void cpu_attn_reshape_and_cache(
  TORCH_CHECK_EQ(key.stride(2), 1);
  TORCH_CHECK_EQ(value.stride(2), 1);

-  const int64_t kv_cache_idx =
-      static_cast<int64_t>(parse_fp8_kv_dtype(kv_cache_dtype));
-  const bool is_fp8 = (kv_cache_idx != 0);
-
-  if (is_fp8) {
-    TORCH_CHECK(key_cache.scalar_type() == at::ScalarType::Byte,
-                "key_cache must be uint8 for FP8 path");
-    TORCH_CHECK(value_cache.scalar_type() == at::ScalarType::Byte,
-                "value_cache must be uint8 for FP8 path");
-    TORCH_CHECK(k_scale > 0, "k_scale must be positive for FP8 path");
-    TORCH_CHECK(v_scale > 0, "v_scale must be positive for FP8 path");
-  }
-
-  const float k_inv = is_fp8 ? 1.0f / static_cast<float>(k_scale) : 0.0f;
-  const float v_inv = is_fp8 ? 1.0f / static_cast<float>(v_scale) : 0.0f;
-
  const int64_t token_num = key.size(0);
-  const int64_t head_num = key.size(1);
-  const int64_t head_dim = key.size(2);
+  const int64_t key_token_num_stride = key.stride(0);
+  const int64_t value_token_num_stride = value.stride(0);
+  const int64_t head_num = value.size(1);
+  const int64_t key_head_num_stride = key.stride(1);
+  const int64_t value_head_num_stride = value.stride(1);
  const int64_t num_blocks = key_cache.size(0);
  const int64_t num_blocks_stride = key_cache.stride(0);
  const int64_t cache_head_num_stride = key_cache.stride(1);
  const int64_t block_size = key_cache.size(2);
  const int64_t block_size_stride = key_cache.stride(2);
+  const int64_t head_dim = key.size(-1);

  cpu_attention::ISA isa_tag = [&]() {
    if (isa == "amx") {
@@ -134,24 +109,16 @@ void cpu_attn_reshape_and_cache(
    }
  }();

-  if (is_fp8) {
-    TORCH_CHECK(isa_tag == cpu_attention::ISA::AMX ||
-                    isa_tag == cpu_attention::ISA::VEC,
-                "FP8 KV cache is only supported on x86 (AMX/VEC) ISA");
-  }
-
  VLLM_DISPATCH_FLOATING_TYPES(
      key.scalar_type(), "cpu_attn_reshape_and_cache", [&]() {
-        CPU_ATTN_DISPATCH(head_dim, isa_tag, kv_cache_idx, [&]() {
-          using kv_t = typename attn_impl::kv_cache_t;
+        CPU_ATTN_DISPATCH(head_dim, isa_tag, [&]() {
          attn_impl::reshape_and_cache(
              key.data_ptr<scalar_t>(), value.data_ptr<scalar_t>(),
-              reinterpret_cast<kv_t*>(key_cache.data_ptr()),
-              reinterpret_cast<kv_t*>(value_cache.data_ptr()),
-              slot_mapping.data_ptr<int64_t>(), token_num, key.stride(0),
-              value.stride(0), head_num, key.stride(1), value.stride(1),
-              num_blocks, num_blocks_stride, cache_head_num_stride, block_size,
-              block_size_stride, k_inv, v_inv);
+              key_cache.data_ptr<scalar_t>(), value_cache.data_ptr<scalar_t>(),
+              slot_mapping.data_ptr<int64_t>(), token_num, key_token_num_stride,
+              value_token_num_stride, head_num, key_head_num_stride,
+              value_head_num_stride, num_blocks, num_blocks_stride,
+              cache_head_num_stride, block_size, block_size_stride);
        });
      });
 }
@@ -170,26 +137,13 @@ void cpu_attention_with_kv_cache(
    const int64_t sliding_window_left, const int64_t sliding_window_right,
    const torch::Tensor& block_table,  // [num_tokens, max_block_num]
    const double softcap, const torch::Tensor& scheduler_metadata,
-    const std::optional<torch::Tensor>& s_aux,  // [num_heads]
-    const double k_scale = 1.0, const double v_scale = 1.0,
-    const std::string& kv_cache_dtype = "auto") {
+    const std::optional<torch::Tensor>& s_aux  // [num_heads]
+) {
  TORCH_CHECK_EQ(query.dim(), 3);
  TORCH_CHECK_EQ(query.stride(2), 1);
  TORCH_CHECK_EQ(key_cache.dim(), 4);
  TORCH_CHECK_EQ(value_cache.dim(), 4);

-  const int64_t kv_cache_idx =
-      static_cast<int64_t>(parse_fp8_kv_dtype(kv_cache_dtype));
-  const bool is_fp8 = (kv_cache_idx != 0);
-  if (is_fp8) {
-    TORCH_CHECK(key_cache.scalar_type() == at::ScalarType::Byte,
-                "key_cache must be uint8 for FP8 path");
-    TORCH_CHECK(value_cache.scalar_type() == at::ScalarType::Byte,
-                "value_cache must be uint8 for FP8 path");
-    TORCH_CHECK(k_scale > 0, "k_scale must be positive for FP8 path");
-    TORCH_CHECK(v_scale > 0, "v_scale must be positive for FP8 path");
-  }
-
  cpu_attention::AttentionInput input;
  input.metadata = reinterpret_cast<cpu_attention::AttentionMetadata*>(
      scheduler_metadata.data_ptr());
@@ -211,32 +165,25 @@ void cpu_attention_with_kv_cache(
  input.block_table = block_table.data_ptr<int32_t>();
  input.alibi_slopes =
      alibi_slopes.has_value() ? alibi_slopes->data_ptr<float>() : nullptr;
+  // For now sink must be bf16
  input.s_aux = s_aux.has_value() ? s_aux->data_ptr<c10::BFloat16>() : nullptr;
  input.scale = scale;
  input.causal = causal;
  input.sliding_window_left = sliding_window_left;
  input.sliding_window_right = sliding_window_right;
  if (input.causal) {
+    // to make boundary calculation easier
    input.sliding_window_right = 0;
  }
-  input.softcap = static_cast<float>(softcap);
-
-  if (is_fp8) {
-    input.k_scale_fp8 = static_cast<float>(k_scale);
-    input.v_scale_fp8 = static_cast<float>(v_scale);
-    TORCH_CHECK(input.metadata->isa == cpu_attention::ISA::AMX ||
-                    input.metadata->isa == cpu_attention::ISA::VEC,
-                "FP8 KV cache is only supported on x86 (AMX/VEC) ISA");
-  }
+  float softcap_fp32 = softcap;
+  input.softcap = softcap_fp32;

  VLLM_DISPATCH_FLOATING_TYPES(
      query.scalar_type(), "cpu_attention_with_kv_cache", [&]() {
-        CPU_ATTN_DISPATCH(
-            query.size(2), input.metadata->isa, kv_cache_idx, [&]() {
-              TORCH_CHECK_EQ(input.block_size % attn_impl::BlockSizeAlignment,
-                             0);
-              cpu_attention::AttentionMainLoop<attn_impl> mainloop;
-              mainloop(&input);
-            });
+        CPU_ATTN_DISPATCH(query.size(2), input.metadata->isa, [&]() {
+          TORCH_CHECK_EQ(input.block_size % attn_impl::BlockSizeAlignment, 0);
+          cpu_attention::AttentionMainLoop<attn_impl> mainloop;
+          mainloop(&input);
+        });
      });
 }
@@ -1,7 +1,6 @@
 #ifndef CPU_ATTN_AMX_HPP
 #define CPU_ATTN_AMX_HPP

-#include "cpu_attn_fp8.hpp"
 #include "cpu_attn_impl.hpp"

 namespace cpu_attention {
@@ -22,10 +21,9 @@ typedef struct __tile_config {
 // 2-2-4 pattern, for 16 < m <= 32
 // TILE 0, 1: load A matrix, row num should be 16, m - 16
 // TILE 2, 3: load B matrix, row num should be 16
-// TILE 4, 5, 6, 7: store results C matrix, row num should be 16, 16,
-// m - 16, m - 16
-// q_buffer_t: A (Q/P) tile type; kv_cache_t: B (K/V cache) tile type.
-template <typename q_buffer_t, typename kv_cache_t>
+// TILE 4, 5, 6, 7: store results C matrix, row num should be 16, 16, m - 16, m
+// - 16
+template <typename kv_cache_t>
 class TileGemm224 {
 public:
  template <AttentionGemmPhase phase, int32_t k_size>
@@ -44,56 +42,13 @@ class TileGemm224 {
  }
 };

-// Dequantize one FP8 tile (AMX_TILE_ROW_NUM rows x 32 cols) to BF16.
-template <typename kv_cache_t>
-FORCE_INLINE void deq_tile_amx(const uint8_t* src, c10::BFloat16* dst) {
-  for (int r = 0; r < AMX_TILE_ROW_NUM; ++r) {
-    if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e4m3fn>) {
-      vec_op::BF16Vec32(src + r * 32, vec_op::fp8_bf16_e4m3_tag{})
-          .save(dst + r * 32);
-    } else {
-      vec_op::BF16Vec32(src + r * 32, vec_op::fp8_bf16_e5m2_tag{})
-          .save(dst + r * 32);
-    }
-  }
-}
-
-// For FP8: dequant src into scratch and return scratch.
-// For BF16: return src directly (scratch is unused; the compiler elides it).
-template <typename kv_cache_t>
-FORCE_INLINE const c10::BFloat16* prepare_b_tile(const kv_cache_t* src,
-                                                 c10::BFloat16* scratch) {
-  if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-    deq_tile_amx<kv_cache_t>(reinterpret_cast<const uint8_t*>(src), scratch);
-    return scratch;
-  } else {
-    return reinterpret_cast<const c10::BFloat16*>(src);
-  }
-}
-
-// Handles both BF16 and FP8 KV cache (2-2-4 pattern).
-template <typename kv_cache_t>
-class TileGemm224<c10::BFloat16, kv_cache_t> {
-  static_assert(std::is_same_v<kv_cache_t, c10::BFloat16> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e5m2>,
-                "kv_cache_t must be BFloat16, Float8_e4m3fn, or Float8_e5m2");
-
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_t, c10::Float8_e5m2>;
-
-  static constexpr int64_t tile_elems = AMX_TILE_BYTES / sizeof(c10::BFloat16);
-  // BF16 path: scratch_elems=1 so the scratch array is eliminated by the
-  // compiler.
-  static constexpr int64_t scratch_elems = fp8_kv ? tile_elems : 1;
-
+template <>
+class TileGemm224<c10::BFloat16> {
 public:
  template <AttentionGemmPhase phase, int32_t k_size>
  FORCE_INLINE static void gemm(const int32_t m_size,
                                c10::BFloat16* __restrict__ a_tile,
-                                kv_cache_t* __restrict__ b_tile,
+                                c10::BFloat16* __restrict__ b_tile,
                                float* __restrict__ c_tile, const int64_t lda,
                                const int64_t ldb, const int64_t ldc,
                                const int32_t block_size,
@@ -101,7 +56,6 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
                                const bool accum_c) {
    const int32_t k_times =
        dynamic_k_size / (AMX_TILE_ROW_NUM * 4 / sizeof(c10::BFloat16));
-
    c10::BFloat16* __restrict__ a_tile_0 = a_tile;
    c10::BFloat16* __restrict__ a_tile_1 = a_tile + lda * AMX_TILE_ROW_NUM;
    const int64_t a_tile_stride = [&]() {
@@ -116,8 +70,8 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
      }
    }();

-    kv_cache_t* __restrict__ b_tile_2 = b_tile;
-    kv_cache_t* __restrict__ b_tile_3 = [&]() {
+    c10::BFloat16* __restrict__ b_tile_2 = b_tile;
+    c10::BFloat16* __restrict__ b_tile_3 = [&]() {
      if constexpr (phase == AttentionGemmPhase::QK) {
        // k_cache is prepacked
        return b_tile + (k_size * AMX_TILE_ROW_BYTES / 4);
@@ -152,16 +106,11 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
      _tile_zero(7);
    }

-    alignas(64) c10::BFloat16 scratch_2[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_3[scratch_elems];
    for (int32_t k = 0; k < k_times; ++k) {
-      const c10::BFloat16* load_2 = prepare_b_tile(b_tile_2, scratch_2);
-      const c10::BFloat16* load_3 = prepare_b_tile(b_tile_3, scratch_3);
-
      _tile_loadd(0, a_tile_0, a_tile_stride);
-      _tile_stream_loadd(2, const_cast<c10::BFloat16*>(load_2), b_tile_stride);
+      _tile_stream_loadd(2, b_tile_2, b_tile_stride);
      _tile_dpbf16ps(4, 0, 2);
-      _tile_stream_loadd(3, const_cast<c10::BFloat16*>(load_3), b_tile_stride);
+      _tile_stream_loadd(3, b_tile_3, b_tile_stride);
      _tile_dpbf16ps(5, 0, 3);
      _tile_loadd(1, a_tile_1, a_tile_stride);
      _tile_dpbf16ps(6, 1, 2);
@@ -205,13 +154,13 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
 };

 // 1-2-2 pattern, for 0 < m <= 16
-// TILE 0, (1): load A matrix, use extra 1 tile for prefetch, row num should
-// be m, m
-// TILE 2, 3, (4, 5): load B matrix, use extra 2 tiles for prefetch, row num
-// should be 16
-// TILE 6, 7: store results C matrix, row num should be m
-// q_buffer_t: A (Q/P) tile type; kv_cache_t: B (K/V cache) tile type.
-template <typename q_buffer_t, typename kv_cache_t>
+// TILE 0, (1): load A matrix, use extra 1 tile for prefetch, row num should be
+// m, m
+// TILE 2, 3, (4, 5): load B matrix, use extra 2 tiles for prefetch, row
+// num should be 16
+// TILE 6, 7, (6, 7): store results C matrix, row num should be
+// m
+template <typename kv_cache_t>
 class TileGemm122 {
 public:
  template <AttentionGemmPhase phase, int32_t k_size>
@@ -230,26 +179,13 @@ class TileGemm122 {
  }
 };

-// Handles both BF16 and FP8 KV cache (1-2-2 pattern).
-template <typename kv_cache_t>
-class TileGemm122<c10::BFloat16, kv_cache_t> {
-  static_assert(std::is_same_v<kv_cache_t, c10::BFloat16> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e5m2>,
-                "kv_cache_t must be BFloat16, Float8_e4m3fn, or Float8_e5m2");
-
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_t, c10::Float8_e5m2>;
-
-  static constexpr int64_t tile_elems = AMX_TILE_BYTES / sizeof(c10::BFloat16);
-  static constexpr int64_t scratch_elems = fp8_kv ? tile_elems : 1;
-
+template <>
+class TileGemm122<c10::BFloat16> {
 public:
  template <AttentionGemmPhase phase, int32_t k_size>
  FORCE_INLINE static void gemm(const int32_t m_size,
                                c10::BFloat16* __restrict__ a_tile,
-                                kv_cache_t* __restrict__ b_tile,
+                                c10::BFloat16* __restrict__ b_tile,
                                float* __restrict__ c_tile, const int64_t lda,
                                const int64_t ldb, const int64_t ldc,
                                const int32_t block_size,
@@ -279,19 +215,21 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
      }
    }();

-    kv_cache_t* __restrict__ b_tile_2 = b_tile;
-    kv_cache_t* __restrict__ b_tile_3 = [&]() {
+    c10::BFloat16* __restrict__ b_tile_2 = b_tile;
+    c10::BFloat16* __restrict__ b_tile_3 = [&]() {
      if constexpr (phase == AttentionGemmPhase::QK) {
+        // k_cache is prepacked
        return b_tile + (k_size * AMX_TILE_ROW_BYTES / 4);
      } else if constexpr (phase == AttentionGemmPhase::PV) {
+        // v_cache is prepacked
        return b_tile + (block_size * AMX_TILE_ROW_BYTES / 4);
      } else {
        TORCH_CHECK(false, "Unreachable");
      }
    }();
-    kv_cache_t* __restrict__ b_tile_4 =
+    c10::BFloat16* __restrict__ b_tile_4 =
        b_tile_2 + AMX_TILE_BYTES / sizeof(c10::BFloat16);
-    kv_cache_t* __restrict__ b_tile_5 =
+    c10::BFloat16* __restrict__ b_tile_5 =
        b_tile_3 + AMX_TILE_BYTES / sizeof(c10::BFloat16);
    int64_t b_stride = AMX_TILE_ROW_BYTES;

@@ -312,25 +250,16 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
      _tile_zero(7);
    }

-    alignas(64) c10::BFloat16 scratch_2[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_3[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_4[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_5[scratch_elems];
    for (int32_t k = 0; k < k_group_times; ++k) {
-      const c10::BFloat16* load_2 = prepare_b_tile(b_tile_2, scratch_2);
-      const c10::BFloat16* load_3 = prepare_b_tile(b_tile_3, scratch_3);
-      const c10::BFloat16* load_4 = prepare_b_tile(b_tile_4, scratch_4);
-      const c10::BFloat16* load_5 = prepare_b_tile(b_tile_5, scratch_5);
-
      _tile_loadd(0, a_tile_0, a_tile_stride);
-      _tile_stream_loadd(2, const_cast<c10::BFloat16*>(load_2), b_stride);
+      _tile_stream_loadd(2, b_tile_2, b_stride);
      _tile_dpbf16ps(6, 0, 2);
-      _tile_stream_loadd(3, const_cast<c10::BFloat16*>(load_3), b_stride);
+      _tile_stream_loadd(3, b_tile_3, b_stride);
      _tile_dpbf16ps(7, 0, 3);
      _tile_loadd(1, a_tile_1, a_tile_stride);
-      _tile_stream_loadd(4, const_cast<c10::BFloat16*>(load_4), b_stride);
+      _tile_stream_loadd(4, b_tile_4, b_stride);
      _tile_dpbf16ps(6, 1, 4);
-      _tile_stream_loadd(5, const_cast<c10::BFloat16*>(load_5), b_stride);
+      _tile_stream_loadd(5, b_tile_5, b_stride);
      _tile_dpbf16ps(7, 1, 5);

      // update ptrs
@@ -350,13 +279,10 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
    }

    if (has_tail) {
-      const c10::BFloat16* load_2 = prepare_b_tile(b_tile_2, scratch_2);
-      const c10::BFloat16* load_3 = prepare_b_tile(b_tile_3, scratch_3);
-
      _tile_loadd(0, a_tile_0, a_tile_stride);
-      _tile_stream_loadd(2, const_cast<c10::BFloat16*>(load_2), b_stride);
+      _tile_stream_loadd(2, b_tile_2, b_stride);
      _tile_dpbf16ps(6, 0, 2);
-      _tile_stream_loadd(3, const_cast<c10::BFloat16*>(load_3), b_stride);
+      _tile_stream_loadd(3, b_tile_3, b_stride);
      _tile_dpbf16ps(7, 0, 3);
    }

@@ -376,25 +302,21 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
    _tile_loadconfig(&config);
  }
 };
-
 }  // namespace

-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e5m2>;
-
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::AMX, scalar_t, head_dim> {
 public:
  using query_t = scalar_t;
  using q_buffer_t = scalar_t;
-  using kv_cache_t = kv_cache_scalar_t;
+  using kv_cache_t = scalar_t;
  using logits_buffer_t = float;
  using partial_output_buffer_t = float;
  using prob_buffer_t = scalar_t;

  constexpr static int64_t BlockSizeAlignment =
-      32;  // AMX_TILE_ROW_NUM = 16 tokens/tile; 32 = 2 tiles
+      AMX_TILE_ROW_BYTES /
+      sizeof(kv_cache_t);  // KV token num unit of QK and PV phases
  constexpr static int64_t HeadDimAlignment =
      2 * (AMX_TILE_ROW_BYTES / 4);  // headdim num unit of PV phase
  constexpr static int64_t MaxQHeadNumPerIteration = 32;
@@ -402,9 +324,6 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
  constexpr static ISA ISAType = ISA::AMX;
  constexpr static bool scale_on_logits = true;

-  float k_scale = 1.0f;
-  float v_scale = 1.0f;
-
 public:
  AttentionImpl() : current_q_head_num_(0) {
    // Use all columns in AMX tiles
@@ -413,50 +332,21 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {

  ~AttentionImpl() { _tile_release(); }

-  void init_from_input(const AttentionInput* input) {
-    if constexpr (fp8_kv) {
-      k_scale = input->k_scale_fp8;
-      v_scale = input->v_scale_fp8;
-    }
-  }
-
-  float get_output_v_scale() const noexcept {
-    if constexpr (fp8_kv) {
-      // AMX dequant places FP8 payload into a BF16 field (exponent bias 127).
-      // Correction = 2^(127 - FP8_bias): E4M3 bias=7 → 2^120, E5M2 bias=15 →
-      // 2^112.
-      constexpr float bias =
-          std::is_same_v<kv_cache_t, c10::Float8_e5m2> ? 0x1p112f : 0x1p120f;
-      return v_scale * bias;
-    }
-    return 1.0f;
-  }
-
  template <template <typename tile_gemm_t> typename attention>
  FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
-    if constexpr (fp8_kv) {
-      // Same bias correction as get_output_v_scale: AMX FP8→BF16 dequant
-      // shifts the exponent bias from FP8 to BF16 (127), so we multiply by
-      // 2^(127-FP8_bias) to recover the true value. E4M3: 2^120, E5M2: 2^112.
-      const float bias =
-          std::is_same_v<kv_cache_t, c10::Float8_e5m2> ? 0x1p112f : 0x1p120f;
-      scale *= k_scale * bias;
-    }
    if (q_head_num > AMX_TILE_ROW_NUM) {
      if (q_head_num != current_q_head_num_) {
        current_q_head_num_ = q_head_num;
-        TileGemm224<q_buffer_t, kv_cache_t>::init_tile_config(q_head_num,
-                                                              amx_tile_config_);
+        TileGemm224<kv_cache_t>::init_tile_config(q_head_num, amx_tile_config_);
      }
-      attention<TileGemm224<q_buffer_t, kv_cache_t>> attention_iteration;
+      attention<TileGemm224<kv_cache_t>> attention_iteration;
      attention_iteration(CPU_ATTENTION_PARAMS);
    } else {
      if (q_head_num != current_q_head_num_) {
        current_q_head_num_ = q_head_num;
-        TileGemm122<q_buffer_t, kv_cache_t>::init_tile_config(q_head_num,
-                                                              amx_tile_config_);
+        TileGemm122<kv_cache_t>::init_tile_config(q_head_num, amx_tile_config_);
      }
-      attention<TileGemm122<q_buffer_t, kv_cache_t>> attention_iteration;
+      attention<TileGemm122<kv_cache_t>> attention_iteration;
      attention_iteration(CPU_ATTENTION_PARAMS);
    }
  }
@@ -521,26 +411,13 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
  // reshape KV to AMX friendly layout
  static void reshape_and_cache(
      const scalar_t* __restrict__ key, const scalar_t* __restrict__ value,
-      kv_cache_t* __restrict__ key_cache, kv_cache_t* __restrict__ value_cache,
+      scalar_t* __restrict__ key_cache, scalar_t* __restrict__ value_cache,
      const int64_t* __restrict__ slot_mapping, const int64_t token_num,
      const int64_t key_token_num_stride, const int64_t value_token_num_stride,
      const int64_t head_num, const int64_t key_head_num_stride,
      const int64_t value_head_num_stride, const int64_t num_blocks,
      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float k_inv = 0.0f, const float v_inv = 0.0f) {
-    if constexpr (fp8_kv) {
-      constexpr auto qfn = select_fp8_quant_fn<kv_cache_t>();
-      reshape_and_cache_fp8_amx_impl<scalar_t, qfn>(
-          key, value, reinterpret_cast<uint8_t*>(key_cache),
-          reinterpret_cast<uint8_t*>(value_cache), slot_mapping, token_num,
-          head_num, head_dim, block_size, key_token_num_stride,
-          key_head_num_stride, value_token_num_stride, value_head_num_stride,
-          num_blocks_stride, cache_head_num_stride, num_blocks_stride,
-          cache_head_num_stride, k_inv, v_inv);
-      return;
-    }
-
+      const int64_t block_size, const int64_t block_size_stride) {
    // For AMX 2D tiles, size of each line is 64 bytes
    constexpr int64_t amx_tile_row_size = AMX_TILE_ROW_BYTES;
    // For AMX B matrix, N always is 16
@@ -549,9 +426,6 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
    // For now suppose block_size is divisible by amx_tile_column_num
    TORCH_CHECK_EQ(block_size % amx_b_tile_k_size, 0);

-    scalar_t* __restrict__ kc = reinterpret_cast<scalar_t*>(key_cache);
-    scalar_t* __restrict__ vc = reinterpret_cast<scalar_t*>(value_cache);
-
 #pragma omp parallel for collapse(2)
    for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
      for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -579,7 +453,8 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
          constexpr int64_t quadword_num_per_group =
              token_num_per_group * quadword_num;
          int32_t* key_cache_start_ptr =
-              reinterpret_cast<int32_t*>(kc + block_idx * num_blocks_stride +
+              reinterpret_cast<int32_t*>(key_cache +
+                                         block_idx * num_blocks_stride +
                                         head_idx * cache_head_num_stride) +
              group_idx * quadword_num_per_group + group_offset;

@@ -608,7 +483,7 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
                                            token_idx * value_token_num_stride +
                                            head_idx * value_head_num_stride;
          scalar_t* value_cache_start_ptr =
-              vc + block_idx * num_blocks_stride +
+              value_cache + block_idx * num_blocks_stride +
              head_idx * cache_head_num_stride +
              sub_group_idx * token_num_per_sub_group * amx_b_tile_n_size +
              sub_group_offset;
@@ -1,214 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#pragma once
-#include <algorithm>
-#include <cmath>
-#include <cstdint>
-#include <limits>
-#include <type_traits>
-
-#include "cpu/utils.hpp"
-
-typedef uint32_t __attribute__((__may_alias__)) u32_alias_t;
-typedef uint16_t __attribute__((__may_alias__)) u16_alias_t;
-typedef float __attribute__((__may_alias__)) f32_alias_t;
-
-// Reference scalar dequant — used to verify vectorized AMX dequant.
-inline float fp8e4m3_to_float_scalar(uint8_t b, float scale) noexcept {
-  // NaN encoding in E4M3
-  if ((b & 0x7F) == 0x7F) return std::numeric_limits<float>::quiet_NaN();
-  uint32_t b_u32 = static_cast<uint32_t>(b);
-  uint32_t sign = (b_u32 & 0x80) << 24;
-  uint32_t payload = (b_u32 & 0x7F) << 20;
-  uint32_t bits = sign | payload;
-  float b_f32_unscaled = *reinterpret_cast<const f32_alias_t*>(&bits);
-  float b_f32_scaled = b_f32_unscaled * scale * 0x1p120f;
-  return b_f32_scaled;
-}
-
-inline uint8_t float_to_fp8e4m3_scalar(float v, float inv_scale) noexcept {
-  v *= inv_scale;
-  constexpr float fp8_max = 448.0f;
-  v = std::max(-fp8_max, std::min(fp8_max, v));
-  if (v == 0.0f) return 0;
-
-  // Inverse mapping of fp8e4m3_to_float_scalar: shift the effective exponent
-  // bias from fp32 (127) back to fp8 e4m3 (7), then pack sign|payload.
-  float v_f32_unscaled = v * 0x1p-120f;
-  uint32_t bits = *reinterpret_cast<const u32_alias_t*>(&v_f32_unscaled);
-  uint8_t sign = static_cast<uint8_t>((bits >> 24) & 0x80);
-  uint8_t payload = static_cast<uint8_t>((bits >> 20) & 0x7F);
-  if (payload == 0) return sign;
-  payload = std::min<uint8_t>(payload, 0x7E);  // keep 0x7F as NaN encoding
-  return static_cast<uint8_t>(sign | payload);
-}
-
-// ---------------------------------------------------------------------------
-// AMX reshape impl — parameterised on the quantisation function.
-// Writes key/value into uint8 FP8 KV cache using the AMX tile-friendly layout.
-// K: halfword-packed (2 FP8 per uint16, token_num_per_group=16).
-// V: sub-group packing (token_num_per_sub_group=2, head_elems_per_group=16).
-// block_size must be divisible by 32.
-// ---------------------------------------------------------------------------
-template <typename scalar_t, uint8_t (*quant_fn)(float, float)>
-inline void reshape_and_cache_fp8_amx_impl(
-    const scalar_t* key_ptr, const scalar_t* value_ptr, uint8_t* key_cache_ptr,
-    uint8_t* value_cache_ptr, const int64_t* slot_ptr, int64_t token_num,
-    int64_t head_num, int64_t head_dim, int64_t block_size, int64_t k_stride0,
-    int64_t k_stride1, int64_t v_stride0, int64_t v_stride1, int64_t kc_stride0,
-    int64_t kc_stride1, int64_t vc_stride0, int64_t vc_stride1, float k_inv,
-    float v_inv) {
-  constexpr int64_t token_num_per_group = 16;  // AMX_TILE_ROW_NUM
-  const int64_t halfword_num = head_dim / 2;   // 2 FP8 per uint16
-  const int64_t halfword_num_per_group = token_num_per_group * halfword_num;
-  constexpr int64_t head_elems_per_group = 16;
-  constexpr int64_t token_num_per_sub_group = 2;  // = 4 / sizeof(BF16)
-  const int64_t group_num = head_dim / head_elems_per_group;
-  const int64_t group_size = block_size * head_elems_per_group;
-
-#pragma omp parallel for collapse(2) schedule(static)
-  for (int64_t tok = 0; tok < token_num; ++tok) {
-    for (int64_t h = 0; h < head_num; ++h) {
-      const int64_t slot = slot_ptr[tok];
-      if (slot < 0) continue;
-      const int64_t block_idx = slot / block_size;
-      const int64_t block_offset = slot % block_size;
-
-      // Key: halfword-packed, 2 FP8 per uint16
-      {
-        const scalar_t* ksrc = key_ptr + tok * k_stride0 + h * k_stride1;
-        const int64_t group_idx = block_offset / token_num_per_group;
-        const int64_t group_offset = block_offset % token_num_per_group;
-        uint16_t* kdst =
-            reinterpret_cast<uint16_t*>(key_cache_ptr + block_idx * kc_stride0 +
-                                        h * kc_stride1) +
-            group_idx * halfword_num_per_group + group_offset;
-        for (int64_t j = 0; j < halfword_num; ++j) {
-          uint8_t fp8_0 = quant_fn(static_cast<float>(ksrc[j * 2]), k_inv);
-          uint8_t fp8_1 = quant_fn(static_cast<float>(ksrc[j * 2 + 1]), k_inv);
-          uint8_t bytes[2] = {fp8_0, fp8_1};
-          uint16_t hw = *reinterpret_cast<const u16_alias_t*>(bytes);
-          kdst[j * token_num_per_group] = hw;
-        }
-      }
-
-      // Value: sub-group packing (token_num_per_sub_group = 2)
-      {
-        const scalar_t* vsrc = value_ptr + tok * v_stride0 + h * v_stride1;
-        const int64_t sub_group_idx = block_offset / token_num_per_sub_group;
-        const int64_t sub_group_offset = block_offset % token_num_per_sub_group;
-        uint8_t* vdst =
-            value_cache_ptr + block_idx * vc_stride0 + h * vc_stride1 +
-            sub_group_idx * token_num_per_sub_group * head_elems_per_group +
-            sub_group_offset;
-        for (int64_t i = 0; i < group_num; ++i) {
-          for (int64_t j = 0; j < head_elems_per_group; ++j)
-            vdst[j * token_num_per_sub_group] =
-                quant_fn(static_cast<float>(vsrc[j]), v_inv);
-          vsrc += head_elems_per_group;
-          vdst += group_size;
-        }
-      }
-    }
-  }
-}
-
-// ---------------------------------------------------------------------------
-// FP8 E5M2 scalar helpers
-// ---------------------------------------------------------------------------
-
-// Reference scalar dequant — used to verify vectorized AMX dequant.
-// FP8 E5M2: s[7] e[6:2] m[1:0], exponent bias = 15 (same as FP16).
-// Byte b → FP16 bits = b << 8 (no bias correction needed).
-inline float fp8e5m2_to_float_scalar(uint8_t b, float scale) noexcept {
-  const uint8_t exp_bits = (b >> 2) & 0x1F;
-  const uint8_t mant_bits = b & 0x03;
-  // NaN: exp=11111, mant!=00
-  if (exp_bits == 0x1F && mant_bits != 0)
-    return std::numeric_limits<float>::quiet_NaN();
-  const uint32_t sign = static_cast<uint32_t>(b & 0x80) << 24;
-  if (exp_bits == 0x1F)
-    return sign ? -std::numeric_limits<float>::infinity()
-                : std::numeric_limits<float>::infinity();
-  if (exp_bits == 0) {  // subnormal: (-1)^s * 2^-14 * mant/4
-    if (mant_bits == 0) return 0.0f;
-    float v = mant_bits * 0x1p-16f;
-    return (sign ? -v : v) * scale;
-  }
-  // Normal: FP32 exp = exp5 - 15 + 127, mantissa top 2 bits
-  uint32_t fp32_bits = sign |
-                       ((static_cast<uint32_t>(exp_bits) - 15 + 127) << 23) |
-                       (static_cast<uint32_t>(mant_bits) << 21);
-  float val = *reinterpret_cast<const f32_alias_t*>(&fp32_bits);
-  return val * scale;
-}
-
-inline uint8_t float_to_fp8e5m2_scalar(float v, float inv_scale) noexcept {
-  v *= inv_scale;
-  constexpr float fp8_e5m2_max = 57344.0f;
-  v = std::max(-fp8_e5m2_max, std::min(fp8_e5m2_max, v));
-  if (v == 0.0f) return 0;
-  uint32_t bits = *reinterpret_cast<const u32_alias_t*>(&v);
-  const uint8_t sign = static_cast<uint8_t>((bits >> 24) & 0x80);
-  const int32_t exp_fp32 = static_cast<int32_t>((bits >> 23) & 0xFF) - 127;
-  const uint8_t mant2 = static_cast<uint8_t>((bits >> 21) & 0x03);
-  if (exp_fp32 < -14) {  // subnormal in E5M2
-    const int shift = -14 - exp_fp32;
-    if (shift + 21 >= 32)
-      return sign;  // underflow: too small for E5M2 subnormal
-    const uint32_t m = (0x800000u | (bits & 0x7FFFFFu)) >> (shift + 21);
-    return sign | static_cast<uint8_t>(std::min<uint32_t>(m, 3u));
-  }
-  const uint8_t exp5 = static_cast<uint8_t>(exp_fp32 + 15);
-  return sign | (exp5 << 2) | mant2;
-}
-
-// ---------------------------------------------------------------------------
-// Select the FP8 quant function at compile time based on kv_cache_t.
-// ---------------------------------------------------------------------------
-template <typename kv_cache_t>
-constexpr auto select_fp8_quant_fn() {
-  if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>)
-    return float_to_fp8e5m2_scalar;
-  else
-    return float_to_fp8e4m3_scalar;
-}
-
-// ---------------------------------------------------------------------------
-// VEC reshape impl — parameterised on the quantisation function.
-// Writes key (column-major) and value (row-major) into uint8 FP8 KV cache.
-// The pragma omp must live outside VLLM_DISPATCH_FLOATING_TYPES because
-// #pragma cannot appear inside variadic macro arguments.
-// ---------------------------------------------------------------------------
-template <typename scalar_t, uint8_t (*quant_fn)(float, float)>
-inline void reshape_and_cache_fp8_vec_impl(
-    const scalar_t* key_ptr, const scalar_t* value_ptr, uint8_t* key_cache_ptr,
-    uint8_t* value_cache_ptr, const int64_t* slot_ptr, int64_t token_num,
-    int64_t head_num, int64_t head_dim, int64_t block_size, int64_t k_stride0,
-    int64_t k_stride1, int64_t v_stride0, int64_t v_stride1, int64_t kc_stride0,
-    int64_t kc_stride1, int64_t vc_stride0, int64_t vc_stride1, float k_inv,
-    float v_inv) {
-#pragma omp parallel for collapse(2) schedule(static)
-  for (int64_t tok = 0; tok < token_num; ++tok) {
-    for (int64_t h = 0; h < head_num; ++h) {
-      const int64_t slot = slot_ptr[tok];
-      if (slot < 0) continue;
-      const int64_t block_idx = slot / block_size;
-      const int64_t block_offset = slot % block_size;
-
-      // Key layout: column-major within block
-      const scalar_t* ksrc = key_ptr + tok * k_stride0 + h * k_stride1;
-      uint8_t* kdst = key_cache_ptr + block_idx * kc_stride0 + h * kc_stride1 +
-                      block_offset;
-      for (int64_t i = 0; i < head_dim; ++i)
-        kdst[i * block_size] = quant_fn(static_cast<float>(ksrc[i]), k_inv);
-
-      // Value layout: row-major within block (contiguous head_dim bytes)
-      const scalar_t* vsrc = value_ptr + tok * v_stride0 + h * v_stride1;
-      uint8_t* vdst = value_cache_ptr + block_idx * vc_stride0 +
-                      h * vc_stride1 + block_offset * head_dim;
-      for (int64_t i = 0; i < head_dim; ++i)
-        vdst[i] = quant_fn(static_cast<float>(vsrc[i]), v_inv);
-    }
-  }
-}
@@ -14,22 +14,8 @@
 namespace cpu_attention {
 enum class ISA { AMX, VEC, VEC16, NEON, VXE };

-// Mirrors csrc/attention/dtype_fp8.cuh Fp8KVCacheDataType exactly.
-enum class Fp8KVCacheDataType {
-  kAuto = 0,
-  kFp8E4M3 = 1,
-  kFp8E5M2 = 2,
-};
-
-struct AttentionInput;
-
-template <ISA isa, typename scalar_t, int64_t head_dim,
-          typename kv_cache_scalar_t = scalar_t>
-class AttentionImpl {
- public:
-  void init_from_input(const AttentionInput*) {}
-  float get_output_v_scale() const noexcept { return 1.0f; }
-};
+template <ISA isa, typename scalar_t, int64_t head_dim>
+class AttentionImpl {};

 struct AttentionWorkItemGroup {
  int32_t req_id;
@@ -161,9 +147,6 @@ struct AttentionMetadata {
      case ISA::NEON:
        ss << "NEON, ";
        break;
-      case ISA::VXE:
-        ss << "VXE, ";
-        break;
    }
    ss << "workitem_group_num: " << workitem_group_num
       << ", reduction_item_num: " << reduction_item_num
@@ -794,9 +777,6 @@ struct AttentionInput {
  int32_t sliding_window_left;
  int32_t sliding_window_right;
  float softcap;
-  // FP8 KV cache scales (used by FP8 attention implementations)
-  float k_scale_fp8 = 1.0f;
-  float v_scale_fp8 = 1.0f;
 };

 #define DEFINE_CPU_ATTENTION_PARAMS                                         \
@@ -1169,11 +1149,7 @@ class AttentionMainLoop {
                       bool use_sink) {
 #ifdef DEFINE_FAST_EXP
      DEFINE_FAST_EXP
-      bool constexpr IsReducedPrecision =
-          std::is_same_v<query_t, c10::BFloat16> ||
-          std::is_same_v<query_t, c10::Half>;
 #endif
-
      using prob_buffer_vec_t = typename VecTypeTrait<prob_buffer_t>::vec_t;
      static_assert(sizeof(prob_buffer_t) <= sizeof(logits_buffer_t));

@@ -1222,17 +1198,8 @@ class AttentionMainLoop {
            vec = vec - max_vec;

            // compute exp
-
-#if defined(DEFINE_FAST_EXP)
-  #ifdef __aarch64__
-            if constexpr (IsReducedPrecision) {
-              vec = fast_exp_f16(vec);
-            } else
-  #endif
-            {
-              vec = fast_exp(vec);
-            }
-
+#ifdef DEFINE_FAST_EXP
+            vec = fast_exp(vec);
            prob_buffer_vec_t output_vec(vec);
            output_vec.save(curr_prob_buffer_iter);
 #else
@@ -1288,11 +1255,7 @@ class AttentionMainLoop {
                       int32_t kv_tile_token_num, float softcap_scale) {
 #ifdef DEFINE_FAST_EXP
      DEFINE_FAST_EXP
-      bool constexpr IsReducedPrecision =
-          std::is_same_v<query_t, c10::BFloat16> ||
-          std::is_same_v<query_t, c10::Half>;
 #endif
-
      float inv_softcap_scale = 1.0 / softcap_scale;
      vec_op::FP32Vec16 softcap_scale_vec(softcap_scale);
      vec_op::FP32Vec16 inv_softcap_scale_vec(inv_softcap_scale);
@@ -1306,15 +1269,8 @@ class AttentionMainLoop {
          vec_op::FP32Vec16 vec(curr_logits_buffer_iter);
          vec = vec * inv_softcap_scale_vec;

-#if defined(DEFINE_FAST_EXP)
-  #ifdef __aarch64__
-          if constexpr (IsReducedPrecision) {
-            vec = fast_exp_f16(vec);
-          } else
-  #endif
-          {
-            vec = fast_exp(vec);
-          }
+#ifdef DEFINE_FAST_EXP
+          vec = fast_exp(vec);
          vec_op::FP32Vec16 inv_vec = ones_vec / vec;
          vec = (vec - inv_vec) / (vec + inv_vec);
 #else
@@ -1391,13 +1347,6 @@ class AttentionMainLoop {
      }

      attention_impl_t attn_impl;
-      constexpr bool fp8_kv = std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                              std::is_same_v<kv_cache_t, c10::Float8_e5m2>;
-      float output_v_scale = 1.0f;
-      if constexpr (fp8_kv) {
-        attn_impl.init_from_input(input);
-        output_v_scale = attn_impl.get_output_v_scale();
-      }

      // general information
      const int32_t q_head_num = input->num_heads;
@@ -1777,7 +1726,7 @@ class AttentionMainLoop {
                               reinterpret_cast<query_t*>(input->output) +
                                   output_buffer_offset,
                               sum_buffer, actual_q_heads_per_kv,
-                               actual_q_token_num, q_head_num, output_v_scale);
+                               actual_q_token_num, q_head_num);
                } else {
                  const int32_t stride =
                      actual_q_heads_per_kv * split_kv_q_token_num_threshold;
@@ -1847,7 +1796,7 @@ class AttentionMainLoop {
              split_output_buffer,
              reinterpret_cast<query_t*>(input->output) + output_buffer_offset,
              split_sum_buffer, actual_q_heads_per_kv, curr_output_token_num,
-              q_head_num, output_v_scale);
+              q_head_num);
        }
      }
    }
@@ -1971,8 +1920,8 @@ class AttentionMainLoop {
                    query_t* __restrict__ curr_output_buffer,
                    float* __restrict__ sum_buffer,
                    const int32_t q_heads_per_kv,
-                    const int32_t actual_q_token_num, const int32_t q_head_num,
-                    const float v_scale = 1.0f) {
+                    const int32_t actual_q_token_num,
+                    const int32_t q_head_num) {
    // final output
    using output_vec_t = typename VecTypeTrait<query_t>::vec_t;

@@ -1986,7 +1935,7 @@ class AttentionMainLoop {
          curr_partial_output_buffer;
      query_t* __restrict__ curr_output_buffer_iter = curr_output_buffer;
      for (int32_t head_idx = 0; head_idx < q_heads_per_kv; ++head_idx) {
-        vec_op::FP32Vec16 inv_sum_scale_vec(v_scale / *curr_sum_buffer);
+        vec_op::FP32Vec16 inv_sum_scale_vec(1.0 / *curr_sum_buffer);

        for (int32_t i = 0; i < group_num_per_head; ++i) {
          vec_op::FP32Vec16 vec(curr_partial_output_buffer_iter);
@@ -248,8 +248,8 @@ class TileGemmNeonFMLA {
 }  // namespace

 // this is similar to "ISA::VEC" at the moment
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::NEON, scalar_t, head_dim, kv_cache_scalar_t> {
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::NEON, scalar_t, head_dim> {
 public:
  using query_t = scalar_t;
  using q_buffer_t = float;
@@ -343,8 +343,7 @@ class AttentionImpl<ISA::NEON, scalar_t, head_dim, kv_cache_scalar_t> {
      const int64_t head_num, const int64_t key_head_num_stride,
      const int64_t value_head_num_stride, const int64_t num_blocks,
      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
+      const int64_t block_size, const int64_t block_size_stride) {
 #pragma omp parallel for collapse(2)
    for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
      for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -389,7 +388,7 @@ class AttentionImpl<ISA::NEON, scalar_t, head_dim, kv_cache_scalar_t> {
 #ifdef ARM_BF16_SUPPORT
 // For BF16 on Arm, reuse the BFMMLA kernels with 32-token alignment.
 template <int64_t head_dim>
-class AttentionImpl<ISA::NEON, c10::BFloat16, head_dim, c10::BFloat16>
+class AttentionImpl<ISA::NEON, c10::BFloat16, head_dim>
    : public AttentionImplNEONBFMMLA<BLOCK_SIZE_ALIGNMENT, ISA::NEON,
                                     head_dim> {};
 #endif
@@ -602,8 +602,7 @@ class AttentionImplNEONBFMMLA {
      [[maybe_unused]] const int64_t num_blocks,
      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
      const int64_t block_size,
-      [[maybe_unused]] const int64_t block_size_stride,
-      const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
+      [[maybe_unused]] const int64_t block_size_stride) {
    const int64_t k_block_stride = (head_dim / TILE_K) * K_INNER_STRIDE;
    const int64_t v_pair_stride =
        (block_size / V_TOKENS_PER_ROW_BLOCK) * V_INNER_STRIDE;
@@ -1,37 +1,11 @@
 #ifndef CPU_ATTN_VEC_HPP
 #define CPU_ATTN_VEC_HPP

-#include "cpu_attn_fp8.hpp"
 #include "cpu_attn_impl.hpp"

 namespace cpu_attention {

 namespace {
-
-// Load 32 kv_cache_t elements starting at ptr and return them as two FP32Vec16s
-// covering the lower 16 and upper 16 positions.
-// For FP8: both halves come from a single BF16Vec32 dequant of 32 bytes.
-// For BF16/FP16/FP32: two separate vector loads at ptr and ptr+16.
-template <typename kv_cache_t>
-FORCE_INLINE std::pair<vec_op::FP32Vec16, vec_op::FP32Vec16> load_b_pair_vec(
-    const kv_cache_t* ptr) {
-  if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e4m3fn>) {
-    // BF16 container, but values are in the FP16 exponent range (bias 15 not
-    // 127).
-    vec_op::BF16Vec32 bf16_b_reg(reinterpret_cast<const uint8_t*>(ptr),
-                                 vec_op::fp8_e4m3_tag{});
-    return {vec_op::FP32Vec16(bf16_b_reg, 0), vec_op::FP32Vec16(bf16_b_reg, 1)};
-  } else if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-    vec_op::BF16Vec32 bf16_b_reg(reinterpret_cast<const uint8_t*>(ptr),
-                                 vec_op::fp8_e5m2_tag{});
-    return {vec_op::FP32Vec16(bf16_b_reg, 0), vec_op::FP32Vec16(bf16_b_reg, 1)};
-  } else {
-    using load_vec_t = typename VecTypeTrait<kv_cache_t>::vec_t;
-    return {vec_op::FP32Vec16(load_vec_t(ptr)),
-            vec_op::FP32Vec16(load_vec_t(ptr + 16))};
-  }
-}
-
 // 8-2-16 pattern, 8 regs for A, 2 regs for B, 16 regs for C, [8, K] @ [k, 32]
 template <typename kv_cache_t>
 class TileGemm82 {
@@ -80,7 +54,10 @@ class TileGemm82 {
                         const int32_t block_size, const int32_t dynamic_k_size,
                         const bool accum_c) {
    static_assert(0 < M && M <= 8);
+    using load_vec_t = typename VecTypeTrait<kv_cache_t>::vec_t;

+    kv_cache_t* __restrict__ curr_b_0 = b_tile;
+    kv_cache_t* __restrict__ curr_b_1 = b_tile + 16;
    float* __restrict__ curr_c_0 = c_tile;
    float* __restrict__ curr_c_1 = c_tile + 16;

@@ -99,14 +76,16 @@ class TileGemm82 {
    }

    float* __restrict__ curr_a = a_tile;
-    kv_cache_t* __restrict__ curr_b = b_tile;
-
    for (int32_t k = 0; k < dynamic_k_size; ++k) {
-      auto [fp32_b_0_reg, fp32_b_1_reg] = load_b_pair_vec(curr_b);
+      load_vec_t b_0_reg(curr_b_0);
+      vec_op::FP32Vec16 fp32_b_0_reg(b_0_reg);
+      load_vec_t b_1_reg(curr_b_1);
+      vec_op::FP32Vec16 fp32_b_1_reg(b_1_reg);

      float* __restrict__ curr_m_a = curr_a;
      vec_op::unroll_loop<int32_t, M>([&](int32_t i) {
-        vec_op::FP32Vec16 a_reg(*curr_m_a);
+        float v = *curr_m_a;
+        vec_op::FP32Vec16 a_reg(v);
        c_regs[i * 2] = c_regs[i * 2] + a_reg * fp32_b_0_reg;
        c_regs[i * 2 + 1] = c_regs[i * 2 + 1] + a_reg * fp32_b_1_reg;

@@ -116,7 +95,8 @@ class TileGemm82 {

      // update
      curr_a += 1;
-      curr_b += ldb;
+      curr_b_0 += ldb;
+      curr_b_1 += ldb;
    }

    vec_op::unroll_loop<int32_t, M>([&](int32_t i) {
@@ -129,20 +109,15 @@ class TileGemm82 {
    });
  }
 };
-
 }  // namespace

 // This is a general but naive implementation based on vector instructions
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e5m2>;
-
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::VEC, scalar_t, head_dim> {
 public:
  using query_t = scalar_t;
  using q_buffer_t = float;
-  using kv_cache_t = kv_cache_scalar_t;
+  using kv_cache_t = scalar_t;
  using logits_buffer_t = float;
  using partial_output_buffer_t = float;
  using prob_buffer_t = float;
@@ -154,45 +129,11 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
  constexpr static int64_t MaxQHeadNumPerIteration = 8;
  constexpr static int64_t HeadDim = head_dim;
  constexpr static ISA ISAType = ISA::VEC;
-  constexpr static bool scale_on_logits = fp8_kv;
-
-  float k_scale = 1.0f;
-  float v_scale = 1.0f;
+  constexpr static bool scale_on_logits = false;  // apply scale on q_buffer

 public:
-  void init_from_input(const AttentionInput* input) {
-    if constexpr (fp8_kv) {
-      k_scale = input->k_scale_fp8;
-      v_scale = input->v_scale_fp8;
-    }
-  }
-
-  float get_output_v_scale() const noexcept {
-    if constexpr (fp8_kv) {
-      // VEC dequant unpacks FP8 into a pseudo-FP16 layout (exponent bias 15).
-      // E4M3 (bias=7) needs correction 2^(15-7) = 2^8; E5M2 bias matches FP16
-      // so no correction.
-      if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-        return v_scale;
-      } else {
-        return v_scale * 0x1p8f;
-      }
-    }
-    return 1.0f;
-  }
-
  template <template <typename tile_gemm_t> typename attention>
  FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
-    if constexpr (fp8_kv) {
-      // Same bias correction as get_output_v_scale: VEC FP8→pseudo-FP16 dequant
-      // uses bias 15; E4M3 (bias=7) needs ×2^8, E5M2 (bias=15) needs no
-      // correction.
-      if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-        scale *= k_scale;
-      } else {
-        scale *= k_scale * 0x1p8f;
-      }
-    }
    attention<TileGemm82<kv_cache_t>> attention_iteration;
    attention_iteration(CPU_ATTENTION_PARAMS);
  }
@@ -220,19 +161,17 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
                              // row-major
  }

-  // Copy q to q_buffer and cast it to fp32.
-  // FP8: QK scale is folded into execute_attention; copy Q unscaled here.
-  void copy_q_heads_tile(scalar_t* __restrict__ src,
-                         float* __restrict__ q_buffer, const int32_t q_num,
-                         const int32_t q_heads_per_kv,
-                         const int64_t q_num_stride,
-                         const int64_t q_head_stride, float scale) {
+  // Copy q to q_buffer and cast it to fp32
+  static void copy_q_heads_tile(
+      scalar_t* __restrict__ src,  // [q_num, q_heads_per_kv, head_size]
+      float* __restrict__ q_buffer, const int32_t q_num,
+      const int32_t q_heads_per_kv, const int64_t q_num_stride,
+      const int64_t q_head_stride, float scale) {
    static_assert(head_dim % 16 == 0);
    constexpr int32_t unroll_size = head_dim / 16;
    using load_vec_t = typename VecTypeTrait<scalar_t>::vec_t;

-    const float effective_scale = fp8_kv ? 1.0f : scale;
-    vec_op::FP32Vec16 scale_vec(effective_scale);
+    vec_op::FP32Vec16 scale_vec(scale);
    for (int32_t q_num_idx = 0; q_num_idx < q_num; ++q_num_idx) {
      for (int32_t q_head_idx = 0; q_head_idx < q_heads_per_kv; ++q_head_idx) {
        scalar_t* __restrict__ curr_q =
@@ -257,26 +196,13 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
  // reshape K as column-major and V as row-major
  static void reshape_and_cache(
      const scalar_t* __restrict__ key, const scalar_t* __restrict__ value,
-      kv_cache_t* __restrict__ key_cache, kv_cache_t* __restrict__ value_cache,
+      scalar_t* __restrict__ key_cache, scalar_t* __restrict__ value_cache,
      const int64_t* __restrict__ slot_mapping, const int64_t token_num,
      const int64_t key_token_num_stride, const int64_t value_token_num_stride,
      const int64_t head_num, const int64_t key_head_num_stride,
      const int64_t value_head_num_stride, const int64_t num_blocks,
      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float k_inv = 0.0f, const float v_inv = 0.0f) {
-    if constexpr (fp8_kv) {
-      constexpr auto qfn = select_fp8_quant_fn<kv_cache_t>();
-      reshape_and_cache_fp8_vec_impl<scalar_t, qfn>(
-          key, value, reinterpret_cast<uint8_t*>(key_cache),
-          reinterpret_cast<uint8_t*>(value_cache), slot_mapping, token_num,
-          head_num, head_dim, block_size, key_token_num_stride,
-          key_head_num_stride, value_token_num_stride, value_head_num_stride,
-          num_blocks_stride, cache_head_num_stride, num_blocks_stride,
-          cache_head_num_stride, k_inv, v_inv);
-      return;
-    }
-
+      const int64_t block_size, const int64_t block_size_stride) {
 #pragma omp parallel for collapse(2)
    for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
      for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -294,9 +220,8 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
                                          token_idx * key_token_num_stride +
                                          head_idx * key_head_num_stride;
          scalar_t* key_cache_start_ptr =
-              reinterpret_cast<scalar_t*>(key_cache) +
-              block_idx * num_blocks_stride + head_idx * cache_head_num_stride +
-              block_offset;
+              key_cache + block_idx * num_blocks_stride +
+              head_idx * cache_head_num_stride + block_offset;

 #pragma GCC unroll 8
          for (int64_t i = 0, j = 0; i < head_dim; ++i, j += block_size) {
@@ -309,9 +234,8 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
                                            token_idx * value_token_num_stride +
                                            head_idx * value_head_num_stride;
          scalar_t* value_cache_start_ptr =
-              reinterpret_cast<scalar_t*>(value_cache) +
-              block_idx * num_blocks_stride + head_idx * cache_head_num_stride +
-              block_offset * head_dim;
+              value_cache + block_idx * num_blocks_stride +
+              head_idx * cache_head_num_stride + block_offset * head_dim;
          std::memcpy(value_cache_start_ptr, value_start_ptr,
                      sizeof(scalar_t) * head_dim);
        }
@@ -319,7 +243,6 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
    }
  }
 };
-
 }  // namespace cpu_attention

 #endif
@@ -116,9 +116,9 @@ class TileGemm161 {
 }  // namespace

 // This is a general but naive implementation based on vector instructions
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::VEC16, scalar_t, head_dim, kv_cache_scalar_t>
-    : public AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::VEC16, scalar_t, head_dim>
+    : public AttentionImpl<ISA::VEC, scalar_t, head_dim> {
 public:
  using query_t = scalar_t;
  using q_buffer_t = float;
@@ -244,8 +244,8 @@ class TileGemmS390X {

 }  // namespace

-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::VXE, scalar_t, head_dim, kv_cache_scalar_t> {
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::VXE, scalar_t, head_dim> {
 public:
  using query_t = scalar_t;
  using q_buffer_t = float;
@@ -342,8 +342,7 @@ class AttentionImpl<ISA::VXE, scalar_t, head_dim, kv_cache_scalar_t> {
      const int64_t head_num, const int64_t key_head_num_stride,
      const int64_t value_head_num_stride, const int64_t num_blocks,
      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
+      const int64_t block_size, const int64_t block_size_stride) {
 #pragma omp parallel for collapse(2)
    for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
      for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -15,9 +15,6 @@ using namespace at::vec;

 namespace vec_op {

-struct fp8_e4m3_tag {};
-struct fp8_e5m2_tag {};
-
 #define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
  AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)  \
@@ -325,9 +322,6 @@ struct BF16Vec32 : public VectorizedRegWrapper<BF16Vec32, 4, c10::BFloat16> {
    reg.val[2] = vec8_data.reg.val[0];
    reg.val[3] = vec8_data.reg.val[0];
  };
-
-  explicit BF16Vec32(const uint8_t*, fp8_e4m3_tag) : Base() {}
-  explicit BF16Vec32(const uint8_t*, fp8_e5m2_tag) : Base() {}
 };

 struct FP32Vec4 : public VectorizedRegWrapper<FP32Vec4, 1, float> {
@@ -1,25 +1,832 @@
 #ifndef CPU_TYPES_RISCV_HPP
 #define CPU_TYPES_RISCV_HPP

-// RISC-V Vector (RVV) CPU type definitions for vLLM.
-//
-// Supports multiple VLENs via compile-time dispatch. The compiler defines
-// __riscv_v_min_vlen from the zvl<N>b extension in -march. The defs header
-// maps VLEN to the correct LMUL suffixes, and the impl header provides
-// VLEN-independent class implementations.
-//
-// To add support for a new VLEN, add the LMUL mapping in
-// cpu_types_riscv_defs.hpp (the impl header needs no changes).
+#include <algorithm>
+#include <cmath>
+#include <cstring>
+#include <iostream>
+#include <limits>
+#include <riscv_vector.h>
+#include <torch/all.h>

-#ifndef __riscv_vector
-  #error "cpu_types_riscv.hpp included in a non-RVV translation unit"
+// ============================================================================
+// Vector Register Type Definitions (VLEN=128 bits)
+// ============================================================================
+
+typedef vfloat16m1_t fixed_vfloat16m1_t
+    __attribute__((riscv_rvv_vector_bits(128)));
+typedef vfloat16m2_t fixed_vfloat16m2_t
+    __attribute__((riscv_rvv_vector_bits(256)));
+
+typedef vfloat32m1_t fixed_vfloat32m1_t
+    __attribute__((riscv_rvv_vector_bits(128)));
+typedef vfloat32m2_t fixed_vfloat32m2_t
+    __attribute__((riscv_rvv_vector_bits(256)));
+typedef vfloat32m4_t fixed_vfloat32m4_t
+    __attribute__((riscv_rvv_vector_bits(512)));
+typedef vfloat32m8_t fixed_vfloat32m8_t
+    __attribute__((riscv_rvv_vector_bits(1024)));
+
+typedef vint32m2_t fixed_vint32m2_t __attribute__((riscv_rvv_vector_bits(256)));
+typedef vint32m4_t fixed_vint32m4_t __attribute__((riscv_rvv_vector_bits(512)));
+
+typedef vuint16m1_t fixed_vuint16m1_t
+    __attribute__((riscv_rvv_vector_bits(128)));
+typedef vuint16m2_t fixed_vuint16m2_t
+    __attribute__((riscv_rvv_vector_bits(256)));
+typedef vuint16m4_t fixed_vuint16m4_t
+    __attribute__((riscv_rvv_vector_bits(512)));
+
+#ifdef RISCV_BF16_SUPPORT
+typedef vbfloat16m1_t fixed_vbfloat16m1_t
+    __attribute__((riscv_rvv_vector_bits(128)));
+typedef vbfloat16m2_t fixed_vbfloat16m2_t
+    __attribute__((riscv_rvv_vector_bits(256)));
+typedef vbfloat16m4_t fixed_vbfloat16m4_t
+    __attribute__((riscv_rvv_vector_bits(512)));
 #endif

-#ifndef __riscv_v_min_vlen
-  #error "compiler did not define __riscv_v_min_vlen; pass -march=...zvl<N>b"
+namespace vec_op {
+
+#ifdef RISCV_BF16_SUPPORT
+  #define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
+    AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
+    AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)  \
+    AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__)
+#else
+  #define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
+    AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
+    AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)
 #endif

-#include "cpu_types_riscv_defs.hpp"
-#include "cpu_types_riscv_impl.hpp"
+#define VLLM_DISPATCH_FLOATING_TYPES(TYPE, NAME, ...) \
+  AT_DISPATCH_SWITCH(TYPE, NAME, VLLM_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__))

-#endif  // CPU_TYPES_RISCV_HPP
+#define FORCE_INLINE __attribute__((always_inline)) inline
+
+namespace {
+template <typename T, T... indexes, typename F>
+constexpr void unroll_loop_item(std::integer_sequence<T, indexes...>, F&& f) {
+  (f(std::integral_constant<T, indexes>{}), ...);
+};
+}  // namespace
+
+template <typename T, T count, typename F,
+          typename = std::enable_if_t<std::is_invocable_v<F, T>>>
+constexpr void unroll_loop(F&& f) {
+  unroll_loop_item(std::make_integer_sequence<T, count>{}, std::forward<F>(f));
+}
+
+template <typename T>
+struct Vec {
+  constexpr static int get_elem_num() { return T::VEC_ELEM_NUM; };
+};
+
+struct FP32Vec8;
+struct FP32Vec16;
+
+// ============================================================================
+// FP16 Implementation
+// ============================================================================
+
+struct FP16Vec8 : public Vec<FP16Vec8> {
+  constexpr static int VEC_ELEM_NUM = 8;
+  fixed_vfloat16m1_t reg;
+
+  explicit FP16Vec8(const void* ptr)
+      : reg(__riscv_vle16_v_f16m1(static_cast<const _Float16*>(ptr),
+                                  VEC_ELEM_NUM)) {};
+
+  explicit FP16Vec8(const FP32Vec8&);
+
+  void save(void* ptr) const {
+    __riscv_vse16_v_f16m1(static_cast<_Float16*>(ptr), reg, VEC_ELEM_NUM);
+  }
+  void save(void* ptr, int elem_num) const {
+    __riscv_vse16_v_f16m1(static_cast<_Float16*>(ptr), reg, elem_num);
+  }
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    ptrdiff_t byte_stride = stride * sizeof(_Float16);
+    __riscv_vsse16_v_f16m1(static_cast<_Float16*>(ptr), byte_stride, reg,
+                           VEC_ELEM_NUM);
+  }
+};
+
+struct FP16Vec16 : public Vec<FP16Vec16> {
+  constexpr static int VEC_ELEM_NUM = 16;
+  fixed_vfloat16m2_t reg;
+
+  explicit FP16Vec16(const void* ptr)
+      : reg(__riscv_vle16_v_f16m2(static_cast<const _Float16*>(ptr),
+                                  VEC_ELEM_NUM)) {};
+
+  explicit FP16Vec16(const FP32Vec16& vec);
+
+  void save(void* ptr) const {
+    __riscv_vse16_v_f16m2(static_cast<_Float16*>(ptr), reg, VEC_ELEM_NUM);
+  }
+  void save(void* ptr, int elem_num) const {
+    __riscv_vse16_v_f16m2(static_cast<_Float16*>(ptr), reg, elem_num);
+  }
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    ptrdiff_t byte_stride = stride * sizeof(_Float16);
+    __riscv_vsse16_v_f16m2(static_cast<_Float16*>(ptr), byte_stride, reg,
+                           VEC_ELEM_NUM);
+  }
+};
+
+// ============================================================================
+// BF16 Implementation
+// ============================================================================
+
+#ifdef RISCV_BF16_SUPPORT
+
+FORCE_INLINE fixed_vuint16m1_t bf16_to_u16(fixed_vbfloat16m1_t v) {
+  return __riscv_vreinterpret_v_bf16m1_u16m1(v);
+}
+FORCE_INLINE fixed_vuint16m2_t bf16_to_u16(fixed_vbfloat16m2_t v) {
+  return __riscv_vreinterpret_v_bf16m2_u16m2(v);
+}
+FORCE_INLINE fixed_vuint16m4_t bf16_to_u16(fixed_vbfloat16m4_t v) {
+  return __riscv_vreinterpret_v_bf16m4_u16m4(v);
+}
+
+struct BF16Vec8 : public Vec<BF16Vec8> {
+  constexpr static int VEC_ELEM_NUM = 8;
+  fixed_vbfloat16m1_t reg;
+
+  explicit BF16Vec8(const void* ptr)
+      : reg(__riscv_vreinterpret_v_u16m1_bf16m1(__riscv_vle16_v_u16m1(
+            reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};
+
+  explicit BF16Vec8(fixed_vbfloat16m1_t data) : reg(data) {};
+  explicit BF16Vec8(const FP32Vec8&);
+
+  void save(void* ptr) const {
+    __riscv_vse16_v_u16m1(reinterpret_cast<uint16_t*>(ptr), bf16_to_u16(reg),
+                          VEC_ELEM_NUM);
+  }
+  void save(void* ptr, int elem_num) const {
+    __riscv_vse16_v_u16m1(reinterpret_cast<uint16_t*>(ptr), bf16_to_u16(reg),
+                          elem_num);
+  }
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
+    __riscv_vsse16_v_u16m1(reinterpret_cast<uint16_t*>(ptr), byte_stride,
+                           bf16_to_u16(reg), VEC_ELEM_NUM);
+  }
+};
+
+struct BF16Vec16 : public Vec<BF16Vec16> {
+  constexpr static int VEC_ELEM_NUM = 16;
+  fixed_vbfloat16m2_t reg;
+
+  explicit BF16Vec16(const void* ptr)
+      : reg(__riscv_vreinterpret_v_u16m2_bf16m2(__riscv_vle16_v_u16m2(
+            reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};
+
+  explicit BF16Vec16(fixed_vbfloat16m2_t data) : reg(data) {};
+  explicit BF16Vec16(const FP32Vec16&);
+
+  void save(void* ptr) const {
+    __riscv_vse16_v_u16m2(reinterpret_cast<uint16_t*>(ptr), bf16_to_u16(reg),
+                          VEC_ELEM_NUM);
+  }
+  void save(void* ptr, int elem_num) const {
+    __riscv_vse16_v_u16m2(reinterpret_cast<uint16_t*>(ptr), bf16_to_u16(reg),
+                          elem_num);
+  }
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
+    __riscv_vsse16_v_u16m2(reinterpret_cast<uint16_t*>(ptr), byte_stride,
+                           bf16_to_u16(reg), VEC_ELEM_NUM);
+  }
+};
+
+struct BF16Vec32 : public Vec<BF16Vec32> {
+  constexpr static int VEC_ELEM_NUM = 32;
+  fixed_vbfloat16m4_t reg;
+
+  explicit BF16Vec32(const void* ptr)
+      : reg(__riscv_vreinterpret_v_u16m4_bf16m4(__riscv_vle16_v_u16m4(
+            reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};
+
+  explicit BF16Vec32(fixed_vbfloat16m4_t data) : reg(data) {};
+
+  explicit BF16Vec32(const BF16Vec8& v) {
+    fixed_vuint16m1_t u16_val = bf16_to_u16(v.reg);
+    fixed_vuint16m4_t u16_combined =
+        __riscv_vcreate_v_u16m1_u16m4(u16_val, u16_val, u16_val, u16_val);
+    reg = __riscv_vreinterpret_v_u16m4_bf16m4(u16_combined);
+  };
+
+  void save(void* ptr) const {
+    __riscv_vse16_v_u16m4(reinterpret_cast<uint16_t*>(ptr), bf16_to_u16(reg),
+                          VEC_ELEM_NUM);
+  }
+  void save(void* ptr, int elem_num) const {
+    __riscv_vse16_v_u16m4(reinterpret_cast<uint16_t*>(ptr), bf16_to_u16(reg),
+                          elem_num);
+  }
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
+    __riscv_vsse16_v_u16m4(reinterpret_cast<uint16_t*>(ptr), byte_stride,
+                           bf16_to_u16(reg), VEC_ELEM_NUM);
+  }
+};
+
+#else
+// ============================================================================
+// BF16 Fallback Implementation (FP32 Simulation)
+// ============================================================================
+
+struct BF16Vec8 : public Vec<BF16Vec8> {
+  constexpr static int VEC_ELEM_NUM = 8;
+  fixed_vfloat32m2_t reg_fp32;
+  explicit BF16Vec8(const void* ptr) {
+    const uint16_t* u16 = static_cast<const uint16_t*>(ptr);
+    float tmp[8];
+    for (int i = 0; i < 8; ++i) {
+      uint32_t v = static_cast<uint32_t>(u16[i]) << 16;
+      std::memcpy(&tmp[i], &v, 4);
+    }
+    reg_fp32 = __riscv_vle32_v_f32m2(tmp, 8);
+  }
+  explicit BF16Vec8(const FP32Vec8&);
+  void save(void* ptr) const {
+    float tmp[8];
+    __riscv_vse32_v_f32m2(tmp, reg_fp32, 8);
+    uint16_t* u16 = static_cast<uint16_t*>(ptr);
+    for (int i = 0; i < 8; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      u16[i] = static_cast<uint16_t>(v >> 16);
+    }
+  }
+  void save(void* ptr, int elem_num) const {
+    float tmp[8];
+    __riscv_vse32_v_f32m2(tmp, reg_fp32, 8);
+    uint16_t* u16 = static_cast<uint16_t*>(ptr);
+    for (int i = 0; i < elem_num; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      u16[i] = static_cast<uint16_t>(v >> 16);
+    }
+  }
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    float tmp[8];
+    __riscv_vse32_v_f32m2(tmp, reg_fp32, 8);
+    uint8_t* u8 = static_cast<uint8_t*>(ptr);
+    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
+    for (int i = 0; i < 8; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      uint16_t val = static_cast<uint16_t>(v >> 16);
+      *reinterpret_cast<uint16_t*>(u8 + i * byte_stride) = val;
+    }
+  }
+};
+
+struct BF16Vec16 : public Vec<BF16Vec16> {
+  constexpr static int VEC_ELEM_NUM = 16;
+  fixed_vfloat32m4_t reg_fp32;
+  explicit BF16Vec16(const void* ptr) {
+    const uint16_t* u16 = static_cast<const uint16_t*>(ptr);
+    float tmp[16];
+    for (int i = 0; i < 16; ++i) {
+      uint32_t v = static_cast<uint32_t>(u16[i]) << 16;
+      std::memcpy(&tmp[i], &v, 4);
+    }
+    reg_fp32 = __riscv_vle32_v_f32m4(tmp, 16);
+  }
+  explicit BF16Vec16(const FP32Vec16&);
+  void save(void* ptr) const {
+    float tmp[16];
+    __riscv_vse32_v_f32m4(tmp, reg_fp32, 16);
+    uint16_t* u16 = static_cast<uint16_t*>(ptr);
+    for (int i = 0; i < 16; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      u16[i] = static_cast<uint16_t>(v >> 16);
+    }
+  }
+  void save(void* ptr, int elem_num) const {
+    float tmp[16];
+    __riscv_vse32_v_f32m4(tmp, reg_fp32, 16);
+    uint16_t* u16 = static_cast<uint16_t*>(ptr);
+    for (int i = 0; i < elem_num; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      u16[i] = static_cast<uint16_t>(v >> 16);
+    }
+  }
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    float tmp[16];
+    __riscv_vse32_v_f32m4(tmp, reg_fp32, 16);
+    uint8_t* u8 = static_cast<uint8_t*>(ptr);
+    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
+    for (int i = 0; i < 16; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      uint16_t val = static_cast<uint16_t>(v >> 16);
+      *reinterpret_cast<uint16_t*>(u8 + i * byte_stride) = val;
+    }
+  }
+};
+
+struct BF16Vec32 : public Vec<BF16Vec32> {
+  constexpr static int VEC_ELEM_NUM = 32;
+  fixed_vfloat32m8_t reg_fp32;
+
+  explicit BF16Vec32(const void* ptr) {
+    const uint16_t* u16 = static_cast<const uint16_t*>(ptr);
+    float tmp[32];
+    for (int i = 0; i < 32; ++i) {
+      uint32_t v = static_cast<uint32_t>(u16[i]) << 16;
+      std::memcpy(&tmp[i], &v, 4);
+    }
+    reg_fp32 = __riscv_vle32_v_f32m8(tmp, 32);
+  }
+
+  explicit BF16Vec32(const BF16Vec8& v) {
+    float tmp_small[8];
+    __riscv_vse32_v_f32m2(tmp_small, v.reg_fp32, 8);
+    float tmp_large[32];
+    for (int i = 0; i < 4; ++i) {
+      std::memcpy(tmp_large + (i * 8), tmp_small, 8 * sizeof(float));
+    }
+    reg_fp32 = __riscv_vle32_v_f32m8(tmp_large, 32);
+  }
+
+  void save(void* ptr) const {
+    float tmp[32];
+    __riscv_vse32_v_f32m8(tmp, reg_fp32, 32);
+    uint16_t* u16 = static_cast<uint16_t*>(ptr);
+    for (int i = 0; i < 32; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      u16[i] = static_cast<uint16_t>(v >> 16);
+    }
+  }
+
+  void save(void* ptr, int elem_num) const {
+    float tmp[32];
+    __riscv_vse32_v_f32m8(tmp, reg_fp32, 32);
+    uint16_t* u16 = static_cast<uint16_t*>(ptr);
+    for (int i = 0; i < elem_num; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      u16[i] = static_cast<uint16_t>(v >> 16);
+    }
+  }
+
+  void save_strided(void* ptr, ptrdiff_t stride) const {
+    float tmp[32];
+    __riscv_vse32_v_f32m8(tmp, reg_fp32, 32);
+    uint8_t* u8 = static_cast<uint8_t*>(ptr);
+    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
+    for (int i = 0; i < 32; ++i) {
+      uint32_t v;
+      std::memcpy(&v, &tmp[i], 4);
+      uint16_t val = static_cast<uint16_t>(v >> 16);
+      *reinterpret_cast<uint16_t*>(u8 + i * byte_stride) = val;
+    }
+  }
+};
+#endif
+
+// ============================================================================
+// FP32 Implementation
+// ============================================================================
+
+struct FP32Vec4 : public Vec<FP32Vec4> {
+  constexpr static int VEC_ELEM_NUM = 4;
+  fixed_vfloat32m1_t reg;
+  explicit FP32Vec4(float v) : reg(__riscv_vfmv_v_f_f32m1(v, VEC_ELEM_NUM)) {};
+  explicit FP32Vec4() : reg(__riscv_vfmv_v_f_f32m1(0.0f, VEC_ELEM_NUM)) {};
+  explicit FP32Vec4(const float* ptr)
+      : reg(__riscv_vle32_v_f32m1(ptr, VEC_ELEM_NUM)) {};
+  explicit FP32Vec4(fixed_vfloat32m1_t data) : reg(data) {};
+  explicit FP32Vec4(const FP32Vec4& data) : reg(data.reg) {};
+  void save(float* ptr) const { __riscv_vse32_v_f32m1(ptr, reg, VEC_ELEM_NUM); }
+  void save(float* ptr, int elem_num) const {
+    __riscv_vse32_v_f32m1(ptr, reg, elem_num);
+  }
+};
+
+struct FP32Vec8 : public Vec<FP32Vec8> {
+  constexpr static int VEC_ELEM_NUM = 8;
+  fixed_vfloat32m2_t reg;
+
+  explicit FP32Vec8(float v) : reg(__riscv_vfmv_v_f_f32m2(v, VEC_ELEM_NUM)) {};
+  explicit FP32Vec8() : reg(__riscv_vfmv_v_f_f32m2(0.0f, VEC_ELEM_NUM)) {};
+  explicit FP32Vec8(const float* ptr)
+      : reg(__riscv_vle32_v_f32m2(ptr, VEC_ELEM_NUM)) {};
+  explicit FP32Vec8(fixed_vfloat32m2_t data) : reg(data) {};
+  explicit FP32Vec8(const FP32Vec8& data) : reg(data.reg) {};
+  explicit FP32Vec8(const FP16Vec8& v)
+      : reg(__riscv_vfwcvt_f_f_v_f32m2(v.reg, VEC_ELEM_NUM)) {};
+  explicit FP32Vec8(fixed_vfloat16m1_t v)
+      : reg(__riscv_vfwcvt_f_f_v_f32m2(v, VEC_ELEM_NUM)) {};
+
+#ifdef RISCV_BF16_SUPPORT
+  explicit FP32Vec8(fixed_vbfloat16m1_t v)
+      : reg(__riscv_vfwcvtbf16_f_f_v_f32m2(v, VEC_ELEM_NUM)) {};
+  explicit FP32Vec8(const BF16Vec8& v)
+      : reg(__riscv_vfwcvtbf16_f_f_v_f32m2(v.reg, VEC_ELEM_NUM)) {};
+#else
+  explicit FP32Vec8(const BF16Vec8& v) : reg(v.reg_fp32) {};
+#endif
+
+  float reduce_sum() const {
+    fixed_vfloat32m1_t scalar = __riscv_vfmv_s_f_f32m1(0.0f, 1);
+    scalar = __riscv_vfredusum_vs_f32m2_f32m1(reg, scalar, VEC_ELEM_NUM);
+    return __riscv_vfmv_f_s_f32m1_f32(scalar);
+  }
+
+  FP32Vec8 operator*(const FP32Vec8& b) const {
+    return FP32Vec8(__riscv_vfmul_vv_f32m2(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec8 operator+(const FP32Vec8& b) const {
+    return FP32Vec8(__riscv_vfadd_vv_f32m2(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec8 operator-(const FP32Vec8& b) const {
+    return FP32Vec8(__riscv_vfsub_vv_f32m2(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec8 operator/(const FP32Vec8& b) const {
+    return FP32Vec8(__riscv_vfdiv_vv_f32m2(reg, b.reg, VEC_ELEM_NUM));
+  }
+
+  FP32Vec8 min(const FP32Vec8& b) const {
+    return FP32Vec8(__riscv_vfmin_vv_f32m2(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec8 max(const FP32Vec8& b) const {
+    return FP32Vec8(__riscv_vfmax_vv_f32m2(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec8 abs() const {
+    return FP32Vec8(__riscv_vfabs_v_f32m2(reg, VEC_ELEM_NUM));
+  }
+
+  FP32Vec8 min(const FP32Vec8& b, int elem_num) const {
+    return FP32Vec8(__riscv_vfmin_vv_f32m2(reg, b.reg, elem_num));
+  }
+  FP32Vec8 max(const FP32Vec8& b, int elem_num) const {
+    return FP32Vec8(__riscv_vfmax_vv_f32m2(reg, b.reg, elem_num));
+  }
+
+  FP32Vec8 clamp(const FP32Vec8& min_v, const FP32Vec8& max_v) const {
+    fixed_vfloat32m2_t temp =
+        __riscv_vfmax_vv_f32m2(min_v.reg, reg, VEC_ELEM_NUM);
+    return FP32Vec8(__riscv_vfmin_vv_f32m2(max_v.reg, temp, VEC_ELEM_NUM));
+  }
+
+  void save(float* ptr) const { __riscv_vse32_v_f32m2(ptr, reg, VEC_ELEM_NUM); }
+  void save(float* ptr, int elem_num) const {
+    __riscv_vse32_v_f32m2(ptr, reg, elem_num);
+  }
+  void save_strided(float* ptr, ptrdiff_t stride) const {
+    ptrdiff_t byte_stride = stride * sizeof(float);
+    __riscv_vsse32_v_f32m2(ptr, byte_stride, reg, VEC_ELEM_NUM);
+  }
+
+  FP32Vec8 exp() const {
+    const float inv_ln2 = 1.44269504088896341f;
+    fixed_vfloat32m2_t x_scaled =
+        __riscv_vfmul_vf_f32m2(reg, inv_ln2, VEC_ELEM_NUM);
+    fixed_vint32m2_t n_int = __riscv_vfcvt_x_f_v_i32m2(x_scaled, VEC_ELEM_NUM);
+    fixed_vfloat32m2_t n_float = __riscv_vfcvt_f_x_v_f32m2(n_int, VEC_ELEM_NUM);
+
+    fixed_vfloat32m2_t r =
+        __riscv_vfsub_vv_f32m2(x_scaled, n_float, VEC_ELEM_NUM);
+
+    fixed_vfloat32m2_t poly =
+        __riscv_vfmv_v_f_f32m2(0.001333355810164f, VEC_ELEM_NUM);
+    poly = __riscv_vfmul_vv_f32m2(poly, r, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(poly, 0.009618129107628f, VEC_ELEM_NUM);
+    poly = __riscv_vfmul_vv_f32m2(poly, r, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(poly, 0.055504108664821f, VEC_ELEM_NUM);
+    poly = __riscv_vfmul_vv_f32m2(poly, r, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(poly, 0.240226506959101f, VEC_ELEM_NUM);
+    poly = __riscv_vfmul_vv_f32m2(poly, r, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(poly, 0.693147180559945f, VEC_ELEM_NUM);
+    poly = __riscv_vfmul_vv_f32m2(poly, r, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(poly, 1.0f, VEC_ELEM_NUM);
+
+    fixed_vint32m2_t biased_exp =
+        __riscv_vadd_vx_i32m2(n_int, 127, VEC_ELEM_NUM);
+    biased_exp = __riscv_vmax_vx_i32m2(biased_exp, 0, VEC_ELEM_NUM);
+    fixed_vint32m2_t exponent_bits =
+        __riscv_vsll_vx_i32m2(biased_exp, 23, VEC_ELEM_NUM);
+    fixed_vfloat32m2_t scale =
+        __riscv_vreinterpret_v_i32m2_f32m2(exponent_bits);
+
+    return FP32Vec8(__riscv_vfmul_vv_f32m2(poly, scale, VEC_ELEM_NUM));
+  }
+
+  FP32Vec8 tanh() const {
+    fixed_vfloat32m2_t x_clamped = __riscv_vfmin_vf_f32m2(
+        __riscv_vfmax_vf_f32m2(reg, -9.0f, VEC_ELEM_NUM), 9.0f, VEC_ELEM_NUM);
+    fixed_vfloat32m2_t x2 =
+        __riscv_vfmul_vf_f32m2(x_clamped, 2.0f, VEC_ELEM_NUM);
+    FP32Vec8 exp_val = FP32Vec8(x2).exp();
+    fixed_vfloat32m2_t num =
+        __riscv_vfsub_vf_f32m2(exp_val.reg, 1.0f, VEC_ELEM_NUM);
+    fixed_vfloat32m2_t den =
+        __riscv_vfadd_vf_f32m2(exp_val.reg, 1.0f, VEC_ELEM_NUM);
+    return FP32Vec8(__riscv_vfdiv_vv_f32m2(num, den, VEC_ELEM_NUM));
+  }
+
+  FP32Vec8 er() const {
+    const float p = 0.3275911f, a1 = 0.254829592f, a2 = -0.284496736f,
+                a3 = 1.421413741f, a4 = -1.453152027f, a5 = 1.061405429f;
+    fixed_vfloat32m2_t abs_x = __riscv_vfabs_v_f32m2(reg, VEC_ELEM_NUM);
+
+    fixed_vfloat32m2_t t = __riscv_vfadd_vf_f32m2(
+        __riscv_vfmul_vf_f32m2(abs_x, p, VEC_ELEM_NUM), 1.0f, VEC_ELEM_NUM);
+    t = __riscv_vfrdiv_vf_f32m2(t, 1.0f, VEC_ELEM_NUM);
+
+    fixed_vfloat32m2_t poly = __riscv_vfmv_v_f_f32m2(a5, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(__riscv_vfmul_vv_f32m2(poly, t, VEC_ELEM_NUM),
+                                  a4, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(__riscv_vfmul_vv_f32m2(poly, t, VEC_ELEM_NUM),
+                                  a3, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(__riscv_vfmul_vv_f32m2(poly, t, VEC_ELEM_NUM),
+                                  a2, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m2(__riscv_vfmul_vv_f32m2(poly, t, VEC_ELEM_NUM),
+                                  a1, VEC_ELEM_NUM);
+    poly = __riscv_vfmul_vv_f32m2(poly, t, VEC_ELEM_NUM);
+
+    fixed_vfloat32m2_t exp_val =
+        FP32Vec8(__riscv_vfneg_v_f32m2(
+                     __riscv_vfmul_vv_f32m2(abs_x, abs_x, VEC_ELEM_NUM),
+                     VEC_ELEM_NUM))
+            .exp()
+            .reg;
+    fixed_vfloat32m2_t res = __riscv_vfrsub_vf_f32m2(
+        __riscv_vfmul_vv_f32m2(poly, exp_val, VEC_ELEM_NUM), 1.0f,
+        VEC_ELEM_NUM);
+
+    vbool16_t mask = __riscv_vmflt_vf_f32m2_b16(reg, 0.0f, VEC_ELEM_NUM);
+    return FP32Vec8(__riscv_vfneg_v_f32m2_m(mask, res, VEC_ELEM_NUM));
+  }
+};
+
+struct FP32Vec16 : public Vec<FP32Vec16> {
+  constexpr static int VEC_ELEM_NUM = 16;
+  fixed_vfloat32m4_t reg;
+
+  explicit FP32Vec16(float v) : reg(__riscv_vfmv_v_f_f32m4(v, VEC_ELEM_NUM)) {};
+  explicit FP32Vec16() : reg(__riscv_vfmv_v_f_f32m4(0.0f, VEC_ELEM_NUM)) {};
+  explicit FP32Vec16(const float* ptr)
+      : reg(__riscv_vle32_v_f32m4(ptr, VEC_ELEM_NUM)) {};
+  explicit FP32Vec16(fixed_vfloat32m4_t data) : reg(data) {};
+  explicit FP32Vec16(const FP32Vec8& data)
+      : reg(__riscv_vcreate_v_f32m2_f32m4(data.reg, data.reg)) {};
+  explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
+  explicit FP32Vec16(const FP16Vec16& v);
+
+#ifdef RISCV_BF16_SUPPORT
+  explicit FP32Vec16(fixed_vbfloat16m2_t v)
+      : reg(__riscv_vfwcvtbf16_f_f_v_f32m4(v, VEC_ELEM_NUM)) {};
+  explicit FP32Vec16(const BF16Vec16& v)
+      : reg(__riscv_vfwcvtbf16_f_f_v_f32m4(v.reg, VEC_ELEM_NUM)) {};
+#else
+  explicit FP32Vec16(const BF16Vec16& v) : reg(v.reg_fp32) {};
+#endif
+
+  FP32Vec16 operator+(const FP32Vec16& b) const {
+    return FP32Vec16(__riscv_vfadd_vv_f32m4(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec16 operator-(const FP32Vec16& b) const {
+    return FP32Vec16(__riscv_vfsub_vv_f32m4(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec16 operator*(const FP32Vec16& b) const {
+    return FP32Vec16(__riscv_vfmul_vv_f32m4(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec16 operator/(const FP32Vec16& b) const {
+    return FP32Vec16(__riscv_vfdiv_vv_f32m4(reg, b.reg, VEC_ELEM_NUM));
+  }
+
+  FP32Vec16 fma(const FP32Vec16& a, const FP32Vec16& b) const {
+    return FP32Vec16(__riscv_vfmacc_vv_f32m4(reg, a.reg, b.reg, VEC_ELEM_NUM));
+  }
+
+  float reduce_sum() const {
+    fixed_vfloat32m1_t scalar = __riscv_vfmv_s_f_f32m1(0.0f, 1);
+    scalar = __riscv_vfredusum_vs_f32m4_f32m1(reg, scalar, VEC_ELEM_NUM);
+    return __riscv_vfmv_f_s_f32m1_f32(scalar);
+  }
+
+  float reduce_max() const {
+    fixed_vfloat32m1_t scalar =
+        __riscv_vfmv_s_f_f32m1(std::numeric_limits<float>::lowest(), 1);
+    scalar = __riscv_vfredmax_vs_f32m4_f32m1(reg, scalar, VEC_ELEM_NUM);
+    return __riscv_vfmv_f_s_f32m1_f32(scalar);
+  }
+
+  float reduce_min() const {
+    fixed_vfloat32m1_t scalar =
+        __riscv_vfmv_s_f_f32m1(std::numeric_limits<float>::max(), 1);
+    scalar = __riscv_vfredmin_vs_f32m4_f32m1(reg, scalar, VEC_ELEM_NUM);
+    return __riscv_vfmv_f_s_f32m1_f32(scalar);
+  }
+
+  template <int group_size>
+  float reduce_sub_sum(int idx) {
+    static_assert(VEC_ELEM_NUM % group_size == 0);
+    const int start = idx * group_size;
+    vuint32m4_t indices = __riscv_vid_v_u32m4(VEC_ELEM_NUM);
+    vbool8_t mask = __riscv_vmand_mm_b8(
+        __riscv_vmsgeu_vx_u32m4_b8(indices, start, VEC_ELEM_NUM),
+        __riscv_vmsltu_vx_u32m4_b8(indices, start + group_size, VEC_ELEM_NUM),
+        VEC_ELEM_NUM);
+    fixed_vfloat32m1_t scalar = __riscv_vfmv_s_f_f32m1(0.0f, 1);
+    scalar =
+        __riscv_vfredusum_vs_f32m4_f32m1_m(mask, reg, scalar, VEC_ELEM_NUM);
+    return __riscv_vfmv_f_s_f32m1_f32(scalar);
+  };
+
+  FP32Vec16 max(const FP32Vec16& b) const {
+    return FP32Vec16(__riscv_vfmax_vv_f32m4(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec16 min(const FP32Vec16& b) const {
+    return FP32Vec16(__riscv_vfmin_vv_f32m4(reg, b.reg, VEC_ELEM_NUM));
+  }
+  FP32Vec16 abs() const {
+    return FP32Vec16(__riscv_vfabs_v_f32m4(reg, VEC_ELEM_NUM));
+  }
+
+  FP32Vec16 clamp(const FP32Vec16& min_v, const FP32Vec16& max_v) const {
+    return FP32Vec16(__riscv_vfmin_vv_f32m4(
+        max_v.reg, __riscv_vfmax_vv_f32m4(min_v.reg, reg, VEC_ELEM_NUM),
+        VEC_ELEM_NUM));
+  }
+
+  void save(float* ptr) const { __riscv_vse32_v_f32m4(ptr, reg, VEC_ELEM_NUM); }
+  void save(float* ptr, int elem_num) const {
+    __riscv_vse32_v_f32m4(ptr, reg, elem_num);
+  }
+  void save_strided(float* ptr, ptrdiff_t stride) const {
+    ptrdiff_t byte_stride = stride * sizeof(float);
+    __riscv_vsse32_v_f32m4(ptr, byte_stride, reg, VEC_ELEM_NUM);
+  }
+
+  FP32Vec16 exp() const {
+    const float inv_ln2 = 1.44269504088896341f;
+    fixed_vfloat32m4_t x_scaled =
+        __riscv_vfmul_vf_f32m4(reg, inv_ln2, VEC_ELEM_NUM);
+    fixed_vint32m4_t n_int = __riscv_vfcvt_x_f_v_i32m4(x_scaled, VEC_ELEM_NUM);
+    fixed_vfloat32m4_t n_float = __riscv_vfcvt_f_x_v_f32m4(n_int, VEC_ELEM_NUM);
+    fixed_vfloat32m4_t r =
+        __riscv_vfsub_vv_f32m4(x_scaled, n_float, VEC_ELEM_NUM);
+
+    fixed_vfloat32m4_t poly =
+        __riscv_vfmv_v_f_f32m4(0.001333355810164f, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, r, VEC_ELEM_NUM),
+                                  0.009618129107628f, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, r, VEC_ELEM_NUM),
+                                  0.055504108664821f, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, r, VEC_ELEM_NUM),
+                                  0.240226506959101f, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, r, VEC_ELEM_NUM),
+                                  0.693147180559945f, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, r, VEC_ELEM_NUM),
+                                  1.0f, VEC_ELEM_NUM);
+
+    fixed_vint32m4_t biased_exp = __riscv_vmax_vx_i32m4(
+        __riscv_vadd_vx_i32m4(n_int, 127, VEC_ELEM_NUM), 0, VEC_ELEM_NUM);
+    fixed_vfloat32m4_t scale = __riscv_vreinterpret_v_i32m4_f32m4(
+        __riscv_vsll_vx_i32m4(biased_exp, 23, VEC_ELEM_NUM));
+
+    return FP32Vec16(__riscv_vfmul_vv_f32m4(poly, scale, VEC_ELEM_NUM));
+  }
+
+  FP32Vec16 tanh() const {
+    fixed_vfloat32m4_t x_clamped = __riscv_vfmin_vf_f32m4(
+        __riscv_vfmax_vf_f32m4(reg, -9.0f, VEC_ELEM_NUM), 9.0f, VEC_ELEM_NUM);
+    FP32Vec16 exp_val =
+        FP32Vec16(__riscv_vfmul_vf_f32m4(x_clamped, 2.0f, VEC_ELEM_NUM)).exp();
+    return FP32Vec16(__riscv_vfdiv_vv_f32m4(
+        __riscv_vfsub_vf_f32m4(exp_val.reg, 1.0f, VEC_ELEM_NUM),
+        __riscv_vfadd_vf_f32m4(exp_val.reg, 1.0f, VEC_ELEM_NUM), VEC_ELEM_NUM));
+  }
+
+  FP32Vec16 er() const {
+    const float p = 0.3275911f, a1 = 0.254829592f, a2 = -0.284496736f,
+                a3 = 1.421413741f, a4 = -1.453152027f, a5 = 1.061405429f;
+    fixed_vfloat32m4_t abs_x = __riscv_vfabs_v_f32m4(reg, VEC_ELEM_NUM);
+    fixed_vfloat32m4_t t = __riscv_vfrdiv_vf_f32m4(
+        __riscv_vfadd_vf_f32m4(__riscv_vfmul_vf_f32m4(abs_x, p, VEC_ELEM_NUM),
+                               1.0f, VEC_ELEM_NUM),
+        1.0f, VEC_ELEM_NUM);
+
+    fixed_vfloat32m4_t poly = __riscv_vfmv_v_f_f32m4(a5, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, t, VEC_ELEM_NUM),
+                                  a4, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, t, VEC_ELEM_NUM),
+                                  a3, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, t, VEC_ELEM_NUM),
+                                  a2, VEC_ELEM_NUM);
+    poly = __riscv_vfadd_vf_f32m4(__riscv_vfmul_vv_f32m4(poly, t, VEC_ELEM_NUM),
+                                  a1, VEC_ELEM_NUM);
+    poly = __riscv_vfmul_vv_f32m4(poly, t, VEC_ELEM_NUM);
+
+    fixed_vfloat32m4_t exp_val =
+        FP32Vec16(__riscv_vfneg_v_f32m4(
+                      __riscv_vfmul_vv_f32m4(abs_x, abs_x, VEC_ELEM_NUM),
+                      VEC_ELEM_NUM))
+            .exp()
+            .reg;
+    fixed_vfloat32m4_t res = __riscv_vfrsub_vf_f32m4(
+        __riscv_vfmul_vv_f32m4(poly, exp_val, VEC_ELEM_NUM), 1.0f,
+        VEC_ELEM_NUM);
+
+    vbool8_t mask = __riscv_vmflt_vf_f32m4_b8(reg, 0.0f, VEC_ELEM_NUM);
+    return FP32Vec16(__riscv_vfneg_v_f32m4_m(mask, res, VEC_ELEM_NUM));
+  }
+};
+
+// ============================================================================
+// Type Traits & Global Helpers
+// ============================================================================
+
+template <typename T>
+struct VecType {
+  using vec_type = void;
+  using vec_t = void;
+};
+
+template <typename T>
+using vec_t = typename VecType<T>::vec_type;
+
+template <>
+struct VecType<float> {
+  using vec_type = FP32Vec8;
+  using vec_t = FP32Vec8;
+};
+template <>
+struct VecType<c10::Half> {
+  using vec_type = FP16Vec8;
+  using vec_t = FP16Vec8;
+};
+template <>
+struct VecType<c10::BFloat16> {
+  using vec_type = BF16Vec8;
+  using vec_t = BF16Vec8;
+};
+
+template <typename T>
+void storeFP32(float v, T* ptr) {
+  *ptr = v;
+}
+template <>
+inline void storeFP32<c10::Half>(float v, c10::Half* ptr) {
+  *reinterpret_cast<_Float16*>(ptr) = static_cast<_Float16>(v);
+}
+
+inline FP16Vec16::FP16Vec16(const FP32Vec16& v) {
+  reg = __riscv_vfncvt_f_f_w_f16m2(v.reg, VEC_ELEM_NUM);
+}
+inline FP16Vec8::FP16Vec8(const FP32Vec8& v) {
+  reg = __riscv_vfncvt_f_f_w_f16m1(v.reg, VEC_ELEM_NUM);
+}
+inline FP32Vec16::FP32Vec16(const FP16Vec16& v) {
+  reg = __riscv_vfwcvt_f_f_v_f32m4(v.reg, VEC_ELEM_NUM);
+}
+inline void fma(FP32Vec16& acc, const FP32Vec16& a, const FP32Vec16& b) {
+  acc = acc.fma(a, b);
+}
+
+#ifdef RISCV_BF16_SUPPORT
+template <>
+inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
+  *ptr = static_cast<__bf16>(v);
+};
+inline BF16Vec8::BF16Vec8(const FP32Vec8& v)
+    : reg(__riscv_vfncvtbf16_f_f_w_bf16m1(v.reg, VEC_ELEM_NUM)) {};
+inline BF16Vec16::BF16Vec16(const FP32Vec16& v)
+    : reg(__riscv_vfncvtbf16_f_f_w_bf16m2(v.reg, VEC_ELEM_NUM)) {};
+#else
+template <>
+inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
+  uint32_t val;
+  std::memcpy(&val, &v, 4);
+  *reinterpret_cast<uint16_t*>(ptr) = static_cast<uint16_t>(val >> 16);
+}
+inline BF16Vec8::BF16Vec8(const FP32Vec8& v) : reg_fp32(v.reg) {}
+inline BF16Vec16::BF16Vec16(const FP32Vec16& v) : reg_fp32(v.reg) {}
+#endif
+
+inline void prefetch(const void* addr) { __builtin_prefetch(addr, 0, 1); }
+
+}  // namespace vec_op
+
+#ifndef CPU_KERNEL_GUARD_IN
+  #define CPU_KERNEL_GUARD_IN(NAME)
+#endif
+
+#ifndef CPU_KERNEL_GUARD_OUT
+  #define CPU_KERNEL_GUARD_OUT(NAME)
+#endif
+
+#endif  // CPU_TYPES_RISCV_HPP
@@ -1,98 +0,0 @@
-#ifndef CPU_TYPES_RISCV_DEFS_HPP
-#define CPU_TYPES_RISCV_DEFS_HPP
-
-// VLEN-to-LMUL mapping for RISC-V Vector extension.
-//
-// LMUL_<N> expands to the LMUL suffix giving N total bits of vector data:
-//   VLEN=128: LMUL_128=m1,  LMUL_256=m2,  LMUL_512=m4,  LMUL_1024=m8
-//   VLEN=256: LMUL_128=mf2, LMUL_256=m1,  LMUL_512=m2,  LMUL_1024=m4
-
-#include <riscv_vector.h>
-
-#if __riscv_v_min_vlen == 128
-  #define LMUL_128 m1
-  #define LMUL_256 m2
-  #define LMUL_512 m4
-  #define LMUL_1024 m8
-  #define BOOL_256 b16
-  #define BOOL_512 b8
-#elif __riscv_v_min_vlen == 256
-  #define LMUL_128 mf2
-  #define LMUL_256 m1
-  #define LMUL_512 m2
-  #define LMUL_1024 m4
-  #define BOOL_256 b32
-  #define BOOL_512 b16
-#else
-  #error "cpu_types_riscv_defs.hpp: unsupported __riscv_v_min_vlen"
-#endif
-
-// Token-paste helpers.
-#define _RVV_P2(a, b) a##b
-#define _RVV_P3(a, b, c) a##b##c
-#define _RVV_P4(a, b, c, d) a##b##c##d
-#define RVVTYPE(base, lmul, suffix) _RVV_P3(base, lmul, suffix)
-#define RVVI(base, lmul) _RVV_P2(base, lmul)
-#define RVVI3(base, lmul, suffix) _RVV_P3(base, lmul, suffix)
-#define RVVI4(a, b, c, d) _RVV_P4(a, b, c, d)
-// For mask intrinsics: RVVIB(base, LMUL_256, BOOL_256) → base##m2##_##b16
-#define _RVV_PB(base, lmul, btype) base##lmul##_##btype
-#define RVVIB(base, lmul, btype) _RVV_PB(base, lmul, btype)
-
-// ---- Semantic fixed-vector typedefs (named by element count) ----
-
-// float16
-typedef RVVTYPE(vfloat16, LMUL_128, _t) fixed_fp16x8_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef RVVTYPE(vfloat16, LMUL_256, _t) fixed_fp16x16_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-
-// float32
-typedef RVVTYPE(vfloat32, LMUL_128, _t) fixed_fp32x4_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef RVVTYPE(vfloat32, LMUL_256, _t) fixed_fp32x8_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-typedef RVVTYPE(vfloat32, LMUL_512, _t) fixed_fp32x16_t
-    __attribute__((riscv_rvv_vector_bits(512)));
-typedef RVVTYPE(vfloat32, LMUL_1024, _t) fixed_fp32x32_t
-    __attribute__((riscv_rvv_vector_bits(1024)));
-
-// int32
-typedef RVVTYPE(vint32, LMUL_256, _t) fixed_i32x8_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-typedef RVVTYPE(vint32, LMUL_512, _t) fixed_i32x16_t
-    __attribute__((riscv_rvv_vector_bits(512)));
-
-// uint16
-typedef RVVTYPE(vuint16, LMUL_128, _t) fixed_u16x8_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef RVVTYPE(vuint16, LMUL_256, _t) fixed_u16x16_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-typedef RVVTYPE(vuint16, LMUL_512, _t) fixed_u16x32_t
-    __attribute__((riscv_rvv_vector_bits(512)));
-
-// bfloat16
-#ifdef RISCV_BF16_SUPPORT
-typedef RVVTYPE(vbfloat16, LMUL_128, _t) fixed_bf16x8_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef RVVTYPE(vbfloat16, LMUL_256, _t) fixed_bf16x16_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-typedef RVVTYPE(vbfloat16, LMUL_512, _t) fixed_bf16x32_t
-    __attribute__((riscv_rvv_vector_bits(512)));
-#endif
-
-// ---- Reduction accumulator type (always m1 = one register of f32) ----
-// Used for scalar reductions; only element [0] is meaningful.
-typedef vfloat32m1_t rvv_f32_accum_t
-    __attribute__((riscv_rvv_vector_bits(__riscv_v_min_vlen)));
-
-// ---- Mask types for f32 elements ----
-#if __riscv_v_min_vlen == 128
-typedef vbool16_t rvv_mask_f32x8_t;
-typedef vbool8_t rvv_mask_f32x16_t;
-#elif __riscv_v_min_vlen == 256
-typedef vbool32_t rvv_mask_f32x8_t;
-typedef vbool16_t rvv_mask_f32x16_t;
-#endif
-
-#endif  // CPU_TYPES_RISCV_DEFS_HPP
@@ -1,905 +0,0 @@
-#ifndef CPU_TYPES_RISCV_IMPL_HPP
-#define CPU_TYPES_RISCV_IMPL_HPP
-
-// Shared implementation of RVV vector-type wrapper classes.
-// This file is VLEN-independent: it uses the semantic type names and
-// RVVI() intrinsic macros from cpu_types_riscv_defs.hpp.
-//
-// DO NOT include this file directly; include cpu_types_riscv.hpp instead.
-
-#include <algorithm>
-#include <cmath>
-#include <cstring>
-#include <iostream>
-#include <limits>
-#include <torch/all.h>
-namespace vec_op {
-
-// BFloat16 is always supported on RISC-V: natively when RISCV_BF16_SUPPORT
-// is defined, otherwise via the FP32-simulation fallback path.
-#define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
-  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
-  AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)  \
-  AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__)
-
-#define VLLM_DISPATCH_FLOATING_TYPES(TYPE, NAME, ...) \
-  AT_DISPATCH_SWITCH(TYPE, NAME, VLLM_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__))
-
-#define FORCE_INLINE __attribute__((always_inline)) inline
-
-namespace {
-template <typename T, T... indexes, typename F>
-constexpr void unroll_loop_item(std::integer_sequence<T, indexes...>, F&& f) {
-  (f(std::integral_constant<T, indexes>{}), ...);
-};
-}  // namespace
-
-template <typename T, T count, typename F,
-          typename = std::enable_if_t<std::is_invocable_v<F, T>>>
-constexpr void unroll_loop(F&& f) {
-  unroll_loop_item(std::make_integer_sequence<T, count>{}, std::forward<F>(f));
-}
-
-template <typename T>
-struct Vec {
-  constexpr static int get_elem_num() { return T::VEC_ELEM_NUM; };
-};
-
-struct FP32Vec8;
-struct FP32Vec16;
-
-// ============================================================================
-// FP16 Implementation
-// ============================================================================
-
-struct FP16Vec8 : public Vec<FP16Vec8> {
-  constexpr static int VEC_ELEM_NUM = 8;
-  fixed_fp16x8_t reg;
-
-  explicit FP16Vec8(const void* ptr)
-      : reg(RVVI(__riscv_vle16_v_f16, LMUL_128)(
-            static_cast<const _Float16*>(ptr), VEC_ELEM_NUM)) {};
-
-  explicit FP16Vec8(const FP32Vec8&);
-
-  void save(void* ptr) const {
-    RVVI(__riscv_vse16_v_f16, LMUL_128)(static_cast<_Float16*>(ptr), reg,
-                                        VEC_ELEM_NUM);
-  }
-  void save(void* ptr, int elem_num) const {
-    RVVI(__riscv_vse16_v_f16, LMUL_128)(static_cast<_Float16*>(ptr), reg,
-                                        elem_num);
-  }
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    ptrdiff_t byte_stride = stride * sizeof(_Float16);
-    RVVI(__riscv_vsse16_v_f16, LMUL_128)(static_cast<_Float16*>(ptr),
-                                         byte_stride, reg, VEC_ELEM_NUM);
-  }
-};
-
-struct FP16Vec16 : public Vec<FP16Vec16> {
-  constexpr static int VEC_ELEM_NUM = 16;
-  fixed_fp16x16_t reg;
-
-  explicit FP16Vec16(const void* ptr)
-      : reg(RVVI(__riscv_vle16_v_f16, LMUL_256)(
-            static_cast<const _Float16*>(ptr), VEC_ELEM_NUM)) {};
-
-  explicit FP16Vec16(const FP32Vec16& vec);
-
-  void save(void* ptr) const {
-    RVVI(__riscv_vse16_v_f16, LMUL_256)(static_cast<_Float16*>(ptr), reg,
-                                        VEC_ELEM_NUM);
-  }
-  void save(void* ptr, int elem_num) const {
-    RVVI(__riscv_vse16_v_f16, LMUL_256)(static_cast<_Float16*>(ptr), reg,
-                                        elem_num);
-  }
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    ptrdiff_t byte_stride = stride * sizeof(_Float16);
-    RVVI(__riscv_vsse16_v_f16, LMUL_256)(static_cast<_Float16*>(ptr),
-                                         byte_stride, reg, VEC_ELEM_NUM);
-  }
-};
-
-// ============================================================================
-// BF16 Implementation
-// ============================================================================
-
-#ifdef RISCV_BF16_SUPPORT
-
-FORCE_INLINE fixed_u16x8_t bf16_to_u16(fixed_bf16x8_t v) {
-  return RVVI4(__riscv_vreinterpret_v_bf16, LMUL_128, _u16, LMUL_128)(v);
-}
-FORCE_INLINE fixed_u16x16_t bf16_to_u16(fixed_bf16x16_t v) {
-  return RVVI4(__riscv_vreinterpret_v_bf16, LMUL_256, _u16, LMUL_256)(v);
-}
-FORCE_INLINE fixed_u16x32_t bf16_to_u16(fixed_bf16x32_t v) {
-  return RVVI4(__riscv_vreinterpret_v_bf16, LMUL_512, _u16, LMUL_512)(v);
-}
-
-struct BF16Vec8 : public Vec<BF16Vec8> {
-  constexpr static int VEC_ELEM_NUM = 8;
-  fixed_bf16x8_t reg;
-
-  explicit BF16Vec8(const void* ptr)
-      : reg(RVVI4(__riscv_vreinterpret_v_u16, LMUL_128, _bf16,
-                  LMUL_128)(RVVI(__riscv_vle16_v_u16, LMUL_128)(
-            reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};
-
-  explicit BF16Vec8(fixed_bf16x8_t data) : reg(data) {};
-  explicit BF16Vec8(const FP32Vec8&);
-
-  void save(void* ptr) const {
-    RVVI(__riscv_vse16_v_u16, LMUL_128)(reinterpret_cast<uint16_t*>(ptr),
-                                        bf16_to_u16(reg), VEC_ELEM_NUM);
-  }
-  void save(void* ptr, int elem_num) const {
-    RVVI(__riscv_vse16_v_u16, LMUL_128)(reinterpret_cast<uint16_t*>(ptr),
-                                        bf16_to_u16(reg), elem_num);
-  }
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
-    RVVI(__riscv_vsse16_v_u16, LMUL_128)(reinterpret_cast<uint16_t*>(ptr),
-                                         byte_stride, bf16_to_u16(reg),
-                                         VEC_ELEM_NUM);
-  }
-};
-
-struct BF16Vec16 : public Vec<BF16Vec16> {
-  constexpr static int VEC_ELEM_NUM = 16;
-  fixed_bf16x16_t reg;
-
-  explicit BF16Vec16(const void* ptr)
-      : reg(RVVI4(__riscv_vreinterpret_v_u16, LMUL_256, _bf16,
-                  LMUL_256)(RVVI(__riscv_vle16_v_u16, LMUL_256)(
-            reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};
-
-  explicit BF16Vec16(fixed_bf16x16_t data) : reg(data) {};
-  explicit BF16Vec16(const FP32Vec16&);
-
-  void save(void* ptr) const {
-    RVVI(__riscv_vse16_v_u16, LMUL_256)(reinterpret_cast<uint16_t*>(ptr),
-                                        bf16_to_u16(reg), VEC_ELEM_NUM);
-  }
-  void save(void* ptr, int elem_num) const {
-    RVVI(__riscv_vse16_v_u16, LMUL_256)(reinterpret_cast<uint16_t*>(ptr),
-                                        bf16_to_u16(reg), elem_num);
-  }
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
-    RVVI(__riscv_vsse16_v_u16, LMUL_256)(reinterpret_cast<uint16_t*>(ptr),
-                                         byte_stride, bf16_to_u16(reg),
-                                         VEC_ELEM_NUM);
-  }
-};
-
-struct BF16Vec32 : public Vec<BF16Vec32> {
-  constexpr static int VEC_ELEM_NUM = 32;
-  fixed_bf16x32_t reg;
-
-  explicit BF16Vec32(const void* ptr)
-      : reg(RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16,
-                  LMUL_512)(RVVI(__riscv_vle16_v_u16, LMUL_512)(
-            reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};
-
-  explicit BF16Vec32(fixed_bf16x32_t data) : reg(data) {};
-
-  explicit BF16Vec32(const BF16Vec8& v) {
-    fixed_u16x8_t u16_val = bf16_to_u16(v.reg);
-    fixed_u16x32_t u16_combined =
-        RVVI4(__riscv_vcreate_v_u16, LMUL_128, _u16, LMUL_512)(
-            u16_val, u16_val, u16_val, u16_val);
-    reg = RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16,
-                LMUL_512)(u16_combined);
-  };
-
-  void save(void* ptr) const {
-    RVVI(__riscv_vse16_v_u16, LMUL_512)(reinterpret_cast<uint16_t*>(ptr),
-                                        bf16_to_u16(reg), VEC_ELEM_NUM);
-  }
-  void save(void* ptr, int elem_num) const {
-    RVVI(__riscv_vse16_v_u16, LMUL_512)(reinterpret_cast<uint16_t*>(ptr),
-                                        bf16_to_u16(reg), elem_num);
-  }
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
-    RVVI(__riscv_vsse16_v_u16, LMUL_512)(reinterpret_cast<uint16_t*>(ptr),
-                                         byte_stride, bf16_to_u16(reg),
-                                         VEC_ELEM_NUM);
-  }
-};
-
-#else
-// ============================================================================
-// BF16 Fallback Implementation (FP32 Simulation)
-// ============================================================================
-
-struct BF16Vec8 : public Vec<BF16Vec8> {
-  constexpr static int VEC_ELEM_NUM = 8;
-  fixed_fp32x8_t reg_fp32;
-  explicit BF16Vec8(const void* ptr) {
-    const uint16_t* u16 = static_cast<const uint16_t*>(ptr);
-    float tmp[8];
-    for (int i = 0; i < 8; ++i) {
-      uint32_t v = static_cast<uint32_t>(u16[i]) << 16;
-      std::memcpy(&tmp[i], &v, 4);
-    }
-    reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_256)(tmp, 8);
-  }
-  explicit BF16Vec8(const FP32Vec8&);
-  void save(void* ptr) const {
-    float tmp[8];
-    RVVI(__riscv_vse32_v_f32, LMUL_256)(tmp, reg_fp32, 8);
-    uint16_t* u16 = static_cast<uint16_t*>(ptr);
-    for (int i = 0; i < 8; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      u16[i] = static_cast<uint16_t>(v >> 16);
-    }
-  }
-  void save(void* ptr, int elem_num) const {
-    float tmp[8];
-    RVVI(__riscv_vse32_v_f32, LMUL_256)(tmp, reg_fp32, 8);
-    uint16_t* u16 = static_cast<uint16_t*>(ptr);
-    for (int i = 0; i < elem_num; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      u16[i] = static_cast<uint16_t>(v >> 16);
-    }
-  }
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    float tmp[8];
-    RVVI(__riscv_vse32_v_f32, LMUL_256)(tmp, reg_fp32, 8);
-    uint8_t* u8 = static_cast<uint8_t*>(ptr);
-    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
-    for (int i = 0; i < 8; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      uint16_t val = static_cast<uint16_t>(v >> 16);
-      *reinterpret_cast<uint16_t*>(u8 + i * byte_stride) = val;
-    }
-  }
-};
-
-struct BF16Vec16 : public Vec<BF16Vec16> {
-  constexpr static int VEC_ELEM_NUM = 16;
-  fixed_fp32x16_t reg_fp32;
-  explicit BF16Vec16(const void* ptr) {
-    const uint16_t* u16 = static_cast<const uint16_t*>(ptr);
-    float tmp[16];
-    for (int i = 0; i < 16; ++i) {
-      uint32_t v = static_cast<uint32_t>(u16[i]) << 16;
-      std::memcpy(&tmp[i], &v, 4);
-    }
-    reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_512)(tmp, 16);
-  }
-  explicit BF16Vec16(const FP32Vec16&);
-  void save(void* ptr) const {
-    float tmp[16];
-    RVVI(__riscv_vse32_v_f32, LMUL_512)(tmp, reg_fp32, 16);
-    uint16_t* u16 = static_cast<uint16_t*>(ptr);
-    for (int i = 0; i < 16; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      u16[i] = static_cast<uint16_t>(v >> 16);
-    }
-  }
-  void save(void* ptr, int elem_num) const {
-    float tmp[16];
-    RVVI(__riscv_vse32_v_f32, LMUL_512)(tmp, reg_fp32, 16);
-    uint16_t* u16 = static_cast<uint16_t*>(ptr);
-    for (int i = 0; i < elem_num; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      u16[i] = static_cast<uint16_t>(v >> 16);
-    }
-  }
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    float tmp[16];
-    RVVI(__riscv_vse32_v_f32, LMUL_512)(tmp, reg_fp32, 16);
-    uint8_t* u8 = static_cast<uint8_t*>(ptr);
-    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
-    for (int i = 0; i < 16; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      uint16_t val = static_cast<uint16_t>(v >> 16);
-      *reinterpret_cast<uint16_t*>(u8 + i * byte_stride) = val;
-    }
-  }
-};
-
-struct BF16Vec32 : public Vec<BF16Vec32> {
-  constexpr static int VEC_ELEM_NUM = 32;
-  fixed_fp32x32_t reg_fp32;
-
-  explicit BF16Vec32(const void* ptr) {
-    const uint16_t* u16 = static_cast<const uint16_t*>(ptr);
-    float tmp[32];
-    for (int i = 0; i < 32; ++i) {
-      uint32_t v = static_cast<uint32_t>(u16[i]) << 16;
-      std::memcpy(&tmp[i], &v, 4);
-    }
-    reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_1024)(tmp, 32);
-  }
-
-  explicit BF16Vec32(const BF16Vec8& v) {
-    float tmp_small[8];
-    RVVI(__riscv_vse32_v_f32, LMUL_256)(tmp_small, v.reg_fp32, 8);
-    float tmp_large[32];
-    for (int i = 0; i < 4; ++i) {
-      std::memcpy(tmp_large + (i * 8), tmp_small, 8 * sizeof(float));
-    }
-    reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_1024)(tmp_large, 32);
-  }
-
-  void save(void* ptr) const {
-    float tmp[32];
-    RVVI(__riscv_vse32_v_f32, LMUL_1024)(tmp, reg_fp32, 32);
-    uint16_t* u16 = static_cast<uint16_t*>(ptr);
-    for (int i = 0; i < 32; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      u16[i] = static_cast<uint16_t>(v >> 16);
-    }
-  }
-
-  void save(void* ptr, int elem_num) const {
-    float tmp[32];
-    RVVI(__riscv_vse32_v_f32, LMUL_1024)(tmp, reg_fp32, 32);
-    uint16_t* u16 = static_cast<uint16_t*>(ptr);
-    for (int i = 0; i < elem_num; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      u16[i] = static_cast<uint16_t>(v >> 16);
-    }
-  }
-
-  void save_strided(void* ptr, ptrdiff_t stride) const {
-    float tmp[32];
-    RVVI(__riscv_vse32_v_f32, LMUL_1024)(tmp, reg_fp32, 32);
-    uint8_t* u8 = static_cast<uint8_t*>(ptr);
-    ptrdiff_t byte_stride = stride * sizeof(uint16_t);
-    for (int i = 0; i < 32; ++i) {
-      uint32_t v;
-      std::memcpy(&v, &tmp[i], 4);
-      uint16_t val = static_cast<uint16_t>(v >> 16);
-      *reinterpret_cast<uint16_t*>(u8 + i * byte_stride) = val;
-    }
-  }
-};
-#endif
-
-// ============================================================================
-// FP32 Implementation
-// ============================================================================
-
-struct FP32Vec4 : public Vec<FP32Vec4> {
-  constexpr static int VEC_ELEM_NUM = 4;
-  fixed_fp32x4_t reg;
-  explicit FP32Vec4(float v)
-      : reg(RVVI(__riscv_vfmv_v_f_f32, LMUL_128)(v, VEC_ELEM_NUM)) {};
-  explicit FP32Vec4()
-      : reg(RVVI(__riscv_vfmv_v_f_f32, LMUL_128)(0.0f, VEC_ELEM_NUM)) {};
-  explicit FP32Vec4(const float* ptr)
-      : reg(RVVI(__riscv_vle32_v_f32, LMUL_128)(ptr, VEC_ELEM_NUM)) {};
-  explicit FP32Vec4(fixed_fp32x4_t data) : reg(data) {};
-  explicit FP32Vec4(const FP32Vec4& data) : reg(data.reg) {};
-  void save(float* ptr) const {
-    RVVI(__riscv_vse32_v_f32, LMUL_128)(ptr, reg, VEC_ELEM_NUM);
-  }
-  void save(float* ptr, int elem_num) const {
-    RVVI(__riscv_vse32_v_f32, LMUL_128)(ptr, reg, elem_num);
-  }
-};
-
-struct FP32Vec8 : public Vec<FP32Vec8> {
-  constexpr static int VEC_ELEM_NUM = 8;
-  fixed_fp32x8_t reg;
-
-  explicit FP32Vec8(float v)
-      : reg(RVVI(__riscv_vfmv_v_f_f32, LMUL_256)(v, VEC_ELEM_NUM)) {};
-  explicit FP32Vec8()
-      : reg(RVVI(__riscv_vfmv_v_f_f32, LMUL_256)(0.0f, VEC_ELEM_NUM)) {};
-  explicit FP32Vec8(const float* ptr)
-      : reg(RVVI(__riscv_vle32_v_f32, LMUL_256)(ptr, VEC_ELEM_NUM)) {};
-  explicit FP32Vec8(fixed_fp32x8_t data) : reg(data) {};
-  explicit FP32Vec8(const FP32Vec8& data) : reg(data.reg) {};
-  explicit FP32Vec8(const FP16Vec8& v)
-      : reg(RVVI(__riscv_vfwcvt_f_f_v_f32, LMUL_256)(v.reg, VEC_ELEM_NUM)) {};
-  explicit FP32Vec8(fixed_fp16x8_t v)
-      : reg(RVVI(__riscv_vfwcvt_f_f_v_f32, LMUL_256)(v, VEC_ELEM_NUM)) {};
-
-#ifdef RISCV_BF16_SUPPORT
-  explicit FP32Vec8(fixed_bf16x8_t v)
-      : reg(RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_256)(v, VEC_ELEM_NUM)) {};
-  explicit FP32Vec8(const BF16Vec8& v)
-      : reg(RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_256)(v.reg, VEC_ELEM_NUM)) {
-        };
-#else
-  explicit FP32Vec8(const BF16Vec8& v) : reg(v.reg_fp32) {};
-#endif
-
-  float reduce_sum() const {
-    rvv_f32_accum_t scalar = __riscv_vfmv_s_f_f32m1(0.0f, 1);
-    scalar = RVVI3(__riscv_vfredusum_vs_f32, LMUL_256, _f32m1)(reg, scalar,
-                                                               VEC_ELEM_NUM);
-    return __riscv_vfmv_f_s_f32m1_f32(scalar);
-  }
-
-  FP32Vec8 operator*(const FP32Vec8& b) const {
-    return FP32Vec8(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_256)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec8 operator+(const FP32Vec8& b) const {
-    return FP32Vec8(
-        RVVI(__riscv_vfadd_vv_f32, LMUL_256)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec8 operator-(const FP32Vec8& b) const {
-    return FP32Vec8(
-        RVVI(__riscv_vfsub_vv_f32, LMUL_256)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec8 operator/(const FP32Vec8& b) const {
-    return FP32Vec8(
-        RVVI(__riscv_vfdiv_vv_f32, LMUL_256)(reg, b.reg, VEC_ELEM_NUM));
-  }
-
-  FP32Vec8 min(const FP32Vec8& b) const {
-    return FP32Vec8(
-        RVVI(__riscv_vfmin_vv_f32, LMUL_256)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec8 max(const FP32Vec8& b) const {
-    return FP32Vec8(
-        RVVI(__riscv_vfmax_vv_f32, LMUL_256)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec8 abs() const {
-    return FP32Vec8(RVVI(__riscv_vfabs_v_f32, LMUL_256)(reg, VEC_ELEM_NUM));
-  }
-
-  FP32Vec8 min(const FP32Vec8& b, int elem_num) const {
-    return FP32Vec8(RVVI(__riscv_vfmin_vv_f32, LMUL_256)(reg, b.reg, elem_num));
-  }
-  FP32Vec8 max(const FP32Vec8& b, int elem_num) const {
-    return FP32Vec8(RVVI(__riscv_vfmax_vv_f32, LMUL_256)(reg, b.reg, elem_num));
-  }
-
-  FP32Vec8 clamp(const FP32Vec8& min_v, const FP32Vec8& max_v) const {
-    fixed_fp32x8_t temp =
-        RVVI(__riscv_vfmax_vv_f32, LMUL_256)(min_v.reg, reg, VEC_ELEM_NUM);
-    return FP32Vec8(
-        RVVI(__riscv_vfmin_vv_f32, LMUL_256)(max_v.reg, temp, VEC_ELEM_NUM));
-  }
-
-  void save(float* ptr) const {
-    RVVI(__riscv_vse32_v_f32, LMUL_256)(ptr, reg, VEC_ELEM_NUM);
-  }
-  void save(float* ptr, int elem_num) const {
-    RVVI(__riscv_vse32_v_f32, LMUL_256)(ptr, reg, elem_num);
-  }
-  void save_strided(float* ptr, ptrdiff_t stride) const {
-    ptrdiff_t byte_stride = stride * sizeof(float);
-    RVVI(__riscv_vsse32_v_f32, LMUL_256)(ptr, byte_stride, reg, VEC_ELEM_NUM);
-  }
-
-  FP32Vec8 exp() const {
-    // Clamp input to prevent NaN: exp(-inf) must return 0, not NaN.
-    // Without clamping, -inf * 0.0 = NaN in the final poly * scale step.
-    // Matches the clamping strategy used by x86 AVX-512 and ARM NEON.
-    constexpr float exp_lo = -87.3365447505f;  // ln(FLT_MIN)
-    constexpr float exp_hi = 88.7228391117f;   // ln(FLT_MAX)
-    fixed_fp32x8_t x = RVVI(__riscv_vfmin_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmax_vf_f32, LMUL_256)(reg, exp_lo, VEC_ELEM_NUM), exp_hi,
-        VEC_ELEM_NUM);
-
-    const float inv_ln2 = 1.44269504088896341f;
-    fixed_fp32x8_t x_scaled =
-        RVVI(__riscv_vfmul_vf_f32, LMUL_256)(x, inv_ln2, VEC_ELEM_NUM);
-    fixed_i32x8_t n_int =
-        RVVI(__riscv_vfcvt_x_f_v_i32, LMUL_256)(x_scaled, VEC_ELEM_NUM);
-    fixed_fp32x8_t n_float =
-        RVVI(__riscv_vfcvt_f_x_v_f32, LMUL_256)(n_int, VEC_ELEM_NUM);
-
-    fixed_fp32x8_t r =
-        RVVI(__riscv_vfsub_vv_f32, LMUL_256)(x_scaled, n_float, VEC_ELEM_NUM);
-
-    fixed_fp32x8_t poly =
-        RVVI(__riscv_vfmv_v_f_f32, LMUL_256)(0.001333355810164f, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, r, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(poly, 0.009618129107628f,
-                                                VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, r, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(poly, 0.055504108664821f,
-                                                VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, r, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(poly, 0.240226506959101f,
-                                                VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, r, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(poly, 0.693147180559945f,
-                                                VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, r, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(poly, 1.0f, VEC_ELEM_NUM);
-
-    fixed_i32x8_t biased_exp =
-        RVVI(__riscv_vadd_vx_i32, LMUL_256)(n_int, 127, VEC_ELEM_NUM);
-    biased_exp =
-        RVVI(__riscv_vmax_vx_i32, LMUL_256)(biased_exp, 0, VEC_ELEM_NUM);
-    fixed_i32x8_t exponent_bits =
-        RVVI(__riscv_vsll_vx_i32, LMUL_256)(biased_exp, 23, VEC_ELEM_NUM);
-    fixed_fp32x8_t scale = RVVI4(__riscv_vreinterpret_v_i32, LMUL_256, _f32,
-                                 LMUL_256)(exponent_bits);
-
-    return FP32Vec8(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, scale, VEC_ELEM_NUM));
-  }
-
-  FP32Vec8 tanh() const {
-    fixed_fp32x8_t x_clamped = RVVI(__riscv_vfmin_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmax_vf_f32, LMUL_256)(reg, -9.0f, VEC_ELEM_NUM), 9.0f,
-        VEC_ELEM_NUM);
-    fixed_fp32x8_t x2 =
-        RVVI(__riscv_vfmul_vf_f32, LMUL_256)(x_clamped, 2.0f, VEC_ELEM_NUM);
-    FP32Vec8 exp_val = FP32Vec8(x2).exp();
-    fixed_fp32x8_t num =
-        RVVI(__riscv_vfsub_vf_f32, LMUL_256)(exp_val.reg, 1.0f, VEC_ELEM_NUM);
-    fixed_fp32x8_t den =
-        RVVI(__riscv_vfadd_vf_f32, LMUL_256)(exp_val.reg, 1.0f, VEC_ELEM_NUM);
-    return FP32Vec8(
-        RVVI(__riscv_vfdiv_vv_f32, LMUL_256)(num, den, VEC_ELEM_NUM));
-  }
-
-  FP32Vec8 er() const {
-    const float p = 0.3275911f, a1 = 0.254829592f, a2 = -0.284496736f,
-                a3 = 1.421413741f, a4 = -1.453152027f, a5 = 1.061405429f;
-    fixed_fp32x8_t abs_x =
-        RVVI(__riscv_vfabs_v_f32, LMUL_256)(reg, VEC_ELEM_NUM);
-
-    fixed_fp32x8_t t = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmul_vf_f32, LMUL_256)(abs_x, p, VEC_ELEM_NUM), 1.0f,
-        VEC_ELEM_NUM);
-    t = RVVI(__riscv_vfrdiv_vf_f32, LMUL_256)(t, 1.0f, VEC_ELEM_NUM);
-
-    fixed_fp32x8_t poly =
-        RVVI(__riscv_vfmv_v_f_f32, LMUL_256)(a5, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, t, VEC_ELEM_NUM), a4,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, t, VEC_ELEM_NUM), a3,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, t, VEC_ELEM_NUM), a2,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, t, VEC_ELEM_NUM), a1,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, t, VEC_ELEM_NUM);
-
-    fixed_fp32x8_t exp_val = FP32Vec8(RVVI(__riscv_vfneg_v_f32, LMUL_256)(
-                                          RVVI(__riscv_vfmul_vv_f32, LMUL_256)(
-                                              abs_x, abs_x, VEC_ELEM_NUM),
-                                          VEC_ELEM_NUM))
-                                 .exp()
-                                 .reg;
-    fixed_fp32x8_t res = RVVI(__riscv_vfrsub_vf_f32, LMUL_256)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_256)(poly, exp_val, VEC_ELEM_NUM), 1.0f,
-        VEC_ELEM_NUM);
-
-    rvv_mask_f32x8_t mask = RVVIB(__riscv_vmflt_vf_f32, LMUL_256, BOOL_256)(
-        reg, 0.0f, VEC_ELEM_NUM);
-    return FP32Vec8(
-        RVVI3(__riscv_vfneg_v_f32, LMUL_256, _m)(mask, res, VEC_ELEM_NUM));
-  }
-};
-
-struct FP32Vec16 : public Vec<FP32Vec16> {
-  constexpr static int VEC_ELEM_NUM = 16;
-  fixed_fp32x16_t reg;
-
-  explicit FP32Vec16(float v)
-      : reg(RVVI(__riscv_vfmv_v_f_f32, LMUL_512)(v, VEC_ELEM_NUM)) {};
-  explicit FP32Vec16()
-      : reg(RVVI(__riscv_vfmv_v_f_f32, LMUL_512)(0.0f, VEC_ELEM_NUM)) {};
-  explicit FP32Vec16(const float* ptr)
-      : reg(RVVI(__riscv_vle32_v_f32, LMUL_512)(ptr, VEC_ELEM_NUM)) {};
-  explicit FP32Vec16(fixed_fp32x16_t data) : reg(data) {};
-  explicit FP32Vec16(const FP32Vec8& data)
-      : 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(const FP16Vec16& v);
-
-#ifdef RISCV_BF16_SUPPORT
-  explicit FP32Vec16(fixed_bf16x16_t v)
-      : reg(RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_512)(v, VEC_ELEM_NUM)) {};
-  explicit FP32Vec16(const BF16Vec16& v)
-      : reg(RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_512)(v.reg, VEC_ELEM_NUM)) {
-        };
-#else
-  explicit FP32Vec16(const BF16Vec16& v) : reg(v.reg_fp32) {};
-#endif
-
-  FP32Vec16 operator+(const FP32Vec16& b) const {
-    return FP32Vec16(
-        RVVI(__riscv_vfadd_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec16 operator-(const FP32Vec16& b) const {
-    return FP32Vec16(
-        RVVI(__riscv_vfsub_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec16 operator*(const FP32Vec16& b) const {
-    return FP32Vec16(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec16 operator/(const FP32Vec16& b) const {
-    return FP32Vec16(
-        RVVI(__riscv_vfdiv_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
-  }
-
-  FP32Vec16 fma(const FP32Vec16& a, const FP32Vec16& b) const {
-    return FP32Vec16(
-        RVVI(__riscv_vfmacc_vv_f32, LMUL_512)(reg, a.reg, b.reg, VEC_ELEM_NUM));
-  }
-
-  float reduce_sum() const {
-    rvv_f32_accum_t scalar = __riscv_vfmv_s_f_f32m1(0.0f, 1);
-    scalar = RVVI3(__riscv_vfredusum_vs_f32, LMUL_512, _f32m1)(reg, scalar,
-                                                               VEC_ELEM_NUM);
-    return __riscv_vfmv_f_s_f32m1_f32(scalar);
-  }
-
-  float reduce_max() const {
-    rvv_f32_accum_t scalar =
-        __riscv_vfmv_s_f_f32m1(std::numeric_limits<float>::lowest(), 1);
-    scalar = RVVI3(__riscv_vfredmax_vs_f32, LMUL_512, _f32m1)(reg, scalar,
-                                                              VEC_ELEM_NUM);
-    return __riscv_vfmv_f_s_f32m1_f32(scalar);
-  }
-
-  float reduce_min() const {
-    rvv_f32_accum_t scalar =
-        __riscv_vfmv_s_f_f32m1(std::numeric_limits<float>::max(), 1);
-    scalar = RVVI3(__riscv_vfredmin_vs_f32, LMUL_512, _f32m1)(reg, scalar,
-                                                              VEC_ELEM_NUM);
-    return __riscv_vfmv_f_s_f32m1_f32(scalar);
-  }
-
-  template <int group_size>
-  float reduce_sub_sum(int idx) {
-    static_assert(VEC_ELEM_NUM % group_size == 0);
-    const int start = idx * group_size;
-    auto indices = RVVI(__riscv_vid_v_u32, LMUL_512)(VEC_ELEM_NUM);
-    rvv_mask_f32x16_t mask = RVVI(__riscv_vmand_mm_, BOOL_512)(
-        RVVIB(__riscv_vmsgeu_vx_u32, LMUL_512, BOOL_512)(indices, start,
-                                                         VEC_ELEM_NUM),
-        RVVIB(__riscv_vmsltu_vx_u32, LMUL_512, BOOL_512)(
-            indices, start + group_size, VEC_ELEM_NUM),
-        VEC_ELEM_NUM);
-    rvv_f32_accum_t scalar = __riscv_vfmv_s_f_f32m1(0.0f, 1);
-    scalar = RVVI3(__riscv_vfredusum_vs_f32, LMUL_512, _f32m1_m)(
-        mask, reg, scalar, VEC_ELEM_NUM);
-    return __riscv_vfmv_f_s_f32m1_f32(scalar);
-  };
-
-  FP32Vec16 max(const FP32Vec16& b) const {
-    return FP32Vec16(
-        RVVI(__riscv_vfmax_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec16 min(const FP32Vec16& b) const {
-    return FP32Vec16(
-        RVVI(__riscv_vfmin_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
-  }
-  FP32Vec16 abs() const {
-    return FP32Vec16(RVVI(__riscv_vfabs_v_f32, LMUL_512)(reg, VEC_ELEM_NUM));
-  }
-
-  FP32Vec16 clamp(const FP32Vec16& min_v, const FP32Vec16& max_v) const {
-    return FP32Vec16(RVVI(__riscv_vfmin_vv_f32, LMUL_512)(
-        max_v.reg,
-        RVVI(__riscv_vfmax_vv_f32, LMUL_512)(min_v.reg, reg, VEC_ELEM_NUM),
-        VEC_ELEM_NUM));
-  }
-
-  void save(float* ptr) const {
-    RVVI(__riscv_vse32_v_f32, LMUL_512)(ptr, reg, VEC_ELEM_NUM);
-  }
-  void save(float* ptr, int elem_num) const {
-    RVVI(__riscv_vse32_v_f32, LMUL_512)(ptr, reg, elem_num);
-  }
-  void save_strided(float* ptr, ptrdiff_t stride) const {
-    ptrdiff_t byte_stride = stride * sizeof(float);
-    RVVI(__riscv_vsse32_v_f32, LMUL_512)(ptr, byte_stride, reg, VEC_ELEM_NUM);
-  }
-
-  FP32Vec16 exp() const {
-    // Clamp input to prevent NaN: exp(-inf) must return 0, not NaN.
-    // Without clamping, -inf * 0.0 = NaN in the final poly * scale step.
-    // Matches the clamping strategy used by x86 AVX-512 and ARM NEON.
-    constexpr float exp_lo = -87.3365447505f;  // ln(FLT_MIN)
-    constexpr float exp_hi = 88.7228391117f;   // ln(FLT_MAX)
-    fixed_fp32x16_t x = RVVI(__riscv_vfmin_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmax_vf_f32, LMUL_512)(reg, exp_lo, VEC_ELEM_NUM), exp_hi,
-        VEC_ELEM_NUM);
-
-    const float inv_ln2 = 1.44269504088896341f;
-    fixed_fp32x16_t x_scaled =
-        RVVI(__riscv_vfmul_vf_f32, LMUL_512)(x, inv_ln2, VEC_ELEM_NUM);
-    fixed_i32x16_t n_int =
-        RVVI(__riscv_vfcvt_x_f_v_i32, LMUL_512)(x_scaled, VEC_ELEM_NUM);
-    fixed_fp32x16_t n_float =
-        RVVI(__riscv_vfcvt_f_x_v_f32, LMUL_512)(n_int, VEC_ELEM_NUM);
-    fixed_fp32x16_t r =
-        RVVI(__riscv_vfsub_vv_f32, LMUL_512)(x_scaled, n_float, VEC_ELEM_NUM);
-
-    fixed_fp32x16_t poly =
-        RVVI(__riscv_vfmv_v_f_f32, LMUL_512)(0.001333355810164f, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, r, VEC_ELEM_NUM),
-        0.009618129107628f, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, r, VEC_ELEM_NUM),
-        0.055504108664821f, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, r, VEC_ELEM_NUM),
-        0.240226506959101f, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, r, VEC_ELEM_NUM),
-        0.693147180559945f, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, r, VEC_ELEM_NUM), 1.0f,
-        VEC_ELEM_NUM);
-
-    fixed_i32x16_t biased_exp = RVVI(__riscv_vmax_vx_i32, LMUL_512)(
-        RVVI(__riscv_vadd_vx_i32, LMUL_512)(n_int, 127, VEC_ELEM_NUM), 0,
-        VEC_ELEM_NUM);
-    fixed_fp32x16_t scale =
-        RVVI4(__riscv_vreinterpret_v_i32, LMUL_512, _f32, LMUL_512)(
-            RVVI(__riscv_vsll_vx_i32, LMUL_512)(biased_exp, 23, VEC_ELEM_NUM));
-
-    return FP32Vec16(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, scale, VEC_ELEM_NUM));
-  }
-
-  FP32Vec16 tanh() const {
-    fixed_fp32x16_t x_clamped = RVVI(__riscv_vfmin_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmax_vf_f32, LMUL_512)(reg, -9.0f, VEC_ELEM_NUM), 9.0f,
-        VEC_ELEM_NUM);
-    FP32Vec16 exp_val = FP32Vec16(RVVI(__riscv_vfmul_vf_f32, LMUL_512)(
-                                      x_clamped, 2.0f, VEC_ELEM_NUM))
-                            .exp();
-    return FP32Vec16(RVVI(__riscv_vfdiv_vv_f32, LMUL_512)(
-        RVVI(__riscv_vfsub_vf_f32, LMUL_512)(exp_val.reg, 1.0f, VEC_ELEM_NUM),
-        RVVI(__riscv_vfadd_vf_f32, LMUL_512)(exp_val.reg, 1.0f, VEC_ELEM_NUM),
-        VEC_ELEM_NUM));
-  }
-
-  FP32Vec16 er() const {
-    const float p = 0.3275911f, a1 = 0.254829592f, a2 = -0.284496736f,
-                a3 = 1.421413741f, a4 = -1.453152027f, a5 = 1.061405429f;
-    fixed_fp32x16_t abs_x =
-        RVVI(__riscv_vfabs_v_f32, LMUL_512)(reg, VEC_ELEM_NUM);
-    fixed_fp32x16_t t = RVVI(__riscv_vfrdiv_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-            RVVI(__riscv_vfmul_vf_f32, LMUL_512)(abs_x, p, VEC_ELEM_NUM), 1.0f,
-            VEC_ELEM_NUM),
-        1.0f, VEC_ELEM_NUM);
-
-    fixed_fp32x16_t poly =
-        RVVI(__riscv_vfmv_v_f_f32, LMUL_512)(a5, VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, t, VEC_ELEM_NUM), a4,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, t, VEC_ELEM_NUM), a3,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, t, VEC_ELEM_NUM), a2,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfadd_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, t, VEC_ELEM_NUM), a1,
-        VEC_ELEM_NUM);
-    poly = RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, t, VEC_ELEM_NUM);
-
-    fixed_fp32x16_t exp_val =
-        FP32Vec16(RVVI(__riscv_vfneg_v_f32, LMUL_512)(
-                      RVVI(__riscv_vfmul_vv_f32, LMUL_512)(abs_x, abs_x,
-                                                           VEC_ELEM_NUM),
-                      VEC_ELEM_NUM))
-            .exp()
-            .reg;
-    fixed_fp32x16_t res = RVVI(__riscv_vfrsub_vf_f32, LMUL_512)(
-        RVVI(__riscv_vfmul_vv_f32, LMUL_512)(poly, exp_val, VEC_ELEM_NUM), 1.0f,
-        VEC_ELEM_NUM);
-
-    rvv_mask_f32x16_t mask = RVVIB(__riscv_vmflt_vf_f32, LMUL_512, BOOL_512)(
-        reg, 0.0f, VEC_ELEM_NUM);
-    return FP32Vec16(
-        RVVI3(__riscv_vfneg_v_f32, LMUL_512, _m)(mask, res, VEC_ELEM_NUM));
-  }
-};
-
-// ============================================================================
-// Type Traits & Global Helpers
-// ============================================================================
-
-template <typename T>
-struct VecType {
-  using vec_type = void;
-  using vec_t = void;
-};
-
-template <typename T>
-using vec_t = typename VecType<T>::vec_type;
-
-template <>
-struct VecType<float> {
-  using vec_type = FP32Vec8;
-  using vec_t = FP32Vec8;
-};
-template <>
-struct VecType<c10::Half> {
-  using vec_type = FP16Vec8;
-  using vec_t = FP16Vec8;
-};
-template <>
-struct VecType<c10::BFloat16> {
-  using vec_type = BF16Vec8;
-  using vec_t = BF16Vec8;
-};
-
-template <typename T>
-void storeFP32(float v, T* ptr) {
-  *ptr = v;
-}
-template <>
-inline void storeFP32<c10::Half>(float v, c10::Half* ptr) {
-  *reinterpret_cast<_Float16*>(ptr) = static_cast<_Float16>(v);
-}
-
-inline FP16Vec16::FP16Vec16(const FP32Vec16& v) {
-  reg = RVVI(__riscv_vfncvt_f_f_w_f16, LMUL_256)(v.reg, VEC_ELEM_NUM);
-}
-inline FP16Vec8::FP16Vec8(const FP32Vec8& v) {
-  reg = RVVI(__riscv_vfncvt_f_f_w_f16, LMUL_128)(v.reg, VEC_ELEM_NUM);
-}
-inline FP32Vec16::FP32Vec16(const FP16Vec16& v) {
-  reg = RVVI(__riscv_vfwcvt_f_f_v_f32, LMUL_512)(v.reg, VEC_ELEM_NUM);
-}
-inline void fma(FP32Vec16& acc, const FP32Vec16& a, const FP32Vec16& b) {
-  acc = acc.fma(a, b);
-}
-
-#ifdef RISCV_BF16_SUPPORT
-template <>
-inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
-  *ptr = static_cast<__bf16>(v);
-};
-inline BF16Vec8::BF16Vec8(const FP32Vec8& v)
-    : reg(RVVI(__riscv_vfncvtbf16_f_f_w_bf16, LMUL_128)(v.reg, VEC_ELEM_NUM)) {
-      };
-inline BF16Vec16::BF16Vec16(const FP32Vec16& v)
-    : reg(RVVI(__riscv_vfncvtbf16_f_f_w_bf16, LMUL_256)(v.reg, VEC_ELEM_NUM)) {
-      };
-#else
-template <>
-inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
-  uint32_t val;
-  std::memcpy(&val, &v, 4);
-  *reinterpret_cast<uint16_t*>(ptr) = static_cast<uint16_t>(val >> 16);
-}
-inline BF16Vec8::BF16Vec8(const FP32Vec8& v) : reg_fp32(v.reg) {}
-inline BF16Vec16::BF16Vec16(const FP32Vec16& v) : reg_fp32(v.reg) {}
-#endif
-
-inline void prefetch(const void* addr) { __builtin_prefetch(addr, 0, 1); }
-
-}  // namespace vec_op
-
-#ifndef CPU_KERNEL_GUARD_IN
-  #define CPU_KERNEL_GUARD_IN(NAME)
-#endif
-
-#ifndef CPU_KERNEL_GUARD_OUT
-  #define CPU_KERNEL_GUARD_OUT(NAME)
-#endif
-
-#endif  // CPU_TYPES_RISCV_IMPL_HPP
@@ -8,9 +8,6 @@
 #include <torch/all.h>
 namespace vec_op {

-struct fp8_e4m3_tag {};
-struct fp8_e5m2_tag {};
-
 #define vec_neg(a) (-(a))
 #define vec_add(a, b) ((a) + (b))
 #define vec_sub(a, b) ((a) - (b))
@@ -244,9 +241,6 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
  explicit BF16Vec32(const BF16Vec8& vec8_data)
      : reg({vec8_data.reg, vec8_data.reg, vec8_data.reg, vec8_data.reg}) {}

-  explicit BF16Vec32(const uint8_t*, fp8_e4m3_tag) : reg{} {}
-  explicit BF16Vec32(const uint8_t*, fp8_e5m2_tag) : reg{} {}
-
  void save(void* ptr) const { *reinterpret_cast<ss16x8x4_t*>(ptr) = reg; }
 };

@@ -11,17 +11,6 @@ static_assert(false, "AVX2 must be supported for the current implementation.");

 namespace vec_op {

-// Tags for FP8 BF16Vec32 constructors (avoid overload collision with
-// BF16Vec32(void*)).
-// VEC path (FP8 → pseudo-FP16 layout, scale correction applied later):
-struct fp8_e4m3_tag {};  // E4M3 → pseudo-FP16; BF16 value = true_E4M3 * 2^-8
-struct fp8_e5m2_tag {};  // E5M2 → FP16 bits directly (same exponent bias=15)
-// AMX path (FP8 → unscaled BF16, no FP32 round-trip):
-// BF16 value = true_E4M3 * 2^-120 (E4M3) or true_E5M2 * 2^-112 (E5M2).
-// Exponent rebiasing is folded into k/v scales by the caller.
-struct fp8_bf16_e4m3_tag {};
-struct fp8_bf16_e5m2_tag {};
-
 #define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)            \
  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__)    \
  AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__) \
@@ -187,50 +176,6 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
                               (__m128i)vec8_data.reg, 2),
            (__m128i)vec8_data.reg, 3)) {}

-  // Decode 32 FP8-E4M3 bytes to pseudo-FP16 layout (stored in the BF16
-  // register).  Result = true_E4M3 * 2^-8; caller applies scale * 2^8.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e4m3_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    __m512i b16 = _mm512_cvtepu8_epi16(b8);
-    __m512i sign =
-        _mm512_slli_epi16(_mm512_and_si512(b16, _mm512_set1_epi16(0x80)), 8);
-    __m512i payload =
-        _mm512_slli_epi16(_mm512_and_si512(b16, _mm512_set1_epi16(0x7F)), 7);
-    reg = _mm512_or_si512(sign, payload);
-  }
-
-  // Decode 32 FP8-E5M2 bytes to FP16 layout.
-  // E5M2 and FP16 share the same 5-bit exponent bias (15), so FP8 byte b maps
-  // directly to FP16 bits by shifting left 8 — no sign/payload reconstruction.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e5m2_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    reg = _mm512_slli_epi16(_mm512_cvtepu8_epi16(b8), 8);
-  }
-
-  // Direct FP8-E4M3 → unscaled BF16 for AMX (no FP32 round-trip).
-  // BF16 value = true_E4M3 * 2^-120; exponent rebiasing folded into k/v scales.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_bf16_e4m3_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    __m512i b16 = _mm512_cvtepu8_epi16(b8);
-    __m512i sign =
-        _mm512_slli_epi16(_mm512_and_si512(b16, _mm512_set1_epi16(0x80)), 8);
-    __m512i payload =
-        _mm512_slli_epi16(_mm512_and_si512(b16, _mm512_set1_epi16(0x7F)), 4);
-    reg = _mm512_or_si512(sign, payload);
-  }
-
-  // Direct FP8-E5M2 → unscaled BF16 for AMX (no FP32 round-trip).
-  // BF16 value = true_E5M2 * 2^-112; exponent rebiasing folded into k/v scales.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_bf16_e5m2_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    __m512i b16 = _mm512_cvtepu8_epi16(b8);
-    __m512i sign =
-        _mm512_slli_epi16(_mm512_and_si512(b16, _mm512_set1_epi16(0x80)), 8);
-    __m512i payload =
-        _mm512_slli_epi16(_mm512_and_si512(b16, _mm512_set1_epi16(0x7F)), 5);
-    reg = _mm512_or_si512(sign, payload);
-  }
-
  void save(void* ptr) const { *reinterpret_cast<__m512i*>(ptr) = reg; }
 };
 #else
@@ -255,77 +200,6 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
            _mm256_castsi128_si256((__m128i)vec8_data.reg),
            (__m128i)vec8_data.reg, 1)) {}

-  // E4M3 decode (AVX2 path) — same bit-layout trick as the AVX512 variant
-  // above.  Result = true_E4M3 * 2^-8; caller applies scale * 2^8.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e4m3_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    __m128i b8_low = _mm256_extracti128_si256(b8, 0);
-    __m128i b8_high = _mm256_extracti128_si256(b8, 1);
-    __m256i b16_low = _mm256_cvtepu8_epi16(b8_low);
-    __m256i b16_high = _mm256_cvtepu8_epi16(b8_high);
-
-    __m256i sign_low = _mm256_slli_epi16(
-        _mm256_and_si256(b16_low, _mm256_set1_epi16(0x80)), 8);
-    __m256i payload_low = _mm256_slli_epi16(
-        _mm256_and_si256(b16_low, _mm256_set1_epi16(0x7F)), 7);
-    __m256i sign_high = _mm256_slli_epi16(
-        _mm256_and_si256(b16_high, _mm256_set1_epi16(0x80)), 8);
-    __m256i payload_high = _mm256_slli_epi16(
-        _mm256_and_si256(b16_high, _mm256_set1_epi16(0x7F)), 7);
-    reg_low = _mm256_or_si256(sign_low, payload_low);
-    reg_high = _mm256_or_si256(sign_high, payload_high);
-  }
-
-  // E5M2 decode (AVX2 path) — b << 8 maps to FP16 bits; see AVX512 variant
-  // above.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e5m2_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    __m128i b8_low = _mm256_extracti128_si256(b8, 0);
-    __m128i b8_high = _mm256_extracti128_si256(b8, 1);
-    reg_low = _mm256_slli_epi16(_mm256_cvtepu8_epi16(b8_low), 8);
-    reg_high = _mm256_slli_epi16(_mm256_cvtepu8_epi16(b8_high), 8);
-  }
-
-  // Direct FP8-E4M3 → unscaled BF16 for AMX (AVX2 path, no FP32 round-trip).
-  // BF16 value = true_E4M3 * 2^-120; exponent rebiasing folded into k/v scales.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_bf16_e4m3_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    __m128i b8_low = _mm256_extracti128_si256(b8, 0);
-    __m128i b8_high = _mm256_extracti128_si256(b8, 1);
-    __m256i b16_low = _mm256_cvtepu8_epi16(b8_low);
-    __m256i b16_high = _mm256_cvtepu8_epi16(b8_high);
-    reg_low = _mm256_or_si256(
-        _mm256_slli_epi16(_mm256_and_si256(b16_low, _mm256_set1_epi16(0x80)),
-                          8),
-        _mm256_slli_epi16(_mm256_and_si256(b16_low, _mm256_set1_epi16(0x7F)),
-                          4));
-    reg_high = _mm256_or_si256(
-        _mm256_slli_epi16(_mm256_and_si256(b16_high, _mm256_set1_epi16(0x80)),
-                          8),
-        _mm256_slli_epi16(_mm256_and_si256(b16_high, _mm256_set1_epi16(0x7F)),
-                          4));
-  }
-
-  // Direct FP8-E5M2 → unscaled BF16 for AMX (AVX2 path, no FP32 round-trip).
-  // BF16 value = true_E5M2 * 2^-112; exponent rebiasing folded into k/v scales.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_bf16_e5m2_tag) {
-    __m256i b8 = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(ptr));
-    __m128i b8_low = _mm256_extracti128_si256(b8, 0);
-    __m128i b8_high = _mm256_extracti128_si256(b8, 1);
-    __m256i b16_low = _mm256_cvtepu8_epi16(b8_low);
-    __m256i b16_high = _mm256_cvtepu8_epi16(b8_high);
-    reg_low = _mm256_or_si256(
-        _mm256_slli_epi16(_mm256_and_si256(b16_low, _mm256_set1_epi16(0x80)),
-                          8),
-        _mm256_slli_epi16(_mm256_and_si256(b16_low, _mm256_set1_epi16(0x7F)),
-                          5));
-    reg_high = _mm256_or_si256(
-        _mm256_slli_epi16(_mm256_and_si256(b16_high, _mm256_set1_epi16(0x80)),
-                          8),
-        _mm256_slli_epi16(_mm256_and_si256(b16_high, _mm256_set1_epi16(0x7F)),
-                          5));
-  }
-
  void save(void* ptr) const {
    _mm256_storeu_si256((__m256i*)ptr, reg_low);
    _mm256_storeu_si256((__m256i*)ptr + 1, reg_high);
@@ -516,11 +390,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
      : reg(_mm512_castsi512_ps(
            _mm512_bslli_epi128(_mm512_cvtepu16_epi32(v.reg), 2))) {}

-  explicit FP32Vec16(const BF16Vec32& v, int upper) {
-    __m256i v_half_i = _mm512_extracti32x8_epi32(v.reg, upper);
-    reg = _mm512_cvtph_ps(v_half_i);
-  }
-
  explicit FP32Vec16(const FP16Vec16& v) : reg(_mm512_cvtph_ps(v.reg)) {}

  explicit FP32Vec16(const FP16Vec8& v) : FP32Vec16(FP32Vec8(v)) {}
@@ -625,14 +494,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
  explicit FP32Vec16(const FP32Vec8& data)
      : reg_low(data.reg), reg_high(data.reg) {}

-  explicit FP32Vec16(const BF16Vec32& v, int upper) {
-    const __m256i& half = upper ? v.reg_high : v.reg_low;
-    __m128i lo = _mm256_extractf128_si256(half, 0);
-    __m128i hi = _mm256_extractf128_si256(half, 1);
-    reg_low = _mm256_cvtph_ps(lo);
-    reg_high = _mm256_cvtph_ps(hi);
-  }
-
  explicit FP32Vec16(const FP16Vec16& v) {
    __m128i low = _mm256_extractf128_si256(v.reg, 0);
    __m128i high = _mm256_extractf128_si256(v.reg, 1);
@@ -22,95 +22,71 @@ ISA_TYPES = {
    "VXE": 4,
 }

-# KV cache index: 0 = auto (same as scalar_t), 1 = fp8_e4m3, 2 = fp8_e5m2
-KV_CACHE_IDX = {
-    "auto": 0,
-    "fp8_e4m3": 1,
-    "fp8_e5m2": 2,
-}
-
-# C++ type for each kv_cache index
-KV_CACHE_CPP_TYPES = {
-    "auto": "scalar_t",
-    "fp8_e4m3": "c10::Float8_e4m3fn",
-    "fp8_e5m2": "c10::Float8_e5m2",
-}
-
 # ISAs supported for head_dims divisible by 32
 ISA_FOR_32 = ["AMX", "NEON", "VEC", "VEC16", "VXE"]

 # ISAs supported for head_dims divisible by 16 only
 ISA_FOR_16 = ["VEC16"]

-# ISAs that support FP8 KV cache (x86 AVX2/AVX-512 required)
-ISA_FOR_FP8 = ["AMX", "VEC"]

-
-def encode_params(head_dim: int, isa_type: str, kv_cache: str = "auto") -> int:
-    """Encode head_dim, ISA type, and KV cache type into a single int64_t."""
+def encode_params(head_dim: int, isa_type: str) -> int:
+    """Encode head_dim and ISA type into a single int64_t."""
    isa_val = ISA_TYPES[isa_type]
-    kv_val = KV_CACHE_IDX[kv_cache]
-    # Encoding: (head_dim << 16) | (kv_cache_idx << 8) | isa_type
-    # This allows head_dim up to 2^48 - 1, 256 KV cache types, and 256 ISA types
-    return (head_dim << 16) | (kv_val << 8) | isa_val
+    # Encoding: (head_dim << 8) | isa_type
+    # This allows head_dim up to 2^56 - 1 and 256 ISA types
+    return (head_dim << 8) | isa_val


-def _make_case(
-    head_dim: int, isa: str, kv_cache: str = "auto", isa_override: str | None = None
-) -> str:
-    """Generate a single switch case line."""
-    encoded = encode_params(head_dim, isa, kv_cache)
-    actual_isa = isa_override if isa_override else isa
-    cpp_type = KV_CACHE_CPP_TYPES[kv_cache]
-    attn_impl = (
-        f"cpu_attention::AttentionImpl<"
-        f"cpu_attention::ISA::{actual_isa}, \\\n"
-        f"                                                       "
-        f"scalar_t, head_dim, {cpp_type}>"
-    )
-    comment = (
-        f"head_dim={head_dim}, isa={isa}"
-        if kv_cache == "auto"
-        else f"head_dim={head_dim}, isa={isa}, kv_cache={kv_cache}"
-    )
-    return (
-        f"""      case {encoded}LL: {{ """
-        f"""/* {comment} */ \\"""
-        f"""
-        constexpr size_t head_dim = {head_dim}; \\"""
-        f"""
-        using attn_impl = {attn_impl}; \\"""
-        f"""
-        return __VA_ARGS__(); \\"""
-        f"""
-      }} \\"""
-    )
-
-
-def generate_cases_for_isa_group(isa_list: list[str], include_fp8: bool = False) -> str:
+def generate_cases_for_isa_group(isa_list: list[str]) -> str:
    """Generate switch cases for a specific ISA group."""
    cases = []

-    # Non-FP8 cases for head_dims divisible by 32
+    # Generate cases for head_dims divisible by 32
    for head_dim in HEAD_DIMS_32:
        for isa in isa_list:
            if isa not in ISA_FOR_32:
                continue
-            cases.append(_make_case(head_dim, isa, "auto"))
+            encoded = encode_params(head_dim, isa)
+            case_str = (
+                f"""      case {encoded}LL: {{ """
+                f"""/* head_dim={head_dim}, isa={isa} */ \\"""
+                f"""
+        constexpr size_t head_dim = {head_dim}; \\"""
+                f"""
+        using attn_impl = cpu_attention::AttentionImpl<"""
+                f"""cpu_attention::ISA::{isa}, \\"""
+                f"""
+                                                       """
+                f"""scalar_t, head_dim>; \\"""
+                f"""
+        return __VA_ARGS__(); \\"""
+                f"""
+      }} \\"""
+            )
+            cases.append(case_str)

-    # Non-FP8 cases for head_dims divisible by 16 only
+    # Generate cases for head_dims divisible by 16 only
    for head_dim in HEAD_DIMS_16:
        for isa in isa_list:
-            cases.append(_make_case(head_dim, isa, "auto", isa_override="VEC16"))
-
-    # FP8 cases: only AMX and VEC, only head_dims divisible by 32
-    if include_fp8:
-        for fp8_type in ("fp8_e4m3", "fp8_e5m2"):
-            for head_dim in HEAD_DIMS_32:
-                for isa in isa_list:
-                    if isa not in ISA_FOR_FP8:
-                        continue
-                    cases.append(_make_case(head_dim, isa, fp8_type))
+            encoded = encode_params(head_dim, isa)
+            case_str = (
+                f"""      case {encoded}LL: {{ """
+                f"""/* head_dim={head_dim}, isa={isa} """
+                f"""(using VEC16) */ \\"""
+                f"""
+        constexpr size_t head_dim = {head_dim}; \\"""
+                f"""
+        using attn_impl = cpu_attention::AttentionImpl<"""
+                f"""cpu_attention::ISA::VEC16, \\"""
+                f"""
+                                                       """
+                f"""scalar_t, head_dim>; \\"""
+                f"""
+        return __VA_ARGS__(); \\"""
+                f"""
+      }} \\"""
+            )
+            cases.append(case_str)

    return "\n".join(cases)

@@ -118,9 +94,8 @@ def generate_cases_for_isa_group(isa_list: list[str], include_fp8: bool = False)
 def generate_helper_function() -> str:
    """Generate helper function to encode parameters."""
    return """
-inline int64_t encode_cpu_attn_params(int64_t head_dim, cpu_attention::ISA isa,
-                                      int64_t kv_cache_idx = 0) {
-  return (head_dim << 16) | (kv_cache_idx << 8) | static_cast<int64_t>(isa);
+inline int64_t encode_cpu_attn_params(int64_t head_dim, cpu_attention::ISA isa) {
+  return (head_dim << 8) | static_cast<int64_t>(isa);
 }
 """

@@ -154,78 +129,87 @@ def generate_header_file() -> str:

    # Generate dispatch macro with conditional compilation for different ISA sets
    header += """
-// Dispatch macro using encoded parameters.
-// KV_CACHE_IDX: Fp8KVCacheDataType enum value (kAuto=0, kFp8E4M3=1, kFp8E5M2=2).
-// FP8 cases (kv_cache_idx != 0) are generated on x86 platforms with AVX2 or
-// AVX-512: BF16Vec32 FP8 constructors have both AVX-512 and AVX2 implementations
-// in cpu_types_x86.hpp. Non-x86 platforms (#else fallback) have fp8=False.
+// Dispatch macro using encoded parameters
 """

-    def _macro_block(guard: str, isa_list: list[str], fp8: bool) -> str:
-        """Return one CPU_ATTN_DISPATCH macro block for a given guard."""
-        enc = (
-            "    int64_t encoded_params = encode_cpu_attn_params("
-            "HEAD_DIM, ISA_TYPE, KV_CACHE_IDX); \\"
-        )
-        cases = generate_cases_for_isa_group(isa_list, include_fp8=fp8)
-        tail = (
-            "\n"
-            "      default: { \\\n"
-            "        TORCH_CHECK(false, "
-            '"Unsupported CPU attention configuration: head_dim=" + \\\n'
-            '                    std::to_string(HEAD_DIM) + " isa=" + \\\n'
-            "                    std::to_string(static_cast<int>(ISA_TYPE))"
-            " + \\\n"
-            '                    " kv_cache_idx=" + '
-            "std::to_string(KV_CACHE_IDX)); \\\n"
-            "      } \\\n"
-            "    } \\\n"
-            "  }()\n\n"
-        )
-        return (
-            f"{guard}\n"
-            "#define CPU_ATTN_DISPATCH(HEAD_DIM, ISA_TYPE, KV_CACHE_IDX, ...) \\\n"
-            "  [&] { \\\n"
-            f"{enc}\n"
-            "    switch (encoded_params) { \\\n"
-            f"{cases}"
-            f"{tail}"
-        )
+    # x86_64 with AMX
+    header += """#if defined(CPU_CAPABILITY_AMXBF16)
+#define CPU_ATTN_DISPATCH(HEAD_DIM, ISA_TYPE, ...) \\
+  [&] { \\
+    int64_t encoded_params = encode_cpu_attn_params(HEAD_DIM, ISA_TYPE); \\
+    switch (encoded_params) { \\
+"""
+    header += generate_cases_for_isa_group(["AMX", "VEC", "VEC16"])
+    header += """
+      default: { \\
+        TORCH_CHECK(false, "Unsupported CPU attention configuration: head_dim=" + \\
+                    std::to_string(HEAD_DIM) + " isa=" + \\
+                    std::to_string(static_cast<int>(ISA_TYPE))); \\
+      } \\
+    } \\
+  }()

-    header += _macro_block(
-        "#if defined(CPU_CAPABILITY_AMXBF16)",
-        ["AMX", "VEC", "VEC16"],
-        fp8=True,
-    )
-    header += _macro_block(
-        "#elif defined(__aarch64__)",
-        ["NEON", "VEC", "VEC16"],
-        fp8=False,
-    )
-    header += _macro_block(
-        "#elif defined(__s390x__)",
-        ["VXE", "VEC", "VEC16"],
-        fp8=False,
-    )
-    header += _macro_block(
-        "#elif defined(__AVX512F__)",
-        ["VEC", "VEC16"],
-        fp8=True,
-    )
-    header += _macro_block(
-        "#elif defined(__AVX2__)",
-        ["VEC", "VEC16"],
-        fp8=False,
-    )
-    header += _macro_block(
-        "#else",
-        ["VEC", "VEC16"],
-        fp8=False,
-    )
-    header += (
-        "#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / __s390x__ */\n\n"
-        "#endif  // CPU_ATTN_DISPATCH_GENERATED_H\n"
-    )
+"""
+
+    # ARM64 with NEON
+    header += """#elif defined(__aarch64__)
+#define CPU_ATTN_DISPATCH(HEAD_DIM, ISA_TYPE, ...) \\
+  [&] { \\
+    int64_t encoded_params = encode_cpu_attn_params(HEAD_DIM, ISA_TYPE); \\
+    switch (encoded_params) { \\
+"""
+    header += generate_cases_for_isa_group(["NEON", "VEC", "VEC16"])
+    header += """
+      default: { \\
+        TORCH_CHECK(false, "Unsupported CPU attention configuration: head_dim=" + \\
+                    std::to_string(HEAD_DIM) + " isa=" + \\
+                    std::to_string(static_cast<int>(ISA_TYPE))); \\
+      } \\
+    } \\
+  }()
+
+"""
+
+    # s390x with VXE
+    header += """#elif defined(__s390x__)
+#define CPU_ATTN_DISPATCH(HEAD_DIM, ISA_TYPE, ...) \\
+  [&] { \\
+    int64_t encoded_params = encode_cpu_attn_params(HEAD_DIM, ISA_TYPE); \\
+    switch (encoded_params) { \\
+"""
+    header += generate_cases_for_isa_group(["VXE", "VEC", "VEC16"])
+    header += """
+      default: { \\
+        TORCH_CHECK(false, "Unsupported CPU attention configuration: head_dim=" + \\
+                    std::to_string(HEAD_DIM) + " isa=" + \\
+                    std::to_string(static_cast<int>(ISA_TYPE))); \\
+      } \\
+    } \\
+  }()
+
+"""
+
+    # Fallback: VEC and VEC16 only
+    header += """#else
+#define CPU_ATTN_DISPATCH(HEAD_DIM, ISA_TYPE, ...) \\
+  [&] { \\
+    int64_t encoded_params = encode_cpu_attn_params(HEAD_DIM, ISA_TYPE); \\
+    switch (encoded_params) { \\
+"""
+    header += generate_cases_for_isa_group(["VEC", "VEC16"])
+    header += """
+      default: { \\
+        TORCH_CHECK(false, "Unsupported CPU attention configuration: head_dim=" + \\
+                    std::to_string(HEAD_DIM) + " isa=" + \\
+                    std::to_string(static_cast<int>(ISA_TYPE))); \\
+      } \\
+    } \\
+  }()
+
+#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / __s390x__ */
+
+#endif  // CPU_ATTN_DISPATCH_GENERATED_H
+"""

    return header

@@ -178,12 +178,7 @@ void rotary_embedding_gptj_impl(

 void rotary_embedding(torch::Tensor& positions, torch::Tensor& query,
                      std::optional<torch::Tensor> key, int64_t head_size,
-                      torch::Tensor& cos_sin_cache, bool is_neox,
-                      int64_t rope_dim_offset, bool inverse) {
-  TORCH_CHECK(rope_dim_offset == 0,
-              "rope_dim_offset != 0 is not supported on CPU");
-  TORCH_CHECK(!inverse, "inverse rotary embedding is not supported on CPU");
-
+                      torch::Tensor& cos_sin_cache, bool is_neox) {
  int num_tokens = positions.numel();
  int rot_dim = cos_sin_cache.size(1);
  int num_heads = query.size(-1) / head_size;
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
khluuandClaude Opus 4.6	5d236e5cf8	[Bugfix] Add MigCudaPlatform to avoid CUDA init before fork NonNvmlCudaPlatform uses torch.cuda for get_device_capability(), which initializes the CUDA context. When this runs at import time (e.g. via requires_fp8 in test utils), it prevents fork-based tests from working ("Cannot re-initialize CUDA in forked subprocess"). MigCudaPlatform inherits from NvmlCudaPlatform (NVML-based capability queries don't init CUDA) and only overrides the memory methods that fail on MIG partitions. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Signed-off-by: khluu <khluu000@gmail.com>	2026-04-16 11:49:08 -07:00
khluu	f683266536	lint Signed-off-by: khluu <khluu000@gmail.com>	2026-04-16 02:35:54 -07:00
khluuandClaude Opus 4.6	8dde2576c0	[Bugfix] Disable expandable_segments on MIG to fix PyTorch NVML assert The previous commit added MIG detection to use NonNvmlCudaPlatform, which fixes vllm's own NVML calls. However, PyTorch's CUDACachingAllocator has a separate NVML code path via the expandable_segments feature that still crashes with "NVML_SUCCESS == r INTERNAL ASSERT FAILED" on MIG devices. This change automatically sets PYTORCH_CUDA_ALLOC_CONF=expandable_segments:False when a MIG device is detected, so PyTorch's allocator falls back to non-NVML memory management and real torch.cuda.OutOfMemoryError surfaces instead of the misleading NVML assert. Signed-off-by: khluu <khluu000@gmail.com> Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-04-15 16:09:29 -07:00
khluuandClaude Opus 4.6	ae60bf5ea7	[Bugfix] Use NonNvmlCudaPlatform on MIG devices to fix NVML masking OOM errors On MIG partitions, NVML device-level queries fail with NVMLError_NoPermission. This causes PyTorch's CUDACachingAllocator (via expandable_segments) to crash with an internal assert before the real torch.cuda.OutOfMemoryError can surface, and causes NVML memory queries to report the full GPU's memory (e.g. 141 GB for H200) instead of the MIG partition's memory (e.g. 16 GB). This change: 1. Adds MIG detection in vllm/platforms/cuda.py that checks CUDA_VISIBLE_DEVICES for MIG UUIDs and probes NVML for permission errors, then falls back to NonNvmlCudaPlatform which uses torch.cuda APIs that correctly report MIG partition memory. 2. Patches test utilities (wait_for_gpu_memory_to_clear, _get_gpu_memory_used) to fall back to torch.cuda.mem_get_info() when NVML fails with permission errors on MIG. Signed-off-by: khluu <khluu000@gmail.com> Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>	2026-04-14 23:44:59 -07:00