benchmark changes

Signed-off-by: Luka Govedič <lgovedic@redhat.com>

.buildkite/hardware_tests/cpu.yaml

@@ -12,7 +12,7 @@ steps:
   - vllm/_custom_ops.py
   - tests/kernels/attention/test_cpu_attn.py
   - tests/kernels/moe/test_cpu_fused_moe.py
-  - tests/kernels/moe/test_cpu_quant_fused_moe.py
+  - tests/kernels/moe/test_cpu_fp8_fused_moe.py
   - tests/kernels/test_onednn.py
   - tests/kernels/test_awq_int4_to_int8.py
   - tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
@@ -21,7 +21,7 @@ steps:
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
       pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
       pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
-      pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
+      pytest -x -v -s tests/kernels/moe/test_cpu_fp8_fused_moe.py
       pytest -x -v -s tests/kernels/test_onednn.py
       pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
       pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py"
@@ -54,20 +54,6 @@ steps:
       pytest -x -v -s tests/models/language/generation -m cpu_model
       pytest -x -v -s tests/models/language/pooling -m cpu_model"
 
-- label: CPU-ModelRunnerV2 Tests
-  depends_on: []
-  device: intel_cpu
-  no_plugin: true
-  soft_fail: true
-  source_file_dependencies:
-  - vllm/v1/worker/cpu/
-  - vllm/v1/worker/gpu/
-  commands:
-    - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
-      uv pip install git+https://github.com/triton-lang/triton-cpu.git@270e696d
-      VLLM_USE_V2_MODEL_RUNNER=1 pytest -x -v -s tests/models/language/generation/test_granite.py -m cpu_model"
-
 - label: CPU-Quantization Model Tests
   depends_on: []
   device: intel_cpu
@@ -75,7 +61,7 @@ steps:
   source_file_dependencies:
   - csrc/cpu/
   - vllm/model_executor/layers/quantization/cpu_wna16.py
-  - vllm/model_executor/layers/quantization/auto_gptq.py
+  - vllm/model_executor/layers/quantization/gptq_marlin.py
   - vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w8a8_int8.py
   - vllm/model_executor/layers/quantization/kernels/scaled_mm/cpu.py
   - vllm/model_executor/layers/quantization/kernels/mixed_precision/cpu.py
@@ -84,7 +70,7 @@ steps:
   - tests/quantization/test_cpu_wna16.py
   commands:
     - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 20m "
       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"
       

.buildkite/hardware_tests/intel.yaml

@@ -8,3 +8,10 @@ steps:
     commands: 
     - bash .buildkite/scripts/hardware_ci/run-hpu-test.sh
 
+  - label: "Intel GPU Test"
+    depends_on: []
+    soft_fail: true
+    device: intel_gpu
+    no_plugin: true
+    commands: 
+    - bash .buildkite/scripts/hardware_ci/run-xpu-test.sh

.buildkite/image_build/image_build.yaml

@@ -6,48 +6,6 @@ steps:
     timeout_in_minutes: 600
     commands:
     - if [[ "$BUILDKITE_BRANCH" == "main" ]]; then .buildkite/image_build/image_build.sh $REGISTRY $REPO $BUILDKITE_COMMIT $BRANCH $IMAGE_TAG $IMAGE_TAG_LATEST; else .buildkite/image_build/image_build.sh $REGISTRY $REPO $BUILDKITE_COMMIT $BRANCH $IMAGE_TAG; fi
-    # Non-root smoke 1: the default (root) image must still be importable
-    # under a non-root UID via `--user 2000:0`. Validates the `vllm` passwd
-    # entry + group-0-writable /home/vllm + uv path cleanup from #31959.
-    # Uses `import vllm` rather than `vllm serve --help` because the latter
-    # instantiates `VllmConfig` which requires a GPU attached to the
-    # container.
-    - docker run --rm --user 2000:0 --entrypoint python3 "$IMAGE_TAG" -c "import vllm; print(vllm.__version__)"
-    # Non-root smoke 2: assert the non-root enabling invariants are baked
-    # into the image. Runs as UID 2000:0 via a shell so we can verify
-    # filesystem perms + passwd/group file state + wrapper presence without
-    # triggering vLLM's GPU-requiring config-init path. The opt-in
-    # `vllm-openai-nonroot` target adds only `USER vllm`, `WORKDIR
-    # /home/vllm`, and an `ENTRYPOINT` override on top of these invariants;
-    # its build correctness is reviewed at the Dockerfile level. Wrapper
-    # logic is covered separately by the pre-commit hook
-    # `test-nonroot-entrypoint` (see .pre-commit-config.yaml).
-    - |
-      docker run --rm --user 2000:0 --entrypoint /bin/sh "$IMAGE_TAG" -ec '
-        if ! getent passwd 2000 | grep -q ^vllm:; then
-          echo FAIL: UID 2000 != vllm
-          exit 1
-        fi
-        if ! id -gn 2>/dev/null | grep -qx root; then
-          echo FAIL: GID 0 not root group
-          exit 1
-        fi
-        touch /home/vllm/.smoke && rm /home/vllm/.smoke
-        touch /opt/uv/cache/.smoke && rm /opt/uv/cache/.smoke
-        if ! test -x /usr/local/bin/vllm-nonroot-entrypoint.sh; then
-          echo FAIL: wrapper missing
-          exit 1
-        fi
-        if ! test -w /etc/passwd; then
-          echo FAIL: /etc/passwd not group-writable
-          exit 1
-        fi
-        if ! test -w /etc/group; then
-          echo FAIL: /etc/group not group-writable
-          exit 1
-        fi
-        echo non-root invariants OK
-      '
     retry:
       automatic:
         - exit_status: -1  # Agent was lost

.buildkite/image_build/image_build_hpu.sh

@@ -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"-hpu) ]]; then

.buildkite/image_build/image_build_xpu.sh

@@ -11,8 +11,8 @@ 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 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
 
 # skip build if image already exists
 if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then

.buildkite/intel_jobs/lora_intel.yaml

@@ -49,7 +49,7 @@ steps:
       'cd tests &&
       export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_fused_moe_lora_kernel.py && 
-      pytest -v -s lora/test_moe_lora_align_sum.py --deselect="tests/lora/test_moe_lora_align_sum.py::test_moe_lora_align_block_size_mixed_base_and_lora[1]"'
+      pytest -v -s lora/test_moe_lora_align_sum.py'
 
 - label: LoRA Punica Kernels
   timeout_in_minutes: 45

.buildkite/intel_jobs/test-intel.yaml

@@ -39,9 +39,7 @@ steps:
         python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096 &&
         python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
         python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 &&
-        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel &&
-        python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192
-        '
+        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel'
   - label: "XPU V1 test"
     depends_on:
       - image-build-xpu
@@ -66,22 +64,3 @@ steps:
         pytest -v -s v1/test_serial_utils.py &&
         pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py &&
         pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py'
-  - label: "XPU server test"
-    depends_on:
-      - image-build-xpu
-    timeout_in_minutes: 30
-    device: intel_gpu
-    no_plugin: true
-    env:
-      REGISTRY: "public.ecr.aws/q9t5s3a7"
-      REPO: "vllm-ci-test-repo"
-    source_file_dependencies:
-      - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml 
-    commands:
-      - >-
-        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-        'pip install av &&
-        cd tests &&
-        pytest -v -s entrypoints/openai/chat_completion/test_audio_in_video.py &&
-        pytest -v -s benchmarks/test_serve_cli.py'

.buildkite/lm-eval-harness/run-lm-eval-chartqa-vllm-vlm-baseline.sh

@@ -2,7 +2,7 @@
 # We can use this script to compute baseline accuracy on chartqa for vllm.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

.buildkite/lm-eval-harness/run-lm-eval-gsm-hf-baseline.sh

@@ -2,7 +2,7 @@
 # We can use this script to compute baseline accuracy on GSM for transformers.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

.buildkite/lm-eval-harness/run-lm-eval-gsm-vllm-baseline.sh

@@ -3,7 +3,7 @@
 # We use this for fp8, which HF does not support.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

.buildkite/lm-eval-harness/run-lm-eval-mmlupro-vllm-baseline.sh

@@ -3,7 +3,7 @@
 # We use this for fp8, which HF does not support.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

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

@@ -31,9 +31,30 @@
     }
   },
   "tests": [
+    {
+      "test_name": "serving_llama8B_tp1_sharegpt",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "sharegpt",
+        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_sharegpt",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "sharegpt",
+        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
+      }
+    },
     {
       "test_name": "serving_llama8B_tp1_random_128_128",
       "server_parameters": {
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "dataset_name": "random",
@@ -41,10 +62,158 @@
         "random-output-len": 128
       }
     },
+    {
+      "test_name": "serving_llama8B_tp2_random_128_128",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_128_128",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp1_random_128_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_random_128_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_128_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp1_random_2048_128",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_random_2048_128",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_2048_128",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp1_random_2048_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_random_2048_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_2048_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 2048
+      }
+    },
     {
       "test_name": "serving_llama8B_int4_tp1_random_128_128",
       "server_parameters": {
-        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4"
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_int4_tp2_random_128_128",
+      "server_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_int4_tp4_random_128_128",
+      "server_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "tensor_parallel_size": 4
       },
       "client_parameters": {
         "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
@@ -56,7 +225,8 @@
     {
       "test_name": "serving_llama8B_int8_tp1_random_128_128",
       "server_parameters": {
-        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8"
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
@@ -66,12 +236,26 @@
       }
     },
     {
-      "test_name": "serving_llama1B_tp1_random_128_128",
+      "test_name": "serving_llama8B_int8_tp2_random_128_128",
       "server_parameters": {
-        "model": "meta-llama/Llama-3.2-1B"
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "tensor_parallel_size": 2
       },
       "client_parameters": {
-        "model": "meta-llama/Llama-3.2-1B",
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_int8_tp4_random_128_128",
+      "server_parameters": {
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
         "dataset_name": "random",
         "random-input-len": 128,
         "random-output-len": 128
@@ -80,7 +264,8 @@
     {
       "test_name": "serving_llama3B_tp1_random_128_128",
       "server_parameters": {
-        "model": "meta-llama/Llama-3.2-3B-Instruct"
+        "model": "meta-llama/Llama-3.2-3B-Instruct",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "meta-llama/Llama-3.2-3B-Instruct",
@@ -89,22 +274,11 @@
         "random-output-len": 128
       }
     },
-    {
-      "test_name": "serving_llama70B_tp1_random_128_128",
-      "server_parameters": {
-        "model": "meta-llama/Llama-3.3-70B-Instruct"
-      },
-      "client_parameters": {
-        "model": "meta-llama/Llama-3.3-70B-Instruct",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
     {
       "test_name": "serving_granite2B_tp1_random_128_128",
       "server_parameters": {
-        "model": "ibm-granite/granite-3.2-2b-instruct"
+        "model": "ibm-granite/granite-3.2-2b-instruct",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "ibm-granite/granite-3.2-2b-instruct",
@@ -116,7 +290,8 @@
     {
       "test_name": "serving_qwen1.7B_tp1_random_128_128",
       "server_parameters": {
-        "model": "Qwen/Qwen3-1.7B"
+        "model": "Qwen/Qwen3-1.7B",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "Qwen/Qwen3-1.7B",
@@ -128,7 +303,8 @@
     {
       "test_name": "serving_qwen4B_tp1_random_128_128",
       "server_parameters": {
-        "model": "Qwen/Qwen3-4B"
+        "model": "Qwen/Qwen3-4B",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "Qwen/Qwen3-4B",
@@ -140,7 +316,8 @@
     {
       "test_name": "serving_qwen8B_tp1_random_128_128",
       "server_parameters": {
-        "model": "Qwen/Qwen3-8B"
+        "model": "Qwen/Qwen3-8B",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "Qwen/Qwen3-8B",
@@ -149,34 +326,11 @@
         "random-output-len": 128
       }
     },
-    {
-      "test_name": "serving_qwen14B_tp1_random_128_128",
-      "server_parameters": {
-        "model": "Qwen/Qwen3-14B"
-      },
-      "client_parameters": {
-        "model": "Qwen/Qwen3-14B",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_qwen30B_tp1_random_128_128",
-      "server_parameters": {
-        "model": "Qwen/Qwen3-30B-A3B"
-      },
-      "client_parameters": {
-        "model": "Qwen/Qwen3-30B-A3B",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
     {
       "test_name": "serving_glm9B_tp1_random_128_128",
       "server_parameters": {
-        "model": "zai-org/glm-4-9b-hf"
+        "model": "zai-org/glm-4-9b-hf",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "zai-org/glm-4-9b-hf",
@@ -188,7 +342,8 @@
     {
       "test_name": "serving_gemma7B_tp1_random_128_128",
       "server_parameters": {
-        "model": "google/gemma-7b"
+        "model": "google/gemma-7b",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "google/gemma-7b",
@@ -196,94 +351,6 @@
         "random-input-len": 128,
         "random-output-len": 128
       }
-    },
-    {
-      "test_name": "serving_gemma3-4b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-3-4b-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-3-4b-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma3-12b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-3-12b-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-3-12b-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma4-4b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-4-E4B-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-4-E4B-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma4-2b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-4-E2B-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-4-E2B-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma4-26b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0,
-        "VLLM_CPU_ATTN_SPLIT_KV": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-4-26B-A4B-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-4-26B-A4B-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_phi4_tp1_random_128_128",
-      "server_parameters": {
-        "model": "microsoft/Phi-4-reasoning"
-      },
-      "client_parameters": {
-        "model": "microsoft/Phi-4-reasoning",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
     }
   ]
 }

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

@@ -1,112 +1,77 @@
-{
-  "defaults": {
-    "qps_list": [
-      "inf"
-    ],
-    "max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
-    "server_parameters": {
-      "model": "meta-llama/Llama-3.1-8B-Instruct",
-      "tensor_parallel_size": 1,
-      "dtype": "bfloat16"
+[
+    {
+        "test_name": "serving_llama8B_tp1_sharegpt",
+        "qps_list": [1, 4, 16, "inf"],
+        "server_parameters": {
+            "model": "meta-llama/Meta-Llama-3.1-8B-Instruct",
+            "tensor_parallel_size": 1,
+            "disable_log_stats": "",
+            "load_format": "dummy"
+        },
+        "client_parameters": {
+            "model": "meta-llama/Meta-Llama-3.1-8B-Instruct",
+            "backend": "vllm",
+            "dataset_name": "sharegpt",
+            "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
+            "temperature": 0,
+            "num_prompts": 200
+        }
     },
-    "client_parameters": {
-      "model": "meta-llama/Llama-3.1-8B-Instruct",
-      "backend": "vllm",
-      "ignore-eos": "",
-      "temperature": 0,
-      "num_prompts": 200
+    {
+        "test_name": "serving_llama70B_tp4_sharegpt",
+        "qps_list": [1, 4, 16, "inf"],
+        "server_parameters": {
+            "model": "meta-llama/Meta-Llama-3.1-70B-Instruct",
+            "tensor_parallel_size": 4,
+            "disable_log_stats": "",
+            "load_format": "dummy"
+        },
+        "client_parameters": {
+            "model": "meta-llama/Meta-Llama-3.1-70B-Instruct",
+            "backend": "vllm",
+            "dataset_name": "sharegpt",
+            "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
+            "temperature": 0,
+            "num_prompts": 200
+        }
+    },
+    {
+        "test_name": "serving_mixtral8x7B_tp2_sharegpt",
+        "qps_list": [1, 4, 16, "inf"],
+        "server_parameters": {
+            "model": "mistralai/Mixtral-8x7B-Instruct-v0.1",
+            "tensor_parallel_size": 2,
+            "disable_log_stats": "",
+            "load_format": "dummy"
+        },
+        "client_parameters": {
+            "model": "mistralai/Mixtral-8x7B-Instruct-v0.1",
+            "backend": "vllm",
+            "dataset_name": "sharegpt",
+            "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
+            "temperature": 0,
+            "num_prompts": 200
+        }
+    },
+    {
+        "test_name": "serving_llama70B_tp4_sharegpt_specdecode",
+        "qps_list": [2],
+        "server_parameters": {
+            "model": "meta-llama/Meta-Llama-3.1-70B-Instruct", 
+            "tensor_parallel_size": 4,
+            "speculative_config": {
+                "model": "turboderp/Qwama-0.5B-Instruct",
+                "num_speculative_tokens": 4,
+                "draft_tensor_parallel_size": 1
+            }
+        },
+        "client_parameters": {
+            "model": "meta-llama/Meta-Llama-3.1-70B-Instruct",
+            "backend": "vllm",
+            "dataset_name": "sharegpt",
+            "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
+            "temperature": 0,
+            "num_prompts": 200 
+        }
     }
-  },
-  "tests": [
-    {
-      "test_name": "serving_llama8B_tp1_sharegpt",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "sharegpt",
-        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
-      }
-    },
-    {
-        "dataset_name": "sharegpt",
-        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp1_random_128_128",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp1_random_128_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 2048
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp1_random_2048_128",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_llama8B_tp1_random_2048_2048",
-      "server_parameters": {
-        "tensor_parallel_size": 1
-      },
-      "client_parameters": {
-        "dataset_name": "random",
-        "random-input-len": 2048,
-        "random-output-len": 2048
-      }
-    },
-    {
-      "test_name": "serving_llama70B_tp4_random_128_128",
-      "server_parameters": {
-        "model": "meta-llama/Llama-3.3-70B-Instruct",
-        "async_scheduling": "",
-	"no_enable_prefix_caching": "",
-        "max_num_batched_tokens": 8192
-      },
-      "client_parameters": {
-        "model": "meta-llama/Llama-3.3-70B-Instruct",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma4-e4b_tp1_random_128_128",
-      "server_parameters": {
-        "model": "google/gemma-4-E4B-it",
-	"enable_auto_tool_choice": "",
-	"tool_call_parser": "gemma4",
-	"chat_template": "examples/tool_chat_template_gemma4.jinja",
-	"reasoning_parser": "gemma4"
-      },
-      "client_parameters": {
-        "model": "google/gemma-4-E4B-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    }
-  ]
-}
+]

.buildkite/release-pipeline.yaml

@@ -7,9 +7,6 @@ env:
   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"
-  MOONCAKE_WHEEL_AARCH64_2_35: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_aarch64.whl"
-  MOONCAKE_WHEEL_AARCH64_2_39: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_39_aarch64.whl"
-  MOONCAKE_WHEEL_X86_64: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_x86_64.whl"
 
 steps:
   - input: "Provide Release version here"
@@ -27,7 +24,7 @@ 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}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda12.9 --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 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64_CU129}\" --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"
@@ -69,7 +66,7 @@ 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}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinux2_28-builder:cuda12.9 --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 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86_CU129}\" --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"
@@ -139,8 +136,6 @@ steps:
               --build-arg CUDA_VERSION=13.0.2 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_35}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 \
               --target vllm-openai \
               --progress plain \
@@ -167,8 +162,6 @@ steps:
               --build-arg CUDA_VERSION=13.0.2 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_35}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 \
               --target vllm-openai \
               --progress plain \
@@ -192,8 +185,6 @@ steps:
               --build-arg CUDA_VERSION=12.9.1 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86_CU129}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_35}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --target vllm-openai \
               --progress plain \
               -f docker/Dockerfile .
@@ -219,8 +210,6 @@ steps:
               --build-arg CUDA_VERSION=12.9.1 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64_CU129}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_35}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --target vllm-openai \
               --progress plain \
               -f docker/Dockerfile .
@@ -245,8 +234,6 @@ steps:
               --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_39}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 \
               --target vllm-openai \
               --progress plain \
@@ -274,8 +261,6 @@ steps:
               --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_39}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 \
               --target vllm-openai \
               --progress plain \
@@ -301,8 +286,6 @@ steps:
               --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86_CU129}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_39}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --target vllm-openai \
               --progress plain \
               -f docker/Dockerfile .
@@ -329,8 +312,6 @@ steps:
               --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
               --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64_CU129}" \
               --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg MOONCAKE_WHEEL_AARCH64="${MOONCAKE_WHEEL_AARCH64_2_39}" \
-              --build-arg MOONCAKE_WHEEL_X86_64="${MOONCAKE_WHEEL_X86_64}" \
               --target vllm-openai \
               --progress plain \
               -f docker/Dockerfile .

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

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

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

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

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

@@ -114,7 +114,8 @@ handle_pytest_exit() {
 # unquoted since they have no spaces and work fine.
 #
 # Already-quoted expressions (containing literal single quotes) are passed
-# through untouched to avoid double-quoting well-formed shell fragments.
+# through untouched to avoid double-quoting values injected by
+# apply_rocm_test_overrides.
 #
 # NOTE: This ONLY fixes -m/-k flags. It cannot recover arbitrary inner
 # double-quotes stripped by the calling shell (see header comment).
@@ -246,6 +247,102 @@ re_quote_pytest_markers() {
   echo "${output% }"
 }
 
+###############################################################################
+# ROCm-specific pytest command rewrites
+#
+# These apply ignore flags and environment overrides for tests that are not
+# yet supported or behave differently on ROCm hardware. Kept as a single
+# function so new exclusions are easy to add in one place.
+###############################################################################
+
+apply_rocm_test_overrides() {
+  local cmds="$1"
+
+  # --- Model registry filter ---
+  if [[ $cmds == *"pytest -v -s models/test_registry.py"* ]]; then
+    cmds=${cmds//"pytest -v -s models/test_registry.py"/"pytest -v -s models/test_registry.py -k 'not BambaForCausalLM and not GritLM and not Mamba2ForCausalLM and not Zamba2ForCausalLM'"}
+  fi
+
+  # --- LoRA: disable custom paged attention ---
+  if [[ $cmds == *"pytest -v -s lora"* ]]; then
+    cmds=${cmds//"pytest -v -s lora"/"pytest -v -s lora"}
+  fi
+
+  # --- Kernel ignores ---
+  if [[ $cmds == *" kernels/core"* ]]; then
+    cmds="${cmds} \
+    --ignore=kernels/core/test_fused_quant_layernorm.py \
+    --ignore=kernels/core/test_permute_cols.py"
+  fi
+
+  if [[ $cmds == *" kernels/attention"* ]]; then
+    cmds="${cmds} \
+    --ignore=kernels/attention/test_attention_selector.py \
+    --ignore=kernels/attention/test_encoder_decoder_attn.py \
+    --ignore=kernels/attention/test_flash_attn.py \
+    --ignore=kernels/attention/test_flashinfer.py \
+    --ignore=kernels/attention/test_prefix_prefill.py \
+    --ignore=kernels/attention/test_cascade_flash_attn.py \
+    --ignore=kernels/attention/test_mha_attn.py \
+    --ignore=kernels/attention/test_lightning_attn.py \
+    --ignore=kernels/attention/test_attention.py"
+  fi
+
+  if [[ $cmds == *" kernels/quantization"* ]]; then
+    cmds="${cmds} \
+    --ignore=kernels/quantization/test_int8_quant.py \
+    --ignore=kernels/quantization/test_machete_mm.py \
+    --ignore=kernels/quantization/test_block_fp8.py \
+    --ignore=kernels/quantization/test_block_int8.py \
+    --ignore=kernels/quantization/test_marlin_gemm.py \
+    --ignore=kernels/quantization/test_cutlass_scaled_mm.py \
+    --ignore=kernels/quantization/test_int8_kernel.py"
+  fi
+
+  if [[ $cmds == *" kernels/mamba"* ]]; then
+    cmds="${cmds} \
+    --ignore=kernels/mamba/test_mamba_mixer2.py \
+    --ignore=kernels/mamba/test_causal_conv1d.py \
+    --ignore=kernels/mamba/test_mamba_ssm_ssd.py"
+  fi
+
+  if [[ $cmds == *" kernels/moe"* ]]; then
+    cmds="${cmds} \
+    --ignore=kernels/moe/test_moe.py \
+    --ignore=kernels/moe/test_cutlass_moe.py"
+  fi
+
+  # --- Entrypoint ignores ---
+  if [[ $cmds == *" entrypoints/openai "* ]]; then
+    cmds=${cmds//" entrypoints/openai "/" entrypoints/openai \
+    --ignore=entrypoints/openai/chat_completion/test_audio.py \
+    --ignore=entrypoints/openai/completion/test_shutdown.py \
+    --ignore=entrypoints/openai/test_completion.py \
+    --ignore=entrypoints/openai/models/test_models.py \
+    --ignore=entrypoints/openai/test_return_tokens_as_ids.py \
+    --ignore=entrypoints/openai/chat_completion/test_root_path.py \
+    --ignore=entrypoints/openai/completion/test_prompt_validation.py "}
+  fi
+
+  if [[ $cmds == *" entrypoints/serve"* ]]; then
+    cmds="${cmds} \
+    --ignore=entrypoints/serve/lora/test_lora_adapters.py"
+  fi
+
+  if [[ $cmds == *" entrypoints/llm "* ]]; then
+    cmds=${cmds//" entrypoints/llm "/" entrypoints/llm \
+    --ignore=entrypoints/llm/test_chat.py \
+    --ignore=entrypoints/llm/test_accuracy.py \
+    --ignore=entrypoints/llm/test_init.py \
+    --ignore=entrypoints/llm/test_prompt_validation.py "}
+  fi
+
+  # Clean up escaped newlines from --ignore appends
+  cmds=$(echo "$cmds" | sed 's/ \\ / /g')
+
+  echo "$cmds"
+}
+
 ###############################################################################
 # Main
 ###############################################################################
@@ -312,6 +409,7 @@ else
   echo "Skipping re-quoting for VLLM_TEST_COMMANDS input"
 fi
 
+commands=$(apply_rocm_test_overrides "$commands")
 echo "Final commands: $commands"
 
 MYPYTHONPATH=".."

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

@@ -352,31 +352,17 @@ if [[ -z "${ZE_AFFINITY_MASK:-}" ]]; then
   echo "Warning: ZE_AFFINITY_MASK is not set. Proceeding without device affinity." >&2
 fi
 
-export CMDS="${commands}"
-export HF_TOKEN ZE_AFFINITY_MASK
-
-{
-  flock 9
-  if ! docker image inspect "${IMAGE}" >/dev/null 2>&1; then
-    echo 'Image missing before container creation, pulling again...'
-    timeout 900 docker pull "${IMAGE}"
-  fi
-
-  docker create \
+docker run \
     --device /dev/dri:/dev/dri \
     --net=host \
     --ipc=host \
     --privileged \
     -v /dev/dri/by-path:/dev/dri/by-path \
-    -v "${HOME}/.cache/huggingface:/root/.cache/huggingface" \
-    --entrypoint='' \
-    -e HF_TOKEN \
-    -e ZE_AFFINITY_MASK \
-    -e CMDS \
+    -v ${HOME}/.cache/huggingface:/root/.cache/huggingface \
+    --entrypoint="" \
+    -e "HF_TOKEN=${HF_TOKEN:-}" \
+    -e "ZE_AFFINITY_MASK=${ZE_AFFINITY_MASK:-}" \
+    -e "CMDS=${commands}" \
     --name "${container_name}" \
-    "${IMAGE}" \
-    bash -c 'set -e; echo "ZE_AFFINITY_MASK is ${ZE_AFFINITY_MASK:-}"; eval "$CMDS"' \
-    >/dev/null
-} 9>/tmp/docker-pull.lock
-
-docker start -a "${container_name}"
+    "${image_name}" \
+    bash -c 'set -e; echo "ZE_AFFINITY_MASK is ${ZE_AFFINITY_MASK:-}"; eval "$CMDS"'

.buildkite/scripts/hardware_ci/run-tpu-v1-test-part2.sh

@@ -61,7 +61,7 @@ echo "Results will be stored in: $RESULTS_DIR"
 echo "--- Installing Python dependencies ---"
 python3 -m pip install --progress-bar off git+https://github.com/thuml/depyf.git \
     && python3 -m pip install --progress-bar off pytest pytest-asyncio tpu-info \
-    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.12" \
+    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.11" \
     && python3 -m pip install --progress-bar off hf-transfer tblib==3.1.0
 echo "--- Python dependencies installed ---"
 

.buildkite/scripts/hardware_ci/run-tpu-v1-test.sh

@@ -61,7 +61,7 @@ echo "Results will be stored in: $RESULTS_DIR"
 echo "--- Installing Python dependencies ---"
 python3 -m pip install --progress-bar off git+https://github.com/thuml/depyf.git \
     && python3 -m pip install --progress-bar off pytest pytest-asyncio tpu-info \
-    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.12" \
+    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.11" \
     && python3 -m pip install --progress-bar off hf-transfer tblib==3.1.0
 echo "--- Python dependencies installed ---"
 

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

@@ -0,0 +1,53 @@
+#!/bin/bash
+
+# This script build the CPU docker image and run the offline inference inside the container.
+# It serves a sanity check for compilation and basic model usage.
+set -ex
+
+image_name="xpu/vllm-ci:${BUILDKITE_COMMIT}"
+container_name="xpu_${BUILDKITE_COMMIT}_$(tr -dc A-Za-z0-9 < /dev/urandom | head -c 10; echo)"
+
+# Try building the docker image
+docker build -t "${image_name}" -f docker/Dockerfile.xpu .
+
+# Setup cleanup
+remove_docker_container() {
+  docker rm -f "${container_name}" || true
+}
+trap remove_docker_container EXIT
+
+# Run the image and test offline inference/tensor parallel
+docker run \
+    --device /dev/dri:/dev/dri \
+    --net=host \
+    --ipc=host \
+    --privileged \
+    -v /dev/dri/by-path:/dev/dri/by-path \
+    --entrypoint="" \
+    -e "HF_TOKEN=${HF_TOKEN}" \
+    -e "ZE_AFFINITY_MASK=${ZE_AFFINITY_MASK}" \
+    --name "${container_name}" \
+    "${image_name}" \
+    bash -c '
+    set -e
+    echo $ZE_AFFINITY_MASK
+    python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager
+    python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 -O3 -cc.cudagraph_mode=NONE
+    python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager -tp 2 --distributed-executor-backend ray
+    python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager -tp 2 --distributed-executor-backend mp
+    python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --attention-backend=TRITON_ATTN
+    python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --quantization fp8
+    python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4  --block-size 64 --enforce-eager --max-model-len 8192
+    python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b  --block-size 64 --enforce-eager -tp 2
+    python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b  --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel
+    python3 examples/basic/offline_inference/generate.py --model OPEA/Qwen2.5-0.5B-Instruct-int4-sym-inc --block-size 64 --enforce-eager --max-model-len 8192
+    cd tests
+    pytest -v -s v1/core --ignore=v1/core/test_reset_prefix_cache_e2e.py --ignore=v1/core/test_scheduler_e2e.py
+    pytest -v -s v1/engine
+    pytest -v -s v1/sample --ignore=v1/sample/test_logprobs.py --ignore=v1/sample/test_logprobs_e2e.py
+    pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py
+    pytest -v -s v1/structured_output
+    pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_tree_attention.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py
+    pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py
+    pytest -v -s v1/test_serial_utils.py
+'

.buildkite/scripts/run-rust-frontend-cargo-ci.sh

@@ -1,156 +0,0 @@
-#!/usr/bin/env bash
-set -euo pipefail
-
-MODE="${1:-}"
-
-if [[ "$MODE" != "style-clippy" && "$MODE" != "test" ]]; then
-  echo "Usage: $0 {style-clippy|test}" >&2
-  exit 2
-fi
-
-ROOT_DIR="$(git rev-parse --show-toplevel)"
-cd "$ROOT_DIR"
-
-export CARGO_TERM_COLOR="${CARGO_TERM_COLOR:-always}"
-export CARGO_HOME="${CARGO_HOME:-$HOME/.cargo}"
-export RUSTUP_HOME="${RUSTUP_HOME:-$HOME/.rustup}"
-export PATH="$CARGO_HOME/bin:$PATH"
-
-log_section() {
-  echo "--- $*"
-}
-
-install_protoc() {
-  if command -v protoc >/dev/null 2>&1; then
-    return
-  fi
-
-  local version="${PROTOC_VERSION:-31.1}"
-  local arch
-  case "$(uname -m)" in
-    x86_64)
-      arch="x86_64"
-      ;;
-    aarch64|arm64)
-      arch="aarch_64"
-      ;;
-    *)
-      echo "Unsupported protoc architecture: $(uname -m)" >&2
-      return 1
-      ;;
-  esac
-
-  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${version}/protoc-${version}-linux-${arch}.zip"
-  local tmp_dir
-  tmp_dir="$(mktemp -d)"
-
-  log_section "Installing protoc ${version}"
-  curl -L --proto '=https' --tlsv1.2 -sSf "$url" -o "$tmp_dir/protoc.zip"
-  mkdir -p "$CARGO_HOME/bin"
-  unzip -q "$tmp_dir/protoc.zip" bin/protoc 'include/*' -d "$CARGO_HOME"
-  chmod +x "$CARGO_HOME/bin/protoc"
-  rm -rf "$tmp_dir"
-}
-
-rust_toolchain() {
-  awk -F '"' '/channel[[:space:]]*=/ { print $2; exit }' rust-toolchain.toml
-}
-
-install_rust_toolchain() {
-  log_section "Installing Rust toolchain"
-  if ! command -v rustup >/dev/null 2>&1; then
-    curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs \
-      | sh -s -- -y --profile minimal --default-toolchain none
-  fi
-
-  local toolchain
-  toolchain="$(rust_toolchain)"
-  rustup toolchain install "$toolchain" --profile minimal --component rustfmt,clippy
-  rustup component add --toolchain "$toolchain" rustfmt clippy
-}
-
-install_cargo_binstall() {
-  if command -v cargo-binstall >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-binstall"
-  curl -L --proto '=https' --tlsv1.2 -sSf \
-    https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
-    | bash
-}
-
-install_cargo_sort() {
-  if command -v cargo-sort >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-sort"
-  install_cargo_binstall
-  cargo binstall --no-confirm cargo-sort
-}
-
-install_cargo_nextest() {
-  if command -v cargo-nextest >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-nextest"
-  install_cargo_binstall
-  cargo binstall --no-confirm --secure cargo-nextest
-}
-
-install_uv() {
-  if command -v uv >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing uv"
-  curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.sh \
-    | env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
-}
-
-run_style_clippy() {
-  install_cargo_sort
-
-  log_section "Checking Rust formatting"
-  cargo fmt --manifest-path rust/Cargo.toml --all -- --check
-
-  log_section "Checking Cargo.toml ordering"
-  cargo sort --workspace --check rust
-
-  log_section "Running clippy"
-  cargo clippy \
-    --manifest-path rust/Cargo.toml \
-    --workspace \
-    --all-targets \
-    --all-features \
-    --locked \
-    -- \
-    -D warnings
-}
-
-run_tests() {
-  install_uv
-  install_cargo_nextest
-
-  log_section "Running cargo nextest"
-  cargo nextest run \
-    --manifest-path rust/Cargo.toml \
-    --workspace \
-    --all-features \
-    --locked \
-    --no-fail-fast
-}
-
-install_protoc
-install_rust_toolchain
-
-case "$MODE" in
-  style-clippy)
-    run_style_clippy
-    ;;
-  test)
-    run_tests
-    ;;
-esac

.buildkite/test-amd.yaml

@@ -139,6 +139,19 @@ steps:
   - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py
   - VLLM_WORKER_MULTIPROC_METHOD=spawn pytest models/multimodal/generation/test_whisper.py -v -s -m 'distributed(num_gpus=2)'
 
+#--------------------------------------------------------  mi250 · benchmarks  ---------------------------------------------------------#
+
+- label: Benchmarks # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
+  agent_pool: mi250_1
+  working_dir: "/vllm-workspace/.buildkite"
+  source_file_dependencies:
+  - benchmarks/
+  - vllm/platforms/rocm.py
+  commands:
+  - bash scripts/run-benchmarks.sh
+
 #----------------------------------------------------------  mi250 · compile  ----------------------------------------------------------#
 
 - label: PyTorch Compilation Unit Tests # TBD
@@ -472,7 +485,7 @@ steps:
   - pytest -v -s model_executor -m '(not slow_test)'
   - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
 
-#------------------------------------------------------  mi250 · models / basic  -------------------------------------------------------#
+#----------------------------------------------------------  mi250 · models  -----------------------------------------------------------#
 
 - label: Basic Models Test (Other CPU) # TBD
   timeout_in_minutes: 180
@@ -533,8 +546,6 @@ steps:
   commands:
   - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
 
-#-----------------------------------------------------  mi250 · models / language  -----------------------------------------------------#
-
 - label: Language Models Test (MTEB) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -579,8 +590,6 @@ steps:
   - pip freeze | grep -E 'torch'
   - pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
-#----------------------------------------------------  mi250 · models / multimodal  ----------------------------------------------------#
-
 - label: Multi-Modal Models (Extended Generation 2) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -789,6 +798,32 @@ steps:
   commands:
   - pytest -v -s v1/attention
 
+- label: V1 Core + KV + Metrics # TBD
+  timeout_in_minutes: 60
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
+  agent_pool: mi250_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  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:
+  - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+  - pytest -v -s -m 'not cpu_test' v1/core
+  - pytest -v -s v1/executor
+  - pytest -v -s v1/kv_offload
+  - pytest -v -s v1/worker
+  - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
+  - pytest -v -s -m 'not cpu_test' v1/metrics
+  - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
+  - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
+
 - label: V1 Sample + Logits # TBD
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -967,6 +1002,18 @@ steps:
 
 #--------------------------------------------------------  mi300 · benchmarks  ---------------------------------------------------------#
 
+- label: Benchmarks # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/.buildkite"
+  source_file_dependencies:
+  - benchmarks/
+  - vllm/platforms/rocm.py
+  commands:
+  - bash scripts/run-benchmarks.sh
+
 - label: Benchmarks CLI Test # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -1232,7 +1279,6 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  optional: true
   fast_check: true
   torch_nightly: true
   working_dir: "/vllm-workspace/tests"
@@ -1253,7 +1299,6 @@ steps:
   agent_pool: mi300_1
   fast_check: true
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -1267,7 +1312,6 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  optional: true
   fast_check: true
   torch_nightly: true
   working_dir: "/vllm-workspace/tests"
@@ -1284,7 +1328,6 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  optional: true
   fast_check: true
   torch_nightly: true
   working_dir: "/vllm-workspace/tests"
@@ -1685,7 +1728,6 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  optional: true
   parallelism: 2
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -1738,13 +1780,12 @@ steps:
   - pytest -v -s -x lora/test_gptoss_tp.py
   - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
 
-#-----------------------------------------------------  mi300 · models / language  -----------------------------------------------------#
+#----------------------------------------------------------  mi300 · models  -----------------------------------------------------------#
 
 - label: Language Models Test (Extended Pooling)  # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -1766,8 +1807,6 @@ steps:
   - pip freeze | grep -E 'torch'
   - pytest -v -s models/language -m 'core_model and (not slow_test)'
 
-#----------------------------------------------------  mi300 · models / multimodal  ----------------------------------------------------#
-
 - label: Multi-Modal Models (Extended Generation 1) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -1873,11 +1912,10 @@ steps:
   - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
   - pytest -v -s models/multimodal/processing/test_tensor_schema.py
 
-- label: Multi-Modal Processor (CPU) %N # TBD
+- label: Multi-Modal Processor (CPU) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  parallelism: 4
   no_gpu: true
   optional: true
   working_dir: "/vllm-workspace/tests"
@@ -1887,9 +1925,7 @@ steps:
   - tests/models/registry.py
   commands:
   - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
-  - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-
-#-----------------------------------------------------  mi300 · models / quantized  -----------------------------------------------------#
+  - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
 
 - label: Quantized Models Test # TBD
   timeout_in_minutes: 180
@@ -1905,31 +1941,7 @@ steps:
   commands:
   - pytest -v -s models/quantization
 
-#--------------------------------------------------  mi300 · models / transformers  ---------------------------------------------------#
-
-- label: Transformers Nightly Models (Shardable) %N # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
-  agent_pool: mi300_1
-  parallelism: 4
-  optional: true
-  working_dir: "/vllm-workspace/"
-  source_file_dependencies:
-  - vllm/model_executor/models/
-  - vllm/model_executor/model_loader/
-  - vllm/multimodal/
-  - vllm/model_executor/layers/
-  - vllm/v1/attention/backends/
-  - vllm/v1/attention/selector.py
-  - vllm/_aiter_ops.py
-  - vllm/platforms/rocm.py
-  - tests/models/
-  commands:
-  - pip install --upgrade git+https://github.com/huggingface/transformers
-  - pytest -v -s tests/models/test_initialization.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-  - pytest -v -s tests/models/multimodal/processing/ --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-
-- label: Transformers Nightly Models (Single) # TBD
+- label: Transformers Nightly Models # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
@@ -1948,7 +1960,9 @@ steps:
   - examples/
   commands:
   - pip install --upgrade git+https://github.com/huggingface/transformers
+  - pytest -v -s tests/models/test_initialization.py
   - pytest -v -s tests/models/test_transformers.py
+  - pytest -v -s tests/models/multimodal/processing/
   - pytest -v -s tests/models/multimodal/test_mapping.py
   - python3 examples/basic/offline_inference/chat.py
   - python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
@@ -2009,7 +2023,7 @@ steps:
   - vllm/model_executor/layers
   - vllm/sampling_metadata.py
   - vllm/v1/sample/
-  - vllm/entrypoints/generate/beam_search/
+  - vllm/beam_search.py
   - tests/samplers
   - tests/conftest.py
   - vllm/_aiter_ops.py
@@ -2397,7 +2411,7 @@ steps:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
   - DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
-- label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs) # TBD
+- label: Hyrbid SSM NixlConnector PD accuracy tests (4 GPUs) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_4
@@ -2599,7 +2613,7 @@ steps:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-rocm.txt --tp-size=8
 
-#-----------------------------------------------------  mi325 · models / language  -----------------------------------------------------#
+#----------------------------------------------------------  mi325 · models  -----------------------------------------------------------#
 
 - label: Language Models Test (Extended Generation) # TBD
   timeout_in_minutes: 180
@@ -2630,8 +2644,6 @@ steps:
   - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
   - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
-#----------------------------------------------------  mi325 · models / multimodal  ----------------------------------------------------#
-
 - label: Multi-Modal Models (Extended Pooling) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
@@ -3051,7 +3063,7 @@ steps:
   commands:
     - pytest -v -s kernels/moe/test_deepep_moe.py
 
-#-----------------------------------------------------  mi355 · models / language  -----------------------------------------------------#
+#----------------------------------------------------------  mi355 · models  -----------------------------------------------------------#
 
 - label: Language Models Test (Extended Generation) # TBD
   timeout_in_minutes: 180
@@ -3119,8 +3131,6 @@ steps:
   - pip freeze | grep -E 'torch'
   - pytest -v -s models/language -m 'core_model and (not slow_test)'
 
-#----------------------------------------------------  mi355 · models / multimodal  ----------------------------------------------------#
-
 - label: Multi-Modal Models (Extended Generation 1) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -3192,8 +3202,6 @@ steps:
   - pytest -v -s models/multimodal/generation/test_memory_leak.py -m core_model
   - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model
 
-#-----------------------------------------------------  mi355 · models / quantized  -----------------------------------------------------#
-
 - label: Quantized Models Test # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]

.buildkite/test_areas/cuda.yaml

@@ -27,5 +27,4 @@ steps:
   - vllm/compilation
   commands:
     - pytest -v -s v1/cudagraph/test_cudagraph_dispatch.py
-    - pytest -v -s v1/cudagraph/test_cudagraph_mode.py
-    - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
+    - pytest -v -s v1/cudagraph/test_cudagraph_mode.py

.buildkite/test_areas/disaggregated.yaml

@@ -88,7 +88,7 @@ steps:
     - 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_spec_decode_test.sh
+    - 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

.buildkite/test_areas/engine.yaml

@@ -36,12 +36,6 @@ steps:
   commands:
     - pytest -v -s v1/engine/test_preprocess_error_handling.py
     - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
-  mirror:
-    amd:
-      device: mi300_1
-      timeout_in_minutes: 40
-      depends_on:
-      - image-build-amd
 
 - label: e2e Scheduling (1 GPU)
   key: e2e-scheduling-1-gpu
@@ -52,15 +46,8 @@ steps:
     - tests/v1/e2e/general/
   commands:
     - pytest -v -s v1/e2e/general/test_async_scheduling.py
-  mirror:
-    amd:
-      device: mi250_1
-      timeout_in_minutes: 60
-      depends_on:
-      - image-build-amd
 
 - label: e2e Core (1 GPU)
-  device: h200_35gb
   key: e2e-core-1-gpu
   timeout_in_minutes: 30
   source_file_dependencies:

.buildkite/test_areas/entrypoints.yaml

@@ -43,12 +43,6 @@ steps:
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
-  mirror:
-    amd:
-      device: mi300_1
-      timeout_in_minutes: 80
-      depends_on:
-      - image-build-amd
 
 - label: Entrypoints Integration (API Server openai - Part 2)
   key: entrypoints-integration-api-server-openai-part-2
@@ -61,12 +55,6 @@ steps:
   commands:
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
   - pytest -v -s entrypoints/test_chat_utils.py
-  mirror:
-    amd:
-      device: mi300_1
-      timeout_in_minutes: 60
-      depends_on:
-      - image-build-amd
 
 - label: Entrypoints Integration (API Server openai - Part 3)
   key: entrypoints-integration-api-server-openai-part-3
@@ -79,16 +67,9 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
-  mirror:
-    amd:
-      device: mi300_1
-      timeout_in_minutes: 60
-      depends_on:
-      - image-build-amd
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion  --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
 
 - label: Entrypoints Integration (API Server 2)
-  device: h200_35gb
   key: entrypoints-integration-api-server-2
   timeout_in_minutes: 130
   working_dir: "/vllm-workspace/tests"
@@ -102,14 +83,8 @@ steps:
   - pytest -v -s entrypoints/serve/instrumentator
   - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
   - pytest -v -s tool_use
-  mirror:
-    amd:
-      device: mi300_1
-      depends_on:
-      - image-build-amd
 
 - label: Entrypoints Integration (Speech to Text)
-  device: h200_35gb
   key: entrypoints-integration-speech_to_text
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"

.buildkite/test_areas/kernels.yaml

@@ -62,19 +62,6 @@ steps:
   commands:
     - pytest -v -s kernels/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
   parallelism: 2
-  mirror:
-    amd:
-      device: mi300_1
-      source_file_dependencies:
-      - csrc/quantization/
-      - vllm/model_executor/layers/quantization
-      - tests/kernels/quantization
-      - tests/kernels/quantization/test_rocm_skinny_gemms.py
-      - vllm/_aiter_ops.py
-      - vllm/platforms/rocm.py
-      - vllm/model_executor/kernels/
-      depends_on:
-      - image-build-amd
 
 - label: Kernels MoE Test %N
   key: kernels-moe-test
@@ -150,12 +137,9 @@ steps:
   - csrc/quantization/fp4/
   - csrc/attention/mla/
   - csrc/quantization/cutlass_w8a8/moe/
-  - vllm/model_executor/layers/fused_moe/experts/cutlass_moe.py
-  - vllm/model_executor/layers/fused_moe/experts/flashinfer_cutlass_moe.py
-  - vllm/model_executor/layers/fused_moe/experts/trtllm_nvfp4_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/nvfp4.py
-  - vllm/model_executor/layers/fused_moe/prepare_finalize/flashinfer_nvlink_one_sided.py
-  - vllm/model_executor/layers/fused_moe/prepare_finalize/flashinfer_nvlink_two_sided.py
+  - vllm/model_executor/layers/fused_moe/cutlass_moe.py
+  - vllm/model_executor/layers/fused_moe/flashinfer_cutlass_moe.py
+  - vllm/model_executor/layers/fused_moe/flashinfer_a2a_prepare_finalize.py
   - vllm/model_executor/layers/quantization/utils/flashinfer_utils.py
   - vllm/v1/attention/backends/flashinfer.py
   - vllm/v1/attention/backends/mla/cutlass_mla.py
@@ -188,7 +172,6 @@ steps:
     - 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
-    - pytest -v -s tests/kernels/moe/test_trtllm_nvfp4_moe.py
     - pytest -v -s tests/kernels/moe/test_cutedsl_moe.py
     # e2e
     - pytest -v -s tests/models/quantization/test_nvfp4.py

.buildkite/test_areas/lm_eval.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: LM Eval Small Models
-  device: h200_35gb
   key: lm-eval-small-models
   timeout_in_minutes: 75
   source_file_dependencies:
@@ -153,7 +152,6 @@ steps:
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-b200.txt
 
 - label: MRCR Eval Small Models
-  device: h200_35gb
   timeout_in_minutes: 30
   source_file_dependencies:
   - tests/evals/mrcr/

.buildkite/test_areas/lora.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: LoRA %N
-  device: h200_35gb
   key: lora
   timeout_in_minutes: 30
   source_file_dependencies:

.buildkite/test_areas/misc.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: V1 Spec Decode
-  device: h200_35gb
   key: v1-spec-decode
   timeout_in_minutes: 30
   source_file_dependencies:
@@ -166,7 +165,6 @@ steps:
   working_dir: "/vllm-workspace/tests" # optional
 
 - label: Examples
-  device: h200_35gb
   key: examples
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/examples"
@@ -237,7 +235,6 @@ steps:
   - bash standalone_tests/python_only_compile.sh
 
 - label: Async Engine, Inputs, Utils, Worker
-  device: h200_35gb
   key: async-engine-inputs-utils-worker
   timeout_in_minutes: 50
   source_file_dependencies:
@@ -367,10 +364,8 @@ steps:
     - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
     - 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]
     - pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
-    - pytest -v -s v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py
   
 - label: Acceptance Length Test (Large Models) # optional
-  device: h200_35gb
   key: acceptance-length-test-large-models
   timeout_in_minutes: 25
   gpu: h100

.buildkite/test_areas/model_runner_v2.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Model Runner V2 Core Tests
-  device: h200_35gb
   key: model-runner-v2-core-tests
   timeout_in_minutes: 45
   source_file_dependencies:
@@ -27,7 +26,6 @@ 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
-  device: h200_35gb
   key: model-runner-v2-examples
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/examples"
@@ -101,7 +99,6 @@ steps:
     - pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
 
 - label: Model Runner V2 Spec Decode
-  device: h200_35gb
   key: model-runner-v2-spec-decode
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"

.buildkite/test_areas/models_basic.yaml

@@ -18,7 +18,6 @@ steps:
     torch_nightly: {}
 
 - label: Basic Models Tests (Extra Initialization) %N
-  device: h200_35gb
   key: basic-models-tests-extra-initialization
   timeout_in_minutes: 45
   source_file_dependencies:
@@ -35,7 +34,6 @@ steps:
     torch_nightly: {}
 
 - label: Basic Models Tests (Other)
-  device: h200_35gb
   key: basic-models-tests-other
   timeout_in_minutes: 45
   source_file_dependencies:
@@ -60,7 +58,6 @@ steps:
     - pytest -v -s models/test_utils.py models/test_vision.py
 
 - label: Transformers Nightly Models
-  device: h200_35gb
   key: transformers-nightly-models
   working_dir: "/vllm-workspace/"
   optional: true
@@ -77,7 +74,6 @@ steps:
     - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
 
 - label: Transformers Backward Compatibility Models Test
-  device: h200_35gb
   key: transformers-backward-compatibility-models-test
   working_dir: "/vllm-workspace/"
   optional: true

.buildkite/test_areas/models_language.yaml

@@ -59,7 +59,6 @@ steps:
       - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
 - label: Language Models Test (Extended Generation) # 80min
-  device: h200_35gb
   key: language-models-test-extended-generation
   timeout_in_minutes: 110
   optional: true
@@ -85,7 +84,6 @@ steps:
     - pytest -v -s models/language/generation_ppl_test
 
 - label: Language Models Test (Extended Pooling)  # 36min
-  device: h200_35gb
   key: language-models-test-extended-pooling
   timeout_in_minutes: 50
   optional: true
@@ -94,12 +92,6 @@ steps:
   - tests/models/language/pooling
   commands:
     - pytest -v -s models/language/pooling -m 'not core_model'
-  mirror:
-    amd:
-      device: mi300_1
-      timeout_in_minutes: 100
-      depends_on:
-      - image-build-amd
 
 - label: Language Models Test (MTEB)
   key: language-models-test-mteb

.buildkite/test_areas/models_multimodal.yaml

@@ -38,7 +38,6 @@ steps:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
-  device: h200_35gb
   key: multi-modal-models-standard-3-llava-qwen2-vl
   timeout_in_minutes: 45
   source_file_dependencies:
@@ -55,7 +54,6 @@ steps:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
-  device: h200_35gb
   key: multi-modal-models-standard-4-other-whisper
   timeout_in_minutes: 45
   source_file_dependencies:
@@ -94,7 +92,6 @@ steps:
     - pytest -v -s models/multimodal/processing/test_tensor_schema.py
 
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
-  device: h200_35gb
   key: multi-modal-accuracy-eval-small-models
   timeout_in_minutes: 70
   working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
@@ -123,7 +120,6 @@ steps:
       - image-build-amd
 
 - label: Multi-Modal Models (Extended Generation 2)
-  device: h200_35gb
   key: multi-modal-models-extended-generation-2
   optional: true
   source_file_dependencies:
@@ -134,7 +130,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)
-  device: h200_35gb
   key: multi-modal-models-extended-generation-3
   optional: true
   source_file_dependencies:

.buildkite/test_areas/pytorch.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: PyTorch Compilation Unit Tests
-  device: h200_35gb
   key: pytorch-compilation-unit-tests
   timeout_in_minutes: 10
   source_file_dependencies:

.buildkite/test_areas/rust_frontend.yaml

@@ -1,107 +0,0 @@
-group: Rust Frontend E2E
-depends_on:
-  - image-build
-steps:
-- label: Rust Frontend OpenAI Coverage
-  timeout_in_minutes: 90
-  device: h200_18gb
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - rust/
-  - vllm/benchmarks/
-  - vllm/entrypoints/openai/
-  - vllm/entrypoints/serve/
-  - vllm/v1/sample/
-  - tests/utils.py
-  - tests/benchmarks/test_serve_cli.py
-  - tests/entrypoints/openai/chat_completion/test_chat_completion.py
-  # - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
-  # - tests/entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py
-  # - tests/entrypoints/openai/completion/test_prompt_validation.py
-  - tests/entrypoints/openai/completion/test_shutdown.py
-  # - tests/entrypoints/openai/test_return_token_ids.py
-  # - tests/entrypoints/openai/test_uds.py
-  - tests/v1/sample/test_logprobs_e2e.py
-  commands:
-  - export VLLM_USE_RUST_FRONTEND=1
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
-  - pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py
-  # - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
-  # - pytest -v -s entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py
-  # - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
-  - pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
-  # - pytest -v -s entrypoints/openai/test_return_token_ids.py
-  # - pytest -v -s entrypoints/openai/test_uds.py
-  - pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
-
-- label: Rust Frontend Serve/Admin Coverage
-  timeout_in_minutes: 60
-  device: h200_18gb
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - rust/
-  - vllm/entrypoints/openai/
-  - vllm/entrypoints/serve/
-  - vllm/v1/engine/
-  - tests/utils.py
-  # - tests/entrypoints/rpc/test_collective_rpc.py
-  - tests/entrypoints/serve/disagg/test_serving_tokens.py
-  - tests/entrypoints/serve/instrumentator/test_basic.py
-  - tests/entrypoints/serve/instrumentator/test_metrics.py
-  # - tests/entrypoints/serve/instrumentator/test_sleep.py
-  commands:
-  - export VLLM_USE_RUST_FRONTEND=1
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  # - pytest -v -s entrypoints/rpc/test_collective_rpc.py
-  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
-  - pytest -v -s entrypoints/serve/disagg/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
-  - pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
-  # - pytest -v -s entrypoints/serve/instrumentator/test_sleep.py
-
-- label: Rust Frontend Core Correctness
-  timeout_in_minutes: 30
-  device: h200_18gb
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - rust/
-  - vllm/entrypoints/openai/
-  - tests/utils.py
-  - tests/entrypoints/openai/correctness/test_lmeval.py
-  commands:
-  - export VLLM_USE_RUST_FRONTEND=1
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
-
-- label: Rust Frontend Tool Use
-  timeout_in_minutes: 60
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - rust/
-  - vllm/entrypoints/openai/
-  - vllm/tool_parsers/
-  - tests/utils.py
-  - tests/tool_use/
-  commands:
-  - export VLLM_USE_RUST_FRONTEND=1
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
-
-- label: Rust Frontend Distributed
-  timeout_in_minutes: 30
-  num_devices: 4
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - rust/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/executor/
-  - vllm/v1/engine/
-  - vllm/v1/worker/
-  - tests/utils.py
-  - tests/v1/distributed/test_internal_lb_dp.py
-  commands:
-  - export VLLM_USE_RUST_FRONTEND=1
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - export NCCL_CUMEM_HOST_ENABLE=0
-  - TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py -k "not 4 and not server_info"

.buildkite/test_areas/rust_frontend_cargo.yaml

@@ -1,30 +0,0 @@
-group: Rust Frontend Cargo
-depends_on: []
-steps:
-- label: Rust Frontend Cargo Style + Clippy
-  key: rust-frontend-cargo-style-clippy
-  depends_on: []
-  timeout_in_minutes: 30
-  device: cpu-medium
-  no_plugin: true
-  source_file_dependencies:
-  - rust/
-  - rust-toolchain.toml
-  - .buildkite/test_areas/rust_frontend_cargo.yaml
-  - .buildkite/scripts/run-rust-frontend-cargo-ci.sh
-  commands:
-  - .buildkite/scripts/run-rust-frontend-cargo-ci.sh style-clippy
-
-- label: Rust Frontend Cargo Tests
-  key: rust-frontend-cargo-tests
-  depends_on: []
-  timeout_in_minutes: 30
-  device: cpu-medium
-  no_plugin: true
-  source_file_dependencies:
-  - rust/
-  - rust-toolchain.toml
-  - .buildkite/test_areas/rust_frontend_cargo.yaml
-  - .buildkite/scripts/run-rust-frontend-cargo-ci.sh
-  commands:
-  - .buildkite/scripts/run-rust-frontend-cargo-ci.sh test

.buildkite/test_areas/samplers.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Samplers Test
-  device: h200_35gb
   key: samplers-test
   timeout_in_minutes: 75
   source_file_dependencies:
@@ -11,7 +10,6 @@ steps:
   - vllm/sampling_metadata.py
   - tests/samplers
   - tests/conftest.py
-  - vllm/entrypoints/generate/beam_search
   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.

.dockerignore

@@ -2,7 +2,6 @@
 /build
 dist
 vllm/*.so
-vllm/vllm-rs
 
 # Byte-compiled / optimized / DLL files
 __pycache__/
@@ -32,4 +31,3 @@ share/python-wheels/
 .installed.cfg
 *.egg
 MANIFEST
-rust/target/

.github/CODEOWNERS

@@ -153,7 +153,9 @@ mkdocs.yaml @hmellor
 /vllm/model_executor/models/deepseek_mtp.py @luccafong
 
 # DeepseekV4-specific files
-/vllm/models/deepseek_v4 @zyongye
+/vllm/v1/attention/ops/deepseek_v4_ops @zyongye
+/vllm/model_executor/layers/deepseek_compressor.py @zyongye
+/vllm/model_executor/layers/deepseek_v4_attention.py @zyongye
 /vllm/model_executor/layers/sparse_attn_indexer.py @zyongye
 
 # Mistral-specific files
@@ -170,13 +172,13 @@ mkdocs.yaml @hmellor
 /vllm/model_executor/layers/fla @ZJY0516 @vadiklyutiy
 
 # ROCm related: specify owner with write access to notify AMD folks for careful code review
-/vllm/**/*rocm* @tjtanaa @dllehr-amd
-/docker/Dockerfile.rocm* @tjtanaa @dllehr-amd
-/vllm/v1/attention/backends/rocm*.py @tjtanaa @dllehr-amd
-/vllm/v1/attention/backends/mla/rocm*.py @tjtanaa @dllehr-amd
-/vllm/v1/attention/ops/rocm*.py @tjtanaa @dllehr-amd
-/vllm/model_executor/layers/fused_moe/rocm*.py @tjtanaa @dllehr-amd
-/csrc/rocm @tjtanaa @dllehr-amd
+/vllm/**/*rocm* @tjtanaa
+/docker/Dockerfile.rocm* @gshtras @tjtanaa
+/vllm/v1/attention/backends/rocm*.py @gshtras @tjtanaa
+/vllm/v1/attention/backends/mla/rocm*.py @gshtras @tjtanaa
+/vllm/v1/attention/ops/rocm*.py @gshtras @tjtanaa
+/vllm/model_executor/layers/fused_moe/rocm*.py @gshtras @tjtanaa
+/csrc/rocm @gshtras @tjtanaa
 /requirements/*rocm* @tjtanaa
 /tests/**/*rocm* @tjtanaa
 /docs/**/*rocm* @tjtanaa

.gitignore

@@ -26,9 +26,6 @@ __pycache__/
 # C extensions
 *.so
 
-# Rust binaries
-vllm/vllm-rs
-
 # Distribution / packaging
 .Python
 build/

.pre-commit-config.yaml

@@ -222,12 +222,6 @@ repos:
     name: Update Dockerfile dependency graph
     entry: tools/pre_commit/update-dockerfile-graph.sh
     language: script
-  - id: test-nonroot-entrypoint
-    name: Test non-root entrypoint wrapper
-    entry: bash docker/entrypoints/test_vllm_nonroot_entrypoint.sh
-    language: system
-    pass_filenames: false
-    files: ^docker/entrypoints/(vllm-nonroot-entrypoint|test_vllm_nonroot_entrypoint)\.sh$
   - id: check-forbidden-imports
     name: Check for forbidden imports
     entry: python tools/pre_commit/check_forbidden_imports.py
@@ -262,32 +256,6 @@ repos:
     entry: python tools/pre_commit/check_boolean_context_manager.py
     language: python
     types: [python]
-  # Rust hooks. These shell out to `cargo`; tools/pre_commit/rust-check.sh
-  # skips with a warning when cargo is not installed.
-  - id: rust-cargo-autoinherit
-    name: Rust - Normalize Cargo manifests with autoinherit
-    entry: tools/pre_commit/rust-check.sh autoinherit --prefer-simple-dotted
-    language: script
-    pass_filenames: false
-    require_serial: true
-    stages: [pre-commit] # Only run locally as Buildkite will cover this
-    files: ^rust/(Cargo\.toml|src/.*/Cargo\.toml)$
-  - id: rust-cargo-sort
-    name: Rust - Sort Cargo manifest sections
-    entry: tools/pre_commit/rust-check.sh sort --workspace
-    language: script
-    pass_filenames: false
-    require_serial: true
-    stages: [pre-commit] # Only run locally as Buildkite will cover this
-    files: ^rust/(Cargo\.toml|src/.*/Cargo\.toml)$
-  - id: rust-cargo-fmt
-    name: Rust - Format code
-    entry: tools/pre_commit/rust-check.sh fmt
-    language: script
-    pass_filenames: false
-    require_serial: true
-    stages: [pre-commit] # Only run locally as Buildkite will cover this
-    files: ^rust/.*(\.rs|Cargo\.toml|rustfmt\.toml)$
   # Keep `suggestion` last
   - id: suggestion
     name: Suggestion

.readthedocs.yaml

@@ -9,7 +9,44 @@ build:
     python: "3.12"
   jobs:
     post_checkout:
-      - bash docs/pre_run_check.sh
+      - |
+        if [ "$READTHEDOCS_VERSION_TYPE" = "external" ]; then
+          MAX_WAIT=300
+          INTERVAL=60
+          ELAPSED=0
+          while :; do
+            RAW=$(curl -sS -w "\n%{http_code}" "https://api.github.com/repos/vllm-project/vllm/commits/${READTHEDOCS_GIT_COMMIT_HASH}/check-runs?check_name=pre-run-check&filter=latest")
+            HTTP_CODE=$(printf %s "$RAW" | tail -n1)
+            BODY=$(printf %s "$RAW" | head -n -1)
+            if [ "$HTTP_CODE" != "200" ]; then
+              echo "GitHub API returned HTTP $HTTP_CODE (likely rate-limited); skipping pre-run-check gate."
+              break
+            fi
+            STATUS=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"status\") or \"\") if r else \"none\")")
+            CONCLUSION=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"conclusion\") or \"\") if r else \"\")")
+            if [ "$STATUS" = "none" ]; then
+              echo "no pre-run-check found for this commit; skipping gate."
+              break
+            fi
+            if [ -n "$CONCLUSION" ]; then
+              echo "pre-run-check conclusion: $CONCLUSION"
+              if [ "$CONCLUSION" = "failure" ] || [ "$CONCLUSION" = "cancelled" ] || [ "$CONCLUSION" = "timed_out" ]; then
+                echo "pre-run-check did not pass; failing docs build."
+                exit 1
+              fi
+              break
+            fi
+            if [ "$ELAPSED" -ge "$MAX_WAIT" ]; then
+              echo "pre-run-check status=$STATUS after ${MAX_WAIT}s; skipping gate."
+              break
+            fi
+            echo "pre-run-check status=$STATUS; waiting ${INTERVAL}s..."
+            sleep "$INTERVAL"
+            ELAPSED=$((ELAPSED + INTERVAL))
+          done
+        else
+          echo "Not a PR build (version type=$READTHEDOCS_VERSION_TYPE); skipping pre-run-check gate."
+        fi
       - git fetch origin main --unshallow --no-tags --filter=blob:none || true
     pre_create_environment:
       - pip install uv

CMakeLists.txt

@@ -311,10 +311,21 @@ set(VLLM_EXT_SRC
   "csrc/attention/paged_attention_v1.cu"
   "csrc/attention/paged_attention_v2.cu"
   "csrc/attention/merge_attn_states.cu"
+  "csrc/attention/vertical_slash_index.cu"
+  "csrc/pos_encoding_kernels.cu"
+  "csrc/activation_kernels.cu"
+  "csrc/layernorm_kernels.cu"
+  "csrc/fused_qknorm_rope_kernel.cu"
+  "csrc/layernorm_quant_kernels.cu"
   "csrc/sampler.cu"
   "csrc/topk.cu"
   "csrc/cuda_view.cu"
+  "csrc/quantization/gptq/q_gemm.cu"
+  "csrc/quantization/w8a8/int8/scaled_quant.cu"
+  "csrc/quantization/w8a8/fp8/common.cu"
+  "csrc/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
   "csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
+  "csrc/quantization/gguf/gguf_kernel.cu"
   "csrc/quantization/activation_kernels.cu"
   "csrc/cuda_utils_kernels.cu"
   "csrc/custom_all_reduce.cu"
@@ -618,38 +629,33 @@ define_extension_target(
 # Setting this variable sidesteps the issue by calling the driver directly.
 target_compile_definitions(_C PRIVATE CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
 
-if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
+# add OR VLLM_GPU_LANG STREQUAL "HIP" here once
+# https://github.com/vllm-project/vllm/issues/35163 is resolved
+if(VLLM_GPU_LANG STREQUAL "CUDA")
   #
   # _C_stable_libtorch extension (ops registered via STABLE_TORCH_LIBRARY)
   #
   set(VLLM_STABLE_EXT_SRC
     "csrc/libtorch_stable/torch_bindings.cpp"
-    "csrc/libtorch_stable/activation_kernels.cu"
-    "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
-    "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
-    "csrc/libtorch_stable/quantization/gptq/q_gemm.cu"
-    "csrc/libtorch_stable/quantization/gguf/gguf_kernel.cu"
-    "csrc/libtorch_stable/pos_encoding_kernels.cu"
-    "csrc/libtorch_stable/fused_qknorm_rope_kernel.cu"
-    "csrc/libtorch_stable/layernorm_kernels.cu"
-    "csrc/libtorch_stable/layernorm_quant_kernels.cu"
-    "csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu")
+    "csrc/cutlass_extensions/common.cpp"
+    "csrc/cuda_utils_kernels.cu"
+    "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
+    "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
+    "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_entry.cu")
 
   if(VLLM_GPU_LANG STREQUAL "CUDA")
     list(APPEND VLLM_STABLE_EXT_SRC
-      "csrc/cuda_utils_kernels.cu"
-      "csrc/cutlass_extensions/common.cpp"
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_entry.cu"
       "csrc/libtorch_stable/permute_cols.cu"
       "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
       "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
       "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu")
+  endif()
 
+  if(VLLM_GPU_LANG STREQUAL "CUDA")
     set_gencode_flags_for_srcs(
       SRCS "${VLLM_STABLE_EXT_SRC}"
       CUDA_ARCHS "${CUDA_ARCHS}")
+  endif()
 
   # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
@@ -1029,9 +1035,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     message(STATUS "Building hadacore")
   endif()
 
-  # if CUDA endif
-  endif()
-
   message(STATUS "Enabling C_stable extension.")
   define_extension_target(
     _C_stable_libtorch
@@ -1051,34 +1054,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   target_compile_definitions(_C_stable_libtorch PRIVATE
     TORCH_TARGET_VERSION=0x020A000000000000ULL)
 
-  # Needed to use cuda/hip APIs from C-shim
-  if(VLLM_GPU_LANG STREQUAL "CUDA")
-    target_compile_definitions(_C_stable_libtorch PRIVATE USE_CUDA)
-    # Needed by CUTLASS kernels
-    target_compile_definitions(_C_stable_libtorch PRIVATE
-      CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
-  elseif(VLLM_GPU_LANG STREQUAL "HIP")
-    target_compile_definitions(_C_stable_libtorch PRIVATE USE_ROCM)
-  endif()
+  # Needed to use cuda APIs from C-shim
+  target_compile_definitions(_C_stable_libtorch PRIVATE
+    USE_CUDA)
 
-  # On ROCm, _C_stable_libtorch calls raw HIP APIs (e.g. hipGetDevice in
-  # get_device_prop()) which must resolve to the same libamdhip64.so that
-  # PyTorch uses.  When PyTorch bundles its own copy (pip/conda wheels),
-  # the raw HIP calls would otherwise resolve to the system ROCm copy,
-  # initializing a second HIP runtime that corrupts device state (wrong
-  # device on DeviceGuard, core dumps on multi-GPU tests).
-  #
-  # If PyTorch doesn't bundle libamdhip64 (built from source against system
-  # ROCm), there is only one copy in the process and no action is needed —
-  # the HIP compiler already links the system libamdhip64 automatically.
-  if(VLLM_GPU_LANG STREQUAL "HIP")
-    find_library(_STABLE_TORCH_AMDHIP64 amdhip64
-      PATHS "${TORCH_INSTALL_PREFIX}/lib" NO_DEFAULT_PATH)
-    if(_STABLE_TORCH_AMDHIP64)
-      message(STATUS "Found PyTorch-bundled libamdhip64 at ${_STABLE_TORCH_AMDHIP64}")
-      target_link_libraries(_C_stable_libtorch PRIVATE ${_STABLE_TORCH_AMDHIP64})
-    endif()
-  endif()
+  # Needed by CUTLASS kernels
+  target_compile_definitions(_C_stable_libtorch PRIVATE
+    CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
 endif()
 
 #
@@ -1240,6 +1222,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
       CUDA_ARCHS "${DSV3_ROUTER_GEMM_ARCHS}")
     list(APPEND VLLM_MOE_EXT_SRC "${DSV3_ROUTER_GEMM_SRC}")
     message(STATUS "Building DSV3 router GEMM kernel for archs: ${DSV3_ROUTER_GEMM_ARCHS}")
+
+    # DeepSeek V4 fused RMSNorm + router GEMV - same arch gating as DSV3.
+    set(DSV4_NORM_ROUTER_GEMM_SRC
+      "csrc/moe/dsv4_norm_router_gemm_entry.cu"
+      "csrc/moe/dsv4_norm_router_gemm_kernel.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${DSV4_NORM_ROUTER_GEMM_SRC}"
+      CUDA_ARCHS "${DSV3_ROUTER_GEMM_ARCHS}")
+    list(APPEND VLLM_MOE_EXT_SRC "${DSV4_NORM_ROUTER_GEMM_SRC}")
+    message(STATUS "Building DSV4 norm+router GEMV kernel for archs: ${DSV3_ROUTER_GEMM_ARCHS}")
   else()
     message(STATUS "Not building DSV3 router GEMM kernel as no compatible archs found"
                    " (requires SM90+ and CUDA >= 12.0)")
@@ -1279,12 +1271,6 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
     WITH_SOABI)
 endif()
 
-# Must run after the last HIP `define_extension_target` so every extension
-# has registered its sources.
-if (VLLM_GPU_LANG STREQUAL "HIP")
-  vllm_finalize_hipify_target()
-endif()
-
 # For CUDA and HIP builds also build the triton_kernels external package.
 if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     include(cmake/external_projects/triton_kernels.cmake)

benchmarks/benchmark_hidden_state_extraction.py

@@ -1,415 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Benchmark hidden state extraction throughput.
-
-Measures two modes:
-  1. Baseline: bulk inference with max_tokens=1, no extraction.
-  2. Extract:  async hidden state extraction via ExampleHiddenStatesConnector
-               with N concurrent clients, each consuming hidden states as
-               soon as their request finishes (overlapping I/O with generation).
-
-Reports tokens/s and prompts/s for each mode.
-
-Usage:
-  python benchmarks/benchmark_hidden_state_extraction.py \
-      --model Qwen/Qwen3-0.6B \
-      --num-prompts 64 \
-      --num-clients 8 \
-      --prompt-len 8192 \
-      --layers 1 2 3 4
-"""
-
-import argparse
-import asyncio
-import time
-from concurrent.futures import ThreadPoolExecutor
-
-import torch
-from transformers import AutoConfig
-
-from vllm import LLM, SamplingParams
-from vllm.config.kv_transfer import KVTransferConfig
-from vllm.distributed.kv_transfer.kv_connector.v1 import (
-    example_hidden_states_connector,
-)
-from vllm.engine.arg_utils import AsyncEngineArgs
-from vllm.sampling_params import RequestOutputKind
-from vllm.v1.engine.async_llm import AsyncLLM
-
-
-def _make_profiler_config(profile_dir: str) -> dict:
-    """Build a profiler_config dict for torch profiling."""
-    return {
-        "profiler": "torch",
-        "torch_profiler_dir": profile_dir,
-        "torch_profiler_with_stack": True,
-    }
-
-
-def make_random_prompts(
-    num_prompts: int, prompt_len: int, vocab_size: int, seed: int = 42
-) -> list[list[int]]:
-    """Generate lists of random token IDs."""
-    # Set seed for reproducibility
-    torch.manual_seed(seed)
-    return [
-        torch.randint(0, vocab_size, (prompt_len,)).tolist() for _ in range(num_prompts)
-    ]
-
-
-def consume_hidden_states(path: str) -> float:
-    """Load hidden states from disk and compute per-position mean.
-
-    Returns a single float: the grand mean of all hidden state values.
-    This forces the benchmark to actually read and reduce the data.
-
-    Uses :func:`load_hidden_states` which acquires a shared flock,
-    blocking (without polling) until the async writer releases its
-    exclusive lock.
-    """
-    obj = example_hidden_states_connector.load_hidden_states(path)
-    hs = obj["hidden_states"]
-    total = hs.mean().item()
-
-    example_hidden_states_connector.cleanup_hidden_states(path)
-
-    return total
-
-
-def run_baseline(
-    model: str,
-    prompts: list[list[int]],
-    extra_args: dict,
-    profile_dir: str | None = None,
-) -> dict:
-    """Baseline: bulk inference, no hidden state extraction."""
-    if profile_dir:
-        extra_args = {
-            **extra_args,
-            "profiler_config": _make_profiler_config(profile_dir),
-        }
-    llm = LLM(
-        model=model,
-        enable_prefix_caching=False,
-        enable_chunked_prefill=False,
-        **extra_args,
-    )
-    sampling_params = SamplingParams(max_tokens=1)
-    prompt_inputs = [{"prompt_token_ids": p} for p in prompts]
-
-    # Warmup
-    llm.generate(prompt_inputs[:4], sampling_params, use_tqdm=False)
-
-    if profile_dir:
-        llm.start_profile()
-
-    t0 = time.perf_counter()
-    outputs = llm.generate(prompt_inputs, sampling_params, use_tqdm=True)
-    elapsed = time.perf_counter() - t0
-
-    if profile_dir:
-        llm.stop_profile()
-
-    total_prompt_tokens = sum(len(o.prompt_token_ids) for o in outputs)
-    num_prompts = len(outputs)
-
-    del llm
-    torch.accelerator.empty_cache()
-
-    return {
-        "mode": "baseline",
-        "elapsed_s": elapsed,
-        "num_prompts": num_prompts,
-        "total_prompt_tokens": total_prompt_tokens,
-        "tokens_per_s": total_prompt_tokens / elapsed,
-        "prompts_per_s": num_prompts / elapsed,
-    }
-
-
-# ---- Async extraction benchmark ----
-
-
-async def _client_loop(
-    engine: AsyncLLM,
-    prompt_queue: asyncio.Queue,
-    consume_pool: ThreadPoolExecutor,
-    results: list[dict],
-    client_id: int,
-):
-    """A single async client: pulls prompts, submits to engine, consumes
-    hidden states as soon as each request finishes."""
-    loop = asyncio.get_event_loop()
-    while True:
-        item = await prompt_queue.get()
-        if item is None:
-            prompt_queue.task_done()
-            break
-        idx, token_ids = item
-
-        request_id = f"req-{idx}"
-        sampling_params = SamplingParams(
-            max_tokens=1,
-            output_kind=RequestOutputKind.FINAL_ONLY,
-        )
-
-        final_output = None
-        async for output in engine.generate(
-            request_id=request_id,
-            prompt={"prompt_token_ids": token_ids},
-            sampling_params=sampling_params,
-        ):
-            if output.finished:
-                final_output = output
-
-        # Consume hidden states on a thread (disk I/O)
-        path = final_output.kv_transfer_params["hidden_states_path"]
-        mean_val = await loop.run_in_executor(consume_pool, consume_hidden_states, path)
-        num_tokens = len(final_output.prompt_token_ids)
-
-        results.append(
-            {
-                "request_id": request_id,
-                "num_prompt_tokens": num_tokens,
-                "mean_hidden_value": mean_val,
-            }
-        )
-        prompt_queue.task_done()
-
-
-async def _run_extraction_async(
-    model: str,
-    prompts: list[list[int]],
-    num_clients: int,
-    layers: list[int],
-    tmpdir: str,
-    extra_args: dict,
-    profile_dir: str | None = None,
-) -> dict:
-    if profile_dir:
-        extra_args = {
-            **extra_args,
-            "profiler_config": _make_profiler_config(profile_dir),
-        }
-    engine_args = AsyncEngineArgs(
-        model=model,
-        enable_prefix_caching=False,
-        enable_chunked_prefill=False,
-        max_num_batched_tokens=40960,
-        max_model_len=40960,
-        speculative_config={
-            "method": "extract_hidden_states",
-            "num_speculative_tokens": 1,
-            "draft_model_config": {
-                "hf_config": {
-                    "eagle_aux_hidden_state_layer_ids": layers,
-                },
-            },
-        },
-        kv_transfer_config=KVTransferConfig(
-            kv_connector="ExampleHiddenStatesConnector",
-            kv_role="kv_producer",
-            kv_connector_extra_config={
-                "shared_storage_path": tmpdir,
-            },
-        ),
-        **extra_args,
-    )
-    engine = AsyncLLM.from_engine_args(engine_args)
-
-    try:
-        # Warmup: run a few prompts sequentially, cleaning up generated files
-        for i in range(min(4, len(prompts))):
-            sp = SamplingParams(max_tokens=1, output_kind=RequestOutputKind.FINAL_ONLY)
-            final_output = None
-            async for output in engine.generate(
-                request_id=f"warmup-{i}",
-                prompt={"prompt_token_ids": prompts[i]},
-                sampling_params=sp,
-            ):
-                if output.finished:
-                    final_output = output
-            if final_output and final_output.kv_transfer_params:
-                path = final_output.kv_transfer_params.get("hidden_states_path")
-                if path:
-                    example_hidden_states_connector.cleanup_hidden_states(path)
-
-        if profile_dir:
-            await engine.start_profile()
-
-        # Fill prompt queue
-        prompt_queue: asyncio.Queue = asyncio.Queue()
-        for idx, token_ids in enumerate(prompts):
-            prompt_queue.put_nowait((idx, token_ids))
-        # Sentinel per client
-        for _ in range(num_clients):
-            prompt_queue.put_nowait(None)
-
-        results: list[dict] = []
-        consume_pool = ThreadPoolExecutor(max_workers=num_clients)
-
-        t0 = time.perf_counter()
-        tasks = [
-            asyncio.create_task(
-                _client_loop(engine, prompt_queue, consume_pool, results, i)
-            )
-            for i in range(num_clients)
-        ]
-        await asyncio.gather(*tasks)
-        elapsed = time.perf_counter() - t0
-
-        consume_pool.shutdown(wait=True)
-
-        if profile_dir:
-            await engine.stop_profile()
-
-        total_prompt_tokens = sum(r["num_prompt_tokens"] for r in results)
-        num_prompts = len(results)
-        mean_hidden = sum(r["mean_hidden_value"] for r in results) / max(
-            len(results), 1
-        )
-
-        return {
-            "mode": "extract",
-            "elapsed_s": elapsed,
-            "num_prompts": num_prompts,
-            "total_prompt_tokens": total_prompt_tokens,
-            "tokens_per_s": total_prompt_tokens / elapsed,
-            "prompts_per_s": num_prompts / elapsed,
-            "mean_hidden_value": mean_hidden,
-        }
-    finally:
-        engine.shutdown()
-
-
-def run_extraction(
-    model: str,
-    prompts: list[list[int]],
-    num_clients: int,
-    layers: list[int],
-    extra_args: dict,
-    profile_dir: str | None = None,
-) -> dict:
-    return asyncio.run(
-        _run_extraction_async(
-            model,
-            prompts,
-            num_clients,
-            layers,
-            "/dev/shm",
-            extra_args,
-            profile_dir=profile_dir,
-        )
-    )
-
-
-def print_results(results: dict):
-    mode = results["mode"]
-    print(f"\n{'=' * 60}")
-    print(f"  {mode.upper()} RESULTS")
-    print(f"{'=' * 60}")
-    print(f"  Prompts:             {results['num_prompts']}")
-    print(f"  Total prompt tokens: {results['total_prompt_tokens']:,}")
-    print(f"  Wall time:           {results['elapsed_s']:.2f}s")
-    print(f"  Tokens/s:            {results['tokens_per_s']:,.0f}")
-    print(f"  Prompts/s:           {results['prompts_per_s']:.2f}")
-    if mode == "extract":
-        print(f"  Mean hidden value:   {results['mean_hidden_value']:.6f}")
-    print(f"{'=' * 60}\n")
-
-
-def main():
-    parser = argparse.ArgumentParser(
-        description="Benchmark hidden state extraction throughput"
-    )
-    parser.add_argument("--model", type=str, required=True)
-    parser.add_argument("--num-prompts", type=int, default=64)
-    parser.add_argument("--num-clients", type=int, default=8)
-    parser.add_argument("--prompt-len", type=int, default=8192)
-    parser.add_argument("--layers", type=int, nargs="+", default=[1, 2, 3, 4])
-    parser.add_argument("--skip-baseline", action="store_true")
-    parser.add_argument("--skip-extract", action="store_true")
-    parser.add_argument("--gpu-memory-utilization", type=float, default=0.9)
-    parser.add_argument("--max-num-batched-tokens", type=int, default=None)
-    parser.add_argument("--max-cudagraph-capture-size", type=int, default=None)
-    parser.add_argument("--max-model-len", type=int, default=None)
-    parser.add_argument("--enforce-eager", action="store_true")
-    parser.add_argument("--load-format", type=str, default=None)
-    parser.add_argument(
-        "--profile",
-        action="store_true",
-        help="Enable torch profiler for both baseline and extraction runs.",
-    )
-    parser.add_argument(
-        "--torch-profiler-dir",
-        type=str,
-        default="./vllm_profile",
-        help="Directory to save torch profiler traces (default: ./vllm_profile).",
-    )
-    parser.add_argument(
-        "--enable-flashinfer-autotune",
-        action="store_true",
-        default=False,
-        help="Enable FlashInfer autotuning (can be slow).",
-    )
-    args = parser.parse_args()
-
-    extra_args = {
-        "gpu_memory_utilization": args.gpu_memory_utilization,
-    }
-    if args.max_model_len is not None:
-        extra_args["max_model_len"] = args.max_model_len
-    if args.max_num_batched_tokens is not None:
-        extra_args["max_num_batched_tokens"] = args.max_num_batched_tokens
-        if args.max_model_len and args.max_num_batched_tokens < args.max_model_len:
-            raise ValueError(
-                "max_num_batched_tokens must be >= max_model_len since chunked prefill"
-                " is not supported by hidden state extraction."
-            )
-    if args.enforce_eager:
-        extra_args["enforce_eager"] = True
-    if args.load_format is not None:
-        extra_args["load_format"] = args.load_format
-    if args.max_cudagraph_capture_size is not None:
-        extra_args["max_cudagraph_capture_size"] = args.max_cudagraph_capture_size
-    extra_args["enable_flashinfer_autotune"] = args.enable_flashinfer_autotune
-
-    # Get vocab size from HF config without loading the full model
-    hf_config = AutoConfig.from_pretrained(args.model, trust_remote_code=True)
-    vocab_size = hf_config.vocab_size
-    prompts = make_random_prompts(args.num_prompts, args.prompt_len, vocab_size)
-    print(
-        f"Generated {args.num_prompts} prompts, "
-        f"{args.prompt_len} tokens each (vocab {vocab_size})"
-    )
-
-    profile_dir = args.torch_profiler_dir if args.profile else None
-    if profile_dir:
-        print(f"Torch profiler enabled, traces will be saved to {profile_dir}/")
-
-    if not args.skip_baseline:
-        baseline_profile_dir = f"{profile_dir}/baseline" if profile_dir else None
-        baseline = run_baseline(
-            args.model, prompts, extra_args, profile_dir=baseline_profile_dir
-        )
-        print_results(baseline)
-
-    if not args.skip_extract:
-        extract_profile_dir = f"{profile_dir}/extract" if profile_dir else None
-        extract = run_extraction(
-            args.model,
-            prompts,
-            args.num_clients,
-            args.layers,
-            extra_args,
-            profile_dir=extract_profile_dir,
-        )
-        print_results(extract)
-
-    if not args.skip_baseline and not args.skip_extract:
-        slowdown = baseline["tokens_per_s"] / extract["tokens_per_s"]
-        print("Extraction slowdown factor: {:.2f}x".format(slowdown))
-
-
-if __name__ == "__main__":
-    main()

benchmarks/benchmark_serving_structured_output.py

@@ -115,39 +115,6 @@ class SampleRequest:
 def sample_requests(
     tokenizer: PreTrainedTokenizerBase, args: argparse.Namespace
 ) -> list[SampleRequest]:
-    def _apply_random_prefix(
-        tokenizer: PreTrainedTokenizerBase,
-        requests: list[SampleRequest],
-        prefix_len: int,
-        seed: int,
-    ) -> list[SampleRequest]:
-        if prefix_len <= 0:
-            return requests
-        rng = np.random.default_rng(seed)
-        vocab_size = tokenizer.vocab_size
-        prohibited = getattr(tokenizer, "all_special_ids", None) or []
-        allowed = np.array([i for i in range(vocab_size) if i not in prohibited])
-        if len(allowed) == 0:
-            return requests
-        prefix_ids = rng.integers(0, len(allowed), size=prefix_len)
-        prefix_token_ids = allowed[prefix_ids].tolist()
-        out = []
-        for req in requests:
-            prompt_ids = tokenizer(req.prompt, add_special_tokens=False).input_ids
-            full_ids = prefix_token_ids + prompt_ids
-            full_prompt = tokenizer.decode(full_ids, skip_special_tokens=False)
-            out.append(
-                SampleRequest(
-                    prompt=full_prompt,
-                    prompt_len=len(tokenizer(full_prompt).input_ids),
-                    expected_output_len=req.expected_output_len,
-                    schema=req.schema,
-                    structure_type=req.structure_type,
-                    completion=req.completion,
-                )
-            )
-        return out
-
     if args.dataset == "json" or args.dataset == "json-unique":
         if args.json_schema_path is None:
             dir_path = os.path.dirname(os.path.realpath(__file__))
@@ -294,9 +261,6 @@ def sample_requests(
                 )
             )
 
-    requests = _apply_random_prefix(
-        tokenizer, requests, args.random_prefix_len, args.seed
-    )
     return requests
 
 
@@ -981,15 +945,6 @@ def create_argument_parser():
         "results in a more uniform arrival of requests.",
     )
     parser.add_argument("--seed", type=int, default=0)
-    parser.add_argument(
-        "--random-prefix-len",
-        type=int,
-        default=0,
-        help=(
-            "Number of prefix tokens to prepend to every prompt. "
-            "The same prefix is used for all prompts to enable prefix caching."
-        ),
-    )
     parser.add_argument(
         "--trust-remote-code",
         action="store_true",

benchmarks/kernels/benchmark_fused_moe_lora_one_shot.py

@@ -1,465 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Benchmark the fused MoE-LoRA fast path (one-shot) vs two-kernel baseline.
-
-The "one_shot" provider goes through `vllm.lora.ops.triton_ops.fused_moe_lora`
-which dispatches to the single-kernel one-shot implementation when
-fully_sharded=False (the prefill default).
-
-The "two_kernel" provider drives `fused_moe_lora_shrink` + `fused_moe_lora_expand`
-directly, bypassing the dispatch and matching the legacy two-kernel path's
-work distribution. This isolates the win from kernel fusion.
-
-Run:
-    .venv/bin/python -m benchmarks.kernels.benchmark_fused_moe_lora_one_shot
-    .venv/bin/python -m benchmarks.kernels.benchmark_fused_moe_lora_one_shot \\
-        --model qwen3moe
-"""
-
-from __future__ import annotations
-
-import argparse
-import os
-import random
-
-import torch
-
-from vllm import _custom_ops as ops
-from vllm.lora.ops.triton_ops import (
-    fused_moe_lora,
-    fused_moe_lora_expand,
-    fused_moe_lora_shrink,
-)
-from vllm.triton_utils import triton
-
-DTYPE = torch.bfloat16
-DEVICE = "cuda"
-
-
-# ----- input fabrication -----------------------------------------------------
-
-
-def _round_up(x: int, base: int) -> int:
-    return ((x + base - 1) // base) * base
-
-
-def _ceildiv(x: int, y: int) -> int:
-    return (x + y - 1) // y
-
-
-def _assign_loras(num_tokens: int, num_sequences: int, max_loras: int) -> torch.Tensor:
-    tokens_per_seq = num_tokens // num_sequences
-    rem = num_tokens % num_sequences
-    out = torch.empty(num_tokens, dtype=torch.int32)
-    start = 0
-    for i in range(num_sequences):
-        end = start + tokens_per_seq + (1 if i < rem else 0)
-        out[start:end] = random.randint(0, max_loras - 1)
-        start = end
-    return out
-
-
-def _assign_experts(num_tokens: int, num_experts: int, top_k: int):
-    expert_indices = torch.empty((num_tokens, top_k), dtype=torch.int32)
-    for i in range(num_tokens):
-        expert_indices[i] = torch.randperm(num_experts)[:top_k]
-    weights = torch.rand((num_tokens, top_k), dtype=torch.float32)
-    weights = weights / weights.sum(dim=1, keepdim=True)
-    return expert_indices, weights
-
-
-def _make_inputs(
-    M: int,
-    K: int,
-    N_per_slice: int,
-    rank: int,
-    num_experts: int,
-    top_k: int,
-    max_loras: int,
-    num_slices: int,
-    block_size_m: int,
-):
-    """Mirrors the production caller's tensor layout."""
-    torch.manual_seed(0)
-    random.seed(0)
-
-    num_sequences = max(1, min(M, 8))
-    topk_ids_cpu, topk_weights_cpu = _assign_experts(M, num_experts, top_k)
-    token_lora_cpu = _assign_loras(M, num_sequences, max_loras)
-    lora_ids_cpu = torch.full((max_loras + 1,), -1, dtype=torch.int32)
-    uniq = torch.unique(token_lora_cpu, sorted=True)
-    lora_ids_cpu[: uniq.size(0)].copy_(uniq)
-
-    topk_ids = topk_ids_cpu.to(DEVICE)
-    topk_weights = topk_weights_cpu.to(device=DEVICE, dtype=DTYPE)
-    token_lora_mapping = token_lora_cpu.to(DEVICE)
-    lora_ids = lora_ids_cpu.to(DEVICE)
-    adapter_enabled = torch.ones(max_loras + 1, dtype=torch.int32, device=DEVICE)
-
-    lora_a = [
-        torch.randn((max_loras, num_experts, rank, K), dtype=DTYPE, device=DEVICE)
-        / max(K, 1) ** 0.5
-        for _ in range(num_slices)
-    ]
-    lora_b = [
-        torch.randn(
-            (max_loras, num_experts, N_per_slice, rank),
-            dtype=DTYPE,
-            device=DEVICE,
-        )
-        / max(rank, 1) ** 0.5
-        for _ in range(num_slices)
-    ]
-    hidden = torch.randn((M, K), dtype=DTYPE, device=DEVICE)
-    out_template = torch.zeros(
-        (M, top_k, num_slices * N_per_slice), dtype=DTYPE, device=DEVICE
-    )
-
-    # Sorted-path metadata (the prefill default).
-    max_pad = topk_ids.numel() + num_experts * (block_size_m - 1)
-    max_pad = _round_up(max_pad, block_size_m)
-    max_blocks = _ceildiv(max_pad, block_size_m)
-    sorted_token_ids = torch.empty(
-        (max_loras * max_pad,), dtype=torch.int32, device=DEVICE
-    )
-    expert_ids = torch.empty(
-        (max_loras * max_blocks,), dtype=torch.int32, device=DEVICE
-    )
-    num_post = torch.empty((max_loras,), dtype=torch.int32, device=DEVICE)
-    ops.moe_lora_align_block_size(
-        topk_ids,
-        token_lora_mapping,
-        num_experts,
-        block_size_m,
-        max_loras,
-        max_pad,
-        max_blocks,
-        sorted_token_ids,
-        expert_ids,
-        num_post,
-        adapter_enabled,
-        lora_ids,
-    )
-    expert_ids = expert_ids.view(max_loras, -1).contiguous()
-    sorted_token_ids = sorted_token_ids.view(max_loras, -1).contiguous()
-    num_active = torch.tensor([max_loras + 1], dtype=torch.int32, device="cpu")
-
-    return dict(
-        hidden=hidden,
-        lora_a=lora_a,
-        lora_b=lora_b,
-        topk_weights=topk_weights,
-        sorted_token_ids=sorted_token_ids,
-        expert_ids=expert_ids,
-        num_post=num_post,
-        token_lora_mapping=token_lora_mapping,
-        lora_ids=lora_ids,
-        num_active=num_active,
-        adapter_enabled=adapter_enabled,
-        out_template=out_template,
-        # bookkeeping
-        M=M,
-        K=K,
-        N_per_slice=N_per_slice,
-        rank=rank,
-        num_experts=num_experts,
-        top_k=top_k,
-        max_loras=max_loras,
-        num_slices=num_slices,
-        block_size_m=block_size_m,
-    )
-
-
-# ----- providers -------------------------------------------------------------
-
-
-def _run_one_shot(inp: dict):
-    """Drive `fused_moe_lora` with fully_sharded=False -> one-shot fast path."""
-    out = inp["out_template"].clone()
-    fused_moe_lora(
-        out,
-        inp["hidden"],
-        inp["lora_a"],
-        inp["lora_b"],
-        inp["topk_weights"],
-        inp["sorted_token_ids"],
-        inp["expert_ids"],
-        inp["num_post"],
-        inp["token_lora_mapping"],
-        inp["rank"],
-        inp["top_k"],
-        inp["lora_ids"],
-        inp["num_active"],
-        inp["adapter_enabled"],
-        inp["block_size_m"],
-        64,
-        32,
-        8,
-        4,
-        3,
-        1,
-        inp["block_size_m"],
-        64,
-        32,
-        8,
-        4,
-        3,
-        1,
-        False,
-        False,
-        0,
-    )
-    return out
-
-
-def _run_two_kernel(inp: dict):
-    """Drive `fused_moe_lora_shrink` + `fused_moe_lora_expand` directly,
-    bypassing the dispatch. Matches the legacy two-kernel work distribution.
-    """
-    M = inp["M"]
-    top_k = inp["top_k"]
-    rank = inp["rank"]
-    num_slices = inp["num_slices"]
-    N_per_slice = inp["N_per_slice"]
-    K = inp["K"]
-    num_experts = inp["num_experts"]
-    block_m = inp["block_size_m"]
-
-    intermediate = torch.zeros((num_slices, M, top_k, rank), dtype=DTYPE, device=DEVICE)
-    out = inp["out_template"].clone()
-    EM = inp["sorted_token_ids"].shape[1]
-    num_tokens = M * top_k
-
-    fused_moe_lora_shrink(
-        intermediate,
-        inp["hidden"],
-        inp["lora_a"],
-        inp["topk_weights"],
-        inp["sorted_token_ids"],
-        inp["expert_ids"],
-        inp["num_post"],
-        inp["token_lora_mapping"],
-        top_k,
-        inp["lora_ids"],
-        inp["adapter_enabled"],
-        torch.device(DEVICE),
-        rank,
-        M,
-        EM,
-        K,
-        num_tokens,
-        num_experts,
-        num_slices,
-        block_m,
-        64,
-        32,
-        8,
-        4,
-        3,
-        1,
-        inp["num_active"],
-        False,
-    )
-    fused_moe_lora_expand(
-        out,
-        intermediate,
-        inp["lora_b"],
-        inp["topk_weights"],
-        inp["sorted_token_ids"],
-        inp["expert_ids"],
-        inp["num_post"],
-        inp["token_lora_mapping"],
-        top_k,
-        inp["lora_ids"],
-        inp["adapter_enabled"],
-        torch.device(DEVICE),
-        rank,
-        M,
-        EM,
-        K,
-        num_tokens,
-        num_experts,
-        num_slices,
-        rank,
-        N_per_slice,
-        block_m,
-        64,
-        32,
-        8,
-        4,
-        3,
-        1,
-        inp["num_active"],
-        False,
-        0,
-    )
-    return out
-
-
-PROVIDER_FNS = {
-    "one_shot": _run_one_shot,
-    "two_kernel": _run_two_kernel,
-}
-
-
-# ----- model presets ---------------------------------------------------------
-
-
-MODEL_PRESETS: dict[str, dict] = {
-    # Mixtral-8x7B style: E=8, top_k=2, hidden=4096, intermediate=14336
-    "mixtral": dict(
-        K=4096,
-        N_per_slice=7168,
-        num_experts=8,
-        top_k=2,
-        max_loras=4,
-        num_slices=2,
-        block_size_m=64,
-    ),
-    # Qwen3-MoE / DeepSeek-V2 style: E=64, top_k=8, hidden=2048, inter=1408
-    "qwen3moe": dict(
-        K=2048,
-        N_per_slice=1408,
-        num_experts=64,
-        top_k=8,
-        max_loras=4,
-        num_slices=2,
-        block_size_m=64,
-    ),
-    # GLM-5.1 (zai-org/GLM-5.1-FP8): E=256, top_k=8, hidden=6144,
-    # moe_intermediate=2048
-    "glm5_1": dict(
-        K=6144,
-        N_per_slice=2048,
-        num_experts=256,
-        top_k=8,
-        max_loras=4,
-        num_slices=2,
-        block_size_m=64,
-    ),
-}
-
-
-M_RANGE = [16, 64, 256, 1024, 4096, 16384]
-RANK_RANGE = [8, 16, 32, 64]
-
-
-def get_benchmark(model: str, max_loras: int | None = None):
-    preset = dict(MODEL_PRESETS[model])
-    if max_loras is not None:
-        preset["max_loras"] = max_loras
-
-    @triton.testing.perf_report(
-        triton.testing.Benchmark(
-            x_names=["M", "rank"],
-            x_vals=[(M, R) for M in M_RANGE for R in RANK_RANGE],
-            line_arg="provider",
-            line_vals=list(PROVIDER_FNS.keys()),
-            line_names=["one_shot (fused)", "two_kernel (legacy)"],
-            styles=[("red", "-"), ("blue", "-")],
-            ylabel="ms",
-            plot_name=f"fused_moe_lora-{model}-loras{preset['max_loras']}",
-            args={"preset": preset},
-        )
-    )
-    def benchmark(M, rank, provider, preset):
-        inp = _make_inputs(
-            M=M,
-            K=preset["K"],
-            N_per_slice=preset["N_per_slice"],
-            rank=rank,
-            num_experts=preset["num_experts"],
-            top_k=preset["top_k"],
-            max_loras=preset["max_loras"],
-            num_slices=preset["num_slices"],
-            block_size_m=preset["block_size_m"],
-        )
-        fn = PROVIDER_FNS[provider]
-        quantiles = [0.5, 0.2, 0.8]
-        ms, min_ms, max_ms = triton.testing.do_bench(
-            lambda: fn(inp), quantiles=quantiles
-        )
-        return ms, max_ms, min_ms
-
-    return benchmark
-
-
-# ----- correctness sanity ---------------------------------------------------
-
-
-def calculate_diff(model: str, M: int, rank: int, max_loras: int | None = None):
-    preset = dict(MODEL_PRESETS[model])
-    if max_loras is not None:
-        preset["max_loras"] = max_loras
-    inp = _make_inputs(
-        M=M,
-        K=preset["K"],
-        N_per_slice=preset["N_per_slice"],
-        rank=rank,
-        num_experts=preset["num_experts"],
-        top_k=preset["top_k"],
-        max_loras=preset["max_loras"],
-        num_slices=preset["num_slices"],
-        block_size_m=preset["block_size_m"],
-    )
-    out_one = _run_one_shot(inp)
-    out_two = _run_two_kernel(inp)
-    max_abs = (out_one.float() - out_two.float()).abs().max().item()
-    print(
-        f"  model={model:<9} M={M:<6} rank={rank:<3}  "
-        f"max|one_shot - two_kernel|={max_abs:.4g}  "
-        f"ref|max|={out_two.float().abs().max().item():.3g}"
-    )
-    if max_abs <= 5e-2:
-        print("  ✅ outputs match within bf16 tolerance")
-    else:
-        print("  ❌ outputs differ beyond expected bf16 noise")
-
-
-# ----- main ------------------------------------------------------------------
-
-
-if __name__ == "__main__":
-    parser = argparse.ArgumentParser()
-    parser.add_argument(
-        "--model",
-        type=str,
-        default="mixtral",
-        choices=list(MODEL_PRESETS.keys()),
-        help="Model preset to sweep",
-    )
-    parser.add_argument(
-        "--save-path",
-        type=str,
-        default="./configs/fused_moe_lora_one_shot/",
-        help="Directory to save benchmark results",
-    )
-    parser.add_argument(
-        "--check-only",
-        action="store_true",
-        help="Run correctness sanity check only, no perf sweep",
-    )
-    parser.add_argument(
-        "--max-loras",
-        type=int,
-        default=None,
-        help="Override max_loras in the model preset (number of LoRA adapters "
-        "active in the batch). Defaults to the preset's value.",
-    )
-    args = parser.parse_args()
-
-    print(f"Correctness check ({args.model}):")
-    calculate_diff(args.model, M=256, rank=32, max_loras=args.max_loras)
-    if args.check_only:
-        raise SystemExit(0)
-
-    effective_max_loras = (
-        args.max_loras
-        if args.max_loras is not None
-        else MODEL_PRESETS[args.model]["max_loras"]
-    )
-    print(f"\nGPU: {torch.cuda.get_device_name()}")
-    print(f"Model preset: {args.model}  max_loras={effective_max_loras}\n")
-    benchmark = get_benchmark(args.model, max_loras=args.max_loras)
-    os.makedirs(args.save_path, exist_ok=True)
-    benchmark.run(print_data=True, save_path=args.save_path)

benchmarks/kernels/benchmark_norm_router_gemm.py

@@ -0,0 +1,183 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Benchmark and correctness check for ``ops.dsv4_norm_router_gemm``.
+
+Two implementations are compared:
+
+  1. ``unfused``   — ``vllm_ops.rms_norm`` then ``ops.dsv3_router_gemm``,
+                     i.e. the current vLLM hot path (two kernel launches).
+  2. ``fused``     — ``ops.dsv4_norm_router_gemm``, the new single-kernel
+                     fused path.
+
+Both produce ``(normed_x: bf16, router_logits: fp32)``.  The correctness
+check verifies that ``fused`` and ``unfused`` agree to within ~1 bf16
+ULP — that is the precision floor for this op.
+"""
+
+import argparse
+
+import torch
+
+from vllm import _custom_ops as vllm_ops
+from vllm.triton_utils import triton
+
+# The fused dsv4_norm_router_gemm kernel is templated only for DSV4-Pro
+# (hidden_size=7168, num_experts=384).  Other shapes fall back to the
+# unfused path on the Python side (NormGatedLinear), so benchmark only
+# the configuration that the fused kernel actually targets.
+HIDDEN_SIZE = 7168
+NUM_EXPERTS_CHOICES = (384,)
+RMS_EPS = 1e-6
+
+
+def unfused_norm_router_gemm(
+    x: torch.Tensor,
+    norm_weight: torch.Tensor,
+    gate_weight: torch.Tensor,
+    eps: float,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    # Call ``_C::rms_norm`` directly (mirroring ``_dsv4_pro_norm_gate``'s
+    # fallback path) so the benchmarked baseline doesn't inherit any
+    # Python wrapper overhead or risk falling through to the native
+    # eager-primitive ``RMSNorm.forward_native`` path.
+    normed = torch.empty_like(x)
+    torch.ops._C.rms_norm(normed, x, norm_weight, eps)
+    logits = vllm_ops.dsv3_router_gemm(normed, gate_weight, torch.float32)
+    return normed, logits
+
+
+def fused_norm_router_gemm(
+    x: torch.Tensor,
+    norm_weight: torch.Tensor,
+    gate_weight: torch.Tensor,
+    eps: float,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    return vllm_ops.dsv4_norm_router_gemm(x, norm_weight, gate_weight, eps)
+
+
+def _make_inputs(num_tokens: int, num_experts: int, hidden_size: int, seed: int = 0):
+    torch.manual_seed(seed)
+    device = "cuda"
+    x = torch.randn(num_tokens, hidden_size, dtype=torch.bfloat16, device=device)
+    norm_w = torch.randn(hidden_size, dtype=torch.bfloat16, device=device)
+    gate_w = torch.randn(num_experts, hidden_size, dtype=torch.bfloat16, device=device)
+    # Down-scale gate_w so the GEMV output stays in a representable range.
+    gate_w = gate_w / float(hidden_size) ** 0.5
+    norm_w = (norm_w * 0.1) + 1.0
+    return x, norm_w, gate_w
+
+
+def calculate_diff(
+    num_tokens: int,
+    num_experts: int,
+    hidden_size: int = HIDDEN_SIZE,
+    normed_atol: float = 2e-3,
+    logits_atol: float = 1e-2,
+    rtol: float = 1e-2,
+) -> None:
+    x, norm_w, gate_w = _make_inputs(num_tokens, num_experts, hidden_size)
+
+    normed_unfused, logits_unfused = unfused_norm_router_gemm(
+        x.clone(), norm_w, gate_w, RMS_EPS
+    )
+    normed_fused, logits_fused = fused_norm_router_gemm(
+        x.clone(), norm_w, gate_w, RMS_EPS
+    )
+
+    def _max_abs(a, b):
+        return (a.float() - b.float()).abs().max().item()
+
+    print(f"\n=== M={num_tokens} E={num_experts} H={hidden_size} ===")
+    print(f"normed_x  |fused - unfused| = {_max_abs(normed_fused, normed_unfused):.3e}")
+    print(f"logits    |fused - unfused| = {_max_abs(logits_fused, logits_unfused):.3e}")
+
+    ok_normed = torch.allclose(
+        normed_fused.float(),
+        normed_unfused.float(),
+        atol=normed_atol,
+        rtol=rtol,
+    )
+    ok_logits = torch.allclose(
+        logits_fused.float(),
+        logits_unfused.float(),
+        atol=logits_atol,
+        rtol=rtol,
+    )
+    if ok_normed and ok_logits:
+        print(
+            f"OK   fused vs unfused within "
+            f"normed_atol={normed_atol:.0e} logits_atol={logits_atol:.0e} "
+            f"rtol={rtol:.0e}"
+        )
+    else:
+        print(
+            f"FAIL normed_ok={ok_normed} logits_ok={ok_logits}; "
+            f"see max-abs values above"
+        )
+
+
+def get_benchmark():
+    # Only num_tokens varies (DSV4-Pro hard-codes E=384); single-axis
+    # sweep yields a clean line plot with M on the x-axis.
+    num_experts = NUM_EXPERTS_CHOICES[0]
+
+    @triton.testing.perf_report(
+        triton.testing.Benchmark(
+            x_names=["num_tokens"],
+            x_vals=list(range(1, 17)),
+            line_arg="provider",
+            line_vals=["unfused", "fused"],
+            line_names=["unfused (rms+dsv3)", "fused (dsv4)"],
+            styles=[("green", "-"), ("red", "-")],
+            ylabel="us",
+            plot_name=f"norm-router-gemm-E{num_experts}-H{HIDDEN_SIZE}",
+            args={},
+        )
+    )
+    def benchmark(num_tokens, provider):
+        x, norm_w, gate_w = _make_inputs(num_tokens, num_experts, HIDDEN_SIZE)
+
+        quantiles = [0.5, 0.2, 0.8]
+        if provider == "unfused":
+            fn = lambda: unfused_norm_router_gemm(  # noqa: E731
+                x, norm_w, gate_w, RMS_EPS
+            )
+        else:
+            fn = lambda: fused_norm_router_gemm(  # noqa: E731
+                x, norm_w, gate_w, RMS_EPS
+            )
+
+        ms, min_ms, max_ms = triton.testing.do_bench(fn, quantiles=quantiles)
+        return 1000 * ms, 1000 * max_ms, 1000 * min_ms
+
+    return benchmark
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument(
+        "--save-path",
+        type=str,
+        default="./configs/norm_router_gemm/",
+    )
+    parser.add_argument(
+        "--skip-bench",
+        action="store_true",
+        help="Run only the correctness check, not the perf sweep.",
+    )
+    args = parser.parse_args()
+
+    # Correctness sweep over the full fast-path range M=1..16.
+    for m in range(1, 17):
+        for e in NUM_EXPERTS_CHOICES:
+            calculate_diff(num_tokens=m, num_experts=e, hidden_size=HIDDEN_SIZE)
+
+    if args.skip_bench:
+        return
+
+    benchmark = get_benchmark()
+    benchmark.run(print_data=True, save_path=args.save_path)
+
+
+if __name__ == "__main__":
+    main()

benchmarks/kernels/benchmark_selective_state_update.py

@@ -1,774 +0,0 @@
-#!/usr/bin/env python3
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Benchmark and tuning script for the Mamba selective_state_update kernel.
-
-Mirrors the fused MoE tuning workflow: sweeps (BLOCK_SIZE_M, num_warps) across
-an effective_batch grid for a given (headdim, dstate, ngroups, cache_dtype) and
-saves the best config per effective_batch to JSON. Generated configs are picked
-up by selective_state_update at runtime.
-
-Usage:
-    python -m benchmarks.kernels.benchmark_selective_state_update \
-        --all-dstates --save-configs --compare
-"""
-
-import argparse
-import json
-import os
-import sys
-from io import StringIO
-from itertools import product
-from typing import Any
-
-import torch
-
-from tests.kernels.mamba.utils import selective_state_update_ref
-from vllm.model_executor.layers.mamba.ops.mamba_ssm import (
-    _CONFIGS_DIR,
-    _canonical_cache_dtype,
-    _get_default_ssm_launch_config,
-    get_ssm_config_file_name,
-    get_ssm_device_name,
-    override_ssm_config,
-    selective_state_update,
-)
-from vllm.triton_utils import triton
-
-# bf16 shares configs with fp16 - same bit width.
-_SSM_CACHE_DTYPE_MAP: dict[str, torch.dtype] = {
-    "float32": torch.float32,
-    "float16": torch.float16,
-    "bfloat16": torch.float16,
-}
-
-_RESULTS_DIR = os.path.dirname(os.path.realpath(__file__))
-
-# ---------------------------------------------------------------------------
-# Tuning search space
-# ---------------------------------------------------------------------------
-
-_BSM_CHOICES_ALL = [4, 8, 16, 32, 64, 128, 256]
-
-NUM_WARPS_CHOICES = [1, 2, 4, 8]
-
-
-def _block_size_m_choices(headdim: int) -> list[int]:
-    """BLOCK_SIZE_M candidates worth sweeping for a given headdim.
-
-    BLOCK_SIZE_M > next_pow2(headdim) wastes >=50% of each tile via masking
-    (offs_m >= dim rows are zeroed out), so we cap the sweep there.
-    """
-    ceiling = 1
-    while ceiling < headdim:
-        ceiling <<= 1
-    return [b for b in _BSM_CHOICES_ALL if b <= ceiling]
-
-
-# Default deployment shapes. effective_batch = batch * nheads scales the
-# kernel grid, so configs transfer across (model, TP) combos sharing
-# (headdim, dstate, cache_dtype).
-DEFAULT_BATCH_SIZES = [1, 8, 16, 32, 64, 128, 256, 512, 1024, 1536, 2048]
-DEFAULT_NHEADS = [128, 256]
-
-ALL_DSTATES = [16, 32, 64, 128, 256]
-
-# Default tuning shape — matches Nemotron-3-Super and Nemotron-3-Nano Mamba layers.
-# Override with CLI flags for other architectures.
-DEFAULT_HEADDIM = 64
-DEFAULT_NGROUPS = 8
-
-
-# ---------------------------------------------------------------------------
-# Benchmark helper
-# ---------------------------------------------------------------------------
-
-
-def _make_inputs(
-    batch: int,
-    nheads: int,
-    dim: int,
-    dstate: int,
-    ngroups: int,
-    dtype: torch.dtype,
-    state_dtype: torch.dtype | None = None,
-    device: str = "cuda",
-):
-    if state_dtype is None:
-        state_dtype = dtype
-    state = torch.randn(batch, nheads, dim, dstate, dtype=state_dtype, device=device)
-    x = torch.randn(batch, nheads, dim, dtype=dtype, device=device)
-    dt = torch.randn(batch, nheads, dim, dtype=dtype, device=device)
-    A = -torch.rand(nheads, dim, dstate, dtype=torch.float32, device=device)
-    B = torch.randn(batch, ngroups, dstate, dtype=dtype, device=device)
-    C = torch.randn(batch, ngroups, dstate, dtype=dtype, device=device)
-    D = torch.randn(nheads, dim, dtype=dtype, device=device)
-    dt_bias = torch.randn(nheads, dim, dtype=dtype, device=device)
-    out = torch.zeros(batch, nheads, dim, dtype=dtype, device=device)
-    return state, x, dt, A, B, C, D, dt_bias, out
-
-
-def benchmark_config(
-    batch: int,
-    nheads: int,
-    dim: int,
-    dstate: int,
-    ngroups: int,
-    block_size_m: int,
-    num_warps_val: int,
-    dtype: torch.dtype,
-    state_dtype: torch.dtype | None = None,
-    num_iters: int = 100,
-    num_warmup: int = 20,
-    graph_batch_size: int = 10,
-) -> float | None:
-    """
-    Time one (BLOCK_SIZE_M, num_warps) config for selective_state_update.
-    Returns elapsed time in microseconds, or None on error.
-
-    Uses CUDA graph capture-and-replay to isolate kernel time from Python
-    eager-mode dispatch / kwarg-resolution overhead, mirroring the timing
-    methodology in benchmarks/kernels/benchmark_moe.py.
-    """
-    state, x, dt, A, B, C, D, dt_bias, out = _make_inputs(
-        batch, nheads, dim, dstate, ngroups, dtype, state_dtype=state_dtype
-    )
-
-    def _call_kernel() -> None:
-        selective_state_update(
-            state,
-            x,
-            dt,
-            A,
-            B,
-            C,
-            D=D,
-            z=None,
-            dt_bias=dt_bias,
-            dt_softplus=True,
-            out=out,
-        )
-
-    try:
-        with override_ssm_config((block_size_m, num_warps_val)):
-            # Eager-mode warmup: triggers Triton autotune / JIT, primes caches.
-            for _ in range(num_warmup):
-                _call_kernel()
-            torch.accelerator.synchronize()
-
-            # Capture graph_batch_size invocations into a CUDA graph so the
-            # timed region runs without Python dispatch overhead per call.
-            graph = torch.cuda.CUDAGraph()
-            with torch.cuda.graph(graph):
-                for _ in range(graph_batch_size):
-                    _call_kernel()
-            torch.accelerator.synchronize()
-
-            # Warmup graph replays (let the runtime stabilize).
-            for _ in range(5):
-                graph.replay()
-            torch.accelerator.synchronize()
-
-            start = torch.cuda.Event(enable_timing=True)
-            end = torch.cuda.Event(enable_timing=True)
-            latencies: list[float] = []
-            for _ in range(num_iters):
-                start.record()
-                graph.replay()
-                end.record()
-                end.synchronize()
-                latencies.append(start.elapsed_time(end))
-            graph.reset()
-        # elapsed_time returns ms; each replay runs graph_batch_size kernels,
-        # so divide by (num_iters * graph_batch_size) and convert ms -> us.
-        return sum(latencies) / (num_iters * graph_batch_size) * 1000
-    except Exception as e:
-        if "OutOfResources" not in str(e):
-            print(
-                f"    Warning: config M={block_size_m},w={num_warps_val} "
-                f"raised {type(e).__name__}: {e}"
-            )
-        return None
-
-
-# ---------------------------------------------------------------------------
-# Tuning loop
-# ---------------------------------------------------------------------------
-
-
-# CUDA grid Y/Z dim limit — both `batch` and `nheads` must fit individually.
-_CUDA_MAX_GRID_DIM = 65535
-
-# Above this, kernel state-offset arithmetic (batch * nheads * headdim * dstate)
-# overflows int32 and the launch raises cudaErrorIllegalAddress.
-# 262144 covers Nemotron Super TP1 BS=2048.
-_MAX_EFFECTIVE_BATCH = 262144
-
-
-def expand_batch_x_nheads(
-    batch_sizes: list[int],
-    nheads_list: list[int],
-    ngroups: int,
-) -> list[tuple[int, int, int]]:
-    """Cross-product batch_sizes × nheads_list → sorted [(effective_batch,
-    batch, nheads)], deduped by effective_batch. Filters pairs that exceed
-    the CUDA grid dim limit, the effective_batch ceiling, or where nheads is
-    not a positive multiple of ngroups.
-    """
-    seen: dict[int, tuple[int, int]] = {}
-    skipped_grid: list[tuple[int, int]] = []
-    skipped_ngroups: list[tuple[int, int]] = []
-    skipped_eb: list[tuple[int, int]] = []
-    for b, n in product(batch_sizes, nheads_list):
-        if b <= 0 or n <= 0:
-            continue
-        if b > _CUDA_MAX_GRID_DIM or n > _CUDA_MAX_GRID_DIM:
-            skipped_grid.append((b, n))
-            continue
-        if n % ngroups != 0:
-            skipped_ngroups.append((b, n))
-            continue
-        if b * n > _MAX_EFFECTIVE_BATCH:
-            skipped_eb.append((b, n))
-            continue
-        seen.setdefault(b * n, (b, n))
-    if skipped_grid:
-        print(
-            f"  Note: skipping (batch, nheads) pairs exceeding CUDA grid dim "
-            f"{_CUDA_MAX_GRID_DIM}: {skipped_grid}"
-        )
-    if skipped_ngroups:
-        print(
-            f"  Note: skipping (batch, nheads) pairs where nheads % ngroups != 0 "
-            f"for ngroups={ngroups}: {skipped_ngroups}"
-        )
-    if skipped_eb:
-        print(
-            f"  Note: skipping (batch, nheads) pairs whose effective_batch "
-            f"exceeds {_MAX_EFFECTIVE_BATCH}: {skipped_eb}"
-        )
-    return sorted((eb, b, n) for eb, (b, n) in seen.items())
-
-
-def tune_dstate(
-    dstate: int,
-    headdim: int,
-    ngroups: int,
-    dtype: torch.dtype,
-    num_iters: int,
-    verbose: bool,
-    active: list[tuple[int, int, int]],
-    state_dtype: torch.dtype | None = None,
-) -> tuple[dict[int, dict], dict[int, dict[tuple[int, int], float]]]:
-    """For each (effective_batch, batch, nheads) in *active*, sweep
-    (BLOCK_SIZE_M, num_warps) and return
-    ({effective_batch: best_config}, {effective_batch: {(bsm, nw): us}}).
-    The second map is the full timing grid, used downstream so we don't
-    re-measure the same config in the comparison phase.
-    """
-    best_per_eb: dict[int, dict] = {}
-    timings: dict[int, dict[tuple[int, int], float]] = {}
-
-    print(f"\n{'=' * 74}")
-    effective_state_dtype = state_dtype if state_dtype is not None else dtype
-    print(
-        f"Tuning  headdim={headdim}  dstate={dstate}  ngroups={ngroups}  "
-        f"dtype={dtype}  ssm_cache_dtype={effective_state_dtype}"
-    )
-    print(f"{'=' * 74}")
-
-    bsm_choices = _block_size_m_choices(headdim)
-    print(f"BSM candidates (capped at next_pow2(headdim={headdim})): {bsm_choices}")
-
-    hdr = f"{'EffBatch':>8} | {'BLOCK_M':>7} | {'warps':>5} | {'us':>10} | note"
-    print(hdr)
-    print("-" * 52)
-
-    for eb, batch, nheads in active:
-        best_time = float("inf")
-        best_cfg: dict = {}
-        eb_timings: dict[tuple[int, int], float] = {}
-
-        for bsm, nw in product(bsm_choices, NUM_WARPS_CHOICES):
-            t = benchmark_config(
-                batch=batch,
-                nheads=nheads,
-                dim=headdim,
-                dstate=dstate,
-                ngroups=ngroups,
-                block_size_m=bsm,
-                num_warps_val=nw,
-                dtype=dtype,
-                state_dtype=state_dtype,
-                num_iters=num_iters,
-            )
-            if t is None:
-                continue
-            eb_timings[(bsm, nw)] = t
-            is_best = t < best_time
-            if is_best:
-                best_time = t
-                best_cfg = {"BLOCK_SIZE_M": bsm, "num_warps": nw}
-            if verbose:
-                marker = " <-- best" if is_best else ""
-                print(f"{eb:>8} | {bsm:>7} | {nw:>5} | {t:>10.2f} |{marker}")
-
-        timings[eb] = eb_timings
-
-        if not best_cfg:
-            print(
-                f"{eb:>8} | {'-':>7} | {'-':>5} | {'-':>10} | "
-                f"no working config (skipped)"
-            )
-            continue
-
-        if not verbose:
-            print(
-                f"{eb:>8} | {best_cfg['BLOCK_SIZE_M']:>7} | "
-                f"{best_cfg['num_warps']:>5} | {best_time:>10.2f} | best"
-            )
-
-        best_per_eb[eb] = best_cfg
-
-    return best_per_eb, timings
-
-
-# ---------------------------------------------------------------------------
-# Correctness validation
-# ---------------------------------------------------------------------------
-
-
-def validate_configs(
-    dstate: int,
-    headdim: int,
-    ngroups: int,
-    tuned: dict[int, dict],
-    active: list[tuple[int, int, int]],
-    dtype: torch.dtype,
-    atol: float = 1e-2,
-    rtol: float = 1e-2,
-    state_dtype: torch.dtype | None = None,
-) -> dict[int, bool]:
-    """
-    For every (effective_batch, batch, nheads) in *active* that has a tuned
-    config, run the kernel with that config and compare against the reference.
-    Returns {effective_batch: passed}.
-    """
-    # Disable TF32 so the reference's matmul matches the Triton kernel's
-    # fp32 accumulation; otherwise large ebs show bf16 rounding mismatches.
-    torch.set_float32_matmul_precision("highest")
-
-    print(f"\n{'=' * 74}")
-    effective_state_dtype = state_dtype if state_dtype is not None else dtype
-    print(
-        f"Validation  headdim={headdim}  dstate={dstate}  ngroups={ngroups}  "
-        f"dtype={dtype}  ssm_cache_dtype={effective_state_dtype}  atol={atol}"
-    )
-    print(f"{'=' * 74}")
-    print(f"{'EffBatch':>8} | {'MaxAbsErr':>12} | {'Status':>8}")
-    print("-" * 36)
-
-    results: dict[int, bool] = {}
-
-    for eb, batch, nheads in active:
-        cfg = tuned.get(eb)
-        if cfg is None:
-            continue
-        state, x, dt, A, B, C, D, dt_bias, out = _make_inputs(
-            batch=batch,
-            nheads=nheads,
-            dim=headdim,
-            dstate=dstate,
-            ngroups=ngroups,
-            dtype=dtype,
-            state_dtype=state_dtype,
-        )
-        # Clone state before GPU kernel modifies it in-place
-        state_ref = state.clone()
-
-        with override_ssm_config((cfg["BLOCK_SIZE_M"], cfg["num_warps"])):
-            selective_state_update(
-                state,
-                x,
-                dt,
-                A,
-                B,
-                C,
-                D=D,
-                z=None,
-                dt_bias=dt_bias,
-                dt_softplus=True,
-                out=out,
-            )
-        torch.accelerator.synchronize()
-        gpu_out = out.detach().cpu()
-
-        # Reference uses the original (unmodified) state
-        # Upcast to fp32 so the reference sums in fp32 (matches the Triton
-        # kernel); summing in bf16 over `dstate` blows up the error.
-        ref_out = (
-            selective_state_update_ref(
-                state_ref.float(),
-                x.float(),
-                dt.float(),
-                A.float(),
-                B.float(),
-                C.float(),
-                D=D.float(),
-                dt_bias=dt_bias.float(),
-                dt_softplus=True,
-            )
-            .to(out.dtype)
-            .cpu()
-        )
-
-        passed = torch.allclose(gpu_out.float(), ref_out.float(), atol=atol, rtol=rtol)
-        max_err = (gpu_out.float() - ref_out.float()).abs().max().item()
-        status = "PASS" if passed else "FAIL"
-        results[eb] = passed
-        print(f"{eb:>8} | {max_err:>12.6f} | {status:>8}")
-
-    n_pass = sum(results.values())
-    n_total = len(results)
-    print(f"\n  {n_pass}/{n_total} configs passed validation for dstate={dstate}")
-    return results
-
-
-# ---------------------------------------------------------------------------
-# Save configs
-# ---------------------------------------------------------------------------
-
-
-def save_configs(
-    headdim: int,
-    dstate: int,
-    cache_dtype: str,
-    configs: dict[int, dict],
-    save_dir: str | None = None,
-) -> str:
-    # bf16 shares configs with fp16, use common filename for both
-    cache_dtype = _canonical_cache_dtype(cache_dtype)
-
-    base_dir = save_dir if save_dir else _CONFIGS_DIR
-    os.makedirs(base_dir, exist_ok=True)
-    file_path = os.path.join(
-        base_dir,
-        get_ssm_config_file_name(headdim, dstate, cache_dtype, get_ssm_device_name()),
-    )
-    # triton_version is informational only, the loader ignores it
-    payload: dict[str, Any] = {
-        "triton_version": triton.__version__,
-        **{str(k): v for k, v in sorted(configs.items())},
-    }
-    with open(file_path, "w") as f:
-        json.dump(payload, f, indent=4)
-    return file_path
-
-
-# ---------------------------------------------------------------------------
-# Comparison table
-# ---------------------------------------------------------------------------
-
-
-def current_heuristic(dstate: int, is_blackwell: bool = False) -> dict:
-    """Return the current hard-coded BLOCK_SIZE_M / num_warps for dstate."""
-    bsm, nw = _get_default_ssm_launch_config(dstate, is_blackwell)
-    return {"BLOCK_SIZE_M": bsm, "num_warps": nw}
-
-
-def compare_heuristic_vs_tuned(
-    dstate: int,
-    headdim: int,
-    ngroups: int,
-    tuned: dict[int, dict],
-    timings: dict[int, dict[tuple[int, int], float]],
-    active: list[tuple[int, int, int]],
-    dtype: torch.dtype,
-    num_iters: int,
-    is_blackwell: bool,
-    state_dtype: torch.dtype | None = None,
-):
-    heur_cfg = current_heuristic(dstate, is_blackwell)
-    heur_key = (heur_cfg["BLOCK_SIZE_M"], heur_cfg["num_warps"])
-
-    print(f"\n{'=' * 74}")
-    print(
-        f"Comparison  headdim={headdim}  dstate={dstate}  "
-        f"ngroups={ngroups}  —  heuristic vs tuned"
-    )
-    print(
-        f"Heuristic: BLOCK_SIZE_M={heur_cfg['BLOCK_SIZE_M']}, "
-        f"num_warps={heur_cfg['num_warps']}"
-    )
-    print(f"{'=' * 74}")
-    hdr = (
-        f"{'EffBatch':>8} | {'Heur(us)':>10} | {'Tuned(us)':>10} | "
-        f"{'Speedup':>8} | Best config"
-    )
-    print(hdr)
-    print("-" * len(hdr))
-
-    for eb, batch, nheads in active:
-        eb_timings = timings.get(eb, {})
-
-        # Heuristic timing: reuse the tuning measurement if the heuristic
-        # config was in the swept grid; otherwise measure it once.
-        t_h = eb_timings.get(heur_key)
-        if t_h is None:
-            t_h = benchmark_config(
-                batch=batch,
-                nheads=nheads,
-                dim=headdim,
-                dstate=dstate,
-                ngroups=ngroups,
-                block_size_m=heur_cfg["BLOCK_SIZE_M"],
-                num_warps_val=heur_cfg["num_warps"],
-                dtype=dtype,
-                state_dtype=state_dtype,
-                num_iters=num_iters,
-            )
-
-        # `tuned[eb]` may be missing if all configs failed in tune_dstate;
-        # in that case fall back to the heuristic so the table still prints.
-        best = tuned.get(eb) or heur_cfg
-        t_t = eb_timings.get((best["BLOCK_SIZE_M"], best["num_warps"]))
-
-        if t_h is None or t_t is None:
-            print(f"{eb:>8} | {'N/A':>10} | {'N/A':>10} | {'N/A':>8} |")
-            continue
-        speedup = t_h / t_t
-        marker = " <--" if speedup > 1.05 else ""
-        print(
-            f"{eb:>8} | {t_h:>10.2f} | {t_t:>10.2f} | "
-            f"{speedup:>7.2f}x | "
-            f"M={best['BLOCK_SIZE_M']},w={best['num_warps']}{marker}"
-        )
-
-
-# ---------------------------------------------------------------------------
-# CLI
-# ---------------------------------------------------------------------------
-
-
-def save_results(device_name: str, output: str, results_file: str | None = None) -> str:
-    """Save the full benchmark output to a results text file."""
-    if results_file is None:
-        results_file = os.path.join(
-            _RESULTS_DIR, f"ssm_benchmark_results_{device_name}.txt"
-        )
-    with open(results_file, "w") as f:
-        f.write(output)
-    return results_file
-
-
-def main():
-    parser = argparse.ArgumentParser(
-        description="Tune selective_state_update kernel for Mamba SSM"
-    )
-    parser.add_argument(
-        "--dstate",
-        type=int,
-        default=128,
-        help="SSM state size to tune for (default: 128)",
-    )
-    parser.add_argument(
-        "--all-dstates",
-        action="store_true",
-        help="Tune all common dstate values: " + str(ALL_DSTATES),
-    )
-    parser.add_argument(
-        "--dtype",
-        type=str,
-        default="bfloat16",
-        choices=["float16", "bfloat16"],
-        help="Activation / input data type (default: bfloat16)",
-    )
-    parser.add_argument(
-        "--mamba-ssm-cache-dtype",
-        type=str,
-        default="float32",
-        choices=list(_SSM_CACHE_DTYPE_MAP.keys()),
-        help="SSM state cache dtype (default: float32)",
-    )
-    parser.add_argument(
-        "--num-iters",
-        type=int,
-        default=100,
-        help="Number of timing iterations (default: 100)",
-    )
-    parser.add_argument(
-        "--save-configs",
-        action="store_true",
-        help=f"Save best configs to JSON in {_CONFIGS_DIR}",
-    )
-    parser.add_argument(
-        "--compare",
-        action="store_true",
-        help="Show comparison table: heuristic vs tuned",
-    )
-    parser.add_argument(
-        "--verbose",
-        action="store_true",
-        help="Print every (BLOCK_SIZE_M, num_warps) result, not just best",
-    )
-    parser.add_argument(
-        "--results-file",
-        type=str,
-        default=None,
-        help="Path to save the benchmark results text file "
-        "(default: ssm_benchmark_results_<device>.txt alongside this script)",
-    )
-    parser.add_argument(
-        "--save-dir",
-        type=str,
-        default=None,
-        help=f"Directory to save JSON configs (default: {_CONFIGS_DIR})",
-    )
-    parser.add_argument(
-        "--headdim",
-        type=int,
-        default=DEFAULT_HEADDIM,
-        help=f"Per-head feature dim (default: {DEFAULT_HEADDIM})",
-    )
-    parser.add_argument(
-        "--ngroups",
-        type=int,
-        default=DEFAULT_NGROUPS,
-        help=f"Number of B/C groups (default: {DEFAULT_NGROUPS})",
-    )
-    parser.add_argument(
-        "--batch-sizes",
-        type=int,
-        nargs="+",
-        default=DEFAULT_BATCH_SIZES,
-        metavar="B",
-        help=f"Decoder batch sizes to sweep (default: {DEFAULT_BATCH_SIZES})",
-    )
-    parser.add_argument(
-        "--nheads",
-        type=int,
-        nargs="+",
-        default=DEFAULT_NHEADS,
-        metavar="N",
-        help=f"Number of heads per rank to sweep (default: {DEFAULT_NHEADS}). "
-        "effective_batch = batch * nheads; cross-product is deduped by eb.",
-    )
-    parser.add_argument(
-        "--validate",
-        action="store_true",
-        help="After tuning, verify each best config against a CPU reference "
-        "implementation. Configs that fail are flagged in the output.",
-    )
-    parser.add_argument(
-        "--atol",
-        type=float,
-        default=1e-2,
-        help="Absolute tolerance for --validate (default: 1e-2)",
-    )
-    args = parser.parse_args()
-
-    dtype = torch.bfloat16 if args.dtype == "bfloat16" else torch.float16
-    state_dtype = _SSM_CACHE_DTYPE_MAP[args.mamba_ssm_cache_dtype]
-    device_name = get_ssm_device_name()
-    cap = torch.cuda.get_device_capability()
-    is_blackwell = cap[0] >= 10
-
-    # Mirror all output to a results file (like Unix tee).
-    buf = StringIO()
-
-    class _Tee:
-        """Writes to both the original stdout and an in-memory buffer."""
-
-        def write(self, s):
-            buf.write(s)
-            sys.__stdout__.write(s)
-
-        def flush(self):
-            sys.__stdout__.flush()
-
-    sys.stdout = _Tee()  # type: ignore[assignment]
-
-    try:
-        print(f"Device : {device_name}  (sm_{cap[0]}{cap[1]})")
-        print(f"Blackwell: {is_blackwell}")
-        print(f"dtype  : {args.dtype}")
-        print(f"ssm_cache_dtype: {args.mamba_ssm_cache_dtype}")
-        print(f"headdim: {args.headdim}")
-        print(f"ngroups: {args.ngroups}")
-        print(f"triton : {triton.__version__}")
-
-        dstates = ALL_DSTATES if args.all_dstates else [args.dstate]
-        active = expand_batch_x_nheads(args.batch_sizes, args.nheads, args.ngroups)
-
-        for dstate in dstates:
-            tuned, timings = tune_dstate(
-                dstate=dstate,
-                headdim=args.headdim,
-                ngroups=args.ngroups,
-                dtype=dtype,
-                num_iters=args.num_iters,
-                verbose=args.verbose,
-                active=active,
-                state_dtype=state_dtype,
-            )
-
-            if args.compare:
-                compare_heuristic_vs_tuned(
-                    dstate=dstate,
-                    headdim=args.headdim,
-                    ngroups=args.ngroups,
-                    tuned=tuned,
-                    timings=timings,
-                    active=active,
-                    dtype=dtype,
-                    num_iters=args.num_iters,
-                    is_blackwell=is_blackwell,
-                    state_dtype=state_dtype,
-                )
-
-            if args.validate:
-                validity = validate_configs(
-                    dstate=dstate,
-                    headdim=args.headdim,
-                    ngroups=args.ngroups,
-                    tuned=tuned,
-                    active=active,
-                    dtype=dtype,
-                    atol=args.atol,
-                    state_dtype=state_dtype,
-                )
-                # Filter out any configs that failed correctness check
-                failed = [eb for eb, ok in validity.items() if not ok]
-                if failed:
-                    print(
-                        f"\n  WARNING: {len(failed)} config(s) failed validation "
-                        f"for dstate={dstate}: effective_batches {failed}"
-                    )
-                    print("  These will NOT be saved even with --save-configs.")
-                    tuned = {
-                        eb: cfg for eb, cfg in tuned.items() if validity.get(eb, True)
-                    }
-
-            if args.save_configs:
-                path = save_configs(
-                    headdim=args.headdim,
-                    dstate=dstate,
-                    cache_dtype=args.mamba_ssm_cache_dtype,
-                    configs=tuned,
-                    save_dir=args.save_dir,
-                )
-                print(f"\nSaved: {path}")
-            else:
-                print(f"\nBest configs for dstate={dstate}:")
-                for eb, cfg in sorted(tuned.items()):
-                    print(f"  effective_batch={eb:>6}: {cfg}")
-                print("\n(Re-run with --save-configs to persist to JSON)")
-    finally:
-        sys.stdout = sys.__stdout__
-        results_path = save_results(device_name, buf.getvalue(), args.results_file)
-        print(f"\nResults saved to: {results_path}")
-
-
-if __name__ == "__main__":
-    main()

benchmarks/kernels/cpu/benchmark_cpu_attn.py

@@ -12,6 +12,7 @@ from vllm._custom_ops import (
     cpu_attn_get_scheduler_metadata,
     cpu_attn_reshape_and_cache,
 )
+from vllm.platforms import CpuArchEnum, current_platform
 from vllm.utils.argparse_utils import FlexibleArgumentParser
 from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
 from vllm.v1.attention.backends.cpu_attn import CPUAttentionBackend, _get_attn_isa
@@ -21,14 +22,15 @@ def get_attn_isa(
     block_size: int | None = None,
     dtype: torch.dtype | None = None,
 ):
-    # Delegate to _get_attn_isa so the fallback path applies the same arch
-    # gating (e.g. RISC-V RVV is only chosen when the build's hardcoded
-    # VLEN=128 kernel is actually present; on VLEN=256 / scalar hosts it
-    # correctly falls through to vec/vec16).
-    return _get_attn_isa(
-        dtype if dtype is not None else torch.bfloat16,
-        block_size if block_size else 32,
-    )
+    if block_size and dtype:
+        return _get_attn_isa(dtype, block_size)
+    else:
+        if current_platform.get_cpu_architecture() == CpuArchEnum.ARM:
+            return "neon"
+        elif torch.cpu._is_amx_tile_supported():
+            return "amx"
+        else:
+            return "vec"
 
 
 # rand number generation takes too much time, cache rand tensors
@@ -233,7 +235,7 @@ if __name__ == "__main__":
     )
     parser.add_argument("--use-sink", action="store_true")
     parser.add_argument(
-        "--isa", type=str, choices=["vec", "neon", "amx", "vec16", "rvv"], default=None
+        "--isa", type=str, choices=["vec", "neon", "amx", "vec16"], default=None
     )
     parser.add_argument("--seed", type=int, default=0)
     parser.add_argument("--iters", type=int, default=20)

benchmarks/kernels/ir/bench_ir_ops.py

@@ -145,9 +145,6 @@ def _bench_one(fn, args, cfg: BenchConfig) -> float:
     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]]]:
@@ -158,11 +155,15 @@ def collect_timings(
         "_".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]
+    provider_strs = [f"{p}-compiled" if op.impls[p].compiled else p for p in providers]
 
     results: dict[str, dict[str, float]] = {c: {} for c in case_names}
-    for provider in providers:
+    for provider, p_str in zip(providers, provider_strs):
         impl = op.impls[provider]
-        desc = f"{op.name} / {provider}"
+        if impl.compiled:
+            impl = impl.compile()
+
+        desc = f"{op.name} / {p_str}"
         for case_name, kwargs in tqdm(
             zip(case_names, shape_configs),
             desc=desc,
@@ -171,11 +172,11 @@ def collect_timings(
         ):
             args = op.generate_inputs(**kwargs)
             if impl.supports_args(*args):
-                results[case_name][provider] = _bench_one(impl.impl_fn, args, cfg)
+                results[case_name][p_str] = _bench_one(impl.impl_fn, args, cfg)
             else:
-                results[case_name][provider] = float("nan")
+                results[case_name][p_str] = float("nan")
 
-    return case_names, providers, results
+    return case_names, provider_strs, results
 
 
 def analyze_results(
@@ -184,7 +185,7 @@ def analyze_results(
     providers: list[str],
     results: dict[str, dict[str, float]],
 ) -> tuple[list[dict[str, str]], list[dict[str, str]], list[str]]:
-    native_col = "native"
+    native_col = "native-compiled" if "native-compiled" in providers else "native"
     non_native = [p for p in providers if p != native_col]
 
     header_cols = ["case"]
@@ -229,7 +230,9 @@ def analyze_results(
         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"\n{p} vs {native_col} ({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})")
@@ -288,10 +291,10 @@ def save_results(
 def parse_args():
     parser = argparse.ArgumentParser(description="Benchmark vLLM IR ops")
     parser.add_argument(
-        "--ops",
+        "op",
         type=str,
-        default=None,
-        help="Comma-separated list of op names to benchmark (substring match)",
+        nargs="*",
+        help="Op name(s) to benchmark",
     )
     parser.add_argument(
         "--no-cuda-graph",
@@ -339,12 +342,11 @@ def main():
     )
     os.makedirs(save_dir, exist_ok=True)
 
-    op_filters = [f.strip() for f in args.ops.split(",")] if args.ops else None
+    op_names = [f.strip() for f in args.ops] if args.op else IrOp.registry.keys()
     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
+    for op_name in op_names:
+        op = IrOp.registry[op_name]
         if not op.has_input_generator:
             print(f"Skipping op '{op.name}': no input generator registered")
             continue
@@ -354,6 +356,11 @@ def main():
                 f"Add it to benchmarks/kernels/ir/shapes.py"
             )
 
+        # Sort descending, so that torch.compile first sees the largest shape
+        # (simulates vLLM usage)
+        shape_configs = SHAPE_CONFIGS[op.name]
+        shape_configs.sort(key=lambda x: x["num_tokens"], reverse=True)
+
         case_names, providers, results = collect_timings(
             op, SHAPE_CONFIGS[op.name], cfg
         )

benchmarks/kernels/ir/shapes.py

@@ -26,4 +26,10 @@ SHAPE_CONFIGS: dict[str, list[dict]] = {
         for d in COMMON_HIDDEN_SIZES
         for n in NUM_TOKENS
     ],
+    "fused_add_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
+    ],
 }

build_rust.sh

@@ -1,44 +0,0 @@
-#!/bin/bash
-# Build the vllm-rs Rust frontend binary and install it into the vllm package.
-# Usage: ./build_rust.sh [--debug]
-#
-# By default builds in release mode. Pass --debug for faster compile times
-# during development.
-
-set -euo pipefail
-
-REPO_ROOT="$(cd "$(dirname "$0")" && pwd)"
-RUST_DIR="$REPO_ROOT/rust"
-TARGET_PATH="${VLLM_RS_TARGET_PATH:-$REPO_ROOT/vllm/vllm-rs}"
-
-# Read the required toolchain from rust-toolchain.toml.
-TOOLCHAIN=$(grep '^channel' "$REPO_ROOT/rust-toolchain.toml" | sed 's/.*= *"\(.*\)"/\1/')
-
-# Ensure rustup and the required toolchain are available.
-if ! command -v rustup &>/dev/null; then
-    echo "rustup not found, installing..."
-    curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --default-toolchain none
-    source "$HOME/.cargo/env"
-fi
-
-if ! rustup run "$TOOLCHAIN" rustc --version &>/dev/null; then
-    echo "Installing Rust toolchain: $TOOLCHAIN"
-    rustup toolchain install "$TOOLCHAIN"
-fi
-
-if [[ "${1:-}" == "--debug" ]]; then
-    PROFILE_ARGS=()
-    PROFILE_DIR="debug"
-else
-    PROFILE_ARGS=(--release)
-    PROFILE_DIR="release"
-fi
-
-cargo +"$TOOLCHAIN" build "${PROFILE_ARGS[@]}" \
-    --manifest-path "$RUST_DIR/Cargo.toml" \
-    --bin vllm-rs \
-    --features native-tls-vendored
-
-mkdir -p "$(dirname "$TARGET_PATH")"
-cp "$RUST_DIR/target/$PROFILE_DIR/vllm-rs" "$TARGET_PATH"
-echo "Installed vllm-rs to $TARGET_PATH"

cmake/cpu_extension.cmake

@@ -195,12 +195,10 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
     endif()
     if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
         message(STATUS "RISC-V RVV VLEN=${VLLM_RVV_VLEN}")
-        # Sources gate FP16/BF16 paths on the compiler-provided
-        # __riscv_zvfh / __riscv_zvfbfmin macros, which GCC and clang
-        # define automatically when those extensions appear in -march.
         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)

cmake/external_projects/triton_kernels.cmake

@@ -1,6 +1,6 @@
 # Install OpenAI triton_kernels from https://github.com/triton-lang/triton/tree/main/python/triton_kernels
 
-set(DEFAULT_TRITON_KERNELS_TAG "v3.5.1")
+set(DEFAULT_TRITON_KERNELS_TAG "v3.6.0")
 
 # Set TRITON_KERNELS_SRC_DIR for use with local development with vLLM. We expect TRITON_KERNELS_SRC_DIR to
 # be directly set to the triton_kernels python directory.

cmake/utils.cmake

@@ -47,17 +47,12 @@ macro (append_cmake_prefix_path PKG EXPR)
   list(APPEND CMAKE_PREFIX_PATH ${_PREFIX_PATH})
 endmacro()
 
-# Resolve hipified output paths for `NAME` into `OUT_SRCS` and register the
-# `.cu` sources with the shared `hipify_all` target. Per-extension hipify
-# targets are unsafe to run in parallel against a shared csrc/ output dir, so
-# accumulation here is paired with a single finalize step.
+#
+# Add a target named `hipify${NAME}` that runs the hipify preprocessor on a set
+# of CUDA source files. The names of the corresponding "hipified" sources are
+# stored in `OUT_SRCS`.
+#
 function (hipify_sources_target OUT_SRCS NAME ORIG_SRCS)
-  if (TARGET hipify_all)
-    message(FATAL_ERROR
-      "hipify_sources_target(${NAME}) called after vllm_finalize_hipify_target. "
-      "Add the new HIP extension before the finalizer call in CMakeLists.txt.")
-  endif()
-
   #
   # Split into C++ and non-C++ (i.e. CUDA) sources.
   #
@@ -78,41 +73,19 @@ function (hipify_sources_target OUT_SRCS NAME ORIG_SRCS)
     list(APPEND HIP_SRCS "${CMAKE_CURRENT_BINARY_DIR}/${SRC}")
   endforeach()
 
-  set_property(GLOBAL APPEND PROPERTY VLLM_HIPIFY_ALL_SRCS ${SRCS})
-  set_property(GLOBAL APPEND PROPERTY VLLM_HIPIFY_ALL_BYPRODUCTS ${HIP_SRCS})
+  set(CSRC_BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/csrc)
+  add_custom_target(
+    hipify${NAME}
+    COMMAND ${Python_EXECUTABLE} ${CMAKE_SOURCE_DIR}/cmake/hipify.py -p ${CMAKE_SOURCE_DIR}/csrc -o ${CSRC_BUILD_DIR} ${SRCS}
+    DEPENDS ${CMAKE_SOURCE_DIR}/cmake/hipify.py ${SRCS}
+    BYPRODUCTS ${HIP_SRCS}
+    COMMENT "Running hipify on ${NAME} extension source files.")
 
   # Swap out original extension sources with hipified sources.
   list(APPEND HIP_SRCS ${CXX_SRCS})
   set(${OUT_SRCS} ${HIP_SRCS} PARENT_SCOPE)
 endfunction()
 
-# Define the single shared `hipify_all` custom target that runs hipify once
-# on the union of every HIP extension's sources. Call after the last HIP
-# `define_extension_target`.
-function (vllm_finalize_hipify_target)
-  if (TARGET hipify_all)
-    return()
-  endif()
-
-  get_property(ALL_SRCS GLOBAL PROPERTY VLLM_HIPIFY_ALL_SRCS)
-  get_property(ALL_BYPRODUCTS GLOBAL PROPERTY VLLM_HIPIFY_ALL_BYPRODUCTS)
-
-  if (NOT ALL_SRCS)
-    return()
-  endif()
-
-  list(REMOVE_DUPLICATES ALL_SRCS)
-  list(REMOVE_DUPLICATES ALL_BYPRODUCTS)
-
-  set(CSRC_BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/csrc)
-  add_custom_target(
-    hipify_all
-    COMMAND ${Python_EXECUTABLE} ${CMAKE_SOURCE_DIR}/cmake/hipify.py -p ${CMAKE_SOURCE_DIR}/csrc -o ${CSRC_BUILD_DIR} ${ALL_SRCS}
-    DEPENDS ${CMAKE_SOURCE_DIR}/cmake/hipify.py ${ALL_SRCS}
-    BYPRODUCTS ${ALL_BYPRODUCTS}
-    COMMENT "Running hipify on all extension source files.")
-endfunction()
-
 #
 # Get additional GPU compiler flags from torch.
 #
@@ -578,7 +551,7 @@ function (define_extension_target MOD_NAME)
 
   if (ARG_LANGUAGE STREQUAL "HIP")
     # Make this target dependent on the hipify preprocessor step.
-    add_dependencies(${MOD_NAME} hipify_all)
+    add_dependencies(${MOD_NAME} hipify${MOD_NAME})
     # Make sure we include the hipified versions of the headers, and avoid conflicts with the ones in the original source folder
     target_include_directories(${MOD_NAME} PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/csrc
       ${ARG_INCLUDE_DIRECTORIES})

csrc/activation_kernels.cu

@@ -1,12 +1,12 @@
-#include <cuda.h>
-#include <torch/csrc/stable/tensor.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <torch/all.h>
+#include <c10/cuda/CUDAGuard.h>
 
 #include <cmath>
 
-#include "../cuda_compat.h"
-#include "../cuda_vec_utils.cuh"
+#include "cuda_compat.h"
+#include "cuda_vec_utils.cuh"
 #include "dispatch_utils.h"
-#include "torch_utils.h"
 
 namespace vllm {
 
@@ -210,68 +210,64 @@ packed_gelu_tanh_kernel(const packed_t& val) {
     return;                                                                    \
   }                                                                            \
   dim3 grid(num_tokens);                                                       \
-  int cc_major = get_device_prop()->major;                                     \
+  int cc_major = at::cuda::getCurrentDeviceProperties()->major;                \
   int support_vec =                                                            \
       (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128)            \
           ? vllm::VecTraits<true>::ARCH_MAX_VEC_SIZE                           \
           : vllm::VecTraits<false>::ARCH_MAX_VEC_SIZE;                         \
-  int vec_size = support_vec / input.element_size();                           \
+  int vec_size = support_vec / at::elementSize(dtype);                         \
   const bool use_vec = (d % vec_size == 0);                                    \
-  const torch::stable::accelerator::DeviceGuard device_guard(                  \
-      input.get_device_index());                                               \
-  const cudaStream_t stream = get_current_cuda_stream();                       \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));            \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();                \
   if (use_vec) {                                                               \
     dim3 block(std::min(d / vec_size, 1024));                                  \
     if (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128) {         \
-      VLLM_STABLE_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {   \
+      VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {          \
         vllm::act_and_mul_kernel<                                              \
             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.mutable_data_ptr<scalar_t>(),                                  \
-            input.const_data_ptr<scalar_t>(), d, LIMIT);                       \
+            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, LIMIT);   \
       });                                                                      \
     } else {                                                                   \
-      VLLM_STABLE_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {   \
+      VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {          \
         vllm::act_and_mul_kernel<                                              \
             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.mutable_data_ptr<scalar_t>(),                                  \
-            input.const_data_ptr<scalar_t>(), d, LIMIT);                       \
+            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, LIMIT);   \
       });                                                                      \
     }                                                                          \
   } else {                                                                     \
     dim3 block(std::min(d, 1024));                                             \
-    VLLM_STABLE_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {     \
+    VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {            \
       vllm::act_and_mul_kernel<                                                \
           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.mutable_data_ptr<scalar_t>(), input.const_data_ptr<scalar_t>(),  \
-          d, LIMIT);                                                           \
+          out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, LIMIT);     \
     });                                                                        \
   }
 
-void silu_and_mul(torch::stable::Tensor& out,    // [..., d]
-                  torch::stable::Tensor& input)  // [..., 2 * d]
+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::stable::Tensor& out,    // [..., d]
-                        torch::stable::Tensor& input,  // [..., 2 * d]
+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);
 }
 
-void mul_and_silu(torch::stable::Tensor& out,    // [..., d]
-                  torch::stable::Tensor& input)  // [..., 2 * d]
+void mul_and_silu(torch::Tensor& out,    // [..., d]
+                  torch::Tensor& input)  // [..., 2 * 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.
@@ -279,15 +275,15 @@ void mul_and_silu(torch::stable::Tensor& out,    // [..., d]
                                 false, false, 0.0f);
 }
 
-void gelu_and_mul(torch::stable::Tensor& out,    // [..., d]
-                  torch::stable::Tensor& input)  // [..., 2 * d]
+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);
 }
 
-void gelu_tanh_and_mul(torch::stable::Tensor& out,    // [..., d]
-                       torch::stable::Tensor& input)  // [..., 2 * d]
+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);
@@ -438,20 +434,19 @@ __global__ void swigluoai_and_mul_kernel(
     return;                                                                    \
   }                                                                            \
   dim3 grid(num_tokens);                                                       \
-  int cc_major = get_device_prop()->major;                                     \
+  int cc_major = at::cuda::getCurrentDeviceProperties()->major;                \
   int support_vec =                                                            \
       (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128)            \
           ? vllm::VecTraits<true>::ARCH_MAX_VEC_SIZE                           \
           : vllm::VecTraits<false>::ARCH_MAX_VEC_SIZE;                         \
-  int vec_size = support_vec / input.element_size();                           \
+  int vec_size = support_vec / at::elementSize(dtype);                         \
   const bool use_vec = (d % vec_size == 0);                                    \
-  const torch::stable::accelerator::DeviceGuard device_guard(                  \
-      input.get_device_index());                                               \
-  const cudaStream_t stream = get_current_cuda_stream();                       \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));            \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();                \
   if (use_vec) {                                                               \
     dim3 block(std::min(d / vec_size, 1024));                                  \
     if (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128) {         \
-      VLLM_STABLE_DISPATCH_FLOATING_TYPES(                                     \
+      VLLM_DISPATCH_FLOATING_TYPES(                                            \
           dtype, "act_and_mul_kernel_with_param", [&] {                        \
             vllm::act_and_mul_kernel_with_param<                               \
                 scalar_t, typename vllm::PackedTypeConverter<scalar_t>::Type,  \
@@ -459,11 +454,11 @@ __global__ void swigluoai_and_mul_kernel(
                 PACKED_KERNEL<                                                 \
                     typename vllm::PackedTypeConverter<scalar_t>::Type>,       \
                 true, true><<<grid, block, 0, stream>>>(                       \
-                out.mutable_data_ptr<scalar_t>(),                              \
-                input.const_data_ptr<scalar_t>(), d, PARAM);                   \
+                out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d,       \
+                PARAM);                                                        \
           });                                                                  \
     } else {                                                                   \
-      VLLM_STABLE_DISPATCH_FLOATING_TYPES(                                     \
+      VLLM_DISPATCH_FLOATING_TYPES(                                            \
           dtype, "act_and_mul_kernel_with_param", [&] {                        \
             vllm::act_and_mul_kernel_with_param<                               \
                 scalar_t, typename vllm::PackedTypeConverter<scalar_t>::Type,  \
@@ -471,49 +466,45 @@ __global__ void swigluoai_and_mul_kernel(
                 PACKED_KERNEL<                                                 \
                     typename vllm::PackedTypeConverter<scalar_t>::Type>,       \
                 true, false><<<grid, block, 0, stream>>>(                      \
-                out.mutable_data_ptr<scalar_t>(),                              \
-                input.const_data_ptr<scalar_t>(), d, PARAM);                   \
+                out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d,       \
+                PARAM);                                                        \
           });                                                                  \
     }                                                                          \
   } else {                                                                     \
     dim3 block(std::min(d, 1024));                                             \
-    VLLM_STABLE_DISPATCH_FLOATING_TYPES(                                       \
-        dtype, "act_and_mul_kernel_with_param", [&] {                          \
-          vllm::act_and_mul_kernel_with_param<                                 \
-              scalar_t, typename vllm::PackedTypeConverter<scalar_t>::Type,    \
-              KERNEL<scalar_t>,                                                \
-              PACKED_KERNEL<                                                   \
-                  typename vllm::PackedTypeConverter<scalar_t>::Type>,         \
-              false><<<grid, block, 0, stream>>>(                              \
-              out.mutable_data_ptr<scalar_t>(),                                \
-              input.const_data_ptr<scalar_t>(), d, PARAM);                     \
-        });                                                                    \
+    VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel_with_param", [&] { \
+      vllm::act_and_mul_kernel_with_param<                                     \
+          scalar_t, typename vllm::PackedTypeConverter<scalar_t>::Type,        \
+          KERNEL<scalar_t>,                                                    \
+          PACKED_KERNEL<typename vllm::PackedTypeConverter<scalar_t>::Type>,   \
+          false><<<grid, block, 0, stream>>>(                                  \
+          out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, PARAM);     \
+    });                                                                        \
   }
 
-#define LAUNCH_SIGLUOAI_AND_MUL(KERNEL, ALPHA, LIMIT)                         \
-  int d = input.size(-1) / 2;                                                 \
-  int64_t num_tokens = input.numel() / input.size(-1);                        \
-  dim3 grid(num_tokens);                                                      \
-  dim3 block(std::min(d, 1024));                                              \
-  const torch::stable::accelerator::DeviceGuard device_guard(                 \
-      input.get_device_index());                                              \
-  const cudaStream_t stream = get_current_cuda_stream();                      \
-  VLLM_STABLE_DISPATCH_FLOATING_TYPES(                                        \
-      input.scalar_type(), "clamp_swiglu_kernel_with_params", [&] {           \
-        vllm::swigluoai_and_mul_kernel<scalar_t, KERNEL<scalar_t>>            \
-            <<<grid, block, 0, stream>>>(out.mutable_data_ptr<scalar_t>(),    \
-                                         input.const_data_ptr<scalar_t>(), d, \
-                                         ALPHA, LIMIT);                       \
+#define LAUNCH_SIGLUOAI_AND_MUL(KERNEL, ALPHA, LIMIT)                          \
+  int d = input.size(-1) / 2;                                                  \
+  int64_t num_tokens = input.numel() / input.size(-1);                         \
+  dim3 grid(num_tokens);                                                       \
+  dim3 block(std::min(d, 1024));                                               \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));            \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();                \
+  VLLM_DISPATCH_FLOATING_TYPES(                                                \
+      input.scalar_type(), "clamp_swiglu_kernel_with_params", [&] {            \
+        vllm::swigluoai_and_mul_kernel<scalar_t, KERNEL<scalar_t>>             \
+            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),             \
+                                         input.data_ptr<scalar_t>(), d, ALPHA, \
+                                         LIMIT);                               \
       });
 
-void fatrelu_and_mul(torch::stable::Tensor& out,    // [..., d],
-                     torch::stable::Tensor& input,  // [..., 2 * d]
+void fatrelu_and_mul(torch::Tensor& out,    // [..., d],
+                     torch::Tensor& input,  // [..., 2 * d]
                      double threshold) {
   LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(
       vllm::fatrelu_kernel, vllm::packed_fatrelu_kernel, threshold);
 }
-void swigluoai_and_mul(torch::stable::Tensor& out,    // [..., d]
-                       torch::stable::Tensor& input,  // [..., 2 * d]
+void swigluoai_and_mul(torch::Tensor& out,    // [..., d]
+                       torch::Tensor& input,  // [..., 2 * d]
                        double alpha, double limit) {
   LAUNCH_SIGLUOAI_AND_MUL(vllm::swigluoai_and_mul, alpha, limit);
 }
@@ -568,46 +559,45 @@ __global__ void activation_kernel(
 }  // namespace vllm
 
 // Launch element-wise activation kernel.
-#define LAUNCH_ACTIVATION_KERNEL(KERNEL)                                       \
-  auto dtype = input.scalar_type();                                            \
-  int d = input.size(-1);                                                      \
-  int64_t num_tokens = input.numel() / input.size(-1);                         \
-  if (num_tokens == 0) {                                                       \
-    return;                                                                    \
-  }                                                                            \
-  dim3 grid(num_tokens);                                                       \
-  int cc_major = get_device_prop()->major;                                     \
-  int support_vec =                                                            \
-      (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128)            \
-          ? vllm::VecTraits<true>::ARCH_MAX_VEC_SIZE                           \
-          : vllm::VecTraits<false>::ARCH_MAX_VEC_SIZE;                         \
-  int vec_size = support_vec / input.element_size();                           \
-  const bool use_vec = (d % vec_size == 0);                                    \
-  const torch::stable::accelerator::DeviceGuard device_guard(                  \
-      input.get_device_index());                                               \
-  const cudaStream_t stream = get_current_cuda_stream();                       \
-  if (use_vec) {                                                               \
-    dim3 block(std::min(d / vec_size, 1024));                                  \
-    if (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128) {         \
-      VLLM_STABLE_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {    \
-        vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, true, true>        \
-            <<<grid, block, 0, stream>>>(out.mutable_data_ptr<scalar_t>(),     \
-                                         input.const_data_ptr<scalar_t>(), d); \
-      });                                                                      \
-    } else {                                                                   \
-      VLLM_STABLE_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {    \
-        vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, true, false>       \
-            <<<grid, block, 0, stream>>>(out.mutable_data_ptr<scalar_t>(),     \
-                                         input.const_data_ptr<scalar_t>(), d); \
-      });                                                                      \
-    }                                                                          \
-  } else {                                                                     \
-    dim3 block(std::min(d, 1024));                                             \
-    VLLM_STABLE_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {      \
-      vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, false>               \
-          <<<grid, block, 0, stream>>>(out.mutable_data_ptr<scalar_t>(),       \
-                                       input.const_data_ptr<scalar_t>(), d);   \
-    });                                                                        \
+#define LAUNCH_ACTIVATION_KERNEL(KERNEL)                                 \
+  auto dtype = input.scalar_type();                                      \
+  int d = input.size(-1);                                                \
+  int64_t num_tokens = input.numel() / input.size(-1);                   \
+  if (num_tokens == 0) {                                                 \
+    return;                                                              \
+  }                                                                      \
+  dim3 grid(num_tokens);                                                 \
+  int cc_major = at::cuda::getCurrentDeviceProperties()->major;          \
+  int support_vec =                                                      \
+      (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128)      \
+          ? vllm::VecTraits<true>::ARCH_MAX_VEC_SIZE                     \
+          : vllm::VecTraits<false>::ARCH_MAX_VEC_SIZE;                   \
+  int vec_size = support_vec / at::elementSize(dtype);                   \
+  const bool use_vec = (d % vec_size == 0);                              \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));      \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();          \
+  if (use_vec) {                                                         \
+    dim3 block(std::min(d / vec_size, 1024));                            \
+    if (CUDA_VERSION >= 12090 && cc_major >= 10 && num_tokens > 128) {   \
+      VLLM_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {     \
+        vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, true, true>  \
+            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),       \
+                                         input.data_ptr<scalar_t>(), d); \
+      });                                                                \
+    } else {                                                             \
+      VLLM_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {     \
+        vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, true, false> \
+            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),       \
+                                         input.data_ptr<scalar_t>(), d); \
+      });                                                                \
+    }                                                                    \
+  } else {                                                               \
+    dim3 block(std::min(d, 1024));                                       \
+    VLLM_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {       \
+      vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, false>         \
+          <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),         \
+                                       input.data_ptr<scalar_t>(), d);   \
+    });                                                                  \
   }
 
 namespace vllm {
@@ -635,20 +625,20 @@ __device__ __forceinline__ T gelu_quick_kernel(const T& x) {
 
 }  // namespace vllm
 
-void gelu_new(torch::stable::Tensor& out,    // [..., d]
-              torch::stable::Tensor& input)  // [..., d]
+void gelu_new(torch::Tensor& out,    // [..., d]
+              torch::Tensor& input)  // [..., d]
 {
   LAUNCH_ACTIVATION_KERNEL(vllm::gelu_new_kernel);
 }
 
-void gelu_fast(torch::stable::Tensor& out,    // [..., d]
-               torch::stable::Tensor& input)  // [..., d]
+void gelu_fast(torch::Tensor& out,    // [..., d]
+               torch::Tensor& input)  // [..., d]
 {
   LAUNCH_ACTIVATION_KERNEL(vllm::gelu_fast_kernel);
 }
 
-void gelu_quick(torch::stable::Tensor& out,    // [..., d]
-                torch::stable::Tensor& input)  // [..., d]
+void gelu_quick(torch::Tensor& out,    // [..., d]
+                torch::Tensor& input)  // [..., d]
 {
   LAUNCH_ACTIVATION_KERNEL(vllm::gelu_quick_kernel);
 }

csrc/attention/dtype_fp8.cuh

@@ -1,7 +1,6 @@
 #pragma once
 
 #include "attention_generic.cuh"
-#include "torch_utils.h"
 
 #include <stdint.h>
 #ifdef ENABLE_FP8
@@ -31,7 +30,7 @@ inline Fp8KVCacheDataType get_fp8_kv_cache_data_type(
   } else if (dtype_str == "fp8_e5m2") {
     return Fp8KVCacheDataType::kFp8E5M2;
   }
-  TORCH_UTILS_CHECK(false, "Unsupported fp8 kv cache data type: ", dtype_str);
+  TORCH_CHECK(false, "Unsupported fp8 kv cache data type: ", dtype_str);
 }
 
 // fp8 vector types for quantization of kv cache

csrc/attention/vertical_slash_index.cu

@@ -0,0 +1,401 @@
+// Copyright (c) Microsoft Corporation.
+// Licensed under the MIT license.
+
+#include <assert.h>
+
+#include <cuda.h>
+
+#include <torch/all.h>
+
+__device__ int64_t save_blocks(int* block_offset, int64_t range_start,
+                               int64_t range_end, int64_t block_size,
+                               int64_t input_block_count, int64_t kv_seqlen) {
+  if (range_start >= kv_seqlen) {
+    return input_block_count;
+  }
+  if (range_end > kv_seqlen) {
+    range_end = kv_seqlen;
+  }
+  int64_t current_block_count = input_block_count;
+  for (int idx = range_start; idx < range_end; idx += block_size) {
+    block_offset[current_block_count++] = idx;
+  }
+  return current_block_count;
+}
+
+__global__ void convert_vertical_slash_indexes_kernel(
+    const int* q_seqlens,         // [BATCH, ]
+    const int* kv_seqlens,        // [BATCH, ]
+    const int* vertical_indexes,  // [BATCH, N_HEADS, NNZ_V]
+    const int* slash_indexes,     // [BATCH, N_HEADS, NNZ_S]
+    int* block_count,             // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* block_offset,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_S]
+    int* column_count,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* column_index,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_V]
+    int64_t N_HEADS, int64_t N_ROWS, int64_t BLOCK_SIZE_M, int64_t BLOCK_SIZE_N,
+    int64_t NNZ_V, int64_t NNZ_S,
+    bool causal  // True for intra, False for succ
+) {
+  const int batch_idx = blockIdx.y;
+  const int head_idx = blockIdx.x;
+  const int group_idx = blockIdx.z;
+
+  int64_t q_seqlen = q_seqlens[batch_idx];
+  int64_t kv_seqlen = kv_seqlens[batch_idx];
+  int64_t block_idx_m = group_idx * blockDim.x + threadIdx.x;
+  int64_t start_m = block_idx_m * BLOCK_SIZE_M;
+  if (start_m >= q_seqlen) {
+    return;
+  }
+  int64_t end_m = start_m + BLOCK_SIZE_M;
+  vertical_indexes += (batch_idx * N_HEADS + head_idx) * NNZ_V;
+  slash_indexes += (batch_idx * N_HEADS + head_idx) * NNZ_S;
+  int64_t row_offset = (batch_idx * N_HEADS + head_idx) * N_ROWS + block_idx_m;
+  block_count += row_offset;
+  block_offset += row_offset * NNZ_S;
+  column_count += row_offset;
+  column_index += row_offset * NNZ_V;
+
+  bool has_slash = true;
+  int64_t tmp_col_cnt = 0, tmp_blk_cnt = 0;
+  int64_t s = 0, v = 0;
+  int64_t v_idx = vertical_indexes[v++];
+  int64_t s_idx = slash_indexes[s++];
+  if (causal) {
+    while (s_idx >= end_m + (kv_seqlen - q_seqlen) && s < NNZ_S) {
+      s_idx = slash_indexes[s++];
+    }
+    if (s_idx > end_m + (kv_seqlen - q_seqlen)) has_slash = false;
+    s_idx = max((kv_seqlen - q_seqlen) + end_m - s_idx, BLOCK_SIZE_M);
+  } else {
+    while (s_idx >= end_m + kv_seqlen && s < NNZ_S) {
+      s_idx = slash_indexes[s++];
+    }
+    if (s_idx > end_m + kv_seqlen) has_slash = false;
+    s_idx = max(kv_seqlen + end_m - s_idx, BLOCK_SIZE_M);
+  }
+
+  int64_t range_start = s_idx - BLOCK_SIZE_M, range_end = s_idx;
+  if (!has_slash) {
+    if (causal) {
+      range_start = (kv_seqlen - q_seqlen) + end_m;
+      range_end = (kv_seqlen - q_seqlen) + end_m + BLOCK_SIZE_N;
+    } else {
+      range_start = kv_seqlen;
+      range_end = kv_seqlen + BLOCK_SIZE_N;
+    }
+  }
+
+  bool slash_finished = false;
+  while (1) {
+    if (v_idx < range_end) {
+      if (v_idx < range_start) {
+        column_index[tmp_col_cnt++] = v_idx;
+      }
+      if (v < NNZ_V) {
+        v_idx = vertical_indexes[v++];
+      } else {
+        if (causal)
+          v_idx = end_m + BLOCK_SIZE_N + (kv_seqlen - q_seqlen);
+        else
+          v_idx = end_m + BLOCK_SIZE_N + kv_seqlen;
+      }
+    } else {
+      if ((s < NNZ_S && causal) ||
+          (s < NNZ_S && !causal && slash_indexes[s] >= start_m)) {
+        if (causal)
+          s_idx = max((kv_seqlen - q_seqlen) + end_m - slash_indexes[s++],
+                      BLOCK_SIZE_M);
+        else
+          s_idx = max(kv_seqlen + end_m - slash_indexes[s++], BLOCK_SIZE_M);
+      } else {
+        if (v == NNZ_V || (v_idx > range_start && causal)) {
+          // add the last vertical if no more slash
+          if (v == NNZ_V && !causal && v_idx < kv_seqlen) {
+            column_index[tmp_col_cnt++] = v_idx;
+          }
+          tmp_blk_cnt = save_blocks(block_offset, range_start, range_end,
+                                    BLOCK_SIZE_N, tmp_blk_cnt, kv_seqlen);
+          break;
+        } else {
+          if (causal) {
+            range_start = (kv_seqlen - q_seqlen) + end_m;
+            range_end = (kv_seqlen - q_seqlen) + end_m + BLOCK_SIZE_N;
+          } else {
+            // if slash_finished but there are vertical left, save current
+            // blocks
+            tmp_blk_cnt = save_blocks(block_offset, range_start, range_end,
+                                      BLOCK_SIZE_N, tmp_blk_cnt, kv_seqlen);
+            range_start = kv_seqlen;
+            range_end = kv_seqlen + BLOCK_SIZE_N;
+          }
+          slash_finished = true;
+        }
+      }
+      if (!slash_finished) {
+        if (s_idx > range_end + BLOCK_SIZE_M) {
+          tmp_blk_cnt = save_blocks(block_offset, range_start, range_end,
+                                    BLOCK_SIZE_N, tmp_blk_cnt, kv_seqlen);
+          range_start = s_idx - BLOCK_SIZE_M;
+          range_end = s_idx;
+        } else if (s_idx > range_end) {
+          range_end += BLOCK_SIZE_M;
+        }
+      }
+    }
+  }
+
+  block_count[0] = tmp_blk_cnt;
+  column_count[0] = tmp_col_cnt;
+}
+
+void convert_vertical_slash_indexes_64x64(
+    const int* q_seqlens,         // [BATCH, ]
+    const int* kv_seqlens,        // [BATCH, ]
+    const int* vertical_indexes,  // [BATCH, N_HEADS, NNZ_V]
+    const int* slash_indexes,     // [BATCH, N_HEADS, NNZ_S]
+    int* block_count,             // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* block_offset,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_S]
+    int* column_count,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* column_index,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_V]
+    int64_t BATCH_SIZE, int64_t N_HEADS, int64_t N_ROWS, int64_t BLOCK_SIZE_M,
+    int64_t BLOCK_SIZE_N, int64_t NNZ_V, int64_t NNZ_S, bool causal) {
+  const int N_THREADS = 64;
+  const dim3 dimBlock(N_THREADS);
+  const dim3 dimGrid(N_HEADS, BATCH_SIZE, (N_ROWS + N_THREADS - 1) / N_THREADS);
+  convert_vertical_slash_indexes_kernel<<<dimGrid, dimBlock>>>(
+      q_seqlens, kv_seqlens, vertical_indexes, slash_indexes, block_count,
+      block_offset, column_count, column_index, N_HEADS, N_ROWS, BLOCK_SIZE_M,
+      BLOCK_SIZE_N, NNZ_V, NNZ_S, causal);
+}
+
+/**
+ * Implements the Algorithm 4 in paper https://arxiv.org/abs/2407.02490.
+ *
+ * This function builds the index of each row of blocks from vertical indices
+ * and slash indices. The vertical indices are treated as points, while the
+ * slash indices are converted as ranges. The output consists of the merged
+ * ranges and separate column indices, where the ranges are represented by
+ * block indices.
+ *
+ * The implementation is referenced from the original MInference repo:
+ * https://github.com/microsoft/MInference/blob/main/csrc/vertical_slash_index.cu.
+ */
+void convert_vertical_slash_indexes(
+    torch::Tensor& block_count,      // [BATCH, N_HEADS, NUM_ROWS]
+    torch::Tensor& block_offset,     // [BATCH, N_HEADS, NUM_ROWS, NNZ_S]
+    torch::Tensor& column_count,     // [BATCH, N_HEADS, NUM_ROWS]
+    torch::Tensor& column_index,     // [BATCH, N_HEADS, NUM_ROWS, NNZ_V]
+    torch::Tensor q_seqlens,         // [BATCH, ]
+    torch::Tensor kv_seqlens,        // [BATCH, ]
+    torch::Tensor vertical_indexes,  // [BATCH, N_HEADS, NNZ_V]
+    torch::Tensor slash_indexes,     // [BATCH, N_HEADS, NNZ_S]
+    int64_t context_size, int64_t block_size_M, int64_t block_size_N,
+    bool causal) {
+  cudaSetDevice(q_seqlens.get_device());
+
+  int batch_size = slash_indexes.size(0);
+  int num_heads = slash_indexes.size(1);
+  int nnz_slash = slash_indexes.size(2);
+  int nnz_vertical = vertical_indexes.size(2);
+  int num_rows = (context_size + block_size_M - 1) / block_size_M;
+
+  convert_vertical_slash_indexes_64x64(
+      q_seqlens.data_ptr<int>(), kv_seqlens.data_ptr<int>(),
+      vertical_indexes.data_ptr<int>(), slash_indexes.data_ptr<int>(),
+      block_count.data_ptr<int>(), block_offset.data_ptr<int>(),
+      column_count.data_ptr<int>(), column_index.data_ptr<int>(), batch_size,
+      num_heads, num_rows, block_size_M, block_size_N, nnz_vertical, nnz_slash,
+      causal);
+}
+
+__global__ void convert_vertical_slash_indexes_kernel_mergehead(
+    const int* q_seqlens,         // [BATCH, ]
+    const int* kv_seqlens,        // [BATCH, ]
+    const int* vertical_indexes,  // [BATCH, N_HEADS, NNZ_V]
+    const int* slash_indexes,     // [BATCH, N_HEADS, NNZ_S]
+    const int* per_head_vertical_topkv, const int* per_head_slash_topkv,
+    int* block_count,   // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* block_offset,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_S]
+    int* column_count,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* column_index,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_V]
+    int64_t N_HEADS, int64_t N_ROWS, int64_t BLOCK_SIZE_M, int64_t BLOCK_SIZE_N,
+    int64_t NNZ_V, int64_t NNZ_S,
+    bool causal  // True for intra, False for succ
+) {
+  const int batch_idx = blockIdx.y;
+  const int head_idx = blockIdx.x;
+  const int group_idx = blockIdx.z;
+
+  int64_t q_seqlen = q_seqlens[batch_idx];
+  int64_t kv_seqlen = kv_seqlens[batch_idx];
+  int64_t block_idx_m = group_idx * blockDim.x + threadIdx.x;
+  int64_t start_m = block_idx_m * BLOCK_SIZE_M;
+  if (start_m >= q_seqlen) {
+    return;
+  }
+  int64_t end_m = start_m + BLOCK_SIZE_M;
+  vertical_indexes += (batch_idx * N_HEADS + head_idx) * NNZ_V;
+  slash_indexes += (batch_idx * N_HEADS + head_idx) * NNZ_S;
+  int64_t row_offset = (batch_idx * N_HEADS + head_idx) * N_ROWS + block_idx_m;
+  block_count += row_offset;
+  block_offset += row_offset * NNZ_S;
+  column_count += row_offset;
+  column_index += row_offset * NNZ_V;
+
+  // MergeHead: each head has it's unique max topk NNZ_V，NNZ_S. (NNZ_V，NNZ_S
+  // above is buffer size, use to compute offset)
+  NNZ_S = per_head_slash_topkv[head_idx];
+  NNZ_V = per_head_vertical_topkv[head_idx];
+
+  bool has_slash = true;
+  int64_t tmp_col_cnt = 0, tmp_blk_cnt = 0;
+  int64_t s = 0, v = 0;
+  int64_t v_idx = vertical_indexes[v++];
+  int64_t s_idx = slash_indexes[s++];
+  if (causal) {
+    while (s_idx >= end_m + (kv_seqlen - q_seqlen) && s < NNZ_S) {
+      s_idx = slash_indexes[s++];
+    }
+    if (s_idx > end_m + (kv_seqlen - q_seqlen)) has_slash = false;
+    s_idx = max((kv_seqlen - q_seqlen) + end_m - s_idx, BLOCK_SIZE_M);
+  } else {
+    while (s_idx >= end_m + kv_seqlen && s < NNZ_S) {
+      s_idx = slash_indexes[s++];
+    }
+    if (s_idx > end_m + kv_seqlen) has_slash = false;
+    s_idx = max(kv_seqlen + end_m - s_idx, BLOCK_SIZE_M);
+  }
+
+  int64_t range_start = s_idx - BLOCK_SIZE_M, range_end = s_idx;
+  if (!has_slash) {
+    if (causal) {
+      range_start = (kv_seqlen - q_seqlen) + end_m;
+      range_end = (kv_seqlen - q_seqlen) + end_m + BLOCK_SIZE_N;
+    } else {
+      range_start = kv_seqlen;
+      range_end = kv_seqlen + BLOCK_SIZE_N;
+    }
+  }
+
+  bool slash_finished = false;
+  while (1) {
+    if (v_idx < range_end) {
+      if (v_idx < range_start) {
+        column_index[tmp_col_cnt++] = v_idx;
+      }
+      if (v < NNZ_V) {
+        v_idx = vertical_indexes[v++];
+      } else {
+        if (causal)
+          v_idx = end_m + BLOCK_SIZE_N + (kv_seqlen - q_seqlen);
+        else
+          v_idx = end_m + BLOCK_SIZE_N + kv_seqlen;
+      }
+    } else {
+      if ((s < NNZ_S && causal) ||
+          (s < NNZ_S && !causal && slash_indexes[s] >= start_m)) {
+        if (causal)
+          s_idx = max((kv_seqlen - q_seqlen) + end_m - slash_indexes[s++],
+                      BLOCK_SIZE_M);
+        else
+          s_idx = max(kv_seqlen + end_m - slash_indexes[s++], BLOCK_SIZE_M);
+      } else {
+        if (v == NNZ_V || (v_idx > range_start && causal)) {
+          // add the last vertical if no more slash
+          if (v == NNZ_V && !causal && v_idx < kv_seqlen) {
+            column_index[tmp_col_cnt++] = v_idx;
+          }
+          tmp_blk_cnt = save_blocks(block_offset, range_start, range_end,
+                                    BLOCK_SIZE_N, tmp_blk_cnt, kv_seqlen);
+          break;
+        } else {
+          if (causal) {
+            range_start = (kv_seqlen - q_seqlen) + end_m;
+            range_end = (kv_seqlen - q_seqlen) + end_m + BLOCK_SIZE_N;
+          } else {
+            // if slash_finished but there are vertical left, save current
+            // blocks
+            tmp_blk_cnt = save_blocks(block_offset, range_start, range_end,
+                                      BLOCK_SIZE_N, tmp_blk_cnt, kv_seqlen);
+            range_start = kv_seqlen;
+            range_end = kv_seqlen + BLOCK_SIZE_N;
+          }
+          slash_finished = true;
+        }
+      }
+      if (!slash_finished) {
+        if (s_idx > range_end + BLOCK_SIZE_M) {
+          tmp_blk_cnt = save_blocks(block_offset, range_start, range_end,
+                                    BLOCK_SIZE_N, tmp_blk_cnt, kv_seqlen);
+          range_start = s_idx - BLOCK_SIZE_M;
+          range_end = s_idx;
+        } else if (s_idx > range_end) {
+          range_end += BLOCK_SIZE_M;
+        }
+      }
+    }
+  }
+
+  block_count[0] = tmp_blk_cnt;
+  column_count[0] = tmp_col_cnt;
+}
+
+void convert_vertical_slash_indexes_64x64_mergehead(
+    const int* q_seqlens,         // [BATCH, ]
+    const int* kv_seqlens,        // [BATCH, ]
+    const int* vertical_indexes,  // [BATCH, N_HEADS, NNZ_V]
+    const int* slash_indexes,     // [BATCH, N_HEADS, NNZ_S]
+    int* per_head_vertical_topkv, int* per_head_slash_topkv,
+    int* block_count,   // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* block_offset,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_S]
+    int* column_count,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M)]
+    int* column_index,  // [BATCH, N_HEADS, cdiv(N_CTX, BLOCK_SIZE_M), NNZ_V]
+    int64_t BATCH_SIZE, int64_t N_HEADS, int64_t N_ROWS, int64_t BLOCK_SIZE_M,
+    int64_t BLOCK_SIZE_N, int64_t NNZ_V, int64_t NNZ_S, bool causal) {
+  const int N_THREADS = 64;
+  const dim3 dimBlock(N_THREADS);
+  const dim3 dimGrid(N_HEADS, BATCH_SIZE, (N_ROWS + N_THREADS - 1) / N_THREADS);
+  convert_vertical_slash_indexes_kernel_mergehead<<<dimGrid, dimBlock>>>(
+      q_seqlens, kv_seqlens, vertical_indexes, slash_indexes,
+      per_head_vertical_topkv, per_head_slash_topkv, block_count, block_offset,
+      column_count, column_index, N_HEADS, N_ROWS, BLOCK_SIZE_M, BLOCK_SIZE_N,
+      NNZ_V, NNZ_S, causal);
+}
+
+/**
+ * Implements the Algorithm 4 in paper https://arxiv.org/abs/2407.02490.
+ *
+ * Like the above convert_vertical_slash_indexes, but with
+ * pre-computed vertical and slash counts.
+ */
+void convert_vertical_slash_indexes_mergehead(
+    torch::Tensor& block_count,            // [BATCH, N_HEADS, NUM_ROWS]
+    torch::Tensor& block_offset,           // [BATCH, N_HEADS, NUM_ROWS, NNZ_S]
+    torch::Tensor& column_count,           // [BATCH, N_HEADS, NUM_ROWS]
+    torch::Tensor& column_index,           // [BATCH, N_HEADS, NUM_ROWS, NNZ_V]
+    torch::Tensor q_seqlens,               // [BATCH, ]
+    torch::Tensor kv_seqlens,              // [BATCH, ]
+    torch::Tensor vertical_indexes,        // [BATCH, N_HEADS, NNZ_V]
+    torch::Tensor slash_indexes,           // [BATCH, N_HEADS, NNZ_S]
+    torch::Tensor vertical_indices_count,  // [N_HEADS, ]
+    torch::Tensor slash_indices_count,     // [N_HEADS, ]
+    int64_t context_size, int64_t block_size_M, int64_t block_size_N,
+    bool causal) {
+  cudaSetDevice(q_seqlens.get_device());
+
+  int batch_size = slash_indexes.size(0);
+  int num_heads = slash_indexes.size(1);
+  int nnz_slash = slash_indexes.size(2);
+  int nnz_vertical = vertical_indexes.size(2);
+  int num_rows = (context_size + block_size_M - 1) / block_size_M;
+
+  convert_vertical_slash_indexes_64x64_mergehead(
+      q_seqlens.data_ptr<int>(), kv_seqlens.data_ptr<int>(),
+      vertical_indexes.data_ptr<int>(), slash_indexes.data_ptr<int>(),
+      vertical_indices_count.data_ptr<int>(),
+      slash_indices_count.data_ptr<int>(), block_count.data_ptr<int>(),
+      block_offset.data_ptr<int>(), column_count.data_ptr<int>(),
+      column_index.data_ptr<int>(), batch_size, num_heads, num_rows,
+      block_size_M, block_size_N, nnz_vertical, nnz_slash, causal);
+}

csrc/cpu/cpu_arch_macros.h

@@ -156,20 +156,4 @@
 
 #endif  // __aarch64__
 
-// RISC-V RVV
-#ifdef __riscv_v
-  #include <riscv_vector.h>
-
-  #ifdef __riscv_zihintpause
-    #define FAST_SPINNING __riscv_pause();
-  #endif
-
-  // FP32Vec16::exp() in cpu_types_riscv.hpp already implements the full
-  // polynomial approximation for RVV, so we simply delegate to it.
-  #define DEFINE_FAST_EXP                             \
-    auto fast_exp = [&](const vec_op::FP32Vec16& vec) \
-                        __attribute__((always_inline)) { return vec.exp(); };
-
-#endif  // __riscv_v
-
 #endif

csrc/cpu/cpu_attn.cpp

@@ -29,8 +29,6 @@ torch::Tensor get_scheduler_metadata(
     isa = cpu_attention::ISA::NEON;
   } else if (isa_hint == "vxe") {
     isa = cpu_attention::ISA::VXE;
-  } else if (isa_hint == "rvv") {
-    isa = cpu_attention::ISA::RVV;
   } else if (isa_hint == "vsx") {
     isa = cpu_attention::ISA::VSX;
   } else {
@@ -133,8 +131,6 @@ void cpu_attn_reshape_and_cache(
       return cpu_attention::ISA::NEON;
     } else if (isa == "vxe") {
       return cpu_attention::ISA::VXE;
-    } else if (isa == "rvv") {
-      return cpu_attention::ISA::RVV;
     } else if (isa == "vsx") {
       return cpu_attention::ISA::VSX;
     } else {

csrc/cpu/cpu_attn_impl.hpp

@@ -12,7 +12,7 @@
 #include "cpu/utils.hpp"
 
 namespace cpu_attention {
-enum class ISA { AMX, VEC, VEC16, NEON, VXE, RVV, VSX };
+enum class ISA { AMX, VEC, VEC16, NEON, VXE, VSX };
 
 // Mirrors csrc/attention/dtype_fp8.cuh Fp8KVCacheDataType exactly.
 enum class Fp8KVCacheDataType {
@@ -164,9 +164,6 @@ struct AttentionMetadata {
       case ISA::VXE:
         ss << "VXE, ";
         break;
-      case ISA::RVV:
-        ss << "RVV, ";
-        break;
       case ISA::VSX:
         ss << "VSX, ";
         break;
@@ -408,19 +405,9 @@ class AttentionScheduler {
     const int64_t cache_size = cpu_utils::get_available_l2_size();
     const int32_t max_num_q_per_iter = input.max_num_q_per_iter;
     const int32_t kv_len_alignment = input.kv_block_alignment;
-    bool has_decode_request = false;
-    bool decode_only_batch = true;
-    for (int32_t req_id = 0; req_id < input.num_reqs; ++req_id) {
-      const int32_t q_token_num =
-          input.query_start_loc[req_id + 1] - input.query_start_loc[req_id];
-      has_decode_request = has_decode_request || (q_token_num == 1);
-      decode_only_batch = decode_only_batch && (q_token_num == 1);
-    }
     int32_t q_head_per_kv = input.num_heads_q / input.num_heads_kv;
-    const bool supports_gqa = q_head_per_kv <= max_num_q_per_iter;
-    const bool use_gqa_fast_path = supports_gqa && decode_only_batch;
-    const bool use_gqa_scratchpad = supports_gqa && has_decode_request;
-    if (!use_gqa_scratchpad) {
+    const bool use_gqa = (max_num_q_per_iter % q_head_per_kv == 0);
+    if (!use_gqa) {
       q_head_per_kv = 1;  // fallback to MHA
     }
     const int32_t min_split_kv_len =
@@ -469,7 +456,7 @@ class AttentionScheduler {
     const int64_t kv_len_per_thread =
         (((total_kv_len / thread_num) + kv_len_alignment - 1) /
          kv_len_alignment) *
-        kv_len_alignment;
+        kv_len_alignment * (use_gqa ? input.num_heads_kv : input.num_heads_q);
     std::vector<AttentionWorkItemGroup> workitems;
     std::vector<ReductionWorkItemGroup> reduce_workitems;
     workitems.reserve(1024);
@@ -690,7 +677,7 @@ class AttentionScheduler {
         metadata_ptr->attention_scratchpad_size_per_thread *
             metadata_ptr->thread_num +
         metadata_ptr->reduction_scratchpad_size_per_kv_head *
-            (use_gqa_fast_path ? input.num_heads_kv : input.num_heads_q);
+            (use_gqa ? input.num_heads_kv : input.num_heads_q);
     cpu_utils::ScratchPadManager::get_scratchpad_manager()->realloc(
         scratchpad_size);
 
@@ -1419,24 +1406,13 @@ class AttentionMainLoop {
       const int32_t q_head_num = input->num_heads;
       const int32_t kv_head_num = input->num_kv_heads;
       const int32_t q_heads_per_kv = q_head_num / kv_head_num;
-      AttentionWorkItemGroup* const workitem_groups =
-          metadata.workitem_groups_ptr;
-      const int32_t* cu_workitem_num_per_thread =
-          metadata.cu_workitem_num_per_thread;
-      ReductionWorkItemGroup* const reduction_items =
-          metadata.reduction_items_ptr;
-      const bool supports_gqa = q_heads_per_kv <= max_q_head_num_per_iter;
-      bool decode_only_batch = true;
-      for (int32_t i = 0; i < metadata.workitem_group_num; ++i) {
-        decode_only_batch =
-            decode_only_batch && (workitem_groups[i].q_token_num == 1);
-      }
-      const bool use_gqa_fast_path = supports_gqa && decode_only_batch;
-      const int32_t actual_kv_head_num =
-          use_gqa_fast_path ? kv_head_num : q_head_num;
-      const int32_t actual_q_heads_per_kv =
-          use_gqa_fast_path ? q_heads_per_kv : 1;
+      const bool use_gqa =
+          (max_q_head_num_per_iter % q_heads_per_kv == 0) ? true : false;
+      const int32_t actual_kv_head_num = use_gqa ? kv_head_num : q_head_num;
+      const int32_t actual_q_heads_per_kv = use_gqa ? q_heads_per_kv : 1;
       TORCH_CHECK_LE(actual_q_heads_per_kv, max_q_head_num_per_iter);
+      const int32_t max_q_token_num_per_iter =
+          max_q_head_num_per_iter / actual_q_heads_per_kv;
       const int64_t q_token_num_stride = input->query_num_tokens_stride;
       const int64_t q_head_num_stride = input->query_num_heads_stride;
       const int64_t kv_cache_head_num_stride = input->cache_num_kv_heads_stride;
@@ -1482,6 +1458,15 @@ class AttentionMainLoop {
               sizeof(q_buffer_t), sizeof(logits_buffer_t),
               sizeof(partial_output_buffer_t), max_q_head_num_per_iter,
               max_q_head_num_per_iter);
+      const int32_t default_q_tile_token_num =
+          default_tile_size / actual_q_heads_per_kv;
+
+      AttentionWorkItemGroup* const workitem_groups =
+          metadata.workitem_groups_ptr;
+      const int32_t* cu_workitem_num_per_thread =
+          metadata.cu_workitem_num_per_thread;
+      ReductionWorkItemGroup* const reduction_items =
+          metadata.reduction_items_ptr;
 
       const int32_t effective_thread_num = metadata.effective_thread_num;
       const int32_t reduction_item_num = metadata.reduction_item_num;
@@ -1525,6 +1510,8 @@ class AttentionMainLoop {
               cu_workitem_num_per_thread[thread_offset + 1] -
               cu_workitem_num_per_thread[thread_offset];
 
+          const int32_t q_head_start_idx = kv_head_idx * actual_q_heads_per_kv;
+
           for (int32_t workitem_group_idx = 0;
                workitem_group_idx < curr_workitem_groups_num;
                ++workitem_group_idx) {
@@ -1539,21 +1526,6 @@ class AttentionMainLoop {
             const int32_t q_token_id_start =
                 current_workitem_group->q_token_id_start;
             const int32_t q_token_num = current_workitem_group->q_token_num;
-            const bool curr_use_gqa =
-                use_gqa_fast_path || (supports_gqa && q_token_num == 1);
-            if (!use_gqa_fast_path && curr_use_gqa &&
-                kv_head_idx % q_heads_per_kv != 0) {
-              continue;
-            }
-            const int32_t curr_q_heads_per_kv =
-                curr_use_gqa ? q_heads_per_kv : 1;
-            const int32_t curr_max_q_token_num_per_iter =
-                max_q_head_num_per_iter / curr_q_heads_per_kv;
-            const int32_t curr_default_q_tile_token_num =
-                default_tile_size / curr_q_heads_per_kv;
-            const int32_t q_head_start_idx =
-                use_gqa_fast_path ? (kv_head_idx * q_heads_per_kv)
-                                  : kv_head_idx;
 
             // taskgroup general information
             const int32_t q_end = input->query_start_loc[current_group_idx + 1];
@@ -1567,7 +1539,7 @@ class AttentionMainLoop {
                              current_workitem_group->local_split_id == 0);
 
             for (int32_t q_token_offset = 0; q_token_offset < q_token_num;
-                 q_token_offset += curr_default_q_tile_token_num) {
+                 q_token_offset += default_q_tile_token_num) {
               bool first_iter_flag[AttentionScheduler::MaxQTileIterNum];
               for (int32_t i = 0; i < AttentionScheduler::MaxQTileIterNum;
                    ++i) {
@@ -1577,9 +1549,9 @@ class AttentionMainLoop {
               const int32_t q_token_start_idx =
                   q_start + q_token_offset + q_token_id_start;
               const int32_t actual_q_token_num = std::min(
-                  curr_default_q_tile_token_num, q_token_num - q_token_offset);
+                  default_q_tile_token_num, q_token_num - q_token_offset);
               const int32_t q_head_tile_size =
-                  actual_q_token_num * curr_q_heads_per_kv;
+                  actual_q_token_num * actual_q_heads_per_kv;
               const int32_t rounded_q_head_tile_size =
                   ((q_head_tile_size + max_q_head_num_per_iter - 1) /
                    max_q_head_num_per_iter) *
@@ -1616,9 +1588,10 @@ class AttentionMainLoop {
                   AttentionScheduler::align_kv_tile_pos(
                       kv_tile_start_pos, kv_tile_end_pos, blocksize_alignment);
 
-              const int32_t curr_kv_head_idx =
-                  use_gqa_fast_path ? kv_head_idx
-                                    : (kv_head_idx / q_heads_per_kv);
+              int32_t curr_kv_head_idx =
+                  use_gqa ? kv_head_idx
+                          : (kv_head_idx /
+                             q_heads_per_kv);  // for GQA disabled case
 
               // std::printf("thread_id: %d, req_id: %d, q_token_start: %d,
               // q_token_end: %d, q_head_start: %d, q_head_end: %d, kv_head_idx:
@@ -1653,12 +1626,12 @@ class AttentionMainLoop {
                   (s_aux != nullptr ? s_aux + q_head_start_idx : nullptr);
 
               // copy the Q tile to q_buffer, the logical layout of q_buffer is
-              // [actual_q_token_num, curr_q_heads_per_kv, head_dim]
+              // [actual_q_token_num, actual_q_heads_per_kv, head_dim]
               {
                 attn_impl.copy_q_heads_tile(
                     q_tile_ptr, q_buffer, actual_q_token_num,
-                    curr_q_heads_per_kv, q_token_num_stride, q_head_num_stride,
-                    scale);
+                    actual_q_heads_per_kv, q_token_num_stride,
+                    q_head_num_stride, scale);
               }
 
               if (use_sink) {
@@ -1672,29 +1645,29 @@ class AttentionMainLoop {
                 float* __restrict__ curr_max_buffer = max_buffer;
                 for (int32_t token_idx = 0; token_idx < actual_q_token_num;
                      ++token_idx) {
-                  for (int32_t head_idx = 0; head_idx < curr_q_heads_per_kv;
+                  for (int32_t head_idx = 0; head_idx < actual_q_heads_per_kv;
                        ++head_idx) {
                     curr_sum_buffer[head_idx] = 1.0f;
                     curr_max_buffer[head_idx] = s_aux_fp32[head_idx];
                   }
 
-                  curr_sum_buffer += curr_q_heads_per_kv;
-                  curr_max_buffer += curr_q_heads_per_kv;
+                  curr_sum_buffer += actual_q_heads_per_kv;
+                  curr_max_buffer += actual_q_heads_per_kv;
                 }
               } else {
                 float* __restrict__ curr_sum_buffer = sum_buffer;
                 float* __restrict__ curr_max_buffer = max_buffer;
                 for (int32_t token_idx = 0; token_idx < actual_q_token_num;
                      ++token_idx) {
-                  for (int32_t head_idx = 0; head_idx < curr_q_heads_per_kv;
+                  for (int32_t head_idx = 0; head_idx < actual_q_heads_per_kv;
                        ++head_idx) {
                     curr_sum_buffer[head_idx] = 0.0f;
                     curr_max_buffer[head_idx] =
                         std::numeric_limits<float>::lowest();
                   }
 
-                  curr_sum_buffer += curr_q_heads_per_kv;
-                  curr_max_buffer += curr_q_heads_per_kv;
+                  curr_sum_buffer += actual_q_heads_per_kv;
+                  curr_max_buffer += actual_q_heads_per_kv;
                 }
               }
 
@@ -1707,17 +1680,16 @@ class AttentionMainLoop {
                     kv_tile_pos_left + kv_tile_size, rounded_kv_tile_end_pos);
                 for (int32_t q_head_tile_token_offset = 0;
                      q_head_tile_token_offset < actual_q_token_num;
-                     q_head_tile_token_offset +=
-                     curr_max_q_token_num_per_iter) {
+                     q_head_tile_token_offset += max_q_token_num_per_iter) {
                   const int32_t q_tile_pos_left =
                       q_tile_start_pos + q_head_tile_token_offset;
                   const int32_t q_tile_token_num =
-                      std::min(curr_max_q_token_num_per_iter,
+                      std::min(max_q_token_num_per_iter,
                                actual_q_token_num - q_head_tile_token_offset);
                   const int32_t q_tile_head_offset =
-                      q_head_tile_token_offset * curr_q_heads_per_kv;
+                      q_head_tile_token_offset * actual_q_heads_per_kv;
                   const int32_t q_tile_head_num =
-                      q_tile_token_num * curr_q_heads_per_kv;
+                      q_tile_token_num * actual_q_heads_per_kv;
                   const int32_t q_tile_pos_right =
                       q_tile_pos_left + q_tile_token_num;
                   const auto [actual_kv_tile_pos_left,
@@ -1727,7 +1699,7 @@ class AttentionMainLoop {
                           q_tile_pos_right, sliding_window_left,
                           sliding_window_right);
                   const int32_t q_iter_idx =
-                      q_head_tile_token_offset / curr_max_q_token_num_per_iter;
+                      q_head_tile_token_offset / max_q_token_num_per_iter;
 
                   if (actual_kv_tile_pos_right <= actual_kv_tile_pos_left) {
                     continue;
@@ -1793,7 +1765,7 @@ class AttentionMainLoop {
                       aligned_actual_kv_tile_pos_left,
                       aligned_actual_kv_tile_pos_right, actual_kv_token_num,
                       kv_cache_block_num_stride, q_tile_head_num,
-                      q_tile_token_num, q_tile_pos_left, curr_q_heads_per_kv,
+                      q_tile_token_num, q_tile_pos_left, actual_q_heads_per_kv,
                       block_size, sliding_window_left, sliding_window_right,
                       scale, softcap_scale, curr_alibi_slopes,
                       first_iter_flag[q_iter_idx], use_sink, debug_info);
@@ -1807,11 +1779,11 @@ class AttentionMainLoop {
                   final_output(partial_q_buffer,
                                reinterpret_cast<query_t*>(input->output) +
                                    output_buffer_offset,
-                               sum_buffer, curr_q_heads_per_kv,
+                               sum_buffer, actual_q_heads_per_kv,
                                actual_q_token_num, q_head_num, output_v_scale);
                 } else {
                   const int32_t stride =
-                      curr_q_heads_per_kv * split_kv_q_token_num_threshold;
+                      actual_q_heads_per_kv * split_kv_q_token_num_threshold;
                   buffer_manager.update(kv_head_idx, total_reduction_split_num,
                                         head_dim, stride, sizeof(float));
                   volatile bool* split_flag_buffer =
@@ -1847,26 +1819,18 @@ class AttentionMainLoop {
           const int32_t curr_split_id = curr_workitem_groups->split_start_id;
           const int32_t curr_split_num = curr_workitem_groups->split_num;
           const int32_t current_group_idx = curr_workitem_groups->req_id;
-          const bool curr_use_gqa =
-              use_gqa_fast_path || (supports_gqa && curr_output_token_num == 1);
-          if (!use_gqa_fast_path && curr_use_gqa &&
-              kv_head_idx % q_heads_per_kv != 0) {
-            continue;
-          }
-          const int32_t curr_q_heads_per_kv = curr_use_gqa ? q_heads_per_kv : 1;
           const int32_t curr_output_head_num =
-              curr_output_token_num * curr_q_heads_per_kv;
+              curr_output_token_num * actual_q_heads_per_kv;
 
           const int32_t q_start = input->query_start_loc[current_group_idx];
           const int32_t q_token_start_idx = q_start + curr_output_token_idx;
-          const int32_t q_head_start_idx =
-              use_gqa_fast_path ? (kv_head_idx * q_heads_per_kv) : kv_head_idx;
+          const int32_t q_head_start_idx = kv_head_idx * actual_q_heads_per_kv;
           size_t output_buffer_offset =
               q_token_start_idx * q_head_num * head_dim +
               q_head_start_idx * head_dim;
 
           const int32_t stride =
-              curr_q_heads_per_kv * split_kv_q_token_num_threshold;
+              actual_q_heads_per_kv * split_kv_q_token_num_threshold;
           buffer_manager.update(kv_head_idx, total_reduction_split_num,
                                 head_dim, stride, sizeof(float));
           volatile bool* split_flag_buffer =
@@ -1885,7 +1849,7 @@ class AttentionMainLoop {
           final_output(
               split_output_buffer,
               reinterpret_cast<query_t*>(input->output) + output_buffer_offset,
-              split_sum_buffer, curr_q_heads_per_kv, curr_output_token_num,
+              split_sum_buffer, actual_q_heads_per_kv, curr_output_token_num,
               q_head_num, output_v_scale);
         }
       }

csrc/cpu/cpu_attn_rvv.hpp

@@ -1,412 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-#ifndef CPU_ATTN_RVV_HPP
-#define CPU_ATTN_RVV_HPP
-
-// RVV attention kernel using VLEN-agnostic RVVI() macros from
-// cpu_types_riscv_defs.hpp.  The Mx8 tile GEMM uses 8 FP32 elements
-// per vector (LMUL_256 bits of FP32 data), which maps to:
-//   VLEN=128: m2 (256 bits = 8 x FP32)
-//   VLEN=256: m1 (256 bits = 8 x FP32)
-// Only VLEN=128 and VLEN=256 are supported; other VLENs (512, 1024)
-// and scalar RISC-V builds fall back to VEC/VEC16.
-#if defined(__riscv_v_min_vlen) && \
-    (__riscv_v_min_vlen == 128 || __riscv_v_min_vlen == 256)
-
-  #include "cpu_attn_impl.hpp"
-  #include "cpu_types_riscv_defs.hpp"
-  #include <riscv_vector.h>
-  #include <type_traits>
-
-namespace cpu_attention {
-
-namespace {
-
-  #define BLOCK_SIZE_ALIGNMENT 32
-  #define HEAD_SIZE_ALIGNMENT 32
-  #define MAX_Q_HEAD_NUM_PER_ITER 16
-
-// ============================================================================
-// B-matrix row loading: load 8 elements as FP32
-// ============================================================================
-
-template <typename kv_cache_t>
-FORCE_INLINE fixed_fp32x8_t load_row8_B_as_f32(const kv_cache_t* p);
-
-template <>
-FORCE_INLINE fixed_fp32x8_t load_row8_B_as_f32<float>(const float* p) {
-  return RVVI(__riscv_vle32_v_f32, LMUL_256)(p, 8);
-}
-
-template <>
-FORCE_INLINE fixed_fp32x8_t load_row8_B_as_f32<c10::Half>(const c10::Half* p) {
-  #ifdef __riscv_zvfh
-  fixed_fp16x8_t h = RVVI(__riscv_vle16_v_f16, LMUL_128)(
-      reinterpret_cast<const _Float16*>(p), 8);
-  return RVVI(__riscv_vfwcvt_f_f_v_f32, LMUL_256)(h, 8);
-  #else
-  alignas(16) float tmp[8];
-  for (int i = 0; i < 8; ++i) {
-    tmp[i] = static_cast<float>(p[i]);
-  }
-  return RVVI(__riscv_vle32_v_f32, LMUL_256)(tmp, 8);
-  #endif
-}
-
-template <>
-FORCE_INLINE fixed_fp32x8_t
-load_row8_B_as_f32<c10::BFloat16>(const c10::BFloat16* p) {
-  #ifdef __riscv_zvfbfmin
-  fixed_bf16x8_t bf = RVVI(__riscv_vle16_v_bf16, LMUL_128)(
-      reinterpret_cast<const __bf16*>(p), 8);
-  return RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_256)(bf, 8);
-  #else
-  fixed_u16x8_t raw = RVVI(__riscv_vle16_v_u16, LMUL_128)(
-      reinterpret_cast<const uint16_t*>(p), 8);
-  fixed_u32x8_t wide = RVVI(__riscv_vzext_vf2_u32, LMUL_256)(raw, 8);
-  fixed_u32x8_t shifted = RVVI(__riscv_vsll_vx_u32, LMUL_256)(wide, 16, 8);
-  return RVVI4(__riscv_vreinterpret_v_u32, LMUL_256, _f32, LMUL_256)(shifted);
-  #endif
-}
-
-// ============================================================================
-// Micro kernel: Mx8 tile, K unrolled by 4, RVV scalar-broadcast FMA
-// ============================================================================
-//
-// RVV has no lane-indexed FMA; instead we load A elements as scalars and
-// use vfmacc_vf (scalar * vector + accumulator).
-//
-// The 8-column tile uses LMUL_256 bits of FP32 data:
-//   VLEN=128: m2 (2 regs per accumulator), M=8 => 18 of 32 regs
-//   VLEN=256: m1 (1 reg  per accumulator), M=8 =>  9 of 32 regs
-
-template <int32_t M, typename kv_cache_t>
-FORCE_INLINE void gemm_micro_rvv_fma_Mx8_Ku4(
-    const float* __restrict A,       // [M x K]
-    const kv_cache_t* __restrict B,  // [K x 8]
-    float* __restrict C,             // [M x 8]
-    int64_t lda, int64_t ldb, int64_t ldc, int32_t K, bool accumulate) {
-  static_assert(1 <= M && M <= 8, "M must be in [1,8]");
-
-  constexpr size_t vl = 8;
-
-  #define ROWS_APPLY(OP) OP(0) OP(1) OP(2) OP(3) OP(4) OP(5) OP(6) OP(7)
-  #define IF_M(i) if constexpr (M > (i))
-
-  #define DECL_A(i) const float* a##i = A + (i) * lda;
-  ROWS_APPLY(DECL_A)
-  #undef DECL_A
-
-  #define DECL_ACC(i) fixed_fp32x8_t acc##i;
-  ROWS_APPLY(DECL_ACC)
-  #undef DECL_ACC
-
-  #define INIT_ACC(i)                                                    \
-    IF_M(i) {                                                            \
-      if (accumulate) {                                                  \
-        acc##i = RVVI(__riscv_vle32_v_f32, LMUL_256)(C + (i) * ldc, vl); \
-      } else {                                                           \
-        acc##i = RVVI(__riscv_vfmv_v_f_f32, LMUL_256)(0.f, vl);          \
-      }                                                                  \
-    }
-  ROWS_APPLY(INIT_ACC)
-  #undef INIT_ACC
-
-  int32_t k = 0;
-
-  for (; k + 3 < K; k += 4) {
-    {
-      fixed_fp32x8_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 0) * ldb);
-  #define STEP_K0(i)                                                          \
-    IF_M(i) {                                                                 \
-      acc##i = RVVI(__riscv_vfmacc_vf_f32, LMUL_256)(acc##i, *(a##i + k + 0), \
-                                                     b, vl);                  \
-    }
-      ROWS_APPLY(STEP_K0)
-  #undef STEP_K0
-    }
-    {
-      fixed_fp32x8_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 1) * ldb);
-  #define STEP_K1(i)                                                          \
-    IF_M(i) {                                                                 \
-      acc##i = RVVI(__riscv_vfmacc_vf_f32, LMUL_256)(acc##i, *(a##i + k + 1), \
-                                                     b, vl);                  \
-    }
-      ROWS_APPLY(STEP_K1)
-  #undef STEP_K1
-    }
-    {
-      fixed_fp32x8_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 2) * ldb);
-  #define STEP_K2(i)                                                          \
-    IF_M(i) {                                                                 \
-      acc##i = RVVI(__riscv_vfmacc_vf_f32, LMUL_256)(acc##i, *(a##i + k + 2), \
-                                                     b, vl);                  \
-    }
-      ROWS_APPLY(STEP_K2)
-  #undef STEP_K2
-    }
-    {
-      fixed_fp32x8_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 3) * ldb);
-  #define STEP_K3(i)                                                          \
-    IF_M(i) {                                                                 \
-      acc##i = RVVI(__riscv_vfmacc_vf_f32, LMUL_256)(acc##i, *(a##i + k + 3), \
-                                                     b, vl);                  \
-    }
-      ROWS_APPLY(STEP_K3)
-  #undef STEP_K3
-    }
-  }
-
-  for (; k < K; ++k) {
-    fixed_fp32x8_t b = load_row8_B_as_f32<kv_cache_t>(B + (int64_t)k * ldb);
-  #define TAIL_ROW(i)                                                        \
-    IF_M(i) {                                                                \
-      acc##i =                                                               \
-          RVVI(__riscv_vfmacc_vf_f32, LMUL_256)(acc##i, *(a##i + k), b, vl); \
-    }
-    ROWS_APPLY(TAIL_ROW)
-  #undef TAIL_ROW
-  }
-
-  #define STORE_ROW(i) \
-    IF_M(i) { RVVI(__riscv_vse32_v_f32, LMUL_256)(C + (i) * ldc, acc##i, vl); }
-  ROWS_APPLY(STORE_ROW)
-  #undef STORE_ROW
-
-  #undef ROWS_APPLY
-  #undef IF_M
-}
-
-// ============================================================================
-// Macro kernel: dispatch M tiles of {8,4,2,1}, step N by 8
-// ============================================================================
-
-template <int32_t N, typename kv_cache_t>
-FORCE_INLINE void gemm_macro_rvv_fma_Mx8_Ku4(const float* __restrict A,
-                                             const kv_cache_t* __restrict B,
-                                             float* __restrict C, int32_t M,
-                                             int32_t K, int64_t lda,
-                                             int64_t ldb, int64_t ldc,
-                                             bool accumulate) {
-  static_assert(N % 8 == 0, "N must be a multiple of 8");
-  for (int32_t m = 0; m < M;) {
-    int32_t mb = (M - m >= 8) ? 8 : (M - m >= 4) ? 4 : (M - m >= 2) ? 2 : 1;
-    const float* Ab = A + m * lda;
-    float* Cb = C + m * ldc;
-
-    for (int32_t n = 0; n < N; n += 8) {
-      const kv_cache_t* Bn = B + n;
-      float* Cn = Cb + n;
-      switch (mb) {
-        case 8:
-          gemm_micro_rvv_fma_Mx8_Ku4<8, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        case 4:
-          gemm_micro_rvv_fma_Mx8_Ku4<4, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        case 2:
-          gemm_micro_rvv_fma_Mx8_Ku4<2, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        default:
-          gemm_micro_rvv_fma_Mx8_Ku4<1, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-      }
-    }
-    m += mb;
-  }
-}
-
-// ============================================================================
-// TileGemm wrapper — plugs into AttentionMainLoop
-// ============================================================================
-
-template <typename kv_cache_t>
-class TileGemmRVV {
- public:
-  template <AttentionGemmPhase phase, int32_t k_size>
-  FORCE_INLINE static void gemm(const int32_t m_size,
-                                float* __restrict__ a_tile,
-                                kv_cache_t* __restrict__ b_tile,
-                                float* __restrict__ c_tile, const int64_t lda,
-                                const int64_t ldb, const int64_t ldc,
-                                const int32_t block_size,
-                                const int32_t dynamic_k_size,
-                                const bool accum_c) {
-    if constexpr (phase == AttentionGemmPhase::QK) {
-      gemm_macro_rvv_fma_Mx8_Ku4<BLOCK_SIZE_ALIGNMENT, kv_cache_t>(
-          a_tile, b_tile, c_tile, m_size, k_size, lda, ldb, ldc, accum_c);
-    } else {
-      gemm_macro_rvv_fma_Mx8_Ku4<HEAD_SIZE_ALIGNMENT, kv_cache_t>(
-          a_tile, b_tile, c_tile, m_size, dynamic_k_size, lda, ldb, ldc,
-          accum_c);
-    }
-  }
-};
-
-}  // namespace
-
-// ============================================================================
-// AttentionImpl<ISA::RVV> — mirrors ISA::NEON specialization
-// ============================================================================
-
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::RVV, scalar_t, head_dim, kv_cache_scalar_t> {
- public:
-  using query_t = scalar_t;
-  using q_buffer_t = float;
-  using kv_cache_t = scalar_t;
-  using logits_buffer_t = float;
-  using partial_output_buffer_t = float;
-  using prob_buffer_t = float;
-
-  constexpr static int64_t BlockSizeAlignment = BLOCK_SIZE_ALIGNMENT;
-  constexpr static int64_t HeadDimAlignment = HEAD_SIZE_ALIGNMENT;
-  constexpr static int64_t MaxQHeadNumPerIteration = MAX_Q_HEAD_NUM_PER_ITER;
-  constexpr static int64_t HeadDim = head_dim;
-  constexpr static ISA ISAType = ISA::RVV;
-  constexpr static bool scale_on_logits = false;
-
-  static_assert(HeadDim % HeadDimAlignment == 0);
-  static_assert(HeadDimAlignment % 8 == 0);
-  static_assert(BlockSizeAlignment % 8 == 0);
-
- public:
-  template <template <typename tile_gemm_t> typename attention>
-  FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
-    attention<TileGemmRVV<kv_cache_t>> attention_iteration;
-    attention_iteration(CPU_ATTENTION_PARAMS);
-  }
-
-  constexpr static int64_t k_cache_token_group_stride(
-      const int32_t block_size) {
-    return BlockSizeAlignment;
-  }
-
-  constexpr static int64_t v_cache_token_group_stride(
-      const int32_t block_size) {
-    return head_dim * BlockSizeAlignment;
-  }
-
-  constexpr static int64_t v_cache_head_group_stride(const int32_t block_size) {
-    return HeadDimAlignment;
-  }
-
-  static 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) {
-    static_assert(head_dim % 16 == 0);
-    constexpr int32_t unroll_size = head_dim / 16;
-    using load_vec_t = typename VecTypeTrait<scalar_t>::vec_t;
-
-    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 =
-            src + q_num_idx * q_num_stride + q_head_idx * q_head_stride;
-        float* __restrict__ curr_q_buffer =
-            q_buffer + q_num_idx * q_heads_per_kv * head_dim +
-            q_head_idx * head_dim;
-
-        vec_op::unroll_loop<int32_t, unroll_size>([&](int32_t i) {
-          load_vec_t vec(curr_q);
-          vec_op::FP32Vec16 fp32_vec(vec);
-          fp32_vec = fp32_vec * scale_vec;
-          fp32_vec.save(curr_q_buffer);
-
-          curr_q += 16;
-          curr_q_buffer += 16;
-        });
-      }
-    }
-  }
-
-  static void reshape_and_cache(
-      const scalar_t* __restrict__ key, const scalar_t* __restrict__ value,
-      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) {
-  #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) {
-        const int64_t pos = slot_mapping[token_idx];
-        if (pos < 0) {
-          continue;
-        }
-
-        const int64_t block_idx = pos / block_size;
-        const int64_t block_offset = pos % block_size;
-        {
-          const scalar_t* key_start_ptr = key +
-                                          token_idx * key_token_num_stride +
-                                          head_idx * key_head_num_stride;
-          scalar_t* key_cache_start_ptr =
-              key_cache + block_idx * num_blocks_stride +
-              head_idx * cache_head_num_stride + block_offset;
-
-          {
-            const ptrdiff_t byte_stride = block_size * sizeof(scalar_t);
-            int64_t i = 0;
-            for (; i < head_dim;) {
-              size_t vl;
-              if constexpr (std::is_same_v<scalar_t, float>) {
-                vl = __riscv_vsetvl_e32m2(head_dim - i);
-                vfloat32m2_t v = __riscv_vle32_v_f32m2(
-                    reinterpret_cast<const float*>(key_start_ptr + i), vl);
-                __riscv_vsse32_v_f32m2(
-                    reinterpret_cast<float*>(key_cache_start_ptr +
-                                             i * block_size),
-                    byte_stride, v, vl);
-              } else {
-                vl = __riscv_vsetvl_e16m1(head_dim - i);
-                vuint16m1_t v = __riscv_vle16_v_u16m1(
-                    reinterpret_cast<const uint16_t*>(key_start_ptr + i), vl);
-                __riscv_vsse16_v_u16m1(
-                    reinterpret_cast<uint16_t*>(key_cache_start_ptr +
-                                                i * block_size),
-                    byte_stride, v, vl);
-              }
-              i += vl;
-            }
-          }
-        }
-        {
-          const scalar_t* value_start_ptr = value +
-                                            token_idx * value_token_num_stride +
-                                            head_idx * value_head_num_stride;
-          scalar_t* value_cache_start_ptr =
-              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);
-        }
-      }
-    }
-  }
-};
-
-}  // namespace cpu_attention
-
-  #undef BLOCK_SIZE_ALIGNMENT
-  #undef HEAD_SIZE_ALIGNMENT
-  #undef MAX_Q_HEAD_NUM_PER_ITER
-
-#endif  // __riscv_v_min_vlen == 128 || 256
-
-#endif  // CPU_ATTN_RVV_HPP

csrc/cpu/cpu_types_riscv_defs.hpp

@@ -71,12 +71,8 @@ typedef RVVTYPE(vuint16, LMUL_256, _t) fixed_u16x16_t
 typedef RVVTYPE(vuint16, LMUL_512, _t) fixed_u16x32_t
     __attribute__((riscv_rvv_vector_bits(512)));
 
-// uint32
-typedef RVVTYPE(vuint32, LMUL_256, _t) fixed_u32x8_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-
 // bfloat16
-#ifdef __riscv_zvfbfmin
+#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

csrc/cpu/cpu_types_riscv_impl.hpp

@@ -15,17 +15,8 @@
 #include <torch/all.h>
 namespace vec_op {
 
-// FP8 KV cache is not supported on RISC-V. These tag types and the
-// corresponding BF16Vec32 stub constructors below exist solely so that
-// templates referencing vec_op::fp8_*_tag in their bodies (e.g. in
-// cpu_attn_vec.hpp) compile under GCC's -Wtemplate-body lookup. The
-// stubs are never instantiated by CPU_ATTN_DISPATCH on __riscv.
-struct fp8_e4m3_tag {};
-struct fp8_e5m2_tag {};
-
-// BFloat16 is always supported on RISC-V: natively when __riscv_zvfbfmin
-// is defined (compiler-provided when -march includes zvfbfmin), otherwise
-// via the FP32-simulation fallback path.
+// 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__)  \
@@ -115,7 +106,7 @@ struct FP16Vec16 : public Vec<FP16Vec16> {
 // BF16 Implementation
 // ============================================================================
 
-#ifdef __riscv_zvfbfmin
+#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);
@@ -194,13 +185,6 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
 
   explicit BF16Vec32(fixed_bf16x32_t data) : reg(data) {};
 
-  // FP8 KV cache stubs: never instantiated on RISC-V (CPU_ATTN_DISPATCH
-  // omits FP8 cases on __riscv); exist only so name lookup succeeds.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e4m3_tag)
-      : BF16Vec32(static_cast<const void*>(ptr)) {}
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e5m2_tag)
-      : BF16Vec32(static_cast<const void*>(ptr)) {}
-
   explicit BF16Vec32(const BF16Vec8& v) {
     fixed_u16x8_t u16_val = bf16_to_u16(v.reg);
     fixed_u16x32_t u16_combined =
@@ -339,13 +323,6 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
     reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_1024)(tmp, 32);
   }
 
-  // FP8 KV cache stubs: never instantiated on RISC-V (CPU_ATTN_DISPATCH
-  // omits FP8 cases on __riscv); exist only so name lookup succeeds.
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e4m3_tag)
-      : BF16Vec32(static_cast<const void*>(ptr)) {}
-  explicit BF16Vec32(const uint8_t* ptr, fp8_e5m2_tag)
-      : BF16Vec32(static_cast<const void*>(ptr)) {}
-
   explicit BF16Vec32(const BF16Vec8& v) {
     float tmp_small[8];
     RVVI(__riscv_vse32_v_f32, LMUL_256)(tmp_small, v.reg_fp32, 8);
@@ -433,7 +410,7 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
   explicit FP32Vec8(fixed_fp16x8_t v)
       : reg(RVVI(__riscv_vfwcvt_f_f_v_f32, LMUL_256)(v, VEC_ELEM_NUM)) {};
 
-#ifdef __riscv_zvfbfmin
+#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)
@@ -631,7 +608,7 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
   explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
   explicit FP32Vec16(const FP16Vec16& v);
 
-#ifdef __riscv_zvfbfmin
+#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)
@@ -891,7 +868,7 @@ inline void fma(FP32Vec16& acc, const FP32Vec16& a, const FP32Vec16& b) {
   acc = acc.fma(a, b);
 }
 
-#ifdef __riscv_zvfbfmin
+#ifdef RISCV_BF16_SUPPORT
 template <>
 inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
   *ptr = static_cast<__bf16>(v);

csrc/cpu/generate_cpu_attn_dispatch.py

@@ -20,8 +20,7 @@ ISA_TYPES = {
     "VEC16": 2,
     "NEON": 3,
     "VXE": 4,
-    "RVV": 5,
-    "VSX": 6,
+    "VSX": 5,
 }
 
 # KV cache index: 0 = auto (same as scalar_t), 1 = fp8_e4m3, 2 = fp8_e5m2
@@ -39,7 +38,7 @@ KV_CACHE_CPP_TYPES = {
 }
 
 # ISAs supported for head_dims divisible by 32
-ISA_FOR_32 = ["AMX", "NEON", "VEC", "VEC16", "VXE", "RVV", "VSX"]
+ISA_FOR_32 = ["AMX", "NEON", "VEC", "VEC16", "VXE", "VSX"]
 
 # ISAs supported for head_dims divisible by 16 only
 ISA_FOR_16 = ["VEC16"]
@@ -150,13 +149,6 @@ def generate_header_file() -> str:
   #include "cpu_attn_vxe.hpp"
 #endif
 
-// cpu_attn_rvv.hpp supports VLEN=128 and VLEN=256 via RVVI() macros.
-// Other VLENs and scalar RISC-V builds skip it entirely.
-#if defined(__riscv) && defined(__riscv_v_min_vlen) && \
-    (__riscv_v_min_vlen == 128 || __riscv_v_min_vlen == 256)
-  #include "cpu_attn_rvv.hpp"
-#endif
-
 #ifdef __powerpc__
   #include "cpu_attn_vsx.hpp"
 #endif
@@ -220,20 +212,6 @@ def generate_header_file() -> str:
         ["VXE", "VEC", "VEC16"],
         fp8=False,
     )
-    # RISC-V with RVV.  cpu_attn_rvv.hpp supports VLEN=128 and VLEN=256
-    # via RVVI() macros.  Builds with a supported VLEN get
-    # RVV+VEC+VEC16; other RISC-V builds fall back to VEC/VEC16 only.
-    header += _macro_block(
-        "#elif defined(__riscv) && defined(__riscv_v_min_vlen) "
-        "&& (__riscv_v_min_vlen == 128 || __riscv_v_min_vlen == 256)",
-        ["RVV", "VEC", "VEC16"],
-        fp8=False,
-    )
-    header += _macro_block(
-        "#elif defined(__riscv)",
-        ["VEC", "VEC16"],
-        fp8=False,
-    )
     header += _macro_block(
         "#elif defined(__powerpc__)",
         ["VSX", "VEC", "VEC16"],
@@ -255,8 +233,8 @@ def generate_header_file() -> str:
         fp8=False,
     )
     header += (
-        "#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / __s390x__ /"
-        " __riscv / __powerpc__ */\n\n"
+        "#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / "
+        "__s390x__ / __powerpc__ */\n\n"
         "#endif  // CPU_ATTN_DISPATCH_GENERATED_H\n"
     )
 

csrc/cpu/sgl-kernels/common.h

@@ -348,13 +348,6 @@ inline int get_cache_blocks<at::Float8_e4m3fn>(int chunk_size) {
   return std::min(MAX_CACHE_BLOCK_SIZE, cache_block_size);
 }
 
-template <>
-inline int get_cache_blocks<uint8_t>(int chunk_size) {
-  // mxfp4 uses bf16 as accumulate type
-  int cache_block_size = get_cache_blocks<at::BFloat16>(chunk_size);
-  return std::min(MAX_CACHE_BLOCK_SIZE, cache_block_size);
-}
-
 // 2d sequential loop in range : [mb0, mb1), [nb0, nb1)
 template <typename T, typename func_t>
 inline void loop_2d(int64_t mb0, int64_t mb1, int64_t nb0, int64_t nb1, int64_t chunk_size, const func_t& f) {

csrc/cpu/sgl-kernels/conv.cpp

@@ -213,7 +213,7 @@ struct tinygemm_kernel<at::BFloat16, K, BLOCK_N, has_bias, has_silu> {
       weight + nb_start * width,                                                             \
       out + bs * seqlen * dim + mb_start * dim + nb_start,                                   \
       has_bias ? bias + nb_start : nullptr,                                                  \
-      has_conv_states ? conv_states + conv_state_index * conv_state_slot_stride + nb_start : nullptr, \
+      has_conv_states ? conv_states + conv_state_index * (K - 1) * dim + nb_start : nullptr, \
       has_initial_states_value,                                                              \
       mb_size,                                                                               \
       dim,                                                                                   \
@@ -233,8 +233,7 @@ void causal_conv1d_fwd_kernel_impl(
     int64_t dim,
     int64_t seqlen,
     int64_t width,
-    int64_t num_seq_blocks,
-    int64_t conv_state_slot_stride) {
+    int64_t num_seq_blocks) {
   // handle 32 x 64 per block
   constexpr int64_t BLOCK_M = block_size_m();
   constexpr int64_t BLOCK_N = block_size_n() * 2;
@@ -283,7 +282,7 @@ void causal_conv1d_fwd_kernel_impl(
     at::parallel_for(0, batch, 0, [&](int64_t begin, int64_t end) {
       for (int64_t bs = begin; bs < end; ++bs) {
         update_conv_state(
-            conv_states + bs * conv_state_slot_stride, input + bs * seqlen * dim, width, dim, seqlen, has_initial_state[bs]);
+            conv_states + bs * (width - 1) * dim, input + bs * seqlen * dim, width, dim, seqlen, has_initial_state[bs]);
       }
     });
   }
@@ -317,8 +316,7 @@ void causal_conv1d_fwd_varlen_kernel_impl(
     int64_t batch,
     int64_t dim,
     int64_t width,
-    int64_t num_seq_blocks,
-    int64_t conv_state_slot_stride) {
+    int64_t num_seq_blocks) {
   // handle 32 x 64 per block
   constexpr int64_t BLOCK_M = block_size_m();
   constexpr int64_t BLOCK_N = block_size_n() * 2;
@@ -368,7 +366,7 @@ void causal_conv1d_fwd_varlen_kernel_impl(
         int32_t seqlen = query_start_loc[bs + 1] - query_start_loc[bs];
         int32_t batch_offset = query_start_loc[bs];
         update_conv_state(
-            conv_states + conv_state_index * conv_state_slot_stride,
+            conv_states + conv_state_index * (width - 1) * dim,
             input + batch_offset * dim,
             width,
             dim,
@@ -391,8 +389,7 @@ void causal_conv1d_update_kernel_impl(
     int64_t batch,
     int64_t dim,
     int64_t seqlen,
-    int64_t width,
-    int64_t conv_state_slot_stride) {
+    int64_t width) {
   // handle 32 x 64 per block
   constexpr int64_t BLOCK_M = block_size_m();
   constexpr int64_t BLOCK_N = block_size_n() * 2;
@@ -433,7 +430,7 @@ void causal_conv1d_update_kernel_impl(
     });
   });
 
-#define CONV_STATE_INDEXR(w) conv_states + conv_state_index*conv_state_slot_stride + (w) * dim
+#define CONV_STATE_INDEXR(w) conv_states + conv_state_index*(width - 1) * dim + (w) * dim
 
   // update conv_states
   at::parallel_for(0, batch, 0, [&](int64_t begin, int64_t end) {
@@ -595,9 +592,6 @@ at::Tensor causal_conv1d_fwd_cpu(
     }
   }
 
-// IMPORTANT: To make the kernal compatible with vLLM KV cache layout 
-  int64_t conv_state_slot_stride = conv_states->stride(0);
-
   // block size for sequence blocks, 32
   constexpr int64_t BLOCK_M = block_size_m();
 
@@ -624,8 +618,7 @@ at::Tensor causal_conv1d_fwd_cpu(
           batch,
           dim,
           width,
-          num_seq_blocks,
-          conv_state_slot_stride);
+          num_seq_blocks);
     } else {
       causal_conv1d_fwd_kernel_impl<scalar_t>(
           out.data_ptr<scalar_t>(),
@@ -640,8 +633,7 @@ at::Tensor causal_conv1d_fwd_cpu(
           dim,
           seqlen,
           width,
-          num_seq_blocks,
-          conv_state_slot_stride);
+          num_seq_blocks);
     }
   });
   return out;
@@ -698,8 +690,6 @@ at::Tensor causal_conv1d_update_cpu(
     conv_states.copy_(conv_states_copy);
   }
 
-  // IMPORTANT: To make the kernal compatible with vLLM KV cache layout 
-  int64_t conv_state_slot_stride = conv_states.stride(0);
   at::Tensor out = at::empty_like(x);
   AT_DISPATCH_REDUCED_FLOATING_TYPES(scalar_type, "causal_conv1d_update_kernel_impl", [&] {
     causal_conv1d_update_kernel_impl<scalar_t>(
@@ -713,8 +703,7 @@ at::Tensor causal_conv1d_update_cpu(
         batch,
         dim,
         seqlen,
-        width,
-        conv_state_slot_stride);
+        width);
   });
   return out;
 }

csrc/cpu/sgl-kernels/fla.cpp

@@ -847,7 +847,6 @@ void fused_sigmoid_gating_delta_rule_update_kernel_impl(
     int64_t v_strideB,
     int64_t v_strideS,
     int64_t v_strideH,
-    int64_t state_slot_stride,
     bool use_qk_l2norm_in_kernel,
     double softplus_threshold) {
   using bVec = at::vec::Vectorized<scalar_t>;
@@ -908,7 +907,7 @@ void fused_sigmoid_gating_delta_rule_update_kernel_impl(
     data_index_init(begin, bi, batch_size, si, seq_len, ni, v_num_heads);
     for (int64_t i = begin; i < end; ++i) {
       int64_t cache_index = indices_ptr[bi];
-      int64_t state_offset = cache_index * state_slot_stride + ni * head_dim * v_head_dim;
+      int64_t state_offset = (cache_index * v_num_heads + ni) * head_dim * v_head_dim;
       float g_val = -std::exp(float(A_log_ptr[ni])) *
                     softplus(float(a_ptr[bi * v_num_heads + ni]) + float(dt_bias_ptr[ni]), softplus_threshold);
       float g_val_exp = std::exp(g_val);
@@ -1322,8 +1321,6 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
   int64_t v_strideB = v.stride(1);
   int64_t v_strideS = v.stride(0);
   int64_t v_strideH = v.stride(2);
-  // IMPORTANT: To make the kernal compatible with vLLM KV cache layout 
-  int64_t state_slot_stride = initial_state_source.stride(0);
   at::Tensor core_attn_out = at::empty({batch_size, seq_len, v_num_heads, v_head_dim}, q.options());
   at::Tensor qk_scale_buf = at::empty({2 * batch_size, seq_len, num_heads}, at::kFloat);
 
@@ -1356,7 +1353,6 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
             v_strideB,
             v_strideS,
             v_strideH,
-            state_slot_stride,
             use_qk_l2norm_in_kernel,
             softplus_threshold);
       });

csrc/cpu/sgl-kernels/gemm.h

@@ -72,17 +72,7 @@ inline int64_t get_row_size(int64_t K, bool use_int8_w8a8) {
   return use_int8_w8a8 ? K + sizeof(int32_t) : K;
 }
 
-enum class CPUAcTMethod : int { silu_and_mul = 0, swiglu = 1 };
-
-constexpr bool operator==(CPUAcTMethod a, int b) {
-  return static_cast<int>(a) == b;
-}
-
-constexpr bool operator==(int a, CPUAcTMethod b) {
-  return a == static_cast<int>(b);
-}
-
-enum class CPUQuantMethod : int64_t { BF16 = 0, INT8_W8A8 = 1, FP8_W8A16 = 2, INT4_W4A8 = 3, MXFP4 = 4 };
+enum class CPUQuantMethod : int64_t { BF16 = 0, INT8_W8A8 = 1, FP8_W8A16 = 2, INT4_W4A8 = 3 };
 
 constexpr bool operator==(CPUQuantMethod a, int64_t b) {
   return static_cast<int64_t>(a) == b;
@@ -109,9 +99,6 @@ inline int64_t get_4bit_block_k_size(int64_t group_size) {
 // pack weight to vnni format
 at::Tensor convert_weight_packed(at::Tensor& weight);
 
-// pack scale to blocked format for mxfp4
-at::Tensor convert_scale_packed(at::Tensor& scale);
-
 // pack weight to vnni format for int4
 std::tuple<at::Tensor, at::Tensor, at::Tensor>
 convert_weight_packed_scale_zp(at::Tensor qweight, at::Tensor qzeros, at::Tensor scales);
@@ -142,9 +129,9 @@ void fused_experts_int8_kernel_impl(
     int64_t topk,
     int64_t num_tokens_post_pad);
 
-// moe implementations for fp8 w8a16 and mxfp4
-template <typename scalar_t, typename packed_t, typename param_t, bool is_mxfp4>
-void fused_experts_fp_kernel_impl(
+// moe implementations for fp8 w8a16
+template <typename scalar_t>
+void fused_experts_fp8_kernel_impl(
     scalar_t* __restrict__ output,
     scalar_t* __restrict__ ic0,
     scalar_t* __restrict__ ic1,
@@ -153,12 +140,10 @@ void fused_experts_fp_kernel_impl(
     scalar_t* __restrict__ B_tmp,
     float* __restrict__ C_tmp,
     const scalar_t* __restrict__ input,
-    const packed_t* __restrict__ packed_w1,
-    const packed_t* __restrict__ packed_w2,
-    const float* __restrict__ w1_bias,
-    const float* __restrict__ w2_bias,
-    const param_t* __restrict__ w1s,
-    const param_t* __restrict__ w2s,
+    const at::Float8_e4m3fn* __restrict__ packed_w1,
+    const at::Float8_e4m3fn* __restrict__ packed_w2,
+    const float* __restrict__ w1s,
+    const float* __restrict__ w2s,
     int64_t block_size_N,
     int64_t block_size_K,
     const float* __restrict__ topk_weights,
@@ -170,11 +155,7 @@ void fused_experts_fp_kernel_impl(
     int64_t K,
     int64_t E,
     int64_t topk,
-    int64_t num_tokens_post_pad,
-    float alpha,
-    float limit,
-    CPUAcTMethod act_func,
-    bool with_bias);
+    int64_t num_tokens_post_pad);
 
 // shared expert implementation for int8 w8a8
 template <typename scalar_t>
@@ -285,7 +266,6 @@ void tinygemm_kernel(
     scalar_t* __restrict__ C,
     scalar_t* __restrict__ Btmp,
     float* __restrict__ Ctmp,
-    const float* __restrict__ Bbias,
     const float* __restrict__ scale,
     int64_t M,
     int64_t N,
@@ -314,26 +294,6 @@ void tinygemm_kernel(
     int64_t ldc,
     bool brg);
 
-// mxfp4
-template <typename scalar_t>
-void tinygemm_kernel(
-    const scalar_t* __restrict__ A,
-    const uint8_t* __restrict__ B,
-    scalar_t* __restrict__ C,
-    scalar_t* __restrict__ Btmp,
-    float* __restrict__ Ctmp,
-    const float* __restrict__ Bbias,
-    const uint8_t* __restrict__ scale,
-    int64_t M,
-    int64_t N,
-    int64_t K,
-    int64_t lda,
-    int64_t ldb,
-    int64_t ldc,
-    bool brg,
-    int64_t block_size_K,
-    bool do_unpack = true);
-
 template <typename scalar_t>
 void tinygemm_kernel(
     scalar_t* C,

csrc/cpu/sgl-kernels/gemm_fp8.cpp

@@ -67,23 +67,6 @@ inline void copy_mul_stub(scalar_t* __restrict__ out, const float* __restrict__
   }
 }
 
-template <>
-inline void
-copy_add_stub(float* __restrict__ out, const float* __restrict__ input, const float* __restrict__ bias, int64_t size) {
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = fVec::size();
-
-  int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    fVec data = fVec::loadu(input + d) + fVec::loadu(bias + d);
-    data.store(out + d);
-  }
-  for (; d < size; ++d) {
-    out[d] = input[d] + bias[d];
-  }
-}
-
 inline void unpack_B(
     at::BFloat16* __restrict__ Btmp,
     const at::Float8_e4m3fn* __restrict__ packed_B,
@@ -360,6 +343,7 @@ struct tinygemm_kernel_nn<at::BFloat16, at::Float8_e4m3fn, float, has_bias, BLOC
     Unroll<ROWS * COLS>{}(storec);
   }
 };
+
 template <int BLOCK_M, int BLOCK_N>
 struct tinygemm_kernel_nn2<at::BFloat16, BLOCK_M, BLOCK_N> {
   static inline void apply(
@@ -934,7 +918,6 @@ void tinygemm_kernel(
     scalar_t* __restrict__ C,
     scalar_t* __restrict__ Btmp,
     float* __restrict__ Ctmp,
-    const float* __restrict__ Bbias,
     const float* __restrict__ scale,
     int64_t M,
     int64_t N,
@@ -945,11 +928,6 @@ void tinygemm_kernel(
     bool brg,
     int64_t block_size_K,
     bool do_unpack) {
-  if (Bbias != nullptr) {
-    tinygemm_kernel<scalar_t, at::Float8_e4m3fn, float, true>(
-        A, B, C, Btmp, Ctmp, scale, Bbias, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
-    return;
-  }
   tinygemm_kernel<scalar_t, at::Float8_e4m3fn, float, false>(
       A, B, C, Btmp, Ctmp, scale, nullptr, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
 }
@@ -971,6 +949,7 @@ void tinygemm_kernel(
     bool brg) {
   tinygemm_kernel2<scalar_t>(A, B, C, Btmp, Ctmp, scale, M, N, K, lda, ldb, ldc, brg);
 }
+
 template <typename scalar_t>
 void tinygemm_kernel(
     const scalar_t* __restrict__ A,
@@ -978,7 +957,6 @@ void tinygemm_kernel(
     scalar_t* __restrict__ C,
     scalar_t* __restrict__ Btmp,
     float* __restrict__ Ctmp,
-    const float* __restrict__ Bbias,
     const uint8_t* __restrict__ scale,
     int64_t M,
     int64_t N,
@@ -989,68 +967,10 @@ void tinygemm_kernel(
     bool brg,
     int64_t block_size_K,
     bool do_unpack) {
-  if (Bbias != nullptr) {
-    tinygemm_kernel<scalar_t, uint8_t, uint8_t, true>(
-        A, B, C, Btmp, Ctmp, scale, Bbias, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
-    return;
-  }
   tinygemm_kernel<scalar_t, uint8_t, uint8_t, false>(
       A, B, C, Btmp, Ctmp, scale, nullptr, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
 }
 
-// tinygemm interface
-template <typename scalar_t>
-void tinygemm_kernel(
-    const scalar_t* __restrict__ A,
-    const at::Float8_e4m3fn* __restrict__ B,
-    float* __restrict__ C,
-    scalar_t* __restrict__ Btmp,
-    const float* __restrict__ Bbias,
-    const float* __restrict__ scale,
-    int64_t M,
-    int64_t N,
-    int64_t K,
-    int64_t lda,
-    int64_t ldb,
-    int64_t ldc,
-    bool brg,
-    int64_t block_size_K,
-    bool do_unpack) {
-  if (Bbias != nullptr) {
-    tinygemm_kernel<scalar_t, at::Float8_e4m3fn, float, true>(
-        A, B, C, Btmp, scale, Bbias, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
-    return;
-  }
-  tinygemm_kernel<scalar_t, at::Float8_e4m3fn, float, false>(
-      A, B, C, Btmp, scale, nullptr, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
-}
-
-template <typename scalar_t>
-void tinygemm_kernel(
-    const scalar_t* __restrict__ A,
-    const uint8_t* __restrict__ B,
-    float* __restrict__ C,
-    scalar_t* __restrict__ Btmp,
-    const float* __restrict__ Bbias,
-    const uint8_t* __restrict__ scale,
-    int64_t M,
-    int64_t N,
-    int64_t K,
-    int64_t lda,
-    int64_t ldb,
-    int64_t ldc,
-    bool brg,
-    int64_t block_size_K,
-    bool do_unpack) {
-  if (Bbias != nullptr) {
-    tinygemm_kernel<scalar_t, uint8_t, uint8_t, true>(
-        A, B, C, Btmp, scale, Bbias, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
-    return;
-  }
-  tinygemm_kernel<scalar_t, uint8_t, uint8_t, false>(
-      A, B, C, Btmp, scale, nullptr, M, N, K, lda, ldb, ldc, brg, block_size_K, do_unpack);
-}
-
 #define INSTANTIATE_TINYGEMM_TEMPLATE(TYPE_A, TYPE_B, TYPE_S) \
   template void tinygemm_kernel<TYPE_A>(                      \
       const TYPE_A* __restrict__ A,                           \
@@ -1058,7 +978,6 @@ void tinygemm_kernel(
       TYPE_A* __restrict__ C,                                 \
       TYPE_A* __restrict__ Btmp,                              \
       float* __restrict__ Ctmp,                               \
-      const float* __restrict__ Bbias,                        \
       const TYPE_S* __restrict__ scale,                       \
       int64_t M,                                              \
       int64_t N,                                              \
@@ -1101,6 +1020,7 @@ inline const float* get_bias_data(const std::optional<at::Tensor>& bias, int64_t
   }
   return nullptr;
 }
+
 // FP8 and MXFP4 WoQ uses the same pattern:
 //   Btmp : [T, BLOCK_N * K]
 //   Ctmp : [T, BLOCK_M * BLOCK_N]

csrc/cpu/sgl-kernels/moe.cpp

@@ -812,7 +812,6 @@ void shared_expert_kernel_impl(
 static inline void check_moe_scales(
     bool use_int8_w8a8,
     bool use_fp8_w8a16,
-    bool use_mxfp4,
     const std::optional<at::Tensor>& w1_scale,
     const std::optional<at::Tensor>& w2_scale,
     const std::optional<std::vector<int64_t>> block_size) {
@@ -826,12 +825,6 @@ static inline void check_moe_scales(
     TORCH_CHECK(block_size.has_value(), "missing block_size for fp8 w8a16.");
     TORCH_CHECK(block_size.value().size() == 2, "expect block_size.size() to be 2.");
   }
-  if (use_mxfp4) {
-    TORCH_CHECK(w1_scale.has_value(), "missing w1_scale for mxfp4.");
-    TORCH_CHECK(w2_scale.has_value(), "missing w2_scale for mxfp4.");
-    TORCH_CHECK(w1_scale.value().scalar_type() == at::kByte, "expect w1_scale to be uint8.");
-    TORCH_CHECK(w2_scale.value().scalar_type() == at::kByte, "expect w2_scale to be uint8.");
-  }
 }
 
 #define CHECK_MOE_SCALES_FP8(DIM0, DIM1)                      \
@@ -846,8 +839,8 @@ static inline void check_moe_scales(
   TORCH_CHECK(w2s.size(DIM1) == div_up(N, block_size_K))
 
 // hidden_states: [M, K]
-// w1: [E, 2N, K] or [E, 2N, K / 2] for uint8
-// w2: [E, K, N] or [E, K, N / 2] for uint8
+// w1: [E, 2N, K]
+// w2: [E, K, N]
 // topk_weights: [M, topk]
 // topk_ids: [M, topk] (int32_t)
 //
@@ -865,10 +858,6 @@ at::Tensor fused_experts_cpu(
     const std::optional<at::Tensor>& w1_zero,
     const std::optional<at::Tensor>& w2_zero,
     const std::optional<std::vector<int64_t>> block_size,
-    const std::optional<at::Tensor>& w1_bias,
-    const std::optional<at::Tensor>& w2_bias,
-    const std::optional<double>& alpha,
-    const std::optional<double>& limit,
     bool is_vnni) {
   auto packed_w1 = is_vnni ? w1 : convert_weight_packed(w1);
   auto packed_w2 = is_vnni ? w2 : convert_weight_packed(w2);
@@ -908,12 +897,8 @@ at::Tensor fused_experts_cpu(
   int64_t topk = topk_weights_.size(1);
 
   // we use int32_t compensation for int8 w8a8
-  int64_t packed_K = moe_comp_method == CPUQuantMethod::MXFP4
-                         ? get_row_size<uint8_t>(K)
-                         : get_row_size(K, moe_comp_method == CPUQuantMethod::INT8_W8A8);
-  int64_t packed_N = moe_comp_method == CPUQuantMethod::MXFP4
-                         ? get_row_size<uint8_t>(N)
-                         : get_row_size(N, moe_comp_method == CPUQuantMethod::INT8_W8A8);
+  int64_t packed_K = get_row_size(K, moe_comp_method == CPUQuantMethod::INT8_W8A8);
+  int64_t packed_N = get_row_size(N, moe_comp_method == CPUQuantMethod::INT8_W8A8);
 
   // check weight shapes
   CHECK_EQ(w2.size(0), E);
@@ -926,7 +911,6 @@ at::Tensor fused_experts_cpu(
   check_moe_scales(
       moe_comp_method == CPUQuantMethod::INT8_W8A8,
       moe_comp_method == CPUQuantMethod::FP8_W8A16,
-      moe_comp_method == CPUQuantMethod::MXFP4,
       w1_scale,
       w2_scale,
       block_size);
@@ -981,7 +965,7 @@ at::Tensor fused_experts_cpu(
   //   5. Aq_tmp : [M, K] or [M * topk, N]
   //   6. As_tmp : [M * topk]
   //
-  // for fp8 w8a16 and mxfp4:
+  // for fp8 w8a16:
   //   7. intermediate_cache0 : [M * topk, 2N]
   //   8. B_tmp : [T, MAX_CACHE_BLOCK_SIZE, BLOCK_N, std::max(K, N)]
   //
@@ -994,7 +978,7 @@ at::Tensor fused_experts_cpu(
   if (moe_comp_method == CPUQuantMethod::INT8_W8A8) {
     buffer_size_nbytes += std::max(M * K, M * topk * N) + M * topk * sizeof(float);
   }
-  if (moe_comp_method == CPUQuantMethod::FP8_W8A16 || moe_comp_method == CPUQuantMethod::MXFP4) {
+  if (moe_comp_method == CPUQuantMethod::FP8_W8A16) {
     buffer_size_nbytes += M * topk * 2 * N * 2 + num_threads * MAX_CACHE_BLOCK_SIZE * BLOCK_N * std::max(K, N) * 2;
   }
   if (moe_comp_method == CPUQuantMethod::INT4_W4A8) {
@@ -1047,11 +1031,9 @@ at::Tensor fused_experts_cpu(
       float* __restrict__ C_tmp = (float*)((void*)(A_tmp + num_threads * BLOCK_M * K));
       scalar_t* __restrict__ intermediate_cache0 = (scalar_t*)((void*)(C_tmp + num_threads * 2 * BLOCK_M * BLOCK_N));
       scalar_t* __restrict__ B_tmp = (scalar_t*)((void*)(intermediate_cache0 + M * topk * 2 * N));
-      bool with_bias = w1_bias.has_value();
-      auto act_func = alpha.has_value() && limit.has_value() ? CPUAcTMethod::swiglu : CPUAcTMethod::silu_and_mul;
 
       CHECK_MOE_SCALES_FP8(1, 2);
-      fused_experts_fp_kernel_impl<scalar_t, at::Float8_e4m3fn, float, false>(
+      fused_experts_fp8_kernel_impl(
           out_hidden_states.data_ptr<scalar_t>(),
           intermediate_cache0,
           intermediate_cache1,
@@ -1062,8 +1044,6 @@ at::Tensor fused_experts_cpu(
           hidden_states.data_ptr<scalar_t>(),
           packed_w1.data_ptr<at::Float8_e4m3fn>(),
           packed_w2.data_ptr<at::Float8_e4m3fn>(),
-          with_bias ? w1_bias.value().data_ptr<float>() : nullptr,
-          with_bias ? w2_bias.value().data_ptr<float>() : nullptr,
           w1s.data_ptr<float>(),
           w2s.data_ptr<float>(),
           block_size_N,
@@ -1077,56 +1057,7 @@ at::Tensor fused_experts_cpu(
           K,
           E,
           topk,
-          num_tokens_post_pad,
-          alpha.has_value() ? float(alpha.value()) : 0,
-          limit.has_value() ? float(limit.value()) : 0,
-          act_func,
-          with_bias);
-    } else if (moe_comp_method == CPUQuantMethod::MXFP4) {
-      scalar_t* __restrict__ A_tmp = (scalar_t*)((void*)(intermediate_cache2 + M * topk * K));
-      float* __restrict__ C_tmp = (float*)((void*)(A_tmp + num_threads * BLOCK_M * K));
-      scalar_t* __restrict__ intermediate_cache0 = (scalar_t*)((void*)(C_tmp + num_threads * 2 * BLOCK_M * BLOCK_N));
-      scalar_t* __restrict__ B_tmp = (scalar_t*)((void*)(intermediate_cache0 + M * topk * 2 * N));
-      bool with_bias = w1_bias.has_value();
-      auto act_func = alpha.has_value() && limit.has_value() ? CPUAcTMethod::swiglu : CPUAcTMethod::silu_and_mul;
-
-      // mxfp4 supports only group size of 32 (2^5)
-      constexpr int64_t group_size = 32;
-      auto w1s = w1_scale.value();
-      auto w2s = w2_scale.value();
-      TORCH_CHECK(w1s.numel() == E * 2 * N * K / group_size, "w1_scale size mismatch");
-      TORCH_CHECK(w2s.numel() == E * K * N / group_size, "w2_scale size mismatch");
-      fused_experts_fp_kernel_impl<scalar_t, uint8_t, uint8_t, true>(
-          out_hidden_states.data_ptr<scalar_t>(),
-          intermediate_cache0,
-          intermediate_cache1,
-          intermediate_cache2,
-          A_tmp,
-          B_tmp,
-          C_tmp,
-          hidden_states.data_ptr<scalar_t>(),
-          packed_w1.data_ptr<uint8_t>(),
-          packed_w2.data_ptr<uint8_t>(),
-          with_bias ? w1_bias.value().data_ptr<float>() : nullptr,
-          with_bias ? w2_bias.value().data_ptr<float>() : nullptr,
-          w1s.data_ptr<uint8_t>(),
-          w2s.data_ptr<uint8_t>(),
-          /*block_size_N*/ 1,
-          /*block_size_K*/ group_size,
-          topk_weights_.data_ptr<float>(),
-          sorted_ids,
-          expert_ids,
-          offsets,
-          M,
-          N,
-          K,
-          E,
-          topk,
-          num_tokens_post_pad,
-          alpha.has_value() ? float(alpha.value()) : 0,
-          limit.has_value() ? float(limit.value()) : 0,
-          act_func,
-          with_bias);
+          num_tokens_post_pad);
     } else if (moe_comp_method == CPUQuantMethod::INT4_W4A8) {
       uint8_t* __restrict__ A_tmp = (uint8_t*)((void*)(intermediate_cache2 + M * topk * K));
       float* __restrict__ C_tmp = (float*)((void*)(A_tmp + num_threads * BLOCK_M * K));
@@ -1254,7 +1185,7 @@ at::Tensor shared_expert_cpu(
   CHECK_EQ(packed_w2.size(1), packed_N);
 
   // check scales
-  check_moe_scales(use_int8_w8a8, use_fp8_w8a16, false, w1_scale, w2_scale, block_size);
+  check_moe_scales(use_int8_w8a8, use_fp8_w8a16, w1_scale, w2_scale, block_size);
 
   at::Tensor out_hidden_states = inplace ? hidden_states : at::empty_like(hidden_states);
 

csrc/cpu/sgl-kernels/moe.h

@@ -176,109 +176,3 @@ inline void silu_and_mul_stub(
     out_vec.store(out + d);
   }
 }
-
-template <typename scalar_t>
-inline void copy_mul_stub(scalar_t* __restrict__ out, const float* __restrict__ input, float weight, int64_t size) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = bVec::size();
-  const fVec weight_vec = fVec(weight);
-  int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    fVec data0 = fVec::loadu(input + d) * weight_vec;
-    fVec data1 = fVec::loadu(input + d + fVec::size()) * weight_vec;
-    bVec out_vec = convert_from_float_ext<scalar_t>(data0, data1);
-    out_vec.store(out + d);
-  }
-  for (; d < size; ++d) {
-    out[d] = static_cast<scalar_t>(input[d] * weight);
-  }
-}
-
-// input = input + input2
-inline void add_bias_stub(float* __restrict__ input, const float* __restrict__ input2, int64_t size) {
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = fVec::size();
-  int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    fVec x_fvec = fVec::loadu(input + d);
-    fVec y_fvec = fVec::loadu(input2 + d);
-    x_fvec = x_fvec + y_fvec;
-    x_fvec.store(input + d);
-  }
-  for (; d < size; ++d) {
-    input[d] = input[d] + input2[d];
-  }
-}
-
-template <typename scalar_t>
-inline void copy_mul_stub(scalar_t* __restrict__ out, const scalar_t* __restrict__ input, float weight, int64_t size) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = bVec::size();
-  const fVec weight_vec = fVec(weight);
-  int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    bVec x = bVec::loadu(input + d);
-    fVec x0, x1;
-    std::tie(x0, x1) = at::vec::convert_to_float(x);
-    x0 = x0 * weight_vec;
-    x1 = x1 * weight_vec;
-    bVec out_vec = convert_from_float_ext<scalar_t>(x0, x1);
-    out_vec.store(out + d);
-  }
-  for (; d < size; ++d) {
-    out[d] = static_cast<scalar_t>(input[d] * weight);
-  }
-}
-
-template <typename scalar_t>
-inline void clamp_sigmoid_and_mul_stub(
-    scalar_t* __restrict__ out,
-    const scalar_t* __restrict__ input,
-    int64_t size,
-    const float alpha,
-    const float limit) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  const fVec one = fVec(1.f);
-  const fVec zero = fVec(0.f);
-  const fVec limit_v = fVec(limit);
-  const fVec nlimit_v = fVec(-limit);
-  const fVec alpha_v = fVec(alpha);
-
-  // no remainder
-#pragma GCC unroll 4
-  for (int64_t d = 0; d < size; d += bVec::size()) {
-    bVec x = bVec::loadu(input + d);
-    fVec x0_, y0_;
-    std::tie(x0_, y0_) = at::vec::convert_to_float(x);
-    float tmp_buffer[fVec::size() * 2];  // 32
-    float tmp_glu[fVec::size()];         // 16
-    float tmp_linear[fVec::size()];      // 16
-    x0_.store(tmp_buffer);
-    y0_.store(tmp_buffer + fVec::size());
-    // interleaved: x[2i] = glu, x[2i+1] = linear
-    for (int j = 0; j < fVec::size(); ++j) {
-      // x0 [0,2,..30]
-      tmp_glu[j] = tmp_buffer[j * 2];
-      // y0 [1,3,...31]
-      tmp_linear[j] = tmp_buffer[j * 2 + 1];
-    }
-    fVec x0 = fVec::loadu(tmp_glu);
-    fVec y0 = fVec::loadu(tmp_linear);
-
-    // clamp
-    x0 = at::vec::minimum(x0, limit_v);
-    y0 = at::vec::minimum(limit_v, at::vec::maximum(nlimit_v, y0));
-    // x * sigmoid(x * alpha)
-    x0 = x0 / (one + (x0 * alpha_v).neg().exp_u20());
-    // (y + 1) * x
-    y0 = y0 + one;
-    x0 = x0 * y0;
-    convert_from_float_and_store<scalar_t>(out + d / 2, x0);
-  }
-}

csrc/cpu/sgl-kernels/moe_fp8.cpp

@@ -7,8 +7,8 @@
 #include "gemm.h"
 #include "moe.h"
 
-template <typename scalar_t, typename packed_t, typename param_t, bool is_mxfp4>
-void fused_experts_fp_kernel_impl(
+template <typename scalar_t>
+void fused_experts_fp8_kernel_impl(
     scalar_t* __restrict__ output,
     scalar_t* __restrict__ ic0,
     scalar_t* __restrict__ ic1,
@@ -17,12 +17,10 @@ void fused_experts_fp_kernel_impl(
     scalar_t* __restrict__ B_tmp,
     float* __restrict__ C_tmp,
     const scalar_t* __restrict__ input,
-    const packed_t* __restrict__ packed_w1,
-    const packed_t* __restrict__ packed_w2,
-    const float* __restrict__ w1_bias,
-    const float* __restrict__ w2_bias,
-    const param_t* __restrict__ w1s,
-    const param_t* __restrict__ w2s,
+    const at::Float8_e4m3fn* __restrict__ packed_w1,
+    const at::Float8_e4m3fn* __restrict__ packed_w2,
+    const float* __restrict__ w1s,
+    const float* __restrict__ w2s,
     int64_t block_size_N,
     int64_t block_size_K,
     const float* __restrict__ topk_weights,
@@ -34,11 +32,7 @@ void fused_experts_fp_kernel_impl(
     int64_t K,
     int64_t E,
     int64_t topk,
-    int64_t num_tokens_post_pad,
-    float alpha,
-    float limit,
-    CPUAcTMethod act_func,
-    bool with_bias) {
+    int64_t num_tokens_post_pad) {
   constexpr int64_t BLOCK_M = block_size_m();
   constexpr int64_t BLOCK_N = block_size_n();
 
@@ -48,20 +42,12 @@ void fused_experts_fp_kernel_impl(
   int64_t scale_size_N = div_up(2 * N, block_size_N);
   int64_t scale_size_K = div_up(K, block_size_K);
   int64_t blocks_n_per_group = block_size_N / BLOCK_N;
-  std::function<int64_t(int64_t)> scale_offset_per_block;
-  if constexpr (is_mxfp4) {
-    scale_offset_per_block = [&](int64_t a) { return a * BLOCK_N; };
-  } else {
-    scale_offset_per_block = [&](int64_t a) { return a / blocks_n_per_group; };
-  }
 
-  const int64_t packed_K = get_row_size<packed_t>(K);
-
-  const int64_t stride_e = 2 * N * packed_K;
-  const int64_t stride_n = packed_K;
+  const int64_t stride_e = 2 * N * K;
+  const int64_t stride_n = K;
 
   int64_t avg_M = std::max(int64_t(1), M * topk / E);
-  const bool use_brgemm = can_use_brgemm<packed_t>(avg_M);
+  const bool use_brgemm = can_use_brgemm<at::Float8_e4m3fn>(avg_M);
 
   int64_t B_tmp_size_per_thread = MAX_CACHE_BLOCK_SIZE * BLOCK_N * std::max(K, N);
 
@@ -71,15 +57,14 @@ void fused_experts_fp_kernel_impl(
     int tid = get_thread_num();
     scalar_t* __restrict__ A = A_tmp + tid * BLOCK_M * K;
 
-    loop_2d<packed_t>(mb0, mb1, nb0, nb1, BLOCK_N * K, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
+    loop_2d<at::Float8_e4m3fn>(mb0, mb1, nb0, nb1, BLOCK_N * K, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
       int64_t n_size = std::min(2 * N - nb * BLOCK_N, BLOCK_N);
 
       // B shape [K, n_size] in vnni format
       int32_t expert_id = expert_ids[mb];
-      const packed_t* __restrict__ B = packed_w1 + expert_id * stride_e + nb * BLOCK_N * stride_n;
-      const param_t* __restrict__ Bs =
-          w1s + expert_id * scale_size_N * scale_size_K + scale_offset_per_block(nb) * scale_size_K;
-      const float* __restrict__ B_bias = with_bias ? w1_bias + expert_id * 2 * N + nb * BLOCK_N : nullptr;
+      const at::Float8_e4m3fn* __restrict__ B = packed_w1 + expert_id * stride_e + nb * BLOCK_N * stride_n;
+      const float* __restrict__ Bs =
+          w1s + expert_id * scale_size_N * scale_size_K + (nb / blocks_n_per_group) * scale_size_K;
 
       // do unpacking for the first row or a new expert
       int32_t pre_expert_id = mb == 0 ? -1 : expert_ids[mb - 1];
@@ -103,7 +88,6 @@ void fused_experts_fp_kernel_impl(
           /*   C            */ ic0 + offset * 2 * N + nb * BLOCK_N,
           /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * K,
           /*   Ctmp         */ C_tmp + tid * 2 * BLOCK_M * BLOCK_N,
-          /*   Bbias        */ B_bias,
           /*   scale        */ Bs,
           /*   M            */ m_size,
           /*   N            */ n_size,
@@ -122,20 +106,12 @@ void fused_experts_fp_kernel_impl(
   });
 
   // stage 1.5: intermediate_cache1 = silu(intermediate_cache0)
-  if (act_func == CPUAcTMethod::silu_and_mul) {
-    at::parallel_for(0, M * topk, 0, [&](int64_t begin, int64_t end) {
-      for (int64_t m = begin; m < end; ++m) {
-        silu_and_mul_stub(ic1 + m * N, ic0 + m * 2 * N, ic0 + m * 2 * N + N, N);
-      }
-    });
-  } else if (act_func == CPUAcTMethod::swiglu) {
-    at::parallel_for(0, M * topk, 0, [&](int64_t begin, int64_t end) {
-      for (int64_t m = begin; m < end; ++m) {
-        clamp_sigmoid_and_mul_stub(ic1 + m * N, ic0 + m * 2 * N, N, alpha, limit);
-        clamp_sigmoid_and_mul_stub(ic1 + m * N + N / 2, ic0 + m * 2 * N + N, N, alpha, limit);
-      }
-    });
-  }
+  at::parallel_for(0, M * topk, 0, [&](int64_t begin, int64_t end) {
+    for (int64_t m = begin; m < end; ++m) {
+      silu_and_mul_stub(ic1 + m * N, ic0 + m * 2 * N, ic0 + m * 2 * N + N, N);
+    }
+  });
+
   // stage 2: intermediate_cache2 = intermediate_cache1 @ w2
   //   w2 : [E, K, N] as [E, OC, IC]
   const int64_t OC = K;  // rename K as OC
@@ -144,16 +120,15 @@ void fused_experts_fp_kernel_impl(
   const int64_t NB2 = div_up(OC, BLOCK_N);
   scale_size_N = div_up(K, block_size_N);
   scale_size_K = div_up(N, block_size_K);
-  const int64_t packed_IC = get_row_size<packed_t>(IC);
-  const int64_t stride_e2 = OC * packed_IC;
-  const int64_t stride_oc = packed_IC;
+  const int64_t stride_e2 = OC * IC;
+  const int64_t stride_oc = IC;
 
   // parallel on [MB2, NB2]
   parallel_2d(MB2, NB2, [&](int64_t mb0, int64_t mb1, int64_t nb0, int64_t nb1) {
     int tid = get_thread_num();
     alignas(64) scalar_t C[BLOCK_M * BLOCK_K];
 
-    loop_2d<packed_t>(mb0, mb1, nb0, nb1, BLOCK_N * IC, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
+    loop_2d<at::Float8_e4m3fn>(mb0, mb1, nb0, nb1, BLOCK_N * IC, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
       int64_t m_size = offsets[mb + 1] - offsets[mb];
       int64_t n_size = std::min(OC - nb * BLOCK_N, BLOCK_N);
 
@@ -164,10 +139,9 @@ void fused_experts_fp_kernel_impl(
 
       // B shape [IC, n_size] in vnni format
       int32_t expert_id = expert_ids[mb];
-      const packed_t* __restrict__ B = packed_w2 + expert_id * stride_e2 + nb * BLOCK_N * stride_oc;
-      const param_t* __restrict__ Bs =
-          w2s + expert_id * scale_size_N * scale_size_K + scale_offset_per_block(nb) * scale_size_K;
-      const float* __restrict__ B_bias = with_bias ? w2_bias + expert_id * OC + nb * BLOCK_N : nullptr;
+      const at::Float8_e4m3fn* __restrict__ B = packed_w2 + expert_id * stride_e2 + nb * BLOCK_N * stride_oc;
+      const float* __restrict__ Bs =
+          w2s + expert_id * scale_size_N * scale_size_K + (nb / blocks_n_per_group) * scale_size_K;
 
       // do unpacking for the first row or a new expert
       int32_t pre_expert_id = mb == 0 ? -1 : expert_ids[mb - 1];
@@ -179,7 +153,6 @@ void fused_experts_fp_kernel_impl(
           /*   C            */ C,
           /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * IC,
           /*   Ctmp         */ C_tmp + tid * 2 * BLOCK_M * BLOCK_N,
-          /*   Bbias        */ B_bias,
           /*   scale        */ Bs,
           /*   M            */ m_size,
           /*   N            */ n_size,
@@ -204,6 +177,7 @@ void fused_experts_fp_kernel_impl(
       at::native::cpublas::brgemm_release();
     }
   });
+
   // stage 3: out = intermediate_cache2.sum(dim=1)
   //   from [M, topk, K] to [M, K]
   at::parallel_for(0, M, 0, [&](int64_t begin, int64_t end) {
@@ -213,43 +187,35 @@ void fused_experts_fp_kernel_impl(
   });
 }
 
-#define INSTANTIATE_MOE_FP_TEMPLATE(TYPE1, TYPE2, TYPE3, IS_MXFP4)           \
-  template void fused_experts_fp_kernel_impl<TYPE1, TYPE2, TYPE3, IS_MXFP4>( \
-      TYPE1* __restrict__ output,                                            \
-      TYPE1* __restrict__ ic0,                                               \
-      TYPE1* __restrict__ ic1,                                               \
-      TYPE1* __restrict__ ic2,                                               \
-      TYPE1* __restrict__ A_tmp,                                             \
-      TYPE1* __restrict__ B_tmp,                                             \
-      float* __restrict__ C_tmp,                                             \
-      const TYPE1* __restrict__ input,                                       \
-      const TYPE2* __restrict__ packed_w1,                                   \
-      const TYPE2* __restrict__ packed_w2,                                   \
-      const float* __restrict__ w1_bias,                                     \
-      const float* __restrict__ w2_bias,                                     \
-      const TYPE3* __restrict__ w1s,                                         \
-      const TYPE3* __restrict__ w2s,                                         \
-      int64_t block_size_N,                                                  \
-      int64_t block_size_K,                                                  \
-      const float* __restrict__ topk_weights,                                \
-      const int32_t* __restrict__ sorted_ids,                                \
-      const int32_t* __restrict__ expert_ids,                                \
-      const int32_t* __restrict__ offsets,                                   \
-      int64_t M,                                                             \
-      int64_t N,                                                             \
-      int64_t K,                                                             \
-      int64_t E,                                                             \
-      int64_t topk,                                                          \
-      int64_t num_tokens_post_pad,                                           \
-      float alpha,                                                           \
-      float limit,                                                           \
-      CPUAcTMethod act_func,                                                 \
-      bool with_bias)
+#define INSTANTIATE_MOE_FP8_TEMPLATE(TYPE)             \
+  template void fused_experts_fp8_kernel_impl<TYPE>(   \
+      TYPE* __restrict__ output,                       \
+      TYPE* __restrict__ ic0,                          \
+      TYPE* __restrict__ ic1,                          \
+      TYPE* __restrict__ ic2,                          \
+      TYPE* __restrict__ A_tmp,                        \
+      TYPE* __restrict__ B_tmp,                        \
+      float* __restrict__ C_tmp,                       \
+      const TYPE* __restrict__ input,                  \
+      const at::Float8_e4m3fn* __restrict__ packed_w1, \
+      const at::Float8_e4m3fn* __restrict__ packed_w2, \
+      const float* __restrict__ w1s,                   \
+      const float* __restrict__ w2s,                   \
+      int64_t block_size_N,                            \
+      int64_t block_size_K,                            \
+      const float* __restrict__ topk_weights,          \
+      const int32_t* __restrict__ sorted_ids,          \
+      const int32_t* __restrict__ expert_ids,          \
+      const int32_t* __restrict__ offsets,             \
+      int64_t M,                                       \
+      int64_t N,                                       \
+      int64_t K,                                       \
+      int64_t E,                                       \
+      int64_t topk,                                    \
+      int64_t num_tokens_post_pad)
 
-INSTANTIATE_MOE_FP_TEMPLATE(at::BFloat16, at::Float8_e4m3fn, float, false);
-INSTANTIATE_MOE_FP_TEMPLATE(at::Half, at::Float8_e4m3fn, float, false);
-INSTANTIATE_MOE_FP_TEMPLATE(at::BFloat16, uint8_t, uint8_t, true);
-INSTANTIATE_MOE_FP_TEMPLATE(at::Half, uint8_t, uint8_t, true);
+INSTANTIATE_MOE_FP8_TEMPLATE(at::BFloat16);
+INSTANTIATE_MOE_FP8_TEMPLATE(at::Half);
 
 template <typename scalar_t>
 void shared_expert_fp8_kernel_impl(
@@ -300,7 +266,6 @@ void shared_expert_fp8_kernel_impl(
           /*   C            */ ic0 + mb * BLOCK_M * 2 * N + nb * BLOCK_N,
           /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * K,
           /*   Ctmp         */ C_tmp + tid * 2 * BLOCK_M * BLOCK_N,
-          /*   Bbias        */ nullptr,
           /*   scale        */ w1s + (nb / blocks_n_per_group) * scale_size_K,
           /*   M            */ m_size,
           /*   N            */ n_size,
@@ -352,7 +317,6 @@ void shared_expert_fp8_kernel_impl(
           /*   C            */ C,
           /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * IC,
           /*   Ctmp         */ C_tmp + tid * 2 * BLOCK_M * BLOCK_N,
-          /*   Bbias        */ nullptr,
           /*   scale        */ w2s + (nb / blocks_n_per_group) * scale_size_K,
           /*   M            */ m_size,
           /*   N            */ n_size,

csrc/cpu/sgl-kernels/vec.h

@@ -11,7 +11,7 @@
 
 #include <ATen/cpu/vec/functional.h>
 #include <ATen/cpu/vec/vec.h>
-#include <immintrin.h>
+
 namespace {
 
 using namespace at::vec;
@@ -21,15 +21,6 @@ inline Vectorized<scalar_t> convert_from_float_ext(const Vectorized<float>& a, c
   return at::vec::convert_from_float<scalar_t>(a, b);
 }
 
-template <typename scalar_t>
-inline void convert_from_float_and_store(scalar_t* out, const Vectorized<float>& a) {
-  float out_buffer[at::vec::Vectorized<float>::size()];
-  a.store(out_buffer);
-  for (int i = 0; i < 16; i++) {
-    out[i] = (scalar_t)out_buffer[i];
-  }
-}
-
 // allow f16, bf16
 template <typename scalar_t, typename std::enable_if_t<is_reduced_floating_point_v<scalar_t>, int> = 1>
 inline std::tuple<Vectorized<float>, Vectorized<float>> load_float_vec2(const scalar_t* __restrict__ data) {
@@ -59,11 +50,6 @@ convert_from_float_ext<at::BFloat16>(const Vectorized<float>& a, const Vectorize
   return (__m512i)(_mm512_cvtne2ps_pbh(__m512(b), __m512(a)));
 }
 
-template <>
-inline void convert_from_float_and_store<at::BFloat16>(at::BFloat16* out, const Vectorized<float>& a) {
-  _mm256_storeu_si256((__m256i*)out, (__m256i)(_mm512_cvtneps_pbh(__m512(a))));
-}
-
 #define CVT_BF16_TO_FP32(a) _mm512_castsi512_ps(_mm512_slli_epi32(_mm512_cvtepu16_epi32(a), 16))
 
 #define CVT_FP16_TO_FP32(a) _mm512_cvtph_ps(a)
@@ -139,6 +125,7 @@ inline __m512bh CVT_FP8_TO_BF16(__m256i a) {
   return cvt_e4m3_bf16_intrinsic_with_denorm(a);
 #endif
 }
+
 // faster version of float8_e4m3fn conversion to bfloat16
 //
 // we mapped cuda implementation from below link and vectorized with avx512:

csrc/cpu/torch_bindings.cpp

@@ -64,8 +64,6 @@ at::Tensor weight_packed_linear(at::Tensor& mat1, at::Tensor& mat2,
 
 at::Tensor convert_weight_packed(at::Tensor& weight);
 
-at::Tensor convert_scale_packed(at::Tensor& scale);
-
 at::Tensor fused_experts_cpu(
     at::Tensor& hidden_states, at::Tensor& w1, at::Tensor& w2,
     at::Tensor& topk_weights, at::Tensor& topk_ids, bool inplace,
@@ -73,11 +71,7 @@ at::Tensor fused_experts_cpu(
     const std::optional<at::Tensor>& w2_scale,
     const std::optional<at::Tensor>& w1_zero,
     const std::optional<at::Tensor>& w2_zero,
-    const std::optional<std::vector<int64_t>> block_size,
-    const std::optional<at::Tensor>& w1_bias,
-    const std::optional<at::Tensor>& w2_bias,
-    const std::optional<double>& alpha, const std::optional<double>& limit,
-    bool is_vnni);
+    const std::optional<std::vector<int64_t>> block_size, bool is_vnni);
 
 at::Tensor int8_scaled_mm_with_quant(at::Tensor& mat1, at::Tensor& mat2,
                                      at::Tensor& scales2,
@@ -104,45 +98,6 @@ at::Tensor int4_scaled_mm_cpu(at::Tensor& x, at::Tensor& w, at::Tensor& w_zeros,
                               at::Tensor& w_scales,
                               std::optional<at::Tensor> bias);
 
-// Adapted from sglang: GDN
-std::tuple<at::Tensor, at::Tensor> chunk_gated_delta_rule_cpu(
-    const at::Tensor& query, const at::Tensor& key, const at::Tensor& value,
-    const at::Tensor& g, const at::Tensor& beta,
-    const at::Tensor& initial_state, bool output_final_state,
-    const at::Tensor& cu_seqlens, bool head_first, bool use_qk_l2norm_in_kernel,
-    double eps = 1e-5);
-
-at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
-    const at::Tensor& A_log, const at::Tensor& dt_bias, const at::Tensor& q,
-    const at::Tensor& k, const at::Tensor& v, const at::Tensor& a,
-    const at::Tensor& b, at::Tensor& initial_state_source,
-    const at::Tensor& initial_state_indices, const at::Tensor& cu_seqlens,
-    bool use_qk_l2norm_in_kernel, double softplus_beta = 1.0,
-    double softplus_threshold = 20.0);
-
-std::tuple<at::Tensor, at::Tensor> fused_gdn_gating_cpu(
-    const at::Tensor& A_log, const at::Tensor& a, const at::Tensor& b,
-    const at::Tensor& dt_bias);
-
-// Adapted from sglang: casual_conv1d kernels
-at::Tensor causal_conv1d_weight_pack(const at::Tensor& weight);
-
-at::Tensor causal_conv1d_fwd_cpu(
-    const at::Tensor& x, const at::Tensor& weight,
-    const std::optional<at::Tensor>& bias,
-    const std::optional<at::Tensor>& conv_states,
-    const std::optional<at::Tensor>& query_start_loc,
-    const std::optional<at::Tensor>& cache_indices,
-    const std::optional<at::Tensor>& has_initial_state, bool silu_activation,
-    int64_t pad_slot_id, bool is_vnni);
-
-at::Tensor causal_conv1d_update_cpu(
-    const at::Tensor& x, const at::Tensor& conv_states,
-    const at::Tensor& weight, const std::optional<at::Tensor>& bias,
-    bool silu_activation, const std::optional<at::Tensor>& cache_seqlens,
-    const std::optional<at::Tensor>& conv_state_indices, int64_t pad_slot_id,
-    bool is_vnni);
-
 void activation_lut_bf16(torch::Tensor& out, torch::Tensor& input,
                          const std::string& activation);
 
@@ -410,15 +365,11 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
   ops.impl("weight_packed_linear", torch::kCPU, &weight_packed_linear);
   ops.def("convert_weight_packed(Tensor! weight) -> Tensor");
   ops.impl("convert_weight_packed", torch::kCPU, &convert_weight_packed);
-  ops.def("convert_scale_packed(Tensor! scale) -> Tensor");
-  ops.impl("convert_scale_packed", torch::kCPU, &convert_scale_packed);
   ops.def(
       "fused_experts_cpu(Tensor hidden_states, Tensor w1, Tensor w2, Tensor "
       "topk_weights, Tensor topk_ids, bool "
       "inplace, int moe_comp_method, Tensor? w1_scale, Tensor? w2_scale, "
-      "Tensor? w1_zero, Tensor? w2_zero, int[]? block_size, "
-      "Tensor? w1_bias, Tensor? w2_bias, float? alpha, float? limit, "
-      "bool is_vnni) -> "
+      "Tensor? w1_zero, Tensor? w2_zero, int[]? block_size, bool is_vnni) -> "
       "Tensor");
   ops.impl("fused_experts_cpu", torch::kCPU, &fused_experts_cpu);
   ops.def(
@@ -446,47 +397,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
       "scales2, SymInt[] block_size, Tensor? bias, ScalarType out_dtype, "
       "bool is_vnni) -> Tensor");
   ops.impl("fp8_scaled_mm_cpu", torch::kCPU, &fp8_scaled_mm_cpu);
-
-  // Adapted from sglang: GDN kernels
-  ops.def(
-      "chunk_gated_delta_rule_cpu(Tensor query, Tensor key, Tensor value, "
-      "Tensor g, Tensor beta, "
-      "Tensor initial_state, bool output_final_state, Tensor cu_seqlens, bool "
-      "head_first, "
-      "bool use_qk_l2norm_in_kernel, float eps=1e-5) -> (Tensor, Tensor)");
-  ops.impl("chunk_gated_delta_rule_cpu", torch::kCPU,
-           &chunk_gated_delta_rule_cpu);
-  ops.def(
-      "fused_sigmoid_gating_delta_rule_update_cpu(Tensor A_log, Tensor "
-      "dt_bias, Tensor q, Tensor k, Tensor v, Tensor "
-      "a, Tensor b, Tensor(a!) initial_state_source, Tensor "
-      "initial_state_indices, Tensor cu_seqlens, bool "
-      "use_qk_l2norm_in_kernel, float softplus_beta=1.0, float "
-      "softplus_threshold=20.0) -> Tensor");
-  ops.impl("fused_sigmoid_gating_delta_rule_update_cpu", torch::kCPU,
-           &fused_sigmoid_gating_delta_rule_update_cpu);
-  ops.def(
-      "fused_gdn_gating_cpu(Tensor A_log, Tensor a, Tensor b, Tensor dt_bias) "
-      "-> (Tensor, Tensor)");
-  ops.impl("fused_gdn_gating_cpu", torch::kCPU, &fused_gdn_gating_cpu);
-
-  // Adapted from sglang: casual_conv1d kernels
-  ops.def("causal_conv1d_weight_pack(Tensor weight) -> Tensor");
-  ops.impl("causal_conv1d_weight_pack", torch::kCPU,
-           &causal_conv1d_weight_pack);
-  ops.def(
-      "causal_conv1d_fwd_cpu(Tensor x, Tensor weight, Tensor? bias, Tensor? "
-      "conv_states, Tensor? query_start_loc,"
-      "Tensor? cache_indices, Tensor? has_initial_state, bool silu_activation, "
-      "int pad_slot_id, bool is_vnni) -> "
-      "Tensor");
-  ops.impl("causal_conv1d_fwd_cpu", torch::kCPU, &causal_conv1d_fwd_cpu);
-  ops.def(
-      "causal_conv1d_update_cpu(Tensor x, Tensor(a!) conv_states, Tensor "
-      "weight, Tensor? bias, bool silu_activation,"
-      "Tensor? cache_seqlens, Tensor? conv_state_indices, int pad_slot_id, "
-      "bool is_vnni) -> Tensor");
-  ops.impl("causal_conv1d_update_cpu", torch::kCPU, &causal_conv1d_update_cpu);
 #endif
 
   // CPU attention kernels

csrc/cuda_vec_utils.cuh

@@ -9,8 +9,6 @@
 
 #ifdef USE_ROCM
   #include <hip/hip_runtime.h>
-  #include <hip/hip_bf16.h>
-  #include <hip/hip_fp16.h>
 #else
   #include <cuda_bf16.h>
   #include <cuda_fp16.h>

csrc/cutlass_extensions/torch_utils.hpp

@@ -1,7 +1,5 @@
 #pragma once
 
-#include "torch_utils.h"
-
 // This header is shared between _C (unstable ABI, used by machete) and
 // _C_stable_libtorch (stable ABI, used by W4A8/sparse). TORCH_TARGET_VERSION
 // is defined only for the stable target, so we switch includes and types
@@ -10,9 +8,13 @@
   #include <torch/csrc/stable/tensor.h>
   #include <torch/headeronly/util/BFloat16.h>
   #include <torch/headeronly/util/Half.h>
+  #include <torch/headeronly/util/shim_utils.h>  // for STD_TORCH_CHECK
 using TorchTensor = torch::stable::Tensor;
+  #define TORCH_UTILS_CHECK STD_TORCH_CHECK
 #else
+  #include <torch/all.h>
 using TorchTensor = torch::Tensor;
+  #define TORCH_UTILS_CHECK TORCH_CHECK
 #endif
 
 #include "cute/layout.hpp"

csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu

@@ -87,12 +87,6 @@ constexpr int kScaleBytesPerToken = kNumQuantBlocks + 1;  // 8 (7 real + 1 pad)
 constexpr int kTokenDataBytes = kNopeDim + kRopeDim * 2;  // 448 + 128 = 576
 constexpr float kFp8Max = 448.0f;
 
-#ifndef USE_ROCM
-// When num_tokens is less than this threshold,
-// run the reduced grid variant on cuda
-constexpr float NUM_TOKEN_CUTOFF = 1024;
-#endif
-
 // Per-warp layout:  32 lanes × 16 elems/lane = 512 elems = HEAD_DIM.
 constexpr int kNumLanes = 32;
 constexpr int kElemsPerLane = kHeadDim / kNumLanes;  // 16
@@ -118,197 +112,6 @@ __device__ __forceinline__ float warpSum(float val) {
   return val;
 }
 
-// ────────────────────────────────────────────────────────────────────────────
-// Per-slot inner pipeline
-// ────────────────────────────────────────────────────────────────────────────
-// Shared by both kernel variants: 1 CTA per (token, head) pair vs. 1 CTA per
-// token
-template <typename scalar_t_in>
-__device__ __forceinline__ void processDeepseekV4Slot(
-    uint4 v0, uint4 v1, int const tokenIdx, int const slotIdx,
-    int const dim_base, int const laneId, int const num_heads_q,
-    float const eps, scalar_t_in* __restrict__ q_inout,
-    uint8_t* __restrict__ k_cache, int64_t const* __restrict__ slot_mapping,
-    int64_t const* __restrict__ position_ids,
-    float const* __restrict__ cos_sin_cache, int const cache_block_size,
-    int const kv_block_stride) {
-  using Converter = vllm::_typeConvert<scalar_t_in>;
-  bool const isKV = (slotIdx == num_heads_q);
-
-  // ── Decode the bf16 → 16 fp32 registers ─────────────────────────────
-  float elements[kElemsPerLane];
-  {
-    typename Converter::packed_hip_type const* p0 =
-        reinterpret_cast<typename Converter::packed_hip_type const*>(&v0);
-    typename Converter::packed_hip_type const* p1 =
-        reinterpret_cast<typename Converter::packed_hip_type const*>(&v1);
-#pragma unroll
-    for (int i = 0; i < 4; i++) {
-      float2 f2 = Converter::convert(p0[i]);
-      elements[2 * i] = f2.x;
-      elements[2 * i + 1] = f2.y;
-    }
-#pragma unroll
-    for (int i = 0; i < 4; i++) {
-      float2 f2 = Converter::convert(p1[i]);
-      elements[8 + 2 * i] = f2.x;
-      elements[8 + 2 * i + 1] = f2.y;
-    }
-  }
-
-  // ── Q branch: RMSNorm (no weight) ───────────────────────────────────
-  if (!isKV) {
-    float sumOfSquares = 0.0f;
-#pragma unroll
-    for (int i = 0; i < kElemsPerLane; i++) {
-      sumOfSquares += elements[i] * elements[i];
-    }
-    sumOfSquares = warpSum<float>(sumOfSquares);
-    float const rms_rcp =
-        rsqrtf(sumOfSquares / static_cast<float>(kHeadDim) + eps);
-#pragma unroll
-    for (int i = 0; i < kElemsPerLane; i++) {
-      elements[i] = elements[i] * rms_rcp;
-    }
-  }
-
-  // ── GPT-J RoPE on dims [NOPE_DIM, HEAD_DIM) ─────────────────────────────
-  // All math in fp32.  cos_sin_cache is loaded as fp32 (its native storage).
-  bool const is_rope_lane = dim_base >= kNopeDim;
-  if (is_rope_lane) {
-    int64_t const pos = position_ids[tokenIdx];
-    constexpr int kHalfRope = kRopeDim / 2;
-    float const* cos_ptr = cos_sin_cache + pos * kRopeDim;
-    float const* sin_ptr = cos_ptr + kHalfRope;
-
-    int const rope_local_base = dim_base - kNopeDim;
-    int const half_base = rope_local_base >> 1;
-
-    // Load phase: 4 vectorized LDGs issue back-to-back.
-    float4 const c0 = *reinterpret_cast<float4 const*>(cos_ptr + half_base);
-    float4 const c1 = *reinterpret_cast<float4 const*>(cos_ptr + half_base + 4);
-    float4 const s0 = *reinterpret_cast<float4 const*>(sin_ptr + half_base);
-    float4 const s1 = *reinterpret_cast<float4 const*>(sin_ptr + half_base + 4);
-    float const cos_arr[8] = {c0.x, c0.y, c0.z, c0.w, c1.x, c1.y, c1.z, c1.w};
-    float const sin_arr[8] = {s0.x, s0.y, s0.z, s0.w, s1.x, s1.y, s1.z, s1.w};
-
-#pragma unroll
-    for (int p = 0; p < kElemsPerLane / 2; p++) {
-      float const x_even = elements[2 * p];
-      float const x_odd = elements[2 * p + 1];
-      elements[2 * p] = x_even * cos_arr[p] - x_odd * sin_arr[p];
-      elements[2 * p + 1] = x_even * sin_arr[p] + x_odd * cos_arr[p];
-    }
-  }
-
-  // ═══════════════════════════════════════════════════════════════════
-  // Q / KV branch dispatch. Restructured as if/else (no early `return`)
-  // so every code path lands at the same exit point — callers own PDL
-  // triggering and per-iteration buffer rotation.
-  // ═══════════════════════════════════════════════════════════════════
-  if (!isKV) {
-    // ── Q: cast back to bf16 and store. ────────────────────────────
-    uint4 out0, out1;
-    typename Converter::packed_hip_type* po0 =
-        reinterpret_cast<typename Converter::packed_hip_type*>(&out0);
-    typename Converter::packed_hip_type* po1 =
-        reinterpret_cast<typename Converter::packed_hip_type*>(&out1);
-#pragma unroll
-    for (int i = 0; i < 4; i++) {
-      po0[i] =
-          Converter::convert(make_float2(elements[2 * i], elements[2 * i + 1]));
-    }
-#pragma unroll
-    for (int i = 0; i < 4; i++) {
-      po1[i] = Converter::convert(
-          make_float2(elements[8 + 2 * i], elements[8 + 2 * i + 1]));
-    }
-    scalar_t_in* dst =
-        q_inout +
-        (static_cast<int64_t>(tokenIdx) * num_heads_q + slotIdx) * kHeadDim +
-        dim_base;
-    *reinterpret_cast<uint4*>(dst) = out0;
-    *reinterpret_cast<uint4*>(dst + 8) = out1;
-  } else {
-    // ── KV: FP8 quant on NoPE + bf16 store on RoPE + cache insert.
-    int64_t const slot_id = slot_mapping[tokenIdx];
-    if (slot_id >= 0) {
-      int64_t const block_idx = slot_id / cache_block_size;
-      int64_t const pos_in_block = slot_id % cache_block_size;
-      uint8_t* block_base =
-          k_cache + block_idx * static_cast<int64_t>(kv_block_stride);
-      uint8_t* token_fp8_ptr = block_base + pos_in_block * kTokenDataBytes;
-      uint8_t* token_bf16_ptr = token_fp8_ptr + kNopeDim;
-      uint8_t* token_scale_ptr =
-          block_base +
-          static_cast<int64_t>(cache_block_size) * kTokenDataBytes +
-          pos_in_block * kScaleBytesPerToken;
-
-#pragma unroll
-      for (int i = 0; i < kElemsPerLane; i++) {
-        elements[i] = Converter::convert(Converter::convert(elements[i]));
-      }
-
-      float local_absmax = 0.0f;
-#pragma unroll
-      for (int i = 0; i < kElemsPerLane; i++) {
-        local_absmax = fmaxf(local_absmax, fabsf(elements[i]));
-      }
-      float const absmax = fmaxf(warp4MaxAbs(local_absmax), 1e-4f);
-      float const exponent = ceilf(log2f(absmax / kFp8Max));
-      float const inv_scale = exp2f(-exponent);
-
-      if (!is_rope_lane) {
-        uint8_t out_bytes[kElemsPerLane];
-#pragma unroll
-        for (int i = 0; i < kElemsPerLane; i++) {
-          float scaled = elements[i] * inv_scale;
-          scaled = fminf(fmaxf(scaled, -kFp8Max), kFp8Max);
-#ifndef USE_ROCM
-          __nv_fp8_storage_t s =
-              __nv_cvt_float_to_fp8(scaled, __NV_SATFINITE, __NV_E4M3);
-          out_bytes[i] = static_cast<uint8_t>(s);
-#else
-          out_bytes[i] = rocm_cvt_float_to_fp8_e4m3(scaled);
-#endif
-        }
-        *reinterpret_cast<uint4*>(token_fp8_ptr + dim_base) =
-            *reinterpret_cast<uint4 const*>(out_bytes);
-
-        if ((laneId & 3) == 0) {
-          int const q_block_idx = laneId >> 2;
-          float encoded = fmaxf(fminf(exponent + 127.0f, 255.0f), 0.0f);
-          token_scale_ptr[q_block_idx] = static_cast<uint8_t>(encoded);
-        }
-        if (laneId == 0) {
-          token_scale_ptr[kNumQuantBlocks] = 0;
-        }
-      } else {
-        uint4 out0, out1;
-        typename Converter::packed_hip_type* po0 =
-            reinterpret_cast<typename Converter::packed_hip_type*>(&out0);
-        typename Converter::packed_hip_type* po1 =
-            reinterpret_cast<typename Converter::packed_hip_type*>(&out1);
-#pragma unroll
-        for (int i = 0; i < 4; i++) {
-          po0[i] = Converter::convert(
-              make_float2(elements[2 * i], elements[2 * i + 1]));
-        }
-#pragma unroll
-        for (int i = 0; i < 4; i++) {
-          po1[i] = Converter::convert(
-              make_float2(elements[8 + 2 * i], elements[8 + 2 * i + 1]));
-        }
-        int const rope_local_base = dim_base - kNopeDim;
-        scalar_t_in* bf16_dst =
-            reinterpret_cast<scalar_t_in*>(token_bf16_ptr) + rope_local_base;
-        *reinterpret_cast<uint4*>(bf16_dst) = out0;
-        *reinterpret_cast<uint4*>(bf16_dst + 8) = out1;
-      }
-    }
-  }
-}
-
 // ────────────────────────────────────────────────────────────────────────────
 // Kernel
 // ────────────────────────────────────────────────────────────────────────────
@@ -346,6 +149,8 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
     return;
   } else {
 #endif
+    using Converter = vllm::_typeConvert<scalar_t_in>;
+
     int const warpsPerBlock = blockDim.x / 32;
     int const warpId = threadIdx.x / 32;
     int const laneId = threadIdx.x % 32;
@@ -371,8 +176,10 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
     // Dim range this lane owns within the 512-wide head.
     int const dim_base = laneId * kElemsPerLane;  // in [0, 512) step 16
 
-    // Two 16-byte loads per thread (8 bf16 each).  Use uint4 as the vector
-    // type; the shared per-slot helper bitcasts to scalar_t_in packed pairs.
+    // ── Load 16 bf16 → 16 fp32 registers (one 16-byte + one 16-byte LDG) ────
+    float elements[kElemsPerLane];
+    float sumOfSquares = 0.0f;
+
     scalar_t_in const* src_ptr;
     if (isKV) {
       src_ptr = kv_in + static_cast<int64_t>(tokenIdx) * kHeadDim + dim_base;
@@ -382,103 +189,196 @@ __global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel(
           dim_base;
       src_ptr = q_inout + q_row_offset;
     }
-    uint4 const v0 = *reinterpret_cast<uint4 const*>(src_ptr);
-    uint4 const v1 = *reinterpret_cast<uint4 const*>(src_ptr + 8);
 
-    processDeepseekV4Slot<scalar_t_in>(
-        v0, v1, tokenIdx, slotIdx, dim_base, laneId, num_heads_q, eps, q_inout,
-        k_cache, slot_mapping, position_ids, cos_sin_cache, cache_block_size,
-        kv_block_stride);
+    // Two 16-byte loads per thread (8 bf16 each).  Use uint4 as the vector
+    // type and bitcast to scalar_t_in packed pairs for conversion.
+    uint4 v0 = *reinterpret_cast<uint4 const*>(src_ptr);
+    uint4 v1 = *reinterpret_cast<uint4 const*>(src_ptr + 8);
 
-#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-    cudaTriggerProgrammaticLaunchCompletion();
-#endif
-#if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
-  }
-#endif
-}
-
-// ────────────────────────────────────────────────────────────────────────────
-// Kernel
-// ────────────────────────────────────────────────────────────────────────────
-//
-// Grid: 1D, gridDim.x = num_tokens_full
-// Block: blockDim.x = 256 threads (8 warps per block) Each
-// warp handles one token, iterating over each head.
-// Q branch (RMSNorm + RoPE, in place) head_slot == num_heads_q
-// KV branch (RoPE + UE8M0 quant + insert)
-//
-template <typename scalar_t_in>
-__global__ void fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid(
-    scalar_t_in* __restrict__ q_inout,  // [N, H, 512] bf16, in place
-    scalar_t_in const* __restrict__ kv_in, uint8_t* __restrict__ k_cache,
-    int64_t const* __restrict__ slot_mapping,
-    int64_t const* __restrict__ position_ids,
-    float const* __restrict__ cos_sin_cache, float const eps,
-    int const num_tokens_full, int const num_tokens_insert,
-    int const num_heads_q, int const cache_block_size,
-    int const kv_block_stride) {
-#if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
-  if constexpr (std::is_same_v<scalar_t_in, c10::BFloat16>) {
-    return;
-  } else {
-#endif
-    int const warpsPerBlock = blockDim.x / 32;
-    int const warpId = threadIdx.x / 32;
-    int const laneId = threadIdx.x % 32;
-
-    int const tokenIdx = blockIdx.x;
-    if (tokenIdx >= num_tokens_full) return;
-
-#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-    cudaGridDependencySynchronize();
-#endif
-
-    int const dim_base = laneId * kElemsPerLane;  // in [0, 512) step 16
-    int const slot_end =
-        (tokenIdx >= num_tokens_insert) ? num_heads_q : (num_heads_q + 1);
-
-    auto src_for_slot = [&](int s) -> scalar_t_in const* {
-      if (s == num_heads_q) {
-        return kv_in + static_cast<int64_t>(tokenIdx) * kHeadDim + dim_base;
+    {
+      typename Converter::packed_hip_type const* p0 =
+          reinterpret_cast<typename Converter::packed_hip_type const*>(&v0);
+      typename Converter::packed_hip_type const* p1 =
+          reinterpret_cast<typename Converter::packed_hip_type const*>(&v1);
+// Each packed_hip_type holds 2 bf16 → 4 packed = 8 elems per uint4.
+#pragma unroll
+      for (int i = 0; i < 4; i++) {
+        float2 f2 = Converter::convert(p0[i]);
+        elements[2 * i] = f2.x;
+        elements[2 * i + 1] = f2.y;
       }
-      return q_inout +
-             (static_cast<int64_t>(tokenIdx) * num_heads_q +
-              static_cast<int64_t>(s)) *
-                 kHeadDim +
-             dim_base;
-    };
+#pragma unroll
+      for (int i = 0; i < 4; i++) {
+        float2 f2 = Converter::convert(p1[i]);
+        elements[8 + 2 * i] = f2.x;
+        elements[8 + 2 * i + 1] = f2.y;
+      }
+    }
 
-    if (warpId < slot_end) {
-      int curr_slot = warpId;
-      scalar_t_in const* src_curr = src_for_slot(curr_slot);
-      uint4 v0_curr = *reinterpret_cast<uint4 const*>(src_curr);
-      uint4 v1_curr = *reinterpret_cast<uint4 const*>(src_curr + 8);
+    // ── Q branch: RMSNorm with no weight (has_weight=False) ─────────────────
+    // Variance + rsqrt + multiply all in fp32, no intermediate bf16 round.
+    // The downstream bf16 round only happens at the final store.
+    if (!isKV) {
+#pragma unroll
+      for (int i = 0; i < kElemsPerLane; i++) {
+        sumOfSquares += elements[i] * elements[i];
+      }
+      sumOfSquares = warpSum<float>(sumOfSquares);
+      float const rms_rcp =
+          rsqrtf(sumOfSquares / static_cast<float>(kHeadDim) + eps);
+#pragma unroll
+      for (int i = 0; i < kElemsPerLane; i++) {
+        elements[i] = elements[i] * rms_rcp;
+      }
+    }
 
-      while (curr_slot < slot_end) {
-        int const next_slot = curr_slot + warpsPerBlock;
-        bool const has_next = (next_slot < slot_end);
+    // ── GPT-J RoPE on dims [NOPE_DIM, HEAD_DIM) ─────────────────────────────
+    // All math in fp32.  cos_sin_cache is loaded as fp32 (its native storage).
+    bool const is_rope_lane = dim_base >= kNopeDim;
+    if (is_rope_lane) {
+      int64_t const pos = position_ids[tokenIdx];
+      constexpr int kHalfRope = kRopeDim / 2;  // 32
+      float const* cos_ptr = cos_sin_cache + pos * kRopeDim;
+      float const* sin_ptr = cos_ptr + kHalfRope;
 
-        // Prefetch src for the next slot
-        uint4 v0_next, v1_next;
-        if (has_next) {
-          scalar_t_in const* src_next = src_for_slot(next_slot);
-          v0_next = *reinterpret_cast<uint4 const*>(src_next);
-          v1_next = *reinterpret_cast<uint4 const*>(src_next + 8);
-        }
+      int const rope_local_base = dim_base - kNopeDim;  // in [0, 64) step 16
+#pragma unroll
+      for (int p = 0; p < kElemsPerLane / 2; p++) {
+        int const pair_dim = rope_local_base + 2 * p;
+        int const half_idx = pair_dim / 2;
+        float const cos_v = VLLM_LDG(cos_ptr + half_idx);
+        float const sin_v = VLLM_LDG(sin_ptr + half_idx);
+        float const x_even = elements[2 * p];
+        float const x_odd = elements[2 * p + 1];
+        elements[2 * p] = x_even * cos_v - x_odd * sin_v;
+        elements[2 * p + 1] = x_even * sin_v + x_odd * cos_v;
+      }
+    }
 
-        processDeepseekV4Slot<scalar_t_in>(
-            v0_curr, v1_curr, tokenIdx, curr_slot, dim_base, laneId,
-            num_heads_q, eps, q_inout, k_cache, slot_mapping, position_ids,
-            cos_sin_cache, cache_block_size, kv_block_stride);
+    // ═══════════════════════════════════════════════════════════════════════
+    // Q branch: cast to bf16 and store back in place.
+    // ═══════════════════════════════════════════════════════════════════════
+    if (!isKV) {
+      uint4 out0, out1;
+      typename Converter::packed_hip_type* po0 =
+          reinterpret_cast<typename Converter::packed_hip_type*>(&out0);
+      typename Converter::packed_hip_type* po1 =
+          reinterpret_cast<typename Converter::packed_hip_type*>(&out1);
+#pragma unroll
+      for (int i = 0; i < 4; i++) {
+        po0[i] = Converter::convert(
+            make_float2(elements[2 * i], elements[2 * i + 1]));
+      }
+#pragma unroll
+      for (int i = 0; i < 4; i++) {
+        po1[i] = Converter::convert(
+            make_float2(elements[8 + 2 * i], elements[8 + 2 * i + 1]));
+      }
+      scalar_t_in* dst =
+          q_inout +
+          (static_cast<int64_t>(tokenIdx) * num_heads_q + slotIdx) * kHeadDim +
+          dim_base;
+      *reinterpret_cast<uint4*>(dst) = out0;
+      *reinterpret_cast<uint4*>(dst + 8) = out1;
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+      cudaTriggerProgrammaticLaunchCompletion();
+#endif
+      return;
+    }
 
-        // ── Buffer rotation: hand the prefetched LDGs to the next iter.
-        v0_curr = v0_next;
-        v1_curr = v1_next;
-        curr_slot = next_slot;
-      }  // while
-    }  // if (warpId < slot_end)
+    // ═══════════════════════════════════════════════════════════════════════
+    // KV branch.
+    // ═══════════════════════════════════════════════════════════════════════
+    int64_t const slot_id = slot_mapping[tokenIdx];
+    if (slot_id < 0) {
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+      cudaTriggerProgrammaticLaunchCompletion();
+#endif
+      return;
+    }
 
+    int64_t const block_idx = slot_id / cache_block_size;
+    int64_t const pos_in_block = slot_id % cache_block_size;
+    uint8_t* block_base =
+        k_cache + block_idx * static_cast<int64_t>(kv_block_stride);
+    uint8_t* token_fp8_ptr = block_base + pos_in_block * kTokenDataBytes;
+    uint8_t* token_bf16_ptr = token_fp8_ptr + kNopeDim;
+    uint8_t* token_scale_ptr =
+        block_base + static_cast<int64_t>(cache_block_size) * kTokenDataBytes +
+        pos_in_block * kScaleBytesPerToken;
+
+    // Round K to bf16 first, matching the unfused reference path where K is
+    // materialized as bf16 before K quantization.  absmax, clamp, and FP8
+    // quant below all run on these bf16-rounded values.
+#pragma unroll
+    for (int i = 0; i < kElemsPerLane; i++) {
+      elements[i] = Converter::convert(Converter::convert(elements[i]));
+    }
+
+    // Per-quant-block absmax must be computed by ALL 32 lanes (warp-collective
+    // shuffle requires full participation).  RoPE lanes contribute garbage,
+    // but their values are gated out below via `!is_rope_lane`.
+    float local_absmax = 0.0f;
+#pragma unroll
+    for (int i = 0; i < kElemsPerLane; i++) {
+      local_absmax = fmaxf(local_absmax, fabsf(elements[i]));
+    }
+    float const absmax = fmaxf(warp4MaxAbs(local_absmax), 1e-4f);
+    float const exponent = ceilf(log2f(absmax / kFp8Max));
+    float const inv_scale = exp2f(-exponent);
+
+    if (!is_rope_lane) {
+      // ── NoPE lane: UE8M0 FP8 quant ───────────────────────────────────────
+      uint8_t out_bytes[kElemsPerLane];
+#pragma unroll
+      for (int i = 0; i < kElemsPerLane; i++) {
+        float scaled = elements[i] * inv_scale;
+        scaled = fminf(fmaxf(scaled, -kFp8Max), kFp8Max);
+#ifndef USE_ROCM
+        __nv_fp8_storage_t s =
+            __nv_cvt_float_to_fp8(scaled, __NV_SATFINITE, __NV_E4M3);
+        out_bytes[i] = static_cast<uint8_t>(s);
+#else
+      out_bytes[i] = rocm_cvt_float_to_fp8_e4m3(scaled);
+#endif
+      }
+      // One 16-byte STG per lane.
+      *reinterpret_cast<uint4*>(token_fp8_ptr + dim_base) =
+          *reinterpret_cast<uint4 const*>(out_bytes);
+
+      // Lane (4k) of each 4-lane group writes the scale byte for block k<7.
+      if ((laneId & 3) == 0) {
+        int const q_block_idx = laneId >> 2;  // 0..6 for NoPE lanes
+        float encoded = fmaxf(fminf(exponent + 127.0f, 255.0f), 0.0f);
+        token_scale_ptr[q_block_idx] = static_cast<uint8_t>(encoded);
+      }
+      // Lane 0 also writes the padding byte at index 7.
+      if (laneId == 0) {
+        token_scale_ptr[kNumQuantBlocks] = 0;  // pad
+      }
+    } else {
+      // ── RoPE lane: cast back to bf16 and store to cache bf16 tail ────────
+      uint4 out0, out1;
+      typename Converter::packed_hip_type* po0 =
+          reinterpret_cast<typename Converter::packed_hip_type*>(&out0);
+      typename Converter::packed_hip_type* po1 =
+          reinterpret_cast<typename Converter::packed_hip_type*>(&out1);
+#pragma unroll
+      for (int i = 0; i < 4; i++) {
+        po0[i] = Converter::convert(
+            make_float2(elements[2 * i], elements[2 * i + 1]));
+      }
+#pragma unroll
+      for (int i = 0; i < 4; i++) {
+        po1[i] = Converter::convert(
+            make_float2(elements[8 + 2 * i], elements[8 + 2 * i + 1]));
+      }
+      int const rope_local_base = dim_base - kNopeDim;  // in [0, 64)
+      scalar_t_in* bf16_dst =
+          reinterpret_cast<scalar_t_in*>(token_bf16_ptr) + rope_local_base;
+      *reinterpret_cast<uint4*>(bf16_dst) = out0;
+      *reinterpret_cast<uint4*>(bf16_dst + 8) = out1;
+    }
 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
     cudaTriggerProgrammaticLaunchCompletion();
 #endif
@@ -530,22 +430,11 @@ void launchFusedDeepseekV4QNormRopeKVRopeQuantInsert(
   config.attrs = attrs;
   config.numAttrs = (sm_version >= 90) ? 1 : 0;
 
-  if (num_tokens_full < NUM_TOKEN_CUTOFF) {
-    cudaLaunchKernelEx(
-        &config, fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel<scalar_t_in>,
-        q_inout, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache, eps,
-        num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
-        kv_block_stride);
-  } else {
-    config.gridDim = dim3(num_tokens_full);
-    cudaLaunchKernelEx(
-        &config,
-        fusedDeepseekV4QNormRopeKVRopeQuantInsertKernelReducedGrid<scalar_t_in>,
-        q_inout, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache, eps,
-        num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
-        kv_block_stride);
-  }
-
+  cudaLaunchKernelEx(
+      &config, fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel<scalar_t_in>,
+      q_inout, kv_in, k_cache, slot_mapping, position_ids, cos_sin_cache, eps,
+      num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
+      kv_block_stride);
 #else
   // ROCm: use standard kernel launch syntax (no PDL/stream serialization)
   // clang-format off
@@ -619,4 +508,4 @@ void fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert(
                 num_tokens_full, num_tokens_insert, num_heads_q,
                 cache_block_size_i, kv_block_stride, stream);
       });
-}
+}

csrc/fused_qknorm_rope_kernel.cu

@@ -18,20 +18,21 @@
 #include <cuda_runtime.h>
 #include <type_traits>
 
-#include "torch_utils.h"
+#include <torch/cuda.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
 
-#include "../async_util.cuh"
-#include "../cuda_compat.h"
-#include "../type_convert.cuh"
+#include "async_util.cuh"
+#include "cuda_compat.h"
 #include "dispatch_utils.h"
+#include "type_convert.cuh"
 
-#define CHECK_TYPE(x, st)                                                  \
-  STD_TORCH_CHECK(x.scalar_type() == st, #x " dtype is ", x.scalar_type(), \
-                  ", while ", st, " is expected")
-#define CHECK_TH_CUDA(x) \
-  STD_TORCH_CHECK(x.is_cuda(), #x " must be a CUDA tensor")
+#define CHECK_TYPE(x, st)                                              \
+  TORCH_CHECK(x.scalar_type() == st, #x " dtype is ", x.scalar_type(), \
+              ", while ", st, " is expected")
+#define CHECK_TH_CUDA(x) TORCH_CHECK(x.is_cuda(), #x " must be a CUDA tensor")
 #define CHECK_CONTIGUOUS(x) \
-  STD_TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
+  TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
 #define CHECK_INPUT(x) \
   CHECK_TH_CUDA(x);    \
   CHECK_CONTIGUOUS(x)
@@ -588,8 +589,8 @@ void launchFusedQKNormRope(void* qkv, int const num_tokens,
       });
       break;
     default:
-      STD_TORCH_CHECK(
-          false, "Unsupported head dimension for fusedQKNormRope: ", head_dim);
+      TORCH_CHECK(false,
+                  "Unsupported head dimension for fusedQKNormRope: ", head_dim);
   }
 }
 
@@ -603,10 +604,10 @@ void launchFusedQKNormRopeNTokenHeads(
     void const* k_weight, void const* cos_sin_cache, bool const interleave,
     int64_t const* position_ids, int const token_heads_per_warp,
     cudaStream_t stream) {
-  STD_TORCH_CHECK(token_heads_per_warp == 1 || token_heads_per_warp == 2 ||
-                      token_heads_per_warp == 4 || token_heads_per_warp == 8,
-                  "token_heads_per_warp must be 1, 2, 4, or 8, got ",
-                  token_heads_per_warp);
+  TORCH_CHECK(token_heads_per_warp == 1 || token_heads_per_warp == 2 ||
+                  token_heads_per_warp == 4 || token_heads_per_warp == 8,
+              "token_heads_per_warp must be 1, 2, 4, or 8, got ",
+              token_heads_per_warp);
 
   // token_heads_per_warp == 1: delegate to the 1-head baseline kernel.
   if (token_heads_per_warp == 1) {
@@ -690,7 +691,7 @@ void launchFusedQKNormRopeNTokenHeads(
         });                                                                  \
         break;                                                               \
       default:                                                               \
-        STD_TORCH_CHECK(false, "Unsupported head dimension: ", head_dim);    \
+        TORCH_CHECK(false, "Unsupported head dimension: ", head_dim);        \
     }                                                                        \
   } while (0)
 
@@ -707,21 +708,19 @@ void launchFusedQKNormRopeNTokenHeads(
 }  // namespace tensorrt_llm::kernels
 
 void fused_qk_norm_rope(
-    torch::stable::Tensor&
-        qkv,              // Combined QKV tensor [num_tokens,
-                          // (num_heads_q+num_heads_k+num_heads_v)*head_dim]
-    int64_t num_heads_q,  // Number of query heads
-    int64_t num_heads_k,  // Number of key heads
-    int64_t num_heads_v,  // Number of value heads
-    int64_t head_dim,     // Dimension per head
-    double eps,           // Epsilon for RMS normalization
-    torch::stable::Tensor& q_weight,  // RMSNorm weights for query [head_dim]
-    torch::stable::Tensor& k_weight,  // RMSNorm weights for key [head_dim]
-    torch::stable::Tensor& cos_sin_cache,  // Cos/sin cache [max_position,
-                                           // head_dim]
-    bool is_neox,  // Whether RoPE is applied in Neox style
-    torch::stable::Tensor& position_ids,  // Position IDs for RoPE [num_tokens]
-    int64_t forced_token_heads_per_warp   // -1 = auto-select, >0 = forced value
+    torch::Tensor& qkv,       // Combined QKV tensor [num_tokens,
+                              // (num_heads_q+num_heads_k+num_heads_v)*head_dim]
+    int64_t num_heads_q,      // Number of query heads
+    int64_t num_heads_k,      // Number of key heads
+    int64_t num_heads_v,      // Number of value heads
+    int64_t head_dim,         // Dimension per head
+    double eps,               // Epsilon for RMS normalization
+    torch::Tensor& q_weight,  // RMSNorm weights for query [head_dim]
+    torch::Tensor& k_weight,  // RMSNorm weights for key [head_dim]
+    torch::Tensor& cos_sin_cache,  // Cos/sin cache [max_position, head_dim]
+    bool is_neox,                  // Whether RoPE is applied in Neox style
+    torch::Tensor& position_ids,   // Position IDs for RoPE [num_tokens]
+    int64_t forced_token_heads_per_warp  // -1 = auto-select, >0 = forced value
 ) {
   // Input validation
   CHECK_INPUT(qkv);
@@ -729,42 +728,40 @@ void fused_qk_norm_rope(
   CHECK_INPUT(q_weight);
   CHECK_INPUT(k_weight);
   CHECK_INPUT(cos_sin_cache);
-  CHECK_TYPE(position_ids, torch::headeronly::ScalarType::Long);
+  CHECK_TYPE(position_ids, torch::kInt64);
 
-  STD_TORCH_CHECK(qkv.dim() == 2,
-                  "QKV tensor must be 2D: [num_tokens, "
-                  "(num_heads_q+num_heads_k+num_heads_v)*head_dim]");
-  STD_TORCH_CHECK(position_ids.dim() == 1,
-                  "Position IDs must be 1D: [num_tokens]");
-  STD_TORCH_CHECK(q_weight.dim() == 1, "Query weights must be 1D: [head_dim]");
-  STD_TORCH_CHECK(k_weight.dim() == 1, "Key weights must be 1D: [head_dim]");
-  STD_TORCH_CHECK(cos_sin_cache.dim() == 2,
-                  "Cos/sin cache must be 2D: [max_position, head_dim]");
-  STD_TORCH_CHECK(q_weight.size(0) == head_dim,
-                  "Query weights size must match head dimension");
-  STD_TORCH_CHECK(k_weight.size(0) == head_dim,
-                  "Key weights size must match head dimension");
+  TORCH_CHECK(qkv.dim() == 2,
+              "QKV tensor must be 2D: [num_tokens, "
+              "(num_heads_q+num_heads_k+num_heads_v)*head_dim]");
+  TORCH_CHECK(position_ids.dim() == 1, "Position IDs must be 1D: [num_tokens]");
+  TORCH_CHECK(q_weight.dim() == 1, "Query weights must be 1D: [head_dim]");
+  TORCH_CHECK(k_weight.dim() == 1, "Key weights must be 1D: [head_dim]");
+  TORCH_CHECK(cos_sin_cache.dim() == 2,
+              "Cos/sin cache must be 2D: [max_position, head_dim]");
+  TORCH_CHECK(q_weight.size(0) == head_dim,
+              "Query weights size must match head dimension");
+  TORCH_CHECK(k_weight.size(0) == head_dim,
+              "Key weights size must match head dimension");
 
-  STD_TORCH_CHECK(cos_sin_cache.size(1) % 2 == 0, "rotary_dim must be even");
-  STD_TORCH_CHECK(cos_sin_cache.size(1) <= head_dim,
-                  "rotary_dim must be less than or equal to head_dim");
+  TORCH_CHECK(cos_sin_cache.size(1) % 2 == 0, "rotary_dim must be even");
+  TORCH_CHECK(cos_sin_cache.size(1) <= head_dim,
+              "rotary_dim must be less than or equal to head_dim");
 
-  STD_TORCH_CHECK(qkv.scalar_type() == q_weight.scalar_type() &&
-                      qkv.scalar_type() == k_weight.scalar_type(),
-                  "qkv, q_weight and k_weight must have the same dtype");
+  TORCH_CHECK(qkv.scalar_type() == q_weight.scalar_type() &&
+                  qkv.scalar_type() == k_weight.scalar_type(),
+              "qkv, q_weight and k_weight must have the same dtype");
 
   int64_t num_tokens = qkv.size(0);
-  STD_TORCH_CHECK(position_ids.size(0) == num_tokens,
-                  "Number of tokens in position_ids must match QKV");
+  TORCH_CHECK(position_ids.size(0) == num_tokens,
+              "Number of tokens in position_ids must match QKV");
 
   int64_t total_heads = num_heads_q + num_heads_k + num_heads_v;
-  STD_TORCH_CHECK(
+  TORCH_CHECK(
       qkv.size(1) == total_heads * head_dim,
       "QKV tensor size must match total number of heads and head dimension");
 
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      qkv.get_device_index());
-  auto stream = get_current_cuda_stream(qkv.get_device_index());
+  auto device_id = qkv.get_device();
+  auto stream = at::cuda::getCurrentCUDAStream(device_id);
 
   // Select token_heads_per_warp: forced value if >0, else auto-select.
   // Auto thresholds are calibrated on SM 9.0 (H100). On other architectures,
@@ -774,7 +771,8 @@ void fused_qk_norm_rope(
     token_heads_per_warp = static_cast<int>(forced_token_heads_per_warp);
   } else {
     token_heads_per_warp = 1;
-    int sm_version = get_device_prop()->major * 10 + get_device_prop()->minor;
+    auto* dev_prop = at::cuda::getDeviceProperties(device_id);
+    int sm_version = dev_prop->major * 10 + dev_prop->minor;
     int64_t total_qk_units = num_tokens * (num_heads_q + num_heads_k);
     if (sm_version == 90) {
       if (head_dim >= 256) {
@@ -797,22 +795,21 @@ void fused_qk_norm_rope(
     }
   }
 
-  VLLM_STABLE_DISPATCH_HALF_TYPES(
-      qkv.scalar_type(), "fused_qk_norm_rope_kernel", [&] {
-        using qkv_scalar_t = scalar_t;
-        VLLM_STABLE_DISPATCH_FLOATING_TYPES(
-            cos_sin_cache.scalar_type(), "fused_qk_norm_rope_kernel", [&] {
-              using cache_scalar_t = scalar_t;
-              tensorrt_llm::kernels::launchFusedQKNormRopeNTokenHeads<
-                  qkv_scalar_t, cache_scalar_t>(
-                  qkv.data_ptr(), static_cast<int>(num_tokens),
-                  static_cast<int>(num_heads_q), static_cast<int>(num_heads_k),
-                  static_cast<int>(num_heads_v), static_cast<int>(head_dim),
-                  static_cast<int>(cos_sin_cache.size(1)),
-                  static_cast<float>(eps), q_weight.data_ptr(),
-                  k_weight.data_ptr(), cos_sin_cache.data_ptr(), !is_neox,
-                  reinterpret_cast<int64_t const*>(position_ids.data_ptr()),
-                  token_heads_per_warp, stream);
-            });
-      });
+  VLLM_DISPATCH_HALF_TYPES(qkv.scalar_type(), "fused_qk_norm_rope_kernel", [&] {
+    using qkv_scalar_t = scalar_t;
+    VLLM_DISPATCH_FLOATING_TYPES(
+        cos_sin_cache.scalar_type(), "fused_qk_norm_rope_kernel", [&] {
+          using cache_scalar_t = scalar_t;
+          tensorrt_llm::kernels::launchFusedQKNormRopeNTokenHeads<
+              qkv_scalar_t, cache_scalar_t>(
+              qkv.data_ptr(), static_cast<int>(num_tokens),
+              static_cast<int>(num_heads_q), static_cast<int>(num_heads_k),
+              static_cast<int>(num_heads_v), static_cast<int>(head_dim),
+              static_cast<int>(cos_sin_cache.size(1)), static_cast<float>(eps),
+              q_weight.data_ptr(), k_weight.data_ptr(),
+              cos_sin_cache.data_ptr(), !is_neox,
+              reinterpret_cast<int64_t const*>(position_ids.data_ptr()),
+              token_heads_per_warp, stream);
+        });
+  });
 }

csrc/layernorm_kernels.cu

@@ -1,12 +1,11 @@
-#include <numeric>
-
-#include "torch_utils.h"
-
-#include "../cub_helpers.h"
-#include "../core/batch_invariant.hpp"
-#include "../type_convert.cuh"
+#include "type_convert.cuh"
 #include "dispatch_utils.h"
-#include "quantization/vectorization_utils.cuh"
+#include "cub_helpers.h"
+#include "core/batch_invariant.hpp"
+#include "libtorch_stable/quantization/vectorization_utils.cuh"
+
+#include <torch/cuda.h>
+#include <c10/cuda/CUDAGuard.h>
 
 namespace vllm {
 
@@ -190,16 +189,16 @@ fused_add_rms_norm_kernel(
 
 }  // namespace vllm
 
-void rms_norm(torch::stable::Tensor& out,     // [..., hidden_size]
-              torch::stable::Tensor& input,   // [..., hidden_size]
-              torch::stable::Tensor& weight,  // [hidden_size]
+void rms_norm(torch::Tensor& out,     // [..., hidden_size]
+              torch::Tensor& input,   // [..., hidden_size]
+              torch::Tensor& weight,  // [hidden_size]
               double epsilon) {
-  STD_TORCH_CHECK(out.is_contiguous());
+  TORCH_CHECK(out.is_contiguous());
   if (input.stride(-1) != 1) {
-    input = torch::stable::contiguous(input);
+    input = input.contiguous();
   }
-  STD_TORCH_CHECK(input.stride(-1) == 1);
-  STD_TORCH_CHECK(weight.is_contiguous());
+  TORCH_CHECK(input.stride(-1) == 1);
+  TORCH_CHECK(weight.is_contiguous());
 
   int hidden_size = input.size(-1);
 
@@ -214,49 +213,45 @@ void rms_norm(torch::stable::Tensor& out,     // [..., hidden_size]
   // For large num_tokens, use smaller blocks to increase SM concurrency.
   const int max_block_size = (num_tokens < 256) ? 1024 : 256;
   dim3 grid(num_tokens);
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      input.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
-  VLLM_STABLE_DISPATCH_RANK234(num_dims, [&] {
-    VLLM_STABLE_DISPATCH_FLOATING_TYPES(
-        input.scalar_type(), "rms_norm_kernel", [&] {
-          const int calculated_vec_size =
-              std::gcd(16 / sizeof(scalar_t), hidden_size);
-          const int block_size =
-              std::min(hidden_size / calculated_vec_size, max_block_size);
-          dim3 block(block_size);
-          VLLM_STABLE_DISPATCH_VEC_SIZE(calculated_vec_size, [&] {
-            vllm::rms_norm_kernel<scalar_t, vec_size, tensor_rank>
-                <<<grid, block, 0, stream>>>(
-                    out.mutable_data_ptr<scalar_t>(),
-                    input.const_data_ptr<scalar_t>(), input_stride_d2,
-                    input_stride_d3, input_stride_d4, input_shape_d2,
-                    input_shape_d3, weight.const_data_ptr<scalar_t>(), epsilon,
-                    num_tokens, hidden_size);
-          });
-        });
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+  VLLM_DISPATCH_RANK234(num_dims, [&] {
+    VLLM_DISPATCH_FLOATING_TYPES(input.scalar_type(), "rms_norm_kernel", [&] {
+      const int calculated_vec_size =
+          std::gcd(16 / sizeof(scalar_t), hidden_size);
+      const int block_size =
+          std::min(hidden_size / calculated_vec_size, max_block_size);
+      dim3 block(block_size);
+      VLLM_DISPATCH_VEC_SIZE(calculated_vec_size, [&] {
+        vllm::rms_norm_kernel<scalar_t, vec_size, tensor_rank>
+            <<<grid, block, 0, stream>>>(
+                out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(),
+                input_stride_d2, input_stride_d3, input_stride_d4,
+                input_shape_d2, input_shape_d3, weight.data_ptr<scalar_t>(),
+                epsilon, num_tokens, hidden_size);
+      });
+    });
   });
 }
 
-#define LAUNCH_FUSED_ADD_RMS_NORM(width)                                \
-  VLLM_STABLE_DISPATCH_FLOATING_TYPES(                                  \
-      input.scalar_type(), "fused_add_rms_norm_kernel", [&] {           \
-        vllm::fused_add_rms_norm_kernel<scalar_t, width>                \
-            <<<grid, block, 0, stream>>>(                               \
-                input.mutable_data_ptr<scalar_t>(), input_stride,       \
-                residual.mutable_data_ptr<scalar_t>(),                  \
-                weight.const_data_ptr<scalar_t>(), epsilon, num_tokens, \
-                hidden_size);                                           \
+#define LAUNCH_FUSED_ADD_RMS_NORM(width)                                    \
+  VLLM_DISPATCH_FLOATING_TYPES(                                             \
+      input.scalar_type(), "fused_add_rms_norm_kernel", [&] {               \
+        vllm::fused_add_rms_norm_kernel<scalar_t, width>                    \
+            <<<grid, block, 0, stream>>>(                                   \
+                input.data_ptr<scalar_t>(), input_stride,                   \
+                residual.data_ptr<scalar_t>(), weight.data_ptr<scalar_t>(), \
+                epsilon, num_tokens, hidden_size);                          \
       });
 
-void fused_add_rms_norm(torch::stable::Tensor& input,     // [..., hidden_size]
-                        torch::stable::Tensor& residual,  // [..., hidden_size]
-                        torch::stable::Tensor& weight,    // [hidden_size]
+void fused_add_rms_norm(torch::Tensor& input,     // [..., hidden_size]
+                        torch::Tensor& residual,  // [..., hidden_size]
+                        torch::Tensor& weight,    // [hidden_size]
                         double epsilon) {
-  STD_TORCH_CHECK(weight.scalar_type() == input.scalar_type());
-  STD_TORCH_CHECK(input.scalar_type() == residual.scalar_type());
-  STD_TORCH_CHECK(residual.is_contiguous());
-  STD_TORCH_CHECK(weight.is_contiguous());
+  TORCH_CHECK(weight.scalar_type() == input.scalar_type());
+  TORCH_CHECK(input.scalar_type() == residual.scalar_type());
+  TORCH_CHECK(residual.is_contiguous());
+  TORCH_CHECK(weight.is_contiguous());
   int hidden_size = input.size(-1);
   int64_t input_stride = input.stride(-2);
   int num_tokens = input.numel() / hidden_size;
@@ -268,9 +263,8 @@ void fused_add_rms_norm(torch::stable::Tensor& input,     // [..., hidden_size]
      hiding on global mem ops. */
   const int max_block_size = (num_tokens < 256) ? 1024 : 256;
   dim3 block(std::min(hidden_size, max_block_size));
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      input.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
   /*If the tensor types are FP16/BF16, try to use the optimized kernel
     with packed + vectorized ops.
     Max optimization is achieved with a width-8 vector of FP16/BF16s

csrc/layernorm_quant_kernels.cu

@@ -5,16 +5,15 @@
  * Currently, only static fp8 quantization is supported.
  */
 
-#include <numeric>
-
-#include "torch_utils.h"
-
-#include "../cub_helpers.h"
-#include "../core/batch_invariant.hpp"
-#include "../quantization/w8a8/fp8/common.cuh"
-#include "../type_convert.cuh"
+#include "type_convert.cuh"
+#include "quantization/w8a8/fp8/common.cuh"
 #include "dispatch_utils.h"
-#include "quantization/vectorization_utils.cuh"
+#include "cub_helpers.h"
+#include "core/batch_invariant.hpp"
+#include "libtorch_stable/quantization/vectorization_utils.cuh"
+
+#include <torch/cuda.h>
+#include <c10/cuda/CUDAGuard.h>
 
 namespace vllm {
 
@@ -203,13 +202,12 @@ fused_add_rms_norm_static_fp8_quant_kernel(
 
 }  // namespace vllm
 
-void rms_norm_static_fp8_quant(
-    torch::stable::Tensor& out,     // [..., hidden_size]
-    torch::stable::Tensor& input,   // [..., hidden_size]
-    torch::stable::Tensor& weight,  // [hidden_size]
-    torch::stable::Tensor& scale,   // [1]
-    double epsilon) {
-  STD_TORCH_CHECK(out.is_contiguous());
+void rms_norm_static_fp8_quant(torch::Tensor& out,     // [..., hidden_size]
+                               torch::Tensor& input,   // [..., hidden_size]
+                               torch::Tensor& weight,  // [hidden_size]
+                               torch::Tensor& scale,   // [1]
+                               double epsilon) {
+  TORCH_CHECK(out.is_contiguous());
   int hidden_size = input.size(-1);
   int input_stride = input.stride(-2);
   int num_tokens = input.numel() / hidden_size;
@@ -217,26 +215,24 @@ void rms_norm_static_fp8_quant(
   // For large num_tokens, use smaller blocks to increase SM concurrency.
   const int max_block_size = (num_tokens < 256) ? 1024 : 256;
   dim3 grid(num_tokens);
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      input.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
-  VLLM_STABLE_DISPATCH_FLOATING_TYPES(
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+  VLLM_DISPATCH_FLOATING_TYPES(
       input.scalar_type(), "rms_norm_kernel_scalar_type", [&] {
-        VLLM_STABLE_DISPATCH_FP8_TYPES(
+        VLLM_DISPATCH_FP8_TYPES(
             out.scalar_type(), "rms_norm_kernel_fp8_type", [&] {
               const int calculated_vec_size =
                   std::gcd(16 / sizeof(scalar_t), hidden_size);
               const int block_size =
                   std::min(hidden_size / calculated_vec_size, max_block_size);
               dim3 block(block_size);
-              VLLM_STABLE_DISPATCH_VEC_SIZE(calculated_vec_size, [&] {
+              VLLM_DISPATCH_VEC_SIZE(calculated_vec_size, [&] {
                 vllm::rms_norm_static_fp8_quant_kernel<scalar_t, fp8_t,
                                                        vec_size>
                     <<<grid, block, 0, stream>>>(
-                        out.mutable_data_ptr<fp8_t>(),
-                        input.const_data_ptr<scalar_t>(), input_stride,
-                        weight.const_data_ptr<scalar_t>(),
-                        scale.const_data_ptr<float>(), epsilon, num_tokens,
+                        out.data_ptr<fp8_t>(), input.data_ptr<scalar_t>(),
+                        input_stride, weight.data_ptr<scalar_t>(),
+                        scale.data_ptr<float>(), epsilon, num_tokens,
                         hidden_size);
               });
             });
@@ -244,32 +240,30 @@ void rms_norm_static_fp8_quant(
 }
 
 #define LAUNCH_FUSED_ADD_RMS_NORM(width)                                     \
-  VLLM_STABLE_DISPATCH_FLOATING_TYPES(                                       \
+  VLLM_DISPATCH_FLOATING_TYPES(                                              \
       input.scalar_type(), "fused_add_rms_norm_kernel_scalar_type", [&] {    \
-        VLLM_STABLE_DISPATCH_FP8_TYPES(                                      \
+        VLLM_DISPATCH_FP8_TYPES(                                             \
             out.scalar_type(), "fused_add_rms_norm_kernel_fp8_type", [&] {   \
               vllm::fused_add_rms_norm_static_fp8_quant_kernel<scalar_t,     \
                                                                width, fp8_t> \
                   <<<grid, block, 0, stream>>>(                              \
-                      out.mutable_data_ptr<fp8_t>(),                         \
-                      input.mutable_data_ptr<scalar_t>(), input_stride,      \
-                      residual.mutable_data_ptr<scalar_t>(),                 \
-                      weight.const_data_ptr<scalar_t>(),                     \
-                      scale.const_data_ptr<float>(), epsilon, num_tokens,    \
-                      hidden_size);                                          \
+                      out.data_ptr<fp8_t>(), input.data_ptr<scalar_t>(),     \
+                      input_stride, residual.data_ptr<scalar_t>(),           \
+                      weight.data_ptr<scalar_t>(), scale.data_ptr<float>(),  \
+                      epsilon, num_tokens, hidden_size);                     \
             });                                                              \
       });
 void fused_add_rms_norm_static_fp8_quant(
-    torch::stable::Tensor& out,       // [..., hidden_size],
-    torch::stable::Tensor& input,     // [..., hidden_size]
-    torch::stable::Tensor& residual,  // [..., hidden_size]
-    torch::stable::Tensor& weight,    // [hidden_size]
-    torch::stable::Tensor& scale,     // [1]
+    torch::Tensor& out,       // [..., hidden_size],
+    torch::Tensor& input,     // [..., hidden_size]
+    torch::Tensor& residual,  // [..., hidden_size]
+    torch::Tensor& weight,    // [hidden_size]
+    torch::Tensor& scale,     // [1]
     double epsilon) {
-  STD_TORCH_CHECK(out.is_contiguous());
-  STD_TORCH_CHECK(residual.is_contiguous());
-  STD_TORCH_CHECK(residual.scalar_type() == input.scalar_type());
-  STD_TORCH_CHECK(weight.scalar_type() == input.scalar_type());
+  TORCH_CHECK(out.is_contiguous());
+  TORCH_CHECK(residual.is_contiguous());
+  TORCH_CHECK(residual.scalar_type() == input.scalar_type());
+  TORCH_CHECK(weight.scalar_type() == input.scalar_type());
   int hidden_size = input.size(-1);
   int input_stride = input.stride(-2);
   int num_tokens = input.numel() / hidden_size;
@@ -281,9 +275,8 @@ void fused_add_rms_norm_static_fp8_quant(
      hiding on global mem ops. */
   const int max_block_size = (num_tokens < 256) ? 1024 : 256;
   dim3 block(std::min(hidden_size, max_block_size));
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      input.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
   /*If the tensor types are FP16/BF16, try to use the optimized kernel
     with packed + vectorized ops.
     Max optimization is achieved with a width-8 vector of FP16/BF16s

csrc/libtorch_stable/dispatch_utils.h

@@ -58,35 +58,6 @@
   THO_DISPATCH_SWITCH(TYPE, NAME,                        \
                       VLLM_STABLE_DISPATCH_CASE_HALF_TYPES(__VA_ARGS__))
 
-// Quant type dispatch (FP8 + INT8)
-#ifdef USE_ROCM
-  #define VLLM_STABLE_DISPATCH_CASE_QUANT_TYPES(...)                  \
-    THO_DISPATCH_CASE(torch::headeronly::ScalarType::Float8_e4m3fn,   \
-                      __VA_ARGS__)                                    \
-    THO_DISPATCH_CASE(torch::headeronly::ScalarType::Float8_e4m3fnuz, \
-                      __VA_ARGS__)                                    \
-    THO_DISPATCH_CASE(torch::headeronly::ScalarType::Char, __VA_ARGS__)
-#else
-  #define VLLM_STABLE_DISPATCH_CASE_QUANT_TYPES(...)                \
-    THO_DISPATCH_CASE(torch::headeronly::ScalarType::Float8_e4m3fn, \
-                      __VA_ARGS__)                                  \
-    THO_DISPATCH_CASE(torch::headeronly::ScalarType::Char, __VA_ARGS__)
-#endif
-
-#define VLLM_STABLE_DISPATCH_QUANT_TYPES(TYPE, NAME, ...) \
-  THO_DISPATCH_SWITCH(TYPE, NAME,                         \
-                      VLLM_STABLE_DISPATCH_CASE_QUANT_TYPES(__VA_ARGS__))
-
-// Group size dispatch (pure C++ if/else, no ATen dependency)
-#define VLLM_STABLE_DISPATCH_GROUP_SIZE(group_size, const_group_size, ...) \
-  if (group_size == 128) {                                                 \
-    constexpr int const_group_size = 128;                                  \
-    __VA_ARGS__();                                                         \
-  } else if (group_size == 64) {                                           \
-    constexpr int const_group_size = 64;                                   \
-    __VA_ARGS__();                                                         \
-  }
-
 // Boolean dispatch
 #define VLLM_STABLE_DISPATCH_BOOL(expr, const_expr, ...) \
   if (expr) {                                            \
@@ -96,56 +67,3 @@
     constexpr bool const_expr = false;                   \
     __VA_ARGS__();                                       \
   }
-
-// Vec size dispatch (pure C++ switch, no ATen dependency)
-#define VLLM_STABLE_DISPATCH_VEC_SIZE(VEC_SIZE, ...) \
-  switch (VEC_SIZE) {                                \
-    case 16: {                                       \
-      constexpr int vec_size = 16;                   \
-      __VA_ARGS__();                                 \
-      break;                                         \
-    }                                                \
-    case 8: {                                        \
-      constexpr int vec_size = 8;                    \
-      __VA_ARGS__();                                 \
-      break;                                         \
-    }                                                \
-    case 4: {                                        \
-      constexpr int vec_size = 4;                    \
-      __VA_ARGS__();                                 \
-      break;                                         \
-    }                                                \
-    case 2: {                                        \
-      constexpr int vec_size = 2;                    \
-      __VA_ARGS__();                                 \
-      break;                                         \
-    }                                                \
-    default: {                                       \
-      constexpr int vec_size = 1;                    \
-      __VA_ARGS__();                                 \
-      break;                                         \
-    }                                                \
-  }
-
-// Tensor rank dispatch (2D, 3D, 4D)
-#define VLLM_STABLE_DISPATCH_RANK234(NUM_DIMS, ...)                          \
-  switch (NUM_DIMS) {                                                        \
-    case 2: {                                                                \
-      constexpr int tensor_rank = 2;                                         \
-      __VA_ARGS__();                                                         \
-      break;                                                                 \
-    }                                                                        \
-    case 3: {                                                                \
-      constexpr int tensor_rank = 3;                                         \
-      __VA_ARGS__();                                                         \
-      break;                                                                 \
-    }                                                                        \
-    case 4: {                                                                \
-      constexpr int tensor_rank = 4;                                         \
-      __VA_ARGS__();                                                         \
-      break;                                                                 \
-    }                                                                        \
-    default:                                                                 \
-      STD_TORCH_CHECK(                                                       \
-          false, "Expects rank 2, 3 or 4 tensors but got unsupported rank"); \
-  }

csrc/libtorch_stable/ops.h

@@ -166,139 +166,3 @@ torch::stable::Tensor awq_dequantize(torch::stable::Tensor _kernel,
 
 torch::stable::Tensor hadacore_transform(torch::stable::Tensor& x,
                                          bool inplace);
-
-// Layernorm kernels (shared CUDA/ROCm)
-void rms_norm(torch::stable::Tensor& out, torch::stable::Tensor& input,
-              torch::stable::Tensor& weight, double epsilon);
-
-void fused_add_rms_norm(torch::stable::Tensor& input,
-                        torch::stable::Tensor& residual,
-                        torch::stable::Tensor& weight, double epsilon);
-
-// Layernorm-quant kernels (shared CUDA/ROCm)
-void rms_norm_static_fp8_quant(torch::stable::Tensor& out,
-                               torch::stable::Tensor& input,
-                               torch::stable::Tensor& weight,
-                               torch::stable::Tensor& scale, double epsilon);
-
-void fused_add_rms_norm_static_fp8_quant(torch::stable::Tensor& out,
-                                         torch::stable::Tensor& input,
-                                         torch::stable::Tensor& residual,
-                                         torch::stable::Tensor& weight,
-                                         torch::stable::Tensor& scale,
-                                         double epsilon);
-
-// Fused layernorm + dynamic per-token quant kernels (shared CUDA/ROCm)
-void rms_norm_dynamic_per_token_quant(
-    torch::stable::Tensor& out, torch::stable::Tensor const& input,
-    torch::stable::Tensor const& weight, torch::stable::Tensor& scales,
-    double const var_epsilon, std::optional<torch::stable::Tensor> scale_ub,
-    std::optional<torch::stable::Tensor> residual);
-
-void rms_norm_per_block_quant(torch::stable::Tensor& out,
-                              torch::stable::Tensor const& input,
-                              torch::stable::Tensor const& weight,
-                              torch::stable::Tensor& scales,
-                              double const var_epsilon,
-                              std::optional<torch::stable::Tensor> scale_ub,
-                              std::optional<torch::stable::Tensor> residual,
-                              int64_t group_size, bool is_scale_transposed);
-
-// Positional encoding kernels (shared CUDA/ROCm)
-void rotary_embedding(torch::stable::Tensor& positions,
-                      torch::stable::Tensor& query,
-                      std::optional<torch::stable::Tensor> key,
-                      int64_t head_size, torch::stable::Tensor& cos_sin_cache,
-                      bool is_neox, int64_t rope_dim_offset, bool inverse);
-
-void fused_qk_norm_rope(torch::stable::Tensor& qkv, int64_t num_heads_q,
-                        int64_t num_heads_k, int64_t num_heads_v,
-                        int64_t head_dim, double eps,
-                        torch::stable::Tensor& q_weight,
-                        torch::stable::Tensor& k_weight,
-                        torch::stable::Tensor& cos_sin_cache, bool is_neox,
-                        torch::stable::Tensor& position_ids,
-                        int64_t forced_token_heads_per_warp);
-
-// Activation kernels (shared CUDA/ROCm)
-void silu_and_mul(torch::stable::Tensor& out, torch::stable::Tensor& input);
-void silu_and_mul_clamp(torch::stable::Tensor& out,
-                        torch::stable::Tensor& input, double limit);
-void mul_and_silu(torch::stable::Tensor& out, torch::stable::Tensor& input);
-void gelu_and_mul(torch::stable::Tensor& out, torch::stable::Tensor& input);
-void gelu_tanh_and_mul(torch::stable::Tensor& out,
-                       torch::stable::Tensor& input);
-void fatrelu_and_mul(torch::stable::Tensor& out, torch::stable::Tensor& input,
-                     double threshold);
-void swigluoai_and_mul(torch::stable::Tensor& out, torch::stable::Tensor& input,
-                       double alpha = 1.702, double limit = 7.0);
-void gelu_new(torch::stable::Tensor& out, torch::stable::Tensor& input);
-void gelu_fast(torch::stable::Tensor& out, torch::stable::Tensor& input);
-void gelu_quick(torch::stable::Tensor& out, torch::stable::Tensor& input);
-
-// INT8 quantization kernels (shared CUDA/ROCm)
-void static_scaled_int8_quant(torch::stable::Tensor& out,
-                              torch::stable::Tensor const& input,
-                              torch::stable::Tensor const& scale,
-                              std::optional<torch::stable::Tensor> const& azp);
-
-void dynamic_scaled_int8_quant(torch::stable::Tensor& out,
-                               torch::stable::Tensor const& input,
-                               torch::stable::Tensor& scales,
-                               std::optional<torch::stable::Tensor> const& azp);
-
-// FP8 quantization kernels (shared CUDA/ROCm)
-void static_scaled_fp8_quant(
-    torch::stable::Tensor& out, torch::stable::Tensor const& input,
-    torch::stable::Tensor const& scale,
-    std::optional<torch::headeronly::IntHeaderOnlyArrayRef> group_shape =
-        std::nullopt);
-
-void dynamic_scaled_fp8_quant(torch::stable::Tensor& out,
-                              torch::stable::Tensor const& input,
-                              torch::stable::Tensor& scale);
-
-void dynamic_per_token_scaled_fp8_quant(
-    torch::stable::Tensor& out, torch::stable::Tensor const& input,
-    torch::stable::Tensor& scale,
-    std::optional<torch::stable::Tensor> const& scale_ub);
-
-// GPTQ kernels (shared CUDA/ROCm)
-torch::stable::Tensor gptq_gemm(torch::stable::Tensor a,
-                                torch::stable::Tensor b_q_weight,
-                                torch::stable::Tensor b_gptq_qzeros,
-                                torch::stable::Tensor b_gptq_scales,
-                                torch::stable::Tensor b_g_idx, bool use_exllama,
-                                bool use_v2_format, int64_t bit);
-
-void gptq_shuffle(torch::stable::Tensor q_weight, torch::stable::Tensor q_perm,
-                  int64_t bit);
-
-// GGML kernels (shared CUDA/ROCm)
-torch::stable::Tensor ggml_dequantize(
-    torch::stable::Tensor W, int64_t type, int64_t m, int64_t n,
-    std::optional<torch::headeronly::ScalarType> const& dtype);
-
-torch::stable::Tensor ggml_mul_mat_vec_a8(torch::stable::Tensor W,
-                                          torch::stable::Tensor X, int64_t type,
-                                          int64_t row);
-
-torch::stable::Tensor ggml_mul_mat_a8(torch::stable::Tensor W,
-                                      torch::stable::Tensor X, int64_t type,
-                                      int64_t row);
-
-torch::stable::Tensor ggml_moe_a8(torch::stable::Tensor X,
-                                  torch::stable::Tensor W,
-                                  torch::stable::Tensor sorted_token_ids,
-                                  torch::stable::Tensor expert_ids,
-                                  torch::stable::Tensor num_tokens_post_padded,
-                                  int64_t type, int64_t row, int64_t top_k,
-                                  int64_t tokens);
-
-torch::stable::Tensor ggml_moe_a8_vec(torch::stable::Tensor X,
-                                      torch::stable::Tensor W,
-                                      torch::stable::Tensor topk_ids,
-                                      int64_t top_k, int64_t type, int64_t row,
-                                      int64_t tokens);
-
-int64_t ggml_moe_get_block_size(int64_t type);

csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu

@@ -38,8 +38,8 @@ namespace vllm {
 // Use UE4M3 by default.
 template <class Type, bool UE8M0_SF = false>
 __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
-    cvt_fp16_to_fp4(int32_t numRows, int32_t numCols, int32_t outputCols,
-                    int32_t num_padded_cols, Type const* __restrict__ in,
+    cvt_fp16_to_fp4(int32_t numRows, int32_t numCols, int32_t num_padded_cols,
+                    Type const* __restrict__ in,
                     float const* __restrict__ SFScale,
                     uint32_t* __restrict__ out, uint32_t* __restrict__ SFout) {
   using PackedVec = vllm::PackedVec<Type, CVT_FP4_PACK16>;
@@ -50,7 +50,7 @@ __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
                 "Vec size is not matched.");
 
   // Precompute SF layout parameter (constant for entire kernel).
-  int32_t const numKTiles = (outputCols + 63) / 64;
+  int32_t const numKTiles = (numCols + 63) / 64;
 
   int sf_m = round_up<int>(numRows, 128);
   int32_t const colIdx = blockDim.x * blockIdx.y + threadIdx.x;
@@ -68,17 +68,16 @@ __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
       PackedVec in_vec;
       int64_t inOffset = rowIdx * (numCols / CVT_FP4_ELTS_PER_THREAD) + colIdx;
 
-      // If we are outside valid columns, feed zeros
-      bool valid_input = (rowIdx < numRows) && (elem_idx < numCols);
-      bool valid_output = (rowIdx < numRows) && (elem_idx < outputCols);
+      // If we are outside valid rows OR outside valid columns -> Use Zeros
+      bool valid = (rowIdx < numRows) && (elem_idx < numCols);
       if constexpr (CVT_FP4_PACK16) {
         ld256_cg_or_zero(reinterpret_cast<u32x8_t&>(in_vec),
                          &reinterpret_cast<const uint32_t*>(in)[inOffset * 8],
-                         valid_input);
+                         valid);
       } else {
         ld128_cg_or_zero(reinterpret_cast<uint4&>(in_vec),
                          &reinterpret_cast<const uint32_t*>(in)[inOffset * 4],
-                         valid_input);
+                         valid);
       }
 
       auto sf_out =
@@ -90,16 +89,16 @@ __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
           cvt_warp_fp16_to_fp4<Type, CVT_FP4_NUM_THREADS_PER_SF, UE8M0_SF>(
               in_vec, global_scale, sf_out);
 
-      if (valid_output) {
+      // We do NOT write output for padding because the 'out' tensor is not
+      // padded.
+      if (valid) {
         if constexpr (CVT_FP4_PACK16) {
-          int64_t outOffset = rowIdx * (outputCols / 8) + colIdx * 2;
+          int64_t outOffset = rowIdx * (numCols / 8) + colIdx * 2;
           uint64_t packed64 =
               (uint64_t(out_val.hi) << 32) | uint64_t(out_val.lo);
           reinterpret_cast<uint64_t*>(out)[outOffset >> 1] = packed64;
         } else {
-          int64_t outOffset =
-              rowIdx * (outputCols / CVT_FP4_ELTS_PER_THREAD) + colIdx;
-          out[outOffset] = out_val;
+          out[inOffset] = out_val;
         }
       }
     }
@@ -110,8 +109,7 @@ __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
 template <class Type, bool UE8M0_SF = false>
 __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
     cvt_fp16_to_fp4_sf_major(int32_t numRows, int32_t numCols,
-                             int32_t outputCols, int32_t sf_n_unpadded,
-                             int32_t num_packed_cols,
+                             int32_t sf_n_unpadded, int32_t num_packed_cols,
                              Type const* __restrict__ in,
                              float const* __restrict__ SFScale,
                              uint32_t* __restrict__ out,
@@ -138,7 +136,7 @@ __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
       PackedVec in_vec;
       int64_t inOffset = rowIdx * (numCols / CVT_FP4_ELTS_PER_THREAD) + colIdx;
 
-      // If we are outside valid columns, feed zeros
+      // If we are outside valid rows OR outside valid columns -> Use Zeros
       bool valid = (rowIdx < numRows) && (elem_idx < numCols);
       if constexpr (CVT_FP4_PACK16) {
         ld256_cg_or_zero(reinterpret_cast<u32x8_t&>(in_vec),
@@ -157,16 +155,16 @@ __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
           cvt_warp_fp16_to_fp4<Type, CVT_FP4_NUM_THREADS_PER_SF, UE8M0_SF>(
               in_vec, global_scale, sf_out);
 
-      if (rowIdx < numRows) {
+      // We do NOT write output for padding because the 'out' tensor is not
+      // padded.
+      if (valid) {
         if constexpr (CVT_FP4_PACK16) {
-          int64_t outOffset = rowIdx * (outputCols / 8) + colIdx * 2;
+          int64_t outOffset = rowIdx * (numCols / 8) + colIdx * 2;
           uint64_t packed64 =
               (uint64_t(out_val.hi) << 32) | uint64_t(out_val.lo);
           reinterpret_cast<uint64_t*>(out)[outOffset >> 1] = packed64;
         } else {
-          int64_t outOffset =
-              rowIdx * (outputCols / CVT_FP4_ELTS_PER_THREAD) + colIdx;
-          out[outOffset] = out_val;
+          out[inOffset] = out_val;
         }
       }
     }
@@ -182,11 +180,8 @@ void scaled_fp4_quant_sm1xxa(torch::stable::Tensor const& output,
                              bool is_sf_swizzled_layout) {
   int32_t m = input.size(0);
   int32_t n = input.size(1);
-  int32_t output_n = output.size(1) * 2;
 
   STD_TORCH_CHECK(n % 16 == 0, "The N dimension must be multiple of 16.");
-  STD_TORCH_CHECK(output_n % 16 == 0,
-                  "The output tensor width must be a multiple of 16.");
   STD_TORCH_CHECK(
       input.scalar_type() == torch::headeronly::ScalarType::Half ||
           input.scalar_type() == torch::headeronly::ScalarType::BFloat16,
@@ -202,7 +197,7 @@ void scaled_fp4_quant_sm1xxa(torch::stable::Tensor const& output,
       input.get_device_index());
   auto stream = get_current_cuda_stream(input.get_device_index());
 
-  int output_sf_n_unpadded = int(output_n / CVT_FP4_SF_VEC_SIZE);
+  int sf_n_unpadded = int(n / CVT_FP4_SF_VEC_SIZE);
 
   // Grid, Block size. Each thread converts 8 values.
   dim3 block(std::min(int(n / ELTS_PER_THREAD), 512));
@@ -210,7 +205,7 @@ void scaled_fp4_quant_sm1xxa(torch::stable::Tensor const& output,
       vllm_runtime_blocks_per_sm(static_cast<int>(block.x));
 
   if (is_sf_swizzled_layout) {
-    int sf_n_int = int(vllm::round_up(output_sf_n_unpadded, 4) / 4);
+    int sf_n_int = int(vllm::round_up(sf_n_unpadded, 4) / 4);
     int32_t num_padded_cols =
         sf_n_int * 4 * CVT_FP4_SF_VEC_SIZE / CVT_FP4_ELTS_PER_THREAD;
 
@@ -225,12 +220,12 @@ void scaled_fp4_quant_sm1xxa(torch::stable::Tensor const& output,
           using cuda_type = vllm::CUDATypeConverter<scalar_t>::Type;
           auto input_ptr = static_cast<cuda_type const*>(input.data_ptr());
           vllm::cvt_fp16_to_fp4<cuda_type, false><<<grid, block, 0, stream>>>(
-              m, n, output_n, num_padded_cols, input_ptr, input_sf_ptr,
+              m, n, num_padded_cols, input_ptr, input_sf_ptr,
               reinterpret_cast<uint32_t*>(output_ptr),
               reinterpret_cast<uint32_t*>(sf_out));
         });
   } else {
-    int num_packed_cols = output_n / CVT_FP4_ELTS_PER_THREAD;
+    int num_packed_cols = n / CVT_FP4_ELTS_PER_THREAD;
     int grid_y = vllm::div_round_up(num_packed_cols, static_cast<int>(block.x));
     int grid_x = std::min(
         m, std::max(1, (multiProcessorCount * numBlocksPerSM) / grid_y));
@@ -242,8 +237,7 @@ void scaled_fp4_quant_sm1xxa(torch::stable::Tensor const& output,
           auto input_ptr = static_cast<cuda_type const*>(input.data_ptr());
           vllm::cvt_fp16_to_fp4_sf_major<cuda_type, false>
               <<<grid, block, 0, stream>>>(
-                  m, n, output_n, output_sf_n_unpadded, num_packed_cols,
-                  input_ptr, input_sf_ptr,
+                  m, n, sf_n_unpadded, num_packed_cols, input_ptr, input_sf_ptr,
                   reinterpret_cast<uint32_t*>(output_ptr),
                   reinterpret_cast<uint32_t*>(sf_out));
         });

csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_kernels.cu

@@ -22,8 +22,6 @@
 
 #include "cutlass/cutlass.h"
 
-#include <type_traits>
-
 #include "cutlass/gemm/collective/collective_builder.hpp"
 #include "cutlass/epilogue/collective/collective_builder.hpp"
 #include "cutlass/gemm/device/gemm_universal_adapter.h"
@@ -32,21 +30,15 @@
 #include "cutlass/util/packed_stride.hpp"
 
 #include "core/math.hpp"
-#include "core/batch_invariant.hpp"
 
 using namespace cute;
 
 #if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
 
-// Configuration for M in (256, inf), also reused for batch-invariant mode
-// to keep a fixed large-M tiling across all batch sizes.
-// Do not change the tile K or tile scheduler here unless you are also
-// updating the batch-invariant behavior; if batch-invariant mode needs a
-// different schedule, add a dedicated batch-invariant config/path instead.
+// Configuration for M in (256, inf)
 struct sm100_fp4_config_default {
   using KernelSchedule = cutlass::gemm::collective::KernelScheduleAuto;
   using EpilogueSchedule = cutlass::epilogue::collective::EpilogueScheduleAuto;
-  using TileScheduler = cutlass::gemm::PersistentScheduler;
   using TileShape = Shape<_256, _256, _256>;
   using ClusterShape = Shape<_2, _1, _1>;
   using PerSmTileShape_MNK = Shape<_128, _256, _256>;
@@ -56,7 +48,6 @@ struct sm100_fp4_config_default {
 struct sm100_fp4_config_M256 {
   using KernelSchedule = cutlass::gemm::collective::KernelScheduleAuto;
   using EpilogueSchedule = cutlass::epilogue::collective::EpilogueScheduleAuto;
-  using TileScheduler = void;
   using TileShape = Shape<_256, _128, _256>;
   using ClusterShape = Shape<_2, _1, _1>;
   using PerSmTileShape_MNK = Shape<_128, _128, _256>;
@@ -66,7 +57,6 @@ struct sm100_fp4_config_M256 {
 struct sm100_fp4_config_M16 {
   using KernelSchedule = cutlass::gemm::collective::KernelScheduleAuto;
   using EpilogueSchedule = cutlass::epilogue::collective::EpilogueScheduleAuto;
-  using TileScheduler = void;
   using TileShape = Shape<_128, _128, _256>;
   using ClusterShape = Shape<_1, _1, _1>;
   using PerSmTileShape_MNK = Shape<_128, _128, _256>;
@@ -108,7 +98,7 @@ struct Fp4GemmSm100 {
           cutlass::epilogue::collective::EpilogueTileAuto, ElementAccumulator,
           ElementAccumulator, ElementC, LayoutCTag, AlignmentC, ElementD,
           LayoutDTag, AlignmentD,
-          typename Config::EpilogueSchedule>::CollectiveOp;
+          cutlass::epilogue::collective::EpilogueScheduleAuto>::CollectiveOp;
 
   using CollectiveMainloop =
       typename cutlass::gemm::collective::CollectiveBuilder<
@@ -117,13 +107,10 @@ struct Fp4GemmSm100 {
           ClusterShape,
           cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(
               sizeof(typename CollectiveEpilogue::SharedStorage))>,
-          typename Config::KernelSchedule>::CollectiveOp;
+          cutlass::gemm::collective::KernelScheduleAuto>::CollectiveOp;
 
-  using TileScheduler = typename Config::TileScheduler;
-  using GemmKernel =
-      cutlass::gemm::kernel::GemmUniversal<Shape<int, int, int, int>,
-                                           CollectiveMainloop,
-                                           CollectiveEpilogue, TileScheduler>;
+  using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+      Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>;
   using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;
   using StrideA = typename Gemm::GemmKernel::StrideA;
   using LayoutA = decltype(cute::make_layout(make_shape(0, 0, 0), StrideA{}));
@@ -218,17 +205,6 @@ void cutlass_fp4_gemm_dispatch(torch::stable::Tensor& D,
                                torch::stable::Tensor const& B_sf,
                                torch::stable::Tensor const& alpha, int64_t m,
                                int64_t n, int64_t k, cudaStream_t stream) {
-  if (vllm::vllm_is_batch_invariant()) {
-    using BiGemm = Fp4GemmSm100<sm100_fp4_config_default, OutType>;
-    static_assert(
-        cute::is_same_v<typename BiGemm::TileScheduler,
-                        cutlass::gemm::PersistentScheduler>,
-        "batch_invariant requires a persistent tile scheduler; stream-K or "
-        "split-K would break numerical invariance");
-    runGemm<BiGemm>(D, A, B, A_sf, B_sf, alpha, m, n, k, stream);
-    return;
-  }
-
   uint32_t const mp2 = std::max(static_cast<uint32_t>(16), next_pow_2(m));
 
   if (mp2 <= 16) {

csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu

@@ -22,8 +22,6 @@
 
 #include "cutlass/cutlass.h"
 
-#include <type_traits>
-
 #include "cutlass/gemm/collective/collective_builder.hpp"
 #include "cutlass/epilogue/collective/collective_builder.hpp"
 #include "cutlass/gemm/device/gemm_universal_adapter.h"
@@ -32,7 +30,6 @@
 #include "cutlass/util/packed_stride.hpp"
 
 #include "core/math.hpp"
-#include "core/batch_invariant.hpp"
 
 using namespace cute;
 
@@ -52,22 +49,12 @@ constexpr auto FLOAT4_E2M1X2 = torch::headeronly::ScalarType::Byte;
 constexpr auto SF_DTYPE = torch::headeronly::ScalarType::Float8_e4m3fn;
 
 struct sm120_fp4_config_M256 {
-  using KernelSchedule = cutlass::gemm::collective::KernelScheduleAuto;
-  using EpilogueSchedule = cutlass::epilogue::collective::EpilogueScheduleAuto;
-  using TileScheduler = void;
   using ClusterShape = Shape<_1, _1, _1>;
   using MmaTileShape = Shape<_128, _128, _128>;
   using PerSmTileShape_MNK = Shape<_128, _128, _128>;
 };
 
 struct sm120_fp4_config_default {
-  // Also used for batch-invariant mode.
-  // Do not change the tile K or tile scheduler here unless you are also
-  // updating the batch-invariant behavior; if batch-invariant mode needs a
-  // different schedule, add a dedicated batch-invariant config/path instead.
-  using KernelSchedule = cutlass::gemm::collective::KernelScheduleAuto;
-  using EpilogueSchedule = cutlass::epilogue::collective::EpilogueScheduleAuto;
-  using TileScheduler = cutlass::gemm::PersistentScheduler;
   using ClusterShape = Shape<_1, _1, _1>;
   using MmaTileShape = Shape<_256, _128, _128>;
   using PerSmTileShape_MNK = Shape<_256, _128, _128>;
@@ -104,7 +91,7 @@ struct Fp4GemmSm120 {
           cutlass::epilogue::collective::EpilogueTileAuto, ElementAccumulator,
           ElementAccumulator, ElementC, LayoutCTag, AlignmentC, ElementD,
           LayoutDTag, AlignmentD,
-          typename Config::EpilogueSchedule>::CollectiveOp;
+          cutlass::epilogue::collective::EpilogueScheduleAuto>::CollectiveOp;
 
   using CollectiveMainloop =
       typename cutlass::gemm::collective::CollectiveBuilder<
@@ -113,13 +100,10 @@ struct Fp4GemmSm120 {
           ClusterShape,
           cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(
               sizeof(typename CollectiveEpilogue::SharedStorage))>,
-          typename Config::KernelSchedule>::CollectiveOp;
+          cutlass::gemm::collective::KernelScheduleAuto>::CollectiveOp;
 
-  using TileScheduler = typename Config::TileScheduler;
-  using GemmKernel =
-      cutlass::gemm::kernel::GemmUniversal<Shape<int, int, int, int>,
-                                           CollectiveMainloop,
-                                           CollectiveEpilogue, TileScheduler>;
+  using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+      Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>;
 
   using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;
 };
@@ -196,41 +180,39 @@ void runGemm(torch::stable::Tensor& D, torch::stable::Tensor const& A,
   CUTLASS_CHECK(gemm.run(arguments, workspace.data_ptr(), stream));
 }
 
-namespace {
-
-// Dispatch function to select appropriate config based on M (file-local;
-// internal linkage avoids clashing with SM100's cutlass_fp4_gemm_dispatch in
-// nvfp4_scaled_mm_kernels.cu).
-template <typename OutType>
-void cutlass_fp4_gemm_dispatch(torch::stable::Tensor& D,
-                               torch::stable::Tensor const& A,
-                               torch::stable::Tensor const& B,
-                               torch::stable::Tensor const& A_sf,
-                               torch::stable::Tensor const& B_sf,
-                               torch::stable::Tensor const& alpha, int m, int n,
-                               int k, cudaStream_t stream) {
-  if (vllm::vllm_is_batch_invariant()) {
-    using BiGemm = Fp4GemmSm120<sm120_fp4_config_default, OutType>;
-    static_assert(
-        cute::is_same_v<typename BiGemm::TileScheduler,
-                        cutlass::gemm::PersistentScheduler>,
-        "batch_invariant requires a persistent tile scheduler; stream-K or "
-        "split-K would break numerical invariance");
-    runGemm<typename BiGemm::Gemm>(D, A, B, A_sf, B_sf, alpha, m, n, k, stream);
-    return;
-  }
-
+void cutlass_fp4_bf16_gemm_dispatch(torch::stable::Tensor& D,
+                                    torch::stable::Tensor const& A,
+                                    torch::stable::Tensor const& B,
+                                    torch::stable::Tensor const& A_sf,
+                                    torch::stable::Tensor const& B_sf,
+                                    torch::stable::Tensor const& alpha, int m,
+                                    int n, int k, cudaStream_t stream) {
   uint32_t const mp2 = std::max(static_cast<uint32_t>(16), next_pow_2(m));
   if (mp2 <= 256) {
-    runGemm<typename Fp4GemmSm120<sm120_fp4_config_M256, OutType>::Gemm>(
+    runGemm<Fp4GemmSm120<sm120_fp4_config_M256, cutlass::bfloat16_t>::Gemm>(
         D, A, B, A_sf, B_sf, alpha, m, n, k, stream);
   } else {
-    runGemm<typename Fp4GemmSm120<sm120_fp4_config_default, OutType>::Gemm>(
+    runGemm<Fp4GemmSm120<sm120_fp4_config_default, cutlass::bfloat16_t>::Gemm>(
         D, A, B, A_sf, B_sf, alpha, m, n, k, stream);
   }
 }
 
-}  // namespace
+void cutlass_fp4_f16_gemm_dispatch(torch::stable::Tensor& D,
+                                   torch::stable::Tensor const& A,
+                                   torch::stable::Tensor const& B,
+                                   torch::stable::Tensor const& A_sf,
+                                   torch::stable::Tensor const& B_sf,
+                                   torch::stable::Tensor const& alpha, int m,
+                                   int n, int k, cudaStream_t stream) {
+  uint32_t const mp2 = std::max(static_cast<uint32_t>(16), next_pow_2(m));
+  if (mp2 <= 256) {
+    runGemm<Fp4GemmSm120<sm120_fp4_config_M256, cutlass::half_t>::Gemm>(
+        D, A, B, A_sf, B_sf, alpha, m, n, k, stream);
+  } else {
+    runGemm<Fp4GemmSm120<sm120_fp4_config_default, cutlass::half_t>::Gemm>(
+        D, A, B, A_sf, B_sf, alpha, m, n, k, stream);
+  }
+}
 
 void cutlass_scaled_fp4_mm_sm120a(torch::stable::Tensor& D,
                                   torch::stable::Tensor const& A,
@@ -293,11 +275,11 @@ void cutlass_scaled_fp4_mm_sm120a(torch::stable::Tensor& D,
   const cudaStream_t stream = get_current_cuda_stream(A.get_device_index());
 
   if (out_dtype == torch::headeronly::ScalarType::BFloat16) {
-    return cutlass_fp4_gemm_dispatch<cutlass::bfloat16_t>(
-        D, A, B, A_sf, B_sf, alpha, m, n, k, stream);
+    return cutlass_fp4_bf16_gemm_dispatch(D, A, B, A_sf, B_sf, alpha, m, n, k,
+                                          stream);
   } else if (out_dtype == torch::headeronly::ScalarType::Half) {
-    return cutlass_fp4_gemm_dispatch<cutlass::half_t>(D, A, B, A_sf, B_sf,
-                                                      alpha, m, n, k, stream);
+    return cutlass_fp4_f16_gemm_dispatch(D, A, B, A_sf, B_sf, alpha, m, n, k,
+                                         stream);
   } else {
     STD_TORCH_CHECK(false, "Unsupported output data type of nvfp4 mm sm120 (",
                     out_dtype, ")");

csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm.cuh

@@ -202,7 +202,7 @@ struct cutlass_3x_gemm_sm120 {
               sizeof(typename CollectiveEpilogue::SharedStorage))>,
           KernelSchedule>::CollectiveOp;
 
-  using GemmKernel = enable_sm120_family<cutlass::gemm::kernel::GemmUniversal<
+  using GemmKernel = enable_sm120_only<cutlass::gemm::kernel::GemmUniversal<
       Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>>;
 };
 

csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8_dispatch.cuh

@@ -72,7 +72,7 @@ struct cutlass_3x_gemm_sm120_custom {
               sizeof(typename CollectiveEpilogue::SharedStorage))>,
           KernelSchedule, void>::CollectiveOp;
 
-  using GemmKernel = enable_sm120_family<cutlass::gemm::kernel::GemmUniversal<
+  using GemmKernel = enable_sm120_only<cutlass::gemm::kernel::GemmUniversal<
       Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>>;
 };
 

csrc/libtorch_stable/torch_bindings.cpp

@@ -266,162 +266,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
   // Hadamard transforms
   // conditionally compiled so impl registration is in source file
   ops.def("hadacore_transform(Tensor! x, bool inplace) -> Tensor");
-
-  // Apply Root Mean Square (RMS) Normalization to the input tensor.
-  ops.def(
-      "rms_norm(Tensor! result, Tensor input, Tensor weight, float epsilon) -> "
-      "()");
-
-  // In-place fused Add and RMS Normalization.
-  ops.def(
-      "fused_add_rms_norm(Tensor! input, Tensor! residual, Tensor weight, "
-      "float epsilon) -> ()");
-
-  // Layernorm-quant
-  // Apply Root Mean Square (RMS) Normalization to the input tensor.
-  ops.def(
-      "rms_norm_static_fp8_quant(Tensor! result, Tensor input, Tensor weight, "
-      "Tensor scale, float epsilon) -> "
-      "()");
-
-  // In-place fused Add and RMS Normalization.
-  ops.def(
-      "fused_add_rms_norm_static_fp8_quant(Tensor! result, Tensor input, "
-      "Tensor! residual, Tensor weight, "
-      "Tensor scale, float epsilon) -> ()");
-
-  // Fused Layernorm + Quant kernels
-  ops.def(
-      "rms_norm_dynamic_per_token_quant(Tensor! result, Tensor input, "
-      "Tensor weight, Tensor! scale, float epsilon, "
-      "Tensor? scale_ub, Tensor!? residual) -> ()");
-
-  // Fused Layernorm + Block quant kernels
-  ops.def(
-      "rms_norm_per_block_quant(Tensor! result, Tensor input, "
-      "Tensor weight, Tensor! scale, float epsilon, "
-      "Tensor? scale_ub, Tensor!? residual, int group_size, "
-      "bool is_scale_transposed) -> ()");
-
-  // Rotary embedding
-  // Apply GPT-NeoX or GPT-J style rotary embedding to query and key.
-  ops.def(
-      "rotary_embedding(Tensor positions, Tensor! query,"
-      "                 Tensor!? key, int head_size,"
-      "                 Tensor cos_sin_cache, bool is_neox, int "
-      "rope_dim_offset=0, bool inverse=False) -> ()");
-
-  // Function for fused QK Norm and RoPE
-  ops.def(
-      "fused_qk_norm_rope(Tensor! qkv, int num_heads_q, "
-      "int num_heads_k, int num_heads_v, int head_dim, float eps, "
-      "Tensor q_weight, Tensor k_weight, Tensor cos_sin_cache, "
-      "bool is_neox, Tensor position_ids, "
-      "int forced_token_heads_per_warp=-1) -> ()");
-
-  // Activation ops
-  // Activation function used in SwiGLU.
-  ops.def("silu_and_mul(Tensor! result, Tensor input) -> ()");
-
-  ops.def("mul_and_silu(Tensor! out, Tensor input) -> ()");
-
-  // SwiGLU activation with input clamping.
-  ops.def(
-      "silu_and_mul_with_clamp(Tensor! result, Tensor input, float limit) "
-      "-> ()");
-
-  // Activation function used in GeGLU with `none` approximation.
-  ops.def("gelu_and_mul(Tensor! out, Tensor input) -> ()");
-
-  // Activation function used in GeGLU with `tanh` approximation.
-  ops.def("gelu_tanh_and_mul(Tensor! out, Tensor input) -> ()");
-
-  // FATReLU implementation.
-  ops.def("fatrelu_and_mul(Tensor! out, Tensor input, float threshold) -> ()");
-
-  ops.def(
-      "swigluoai_and_mul(Tensor! out, Tensor input, float alpha=1.702, float "
-      "limit=7.0) "
-      "-> ()");
-
-  // GELU implementation used in GPT-2.
-  ops.def("gelu_new(Tensor! out, Tensor input) -> ()");
-
-  // Approximate GELU implementation.
-  ops.def("gelu_fast(Tensor! out, Tensor input) -> ()");
-
-  // Quick GELU implementation.
-  ops.def("gelu_quick(Tensor! out, Tensor input) -> ()");
-
-  // Compute int8 quantized tensor for given scaling factor.
-  ops.def(
-      "static_scaled_int8_quant(Tensor! result, Tensor input, Tensor scale,"
-      "Tensor? azp) -> ()");
-
-  // Compute int8 quantized tensor and scaling factor
-  ops.def(
-      "dynamic_scaled_int8_quant(Tensor! result, Tensor input, Tensor! scale, "
-      "Tensor!? azp) -> ()");
-
-  // Compute FP8 quantized tensor for given scaling factor.
-  // Supports per-tensor, per-channel, per-token, and arbitrary 2D group
-  // scaling. Optional group_m/group_n specify the group shape explicitly;
-  // required for 1D scales to disambiguate per-channel vs per-token.
-  ops.def(
-      "static_scaled_fp8_quant(Tensor! result, Tensor input, Tensor scale, "
-      "int[]? group_shape=None) -> ()");
-
-  // Compute dynamic-per-tensor FP8 quantized tensor and scaling factor.
-  ops.def(
-      "dynamic_scaled_fp8_quant(Tensor! result, Tensor input, Tensor! scale) "
-      "-> "
-      "()");
-
-  // Compute dynamic-per-token FP8 quantized tensor and scaling factor.
-  ops.def(
-      "dynamic_per_token_scaled_fp8_quant(Tensor! result, Tensor input, "
-      "Tensor! scale, Tensor? scale_ub) -> "
-      "()");
-
-  // Quantized GEMM for GPTQ.
-  // Note: even though the C++ inferred schema is correct for this op, it seems
-  // to prevent the meta function registry.
-  ops.def(
-      "gptq_gemm(Tensor a, Tensor b_q_weight, Tensor b_gptq_qzeros, "
-      "Tensor b_gptq_scales, Tensor b_g_idx, bool use_exllama, bool "
-      "use_v2_format, int bit) "
-      "-> Tensor");
-
-  // Post processing for GPTQ.
-  ops.def("gptq_shuffle(Tensor! q_weight, Tensor q_perm, int bit) -> ()");
-
-  // Dequantization for GGML.
-  ops.def(
-      "ggml_dequantize(Tensor W, int type, SymInt m, SymInt n, ScalarType? "
-      "dtype) -> Tensor");
-
-  // mmvq kernel for GGML.
-  ops.def(
-      "ggml_mul_mat_vec_a8(Tensor W, Tensor X, int type, SymInt row) "
-      "-> Tensor");
-
-  // mmq kernel for GGML.
-  ops.def(
-      "ggml_mul_mat_a8(Tensor W, Tensor X, int type, SymInt row) -> Tensor");
-
-  // moe kernel for GGML.
-  ops.def(
-      "ggml_moe_a8(Tensor X, Tensor W, "
-      "Tensor sorted_token_ids, Tensor expert_ids, Tensor "
-      "num_tokens_post_padded, "
-      "int type, SymInt row, SymInt top_k, SymInt tokens) -> Tensor");
-
-  ops.def(
-      "ggml_moe_a8_vec(Tensor X, Tensor W, "
-      "Tensor topk_ids, int top_k, "
-      "int type, SymInt row, SymInt tokens) -> Tensor");
-
-  ops.def("ggml_moe_get_block_size(int type) -> int");
 }
 
 STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
@@ -468,57 +312,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
   // AllSpark ops: conditionally compiled so impl registrations are in source
   // files (allspark_repack.cu and allspark_qgemm_w8a16.cu)
 #endif
-
-  // Layernorm kernels (shared CUDA/ROCm)
-  ops.impl("rms_norm", TORCH_BOX(&rms_norm));
-  ops.impl("fused_add_rms_norm", TORCH_BOX(&fused_add_rms_norm));
-
-  // Layernorm-quant kernels (shared CUDA/ROCm)
-  ops.impl("rms_norm_static_fp8_quant", TORCH_BOX(&rms_norm_static_fp8_quant));
-  ops.impl("fused_add_rms_norm_static_fp8_quant",
-           TORCH_BOX(&fused_add_rms_norm_static_fp8_quant));
-
-  // Fused layernorm + dynamic per-token quant kernels (shared CUDA/ROCm)
-  ops.impl("rms_norm_dynamic_per_token_quant",
-           TORCH_BOX(&rms_norm_dynamic_per_token_quant));
-  ops.impl("rms_norm_per_block_quant", TORCH_BOX(&rms_norm_per_block_quant));
-
-  // Positional encoding kernels (shared CUDA/ROCm)
-  ops.impl("rotary_embedding", TORCH_BOX(&rotary_embedding));
-  ops.impl("fused_qk_norm_rope", TORCH_BOX(&fused_qk_norm_rope));
-
-  // Activation kernels (shared CUDA/ROCm)
-  ops.impl("silu_and_mul", TORCH_BOX(&silu_and_mul));
-  ops.impl("mul_and_silu", TORCH_BOX(&mul_and_silu));
-  ops.impl("gelu_and_mul", TORCH_BOX(&gelu_and_mul));
-  ops.impl("gelu_tanh_and_mul", TORCH_BOX(&gelu_tanh_and_mul));
-  ops.impl("fatrelu_and_mul", TORCH_BOX(&fatrelu_and_mul));
-  ops.impl("swigluoai_and_mul", TORCH_BOX(&swigluoai_and_mul));
-  ops.impl("gelu_new", TORCH_BOX(&gelu_new));
-  ops.impl("gelu_fast", TORCH_BOX(&gelu_fast));
-  ops.impl("gelu_quick", TORCH_BOX(&gelu_quick));
-  ops.impl("silu_and_mul_with_clamp", TORCH_BOX(&silu_and_mul_clamp));
-
-  // INT8 quantization kernels
-  ops.impl("static_scaled_int8_quant", TORCH_BOX(&static_scaled_int8_quant));
-  ops.impl("dynamic_scaled_int8_quant", TORCH_BOX(&dynamic_scaled_int8_quant));
-
-  // FP8 quantization kernels
-  ops.impl("static_scaled_fp8_quant", TORCH_BOX(&static_scaled_fp8_quant));
-  ops.impl("dynamic_scaled_fp8_quant", TORCH_BOX(&dynamic_scaled_fp8_quant));
-  ops.impl("dynamic_per_token_scaled_fp8_quant",
-           TORCH_BOX(&dynamic_per_token_scaled_fp8_quant));
-
-  // GPTQ kernels
-  ops.impl("gptq_gemm", TORCH_BOX(&gptq_gemm));
-  ops.impl("gptq_shuffle", TORCH_BOX(&gptq_shuffle));
-
-  // GGML kernels
-  ops.impl("ggml_dequantize", TORCH_BOX(&ggml_dequantize));
-  ops.impl("ggml_mul_mat_vec_a8", TORCH_BOX(&ggml_mul_mat_vec_a8));
-  ops.impl("ggml_mul_mat_a8", TORCH_BOX(&ggml_mul_mat_a8));
-  ops.impl("ggml_moe_a8", TORCH_BOX(&ggml_moe_a8));
-  ops.impl("ggml_moe_a8_vec", TORCH_BOX(&ggml_moe_a8_vec));
 }
 
 // These capability-check functions take only primitive args (no tensors), so
@@ -536,9 +329,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, ops) {
   ops.impl("cutlass_scaled_mm_supports_fp4",
            TORCH_BOX(&cutlass_scaled_mm_supports_fp4));
 #endif
-
-  // GGML block size lookup (no tensor args)
-  ops.impl("ggml_moe_get_block_size", TORCH_BOX(&ggml_moe_get_block_size));
 }
 
 REGISTER_EXTENSION(_C_stable_libtorch)

csrc/libtorch_stable/torch_utils.h

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

csrc/moe/dsv3_router_gemm_bf16_out.cu

@@ -182,7 +182,7 @@ void invokeRouterGemmBf16Output(__nv_bfloat16* output, T const* mat_a,
   config.stream = stream;
   cudaLaunchAttribute attrs[1];
   attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-  attrs[0].val.programmaticStreamSerializationAllowed = 1;
+  attrs[0].val.programmaticStreamSerializationAllowed = getEnvEnablePDL();
   config.numAttrs = 1;
   config.attrs = attrs;
   cudaLaunchKernelEx(

csrc/moe/dsv3_router_gemm_float_out.cu

@@ -182,7 +182,7 @@ void invokeRouterGemmFloatOutput(float* output, T const* mat_a, T const* mat_b,
   config.stream = stream;
   cudaLaunchAttribute attrs[1];
   attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-  attrs[0].val.programmaticStreamSerializationAllowed = 1;
+  attrs[0].val.programmaticStreamSerializationAllowed = getEnvEnablePDL();
   config.numAttrs = 1;
   config.attrs = attrs;
   cudaLaunchKernelEx(

csrc/moe/dsv3_router_gemm_utils.h

@@ -29,3 +29,15 @@ inline int getSMVersion() {
   auto* props = at::cuda::getCurrentDeviceProperties();
   return props->major * 10 + props->minor;
 }
+
+inline bool getEnvEnablePDL() {
+  static std::once_flag flag;
+  static bool enablePDL = false;
+  std::call_once(flag, [&]() {
+    if (getSMVersion() >= 90) {
+      const char* env = std::getenv("TRTLLM_ENABLE_PDL");
+      enablePDL = env && env[0] == '1' && env[1] == '\0';
+    }
+  });
+  return enablePDL;
+}

csrc/moe/dsv4_norm_router_gemm.h

@@ -0,0 +1,30 @@
+/*
+ * Fused RMSNorm + router GEMV for DeepSeek V4.
+ *
+ * Computes in a single kernel:
+ *   normed_x[m,k]   = x[m,k] * rsqrt(mean(x[m]^2) + eps) * norm_weight[k]
+ *   router_logits[m,n] = sum_k(normed_x[m,k] * gate_weight[n,k])
+ *
+ * The GEMV body mirrors the algorithm in csrc/moe/dsv3_router_gemm_*.cu
+ * (warp butterfly + smem cross-warp reduction, fp32 accumulation, PDL on
+ * SM90+).  Blocks 0..kNumTokens-1 each materialize one token's normed_x
+ * row to global memory using the algebraic identity
+ *      logits[m,n] = rsqrt[m] * sum_k(x[m,k] * nw[k] * gw[n,k])
+ * which lets every block produce its column of logits before normed_x
+ * exists in gmem.
+ *
+ * Logits output is fp32 only — DeepSeek V4 router gate is hard-coded to
+ * fp32 (vllm/model_executor/models/deepseek_v4.py:749).
+ */
+
+#pragma once
+
+#include <cuda_bf16.h>
+#include <cuda_runtime.h>
+
+#include "dsv3_router_gemm_utils.h"
+
+template <typename T, int kNumTokens, int kNumExperts, int kHiddenDim>
+void invokeNormRouterGemm(float* logits, __nv_bfloat16* normed_x, T const* x,
+                          T const* norm_weight, T const* gate_weight, float eps,
+                          cudaStream_t stream);

csrc/moe/dsv4_norm_router_gemm_entry.cu

@@ -0,0 +1,130 @@
+/*
+ * TORCH op entry for the fused RMSNorm + router GEMV kernel
+ * (DeepSeek V4 Pro).  This op is DSV4-Pro-specific: the kernel is
+ * instantiated only for ``num_experts == 384`` and ``hidden_dim ==
+ * 7168``.  Other configurations (e.g. DSV4-Flash with H=4096) must
+ * fall back to the unfused ``rms_norm`` + ``dsv3_router_gemm`` path.
+ */
+
+#include <ATen/ATen.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <torch/all.h>
+
+#include <cuda_bf16.h>
+#include <cuda_runtime.h>
+
+#include "core/registration.h"
+#include "dsv4_norm_router_gemm.h"
+
+namespace {
+
+// DSV4-Pro hard-coded shape constants.  Renamed from the earlier
+// ``kKimiK2NumExperts`` to avoid the misleading impression that this
+// kernel targets Kimi K2 — 384 happens to match Kimi K2's gate but the
+// intent here is DSV4-Pro.
+constexpr int kDsv4NumExperts = 384;
+constexpr int kDsv4HiddenDim = 7168;
+
+template <int kBegin, int kEnd>
+struct LoopUnroller {
+  static void unroll(int num_tokens, float* logits, __nv_bfloat16* normed_x,
+                     __nv_bfloat16 const* x, __nv_bfloat16 const* norm_weight,
+                     __nv_bfloat16 const* gate_weight, float eps,
+                     cudaStream_t stream) {
+    if (num_tokens == kBegin) {
+      invokeNormRouterGemm<__nv_bfloat16, kBegin, kDsv4NumExperts,
+                           kDsv4HiddenDim>(logits, normed_x, x, norm_weight,
+                                           gate_weight, eps, stream);
+    } else {
+      LoopUnroller<kBegin + 1, kEnd>::unroll(num_tokens, logits, normed_x, x,
+                                             norm_weight, gate_weight, eps,
+                                             stream);
+    }
+  }
+};
+
+template <int kEnd>
+struct LoopUnroller<kEnd, kEnd> {
+  static void unroll(int num_tokens, float* logits, __nv_bfloat16* normed_x,
+                     __nv_bfloat16 const* x, __nv_bfloat16 const* norm_weight,
+                     __nv_bfloat16 const* gate_weight, float eps,
+                     cudaStream_t stream) {
+    if (num_tokens == kEnd) {
+      invokeNormRouterGemm<__nv_bfloat16, kEnd, kDsv4NumExperts,
+                           kDsv4HiddenDim>(logits, normed_x, x, norm_weight,
+                                           gate_weight, eps, stream);
+    } else {
+      throw std::invalid_argument(
+          "Invalid num_tokens, only supports 1 to 16 for "
+          "dsv4_norm_router_gemm");
+    }
+  }
+};
+
+}  // namespace
+
+void dsv4_norm_router_gemm(at::Tensor& logits,    // [num_tokens, E] fp32
+                           at::Tensor& normed_x,  // [num_tokens, H] bf16
+                           at::Tensor const& x,   // [num_tokens, H] bf16
+                           at::Tensor const& norm_weight,  // [H] bf16
+                           at::Tensor const& gate_weight,  // [E, H] bf16
+                           double eps) {
+  TORCH_CHECK(x.dim() == 2 && norm_weight.dim() == 1 && gate_weight.dim() == 2,
+              "x must be 2D, norm_weight 1D, gate_weight 2D");
+  TORCH_CHECK(logits.dim() == 2 && normed_x.dim() == 2,
+              "logits and normed_x must be 2D");
+
+  int const num_tokens = x.size(0);
+  int const hidden_dim = x.size(1);
+  int const num_experts = gate_weight.size(0);
+
+  TORCH_CHECK(hidden_dim == kDsv4HiddenDim,
+              "Expected hidden_dim=", kDsv4HiddenDim,
+              " (DSV4-Pro), but got hidden_dim=", hidden_dim);
+  TORCH_CHECK(gate_weight.size(1) == hidden_dim,
+              "gate_weight.shape[1] must equal x.shape[1]");
+  TORCH_CHECK(norm_weight.size(0) == hidden_dim,
+              "norm_weight.shape[0] must equal x.shape[1]");
+  TORCH_CHECK(num_experts == kDsv4NumExperts,
+              "Expected num_experts=", kDsv4NumExperts,
+              " (DSV4-Pro), but got num_experts=", num_experts);
+  TORCH_CHECK(num_tokens >= 1 && num_tokens <= 16,
+              "num_tokens must be in [1, 16] for dsv4_norm_router_gemm");
+
+  TORCH_CHECK(x.dtype() == at::kBFloat16, "x must be bf16");
+  TORCH_CHECK(norm_weight.dtype() == at::kBFloat16, "norm_weight must be bf16");
+  TORCH_CHECK(gate_weight.dtype() == at::kBFloat16, "gate_weight must be bf16");
+  TORCH_CHECK(normed_x.dtype() == at::kBFloat16, "normed_x must be bf16");
+  TORCH_CHECK(logits.dtype() == at::kFloat,
+              "logits must be float32 (DSV4 router output is hard-coded fp32)");
+
+  TORCH_CHECK(normed_x.size(0) == num_tokens && normed_x.size(1) == hidden_dim,
+              "normed_x must be [num_tokens, hidden_dim]");
+  TORCH_CHECK(logits.size(0) == num_tokens && logits.size(1) == num_experts,
+              "logits must be [num_tokens, num_experts]");
+
+  TORCH_CHECK(x.is_contiguous() && norm_weight.is_contiguous() &&
+                  gate_weight.is_contiguous() && normed_x.is_contiguous() &&
+                  logits.is_contiguous(),
+              "all tensors must be contiguous");
+
+  auto const sm = getSMVersion();
+  TORCH_CHECK(sm >= 90 && sm <= 103,
+              "dsv4_norm_router_gemm requires SM_90 <= CUDA ARCH <= SM_103");
+
+  cudaStream_t const stream = at::cuda::getCurrentCUDAStream();
+
+  auto* logits_ptr = reinterpret_cast<float*>(logits.mutable_data_ptr());
+  auto* nx_ptr = reinterpret_cast<__nv_bfloat16*>(normed_x.mutable_data_ptr());
+  auto* x_ptr = reinterpret_cast<__nv_bfloat16 const*>(x.data_ptr());
+  auto* nw_ptr = reinterpret_cast<__nv_bfloat16 const*>(norm_weight.data_ptr());
+  auto* gw_ptr = reinterpret_cast<__nv_bfloat16 const*>(gate_weight.data_ptr());
+  float const eps_f = static_cast<float>(eps);
+
+  LoopUnroller<1, 16>::unroll(num_tokens, logits_ptr, nx_ptr, x_ptr, nw_ptr,
+                              gw_ptr, eps_f, stream);
+}
+
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+  m.impl("dsv4_norm_router_gemm", &dsv4_norm_router_gemm);
+}

csrc/moe/dsv4_norm_router_gemm_kernel.cu

@@ -0,0 +1,249 @@
+/*
+ * Fused RMSNorm + router GEMV for DeepSeek V4 (logits are fp32; bf16
+ * output is unsupported because DSV4 hard-codes fp32 logits).  See
+ * dsv4_norm_router_gemm.h for the math.
+ *
+ * The GEMV body mirrors csrc/moe/dsv3_router_gemm_float_out.cu (warp
+ * butterfly reduction + smem cross-warp reduction, fp32 accumulation,
+ * 128-thread block, PDL on SM90+).  RMSNorm is folded into the same
+ * pass via the identity
+ *   logits[m,n] = rsqrt[m] * sum_k(x[m,k] * nw[k] * gw[n,k])
+ * so x is read exactly once per block during the GEMV phase.  Blocks
+ * 0..kNumTokens-1 each materialize one row of normed_x for downstream
+ * experts / shared_experts to consume.
+ */
+
+#include <ATen/ATen.h>
+#include <ATen/cuda/CUDAContext.h>
+
+#include <cuda_bf16.h>
+#include <cuda_runtime.h>
+
+#include "dsv4_norm_router_gemm.h"
+
+namespace {
+
+// Convert 8 bf16 values packed in uint4 into 8 floats. Mirrors the helper
+// in dsv3_router_gemm_float_out.cu (kept local so the dsv3 file stays
+// untouched).
+template <int VPT>
+__device__ __forceinline__ void bf16_uint4_to_float8(uint4 const& vec,
+                                                     float* dst) {
+  __nv_bfloat16* bf16_ptr =
+      reinterpret_cast<__nv_bfloat16*>(const_cast<uint4*>(&vec));
+#pragma unroll
+  for (int i = 0; i < VPT; i++) {
+    dst[i] = __bfloat162float(bf16_ptr[i]);
+  }
+}
+
+template <typename T, int kBlockSize, int VPT, int kNumTokens, int kNumExperts,
+          int kHiddenDim>
+__global__ __launch_bounds__(128, 1) void norm_router_gemm_kernel(
+    float* __restrict__ logits, __nv_bfloat16* __restrict__ normed_x,
+    T const* __restrict__ x, T const* __restrict__ norm_weight,
+    T const* __restrict__ gate_weight, float eps) {
+  static_assert(kBlockSize == 128, "kernel assumes blockDim.x == 128");
+  static_assert(kHiddenDim % (VPT * kBlockSize) == 0,
+                "kHiddenDim must be a multiple of VPT * kBlockSize");
+
+  int const n_idx = blockIdx.x;
+  int const tid = threadIdx.x;
+  constexpr int kWarpSize = 32;
+  constexpr int kNumWarps = kBlockSize / kWarpSize;
+  constexpr int k_elems_per_iter = VPT * kBlockSize;
+  constexpr int k_iterations = kHiddenDim / k_elems_per_iter;
+
+  T const* gw_col = gate_weight + n_idx * kHiddenDim;
+
+  // Per-thread accumulators — fp32 throughout, matching dsv3 / layernorm.
+  float partial[kNumTokens] = {};
+  float ss[kNumTokens] = {};
+
+  // Cross-warp reduction scratch.
+  __shared__ float sm_partial[kNumTokens][kNumWarps];
+  __shared__ float sm_ss[kNumTokens][kNumWarps];
+  __shared__ float s_rsqrt[kNumTokens];
+
+  int k_bases[k_iterations];
+#pragma unroll
+  for (int ki = 0; ki < k_iterations; ki++) {
+    k_bases[ki] = ki * k_elems_per_iter + tid * VPT;
+  }
+
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+  asm volatile("griddepcontrol.wait;");
+#endif
+
+  // ---- Phase 1: single pass over x, accumulate partial GEMV and ss. ----
+#pragma unroll
+  for (int ki = 0; ki < k_iterations; ki++) {
+    int const k_base = k_bases[ki];
+
+    uint4 nw_vec = *reinterpret_cast<uint4 const*>(norm_weight + k_base);
+    float nw_f[VPT];
+    bf16_uint4_to_float8<VPT>(nw_vec, nw_f);
+
+    uint4 b_vec = *reinterpret_cast<uint4 const*>(gw_col + k_base);
+    float b_f[VPT];
+    bf16_uint4_to_float8<VPT>(b_vec, b_f);
+
+#pragma unroll
+    for (int m = 0; m < kNumTokens; m++) {
+      uint4 a_vec =
+          *reinterpret_cast<uint4 const*>(x + m * kHiddenDim + k_base);
+      float a_f[VPT];
+      bf16_uint4_to_float8<VPT>(a_vec, a_f);
+
+#pragma unroll
+      for (int k = 0; k < VPT; k++) {
+        float a = a_f[k];
+        ss[m] += a * a;
+        partial[m] += a * nw_f[k] * b_f[k];
+      }
+    }
+  }
+
+  // ---- Phase 2: warp butterfly reduction for both ss[] and partial[]. ----
+  int const warpId = tid / kWarpSize;
+  int const laneId = tid % kWarpSize;
+
+#pragma unroll
+  for (int m = 0; m < kNumTokens; m++) {
+    float p = partial[m];
+    float s = ss[m];
+
+    p += __shfl_xor_sync(0xffffffff, p, 16);
+    s += __shfl_xor_sync(0xffffffff, s, 16);
+    p += __shfl_xor_sync(0xffffffff, p, 8);
+    s += __shfl_xor_sync(0xffffffff, s, 8);
+    p += __shfl_xor_sync(0xffffffff, p, 4);
+    s += __shfl_xor_sync(0xffffffff, s, 4);
+    p += __shfl_xor_sync(0xffffffff, p, 2);
+    s += __shfl_xor_sync(0xffffffff, s, 2);
+    p += __shfl_xor_sync(0xffffffff, p, 1);
+    s += __shfl_xor_sync(0xffffffff, s, 1);
+
+    if (laneId == 0) {
+      sm_partial[m][warpId] = p;
+      sm_ss[m][warpId] = s;
+    }
+  }
+
+  __syncthreads();
+
+  // ---- Phase 3: tid 0 finalises the reduction, writes logits, stashes
+  //               rsqrt[m] in smem for phase 4. ----
+  if (tid == 0) {
+#pragma unroll
+    for (int m = 0; m < kNumTokens; m++) {
+      float p_sum = 0.0f;
+      float s_sum = 0.0f;
+#pragma unroll
+      for (int w = 0; w < kNumWarps; w++) {
+        p_sum += sm_partial[m][w];
+        s_sum += sm_ss[m][w];
+      }
+      // Order matches layernorm_kernels.cu: rsqrtf(variance / H + eps).
+      // Use division (not multiply-by-reciprocal) to avoid an extra ULP
+      // mismatch with the reference RMSNorm.
+      float rs = rsqrtf(s_sum / static_cast<float>(kHiddenDim) + eps);
+      s_rsqrt[m] = rs;
+      logits[m * kNumExperts + n_idx] = p_sum * rs;
+    }
+  }
+
+  __syncthreads();
+
+  // ---- Phase 4: spread normed_x writes across blocks 0..kNumTokens-1.
+  //              Each writer block handles exactly one token row,
+  //              avoiding the long tail of block 0 doing all M rows.
+  //              Every block has every token's rsqrt[] in s_rsqrt
+  //              already (computed independently in phase 3), so no
+  //              cross-block synchronization is required. ----
+  if (n_idx < kNumTokens) {
+    int const m_writer = n_idx;
+    float const rs = s_rsqrt[m_writer];
+    __nv_bfloat16 const* x_row = x + m_writer * kHiddenDim;
+    __nv_bfloat16* normed_row = normed_x + m_writer * kHiddenDim;
+
+#pragma unroll
+    for (int ki = 0; ki < k_iterations; ki++) {
+      int const k_base = k_bases[ki];
+
+      uint4 nw_vec = *reinterpret_cast<uint4 const*>(norm_weight + k_base);
+      float nw_f[VPT];
+      bf16_uint4_to_float8<VPT>(nw_vec, nw_f);
+
+      uint4 a_vec = *reinterpret_cast<uint4 const*>(x_row + k_base);
+      float a_f[VPT];
+      bf16_uint4_to_float8<VPT>(a_vec, a_f);
+
+      uint4 normed_vec;
+      __nv_bfloat16* np = reinterpret_cast<__nv_bfloat16*>(&normed_vec);
+#pragma unroll
+      for (int k = 0; k < VPT; k++) {
+        np[k] = __float2bfloat16(a_f[k] * rs * nw_f[k]);
+      }
+      *reinterpret_cast<uint4*>(normed_row + k_base) = normed_vec;
+    }
+  }
+
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+  asm volatile("griddepcontrol.launch_dependents;");
+#endif
+}
+
+}  // namespace
+
+template <typename T, int kNumTokens, int kNumExperts, int kHiddenDim>
+void invokeNormRouterGemm(float* logits, __nv_bfloat16* normed_x, T const* x,
+                          T const* norm_weight, T const* gate_weight, float eps,
+                          cudaStream_t stream) {
+  constexpr int VPT = 16 / sizeof(T);
+  constexpr int kBlockSize = 128;
+
+  cudaLaunchConfig_t config;
+  config.gridDim = kNumExperts;
+  config.blockDim = kBlockSize;
+  config.dynamicSmemBytes = 0;
+  config.stream = stream;
+
+  cudaLaunchAttribute attrs[1];
+  attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
+  attrs[0].val.programmaticStreamSerializationAllowed = 1;
+  config.numAttrs = 1;
+  config.attrs = attrs;
+
+  cudaLaunchKernelEx(&config,
+                     norm_router_gemm_kernel<T, kBlockSize, VPT, kNumTokens,
+                                             kNumExperts, kHiddenDim>,
+                     logits, normed_x, x, norm_weight, gate_weight, eps);
+}
+
+// Template instantiations — DSV4-Pro is the only supported configuration:
+// num_experts=384, hidden_dim=7168.  Other shapes (e.g. DSV4-Flash with
+// hidden_dim=4096) fall back to the unfused path on the Python side.
+#define INSTANTIATE(M)                                                    \
+  template void invokeNormRouterGemm<__nv_bfloat16, M, 384, 7168>(        \
+      float*, __nv_bfloat16*, __nv_bfloat16 const*, __nv_bfloat16 const*, \
+      __nv_bfloat16 const*, float, cudaStream_t);
+
+INSTANTIATE(1)
+INSTANTIATE(2)
+INSTANTIATE(3)
+INSTANTIATE(4)
+INSTANTIATE(5)
+INSTANTIATE(6)
+INSTANTIATE(7)
+INSTANTIATE(8)
+INSTANTIATE(9)
+INSTANTIATE(10)
+INSTANTIATE(11)
+INSTANTIATE(12)
+INSTANTIATE(13)
+INSTANTIATE(14)
+INSTANTIATE(15)
+INSTANTIATE(16)
+
+#undef INSTANTIATE

csrc/moe/marlin_moe_wna16/generate_kernels.py

@@ -15,11 +15,11 @@ SUPPORT_SM80 = False
 for arch in sys.argv[1].split(","):
     arch = arch[: arch.index(".") + 2].replace(".", "")
     arch = int(arch)
-    # SM89 and the SM12x family (SM120 RTX 5090, SM121 DGX Spark GB10)
-    # fully support mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32.
+    # only SM89 and SM120 fully support
+    # mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32.
     # SM90 and SM100 can use this PTX, but it’s simulated
     # with FP16 MMA, so it cannot achieve any acceleration.
-    if arch == 89 or arch // 10 == 12:
+    if arch in [89, 120]:
         SUPPORT_FP8 = True
     if arch >= 80:
         SUPPORT_SM80 = True

csrc/moe/marlin_moe_wna16/ops.cu

@@ -448,8 +448,8 @@ void marlin_mm(const void* A, const void* B, void* C, void* C_tmp, void* b_bias,
                 "FP8 only support Ada Lovelace or newer GPUs.");
     TORCH_CHECK(
         major_capability * 10 + minor_capability == 89 ||
-            major_capability == 12,
-        "Marlin W4A8-FP8 only support SM89 or SM12x device (It is slower than "
+            major_capability * 10 + minor_capability == 120,
+        "Marlin W4A8-FP8 only support SM89 or SM120 device (It is slower than "
         "Marlin W4A16 on other devices).");
   }