remove torch 2.9, 2.10 workround

Signed-off-by: yewentao256 <zhyanwentao@126.com>

.buildkite/ci_config.yaml

@@ -8,7 +8,6 @@ run_all_patterns:
   - "CMakeLists.txt"
   - "requirements/common.txt"
   - "requirements/cuda.txt"
-  - "requirements/kv_connectors.txt"
   - "requirements/build/cuda.txt"
   - "requirements/test/cuda.txt"
   - "setup.py"

.buildkite/hardware_tests/cpu.yaml

@@ -12,19 +12,15 @@ 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/test_onednn.py
   - tests/kernels/test_awq_int4_to_int8.py
-  - tests/kernels/quantization/test_cpu_fp8_scaled_mm.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/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/test_onednn.py
-      pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
-      pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py"
+      pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py"
 
 - label: CPU-Compatibility Tests
   depends_on: []
@@ -61,16 +57,15 @@ 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
-  - vllm/model_executor/layers/fused_moe/experts/cpu_moe.py
   - tests/quantization/test_compressed_tensors.py
   - 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/intel_jobs/lora_intel.yaml

@@ -18,18 +18,17 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_layers.py &&
       pytest -v -s lora/test_lora_checkpoints.py &&
-      pytest -v -s lora/test_lora_functions.py &&
+      (pytest -v -s lora/test_lora_functions.py --deselect="tests/lora/test_lora_functions.py::test_lora_functions_sync" --deselect="tests/lora/test_lora_functions.py::test_lora_functions_async" || true) &&
       pytest -v -s lora/test_lora_huggingface.py &&
       pytest -v -s lora/test_lora_manager.py &&
       pytest -v -s lora/test_lora_utils.py &&
       pytest -v -s lora/test_peft_helper.py &&
       pytest -v -s lora/test_resolver.py &&
       pytest -v -s lora/test_utils.py &&
-      pytest -v -s lora/test_add_lora.py  &&
-      pytest -v -s lora/test_worker.py'
+      (pytest -v -s lora/test_add_lora.py --deselect="tests/lora/test_add_lora.py::test_add_lora" || true) &&
+      (pytest -v -s lora/test_worker.py --deselect="tests/lora/test_worker.py::test_worker_apply_lora" || true)'
 
 - label: LoRA Fused/MoE Kernels
   timeout_in_minutes: 45
@@ -47,9 +46,8 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       '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
@@ -67,9 +65,8 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       set -o pipefail &&
-      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-3-43264-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype1-1-2049-64-128-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]"'
+      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-2-2049-64-32-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype1-2-64000-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype1-1-102656-32-4-4]"'
 
 - label: LoRA Punica FP8/XPU Ops
   timeout_in_minutes: 45
@@ -87,7 +84,6 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_punica_ops_fp8.py &&
       pytest -v -s lora/test_punica_xpu_ops.py'
 
@@ -107,12 +103,10 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       (pytest -v -s lora/test_mixtral.py --deselect="tests/lora/test_mixtral.py::test_mixtral_lora[4]" || true) &&
       pytest -v -s lora/test_quant_model.py --deselect="tests/lora/test_quant_model.py::test_quant_model_lora[model0]" --deselect="tests/lora/test_quant_model.py::test_quant_model_lora[model1]" --deselect="tests/lora/test_quant_model.py::test_quant_model_tp_equality[model0]" &&
-      pytest -v -s lora/test_transformers_model.py &&
-      pytest -v -s lora/test_chatglm3_tp.py &&
-      pytest -s -v lora/test_minicpmv_tp.py'
+      pytest -v -s lora/test_qwen35_densemodel_lora.py &&
+      pytest -v -s lora/test_transformers_model.py'
 
 - label: LoRA Multimodal
   timeout_in_minutes: 45
@@ -130,6 +124,6 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_default_mm_loras.py && 
+      (pytest -v -s lora/test_qwen3_unembed.py || true) &&
       pytest -v -s lora/test_whisper.py'

.buildkite/intel_jobs/misc_intel.yaml

@@ -49,7 +49,7 @@ steps:
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       cd tests &&
-      pytest -v -s v1/logits_processors --ignore=v1/logits_processors/test_custom_online.py --ignore=v1/logits_processors/test_custom_offline.py &&
+      pytest -v -s v1/logits_processors &&
       pytest -v -s v1/test_oracle.py &&
       pytest -v -s v1/test_request.py &&
       pytest -v -s v1/test_outputs.py'

.buildkite/intel_jobs/test-intel.yaml

@@ -36,7 +36,6 @@ steps:
         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 facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8 &&
-        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'
@@ -62,5 +61,5 @@ steps:
         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/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'
+        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'

.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/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,11 +24,10 @@ 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"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -41,11 +37,10 @@ steps:
         agents:
           queue: arm64_cpu_queue_release
         commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04  --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
           - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -59,7 +54,6 @@ steps:
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
           - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -69,11 +63,10 @@ steps:
         agents:
           queue: cpu_queue_release
         commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86_CU129}\" --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"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -83,11 +76,10 @@ steps:
         agents:
           queue: cpu_queue_release
         commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinux2_28-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
           - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -101,7 +93,6 @@ steps:
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
           - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -139,8 +130,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 \
@@ -149,7 +138,6 @@ steps:
           # re-tag to default image tag and push, just in case arm64 build fails
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"'
 
       - label: "Build release image - aarch64 - CUDA 13.0"
         depends_on: ~
@@ -167,14 +155,11 @@ 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 \
               -f docker/Dockerfile .
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"'
 
       - label: "Build release image - x86_64 - CUDA 12.9"
         depends_on: ~
@@ -192,8 +177,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 .
@@ -201,7 +184,6 @@ steps:
           # re-tag to default image tag and push, just in case arm64 build fails
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"'
 
       - label: "Build release image - aarch64 - CUDA 12.9"
         depends_on: ~
@@ -219,13 +201,10 @@ 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 .
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"'
 
       - label: "Build release image - x86_64 - CUDA 13.0 - Ubuntu 24.04"
         depends_on: ~
@@ -245,8 +224,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 \
@@ -254,7 +231,6 @@ steps:
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"'
 
       - label: "Build release image - aarch64 - CUDA 13.0 - Ubuntu 24.04"
         depends_on: ~
@@ -274,14 +250,11 @@ 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 \
               -f docker/Dockerfile .
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"'
 
       - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
         depends_on: ~
@@ -301,15 +274,12 @@ 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 .
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'
 
       - label: "Build release image - aarch64 - CUDA 12.9 - Ubuntu 24.04"
         depends_on: ~
@@ -329,20 +299,16 @@ 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 .
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'
 
       - block: "Build release image for x86_64 CPU"
         key: block-cpu-release-image-build
         depends_on: ~
 
       - label: "Build release image - x86_64 - CPU"
-        key: build-cpu-release-image-x86
         depends_on:
           - block-cpu-release-image-build
           - input-release-version
@@ -353,7 +319,6 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --build-arg VLLM_CPU_X86=true --tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version) --tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:latest --progress plain --target vllm-openai -f docker/Dockerfile.cpu ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:latest"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -362,8 +327,7 @@ steps:
         depends_on: ~
 
       - label: "Build release image - arm64 - CPU"
-        key: build-cpu-release-image-arm64
-        depends_on:
+        depends_on: 
           - block-arm64-cpu-release-image-build
           - input-release-version
         agents:
@@ -373,7 +337,6 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version) --tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:latest --progress plain --target vllm-openai -f docker/Dockerfile.cpu ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:latest"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -391,7 +354,15 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"'
+
+      - label: "Annotate release workflow - CUDA 13.0"
+        depends_on:
+          - create-multi-arch-manifest
+        id: annotate-release-workflow
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/annotate-release.sh"
 
       - label: "Create multi-arch manifest - CUDA 12.9"
         depends_on:
@@ -404,7 +375,6 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"'
 
       - label: "Create multi-arch manifest - CUDA 13.0 - Ubuntu 24.04"
         depends_on:
@@ -417,7 +387,6 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-ubuntu2404 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"'
 
       - label: "Create multi-arch manifest - CUDA 12.9 - Ubuntu 24.04"
         depends_on:
@@ -430,7 +399,6 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129-ubuntu2404 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"'
 
       - label: "Publish nightly multi-arch image to DockerHub"
         depends_on:
@@ -468,6 +436,24 @@ steps:
           DOCKER_BUILDKIT: "1"
           DOCKERHUB_USERNAME: "vllmbot"
 
+  - group: "Publish wheels"
+    key: "publish-wheels"
+    steps:
+      - block: "Confirm update release wheels to PyPI (experimental, use with caution)?"
+        key: block-upload-release-wheels
+        depends_on:
+          - input-release-version
+          - build-wheels
+
+      - label: "Upload release wheels to PyPI"
+        depends_on:
+          - block-upload-release-wheels
+        id: upload-release-wheels
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/upload-release-wheels-pypi.sh"
+
   # =============================================================================
   # ROCm Release Pipeline (x86_64 only)
   # =============================================================================
@@ -581,7 +567,7 @@ steps:
           echo ""
           echo " Build complete - Image and wheels cached"
         fi
-
+        
     artifact_paths:
       - "artifacts/rocm-base-wheels/*.whl"
     env:
@@ -737,7 +723,7 @@ steps:
       - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
     env:
       S3_BUCKET: "vllm-wheels"
-      VARIANT: "rocm722"
+      VARIANT: "rocm721"
 
   # ROCm Job 6: Build ROCm Release Docker Image
   - label: ":docker: Build release image - x86_64 - ROCm"
@@ -797,7 +783,7 @@ steps:
         
         # Push to ECR
         docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm
-
+        
         echo ""
         echo " Successfully built and pushed ROCm release image"
         echo "   Image: public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm"
@@ -824,60 +810,3 @@ steps:
     env:
       DOCKER_BUILDKIT: "1"
       DOCKERHUB_USERNAME: "vllmbot"
-
-  # =============================================================================
-  # Publish to DockerHub and PyPI (at the end so all builds complete first)
-  # =============================================================================
-
-  - block: "Publish release images to DockerHub"
-    key: block-publish-release-images
-    depends_on:
-      - create-multi-arch-manifest
-      - create-multi-arch-manifest-cuda-12-9
-      - create-multi-arch-manifest-ubuntu2404
-      - create-multi-arch-manifest-cuda-12-9-ubuntu2404
-      - build-rocm-release-image
-      - input-release-version
-      # Wait for CPU builds if their block steps were unblocked, so publish
-      # doesn't race the in-progress CPU build. allow_failure lets publish
-      # proceed when the operator legitimately leaves the CPU block steps
-      # unblocked or the CPU build fails.
-      - step: build-cpu-release-image-x86
-        allow_failure: true
-      - step: build-cpu-release-image-arm64
-        allow_failure: true
-    if: build.env("NIGHTLY") != "1"
-
-  - label: "Publish release images to DockerHub"
-    depends_on:
-      - block-publish-release-images
-    key: publish-release-images-dockerhub
-    agents:
-      queue: small_cpu_queue_release
-    commands:
-      - "bash .buildkite/scripts/publish-release-images.sh"
-    plugins:
-      - docker-login#v3.0.0:
-          username: vllmbot
-          password-env: DOCKERHUB_TOKEN
-    env:
-      DOCKER_BUILDKIT: "1"
-      DOCKERHUB_USERNAME: "vllmbot"
-
-  - group: "Publish wheels"
-    key: "publish-wheels"
-    steps:
-      - block: "Confirm update release wheels to PyPI (experimental, use with caution)?"
-        key: block-upload-release-wheels
-        depends_on:
-          - input-release-version
-          - build-wheels
-
-      - label: "Upload release wheels to PyPI"
-        depends_on:
-          - block-upload-release-wheels
-        id: upload-release-wheels
-        agents:
-          queue: small_cpu_queue_release
-        commands:
-          - "bash .buildkite/scripts/upload-release-wheels-pypi.sh"

.buildkite/scripts/annotate-build-artifact.sh

@@ -1,9 +0,0 @@
-#!/bin/bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Append a build artifact line to the Buildkite annotation.
-# Usage: annotate-build-artifact.sh <label> <value>
-set -e
-echo "- **${1}**: \`${2}\`" | \
-  buildkite-agent annotate --append --style 'info' --context 'release-artifacts'

.buildkite/scripts/annotate-release.sh

@@ -0,0 +1,119 @@
+#!/bin/bash
+
+set -ex
+
+# Get release version, default to 1.0.0.dev for nightly/per-commit builds
+RELEASE_VERSION=$(buildkite-agent meta-data get release-version 2>/dev/null | sed 's/^v//')
+if [ -z "${RELEASE_VERSION}" ]; then
+  RELEASE_VERSION="1.0.0.dev"
+fi
+
+ROCM_BASE_CACHE_KEY=$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
+
+buildkite-agent annotate --style 'info' --context 'release-workflow' << EOF
+To download the wheel (by commit):
+\`\`\`
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_35_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_35_aarch64.whl .
+
+(Optional) For CUDA 12.9:
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu129-cp38-abi3-manylinux_2_31_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu129-cp38-abi3-manylinux_2_31_aarch64.whl .
+
+(Optional) For CPU:
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cpu-cp38-abi3-manylinux_2_35_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cpu-cp38-abi3-manylinux_2_35_aarch64.whl .
+\`\`\`
+
+
+To download and upload the image:
+
+\`\`\`
+# Download images:
+
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu129
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu129
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm
+docker pull public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:v${RELEASE_VERSION}
+docker pull public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:v${RELEASE_VERSION}
+
+# Tag and push images:
+
+## CUDA
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64 vllm/vllm-openai:x86_64
+docker tag vllm/vllm-openai:x86_64 vllm/vllm-openai:latest-x86_64
+docker tag vllm/vllm-openai:x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
+docker push vllm/vllm-openai:latest-x86_64
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu129 vllm/vllm-openai:x86_64-cu129
+docker tag vllm/vllm-openai:x86_64-cu129 vllm/vllm-openai:latest-x86_64-cu129
+docker tag vllm/vllm-openai:x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
+docker push vllm/vllm-openai:latest-x86_64-cu129
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64 vllm/vllm-openai:aarch64
+docker tag vllm/vllm-openai:aarch64 vllm/vllm-openai:latest-aarch64
+docker tag vllm/vllm-openai:aarch64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
+docker push vllm/vllm-openai:latest-aarch64
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu129 vllm/vllm-openai:aarch64-cu129
+docker tag vllm/vllm-openai:aarch64-cu129 vllm/vllm-openai:latest-aarch64-cu129
+docker tag vllm/vllm-openai:aarch64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
+docker push vllm/vllm-openai:latest-aarch64-cu129
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
+
+## ROCm
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT} vllm/vllm-openai-rocm:latest
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT} vllm/vllm-openai-rocm:v${RELEASE_VERSION}
+docker push vllm/vllm-openai-rocm:latest
+docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base vllm/vllm-openai-rocm:latest-base
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
+docker push vllm/vllm-openai-rocm:latest-base
+docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
+
+## CPU
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:v${RELEASE_VERSION} vllm/vllm-openai-cpu:x86_64
+docker tag vllm/vllm-openai-cpu:x86_64 vllm/vllm-openai-cpu:latest-x86_64
+docker tag vllm/vllm-openai-cpu:x86_64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
+docker push vllm/vllm-openai-cpu:latest-x86_64
+docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:v${RELEASE_VERSION} vllm/vllm-openai-cpu:arm64
+docker tag vllm/vllm-openai-cpu:arm64 vllm/vllm-openai-cpu:latest-arm64
+docker tag vllm/vllm-openai-cpu:arm64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
+docker push vllm/vllm-openai-cpu:latest-arm64
+docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
+
+# Create multi-arch manifest:
+
+docker manifest rm vllm/vllm-openai:latest
+docker manifest create vllm/vllm-openai:latest vllm/vllm-openai:latest-x86_64 vllm/vllm-openai:latest-aarch64
+docker manifest create vllm/vllm-openai:v${RELEASE_VERSION} vllm/vllm-openai:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
+docker manifest push vllm/vllm-openai:latest
+docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}
+
+docker manifest rm vllm/vllm-openai:latest-cu129
+docker manifest create vllm/vllm-openai:latest-cu129 vllm/vllm-openai:latest-x86_64-cu129 vllm/vllm-openai:latest-aarch64-cu129
+docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
+docker manifest push vllm/vllm-openai:latest-cu129
+docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu129
+
+docker manifest rm vllm/vllm-openai-cpu:latest || true
+docker manifest create vllm/vllm-openai-cpu:latest vllm/vllm-openai-cpu:latest-x86_64 vllm/vllm-openai-cpu:latest-arm64
+docker manifest create vllm/vllm-openai-cpu:v${RELEASE_VERSION} vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
+docker manifest push vllm/vllm-openai-cpu:latest
+docker manifest push vllm/vllm-openai-cpu:v${RELEASE_VERSION}
+\`\`\`
+EOF

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

@@ -1,55 +0,0 @@
-#!/bin/bash
-# Usage: ./ci-fetch-log.sh <buildkite_job_url> [output_file]
-#        ./ci-fetch-log.sh <build_number> <job_uuid> [output_file]
-#
-# Downloads the raw log for a Buildkite job from the public, unauthenticated
-# /organizations/<org>/pipelines/<pipeline>/builds/<n>/jobs/<uuid>/download
-# endpoint, then strips ANSI/timestamps via ci-clean-log.sh.
-#
-# Find <build_number> and <job_uuid> via:
-#   gh pr checks <PR> --repo vllm-project/vllm
-# Each failing row's URL is .../builds/<build_number>#<job_uuid>.
-
-set -euo pipefail
-
-ORG="vllm"
-PIPELINE="ci"
-
-usage() {
-    echo "Usage: $0 <buildkite_job_url> [output_file]"
-    echo "       $0 <build_number> <job_uuid> [output_file]"
-    exit 1
-}
-
-if [ $# -lt 1 ]; then usage; fi
-
-if [[ "$1" == https://* ]]; then
-    BUILD=$(echo "$1" | sed -nE 's#.*/builds/([0-9]+).*#\1#p')
-    JOB=$(echo "$1" | grep -oE '[0-9a-f]{8}-[0-9a-f-]+' | head -n 1)
-    OUT="${2:-ci-${BUILD}-${JOB:0:8}.log}"
-else
-    if [ $# -lt 2 ]; then usage; fi
-    BUILD="$1"
-    JOB="$2"
-    OUT="${3:-ci-${BUILD}-${JOB:0:8}.log}"
-fi
-
-if [ -z "$BUILD" ] || [ -z "$JOB" ]; then
-    echo "Could not parse build number or job UUID from: $1" >&2
-    usage
-fi
-
-COOKIES=$(mktemp)
-trap 'rm -f "$COOKIES"' EXIT
-
-# Buildkite issues a session cookie on first hit; subsequent /download needs it.
-curl -fsSL -c "$COOKIES" -A "vllm-ci-fetch-log" \
-    "https://buildkite.com/${ORG}/${PIPELINE}/builds/${BUILD}" -o /dev/null
-
-curl -fsSL -b "$COOKIES" -A "vllm-ci-fetch-log" \
-    "https://buildkite.com/organizations/${ORG}/pipelines/${PIPELINE}/builds/${BUILD}/jobs/${JOB}/download" \
-    -o "$OUT"
-
-bash "$(dirname "$0")/ci-clean-log.sh" "$OUT"
-
-echo "$OUT"

.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
 ###############################################################################
@@ -281,11 +378,9 @@ HF_MOUNT="/root/.cache/huggingface"
 # double-quotes will have been stripped by the calling shell.
 if [[ -n "${VLLM_TEST_COMMANDS:-}" ]]; then
   commands="${VLLM_TEST_COMMANDS}"
-  commands_source="env"
   echo "Commands sourced from VLLM_TEST_COMMANDS (quoting preserved)"
 else
   commands="$*"
-  commands_source="argv"
   if [[ -z "$commands" ]]; then
     echo "Error: No test commands provided." >&2
     echo "Usage:" >&2
@@ -302,16 +397,11 @@ fi
 
 echo "Raw commands: $commands"
 
-# Only try to repair stripped pytest -m/-k quoting in legacy argv mode.
-# VLLM_TEST_COMMANDS preserves inner quoting already, and re-quoting that path
-# can corrupt embedded echo strings or otherwise well-formed shell fragments.
-if [[ "$commands_source" == "argv" ]]; then
-  commands=$(re_quote_pytest_markers "$commands")
-  echo "After re-quoting: $commands"
-else
-  echo "Skipping re-quoting for VLLM_TEST_COMMANDS input"
-fi
+# Fix quoting before ROCm overrides (so overrides see correct structure)
+commands=$(re_quote_pytest_markers "$commands")
+echo "After re-quoting: $commands"
 
+commands=$(apply_rocm_test_overrides "$commands")
 echo "Final commands: $commands"
 
 MYPYTHONPATH=".."

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

@@ -31,21 +31,6 @@ function cpu_tests() {
     set -e
     pip list"
 
-  # Run kernel tests
-  docker exec cpu-test bash -c "
-    set -e
-    pytest -x -v -s tests/kernels/test_onednn.py
-    pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
-    pytest -x -v -s tests/kernels/core/test_cpu_activation.py
-    pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic"
-
-  # skip tests requiring model downloads if HF_TOKEN is not set
-  # due to rate-limits
-  if [ -z "$HF_TOKEN" ]; then
-    echo "Warning: HF_TOKEN is not set. Skipping tests that require model downloads."
-    return
-  fi
-
   # offline inference
   docker exec cpu-test bash -c "
     set -e
@@ -61,6 +46,13 @@ function cpu_tests() {
     set -e
     pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs"
 
+  # Run kernel tests
+  docker exec cpu-test bash -c "
+    set -e
+    pytest -x -v -s tests/kernels/test_onednn.py
+    pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
+    pytest -x -v -s tests/kernels/core/test_cpu_activation.py
+    pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic"
 
   # basic online serving
   docker exec cpu-test bash -c '
@@ -75,21 +67,6 @@ function cpu_tests() {
       --num-prompts 20 \
       --endpoint /v1/completions
     kill -s SIGTERM $server_pid &'
-
-  # smoke test for Gated DeltaNet
-  docker exec cpu-test bash -c '
-    set -e
-    VLLM_CPU_OMP_THREADS_BIND=$E2E_OMP_THREADS vllm serve Qwen/Qwen3.5-0.8B --max-model-len 2048 &
-    server_pid=$!
-    timeout 600 bash -c "until curl localhost:8000/v1/models; do sleep 1; done" || exit 1
-    vllm bench serve \
-      --backend vllm \
-      --dataset-name random \
-      --model Qwen/Qwen3.5-0.8B \
-      --num-prompts 20 \
-      --endpoint /v1/completions
-    kill -s SIGTERM $server_pid &'
-
 }
 
 # All of CPU tests are expected to be finished less than 40 mins.

.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 ---"
 
@@ -136,6 +136,8 @@ run_and_track_test 3 "test_accuracy.py::test_lm_eval_accuracy_v1_engine" \
     "python3 -m pytest -s -v /workspace/vllm/tests/entrypoints/llm/test_accuracy.py::test_lm_eval_accuracy_v1_engine"
 run_and_track_test 4 "test_quantization_accuracy.py" \
     "python3 -m pytest -s -v /workspace/vllm/tests/tpu/test_quantization_accuracy.py"
+run_and_track_test 5 "examples/offline_inference/tpu.py" \
+    "python3 /workspace/vllm/examples/offline_inference/tpu.py"
 run_and_track_test 6 "test_tpu_model_runner.py" \
     "python3 -m pytest -s -v /workspace/vllm/tests/v1/tpu/worker/test_tpu_model_runner.py"
 run_and_track_test 7 "test_sampler.py" \

.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/publish-release-images.sh

@@ -1,180 +0,0 @@
-#!/bin/bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Publish release Docker images from ECR to DockerHub.
-# Pulls per-arch images, tags with latest and versioned tags, pushes them,
-# then creates and pushes multi-arch manifests.
-
-set -euo pipefail
-
-RELEASE_VERSION=$(buildkite-agent meta-data get release-version --default "" | sed 's/^v//')
-if [ -z "${RELEASE_VERSION}" ]; then
-  echo "ERROR: release-version metadata not set"
-  exit 1
-fi
-
-COMMIT="$BUILDKITE_COMMIT"
-ROCM_BASE_CACHE_KEY=$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
-
-echo "========================================"
-echo "Publishing release images v${RELEASE_VERSION}"
-echo "  Commit: ${COMMIT}"
-echo "  ROCm base cache key: ${ROCM_BASE_CACHE_KEY}"
-echo "========================================"
-
-# Login to ECR to pull staging images
-aws ecr-public get-login-password --region us-east-1 | \
-  docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-
-# ---- CUDA (default: 13.0) ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64 vllm/vllm-openai:latest-x86_64
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
-docker push vllm/vllm-openai:latest-x86_64
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64 vllm/vllm-openai:latest-aarch64
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
-docker push vllm/vllm-openai:latest-aarch64
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
-
-docker manifest rm vllm/vllm-openai:latest || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION} || true
-docker manifest create vllm/vllm-openai:latest vllm/vllm-openai:latest-x86_64 vllm/vllm-openai:latest-aarch64
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION} vllm/vllm-openai:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
-docker manifest push vllm/vllm-openai:latest
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}
-
-# ---- CUDA 12.9 ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129 vllm/vllm-openai:latest-x86_64-cu129
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
-docker push vllm/vllm-openai:latest-x86_64-cu129
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129 vllm/vllm-openai:latest-aarch64-cu129
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-docker push vllm/vllm-openai:latest-aarch64-cu129
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-
-docker manifest rm vllm/vllm-openai:latest-cu129 || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION}-cu129 || true
-docker manifest create vllm/vllm-openai:latest-cu129 vllm/vllm-openai:latest-x86_64-cu129 vllm/vllm-openai:latest-aarch64-cu129
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-docker manifest push vllm/vllm-openai:latest-cu129
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu129
-
-# ---- Ubuntu 24.04 (CUDA 13.0) ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-ubuntu2404
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-ubuntu2404 vllm/vllm-openai:latest-x86_64-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-ubuntu2404
-docker push vllm/vllm-openai:latest-x86_64-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-ubuntu2404 vllm/vllm-openai:latest-aarch64-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-ubuntu2404
-docker push vllm/vllm-openai:latest-aarch64-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-ubuntu2404
-
-docker manifest rm vllm/vllm-openai:latest-ubuntu2404 || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION}-ubuntu2404 || true
-docker manifest create vllm/vllm-openai:latest-ubuntu2404 vllm/vllm-openai:latest-x86_64-ubuntu2404 vllm/vllm-openai:latest-aarch64-ubuntu2404
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-ubuntu2404
-docker manifest push vllm/vllm-openai:latest-ubuntu2404
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-ubuntu2404
-
-# ---- Ubuntu 24.04 (CUDA 12.9) ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129-ubuntu2404
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129-ubuntu2404 vllm/vllm-openai:latest-x86_64-cu129-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129-ubuntu2404
-docker push vllm/vllm-openai:latest-x86_64-cu129-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129-ubuntu2404 vllm/vllm-openai:latest-aarch64-cu129-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129-ubuntu2404
-docker push vllm/vllm-openai:latest-aarch64-cu129-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129-ubuntu2404
-
-docker manifest rm vllm/vllm-openai:latest-cu129-ubuntu2404 || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION}-cu129-ubuntu2404 || true
-docker manifest create vllm/vllm-openai:latest-cu129-ubuntu2404 vllm/vllm-openai:latest-x86_64-cu129-ubuntu2404 vllm/vllm-openai:latest-aarch64-cu129-ubuntu2404
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129-ubuntu2404
-docker manifest push vllm/vllm-openai:latest-cu129-ubuntu2404
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu129-ubuntu2404
-
-# ---- ROCm ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm vllm/vllm-openai-rocm:latest
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm vllm/vllm-openai-rocm:v${RELEASE_VERSION}
-docker push vllm/vllm-openai-rocm:latest
-docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base vllm/vllm-openai-rocm:latest-base
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
-docker push vllm/vllm-openai-rocm:latest-base
-docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
-
-# ---- CPU ----
-# CPU images are behind separate block steps and may not have been built.
-# All-or-nothing: inspect both arches first, then either publish everything
-# (per-arch + multi-arch manifest) or skip everything. Publishing only one
-# arch would leave `:latest-x86_64` pointing at the new release while the
-# `:latest` multi-arch manifest still resolves to the previous release.
-
-CPU_X86_TAG=public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:v${RELEASE_VERSION}
-CPU_ARM_TAG=public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:v${RELEASE_VERSION}
-
-CPU_X86_AVAILABLE=false
-CPU_ARM_AVAILABLE=false
-docker manifest inspect "${CPU_X86_TAG}" >/dev/null 2>&1 && CPU_X86_AVAILABLE=true
-docker manifest inspect "${CPU_ARM_TAG}" >/dev/null 2>&1 && CPU_ARM_AVAILABLE=true
-
-if [ "$CPU_X86_AVAILABLE" = "true" ] && [ "$CPU_ARM_AVAILABLE" = "true" ]; then
-  docker pull "${CPU_X86_TAG}"
-  docker tag "${CPU_X86_TAG}" vllm/vllm-openai-cpu:latest-x86_64
-  docker tag "${CPU_X86_TAG}" vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
-  docker push vllm/vllm-openai-cpu:latest-x86_64
-  docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
-
-  docker pull "${CPU_ARM_TAG}"
-  docker tag "${CPU_ARM_TAG}" vllm/vllm-openai-cpu:latest-arm64
-  docker tag "${CPU_ARM_TAG}" vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
-  docker push vllm/vllm-openai-cpu:latest-arm64
-  docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
-
-  docker manifest rm vllm/vllm-openai-cpu:latest || true
-  docker manifest rm vllm/vllm-openai-cpu:v${RELEASE_VERSION} || true
-  docker manifest create vllm/vllm-openai-cpu:latest vllm/vllm-openai-cpu:latest-x86_64 vllm/vllm-openai-cpu:latest-arm64
-  docker manifest create vllm/vllm-openai-cpu:v${RELEASE_VERSION} vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
-  docker manifest push vllm/vllm-openai-cpu:latest
-  docker manifest push vllm/vllm-openai-cpu:v${RELEASE_VERSION}
-elif [ "$CPU_X86_AVAILABLE" = "false" ] && [ "$CPU_ARM_AVAILABLE" = "false" ]; then
-  echo "WARNING: Neither CPU image found in ECR, skipping CPU publish (ensure block-cpu-release-image-build and block-arm64-cpu-release-image-build were unblocked and the builds finished pushing)"
-else
-  # Partial state: one arch built, the other did not. Fail loudly rather than
-  # ship a Docker Hub state where `:latest-${arch}` and `:latest` (multi-arch)
-  # disagree on which release they point at.
-  echo "ERROR: Partial CPU build detected (x86_64=${CPU_X86_AVAILABLE}, arm64=${CPU_ARM_AVAILABLE})."
-  echo "       Refusing to publish to avoid split-tag drift between per-arch and multi-arch tags."
-  echo "       Re-run the missing CPU build and retry, or manually publish if a single-arch release is intended."
-  exit 1
-fi
-
-echo ""
-echo "Successfully published release images for v${RELEASE_VERSION}"

.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh

@@ -51,7 +51,6 @@ vllm serve "$MODEL" \
   --offload-num-in-group 2 \
   --offload-prefetch-step 1 \
   --offload-params w13_weight w2_weight \
-  --generation-config vllm \
   --port "$PORT" \
   ${EXTRA_ARGS+"${EXTRA_ARGS[@]}"} &
 SERVER_PID=$!

.buildkite/scripts/upload-release-wheels-pypi.sh

@@ -39,17 +39,10 @@ fi
 
 set -x # avoid printing secrets above
 
-# install uv if not already available
-if ! command -v uv &> /dev/null; then
-  curl -LsSf https://astral.sh/uv/install.sh | UV_VERSION=0.11.14 sh
-  export PATH="$HOME/.local/bin:$PATH"
-fi
-
-# install twine and sdist build prerequisites using uv with Python 3.12
-uv venv --python 3.12 /tmp/vllm-release-env
+# install twine from pypi
+python3 -m venv /tmp/vllm-release-env
 source /tmp/vllm-release-env/bin/activate
-uv pip install twine
-uv pip install -r requirements/build/cuda.txt
+pip install twine
 python3 -m twine --version
 
 # copy release wheels to local directory

.buildkite/test-amd.yaml

@@ -230,6 +230,7 @@ steps:
   - tests/entrypoints/llm/test_collective_rpc.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s entrypoints/llm/test_collective_rpc.py
   - pytest -v -s ./compile/fullgraph/test_basic_correctness.py
   - pytest -v -s ./compile/test_wrapper.py
@@ -271,6 +272,7 @@ steps:
   - tests/v1/worker/test_worker_memory_snapshot.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - VLLM_TEST_SAME_HOST=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - VLLM_TEST_SAME_HOST=1 VLLM_TEST_WITH_DEFAULT_DEVICE_SET=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s v1/shutdown
@@ -394,8 +396,8 @@ steps:
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # Features demo
     - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
@@ -460,7 +462,7 @@ steps:
   - tests/lora
   - vllm/platforms/rocm.py
   commands:
-  - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_qwen3_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py
+  - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_llm_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py
 
 #------------------------------------------------------  mi250 · model_executor  -------------------------------------------------------#
 
@@ -588,6 +590,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - pip freeze | grep -E 'torch'
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
 - label: Multi-Modal Models (Extended Generation 2) # TBD
@@ -618,7 +621,6 @@ steps:
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
   agent_pool: mi250_1
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -798,6 +800,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]
@@ -836,6 +864,7 @@ steps:
   - tests/entrypoints/openai/test_multi_api_servers.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
@@ -854,7 +883,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - ATTENTION_BACKEND=ROCM_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
+  - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
 
 - label: V1 e2e (2 GPUs) # TBD
   timeout_in_minutes: 180
@@ -901,9 +930,7 @@ steps:
   - tests/renderers
   - tests/standalone_tests/lazy_imports.py
   - tests/tokenizers_
-  - tests/reasoning
   - tests/tool_parsers
-  - tests/parser
   - tests/transformers_utils
   - tests/config
   commands:
@@ -915,9 +942,8 @@ steps:
   - pytest -v -s -m 'cpu_test' multimodal
   - pytest -v -s renderers
   - pytest -v -s tokenizers_
-  - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
+  - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py --ignore=reasoning/test_gemma4_reasoning_parser.py
   - pytest -v -s tool_parsers
-  - pytest -v -s parser
   - pytest -v -s transformers_utils
   - pytest -v -s config
 
@@ -1076,13 +1102,12 @@ steps:
   - vllm/compilation/
   - vllm/model_executor/layers
   - tests/compile/passes/distributed/
-  - tests/compile/fusions_e2e/
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
   - export VLLM_TEST_CLEAN_GPU_MEMORY=1
   - VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/passes/distributed/test_async_tp.py
-  - pytest -v -s tests/compile/fusions_e2e/test_tp2_ar_rms.py::test_tp2_ar_rms_fusions
+  - pytest -v -s tests/compile/passes/distributed/test_sequence_parallelism.py
 
 #-----------------------------------------------------------  mi300 · cuda  ------------------------------------------------------------#
 
@@ -1147,6 +1172,7 @@ steps:
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s tests/distributed/test_context_parallel.py
   - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=allgather_reducescatter --disable-nccl-for-dp-synchronization
 
@@ -1160,6 +1186,7 @@ steps:
   source_file_dependencies:
   - vllm/
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s distributed/test_custom_all_reduce.py
   - torchrun --nproc_per_node=2 distributed/test_ca_buffer_sharing.py
   - TARGET_TEST_SUITE=A100 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
@@ -1179,6 +1206,7 @@ steps:
   - tests/examples/features/data_parallel/data_parallel_offline.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - torchrun --nproc-per-node=4 distributed/test_torchrun_example.py
   - PP_SIZE=2 torchrun --nproc-per-node=4 distributed/test_torchrun_example.py
   - TP_SIZE=4 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
@@ -1224,6 +1252,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - export VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s distributed/test_ray_v2_executor.py
   - pytest -v -s distributed/test_ray_v2_executor_e2e.py
   - pytest -v -s distributed/test_pipeline_parallel.py -k "ray"
@@ -1245,6 +1274,7 @@ steps:
   - vllm/v1/worker/gpu_worker.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
 
 #--------------------------------------------------------  mi300 · entrypoints  --------------------------------------------------------#
@@ -1253,7 +1283,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"
@@ -1274,7 +1303,6 @@ steps:
   agent_pool: mi300_1
   fast_check: true
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -1288,7 +1316,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"
@@ -1299,13 +1326,13 @@ steps:
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
+  - pytest -v -s entrypoints/openai/speech_to_text/
   - pytest -v -s entrypoints/test_chat_utils.py
 
 - label: Entrypoints Integration (API Server openai - Part 3) # TBD
   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"
@@ -1315,21 +1342,7 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
-
-- label: Entrypoints Integration (Speech to Text) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
-  agent_pool: mi300_1
-  fast_check: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/speech_to_text
-  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/speech_to_text
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
 
 - label: Entrypoints Integration (LLM) # TBD
   timeout_in_minutes: 180
@@ -1642,8 +1655,8 @@ steps:
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # Features demo
     - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
@@ -1706,7 +1719,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:
@@ -1754,7 +1766,7 @@ steps:
   - export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
   - pytest -v -s -x lora/test_chatglm3_tp.py
   - pytest -v -s -x lora/test_llama_tp.py
-  - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
+  - pytest -v -s -x lora/test_llm_with_multi_loras.py
   - pytest -v -s -x lora/test_olmoe_tp.py
   - pytest -v -s -x lora/test_gptoss_tp.py
   - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
@@ -1765,7 +1777,6 @@ steps:
   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/
@@ -1791,17 +1802,15 @@ steps:
   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/
   - tests/models/multimodal/generation
   - tests/models/multimodal/test_mapping.py
   commands:
-  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@rocm-7.0-v2.3.0'
-  - 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
-
+  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+  - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
+  - pytest -v -s models/multimodal/test_mapping.py
 
 - label: Multi-Modal Models (Extended Generation 2) # TBD
   timeout_in_minutes: 180
@@ -1813,10 +1822,8 @@ steps:
   - vllm/
   - tests/models/multimodal/generation
   commands:
-  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@rocm-7.0-v2.3.0'
-  - 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 '(not core_model) and (not hybrid_model)'
-
+  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+  - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
 
 - label: Multi-Modal Models (Extended Generation 3) # TBD
   timeout_in_minutes: 180
@@ -1836,7 +1843,6 @@ steps:
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2197,6 +2203,7 @@ steps:
   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/
@@ -2273,6 +2280,7 @@ steps:
   - tests/entrypoints/openai/test_multi_api_servers.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
@@ -2292,6 +2300,7 @@ steps:
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
   - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
   - pytest -v -s tests/v1/distributed/test_dbo.py
@@ -2354,6 +2363,7 @@ steps:
   - tests/distributed/test_utils
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
   - TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
   - TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
@@ -2483,6 +2493,7 @@ steps:
   - tests/entrypoints/llm/test_collective_rpc.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s entrypoints/llm/test_collective_rpc.py
   - pytest -v -s ./compile/fullgraph/test_basic_correctness.py
   - pytest -v -s ./compile/test_wrapper.py
@@ -2507,6 +2518,7 @@ steps:
   - tests/v1/worker/test_worker_memory_snapshot.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - VLLM_TEST_SAME_HOST=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - VLLM_TEST_SAME_HOST=1 VLLM_TEST_WITH_DEFAULT_DEVICE_SET=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s v1/shutdown
@@ -2527,6 +2539,7 @@ steps:
   - tests/distributed/test_multiproc_executor.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s compile/fullgraph/test_basic_correctness.py
   - pytest -v -s distributed/test_pynccl.py
   - pytest -v -s distributed/test_events.py
@@ -2614,7 +2627,6 @@ steps:
   agent_pool: mi325_1
   torch_nightly: true
   parallelism: 2
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2640,7 +2652,6 @@ steps:
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
   agent_pool: mi325_1
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2706,6 +2717,7 @@ steps:
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s tests/distributed/test_context_parallel.py
   - pytest -v -s tests/v1/distributed/test_dbo.py
 
@@ -2736,7 +2748,6 @@ steps:
   agent_pool: mi355_1
   fast_check: true
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2752,7 +2763,6 @@ steps:
   agent_pool: mi355_1
   fast_check: true
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2761,6 +2771,7 @@ steps:
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
+  - pytest -v -s entrypoints/openai/speech_to_text/
   - pytest -v -s entrypoints/test_chat_utils.py
 
 - label: Entrypoints Integration (API Server openai - Part 3) # TBD
@@ -2776,21 +2787,7 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
-
-- label: Entrypoints Integration (Speech to Text) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi355]
-  agent_pool: mi355_1
-  fast_check: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/speech_to_text
-  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/speech_to_text
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
 
 - label: Entrypoints Integration (Pooling) # TBD
   timeout_in_minutes: 180
@@ -2941,8 +2938,8 @@ steps:
   - python3 pooling/embed/vision_embedding_offline.py --seed 0
   # Features demo
   - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-  - python3 deployment/llm_engine_example.py
-  - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+  - python3 offline_inference/llm_engine_example.py
+  - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
   - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
   - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
@@ -3054,7 +3051,7 @@ steps:
   - vllm/
   - tests/models/language/generation
   commands:
-  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@rocm-7.0-v2.3.0'
+  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr'
   - 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 '(not core_model) and (not hybrid_model)'
 
@@ -3062,7 +3059,6 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -3262,7 +3258,6 @@ steps:
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -3288,7 +3283,6 @@ steps:
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -3329,7 +3323,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - ATTENTION_BACKEND=ROCM_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
+  - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
 
 - label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
   timeout_in_minutes: 180

.buildkite/test_areas/attention.yaml

@@ -17,7 +17,7 @@ steps:
 - label: V1 attention (B200)
   key: v1-attention-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
     - vllm/config/attention.py
     - vllm/model_executor/layers/attention

.buildkite/test_areas/benchmarks.yaml

@@ -14,7 +14,7 @@ steps:
 
 - label: Attention Benchmarks Smoke Test (B200)
   key: attention-benchmarks-smoke-test-b200
-  device: b200-k8s
+  device: b200
   num_gpus: 2
   optional: true
   working_dir: "/vllm-workspace/"

.buildkite/test_areas/compile.yaml

@@ -43,7 +43,7 @@ steps:
   key: asynctp-correctness-tests-b200
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   commands:
@@ -68,7 +68,7 @@ steps:
   key: fusion-and-compile-unit-tests-2xb200
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
   - csrc/quantization/fp4/
   - vllm/model_executor/layers/quantization/
@@ -137,7 +137,7 @@ steps:
   key: fusion-e2e-config-sweep-b200
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   num_devices: 1
   optional: true
   commands:
@@ -209,7 +209,7 @@ steps:
   key: fusion-e2e-tp2-b200
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   num_devices: 2
   source_file_dependencies:
     - csrc/quantization/

.buildkite/test_areas/cuda.yaml

@@ -7,11 +7,7 @@ steps:
   timeout_in_minutes: 15
   device: h200_18gb
   source_file_dependencies:
-  - vllm/envs.py
-  - vllm/logger.py
-  - vllm/platforms/
-  - vllm/plugins/
-  - vllm/utils/
+  - vllm/
   - tests/cuda
   commands:
     - pytest -v -s cuda/test_cuda_context.py
@@ -27,5 +23,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

@@ -8,7 +8,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
@@ -19,7 +19,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
@@ -31,7 +31,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
@@ -43,7 +43,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
@@ -51,11 +51,11 @@ steps:
 
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
   key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
@@ -67,7 +67,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
@@ -83,7 +83,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - vllm/v1/worker/kv_connector_model_runner_mixin.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
@@ -96,7 +96,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading/

.buildkite/test_areas/distributed.yaml

@@ -212,7 +212,7 @@ steps:
 
 - label: Distributed Tests (2 GPUs)(B200)
   key: distributed-tests-2-gpus-b200
-  device: b200-k8s
+  device: b200
   optional: true
   working_dir: "/vllm-workspace/"
   num_devices: 2

.buildkite/test_areas/e2e_integration.yaml

@@ -25,7 +25,7 @@ steps:
 - label: Qwen3-30B-A3B-FP8-block Accuracy (B200)
   key: qwen3-30b-a3b-fp8-block-accuracy-b200
   timeout_in_minutes: 60
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   working_dir: "/vllm-workspace"

.buildkite/test_areas/engine.yaml

@@ -7,25 +7,14 @@ steps:
   timeout_in_minutes: 15
   device: h200_18gb
   source_file_dependencies:
-  - vllm/compilation/
-  - vllm/config/
-  - vllm/engine/
-  - vllm/entrypoints/logger.py
-  - vllm/envs.py
-  - vllm/logger.py
-  - vllm/logging_utils/
-  - vllm/platforms/
-  - vllm/sequence.py
-  - vllm/triton_utils/
-  - vllm/utils/
+  - vllm/
   - tests/engine
   - tests/test_sequence
   - tests/test_config
   - tests/test_logger
   - tests/test_vllm_port
-  - tests/test_jit_monitor.py
   commands:
-  - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py test_jit_monitor.py
+  - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
 
 - label: Engine (1 GPU)
   key: engine-1-gpu
@@ -36,12 +25,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,12 +35,6 @@ 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)
   key: e2e-core-1-gpu
@@ -74,27 +51,16 @@ steps:
   optional: true
   num_devices: 2
   source_file_dependencies:
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/logger.py
-    - vllm/logging_utils/
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/sampling_params.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
-    - tests/v1/e2e/spec_decode
+    - vllm/
+    - tests/v1/e2e
   commands:
     # Only run tests that need exactly 2 GPUs
     - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
+  mirror:
+    amd:
+      device: mi325_2
+      depends_on:
+      - image-build-amd
 
 - label: V1 e2e (4 GPUs)
   key: v1-e2e-4-gpus
@@ -102,27 +68,16 @@ steps:
   optional: true
   num_devices: 4
   source_file_dependencies:
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/logger.py
-    - vllm/logging_utils/
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/sampling_params.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
-    - tests/v1/e2e/spec_decode
+    - vllm/
+    - tests/v1/e2e
   commands:
     # Only run tests that need 4 GPUs
     - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle_correctness_heavy"
+  mirror:
+    amd:
+      device: mi325_4
+      depends_on:
+      - image-build-amd
 
 - label: V1 e2e (4xH100)
   key: v1-e2e-4xh100

.buildkite/test_areas/entrypoints.yaml

@@ -11,7 +11,7 @@ steps:
   - tests/entrypoints/
   commands:
   - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text
+  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling
 
 - label: Entrypoints Integration (LLM)
   key: entrypoints-integration-llm
@@ -26,11 +26,6 @@ steps:
   - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
   - pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
   - pytest -v -s entrypoints/offline_mode # Needs to avoid interference with other tests
-  mirror:
-    amd:
-      device: mi300_1
-      depends_on:
-      - image-build-amd
 
 - label: Entrypoints Integration (API Server openai - Part 1)
   key: entrypoints-integration-api-server-openai-part-1
@@ -45,11 +40,11 @@ steps:
   - 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
+      device: mi325_1
       depends_on:
       - image-build-amd
 
+
 - label: Entrypoints Integration (API Server openai - Part 2)
   key: entrypoints-integration-api-server-openai-part-2
   timeout_in_minutes: 50
@@ -60,11 +55,11 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
+  - pytest -v -s entrypoints/openai/speech_to_text/
   - pytest -v -s entrypoints/test_chat_utils.py
   mirror:
     amd:
-      device: mi300_1
-      timeout_in_minutes: 60
+      device: mi325_1
       depends_on:
       - image-build-amd
 
@@ -79,13 +74,7 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
-  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/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
 
 - label: Entrypoints Integration (API Server 2)
   key: entrypoints-integration-api-server-2
@@ -101,22 +90,6 @@ 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)
-  key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 50
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/speech_to_text
-  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/speech_to_text
 
 - label: Entrypoints Integration (Pooling)
   key: entrypoints-integration-pooling
@@ -147,5 +120,5 @@ steps:
   - csrc/
   - vllm/entrypoints/openai/
   - vllm/model_executor/models/whisper.py
-  commands: # LMEval
+  commands: # LMEval+Transcription WER check
   - pytest -s entrypoints/openai/correctness/

.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
@@ -104,7 +91,6 @@ steps:
 
 - label: Kernels KDA Test
   timeout_in_minutes: 20
-  device: h200_18gb
   source_file_dependencies:
   - vllm/model_executor/layers/fla/ops/kda.py
   - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
@@ -119,11 +105,7 @@ steps:
   device: h100
   num_devices: 1
   source_file_dependencies:
-  - cmake/external_projects/deepgemm.cmake
   - tools/install_deepgemm.sh
-  - tools/build_deepgemm_C.py
-  - tools/setup_deepgemm_pythons.sh
-  - tools/check_wheel_deepgemm.py
   - vllm/utils/deep_gemm.py
   - vllm/model_executor/layers/fused_moe
   - vllm/model_executor/layers/quantization
@@ -133,7 +115,6 @@ steps:
   - tests/kernels/attention/test_deepgemm_attention.py
   - tests/quantization/test_cutlass_w4a16.py
   commands:
-    - python3 ../tools/check_wheel_deepgemm.py
     - pytest -v -s kernels/quantization/test_block_fp8.py
     - pytest -v -s kernels/moe/test_deepgemm.py
     - pytest -v -s kernels/moe/test_batched_deepgemm.py
@@ -144,7 +125,7 @@ steps:
   key: kernels-b200
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   # optional: true
   source_file_dependencies:
   - csrc/quantization/fp4/
@@ -231,7 +212,7 @@ steps:
 - label: Kernels Fp4 MoE Test (B200)
   key: kernels-fp4-moe-test-b200
   timeout_in_minutes: 60
-  device: b200-k8s
+  device: b200
   num_devices: 1
   optional: true
   commands:
@@ -261,7 +242,7 @@ steps:
 - label: Kernels FusedMoE Layer Test (2 B200s)
   key: kernels-fusedmoe-layer-test-2-b200s
   timeout_in_minutes: 90
-  device: b200-k8s
+  device: b200
   num_devices: 2
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/

.buildkite/test_areas/lm_eval.yaml

@@ -40,7 +40,7 @@ steps:
 - label: LM Eval Small Models (B200)
   key: lm-eval-small-models-b200
   timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
   - csrc/
@@ -48,22 +48,10 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt
 
-- label: LM Eval Large Models (B200, EP)
-  key: lm-eval-large-models-b200-ep
-  timeout_in_minutes: 120
-  device: b200-k8s
-  optional: true
-  num_devices: 2
-  source_file_dependencies:
-  - csrc/
-  - vllm/model_executor/layers/quantization
-  commands:
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell-ep.txt
-
 - label: LM Eval Qwen3.5 Models (B200)
   key: lm-eval-qwen3-5-models-b200
   timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   source_file_dependencies:
@@ -96,7 +84,7 @@ steps:
   
 - label: MoE Refactor Integration Test (B200 - TEMPORARY)
   key: moe-refactor-integration-test-b200-temporary
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   commands:
@@ -104,7 +92,7 @@ steps:
 
 - label: MoE Refactor Integration Test (B200 DP - TEMPORARY)
   key: moe-refactor-integration-test-b200-dp-temporary
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   commands:
@@ -114,7 +102,6 @@ steps:
 - label: LM Eval TurboQuant KV Cache
   key: lm-eval-turboquant-kv-cache
   timeout_in_minutes: 75
-  device: h200_18gb
   source_file_dependencies:
   - vllm/model_executor/layers/quantization/turboquant/
   - vllm/v1/attention/backends/turboquant_attn.py
@@ -140,7 +127,7 @@ steps:
 - label: GPQA Eval (GPT-OSS) (B200)
   key: gpqa-eval-gpt-oss-b200
   timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   source_file_dependencies:
@@ -150,10 +137,3 @@ steps:
   commands:
     - uv pip install --system 'gpt-oss[eval]==0.0.5'
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-b200.txt
-
-- label: MRCR Eval Small Models
-  timeout_in_minutes: 30
-  source_file_dependencies:
-  - tests/evals/mrcr/
-  commands:
-    - pytest -s -v evals/mrcr/test_mrcr_correctness.py --config-list-file=evals/mrcr/configs/models-small.txt

.buildkite/test_areas/lora.yaml

@@ -9,7 +9,7 @@ steps:
   - vllm/lora
   - tests/lora
   commands:
-    - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_qwen3_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py 
+    - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_llm_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py 
   parallelism: 4
 
 
@@ -19,7 +19,6 @@ steps:
   num_devices: 4
   source_file_dependencies:
   - vllm/lora
-  - vllm/model_executor/layers/fused_moe/
   - tests/lora
   commands:
     # FIXIT: find out which code initialize cuda before running the test
@@ -31,7 +30,7 @@ steps:
     # requires multi-GPU testing for validation.
     - pytest -v -s -x lora/test_chatglm3_tp.py
     - pytest -v -s -x lora/test_llama_tp.py
-    - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
+    - pytest -v -s -x lora/test_llm_with_multi_loras.py
     - pytest -v -s -x lora/test_olmoe_tp.py
     - pytest -v -s -x lora/test_gptoss_tp.py
     - pytest -v -s -x lora/test_qwen35_densemodel_lora.py

.buildkite/test_areas/misc.yaml

@@ -6,37 +6,24 @@ steps:
   key: v1-spec-decode
   timeout_in_minutes: 30
   source_file_dependencies:
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/inputs/
-    - vllm/model_executor/
-    - vllm/platforms/
-    - vllm/sampling_params.py
-    - vllm/transformers_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/v1/spec_decode
   commands:
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
     # TODO: create another `optional` test group for slow tests
     - pytest -v -s -m 'not slow_test' v1/spec_decode
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
 
 - label: V1 Sample + Logits
   key: v1-sample-logits
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/inputs/
-    - vllm/logger.py
-    - vllm/model_executor/
-    - vllm/platforms/
-    - vllm/sampling_params.py
-    - vllm/transformers_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/v1/sample
     - tests/v1/logits_processors
     - tests/v1/test_oracle.py
@@ -51,7 +38,7 @@ steps:
     - pytest -v -s v1/test_outputs.py
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
@@ -59,27 +46,10 @@ steps:
   key: v1-core-kv-metrics
   timeout_in_minutes: 30
   source_file_dependencies:
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/entrypoints/pooling/
-    - vllm/inputs/
-    - vllm/lora/
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/outputs.py
-    - vllm/platforms/
-    - vllm/pooling_params.py
-    - vllm/profiler/
-    - vllm/sampling_params.py
-    - vllm/tokenizers/
-    - vllm/transformers_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/v1/core
     - tests/v1/executor
     - tests/v1/kv_offload
-    - tests/v1/simple_kv_offload
     - tests/v1/worker
     - tests/v1/kv_connector/unit
     - tests/v1/metrics
@@ -91,34 +61,24 @@ steps:
     - pytest -v -s -m 'not cpu_test' v1/core
     - pytest -v -s v1/executor
     - pytest -v -s v1/kv_offload
-    - pytest -v -s v1/simple_kv_offload
     - pytest -v -s v1/worker
     - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/metrics
     # Integration test for streaming correctness (requires special branch).
     - 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
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
 
 - label: V1 Others (CPU)
   key: v1-others-cpu
   depends_on:
     - image-build-cpu
   source_file_dependencies:
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/inputs/
-    - vllm/lora/
-    - vllm/multimodal/
-    - vllm/outputs.py
-    - vllm/platforms/
-    - vllm/pooling_params.py
-    - vllm/profiler/
-    - vllm/sampling_params.py
-    - vllm/tokenizers/
-    - vllm/transformers_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/v1
   device: cpu-small
   commands:
@@ -129,35 +89,12 @@ steps:
     - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'cpu_test' v1/metrics
 
-- label: Extract Hidden States Integration
-  key: extract-hidden-states-integration
-  timeout_in_minutes: 20
-  device: h200_18gb
-  source_file_dependencies:
-    - vllm/v1/spec_decode/extract_hidden_states.py
-    - vllm/model_executor/models/extract_hidden_states.py
-    - vllm/transformers_utils/configs/extract_hidden_states.py
-    - tests/v1/kv_connector/extract_hidden_states_integration
-  commands:
-    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-    - pytest -v -s v1/kv_connector/extract_hidden_states_integration
-
 - label: Regression
   key: regression
   timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/inputs/
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/sampling_params.py
-  - vllm/transformers_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/test_regression
   commands:
   - pip install modelscope
@@ -190,8 +127,8 @@ steps:
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # for features demo
     - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
     # https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
@@ -201,18 +138,7 @@ steps:
   timeout_in_minutes: 20
   num_devices: 2
   source_file_dependencies:
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/inputs/
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/sampling_params.py
-  - vllm/tracing/
-  - vllm/transformers_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/v1/tracing
   commands:
   - "pip install \
@@ -225,9 +151,7 @@ steps:
 - label: Python-only Installation
   key: python-only-installation
   depends_on: ~
-  optional: true
   timeout_in_minutes: 20
-  device: h200_18gb
   source_file_dependencies:
   - tests/standalone_tests/python_only_compile.sh
   - setup.py
@@ -238,19 +162,7 @@ steps:
   key: async-engine-inputs-utils-worker
   timeout_in_minutes: 50
   source_file_dependencies:
-  - vllm/assets/
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/inputs/
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/sampling_params.py
-  - vllm/tokenizers/
-  - vllm/transformers_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/detokenizer
   - tests/multimodal
   - tests/utils_
@@ -265,30 +177,7 @@ steps:
   - image-build-cpu
   timeout_in_minutes: 30
   source_file_dependencies:
-  - vllm/assets/
-  - vllm/config/
-  - vllm/engine/arg_utils.py
-  - vllm/entrypoints/chat_utils.py
-  - vllm/entrypoints/mcp/
-  - vllm/entrypoints/openai/chat_completion/protocol.py
-  - vllm/entrypoints/openai/engine/protocol.py
-  - vllm/envs.py
-  - vllm/exceptions.py
-  - vllm/inputs/
-  - vllm/model_executor/layers/quantization/quark/
-  - vllm/multimodal/
-  - vllm/outputs.py
-  - vllm/platforms/
-  - vllm/pooling_params.py
-  - vllm/ray/
-  - vllm/reasoning/
-  - vllm/renderers/
-  - vllm/sampling_params.py
-  - vllm/tokenizers/
-  - vllm/tool_parsers/
-  - vllm/transformers_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/test_inputs.py
   - tests/test_outputs.py
   - tests/test_pooling_params.py
@@ -299,7 +188,6 @@ steps:
   - tests/tokenizers_
   - tests/reasoning
   - tests/tool_parsers
-  - tests/parser
   - tests/transformers_utils
   - tests/config
   device: cpu-small
@@ -312,26 +200,11 @@ steps:
   - pytest -v -s -m 'cpu_test' multimodal
   - pytest -v -s renderers
   - pytest -v -s tokenizers_
-  - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
+  - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py --ignore=reasoning/test_gemma4_reasoning_parser.py
   - pytest -v -s tool_parsers
-  - pytest -v -s parser
   - pytest -v -s transformers_utils
   - pytest -v -s config
 
-- label: Batch Invariance (A100)
-  key: batch-invariance-a100
-  timeout_in_minutes: 30
-  device: a100
-  source_file_dependencies:
-    - vllm/v1/attention
-    - vllm/model_executor/layers
-    - tests/v1/determinism/
-  commands:
-    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-    - pip install pytest-timeout pytest-forked
-    - pytest -v -s v1/determinism/test_batch_invariance.py
-    - 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]
-
 - label: Batch Invariance (H100)
   key: batch-invariance-h100
   timeout_in_minutes: 30
@@ -351,7 +224,7 @@ steps:
 - label: Batch Invariance (B200)
   key: batch-invariance-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
     - vllm/v1/attention
     - vllm/model_executor/layers

.buildkite/test_areas/model_runner_v2.yaml

@@ -37,7 +37,7 @@ steps:
     - examples/generate/multimodal/
     - examples/features/
     - examples/pooling/embed/vision_embedding_offline.py
-    - examples/features/tensorize_vllm_model.py
+    - examples/others/tensorize_vllm_model.py
   commands:
     - set -x
     - export VLLM_USE_V2_MODEL_RUNNER=1
@@ -55,8 +55,8 @@ steps:
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # for features demo
     - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
     # https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
@@ -106,12 +106,13 @@ steps:
   - vllm/v1/worker/gpu/
   - vllm/v1/worker/gpu_worker.py
   - tests/v1/spec_decode/test_max_len.py
-  - tests/v1/spec_decode/test_rejection_sampler_utils.py
+  - tests/v1/spec_decode/test_probabilistic_rejection_sampler_utils.py
+  - tests/v1/spec_decode/test_synthetic_rejection_sampler_utils.py
   - tests/v1/e2e/spec_decode/test_spec_decode.py
   commands:
   - set -x
   - export VLLM_USE_V2_MODEL_RUNNER=1
   - pytest -v -s v1/spec_decode/test_max_len.py -k "eagle or mtp"
-  - pytest -v -s v1/spec_decode/test_rejection_sampler_utils.py
+  - pytest -v -s v1/spec_decode/test_probabilistic_rejection_sampler_utils.py
   - pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py
   - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle or mtp"

.buildkite/test_areas/models_basic.yaml

@@ -5,7 +5,6 @@ steps:
 - label: Basic Models Tests (Initialization)
   key: basic-models-tests-initialization
   timeout_in_minutes: 45
-  device: h200_18gb
   torch_nightly: true
   source_file_dependencies:
   - vllm/
@@ -43,6 +42,12 @@ steps:
   - tests/models/test_registry.py
   commands:
     - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
+
 
 - label: Basic Models Test (Other CPU) # 5min
   key: basic-models-test-other-cpu

.buildkite/test_areas/models_language.yaml

@@ -5,7 +5,6 @@ steps:
 - label: Language Models Tests (Standard)
   key: language-models-tests-standard
   timeout_in_minutes: 25
-  device: h200_18gb
   source_file_dependencies:
   - vllm/
   - tests/models/language
@@ -49,14 +48,6 @@ steps:
   parallelism: 2
   mirror:
     torch_nightly: {}
-    amd:
-      device: mi300_1
-      depends_on:
-      - image-build-amd
-      commands:
-      - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr'
-      - 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
 
 - label: Language Models Test (Extended Generation) # 80min
   key: language-models-test-extended-generation
@@ -71,6 +62,15 @@ steps:
     - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
     - 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 '(not core_model) and (not hybrid_model)'
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
+      commands:
+      - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr'
+      - 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 '(not core_model) and (not hybrid_model)'
 
 - label: Language Models Test (PPL)
   key: language-models-test-ppl
@@ -94,8 +94,7 @@ steps:
     - pytest -v -s models/language/pooling -m 'not core_model'
   mirror:
     amd:
-      device: mi300_1
-      timeout_in_minutes: 100
+      device: mi325_1
       depends_on:
       - image-build-amd
 

.buildkite/test_areas/models_multimodal.yaml

@@ -15,7 +15,7 @@ steps:
     - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
@@ -33,7 +33,7 @@ steps:
     - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
@@ -49,7 +49,7 @@ steps:
     - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
@@ -64,6 +64,11 @@ steps:
     - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing
     - pytest 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  # Otherwise, mp_method="spawn" doesn't work
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
 
 - label: Multi-Modal Processor (CPU)
   key: multi-modal-processor-cpu
@@ -115,7 +120,7 @@ steps:
     - pytest -v -s models/multimodal/test_mapping.py
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 

.buildkite/test_areas/pytorch.yaml

@@ -6,30 +6,7 @@ steps:
   key: pytorch-compilation-unit-tests
   timeout_in_minutes: 10
   source_file_dependencies:
-    - vllm/__init__.py
-    - vllm/_aiter_ops.py
-    - vllm/_custom_ops.py
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/env_override.py
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/ir/
-    - vllm/kernels/
-    - vllm/logger.py
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/plugins/
-    - vllm/sampling_params.py
-    - vllm/sequence.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/compile
   commands:
   # Run unit tests defined directly under compile/,
@@ -47,30 +24,7 @@ steps:
   device: h100
   num_devices: 1
   source_file_dependencies:
-    - vllm/__init__.py
-    - vllm/_aiter_ops.py
-    - vllm/_custom_ops.py
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/env_override.py
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/ir/
-    - vllm/kernels/
-    - vllm/logger.py
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/plugins/
-    - vllm/sampling_params.py
-    - vllm/sequence.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/compile/h100/
   commands:
   - "find compile/h100/ -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
@@ -79,30 +33,7 @@ steps:
   key: pytorch-compilation-passes-unit-tests
   timeout_in_minutes: 20
   source_file_dependencies:
-    - vllm/__init__.py
-    - vllm/_aiter_ops.py
-    - vllm/_custom_ops.py
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/env_override.py
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/ir/
-    - vllm/kernels/
-    - vllm/logger.py
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/plugins/
-    - vllm/sampling_params.py
-    - vllm/sequence.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/compile/passes
   commands:
   - pytest -s -v compile/passes --ignore compile/passes/distributed
@@ -111,30 +42,7 @@ steps:
   key: pytorch-fullgraph-smoke-test
   timeout_in_minutes: 35
   source_file_dependencies:
-  - vllm/__init__.py
-  - vllm/_aiter_ops.py
-  - vllm/_custom_ops.py
-  - vllm/compilation/
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/env_override.py
-  - vllm/envs.py
-  - vllm/forward_context.py
-  - vllm/inputs/
-  - vllm/ir/
-  - vllm/kernels/
-  - vllm/logger.py
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/plugins/
-  - vllm/sampling_params.py
-  - vllm/sequence.py
-  - vllm/transformers_utils/
-  - vllm/triton_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/compile
   commands:
   # Run smoke tests under fullgraph directory, except test_full_graph.py
@@ -148,30 +56,7 @@ steps:
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
-  - vllm/__init__.py
-  - vllm/_aiter_ops.py
-  - vllm/_custom_ops.py
-  - vllm/compilation/
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/env_override.py
-  - vllm/envs.py
-  - vllm/forward_context.py
-  - vllm/inputs/
-  - vllm/ir/
-  - vllm/kernels/
-  - vllm/logger.py
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/plugins/
-  - vllm/sampling_params.py
-  - vllm/sequence.py
-  - vllm/transformers_utils/
-  - vllm/triton_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/compile
   commands:
     # fp8 kv scales not supported on sm89, tested on Blackwell instead

.buildkite/test_areas/quantization.yaml

@@ -25,7 +25,7 @@ steps:
   key: quantized-moe-test-b200
   timeout_in_minutes: 60
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
   - tests/quantization/test_blackwell_moe.py
   - vllm/model_executor/models/deepseek_v2.py

.buildkite/test_areas/samplers.yaml

@@ -17,7 +17,7 @@ steps:
     - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
   mirror:
     amd:
-      device: mi250_1
+      device: mi325_1
       depends_on:
       - image-build-amd
       commands:

.buildkite/test_areas/spec_decode.yaml

@@ -16,7 +16,7 @@ steps:
 - label: Spec Decode Eagle Nightly B200
   key: spec-decode-eagle-nightly-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -40,7 +40,7 @@ steps:
 - label: Spec Decode Speculators + MTP Nightly B200
   key: spec-decode-speculators-mtp-nightly-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -75,7 +75,7 @@ steps:
 - label: Spec Decode Draft Model Nightly B200
   key: spec-decode-draft-model-nightly-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -84,23 +84,23 @@ steps:
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
 
-- label: Speculators Correctness
-  key: speculators-correctness
-  timeout_in_minutes: 60
+- label: DFlash Speculators Correctness
+  key: dflash-speculators-correctness
+  timeout_in_minutes: 30
   device: h100
   optional: true
   num_devices: 1
   source_file_dependencies:
     - vllm/v1/spec_decode/
     - vllm/model_executor/models/qwen3_dflash.py
-    - tests/v1/spec_decode/test_speculators_correctness.py
+    - tests/v1/spec_decode/test_speculators_dflash.py
   commands:
     - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
-    - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
+    - pytest -v -s v1/spec_decode/test_speculators_dflash.py -m slow_test
 
 - label: Spec Decode MTP hybrid (B200)
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
     - vllm/v1/spec_decode/

.github/CODEOWNERS

@@ -6,8 +6,8 @@
 /vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery @xuechendi
 /vllm/lora @jeejeelee
 /vllm/model_executor/layers/attention @LucasWilkinson @MatthewBonanni
-/vllm/model_executor/layers/fused_moe @mgoin @pavanimajety @zyongye
-/vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety @zyongye
+/vllm/model_executor/layers/fused_moe @mgoin @pavanimajety
+/vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety
 /vllm/model_executor/layers/mamba @tdoublep @tomeras91
 /vllm/model_executor/layers/mamba/gdn_linear_attn.py @tdoublep @ZJY0516 @vadiklyutiy
 /vllm/model_executor/layers/rotary_embedding.py @vadiklyutiy
@@ -18,8 +18,7 @@
 /vllm/kernels/helion @ProExpertProg @zou3519
 /vllm/multimodal @DarkLight1337 @ywang96 @NickLucche @tjtanaa
 /vllm/vllm_flash_attn @LucasWilkinson @MatthewBonanni
-/CMakeLists.txt @tlrmchlsmth @LucasWilkinson @Harry-Chen
-/cmake @tlrmchlsmth @LucasWilkinson @Harry-Chen
+CMakeLists.txt @tlrmchlsmth @LucasWilkinson
 
 # Any change to the VllmConfig changes can have a large user-facing impact,
 # so spam a lot of people
@@ -31,8 +30,8 @@
 /vllm/entrypoints/cli @hmellor @mgoin @DarkLight1337 @russellb
 /vllm/entrypoints/mcp @heheda12345
 /vllm/entrypoints/openai @aarnphm @chaunceyjiang @DarkLight1337 @russellb
-/vllm/entrypoints/speech_to_text/realtime @njhill
-/vllm/entrypoints/speech_to_text @NickLucche
+/vllm/entrypoints/openai/realtime @njhill
+/vllm/entrypoints/openai/speech_to_text @NickLucche
 /vllm/entrypoints/pooling @noooop
 /vllm/entrypoints/sagemaker @DarkLight1337
 /vllm/entrypoints/serve @njhill
@@ -71,10 +70,6 @@
 /vllm/v1/worker/gpu @WoosukKwon @njhill
 /vllm/v1/worker/gpu/kv_connector.py @orozery
 
-# CI & building
-/.buildkite @Harry-Chen
-/docker/Dockerfile @Harry-Chen
-
 # Test ownership
 /.buildkite/lm-eval-harness @mgoin 
 /tests/distributed/test_multi_node_assignment.py @youkaichao
@@ -82,11 +77,11 @@
 /tests/distributed/test_same_node.py @youkaichao
 /tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche
 /tests/evals @mgoin @vadiklyutiy
-/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye
+/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256
 /tests/kernels/ir @ProExpertProg @tjtanaa
 /tests/models @DarkLight1337 @ywang96
 /tests/multimodal @DarkLight1337 @ywang96 @NickLucche
-/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye
+/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety
 /tests/test_inputs.py @DarkLight1337 @ywang96
 /tests/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
 /tests/v1/structured_output @mgoin @russellb @aarnphm
@@ -152,12 +147,6 @@ mkdocs.yaml @hmellor
 # MTP-specific files
 /vllm/model_executor/models/deepseek_mtp.py @luccafong
 
-# DeepseekV4-specific files
-/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
 /vllm/model_executor/models/mistral*.py @patrickvonplaten
 /vllm/model_executor/models/mixtral*.py @patrickvonplaten

.github/mergify.yml

@@ -477,7 +477,9 @@ pull_request_rules:
   conditions:
     - label != stale
     - or:
-      - files~=^examples/disaggregated/
+      - files~=^examples/online_serving/disaggregated[^/]*/.*
+      - files~=^examples/offline_inference/disaggregated[^/]*/.*
+      - files~=^examples/others/lmcache/
       - files~=^tests/v1/kv_connector/
       - files~=^vllm/distributed/kv_transfer/
       - title~=(?i)\bP/?D\b

.github/workflows/scripts/build.sh

@@ -9,9 +9,6 @@ PATH=${cuda_home}/bin:$PATH
 LD_LIBRARY_PATH=${cuda_home}/lib64:$LD_LIBRARY_PATH
 
 # Install requirements
-if [ "$(echo $2 | cut -d. -f1)" = "12" ]; then
-    sed -i 's/^nvidia-cutlass-dsl\[cu13\]>=/nvidia-cutlass-dsl>=/' requirements/cuda.txt
-fi
 $python_executable -m pip install -r requirements/build/cuda.txt -r requirements/cuda.txt
 
 # Limit the number of parallel jobs to avoid OOM

.pre-commit-config.yaml

@@ -131,16 +131,6 @@ repos:
         --python-version, "3.12",
       ]
       files: ^requirements/(common|xpu|test/xpu)\.(in|txt)$
-    - id: pip-compile
-      alias: pip-compile-docs
-      name: pip-compile-docs
-      args: [
-        requirements/docs.in,
-        -o, requirements/docs.txt,
-        --python-platform, x86_64-manylinux_2_28,
-        --python-version, "3.12",
-      ]
-      files: ^requirements/docs\.(in|txt)$
 - repo: local
   hooks:
   - id: format-torch-nightly-test

.readthedocs.yaml

@@ -9,7 +9,7 @@ build:
     python: "3.12"
   jobs:
     post_checkout:
-      - bash docs/pre_run_check.sh
+      # - bash docs/maybe_skip_pr_build.sh
       - git fetch origin main --unshallow --no-tags --filter=blob:none || true
     pre_create_environment:
       - pip install uv

CMakeLists.txt

@@ -13,12 +13,8 @@ cmake_minimum_required(VERSION 3.26)
 # cmake --install . --component _C
 project(vllm_extensions LANGUAGES CXX)
 
-set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
-set(CMAKE_CUDA_STANDARD 20)
-set(CMAKE_CUDA_STANDARD_REQUIRED ON)
-set(CMAKE_HIP_STANDARD 20)
-set(CMAKE_HIP_STANDARD_REQUIRED ON)
 
 
 # CUDA by default, can be overridden by using -DVLLM_TARGET_DEVICE=... (used by setup.py)
@@ -109,24 +105,6 @@ else()
   set(CUDA_SUPPORTED_ARCHS "7.0;7.5;8.0;8.6;8.7;8.9;9.0")
 endif()
 
-#
-# spinloop extension (pure CXX; must stay above the non-CUDA device branch so
-# CPU builds define the target before the early return)
-#
-set(VLLM_SPINLOOP_EXT_SRC "csrc/spinloop.cpp")
-set(SPINLOOP_COMPILE_FLAGS "")
-if(CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|amd64")
-  list(APPEND SPINLOOP_COMPILE_FLAGS "-mmwaitx")
-endif()
-define_extension_target(
-  spinloop
-  DESTINATION vllm
-  LANGUAGE CXX
-  SOURCES ${VLLM_SPINLOOP_EXT_SRC}
-  COMPILE_FLAGS ${SPINLOOP_COMPILE_FLAGS}
-  USE_SABI 3.11
-  WITH_SOABI)
-
 #
 # Forward the non-CUDA device extensions to external CMake scripts.
 #
@@ -311,6 +289,7 @@ 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"
@@ -328,12 +307,12 @@ set(VLLM_EXT_SRC
   "csrc/quantization/activation_kernels.cu"
   "csrc/cuda_utils_kernels.cu"
   "csrc/custom_all_reduce.cu"
-  "csrc/torch_bindings.cpp"
-  "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
+  "csrc/torch_bindings.cpp")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   list(APPEND VLLM_EXT_SRC
-    "csrc/minimax_reduce_rms_kernel.cu")
+    "csrc/minimax_reduce_rms_kernel.cu"
+    "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
 
   SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")
 
@@ -368,6 +347,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   FetchContent_MakeAvailable(cutlass)
 
   list(APPEND VLLM_EXT_SRC
+    "csrc/quantization/awq/gemm_kernels.cu"
     "csrc/cutlass_extensions/common.cpp")
 
   set_gencode_flags_for_srcs(
@@ -499,6 +479,46 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
                    " in CUDA target architectures")
   endif()
 
+  # Only build AllSpark kernels if we are building for at least some compatible archs.
+  cuda_archs_loose_intersection(ALLSPARK_ARCHS "8.0;8.6;8.7;8.9" "${CUDA_ARCHS}")
+  if (ALLSPARK_ARCHS)
+    set(ALLSPARK_SRCS
+       "csrc/quantization/gptq_allspark/allspark_repack.cu"
+       "csrc/quantization/gptq_allspark/allspark_qgemm_w8a16.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${ALLSPARK_SRCS}"
+      CUDA_ARCHS "${ALLSPARK_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${ALLSPARK_SRCS}")
+    message(STATUS "Building AllSpark kernels for archs: ${ALLSPARK_ARCHS}")
+  else()
+    message(STATUS "Not building AllSpark kernels as no compatible archs found"
+                   " in CUDA target architectures")
+  endif()
+
+  # CUTLASS MLA Archs and flags
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
+  endif()
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
+    set(SRCS
+      "csrc/attention/mla/sm100_cutlass_mla_kernel.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${SRCS}"
+      CUDA_ARCHS "${MLA_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${SRCS}")
+    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
+    # Add MLA-specific include directories only to MLA source files
+    set_source_files_properties(${SRCS}
+      PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
+    message(STATUS "Building CUTLASS MLA for archs: ${MLA_ARCHS}")
+  else()
+    message(STATUS "Not building CUTLASS MLA as no compatible archs were found.")
+    # clear MLA_ARCHS
+    set(MLA_ARCHS)
+  endif()
+
   # Expert-specialization MXFP8 blockscaled grouped kernels (SM100+).
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
     cuda_archs_loose_intersection(ES_MXFP8_GROUPED_MM_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
@@ -526,6 +546,24 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
   endif()
 
+  # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
+  endif()
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
+    set(DSV3_FUSED_A_GEMM_SRC "csrc/dsv3_fused_a_gemm.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${DSV3_FUSED_A_GEMM_SRC}"
+      CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
+    list(APPEND VLLM_EXT_SRC ${DSV3_FUSED_A_GEMM_SRC})
+    message(STATUS "Building dsv3_fused_a_gemm for archs: ${DSV3_FUSED_A_GEMM_ARCHS}")
+  else()
+    message(STATUS "Not building dsv3_fused_a_gemm as no compatible archs found "
+                   "in CUDA target architectures.")
+  endif()
+
   #
   # Machete kernels
 
@@ -597,6 +635,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   endif()
 
 
+  # Hadacore kernels
+  cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
+  if(HADACORE_ARCHS)
+    set(SRCS "csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${SRCS}"
+      CUDA_ARCHS "${HADACORE_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${SRCS}")
+    message(STATUS "Building hadacore")
+  endif()
 
 # if CUDA endif
 endif()
@@ -646,8 +694,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     list(APPEND VLLM_STABLE_EXT_SRC
       "csrc/libtorch_stable/permute_cols.cu"
       "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
-      "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
-      "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu")
+      "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu")
   endif()
 
   if(VLLM_GPU_LANG STREQUAL "CUDA")
@@ -656,40 +703,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
       CUDA_ARCHS "${CUDA_ARCHS}")
   endif()
 
-  # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
-  endif()
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
-    set(SRCS "csrc/libtorch_stable/dsv3_fused_a_gemm.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building dsv3_fused_a_gemm for archs: ${DSV3_FUSED_A_GEMM_ARCHS}")
-  else()
-    message(STATUS "Not building dsv3_fused_a_gemm as no compatible archs found "
-                   "in CUDA target architectures.")
-  endif()
-
-  # Only build AllSpark kernels if we are building for at least some compatible archs.
-  cuda_archs_loose_intersection(ALLSPARK_ARCHS "8.0;8.6;8.7;8.9" "${CUDA_ARCHS}")
-  if (ALLSPARK_ARCHS)
-    set(SRCS
-       "csrc/libtorch_stable/quantization/gptq_allspark/allspark_repack.cu"
-       "csrc/libtorch_stable/quantization/gptq_allspark/allspark_qgemm_w8a16.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${ALLSPARK_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building AllSpark kernels for archs: ${ALLSPARK_ARCHS}")
-  else()
-    message(STATUS "Not building AllSpark kernels as no compatible archs found"
-                   " in CUDA target architectures")
-  endif()
-
   #
   # CUTLASS scaled_mm kernels (moved from _C to _C_stable_libtorch)
   #
@@ -999,41 +1012,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
   endif()
 
-  # CUTLASS MLA Archs and flags
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
-  endif()
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
-    set(SRCS
-      "csrc/libtorch_stable/attention/mla/sm100_cutlass_mla_kernel.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${MLA_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
-    # Add MLA-specific include directories only to MLA source files
-    set_source_files_properties(${SRCS}
-      PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
-    message(STATUS "Building CUTLASS MLA for archs: ${MLA_ARCHS}")
-  else()
-    message(STATUS "Not building CUTLASS MLA as no compatible archs were found.")
-    # clear MLA_ARCHS
-    set(MLA_ARCHS)
-  endif()
-
-  # Hadacore kernels
-  cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
-  if(HADACORE_ARCHS)
-    set(SRCS "csrc/libtorch_stable/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${HADACORE_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building hadacore")
-  endif()
-
   message(STATUS "Enabling C_stable extension.")
   define_extension_target(
     _C_stable_libtorch
@@ -1069,13 +1047,14 @@ endif()
 set(VLLM_MOE_EXT_SRC
   "csrc/moe/torch_bindings.cpp"
   "csrc/moe/moe_align_sum_kernels.cu"
-  "csrc/moe/topk_softmax_kernels.cu"
-  "csrc/moe/topk_softplus_sqrt_kernels.cu")
+  "csrc/moe/topk_softmax_kernels.cu")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   list(APPEND VLLM_MOE_EXT_SRC
     "csrc/moe/moe_wna16.cu"
-    "csrc/moe/grouped_topk_kernels.cu")
+    "csrc/moe/grouped_topk_kernels.cu"
+    "csrc/moe/router_gemm.cu"
+    "csrc/moe/topk_softplus_sqrt_kernels.cu")
 endif()
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
@@ -1221,16 +1200,6 @@ 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)")

benchmarks/attention_benchmarks/configs/mla_decode.yaml

@@ -53,7 +53,6 @@ backends:
   - FLASHINFER_MLA
   - FLASH_ATTN_MLA  # Hopper only
   - FLASHMLA        # Hopper only
-  - TOKENSPEED_MLA  # Blackwell + R1 dims + FP8 KV (use --kv-cache-dtype fp8)
 
 device: "cuda:0"
 repeats: 100

benchmarks/attention_benchmarks/configs/mla_prefill.yaml

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

benchmarks/attention_benchmarks/mla_runner.py

@@ -29,7 +29,6 @@ from vllm.config import (
     VllmConfig,
     set_current_vllm_config,
 )
-from vllm.v1.attention.backends.mla.prefill.registry import MLAPrefillBackendEnum
 
 # ============================================================================
 # VllmConfig Creation
@@ -80,8 +79,8 @@ def create_minimal_vllm_config(
         index_topk: Optional topk value for sparse MLA backends. If provided,
                     the config will include index_topk for sparse attention.
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4", "flashinfer",
-                        "trtllm"). Configures the attention config to force
-                        the specified prefill backend.
+                        "cudnn", "trtllm"). Configures the attention config to
+                        force the specified prefill backend.
 
     Returns:
         VllmConfig for benchmarking
@@ -180,13 +179,19 @@ def create_minimal_vllm_config(
 
     if prefill_backend is not None:
         prefill_cfg = get_prefill_backend_config(prefill_backend)
-        vllm_config.attention_config.mla_prefill_backend = prefill_cfg[
-            "mla_prefill_backend"
-        ]
         if prefill_cfg["flash_attn_version"] is not None:
             vllm_config.attention_config.flash_attn_version = prefill_cfg[
                 "flash_attn_version"
             ]
+        vllm_config.attention_config.disable_flashinfer_prefill = prefill_cfg[
+            "disable_flashinfer_prefill"
+        ]
+        vllm_config.attention_config.use_cudnn_prefill = prefill_cfg[
+            "use_cudnn_prefill"
+        ]
+        vllm_config.attention_config.use_trtllm_ragged_deepseek_prefill = prefill_cfg[
+            "use_trtllm_ragged_deepseek_prefill"
+        ]
 
     return vllm_config
 
@@ -201,27 +206,39 @@ def create_minimal_vllm_config(
 _PREFILL_BACKEND_CONFIG: dict[str, dict] = {
     "fa2": {
         "flash_attn_version": 2,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "fa3": {
         "flash_attn_version": 3,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "fa4": {
         "flash_attn_version": 4,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "flashinfer": {
         "flash_attn_version": None,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASHINFER,
+        "disable_flashinfer_prefill": False,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
+    },
+    "cudnn": {
+        "flash_attn_version": None,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": True,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "trtllm": {
         "flash_attn_version": None,
-        "mla_prefill_backend": MLAPrefillBackendEnum.TRTLLM_RAGGED,
-    },
-    "tokenspeed": {
-        "flash_attn_version": None,
-        "mla_prefill_backend": MLAPrefillBackendEnum.TOKENSPEED_MLA,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": True,
     },
 }
 
@@ -608,21 +625,6 @@ def _create_backend_impl(
     # Create mock layer
     layer = MockLayer(device, impl=impl, kv_cache_spec=kv_cache_spec)
 
-    # Attach a prefill backend (MLAAttention does this in __init__; the metadata
-    # builder reads layer.prefill_backend from static_forward_context).
-    from vllm.v1.attention.backends.mla.prefill import get_mla_prefill_backend
-
-    prefill_backend_cls = get_mla_prefill_backend(vllm_config)
-    layer.prefill_backend = prefill_backend_cls(
-        num_heads=mla_dims["num_q_heads"],
-        scale=(mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"]) ** -0.5,
-        kv_lora_rank=mla_dims["kv_lora_rank"],
-        qk_nope_head_dim=mla_dims["qk_nope_head_dim"],
-        qk_rope_head_dim=mla_dims["qk_rope_head_dim"],
-        v_head_dim=mla_dims["v_head_dim"],
-        vllm_config=vllm_config,
-    )
-
     # Create builder instance if needed
     builder_instance = None
     if builder_class:
@@ -959,6 +961,19 @@ def _run_mla_benchmark_batched(
     results = []
 
     with set_current_vllm_config(vllm_config):
+        # Clear cached prefill backend detection functions so they re-evaluate
+        # with the current VllmConfig. These are @functools.cache decorated and
+        # would otherwise return stale results from a previous backend's config.
+        from vllm.model_executor.layers.attention.mla_attention import (
+            use_cudnn_prefill,
+            use_flashinfer_prefill,
+            use_trtllm_ragged_deepseek_prefill,
+        )
+
+        use_flashinfer_prefill.cache_clear()
+        use_cudnn_prefill.cache_clear()
+        use_trtllm_ragged_deepseek_prefill.cache_clear()
+
         # Create backend impl, layer, builder, and indexer (reused across benchmarks)
         impl, layer, builder_instance, indexer = _create_backend_impl(
             backend_cfg,
@@ -970,36 +985,38 @@ def _run_mla_benchmark_batched(
             kv_cache_dtype=kv_cache_dtype,
         )
 
-        # Verify the actual prefill backend matches what was requested. The
-        # selector + impl construction already raise on misuse; here we just
-        # check the resolved class against the requested name as a sanity guard.
+        # Verify the actual prefill backend matches what was requested
         if prefill_backend is not None:
-            expected_class = {
-                "fa2": "FlashAttnPrefillBackend",
-                "fa3": "FlashAttnPrefillBackend",
-                "fa4": "FlashAttnPrefillBackend",
-                "flashinfer": "FlashInferPrefillBackend",
-                "trtllm": "TrtllmRaggedPrefillBackend",
-                "tokenspeed": "TokenspeedMLAPrefillBackend",
-            }.get(prefill_backend)
-            actual_class = type(getattr(layer, "prefill_backend", None)).__name__
-            if expected_class and actual_class != expected_class:
-                raise RuntimeError(
-                    f"Prefill backend '{prefill_backend}' requested "
-                    f"{expected_class}, got {actual_class}. Check "
-                    f"attention_config plumbing or installed deps."
-                )
-            if prefill_backend in {"fa2", "fa3", "fa4"}:
-                fa_version = int(prefill_backend[2:])
-                actual_fa_version = getattr(
-                    layer.prefill_backend, "vllm_flash_attn_version", None
-                )
+            prefill_cfg = get_prefill_backend_config(prefill_backend)
+            fa_version = prefill_cfg["flash_attn_version"]
+
+            if fa_version is not None:
+                # FA backend: verify the impl's FA version
+                actual_fa_version = getattr(impl, "vllm_flash_attn_version", None)
                 if actual_fa_version != fa_version:
                     raise RuntimeError(
                         f"Prefill backend '{prefill_backend}' requested FA "
-                        f"version {fa_version}, got "
-                        f"{actual_fa_version} on {actual_class}."
+                        f"version {fa_version}, but the impl is using FA "
+                        f"version {actual_fa_version}. Check "
+                        f"vllm/v1/attention/backends/fa_utils.py."
                     )
+            else:
+                # Non-FA backend: verify the builder picked the right path
+                expected_flags = {
+                    "flashinfer": "_use_fi_prefill",
+                    "cudnn": "_use_cudnn_prefill",
+                    "trtllm": "_use_trtllm_ragged_prefill",
+                }
+                flag_name = expected_flags.get(prefill_backend)
+                if flag_name and not getattr(builder_instance, flag_name, False):
+                    raise RuntimeError(
+                        f"Prefill backend '{prefill_backend}' was requested "
+                        f"but the metadata builder did not enable it. This "
+                        f"usually means a dependency is missing (e.g., "
+                        f"flashinfer not installed) or the platform doesn't "
+                        f"support it."
+                    )
+
         # Run each benchmark with the shared impl
         for config, threshold, num_splits in configs_with_params:
             # Set threshold for this benchmark (FlashAttn/FlashMLA only)

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_moe.py

@@ -27,10 +27,10 @@ from vllm.model_executor.layers.fused_moe.config import (
     RoutingMethodType,
     _get_config_dtype_str,
 )
-from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
+from vllm.model_executor.layers.fused_moe.fused_moe import *
+from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
     TritonOrDeepGemmExperts,
 )
-from vllm.model_executor.layers.fused_moe.fused_moe import *
 from vllm.transformers_utils.config import get_config
 from vllm.triton_utils import triton
 from vllm.utils.argparse_utils import FlexibleArgumentParser

benchmarks/kernels/benchmark_norm_router_gemm.py

@@ -1,183 +0,0 @@
-# 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/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/multi_turn/benchmark_serving_multi_turn.py

@@ -1473,12 +1473,6 @@ async def main() -> None:
         "(for example: --warmup-percentages=0%%,50%%)",
     )
 
-    parser.add_argument(
-        "--trust-remote-code",
-        action="store_true",
-        help="Trust remote code when loading the tokenizer.",
-    )
-
     args = parser.parse_args()
 
     logger.info(args)
@@ -1521,9 +1515,7 @@ async def main() -> None:
     np.random.seed(args.seed)
 
     logger.info("Loading tokenizer")
-    tokenizer = AutoTokenizer.from_pretrained(
-        args.model, trust_remote_code=args.trust_remote_code
-    )
+    tokenizer = AutoTokenizer.from_pretrained(args.model)
 
     await get_server_info(args.url)
 

cmake/cpu_extension.cmake

@@ -1,7 +1,7 @@
 include(FetchContent)
 
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
-set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_EXTENSIONS ON)
 set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
 
@@ -32,23 +32,18 @@ else()
         "-DVLLM_CPU_EXTENSION")
 
     # locate PyTorch's libgomp (e.g. site-packages/torch.libs/libgomp-947d5fa1.so.1.0.0)
-    # and create a local shim dir with it. When PyTorch is built from source or packaged
-    # by a distro (common on RISC-V, s390x, Fedora/RHEL aarch64), no vendored libgomp
-    # exists and the shim dir is empty; fall back to the system libgomp in that case.
+    # and create a local shim dir with it
     vllm_prepare_torch_gomp_shim(VLLM_TORCH_GOMP_SHIM_DIR)
 
-    if(VLLM_TORCH_GOMP_SHIM_DIR)
-        find_library(OPEN_MP
-            NAMES gomp
-            PATHS "${VLLM_TORCH_GOMP_SHIM_DIR}"
-            NO_DEFAULT_PATH
-            REQUIRED
-        )
-        # Use the same libgomp as PyTorch at runtime
+    find_library(OPEN_MP
+        NAMES gomp
+        PATHS ${VLLM_TORCH_GOMP_SHIM_DIR}
+        NO_DEFAULT_PATH
+        REQUIRED
+    )
+    # Set LD_LIBRARY_PATH to include the shim dir at build time to use the same libgomp as PyTorch
+    if (OPEN_MP)
         set(ENV{LD_LIBRARY_PATH} "${VLLM_TORCH_GOMP_SHIM_DIR}:$ENV{LD_LIBRARY_PATH}")
-    else()
-        # Fall back to system / toolchain libgomp
-        find_library(OPEN_MP NAMES gomp REQUIRED)
     endif()
 endif()
 
@@ -195,12 +190,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)
@@ -328,6 +321,14 @@ if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND
     set(ONEDNN_VERBOSE "ON")
     set(CMAKE_POLICY_DEFAULT_CMP0077 NEW)
 
+    # TODO: Refactor this
+    if (ENABLE_X86_ISA)
+        # Note: only enable oneDNN for AVX512
+        list(APPEND DNNL_COMPILE_FLAGS ${CXX_COMPILE_FLAGS_AVX512})
+    else()
+        list(APPEND DNNL_COMPILE_FLAGS ${CXX_COMPILE_FLAGS})
+    endif()
+
     set(VLLM_BUILD_TYPE ${CMAKE_BUILD_TYPE})
     set(CMAKE_BUILD_TYPE "Release") # remove oneDNN debug symbols to reduce size
     FetchContent_MakeAvailable(oneDNN)
@@ -340,14 +341,8 @@ if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND
         PRIVATE ${oneDNN_SOURCE_DIR}/src
     )
     target_link_libraries(dnnl_ext dnnl torch)
-    if (ENABLE_X86_ISA)
-        target_compile_options(dnnl_ext PRIVATE ${CXX_COMPILE_FLAGS_AVX2} -fPIC)
-    else()
-        target_compile_options(dnnl_ext PRIVATE ${CXX_COMPILE_FLAGS} -fPIC)
-    endif()
+    target_compile_options(dnnl_ext PRIVATE ${DNNL_COMPILE_FLAGS} -fPIC)
     list(APPEND LIBS dnnl_ext)
-
-
     set(USE_ONEDNN ON)
 else()
     set(USE_ONEDNN OFF)
@@ -411,15 +406,12 @@ endif()
 
 if (ENABLE_X86_ISA)
     set(VLLM_EXT_SRC_SGL
-        "csrc/cpu/sgl-kernels/fla.cpp"
-        "csrc/cpu/sgl-kernels/conv.cpp"
         "csrc/cpu/sgl-kernels/gemm.cpp"
         "csrc/cpu/sgl-kernels/gemm_int8.cpp"
         "csrc/cpu/sgl-kernels/gemm_fp8.cpp"
         "csrc/cpu/sgl-kernels/gemm_int4.cpp"
         "csrc/cpu/sgl-kernels/moe.cpp"
         "csrc/cpu/sgl-kernels/moe_int8.cpp"
-        "csrc/cpu/sgl-kernels/moe_int4.cpp"
         "csrc/cpu/sgl-kernels/moe_fp8.cpp")
 
     set(VLLM_EXT_SRC_AVX512
@@ -438,11 +430,10 @@ if (ENABLE_X86_ISA)
         "csrc/cpu/pos_encoding.cpp"
         "csrc/moe/dynamic_4bit_int_moe_cpu.cpp") 
 
-    set(VLLM_EXT_SRC_AVX2
+    set(VLLM_EXT_SRC_AVX2 
         "csrc/cpu/utils.cpp"
         "csrc/cpu/spec_decode_utils.cpp"
         "csrc/cpu/cpu_attn.cpp"
-        "csrc/cpu/dnnl_kernels.cpp"
         "csrc/cpu/torch_bindings.cpp"
         # TODO: Remove these files
         "csrc/cpu/activation.cpp"
@@ -457,7 +448,7 @@ if (ENABLE_X86_ISA)
 
     set(_C_LIBS numa dnnl_ext)
     set(_C_AVX512_LIBS numa dnnl_ext)
-    set(_C_AVX2_LIBS numa dnnl_ext)
+    set(_C_AVX2_LIBS numa)
 
     # AMX + AVX512F + AVX512BF16 + AVX512VNNI
     define_extension_target(

cmake/external_projects/deepgemm.cmake

@@ -53,80 +53,49 @@ cuda_archs_loose_intersection(DEEPGEMM_ARCHS
 if(DEEPGEMM_ARCHS)
   message(STATUS "DeepGEMM CUDA architectures: ${DEEPGEMM_ARCHS}")
 
-  #
-  # DeepGEMM integration notes
-  # --------------------------
-  # We vendor DeepGEMM into vllm/third_party/deep_gemm/ and bundle a
-  # `_C.cpython-X.Y-*.so` for every CPython in `requires-python`. The
-  # per-Python build is delegated to tools/build_deepgemm_C.py.
-  #
-  # Why per-Python: DeepGEMM's binding uses PYBIND11_MODULE, which links
-  # private CPython symbols — a single `_C.abi3.so` is not viable today
-  # (see #41476 / #41512 for the failed attempt).
-  #
-  # TODOs (tracked in vllm-project/vllm#42431):
-  #   - Replace DeepGEMM's pybind11 binding with a TORCH_LIBRARY + shim
-  #     binding (cf. vllm-flash-attention/csrc/common/pytorch_shim.h) to
-  #     collapse to one `_C.abi3.so`. Needs either an upstream change or
-  #     a maintained binding fork in vLLM.
-  #   - AOT-compile DeepGEMM's CUDA kernels instead of runtime JIT to drop
-  #     the vendored CUTLASS/CCCL headers and the CUDA-toolkit-at-runtime
-  #     requirement.
-  #
+  find_package(CUDAToolkit REQUIRED)
 
-  # DEEPGEMM_PYTHON_INTERPRETERS: ":"-separated target Python paths.
-  # Empty/unset → fall back to the build interpreter (editable installs).
-  # (Empty-but-set env vars test as DEFINED in cmake — treat as unset.)
-  if(NOT "$ENV{DEEPGEMM_PYTHON_INTERPRETERS}" STREQUAL "")
-    string(REPLACE ":" ";" _dg_pythons "$ENV{DEEPGEMM_PYTHON_INTERPRETERS}")
-  else()
-    set(_dg_pythons "${Python_EXECUTABLE}")
-  endif()
-  message(STATUS "DeepGEMM _C will be built for: ${_dg_pythons}")
+  #
+  # Build the _C pybind11 extension from DeepGEMM's C++ source.
+  # This is a CXX-only module — CUDA kernels are JIT-compiled at runtime.
+  #
+  Python_add_library(_deep_gemm_C MODULE WITH_SOABI
+    "${deepgemm_SOURCE_DIR}/csrc/python_api.cpp")
 
-  # add_custom_command does no implicit header scanning; glob explicitly so
-  # header-only edits in DeepGEMM/cutlass/fmt re-trigger the rebuild.
-  file(GLOB_RECURSE _dg_headers
-    "${deepgemm_SOURCE_DIR}/csrc/*.h"
-    "${deepgemm_SOURCE_DIR}/csrc/*.hpp"
-    "${deepgemm_SOURCE_DIR}/deep_gemm/include/*.h"
-    "${deepgemm_SOURCE_DIR}/deep_gemm/include/*.hpp"
-    "${deepgemm_SOURCE_DIR}/deep_gemm/include/*.cuh")
+  # The pybind11 module name must be _C to match DeepGEMM's Python imports.
+  set_target_properties(_deep_gemm_C PROPERTIES OUTPUT_NAME "_C")
 
-  set(_dg_markers)
-  set(_dg_seen_soabis)
-  foreach(_pybin IN LISTS _dg_pythons)
-    execute_process(
-      COMMAND "${_pybin}" -c
-        "import sysconfig; print(sysconfig.get_config_var('SOABI'))"
-      OUTPUT_VARIABLE _dg_soabi
-      OUTPUT_STRIP_TRAILING_WHITESPACE
-      COMMAND_ERROR_IS_FATAL ANY)
-    # Dedup interpreters that resolve to the same CPython.
-    if(_dg_soabi IN_LIST _dg_seen_soabis)
-      continue()
-    endif()
-    list(APPEND _dg_seen_soabis "${_dg_soabi}")
-    set(_dg_dir "${CMAKE_CURRENT_BINARY_DIR}/deepgemm_C_${_dg_soabi}")
-    set(_dg_marker "${_dg_dir}/.built")
-    add_custom_command(
-      OUTPUT "${_dg_marker}"
-      COMMAND "${Python_EXECUTABLE}"
-              "${CMAKE_SOURCE_DIR}/tools/build_deepgemm_C.py"
-              "${deepgemm_SOURCE_DIR}" "${_dg_dir}" "${_pybin}"
-      COMMAND "${CMAKE_COMMAND}" -E touch "${_dg_marker}"
-      DEPENDS "${CMAKE_SOURCE_DIR}/tools/build_deepgemm_C.py"
-              "${deepgemm_SOURCE_DIR}/csrc/python_api.cpp"
-              ${_dg_headers}
-      COMMENT "Building DeepGEMM _C for ${_pybin}"
-      VERBATIM)
-    list(APPEND _dg_markers "${_dg_marker}")
-    install(DIRECTORY "${_dg_dir}/"
-      DESTINATION vllm/third_party/deep_gemm
-      COMPONENT _deep_gemm_C
-      FILES_MATCHING PATTERN "_C.cpython-*.so")
-  endforeach()
-  add_custom_target(_deep_gemm_C ALL DEPENDS ${_dg_markers})
+  target_compile_definitions(_deep_gemm_C PRIVATE
+    "-DTORCH_EXTENSION_NAME=_C")
+
+  target_include_directories(_deep_gemm_C PRIVATE
+    "${deepgemm_SOURCE_DIR}/csrc"
+    "${deepgemm_SOURCE_DIR}/deep_gemm/include"
+    "${deepgemm_SOURCE_DIR}/third-party/cutlass/include"
+    "${deepgemm_SOURCE_DIR}/third-party/cutlass/tools/util/include"
+    "${deepgemm_SOURCE_DIR}/third-party/fmt/include")
+
+  target_compile_options(_deep_gemm_C PRIVATE
+    $<$<COMPILE_LANGUAGE:CXX>:-std=c++17>
+    $<$<COMPILE_LANGUAGE:CXX>:-O3>
+    $<$<COMPILE_LANGUAGE:CXX>:-Wno-psabi>
+    $<$<COMPILE_LANGUAGE:CXX>:-Wno-deprecated-declarations>)
+
+  # torch_python is required because DeepGEMM uses pybind11 type casters
+  # for at::Tensor (via PYBIND11_MODULE), unlike vLLM's own extensions which
+  # use torch::Library custom ops.
+  find_library(TORCH_PYTHON_LIBRARY torch_python
+    PATHS "${TORCH_INSTALL_PREFIX}/lib"
+    REQUIRED)
+
+  target_link_libraries(_deep_gemm_C PRIVATE
+    torch ${TORCH_LIBRARIES} "${TORCH_PYTHON_LIBRARY}"
+    CUDA::cudart CUDA::nvrtc)
+
+  # Install the shared library into the vendored package directory
+  install(TARGETS _deep_gemm_C
+    LIBRARY DESTINATION vllm/third_party/deep_gemm
+    COMPONENT _deep_gemm_C)
 
   #
   # Vendor DeepGEMM Python package files

cmake/external_projects/vllm_flash_attn.cmake

@@ -39,7 +39,7 @@ else()
   FetchContent_Declare(
           vllm-flash-attn
           GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG bce29425653ec0fbc579d329883030e832d15ada
+          GIT_TAG f5bc33cfc02c744d24a2e9d50e6db656de40611c
           GIT_PROGRESS TRUE
           # Don't share the vllm-flash-attn build between build types
           BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn

csrc/attention/mla/sm100_cutlass_mla_kernel.cu

@@ -18,12 +18,13 @@ limitations under the License.
  * Taken from SGLANG PR https://github.com/sgl-project/sglang/pull/6929
  * by Alcanderian JieXin Liang
  */
-#include "libtorch_stable/torch_utils.h"
-
-#include <torch/csrc/stable/library.h>
+#include "core/registration.h"
 
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
 #include <cutlass/cutlass.h>
 #include <cutlass/kernel_hardware_info.h>
+#include <torch/all.h>
 
 #include <cute/tensor.hpp>
 #include <iostream>
@@ -34,27 +35,27 @@ limitations under the License.
 // clang-format off
 #if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
 void sm100_cutlass_mla_decode(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    torch::Tensor const& out,
+    torch::Tensor const& lse,
+    torch::Tensor const& q_nope,
+    torch::Tensor const& q_pe,
+    torch::Tensor const& kv_c_and_k_pe_cache,
+    torch::Tensor const& seq_lens,
+    torch::Tensor const& page_table,
+    torch::Tensor const& workspace,
     double sm_scale,
     int64_t num_kv_splits) {
-  STD_TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
+  TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
 }
 int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_batches, int64_t sm_count, int64_t num_kv_splits) {
-  STD_TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
+  TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
 }
 #else
 
 #define CUTLASS_CHECK(status)                                                       \
   {                                                                                 \
     cutlass::Status error = status;                                                 \
-    STD_TORCH_CHECK(error == cutlass::Status::kSuccess, cutlassGetStatusString(error)); \
+    TORCH_CHECK(error == cutlass::Status::kSuccess, cutlassGetStatusString(error)); \
   }
 
 using namespace cute;
@@ -99,23 +100,23 @@ struct MlaSm100 {
 
 template <typename T>
 typename T::Fmha::Arguments args_from_options(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
+    at::Tensor const& out,
+    at::Tensor const& lse,
+    at::Tensor const& q_nope,
+    at::Tensor const& q_pe,
+    at::Tensor const& kv_c_and_k_pe_cache,
+    at::Tensor const& seq_lens,
+    at::Tensor const& page_table,
     double sm_scale,
     int64_t num_kv_splits) {
   cutlass::KernelHardwareInfo hw_info;
-  hw_info.device_id = q_nope.get_device_index();
+  hw_info.device_id = q_nope.device().index();
   hw_info.sm_count = cutlass::KernelHardwareInfo::query_device_multiprocessor_count(hw_info.device_id);
 
-  int batches = q_nope.size(0);
-  int page_count_per_seq = page_table.size(1);
-  int page_count_total = kv_c_and_k_pe_cache.size(0);
-  int page_size = kv_c_and_k_pe_cache.size(1);
+  int batches = q_nope.sizes()[0];
+  int page_count_per_seq = page_table.sizes()[1];
+  int page_count_total = kv_c_and_k_pe_cache.sizes()[0];
+  int page_size = kv_c_and_k_pe_cache.sizes()[1];
   int max_seq_len = page_size * page_count_per_seq;
   using TileShapeH = typename T::TileShapeH;
   using TileShapeD = typename T::TileShapeD;
@@ -185,14 +186,14 @@ typename T::Fmha::Arguments args_from_options(
 
 template <typename Element, typename ElementOut, bool IsPaged128, typename PersistenceOption>
 void runMla(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    at::Tensor const& out,
+    at::Tensor const& lse,
+    at::Tensor const& q_nope,
+    at::Tensor const& q_pe,
+    at::Tensor const& kv_c_and_k_pe_cache,
+    at::Tensor const& seq_lens,
+    at::Tensor const& page_table,
+    at::Tensor const& workspace,
     double sm_scale,
     int64_t num_kv_splits,
     cudaStream_t stream) {
@@ -219,37 +220,37 @@ void runMla(
   }()
 
 void sm100_cutlass_mla_decode(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    torch::Tensor const& out,
+    torch::Tensor const& lse,
+    torch::Tensor const& q_nope,
+    torch::Tensor const& q_pe,
+    torch::Tensor const& kv_c_and_k_pe_cache,
+    torch::Tensor const& seq_lens,
+    torch::Tensor const& page_table,
+    torch::Tensor const& workspace,
     double sm_scale,
     int64_t num_kv_splits) {
-  auto in_dtype = q_nope.scalar_type();
-  torch::stable::accelerator::DeviceGuard device_guard(q_nope.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream(q_nope.get_device_index());
-  const int page_size = kv_c_and_k_pe_cache.size(1);
-
+  auto in_dtype = q_nope.dtype();
+  at::cuda::CUDAGuard device_guard{(char)q_nope.get_device()};
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream(q_nope.get_device());
+  const int page_size = kv_c_and_k_pe_cache.sizes()[1];
+  
   // NOTE(alcanderian): IsPersistent has bug with manual split_kv.
   // Kernel will hang if batch is too large with large num_kv_splits. (for example bs=8, num_kv_splits=8)
   // Maybe per batch split kv will fix this.
   DISPATCH_BOOL(page_size == 128, IsPaged128, [&] {
     DISPATCH_BOOL(num_kv_splits <= 1, NotManualSplitKV, [&] {
-      if (in_dtype == torch::headeronly::ScalarType::Half) {
+      if (in_dtype == at::ScalarType::Half) {
         runMla<cutlass::half_t, cutlass::half_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
           out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
-      } else if (in_dtype == torch::headeronly::ScalarType::BFloat16) {
+      } else if (in_dtype == at::ScalarType::BFloat16) {
         runMla<cutlass::bfloat16_t, cutlass::bfloat16_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
           out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
-      } else if (in_dtype == torch::headeronly::ScalarType::Float8_e4m3fn) {
+      } else if (in_dtype == at::ScalarType::Float8_e4m3fn) {
         runMla<cutlass::float_e4m3_t, cutlass::bfloat16_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
           out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
       } else {
-        STD_TORCH_CHECK(false, "Unsupported input data type of MLA");
+        TORCH_CHECK(false, "Unsupported input data type of MLA");
       }
       return true;
     });
@@ -279,12 +280,12 @@ int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_ba
 
 #endif
 
-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("sm100_cutlass_mla_decode", TORCH_BOX(&sm100_cutlass_mla_decode));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+  m.impl("sm100_cutlass_mla_decode", &sm100_cutlass_mla_decode);
 }
 
-STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, m) {
-  m.impl("sm100_cutlass_mla_get_workspace_size", TORCH_BOX(&sm100_cutlass_mla_get_workspace_size));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CatchAll, m) {
+  m.impl("sm100_cutlass_mla_get_workspace_size", &sm100_cutlass_mla_get_workspace_size);
 }
 
 // clang-format on

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/cache.h

@@ -12,8 +12,7 @@ void swap_blocks(torch::Tensor& src, torch::Tensor& dst,
 
 void swap_blocks_batch(const torch::Tensor& src_ptrs,
                        const torch::Tensor& dst_ptrs,
-                       const torch::Tensor& sizes,
-                       bool is_src_access_order_any);
+                       const torch::Tensor& sizes);
 
 void reshape_and_cache(torch::Tensor& key, torch::Tensor& value,
                        torch::Tensor& key_cache, torch::Tensor& value_cache,

csrc/cache_kernels.cu

@@ -77,8 +77,7 @@ void swap_blocks(torch::Tensor& src, torch::Tensor& dst,
 
 void swap_blocks_batch(const torch::Tensor& src_ptrs,
                        const torch::Tensor& dst_ptrs,
-                       const torch::Tensor& sizes,
-                       bool is_src_access_order_any) {
+                       const torch::Tensor& sizes) {
   TORCH_CHECK(src_ptrs.device().is_cpu(), "src_ptrs must be on CPU");
   TORCH_CHECK(dst_ptrs.device().is_cpu(), "dst_ptrs must be on CPU");
   TORCH_CHECK(sizes.device().is_cpu(), "sizes must be on CPU");
@@ -98,13 +97,13 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
 
   const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
 
-  // Use cuMemcpyBatchAsync / hipMemcpyBatchAsync to submit all copies in a
-  // single driver call, amortizing per-copy submission overhead. int64_t
-  // and CUdeviceptr/void*/size_t are all 8 bytes on 64-bit platforms, so we
-  // reinterpret_cast the tensor data directly to avoid copies.
+  // Use cuMemcpyBatchAsync (CUDA 12.8+) to submit all copies in a single
+  // driver call, amortizing per-copy submission overhead.
+  // int64_t and CUdeviceptr/size_t are both 8 bytes on 64-bit platforms,
+  // so we reinterpret_cast the tensor data directly to avoid copies.
+  static_assert(sizeof(CUdeviceptr) == sizeof(int64_t));
   static_assert(sizeof(size_t) == sizeof(int64_t));
 #if !defined(USE_ROCM) && defined(CUDA_VERSION) && CUDA_VERSION >= 12080
-  static_assert(sizeof(CUdeviceptr) == sizeof(int64_t));
   // Resolve cuMemcpyBatchAsync at runtime via cuGetProcAddress so that
   // binaries compiled with CUDA 12.8+ still work on older drivers, and
   // we avoid the CUDA 13.0 header remapping (#define to _v2 signature).
@@ -125,12 +124,7 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
 
   if (batch_fn != nullptr) {
     CUmemcpyAttributes attr = {};
-    // ANY lets the DMA engine prefetch source bytes out of stream order,
-    // which is only safe when no GPU stream is concurrently writing the
-    // source.
-    attr.srcAccessOrder = is_src_access_order_any
-                              ? CU_MEMCPY_SRC_ACCESS_ORDER_ANY
-                              : CU_MEMCPY_SRC_ACCESS_ORDER_STREAM;
+    attr.srcAccessOrder = CU_MEMCPY_SRC_ACCESS_ORDER_STREAM;
     size_t attrs_idx = 0;
     size_t fail_idx = 0;
     CUresult result = batch_fn(reinterpret_cast<CUdeviceptr*>(dst_data),
@@ -140,30 +134,12 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
                                &fail_idx, static_cast<CUstream>(stream));
     TORCH_CHECK(result == CUDA_SUCCESS, "cuMemcpyBatchAsync failed at index ",
                 fail_idx, " with error ", result);
-    return;
-  }
-#elif defined(USE_ROCM) && defined(HIP_VERSION) && HIP_VERSION >= 70100000
-  // ROCm 7.1+ exposes hipMemcpyBatchAsync. The 7.2.1 implementation early-
-  // returns hipErrorNotSupported whenever numAttrs > 0 (see ROCm/clr @
-  // rocm-7.2.1 hipamd/src/hip_memory.cpp:2819-2822), so call with
-  // numAttrs=0.
-  {
-    hipMemcpyAttributes attr = {};
-    size_t attrs_idx = 0;
-    size_t fail_idx = 0;
-    hipError_t result = hipMemcpyBatchAsync(
-        reinterpret_cast<void**>(dst_data), reinterpret_cast<void**>(src_data),
-        reinterpret_cast<size_t*>(size_data), static_cast<size_t>(n), &attr,
-        &attrs_idx, 0, &fail_idx, static_cast<hipStream_t>(stream));
-    TORCH_CHECK(result == hipSuccess, "hipMemcpyBatchAsync failed at index ",
-                fail_idx, " with error ", result);
-    return;
-  }
+  } else
 #endif
   {
-    // Fallback for CUDA < 12.8, older CUDA drivers, and ROCm < 7.1:
-    // individual async copies. cudaMemcpyDefault lets the driver infer
-    // direction from pointer types.
+    // Fallback for CUDA < 12.8, older drivers, and ROCm:
+    // individual async copies.
+    // cudaMemcpyDefault lets the driver infer direction from pointer types.
     for (int64_t i = 0; i < n; i++) {
       cudaMemcpyAsync(reinterpret_cast<void*>(dst_data[i]),
                       reinterpret_cast<void*>(src_data[i]),

csrc/cache_kernels_fused.cu

@@ -21,33 +21,28 @@ namespace vllm {
 
 // NOTE Be EXTRA careful with raw_kv_scalar_t, for __half and __nv_bfloat16 it's
 // using u16 as the backing type.
-template <typename qk_t, typename cos_sin_t, bool IS_NEOX,
-          typename raw_kv_scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt>
+template <typename qk_t, bool IS_NEOX, typename raw_kv_scalar_t,
+          typename cache_t, Fp8KVCacheDataType kv_dt>
 __global__ void concat_and_cache_mla_rope_fused_kernel(
     const int64_t* __restrict__ positions,  // [num_tokens]
     qk_t* __restrict__ q_pe,        // [num_tokens, num_q_heads, rot_dim]
     qk_t* __restrict__ k_pe,        // [num_tokens, rot_dim]
     const qk_t* __restrict__ kv_c,  // [num_tokens, kv_lora_rank]
-    const cos_sin_t* __restrict__ rope_cos_sin_cache,  // [max_position, 2,
-                                                       // rot_dim // 2]
+    const qk_t* __restrict__ rope_cos_sin_cache,  // [max_position, 2,
+                                                  // rot_dim // 2]
     const int rot_dim, const int64_t q_pe_stride_token,
     const int64_t q_pe_stride_head, const int64_t k_pe_stride,
     const int64_t kv_c_stride, const int num_q_heads,
     cache_t* __restrict__ kv_cache,  // [num_blocks, block_size, (kv_lora_rank +
                                      // rot_dim)]
-    const int64_t* __restrict__ slot_mapping,  // [num_tokens]
+    const int64_t* __restrict__ kv_cache_slot_mapping,  // [num_tokens]
     const int block_stride, const int entry_stride, const int kv_lora_rank,
     const int block_size, const float* kv_cache_quant_scale) {
   // Each thread block is responsible for one token.
   const int64_t token_idx = blockIdx.x;
-  const int64_t slot_idx = slot_mapping[token_idx];
-  // NOTE: slot_idx can be -1 if the token is padded
-  if (slot_idx < 0) {
-    return;
-  }
   const int64_t pos = positions[token_idx];
 
-  const cos_sin_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;
+  const qk_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;
 
   const int embed_dim = rot_dim / 2;
 
@@ -59,8 +54,8 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
 
     // NOTE: Would be nice to have interleaved sin/cos so we could just load
     // both at the same time.
-    qk_t cos = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx));
-    qk_t sin = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim));
+    qk_t cos = VLLM_LDG(cos_sin_ptr + pair_idx);
+    qk_t sin = VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim);
 
     qk_t* q_pe_head_ptr =
         q_pe + token_idx * q_pe_stride_token + head_idx * q_pe_stride_head;
@@ -86,15 +81,21 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
     q_pe_head_ptr[pair_idx_y] = y_dst;
   }
 
+  const int64_t slot_idx = kv_cache_slot_mapping[token_idx];
   const int64_t block_idx = slot_idx / block_size;
   const int64_t entry_idx = slot_idx % block_size;
 
+  // NOTE: slot_idx can be -1 if the token is padded
+  if (slot_idx < 0) {
+    return;
+  }
+
   // K with 1 HEAD
   for (int i = threadIdx.x; i < embed_dim; i += blockDim.x) {
     int pair_idx = i;
 
-    qk_t cos = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx));
-    qk_t sin = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim));
+    qk_t cos = VLLM_LDG(cos_sin_ptr + pair_idx);
+    qk_t sin = VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim);
 
     qk_t* k_pe_head_ptr = k_pe + token_idx * k_pe_stride;
 
@@ -164,43 +165,36 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
 
 }  // namespace vllm
 
-#define CALL_CONCAT_AND_CACHE_MLA_ROPE_FUSED(RAW_KV_T, CACHE_T, KV_DTYPE)     \
-  do {                                                                        \
-    VLLM_DISPATCH_FLOATING_TYPES(q_pe.scalar_type(), "qk_scalar_type", [&] {  \
-      using qk_t = scalar_t;                                                  \
-      VLLM_DISPATCH_FLOATING_TYPES(                                           \
-          rope_cos_sin_cache.scalar_type(), "rope_cos_sin_cache_scalar_type", \
-          [&] {                                                               \
-            using cos_sin_t = scalar_t;                                       \
-            if (rope_is_neox) {                                               \
-              vllm::concat_and_cache_mla_rope_fused_kernel<                   \
-                  qk_t, cos_sin_t, true, RAW_KV_T, CACHE_T, KV_DTYPE>         \
-                  <<<grid, block, 0, stream>>>(                               \
-                      positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),   \
-                      k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),           \
-                      rope_cos_sin_cache.data_ptr<cos_sin_t>(), rot_dim,      \
-                      q_pe_stride_token, q_pe_stride_head, k_pe_stride,       \
-                      kv_c_stride, num_q_heads,                               \
-                      reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),        \
-                      slot_mapping.data_ptr<int64_t>(), block_stride,         \
-                      entry_stride, kv_lora_rank, block_size,                 \
-                      kv_cache_quant_scale.data_ptr<float>());                \
-            } else {                                                          \
-              vllm::concat_and_cache_mla_rope_fused_kernel<                   \
-                  qk_t, cos_sin_t, false, RAW_KV_T, CACHE_T, KV_DTYPE>        \
-                  <<<grid, block, 0, stream>>>(                               \
-                      positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),   \
-                      k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),           \
-                      rope_cos_sin_cache.data_ptr<cos_sin_t>(), rot_dim,      \
-                      q_pe_stride_token, q_pe_stride_head, k_pe_stride,       \
-                      kv_c_stride, num_q_heads,                               \
-                      reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),        \
-                      slot_mapping.data_ptr<int64_t>(), block_stride,         \
-                      entry_stride, kv_lora_rank, block_size,                 \
-                      kv_cache_quant_scale.data_ptr<float>());                \
-            }                                                                 \
-          });                                                                 \
-    });                                                                       \
+#define CALL_CONCAT_AND_CACHE_MLA_ROPE_FUSED(RAW_KV_T, CACHE_T, KV_DTYPE)      \
+  do {                                                                         \
+    VLLM_DISPATCH_FLOATING_TYPES(q_pe.scalar_type(), "qk_scalar_type", [&] {   \
+      using qk_t = scalar_t;                                                   \
+      if (rope_is_neox) {                                                      \
+        vllm::concat_and_cache_mla_rope_fused_kernel<qk_t, true, RAW_KV_T,     \
+                                                     CACHE_T, KV_DTYPE>        \
+            <<<grid, block, 0, stream>>>(                                      \
+                positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),          \
+                k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),                  \
+                rope_cos_sin_cache.data_ptr<qk_t>(), rot_dim,                  \
+                q_pe_stride_token, q_pe_stride_head, k_pe_stride, kv_c_stride, \
+                num_q_heads, reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),  \
+                kv_cache_slot_mapping.data_ptr<int64_t>(), block_stride,       \
+                entry_stride, kv_lora_rank, block_size,                        \
+                kv_cache_quant_scale.data_ptr<float>());                       \
+      } else {                                                                 \
+        vllm::concat_and_cache_mla_rope_fused_kernel<qk_t, false, RAW_KV_T,    \
+                                                     CACHE_T, KV_DTYPE>        \
+            <<<grid, block, 0, stream>>>(                                      \
+                positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),          \
+                k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),                  \
+                rope_cos_sin_cache.data_ptr<qk_t>(), rot_dim,                  \
+                q_pe_stride_token, q_pe_stride_head, k_pe_stride, kv_c_stride, \
+                num_q_heads, reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),  \
+                kv_cache_slot_mapping.data_ptr<int64_t>(), block_stride,       \
+                entry_stride, kv_lora_rank, block_size,                        \
+                kv_cache_quant_scale.data_ptr<float>());                       \
+      }                                                                        \
+    });                                                                        \
   } while (false)
 
 // Executes RoPE on q_pe and k_pe, then writes k_pe and kv_c in the kv cache.
@@ -214,52 +208,43 @@ void concat_and_cache_mla_rope_fused(
     torch::Tensor& kv_c,                // [num_tokens, kv_lora_rank]
     torch::Tensor& rope_cos_sin_cache,  // [max_position, rot_dim]
     bool rope_is_neox,
-    torch::Tensor& slot_mapping,  // [num_tokens] or [num_actual_tokens]
+    torch::Tensor&
+        kv_cache_slot_mapping,  // [num_tokens] or [num_actual_tokens]
     torch::Tensor&
         kv_cache,  // [num_blocks, block_size, (kv_lora_rank + rot_dim)]
     const std::string& kv_cache_dtype, torch::Tensor& kv_cache_quant_scale) {
-  // NOTE(woosuk): In vLLM V1, query/key/position.size(0) can be different from
-  // slot_mapping.size(0) because of padding for CUDA graphs.
-  // In vLLM V0, key.size(0) is always equal to slot_mapping.size(0) because
-  // both include padding.
-  // In vLLM V1, however, key.size(0) can be larger than slot_mapping.size(0)
-  // since key includes padding for CUDA graphs, while slot_mapping does not.
-  // In this case, slot_mapping.size(0) represents the actual number of tokens
-  // before padding.
-  // For compatibility with both cases, we use slot_mapping.size(0) as the
-  // number of tokens.
-  int num_tokens = slot_mapping.size(0);
-  int num_padded_tokens = q_pe.size(0);
-  TORCH_CHECK_GE(num_padded_tokens, num_tokens);
+  const int64_t num_tokens = q_pe.size(0);
 
   const int num_q_heads = q_pe.size(1);
   const int rot_dim = q_pe.size(2);
   const int kv_lora_rank = kv_c.size(1);
 
-  TORCH_CHECK_EQ(positions.size(0), num_padded_tokens);
+  TORCH_CHECK(positions.size(0) >=
+              num_tokens);  // CUDA Graphs might pad this for us
   TORCH_CHECK_EQ(positions.dim(), 1);
   TORCH_CHECK_EQ(positions.scalar_type(), c10::ScalarType::Long);
 
-  TORCH_CHECK_EQ(q_pe.dim(), 3);
-  TORCH_CHECK_EQ(q_pe.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(q_pe.size(0), num_tokens);
   TORCH_CHECK_EQ(q_pe.size(1), num_q_heads);
   TORCH_CHECK_EQ(q_pe.size(2), rot_dim);
+  TORCH_CHECK_EQ(q_pe.dim(), 3);
 
-  TORCH_CHECK_EQ(k_pe.dim(), 2);
-  TORCH_CHECK_EQ(k_pe.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(k_pe.size(0), num_tokens);
   TORCH_CHECK_EQ(k_pe.size(1), rot_dim);
+  TORCH_CHECK_EQ(k_pe.dim(), 2);
   TORCH_CHECK_EQ(k_pe.scalar_type(), q_pe.scalar_type());
 
-  TORCH_CHECK_EQ(kv_c.dim(), 2);
-  TORCH_CHECK_EQ(kv_c.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(kv_c.size(0), num_tokens);
   TORCH_CHECK_EQ(kv_c.size(1), kv_lora_rank);
+  TORCH_CHECK_EQ(kv_c.dim(), 2);
   TORCH_CHECK_EQ(kv_c.scalar_type(), q_pe.scalar_type());
   TORCH_CHECK_EQ(kv_c.dtype(), q_pe.dtype());
 
   TORCH_CHECK_EQ(rope_cos_sin_cache.size(1), rot_dim);
+  TORCH_CHECK_EQ(rope_cos_sin_cache.scalar_type(), q_pe.scalar_type());
 
-  TORCH_CHECK_EQ(slot_mapping.size(0), num_tokens);
-  TORCH_CHECK_EQ(slot_mapping.scalar_type(), c10::ScalarType::Long);
+  TORCH_CHECK_EQ(kv_cache_slot_mapping.size(0), num_tokens);
+  TORCH_CHECK_EQ(kv_cache_slot_mapping.scalar_type(), c10::ScalarType::Long);
 
   TORCH_CHECK_EQ(kv_cache.size(2), kv_lora_rank + rot_dim);
   TORCH_CHECK_EQ(kv_cache.dim(), 3);

csrc/core/batch_invariant.hpp

@@ -1,6 +1,7 @@
 #pragma once
 #include <cstdlib>
 #include <string>
+#include <cctype>
 
 namespace vllm {
 

csrc/core/scalar_type.hpp

@@ -1,13 +1,7 @@
 #pragma once
 
-#include <cstdint>
-#include <string>
-#include <tuple>
-#include <utility>
-#include <variant>
-
-// For STD_TORCH_CHECK
-#include <torch/headeronly/util/Exception.h>
+// For TORCH_CHECK
+#include <torch/library.h>
 
 namespace vllm {
 
@@ -51,7 +45,7 @@ class ScalarType {
   // IEEE 754 compliant floating point type
   static constexpr ScalarType float_IEEE754(uint8_t exponent,
                                             uint8_t mantissa) {
-    STD_TORCH_CHECK(mantissa > 0 && exponent > 0);
+    TORCH_CHECK(mantissa > 0 && exponent > 0);
     return ScalarType(exponent, mantissa, true, 0, false, NAN_IEEE_754);
   }
 
@@ -59,12 +53,11 @@ class ScalarType {
   static constexpr ScalarType float_(uint8_t exponent, uint8_t mantissa,
                                      bool finite_values_only,
                                      NanRepr nan_repr) {
-    STD_TORCH_CHECK(nan_repr < NAN_REPR_ID_MAX, "Invalid NanRepr");
-    STD_TORCH_CHECK(mantissa > 0 && exponent > 0);
-    STD_TORCH_CHECK(
-        nan_repr != NAN_IEEE_754,
-        "use `float_IEEE754` constructor for floating point types that "
-        "follow IEEE 754 conventions");
+    TORCH_CHECK(nan_repr < NAN_REPR_ID_MAX, "Invalid NanRepr");
+    TORCH_CHECK(mantissa > 0 && exponent > 0);
+    TORCH_CHECK(nan_repr != NAN_IEEE_754,
+                "use `float_IEEE754` constructor for floating point types that "
+                "follow IEEE 754 conventions");
     return ScalarType(exponent, mantissa, true, 0, finite_values_only,
                       nan_repr);
   }
@@ -183,8 +176,8 @@ class ScalarType {
 
  private:
   double _floating_point_max() const {
-    STD_TORCH_CHECK(mantissa <= 52 && exponent <= 11,
-                    "Cannot represent max/min as a double for type ", str());
+    TORCH_CHECK(mantissa <= 52 && exponent <= 11,
+                "Cannot represent max/min as a double for type ", str());
 
     uint64_t max_mantissa = (uint64_t(1) << mantissa) - 1;
     if (nan_repr == NAN_EXTD_RANGE_MAX_MIN) {
@@ -193,8 +186,8 @@ class ScalarType {
 
     uint64_t max_exponent = (uint64_t(1) << exponent) - 2;
     if (nan_repr == NAN_EXTD_RANGE_MAX_MIN || nan_repr == NAN_NONE) {
-      STD_TORCH_CHECK(exponent < 11,
-                      "Cannot represent max/min as a double for type ", str());
+      TORCH_CHECK(exponent < 11,
+                  "Cannot represent max/min as a double for type ", str());
       max_exponent += 1;
     }
 
@@ -223,17 +216,16 @@ class ScalarType {
     if (is_floating_point()) {
       return {_floating_point_max()};
     } else {
-      STD_TORCH_CHECK(size_bits() < 64 || size_bits() == 64 && is_signed(),
-                      "Cannot represent max as a int64_t");
+      TORCH_CHECK(size_bits() < 64 || size_bits() == 64 && is_signed(),
+                  "Cannot represent max as a int64_t");
       return {(int64_t(1) << mantissa) - 1};
     }
   }
 
   constexpr std::variant<int64_t, double> _raw_min() const {
     if (is_floating_point()) {
-      STD_TORCH_CHECK(
-          is_signed(),
-          "We currently assume all floating point types are signed");
+      TORCH_CHECK(is_signed(),
+                  "We currently assume all floating point types are signed");
       constexpr uint64_t sign_bit_double = (uint64_t(1) << 63);
 
       double max = _floating_point_max();
@@ -241,8 +233,8 @@ class ScalarType {
       uint64_t min_raw = max_raw | sign_bit_double;
       return {*reinterpret_cast<double*>(&min_raw)};
     } else {
-      STD_TORCH_CHECK(!is_signed() || size_bits() <= 64,
-                      "Cannot represent min as a int64_t");
+      TORCH_CHECK(!is_signed() || size_bits() <= 64,
+                  "Cannot represent min as a int64_t");
       if (is_signed()) {
         // set the top bit to 1 (i.e. INT64_MIN) and the rest to 0
         // then perform an arithmetic shift right to set all the bits above

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,10 +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 {
     TORCH_CHECK(false, "Unsupported CPU attention ISA hint: " + isa_hint);
   }
@@ -133,10 +129,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 {
       TORCH_CHECK(false, "Invalid ISA type: " + isa);
     }

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 };
 
 // Mirrors csrc/attention/dtype_fp8.cuh Fp8KVCacheDataType exactly.
 enum class Fp8KVCacheDataType {
@@ -164,12 +164,6 @@ struct AttentionMetadata {
       case ISA::VXE:
         ss << "VXE, ";
         break;
-      case ISA::RVV:
-        ss << "RVV, ";
-        break;
-      case ISA::VSX:
-        ss << "VSX, ";
-        break;
     }
     ss << "workitem_group_num: " << workitem_group_num
        << ", reduction_item_num: " << reduction_item_num
@@ -459,7 +453,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);

csrc/cpu/cpu_attn_rvv.hpp

@@ -1,445 +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
-
-// This kernel is currently hardcoded to VLEN=128 (m1/m2 intrinsics, vl=8).
-// The fixed-width typedefs below use `riscv_rvv_vector_bits(128)`, which
-// only matches `vfloat16m1_t`/`vuint16m1_t` register layout when VLEN==128;
-// at VLEN>=256 those typedefs fail to compile.  Scalar RISC-V builds
-// (-march=rv64gc) additionally don't have <riscv_vector.h>.  For both
-// cases we omit the file entirely and let the dispatcher fall back to the
-// scalar VEC / VEC16 implementations.  TODO: migrate to RVVI() macros +
-// semantic names in cpu_types_riscv_defs.hpp to support VLEN>=256 natively.
-#if defined(__riscv_v_min_vlen) && __riscv_v_min_vlen == 128
-
-  #include "cpu_attn_impl.hpp"
-  #include <riscv_vector.h>
-  #include <type_traits>
-
-namespace cpu_attention {
-
-namespace {
-
-// File-local concrete-LMUL typedefs.  The shared _defs.hpp exposes
-// VLEN-independent semantic names (fixed_fp32x8_t, fixed_fp16x8_t, ...),
-// but this kernel is currently hardcoded to VLEN=128 (m1/m2 intrinsics),
-// so keep the legacy concrete aliases scoped to this file.
-typedef vfloat16m1_t fixed_vfloat16m1_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef vfloat32m2_t fixed_vfloat32m2_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-typedef vuint16m1_t fixed_vuint16m1_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef vuint32m2_t fixed_vuint32m2_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-  #ifdef __riscv_zvfbfmin
-typedef vbfloat16m1_t fixed_vbfloat16m1_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-  #endif
-
-  #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 (using m2 LMUL at VLEN=128)
-// ============================================================================
-
-template <typename kv_cache_t>
-FORCE_INLINE fixed_vfloat32m2_t load_row8_B_as_f32(const kv_cache_t* p);
-
-template <>
-FORCE_INLINE fixed_vfloat32m2_t load_row8_B_as_f32<float>(const float* p) {
-  return __riscv_vle32_v_f32m2(p, 8);
-}
-
-template <>
-FORCE_INLINE fixed_vfloat32m2_t
-load_row8_B_as_f32<c10::Half>(const c10::Half* p) {
-  #ifdef __riscv_zvfh
-  fixed_vfloat16m1_t h =
-      __riscv_vle16_v_f16m1(reinterpret_cast<const _Float16*>(p), 8);
-  return __riscv_vfwcvt_f_f_v_f32m2(h, 8);
-  #else
-  // Fallback for hardware without Zvfh: scalar half->float conversion.
-  // c10::Half provides operator float() so this is correct on any RVV CPU
-  // that has only the base V extension.  Slower than the Zvfh path, but
-  // keeps the kernel buildable on Zvfhmin-only / no-fp16 hardware.
-  alignas(16) float tmp[8];
-  for (int i = 0; i < 8; ++i) {
-    tmp[i] = static_cast<float>(p[i]);
-  }
-  return __riscv_vle32_v_f32m2(tmp, 8);
-  #endif
-}
-
-template <>
-FORCE_INLINE fixed_vfloat32m2_t
-load_row8_B_as_f32<c10::BFloat16>(const c10::BFloat16* p) {
-  #ifdef __riscv_zvfbfmin
-  fixed_vbfloat16m1_t bf =
-      __riscv_vle16_v_bf16m1(reinterpret_cast<const __bf16*>(p), 8);
-  return __riscv_vfwcvtbf16_f_f_v_f32m2(bf, 8);
-  #else
-  // Fallback: load as uint16, zero-extend to uint32, shift left by 16
-  fixed_vuint16m1_t raw =
-      __riscv_vle16_v_u16m1(reinterpret_cast<const uint16_t*>(p), 8);
-  fixed_vuint32m2_t wide = __riscv_vzext_vf2_u32m2(raw, 8);
-  fixed_vuint32m2_t shifted = __riscv_vsll_vx_u32m2(wide, 16, 8);
-  return __riscv_vreinterpret_v_u32m2_f32m2(shifted);
-  #endif
-}
-
-// ============================================================================
-// Micro kernel: Mx8 tile, K unrolled by 4, RVV scalar-broadcast FMA
-// ============================================================================
-//
-// NEON uses vfmaq_laneq_f32 (lane-indexed FMA from a preloaded A vector).
-// RVV has no lane-indexed FMA; instead we load A elements as scalars and
-// use __riscv_vfmacc_vf (scalar * vector + accumulator), which is equally
-// efficient and avoids the need for vrgather/vslidedown.
-//
-// At VLEN=128, m2 holds 8 x FP32, matching the 8-column tile width.
-// Register budget: M accumulators (m2 each) + 1 B temp = 2M+2 regs.
-// M=8 => 18 regs out of 32 available — no spills.
-
-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;
-
-  // helpers for per-M codegen
-  #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))
-
-    // A row base pointers
-  #define DECL_A(i) const float* a##i = A + (i) * lda;
-  ROWS_APPLY(DECL_A)
-  #undef DECL_A
-
-    // declare one m2 accumulator per row
-  #define DECL_ACC(i) fixed_vfloat32m2_t acc##i;
-  ROWS_APPLY(DECL_ACC)
-  #undef DECL_ACC
-
-    // initialize accumulators
-  #define INIT_ACC(i)                                      \
-    IF_M(i) {                                              \
-      if (accumulate) {                                    \
-        acc##i = __riscv_vle32_v_f32m2(C + (i) * ldc, vl); \
-      } else {                                             \
-        acc##i = __riscv_vfmv_v_f_f32m2(0.f, vl);          \
-      }                                                    \
-    }
-  ROWS_APPLY(INIT_ACC)
-  #undef INIT_ACC
-
-  int32_t k = 0;
-
-  // K unrolled by 4
-  for (; k + 3 < K; k += 4) {
-    // k + 0
-    {
-      fixed_vfloat32m2_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 = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 0), b, vl); \
-    }
-      ROWS_APPLY(STEP_K0)
-  #undef STEP_K0
-    }
-    // k + 1
-    {
-      fixed_vfloat32m2_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 = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 1), b, vl); \
-    }
-      ROWS_APPLY(STEP_K1)
-  #undef STEP_K1
-    }
-    // k + 2
-    {
-      fixed_vfloat32m2_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 = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 2), b, vl); \
-    }
-      ROWS_APPLY(STEP_K2)
-  #undef STEP_K2
-    }
-    // k + 3
-    {
-      fixed_vfloat32m2_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 = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 3), b, vl); \
-    }
-      ROWS_APPLY(STEP_K3)
-  #undef STEP_K3
-    }
-  }
-
-  // K tail
-  for (; k < K; ++k) {
-    fixed_vfloat32m2_t b = load_row8_B_as_f32<kv_cache_t>(B + (int64_t)k * ldb);
-  #define TAIL_ROW(i) \
-    IF_M(i) { acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k), b, vl); }
-    ROWS_APPLY(TAIL_ROW)
-  #undef TAIL_ROW
-  }
-
-    // store accumulators to C
-  #define STORE_ROW(i) \
-    IF_M(i) { __riscv_vse32_v_f32m2(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;
-        {
-          // Write Key (transpose to column-major: [head_dim, 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;
-
-          // Strided vector store for efficient transpose.
-          // Load contiguous key elements, store with stride = block_size.
-          {
-            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 {
-                // Half and BFloat16 are both 16-bit types
-                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;
-            }
-          }
-        }
-        {
-          // Write Value (row-major: [block_size, head_dim])
-          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
-
-#endif  // CPU_ATTN_RVV_HPP

csrc/cpu/cpu_attn_vec.hpp

@@ -27,8 +27,8 @@ FORCE_INLINE std::pair<vec_op::FP32Vec16, vec_op::FP32Vec16> load_b_pair_vec(
     return {vec_op::FP32Vec16(bf16_b_reg, 0), vec_op::FP32Vec16(bf16_b_reg, 1)};
   } else {
     using load_vec_t = typename VecTypeTrait<kv_cache_t>::vec_t;
-    return std::make_pair(vec_op::FP32Vec16(load_vec_t(ptr)),
-                          vec_op::FP32Vec16(load_vec_t(ptr + 16)));
+    return {vec_op::FP32Vec16(load_vec_t(ptr)),
+            vec_op::FP32Vec16(load_vec_t(ptr + 16))};
   }
 }
 

csrc/cpu/cpu_attn_vsx.hpp

@@ -1,359 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#ifndef CPU_ATTN_VSX_HPP
-#define CPU_ATTN_VSX_HPP
-
-#include "cpu_attn_impl.hpp"
-#include <altivec.h>
-#include <type_traits>
-
-namespace cpu_attention {
-
-namespace {
-
-// ppc64le Vector = 16 bytes (128 bits)
-#define BLOCK_SIZE_ALIGNMENT 32
-#define HEAD_SIZE_ALIGNMENT 32
-#define MAX_Q_HEAD_NUM_PER_ITER 16
-
-template <typename kv_cache_t>
-FORCE_INLINE void load_row8_B_as_f32(const kv_cache_t* p, __vector float& b0,
-                                     __vector float& b1);
-
-// [1] Float Specialization
-template <>
-FORCE_INLINE void load_row8_B_as_f32<float>(const float* p, __vector float& b0,
-                                            __vector float& b1) {
-  b0 = vec_xl(0, const_cast<float*>(p));
-  b1 = vec_xl(0, const_cast<float*>(p + 4));
-}
-
-// [2] BFloat16 Specialization (Little Endian ppc64le)
-// On ppc64le (LE): BF16 bits should land in the HIGH 16 bits of each float32.
-// Byte layout of float32 on LE: [byte0(LSB), byte1, byte2, byte3(MSB)]
-// We need BF16 in bytes2-3 (high half) with bytes0-1 zeroed.
-// vec_mergeh on LE interleaves elements 0..3: result_i = {a[i], b[i]}
-// So vec_mergeh(zeros_u16, raw_u16) gives for each uint16 pair:
-//   uint16[2i]   = zeros[i]  -> low 16 bits of uint32  -> zeroed mantissa LSBs
-//   uint16[2i+1] = raw[i]    -> high 16 bits of uint32 -> BF16 bits
-// Cast to float32 gives exactly (bf16_bits << 16) per element.
-template <>
-FORCE_INLINE void load_row8_B_as_f32<c10::BFloat16>(const c10::BFloat16* p,
-                                                    __vector float& b0,
-                                                    __vector float& b1) {
-  __vector unsigned short raw = vec_xl(
-      0, reinterpret_cast<unsigned short*>(const_cast<c10::BFloat16*>(p)));
-  __vector unsigned short zeros = vec_splat_u16(0);
-
-  // LE: zeros in low 16 bits, raw in high 16 bits → bf16 << 16 == float32
-  b0 = (__vector float)vec_mergeh(zeros, raw);
-  b1 = (__vector float)vec_mergel(zeros, raw);
-}
-
-// Note: c10::Half (FP16) is not supported on PowerPC architecture
-
-template <int32_t M, typename kv_cache_t>
-FORCE_INLINE void gemm_micro_ppc64le_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]");
-
-#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))
-
-  // 1. Define A pointers
-#define DECL_A(i) const float* a##i = A + (i) * lda;
-  ROWS_APPLY(DECL_A)
-#undef DECL_A
-
-  // 2. Define Accumulators (2 vectors covers 8 columns)
-#define DECL_ACC(i) __vector float acc##i##_0, acc##i##_1;
-  ROWS_APPLY(DECL_ACC)
-#undef DECL_ACC
-
-  // 3. Initialize Accumulators (Load C or Zero)
-#define INIT_ACC(i)                                                  \
-  IF_M(i) {                                                          \
-    if (accumulate) {                                                \
-      acc##i##_0 = vec_xl(0, const_cast<float*>(C + (i) * ldc + 0)); \
-      acc##i##_1 = vec_xl(0, const_cast<float*>(C + (i) * ldc + 4)); \
-    } else {                                                         \
-      acc##i##_0 = vec_splats(0.0f);                                 \
-      acc##i##_1 = vec_splats(0.0f);                                 \
-    }                                                                \
-  }
-  ROWS_APPLY(INIT_ACC)
-#undef INIT_ACC
-
-  int32_t k = 0;
-
-  for (; k + 3 < K; k += 4) {
-    // Load 4 values of A for each Row M: A[k...k+3]
-#define LOAD_A4(i)        \
-  __vector float a##i##v; \
-  IF_M(i) a##i##v = vec_xl(0, const_cast<float*>(a##i + k));
-    ROWS_APPLY(LOAD_A4)
-#undef LOAD_A4
-
-    // FMA for specific lane L of A
-    // ppc64le: vec_madd(b, vec_splat(a, lane), acc)
-#define FMAS_LANE(i, aiv, L)                        \
-  IF_M(i) {                                         \
-    __vector float a_broad = vec_splat(aiv, L);     \
-    acc##i##_0 = vec_madd(b0, a_broad, acc##i##_0); \
-    acc##i##_1 = vec_madd(b1, a_broad, acc##i##_1); \
-  }
-
-    // Unroll K=0..3
-    {
-      __vector float b0, b1;
-      load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 0) * ldb, b0, b1);
-#define STEP_K0(i) FMAS_LANE(i, a##i##v, 0)
-      ROWS_APPLY(STEP_K0)
-#undef STEP_K0
-    }
-    {
-      __vector float b0, b1;
-      load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 1) * ldb, b0, b1);
-#define STEP_K1(i) FMAS_LANE(i, a##i##v, 1)
-      ROWS_APPLY(STEP_K1)
-#undef STEP_K1
-    }
-    {
-      __vector float b0, b1;
-      load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 2) * ldb, b0, b1);
-#define STEP_K2(i) FMAS_LANE(i, a##i##v, 2)
-      ROWS_APPLY(STEP_K2)
-#undef STEP_K2
-    }
-    {
-      __vector float b0, b1;
-      load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 3) * ldb, b0, b1);
-#define STEP_K3(i) FMAS_LANE(i, a##i##v, 3)
-      ROWS_APPLY(STEP_K3)
-#undef STEP_K3
-    }
-#undef FMAS_LANE
-  }
-
-  for (; k < K; ++k) {
-    __vector float b0, b1;
-    load_row8_B_as_f32<kv_cache_t>(B + (int64_t)k * ldb, b0, b1);
-#define TAIL_ROW(i)                              \
-  IF_M(i) {                                      \
-    __vector float ai = vec_splats(*(a##i + k)); \
-    acc##i##_0 = vec_madd(b0, ai, acc##i##_0);   \
-    acc##i##_1 = vec_madd(b1, ai, acc##i##_1);   \
-  }
-    ROWS_APPLY(TAIL_ROW)
-#undef TAIL_ROW
-  }
-
-#define STORE_ROW(i)                           \
-  IF_M(i) {                                    \
-    vec_xst(acc##i##_0, 0, C + (i) * ldc + 0); \
-    vec_xst(acc##i##_1, 0, C + (i) * ldc + 4); \
-  }
-  ROWS_APPLY(STORE_ROW)
-#undef STORE_ROW
-
-#undef ROWS_APPLY
-#undef IF_M
-}
-
-template <int32_t N, typename kv_cache_t>
-FORCE_INLINE void gemm_macro_ppc64le_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_ppc64le_Mx8_Ku4<8, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        case 4:
-          gemm_micro_ppc64le_Mx8_Ku4<4, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        case 2:
-          gemm_micro_ppc64le_Mx8_Ku4<2, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        default:
-          gemm_micro_ppc64le_Mx8_Ku4<1, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-      }
-    }
-    m += mb;
-  }
-}
-
-template <typename kv_cache_t>
-class TileGemmPPC64 {
- 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_ppc64le_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_ppc64le_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
-
-template <typename scalar_t, int64_t head_dim>
-class AttentionImpl<ISA::VSX, scalar_t, head_dim> {
- 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::VSX;
-  constexpr static bool scale_on_logits =
-      false;  // Scale is applied to Q during copy
-
- public:
-  AttentionImpl() {}
-
-  template <template <typename tile_gemm_t> typename attention>
-  FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
-    attention<TileGemmPPC64<kv_cache_t>> attention_iteration;
-    attention_iteration(CPU_ATTENTION_PARAMS);
-  }
-
-  // Strides for Memory Layout
-  constexpr static int64_t k_cache_token_group_stride(
-      const int32_t block_size) {
-    return BlockSizeAlignment;  // [head_dim, block_size] layout
-  }
-
-  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) {
-    __vector float scale_vec = vec_splats(scale);
-    constexpr bool is_bf16 = std::is_same<scalar_t, c10::BFloat16>::value;
-
-    for (int32_t i = 0; i < q_num; ++i) {
-      for (int32_t h = 0; h < q_heads_per_kv; ++h) {
-        scalar_t* curr_src = src + i * q_num_stride + h * q_head_stride;
-        float* curr_dst =
-            q_buffer + i * q_heads_per_kv * head_dim + h * head_dim;
-
-        int32_t d = 0;
-        for (; d <= head_dim - 8; d += 8) {
-          __vector float v0, v1;
-          load_row8_B_as_f32<scalar_t>(curr_src + d, v0, v1);
-
-          v0 = vec_mul(v0, scale_vec);
-          v1 = vec_mul(v1, scale_vec);
-
-          vec_xst(v0, 0, curr_dst + d);
-          vec_xst(v1, 0, curr_dst + d + 4);
-        }
-
-        for (; d < head_dim; ++d) {
-          float val = static_cast<float>(curr_src[d]);
-          curr_dst[d] = val * scale;
-        }
-      }
-    }
-  }
-
-  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) {
-    // k_inv and v_inv are unused on VSX: FP8 KV cache is not supported on
-    // PowerPC. The parameters are present to match the common interface.
-#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_src = key + token_idx * key_token_num_stride +
-                                    head_idx * key_head_num_stride;
-          scalar_t* key_dst = key_cache + block_idx * num_blocks_stride +
-                              head_idx * cache_head_num_stride + block_offset;
-
-          for (int64_t i = 0, j = 0; i < head_dim; ++i, j += block_size) {
-            key_dst[j] = key_src[i];
-          }
-        }
-
-        {
-          const scalar_t* val_src = value + token_idx * value_token_num_stride +
-                                    head_idx * value_head_num_stride;
-          scalar_t* val_dst = value_cache + block_idx * num_blocks_stride +
-                              head_idx * cache_head_num_stride +
-                              block_offset * head_dim;
-
-          std::memcpy(val_dst, val_src, sizeof(scalar_t) * head_dim);
-        }
-      }
-    }
-  }
-};
-
-}  // namespace cpu_attention
-
-#undef BLOCK_SIZE_ALIGNMENT
-#undef HEAD_SIZE_ALIGNMENT
-#undef MAX_Q_HEAD_NUM_PER_ITER
-
-#endif  // CPU_ATTN_VSX_HPP

csrc/cpu/cpu_types_arm.hpp

@@ -486,10 +486,6 @@ struct FP32Vec16 : public VectorizedRegWrapper<FP32Vec16, 4, float> {
 
   explicit FP32Vec16(const BF16Vec8& v) : FP32Vec16(FP32Vec8(v)) {};
 
-  // FP8 stub: dead code on ARM (fp8 KV cache is x86-only), needed for
-  // load_b_pair_vec template to compile on all platforms.
-  explicit FP32Vec16(const BF16Vec32&, int) : Base() {}
-
   explicit FP32Vec16(const FP16Vec16& v) {
     reg.val[0] = Vectorized<float>(vcvt_f32_f16(vget_low_f16(v.reg.val[0])));
     reg.val[1] = Vectorized<float>(vcvt_f32_f16(vget_high_f16(v.reg.val[0])));

csrc/cpu/cpu_types_riscv_defs.hpp

@@ -72,7 +72,7 @@ typedef RVVTYPE(vuint16, LMUL_512, _t) fixed_u16x32_t
     __attribute__((riscv_rvv_vector_bits(512)));
 
 // 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/cpu_types_scalar.hpp

@@ -6,9 +6,6 @@
 
 namespace vec_op {
 
-struct fp8_e4m3_tag {};
-struct fp8_e5m2_tag {};
-
 #define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)            \
   AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__)    \
   AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__) \
@@ -148,9 +145,6 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
   }
 
   void save(void* ptr) const { *reinterpret_cast<f16x32_t*>(ptr) = reg; }
-
-  explicit BF16Vec32(const uint8_t*, fp8_e4m3_tag) : reg{} {}
-  explicit BF16Vec32(const uint8_t*, fp8_e5m2_tag) : reg{} {}
 };
 
 struct FP32Vec4 : public Vec<FP32Vec4> {
@@ -308,10 +302,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
 
   FP32Vec16(const BF16Vec8& v) : FP32Vec16(FP32Vec8(v)) {};
 
-  // FP8 stub: dead code on scalar path (fp8 KV cache is x86-only), needed for
-  // load_b_pair_vec template to compile on all platforms.
-  explicit FP32Vec16(const BF16Vec32&, int) : reg{} {}
-
   FP32Vec16 operator*(const FP32Vec16& b) const {
     f32x16_t ret;
     unroll_loop<int, VEC_ELEM_NUM>(

csrc/cpu/cpu_types_vsx.hpp

@@ -9,10 +9,6 @@
 
 namespace vec_op {
 
-// FP8 tag types for tag dispatch (see cpu_attn_vec.hpp)
-struct fp8_e4m3_tag {};
-struct fp8_e5m2_tag {};
-
 // FIXME: FP16 is not fully supported in Torch-CPU
 #define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
   AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
@@ -146,9 +142,6 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
       : reg({vec8_data.reg, vec8_data.reg, vec8_data.reg, vec8_data.reg}) {}
 
   void save(void* ptr) const { *reinterpret_cast<ss16x8x4_t*>(ptr) = reg; }
-
-  explicit BF16Vec32(const uint8_t*, fp8_e4m3_tag) : reg{} {}
-  explicit BF16Vec32(const uint8_t*, fp8_e5m2_tag) : reg{} {}
 };
 
 struct FP32Vec4 : public Vec<FP32Vec4> {
@@ -411,10 +404,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
 
   explicit FP32Vec16(const BF16Vec8& v) : FP32Vec16(FP32Vec8(v)) {}
 
-  // FP8 stub: dead code on PowerPC (fp8 KV cache is x86-only), needed for
-  // load_b_pair_vec template to compile on all platforms.
-  explicit FP32Vec16(const BF16Vec32&, int) : reg{} {}
-
   explicit FP32Vec16(const INT32Vec16& v) {
     reg.val[0] = vec_ctf(v.reg.val[0], 0);
     reg.val[1] = vec_ctf(v.reg.val[1], 0);

csrc/cpu/cpu_types_vxe.hpp

@@ -688,10 +688,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
 
   explicit FP32Vec16(const BF16Vec8& v) : FP32Vec16(FP32Vec8(v)) {}
 
-  // FP8 stub: dead code on s390x (fp8 KV cache is x86-only), needed for
-  // load_b_pair_vec template to compile on all platforms.
-  explicit FP32Vec16(const BF16Vec32&, int) : reg{} {}
-
   FP32Vec16 operator*(const FP32Vec16& b) const {
     return FP32Vec16(f32x4x4_t({vec_mul(reg.val[0], b.reg.val[0]),
                                 vec_mul(reg.val[1], b.reg.val[1]),

csrc/cpu/cpu_types_x86.hpp

@@ -122,17 +122,9 @@ struct FP16Vec16 : public Vec<FP16Vec16> {
   void save(void* ptr) const { _mm256_storeu_si256((__m256i*)ptr, reg); }
 
   void save(void* ptr, const int elem_num) const {
-#ifdef __AVX512BW__
     constexpr uint32_t M = 0xFFFFFFFF;
     __mmask16 mask = _cvtu32_mask16(M >> (32 - elem_num));
     _mm256_mask_storeu_epi16(ptr, mask, reg);
-#else
-    // Fallback for lack of 16-bit masked store
-    int16_t tmp[VEC_ELEM_NUM];
-    _mm256_storeu_si256((__m256i*)tmp, reg);
-    for (int i = 0; i < elem_num; ++i)
-      reinterpret_cast<int16_t*>(ptr)[i] = tmp[i];
-#endif
   }
 };
 
@@ -169,17 +161,9 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
   void save(void* ptr) const { _mm256_storeu_si256((__m256i*)ptr, reg); }
 
   void save(void* ptr, const int elem_num) const {
-#ifdef __AVX512BW__
     constexpr uint32_t M = 0xFFFFFFFF;
     __mmask16 mask = _cvtu32_mask16(M >> (32 - elem_num));
     _mm256_mask_storeu_epi16(ptr, mask, reg);
-#else
-    // Fallback for lack of 16-bit masked store
-    int16_t tmp[VEC_ELEM_NUM];
-    _mm256_storeu_si256((__m256i*)tmp, reg);
-    for (int i = 0; i < elem_num; ++i)
-      reinterpret_cast<int16_t*>(ptr)[i] = tmp[i];
-#endif
   }
 };
 
@@ -263,12 +247,13 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
   explicit BF16Vec32(__m256i low, __m256i high)
       : reg_low(low), reg_high(high) {}
 
-  explicit BF16Vec32()
-      : reg_low(_mm256_setzero_si256()), reg_high(_mm256_setzero_si256()) {}
-
   explicit BF16Vec32(BF16Vec8& vec8_data)
-      : reg_low(_mm256_broadcastsi128_si256((__m128i)vec8_data.reg)),
-        reg_high(_mm256_broadcastsi128_si256((__m128i)vec8_data.reg)) {}
+      : reg_low((__m256i)_mm256_inserti32x4(
+            _mm256_castsi128_si256((__m128i)vec8_data.reg),
+            (__m128i)vec8_data.reg, 1)),
+        reg_high((__m256i)_mm256_inserti32x4(
+            _mm256_castsi128_si256((__m128i)vec8_data.reg),
+            (__m128i)vec8_data.reg, 1)) {}
 
   // E4M3 decode (AVX2 path) — same bit-layout trick as the AVX512 variant
   // above.  Result = true_E4M3 * 2^-8; caller applies scale * 2^8.
@@ -689,11 +674,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
                      _mm256_sub_ps(reg_high, b.reg_high));
   }
 
-  FP32Vec16 operator-() const {
-    const __m256 neg = _mm256_set1_ps(-0.0f);
-    return FP32Vec16(_mm256_xor_ps(reg_low, neg), _mm256_xor_ps(reg_high, neg));
-  }
-
   FP32Vec16 operator/(const FP32Vec16& b) const {
     return FP32Vec16(_mm256_div_ps(reg_low, b.reg_low),
                      _mm256_div_ps(reg_high, b.reg_high));
@@ -759,85 +739,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
     _mm256_storeu_ps(ptr, reg_low);
     _mm256_storeu_ps(ptr + 8, reg_high);
   }
-
-  void save(float* ptr, const int elem_num) const {
-    // Partial store: cmpgt produces a sign-bit mask (0xFFFFFFFF/0 per lane)
-    // for the first elem_num lanes, applied across the two 8-wide halves.
-    if (elem_num <= 8) {
-      __m256i mask =
-          _mm256_cmpgt_epi32(_mm256_set1_epi32(elem_num),
-                             _mm256_setr_epi32(0, 1, 2, 3, 4, 5, 6, 7));
-      _mm256_maskstore_ps(ptr, mask, reg_low);
-    } else {
-      _mm256_storeu_ps(ptr, reg_low);
-      __m256i mask =
-          _mm256_cmpgt_epi32(_mm256_set1_epi32(elem_num - 8),
-                             _mm256_setr_epi32(0, 1, 2, 3, 4, 5, 6, 7));
-      _mm256_maskstore_ps(ptr + 8, mask, reg_high);
-    }
-  }
-
-  FP32Vec16 clamp(const FP32Vec16& min, const FP32Vec16& max) const {
-    return FP32Vec16(
-        _mm256_min_ps(max.reg_low, _mm256_max_ps(min.reg_low, reg_low)),
-        _mm256_min_ps(max.reg_high, _mm256_max_ps(min.reg_high, reg_high)));
-  }
-
-  FP32Vec16 abs() const {
-    const __m256 sign_mask = _mm256_set1_ps(-0.0f);
-    return FP32Vec16(_mm256_andnot_ps(sign_mask, reg_low),
-                     _mm256_andnot_ps(sign_mask, reg_high));
-  }
-
-  FP32Vec16 min(const FP32Vec16& b) const {
-    return FP32Vec16(_mm256_min_ps(reg_low, b.reg_low),
-                     _mm256_min_ps(reg_high, b.reg_high));
-  }
-
-  // Partial element-wise min over the first elem_num lanes only (tail path).
-  // Scalar via AliasReg: AVX2 has no masked vminps, so we spill, loop, reload.
-  FP32Vec16 min(const FP32Vec16& b, const int elem_num) const {
-    AliasReg ar_this_low, ar_this_high, ar_b_low, ar_b_high;
-    ar_this_low.reg = reg_low;
-    ar_this_high.reg = reg_high;
-    ar_b_low.reg = b.reg_low;
-    ar_b_high.reg = b.reg_high;
-    for (int i = 0; i < elem_num && i < 8; ++i)
-      ar_this_low.values[i] =
-          std::min(ar_this_low.values[i], ar_b_low.values[i]);
-    for (int i = 0; i < elem_num - 8 && i < 8; ++i)
-      ar_this_high.values[i] =
-          std::min(ar_this_high.values[i], ar_b_high.values[i]);
-    return FP32Vec16(ar_this_low.reg, ar_this_high.reg);
-  }
-
-  // Partial element-wise max over the first elem_num lanes only (tail path).
-  // Scalar via AliasReg: AVX2 has no masked vmaxps, so we spill, loop, reload.
-  FP32Vec16 max(const FP32Vec16& b, const int elem_num) const {
-    AliasReg ar_this_low, ar_this_high, ar_b_low, ar_b_high;
-    ar_this_low.reg = reg_low;
-    ar_this_high.reg = reg_high;
-    ar_b_low.reg = b.reg_low;
-    ar_b_high.reg = b.reg_high;
-    for (int i = 0; i < elem_num && i < 8; ++i)
-      ar_this_low.values[i] =
-          std::max(ar_this_low.values[i], ar_b_low.values[i]);
-    for (int i = 0; i < elem_num - 8 && i < 8; ++i)
-      ar_this_high.values[i] =
-          std::max(ar_this_high.values[i], ar_b_high.values[i]);
-    return FP32Vec16(ar_this_low.reg, ar_this_high.reg);
-  }
-
-  float reduce_min() const {
-    __m256 v = _mm256_min_ps(reg_low, reg_high);
-    __m256 v_shuffled = _mm256_permute_ps(v, 0b00001011);
-    __m256 v_min = _mm256_min_ps(v, v_shuffled);
-    v_shuffled = _mm256_permute_ps(v_min, 0b00000001);
-    v_min = _mm256_min_ps(v_min, v_shuffled);
-    v_shuffled = _mm256_permute2f128_ps(v_min, v_min, 0b00000001);
-    v_min = _mm256_min_ps(v_min, v_shuffled);
-    return _mm256_cvtss_f32(v_min);
-  }
 };
 #endif
 
@@ -890,34 +791,6 @@ struct INT8Vec64 : public Vec<INT8Vec64> {
   // non-temporal save
   void nt_save(int8_t* ptr) { _mm512_stream_si512((__m512i*)ptr, reg); }
 };
-#else
-struct INT8Vec16 : public Vec<INT8Vec16> {
-  constexpr static int VEC_ELEM_NUM = 16;
-  union AliasReg {
-    __m128i reg;
-    int8_t values[VEC_ELEM_NUM];
-  };
-
-  __m128i reg;
-
-  explicit INT8Vec16(const FP32Vec16& vec) {
-    __m256i lo_i32 = _mm256_cvtps_epi32(vec.reg_low);
-    __m256i hi_i32 = _mm256_cvtps_epi32(vec.reg_high);
-    __m256i packed16 = _mm256_packs_epi32(lo_i32, hi_i32);
-    packed16 = _mm256_permute4x64_epi64(packed16, 0xD8);
-    __m256i packed8 = _mm256_packs_epi16(packed16, _mm256_setzero_si256());
-    packed8 = _mm256_permute4x64_epi64(packed8, 0xD8);
-    reg = _mm256_castsi256_si128(packed8);
-  }
-
-  void save(int8_t* ptr) const { _mm_storeu_si128((__m128i*)ptr, reg); }
-
-  void save(int8_t* ptr, const int elem_num) const {
-    AliasReg ar;
-    ar.reg = reg;
-    for (int i = 0; i < elem_num; ++i) ptr[i] = ar.values[i];
-  }
-};
 #endif
 
 template <typename T>

csrc/cpu/dnnl_kernels.cpp

@@ -215,7 +215,7 @@ void dynamic_quant_epilogue(const float* input, scalar_t* output,
         float zp_scale_val = a_scale[i] * static_cast<float>(azp[i]);
         token_zp_scale_vec = cvt_vec_t(zp_scale_val);
       }
-      for (; j < hidden_size - vec_elem_num; j += vec_elem_num) {
+      for (; j < hidden_size - vec_elem_num; ++j) {
         cvt_vec_t elems_fp32(input_ptr + j);
         elems_fp32 = elems_fp32 * token_scale_vec;
         if constexpr (AZP) {

csrc/cpu/generate_cpu_attn_dispatch.py

@@ -20,8 +20,6 @@ ISA_TYPES = {
     "VEC16": 2,
     "NEON": 3,
     "VXE": 4,
-    "RVV": 5,
-    "VSX": 6,
 }
 
 # KV cache index: 0 = auto (same as scalar_t), 1 = fp8_e4m3, 2 = fp8_e5m2
@@ -39,7 +37,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"]
 
 # ISAs supported for head_dims divisible by 16 only
 ISA_FOR_16 = ["VEC16"]
@@ -150,19 +148,6 @@ def generate_header_file() -> str:
   #include "cpu_attn_vxe.hpp"
 #endif
 
-// cpu_attn_rvv.hpp is hardcoded to VLEN==128 (m1/m2 intrinsics, vl=8) and
-// itself includes <riscv_vector.h>, which is unavailable on scalar
-// (-march=rv64gc) builds. Gate the include the same way as the dispatch
-// macro below, so non-128 / scalar RISC-V builds skip it entirely.
-#if defined(__riscv) && defined(__riscv_v_min_vlen) && \
-    __riscv_v_min_vlen == 128
-  #include "cpu_attn_rvv.hpp"
-#endif
-
-#ifdef __powerpc__
-  #include "cpu_attn_vsx.hpp"
-#endif
-
 """
 
     header += generate_helper_function()
@@ -222,28 +207,6 @@ def generate_header_file() -> str:
         ["VXE", "VEC", "VEC16"],
         fp8=False,
     )
-    # RISC-V with RVV.  cpu_attn_rvv.hpp is hardcoded to VLEN==128
-    # (riscv_rvv_vector_bits(128) typedefs + vl=8 m1/m2 intrinsics), so
-    # we split the dispatch into two top-level branches: VLEN==128 builds
-    # get the full RVV+VEC+VEC16 case set, other VLEN builds get a
-    # VEC/VEC16-only fallback.  Preprocessor directives cannot appear
-    # inside a #define body, so this duplication is necessary.
-    header += _macro_block(
-        "#elif defined(__riscv) && defined(__riscv_v_min_vlen) "
-        "&& __riscv_v_min_vlen == 128",
-        ["RVV", "VEC", "VEC16"],
-        fp8=False,
-    )
-    header += _macro_block(
-        "#elif defined(__riscv)",
-        ["VEC", "VEC16"],
-        fp8=False,
-    )
-    header += _macro_block(
-        "#elif defined(__powerpc__)",
-        ["VSX", "VEC", "VEC16"],
-        fp8=False,
-    )
     header += _macro_block(
         "#elif defined(__AVX512F__)",
         ["VEC", "VEC16"],
@@ -260,8 +223,7 @@ def generate_header_file() -> str:
         fp8=False,
     )
     header += (
-        "#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / __s390x__ /"
-        " __riscv / __powerpc__ */\n\n"
+        "#endif  /* CPU_CAPABILITY_AMXBF16 / __aarch64__ / __s390x__ */\n\n"
         "#endif  // CPU_ATTN_DISPATCH_GENERATED_H\n"
     )
 

csrc/cpu/sgl-kernels/common.h

@@ -1,12 +1,13 @@
 // Adapted from
 // https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
 
-// clang-format off
-
 #pragma once
 
 #include <ATen/ATen.h>
 #include <ATen/Parallel.h>
+#include <ATen/record_function.h>
+
+// clang-format off
 
 #if defined(_OPENMP)
 #include <omp.h>
@@ -15,157 +16,40 @@
 namespace {
 
 // dispatch bool
-#define AT_DISPATCH_BOOL(BOOL_V, BOOL_NAME, ...) \
-  [&] {                                          \
-    if (BOOL_V) {                                \
-      constexpr bool BOOL_NAME = true;           \
-      return __VA_ARGS__();                      \
-    } else {                                     \
-      constexpr bool BOOL_NAME = false;          \
-      return __VA_ARGS__();                      \
-    }                                            \
+#define AT_DISPATCH_BOOL(BOOL_V, BOOL_NAME, ...)                                 \
+  [&] {                                                                          \
+    if (BOOL_V) {                                                                \
+      constexpr bool BOOL_NAME = true;                                           \
+      return __VA_ARGS__();                                                      \
+    } else {                                                                     \
+      constexpr bool BOOL_NAME = false;                                          \
+      return __VA_ARGS__();                                                      \
+    }                                                                            \
   }()
 
-#define AT_DISPATCH_BOOL2(BOOL_V1, BOOL_NAME1, BOOL_V2, BOOL_NAME2, ...) \
-  [&] {                                                                  \
-    if (BOOL_V1) {                                                       \
-      constexpr bool BOOL_NAME1 = true;                                  \
-      if (BOOL_V2) {                                                     \
-        constexpr bool BOOL_NAME2 = true;                                \
-        return __VA_ARGS__();                                            \
-      } else {                                                           \
-        constexpr bool BOOL_NAME2 = false;                               \
-        return __VA_ARGS__();                                            \
-      }                                                                  \
-    } else {                                                             \
-      constexpr bool BOOL_NAME1 = false;                                 \
-      if (BOOL_V2) {                                                     \
-        constexpr bool BOOL_NAME2 = true;                                \
-        return __VA_ARGS__();                                            \
-      } else {                                                           \
-        constexpr bool BOOL_NAME2 = false;                               \
-        return __VA_ARGS__();                                            \
-      }                                                                  \
-    }                                                                    \
-  }()
-
-// dispatch: bfloat16, float16, int8_t, fp8_e4m3, uint8_t(mxfp4/int4)
-#define CPU_DISPATCH_PACKED_TYPES(TYPE, ...)                     \
-  [&] {                                                          \
-    switch (TYPE) {                                              \
-      case at::ScalarType::BFloat16: {                           \
-        using packed_t = at::BFloat16;                           \
-        return __VA_ARGS__();                                    \
-      }                                                          \
-      case at::ScalarType::Half: {                               \
-        using packed_t = at::Half;                               \
-        return __VA_ARGS__();                                    \
-      }                                                          \
-      case at::ScalarType::Char: {                               \
-        using packed_t = int8_t;                                 \
-        return __VA_ARGS__();                                    \
-      }                                                          \
-      case at::ScalarType::Float8_e4m3fn: {                      \
-        using packed_t = at::Float8_e4m3fn;                      \
-        return __VA_ARGS__();                                    \
-      }                                                          \
-      case at::ScalarType::Byte: {                               \
-        using packed_t = uint8_t;                                \
-        return __VA_ARGS__();                                    \
-      }                                                          \
-      default:                                                   \
-        TORCH_CHECK(false, "Unsupported floating data type.\n"); \
-    }                                                            \
-  }()
-
-// Helper MICRO for CPU_DISPATCH_FLOATING_TYPES_EXT:
-//   TYPE1: the primary dtype (input, output, weight);
-//   TYPE2: defined as PARAM_T input
-#define CPU_DISPATCH_TYPE1_WITH_PARAM(TYPE1, PARAM_T, ...)   \
-  switch (TYPE1) {                                           \
-    case at::ScalarType::BFloat16: {                         \
-      using scalar_t = at::BFloat16;                         \
-      using param_t = PARAM_T;                               \
-      return __VA_ARGS__();                                  \
-    }                                                        \
-    case at::ScalarType::Half: {                             \
-      using scalar_t = at::Half;                             \
-      using param_t = PARAM_T;                               \
-      return __VA_ARGS__();                                  \
-    }                                                        \
-    case at::ScalarType::Float: {                            \
-      using scalar_t = float;                                \
-      using param_t = PARAM_T;                               \
-      return __VA_ARGS__();                                  \
-    }                                                        \
-    default:                                                 \
-      TORCH_CHECK(false, "Unsupported floating data type."); \
-  }
-
-// Helper MICRO for CPU_DISPATCH_REDUCED_FLOATING_TYPES_EXT:
-//   TYPE1: the primary dtype (input, output, weight);
-//   TYPE2: defined as PARAM_T input
-#define CPU_DISPATCH_TYPE1_WITH_PARAM_REDUCED(TYPE1, PARAM_T, ...) \
-  switch (TYPE1) {                                                 \
-    case at::ScalarType::BFloat16: {                               \
-      using scalar_t = at::BFloat16;                               \
-      using param_t = PARAM_T;                                     \
-      return __VA_ARGS__();                                        \
-    }                                                              \
-    case at::ScalarType::Half: {                                   \
-      using scalar_t = at::Half;                                   \
-      using param_t = PARAM_T;                                     \
-      return __VA_ARGS__();                                        \
-    }                                                              \
-    default:                                                       \
-      TORCH_CHECK(false, "Unsupported floating data type.");       \
-  }
-
-// Helper MICRO for CPU_DISPATCH_REDUCED_FLOATING_TYPES_EXT:
-//   TYPE1: the dtype both for scalar_t and param_t
-#define CPU_DISPATCH_TYPE1_WITH_SAME_PARAM_REDUCED(TYPE1, ...)       \
-  switch (TYPE1) {                                                   \
-    case at::ScalarType::BFloat16: {                                 \
-      using scalar_t = at::BFloat16;                                 \
-      using param_t = at::BFloat16;                                  \
-      return __VA_ARGS__();                                          \
-    }                                                                \
-    case at::ScalarType::Half: {                                     \
-      using scalar_t = at::Half;                                     \
-      using param_t = at::Half;                                      \
-      return __VA_ARGS__();                                          \
-    }                                                                \
-    default:                                                         \
-      TORCH_CHECK(false, "Unsupported reduced floating data type."); \
-  }
-
-// dispatch with mixed dtypes (TYPE1, TYPE2):
-//   TYPE1: the primary dtype (input, output, weight);
-//   TYPE2: the secondary dtype (bias, etc.).
-#define CPU_DISPATCH_FLOATING_TYPES_EXT(TYPE1, TYPE2, ...)            \
-  [&] {                                                               \
-    if (TYPE2 == at::kFloat) {                                        \
-      CPU_DISPATCH_TYPE1_WITH_PARAM(TYPE1, float, __VA_ARGS__)        \
-    } else if (TYPE2 == at::ScalarType::BFloat16) {                   \
-      CPU_DISPATCH_TYPE1_WITH_PARAM(TYPE1, at::BFloat16, __VA_ARGS__) \
-    } else if (TYPE2 == at::ScalarType::Half) {                       \
-      CPU_DISPATCH_TYPE1_WITH_PARAM(TYPE1, at::Half, __VA_ARGS__)     \
-    } else {                                                          \
-      TORCH_CHECK(false, "Unsupported floating data type.");          \
-    }                                                                 \
-  }()
-
-// dispatch with mixed dtypes (reduced one, no float for TYPE1) (TYPE1, TYPE2):
-//   TYPE1: the primary dtype (input, output, weight);
-//   TYPE2: the secondary dtype (bias, etc.).
-#define CPU_DISPATCH_REDUCED_FLOATING_TYPES_EXT(TYPE1, TYPE2, ...)     \
-  [&] {                                                                \
-    if (TYPE2 == at::kFloat) {                                         \
-      CPU_DISPATCH_TYPE1_WITH_PARAM_REDUCED(TYPE1, float, __VA_ARGS__) \
-    } else {                                                           \
-      TORCH_CHECK(TYPE1 == TYPE2);                                     \
-      CPU_DISPATCH_TYPE1_WITH_SAME_PARAM_REDUCED(TYPE1, __VA_ARGS__)   \
-    }                                                                  \
+// dispatch: bfloat16, float16, int8_t, fp8_e4m3
+#define CPU_DISPATCH_PACKED_TYPES(TYPE, ...)                                    \
+  [&] {                                                                         \
+    switch (TYPE) {                                                             \
+      case at::ScalarType::BFloat16 : {                                         \
+        using packed_t = at::BFloat16;                                          \
+        return __VA_ARGS__();                                                   \
+      }                                                                         \
+      case at::ScalarType::Half: {                                              \
+        using packed_t = at::Half;                                              \
+        return __VA_ARGS__();                                                   \
+      }                                                                         \
+      case at::ScalarType::Char : {                                             \
+        using packed_t = int8_t;                                                \
+        return __VA_ARGS__();                                                   \
+      }                                                                         \
+      case at::ScalarType::Float8_e4m3fn : {                                    \
+        using packed_t = at::Float8_e4m3fn;                                     \
+        return __VA_ARGS__();                                                   \
+      }                                                                         \
+      default:                                                                  \
+        TORCH_CHECK(false, "Unsupported floating data type.\n");                \
+    }                                                                           \
   }()
 
 #define UNUSED(x) (void)(x)
@@ -186,51 +70,13 @@ namespace {
 #define CHECK_DIM(d, x) TORCH_CHECK(x.dim() == d, #x " must be a " #d "D tensor")
 
 #define CHECK_EQ(a, b) TORCH_CHECK((a) == (b), "CHECK_EQ(" #a ", " #b ") failed. ", a, " vs ", b)
-#define CHECK_GT(a, b) TORCH_CHECK((a) > (b), "CHECK_GT(" #a ", " #b ") failed. ", a, " vs ", b)
-#define CHECK_GE(a, b) TORCH_CHECK((a) >= (b), "CHECK_GE(" #a ", " #b ") failed. ", a, " vs ", b)
 
-template <bool is_only_lastdim_contiguous>
-static inline void CHECK_INPUT_SHAPE_DTYPE(const at::Tensor& tensor, const at::IntArrayRef sizes, at::ScalarType st) {
-  TORCH_CHECK(tensor.sizes() == sizes, "Input tensor shape mismatch: expected ", sizes, ", got ", tensor.sizes());
-  TORCH_CHECK(tensor.scalar_type() == st, "Input tensor dtype mismatch");
-  if constexpr (is_only_lastdim_contiguous) {
-    CHECK_LAST_DIM_CONTIGUOUS_INPUT(tensor);
-  } else {
-    CHECK_INPUT(tensor);
-  }
-}
-
-// [NB] Parallel Routines
-//
-//  * at::parallel_for - applies for most of generic use cases, this will be compiled
-//                       against openmp in default torch release.
-//
-//  * parallel_for     - same function as above, can choose payload partition scheme in
-//                       balance211.
-//
-//  * parallel_2d      - parallel for 2 dimensions, used in GEMM, etc.
-//                       this one will do payload balance across 2 dimensions.
-//
-
-// grain size for each thread
+// parallel routines
 constexpr int GRAIN_SIZE = 1024;
 
 template <typename T, typename std::enable_if<std::is_integral<T>::value, int>::type = 0>
-inline T div_up(T x, T y) {
-  return (x + y - 1) / y;
-}
+inline T div_up(T x, T y) { return (x + y - 1) / y; }
 
-// you can only use at::get_thread_num() with at::parallel_for()
-// as it is lazy initialized, otherwise it will always return 0.
-inline int get_thread_num() {
-#if defined(_OPENMP)
-  return omp_get_thread_num();
-#else
-  return 0;
-#endif
-}
-
-// balance payload across each thread
 template <typename T>
 inline void balance211(T n, T nth, T ith, T& n_start, T& n_end) {
 #if 0
@@ -248,10 +94,10 @@ inline void balance211(T n, T nth, T ith, T& n_start, T& n_end) {
     }
     n_end += n_start;
 #else
-  // pytorch aten partition pattern
-  T n_my = div_up(n, nth);
-  n_start = ith * n_my;
-  n_end = std::min(n_start + n_my, n);
+    // pytorch aten partition pattern
+    T n_my = div_up(n, nth);
+    n_start = ith * n_my;
+    n_end = std::min(n_start + n_my, n);
 #endif
 }
 
@@ -259,15 +105,23 @@ template <typename func_t>
 inline void parallel_for(int n, const func_t& f) {
 #if defined(_OPENMP)
 #pragma omp parallel
-  {
+{
     int nth = omp_get_num_threads();
     int ith = omp_get_thread_num();
     int tbegin, tend;
     balance211(n, nth, ith, tbegin, tend);
     f(tbegin, tend);
-  }
+}
 #else
-  f(0, n);
+    f(0, n);
+#endif
+}
+
+inline int get_thread_num() {
+#if defined(_OPENMP)
+  return omp_get_thread_num();
+#else
+  return 0;
 #endif
 }
 
@@ -283,6 +137,7 @@ int inline adjust_num_threads(int m) {
 
 template <typename func_t>
 inline void parallel_2d(int m, int n, const func_t& f) {
+
   // make sure we have even num_threads
   int nth = adjust_num_threads(m);
 
@@ -310,66 +165,31 @@ inline void parallel_2d(int m, int n, const func_t& f) {
 
 #if defined(_OPENMP)
 #pragma omp parallel num_threads(nth)
-  {
-    int ith = omp_get_thread_num();
-    int ith_m = ith / nth_n;
-    int ith_n = ith % nth_n;
+{
+  int ith = omp_get_thread_num();
+  int ith_m = ith / nth_n;
+  int ith_n = ith % nth_n;
 
-    int thread_block_m = div_up(m, nth_m);
-    int thread_block_n = div_up(n, nth_n);
+  int thread_block_m = div_up(m, nth_m);
+  int thread_block_n = div_up(n, nth_n);
 
-    int begin_m = ith_m * thread_block_m;
-    int end_m = std::min(m, begin_m + thread_block_m);
-    int begin_n = ith_n * thread_block_n;
-    int end_n = std::min(n, begin_n + thread_block_n);
+  int begin_m = ith_m * thread_block_m;
+  int end_m = std::min(m, begin_m + thread_block_m);
+  int begin_n = ith_n * thread_block_n;
+  int end_n = std::min(n, begin_n + thread_block_n);
 
-    f(begin_m, end_m, begin_n, end_n);
-  }
+  f(begin_m, end_m, begin_n, end_n);
+}
 #else
   f(0, m, 0, n);
 #endif
 }
 
-// limit max cache blocks
-// when we need to do pre-unpack for weights, e.g. fp8
-#define MAX_CACHE_BLOCK_SIZE 4
-
 template <typename T>
-inline int get_cache_blocks(int chunk_size) {
+int get_cache_blocks(int BLOCK_SIZE, int K) {
   // L2 2MB and ratio of 50%
   const int L2_size = 2048 * 1024 >> 1;
-  return std::max(1, int(L2_size / (chunk_size * sizeof(T))));
-}
-
-template <>
-inline int get_cache_blocks<at::Float8_e4m3fn>(int chunk_size) {
-  // fp8 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);
-}
-
-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) {
-  // get number of blocks for L2 in most inner loop
-  int64_t cache_blocks_nb = get_cache_blocks<T>(chunk_size);
-
-  // loop order: [NB / cache_blocks_nb, MB, cache_blocks_nb]
-  // TODO: implement reverse order of [MB / cache_blocks_mb, NB, cache_blocks_mb]
-  for (int64_t nbb = nb0; nbb < nb1; nbb += cache_blocks_nb) {
-    for (int64_t mb = mb0; mb < mb1; ++mb) {
-      for (int64_t nb = nbb; nb < std::min(nbb + cache_blocks_nb, nb1); ++nb) {
-        f(mb, nb, nb - nbb);
-      }
-    }
-  }
+  return std::max(1, int(L2_size / (BLOCK_SIZE * K * sizeof(T))));
 }
 
 // data indexing for dimension collapse
@@ -423,10 +243,4 @@ struct Unroll<1> {
   }
 };
 
-// conditional data ptr for optional tensor
-template <typename T>
-inline T* conditional_data_ptr(const std::optional<at::Tensor>& opt) {
-  return opt.has_value() ? opt.value().data_ptr<T>() : nullptr;
-}
-
-}  // anonymous namespace
+} // anonymous namespace

csrc/cpu/sgl-kernels/conv.cpp

@@ -1,720 +0,0 @@
-// Adapted from
-// https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
-
-// clang-format off
-
-#include "common.h"
-#include "gemm.h"
-#include "vec.h"
-
-namespace {
-
-template <typename scalar_t>
-inline void copy_stub(scalar_t* __restrict__ y, const scalar_t* __restrict__ x, int64_t size) {
-  using Vec = at::vec::Vectorized<scalar_t>;
-  const bool is_padding = (x == nullptr);
-  for (int64_t d = 0; d < size; d += Vec::size()) {
-    Vec data_vec = is_padding ? Vec(0.f) : Vec::loadu(x + d);
-    data_vec.store(y + d);
-  }
-}
-
-// no remainder
-template <typename scalar_t>
-void inline update_conv_state(
-    scalar_t* __restrict__ conv_states,
-    const scalar_t* __restrict__ input,
-    int64_t width,
-    int64_t dim,
-    int64_t seqlen,
-    bool has_initial_states) {
-  // width for `conv_states`
-  int64_t width1 = width - 1;
-  int64_t w = 0;
-  for (; w < width1 - seqlen; ++w) {
-    scalar_t* y = conv_states + w * dim;
-    const scalar_t* x = has_initial_states ? conv_states + (w + seqlen) * dim : nullptr;
-    copy_stub(y, x, dim);
-  }
-  for (; w < width1; ++w) {
-    scalar_t* y = conv_states + w * dim;
-    const scalar_t* x = input + (w + seqlen - width1) * dim;
-    copy_stub(y, x, dim);
-  }
-}
-
-// A : [M, BLOCK_N]
-// B : [BLOCK_N, K], prepacked as [K/2, BLOCK_N, 2]
-// C : [M, BLOCK_N]
-// bias : [BLOCK_N]
-//
-// lda : leading dimension of `input` and `out`
-//
-template <typename scalar_t, int K, int BLOCK_N, bool has_bias, bool has_silu>
-struct tinygemm_kernel {
-  static inline void apply(
-      const scalar_t* __restrict__ A,
-      const scalar_t* __restrict__ B,
-      scalar_t* __restrict__ C,
-      const scalar_t* __restrict__ bias,
-      const scalar_t* __restrict__ conv_states,
-      bool has_initial_state,
-      int64_t M,
-      int64_t lda,
-      bool is_first_token) {
-    TORCH_CHECK(false, "tinygemm_kernel_nn: scalar path not implemented!");
-  }
-};
-
-#if defined(CPU_CAPABILITY_AVX512)
-template <int K, int BLOCK_N, bool has_bias, bool has_silu>
-struct tinygemm_kernel<at::BFloat16, K, BLOCK_N, has_bias, has_silu> {
-  static inline void apply(
-      const at::BFloat16* __restrict__ A,
-      const at::BFloat16* __restrict__ B,
-      at::BFloat16* __restrict__ C,
-      const at::BFloat16* __restrict__ bias,
-      const at::BFloat16* __restrict__ conv_states,
-      bool has_initial_state,
-      int64_t M,
-      int64_t lda,
-      bool is_first_token) {
-    assert(K == 4);
-    constexpr int ROWS = K;
-    constexpr int COLS = BLOCK_N / block_size_n();
-
-    // leading dimension size for b for next block [K/2, 32, 2]
-    constexpr int ldb = block_size_n() * K;
-
-    __m512bh va[ROWS * COLS];
-    __m512bh vb[ROWS * COLS];
-    __m512 vc[COLS * 2];
-
-    // k: {-3, -2, -1} -> {0, 1, 2}
-    auto set_conv_states = [&](int k, int col) -> __m512i {
-      return has_initial_state ? _mm512_loadu_si512(conv_states + (k + K - 1) * lda + col * 32)
-                               : _mm512_setzero_si512();
-    };
-
-#define MM512_LOAD_A(idx)                                                 \
-  ((idx) < 0 && is_first_token) ? (__m512bh)(set_conv_states((idx), col)) \
-                                : (__m512bh)(_mm512_loadu_si512(A + (idx) * lda + col * 32))
-
-#define MM512_PACK_A(ap, bp, a, b)                       \
-  do {                                                   \
-    __m512i r0 = (__m512i)(a);                           \
-    __m512i r1 = (__m512i)(b);                           \
-    __m512i d0 = _mm512_unpacklo_epi16(r0, r1);          \
-    __m512i d1 = _mm512_unpackhi_epi16(r0, r1);          \
-    r0 = _mm512_shuffle_i32x4(d0, d1, 0x88);             \
-    r1 = _mm512_shuffle_i32x4(d0, d1, 0xdd);             \
-    (ap) = (__m512bh)_mm512_shuffle_i32x4(r0, r1, 0x88); \
-    (bp) = (__m512bh)_mm512_shuffle_i32x4(r0, r1, 0xdd); \
-  } while (0)
-
-    // step 0 : preload a at time step [-3][-2][-1]
-    auto preloada = [&](auto i) {
-      constexpr int col = i;
-      int64_t m = 0;
-      va[1 * COLS + col] = MM512_LOAD_A(m - 3);
-      va[2 * COLS + col] = MM512_LOAD_A(m - 2);
-      va[3 * COLS + col] = MM512_LOAD_A(m - 1);
-    };
-    Unroll<COLS>{}(preloada);
-
-    auto loada = [&](auto i, int64_t m) {
-      constexpr int col = i;
-      // update previous time step
-      va[0 * COLS + col] = va[1 * COLS + col];
-      va[1 * COLS + col] = va[2 * COLS + col];
-      va[2 * COLS + col] = va[3 * COLS + col];
-      // load current time step
-      va[3 * COLS + col] = MM512_LOAD_A(m);
-    };
-
-    // step 1 : load weight for just once
-    auto loadb = [&](auto i) {
-      constexpr int row = i / COLS;
-      constexpr int col = i % COLS;
-      vb[row * COLS + col] = (__m512bh)(_mm512_loadu_si512(B + col * ldb + row * 32));
-    };
-    Unroll<ROWS * COLS>{}(loadb);
-
-    // [NB] accumulates 4x32 bfloat16 blocks
-    //
-    //   +------------+------------+
-    //   |    col0    |    col1    |
-    //   +------------+------------+
-    //   |  va0  va1  |  va0  va1  |
-    //   |  va2  va3  |  va2  va3  |
-    //   +------------+------------+
-    //   |  vc0  vc1  |  vc0  vc1  |
-    //   +------------+------------+
-    //
-    //  * va and vb shares the same memory layout
-    //  * block_n 32 with 4 rows equals to 4 registers
-    //  * 37 uops with avx512bf16 v.s. 57 uops with avx512f
-    //
-    auto compute = [&](auto i) {
-      constexpr int col = i;
-
-      // init accumulators
-      if constexpr (has_bias) {
-        __m512i b16 = _mm512_loadu_si512(reinterpret_cast<const __m512i*>(bias + col * 32));
-        vc[col * 2 + 0] = CVT_BF16_TO_FP32(_mm512_extracti32x8_epi32(b16, 0));
-        vc[col * 2 + 1] = CVT_BF16_TO_FP32(_mm512_extracti32x8_epi32(b16, 1));
-      } else {
-        vc[col * 2 + 0] = _mm512_set1_ps(0.f);
-        vc[col * 2 + 1] = _mm512_set1_ps(0.f);
-      }
-
-      // convert to vnni2 format
-      __m512bh va0, va1, va2, va3;
-      MM512_PACK_A(va0, va1, va[0 * COLS + col], va[1 * COLS + col]);
-      MM512_PACK_A(va2, va3, va[2 * COLS + col], va[3 * COLS + col]);
-
-      // accumulate
-      vc[col * 2 + 0] = _mm512_dpbf16_ps(vc[col * 2 + 0], va0, vb[0 * COLS + col]);
-      vc[col * 2 + 0] = _mm512_dpbf16_ps(vc[col * 2 + 0], va2, vb[2 * COLS + col]);
-      vc[col * 2 + 1] = _mm512_dpbf16_ps(vc[col * 2 + 1], va1, vb[1 * COLS + col]);
-      vc[col * 2 + 1] = _mm512_dpbf16_ps(vc[col * 2 + 1], va3, vb[3 * COLS + col]);
-    };
-
-    using fVec = at::vec::Vectorized<float>;
-    using bVec = at::vec::Vectorized<at::BFloat16>;
-    const fVec one = fVec(1.f);
-    auto storec = [&](auto i, int64_t m) {
-      constexpr int col = i;
-      fVec x0 = fVec(vc[col * 2 + 0]);
-      fVec x1 = fVec(vc[col * 2 + 1]);
-      if constexpr (has_silu) {
-        x0 = x0 / (one + x0.neg().exp_u20());
-        x1 = x1 / (one + x1.neg().exp_u20());
-      }
-      bVec out_vec = convert_from_float_ext<at::BFloat16>(x0, x1);
-      out_vec.store(C + m * lda + col * 32);
-    };
-
-    for (int64_t m = 0; m < M; ++m) {
-      // step 3.a : load a at current time step
-      Unroll<COLS>{}(loada, m);
-      // step 3.b : accumulate for window size (4)
-      Unroll<COLS>{}(compute);
-      // step 3.c : store c at current time step
-      Unroll<COLS>{}(storec, m);
-    }
-  }
-};
-#endif
-
-#define LAUNCH_TINYGEMM_KERNEL(K, NB_SIZE)                                                   \
-  tinygemm_kernel<scalar_t, K, NB_SIZE, has_bias, has_silu>::apply(                          \
-      input + bs * seqlen * dim + mb_start * dim + nb_start,                                 \
-      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_initial_states_value,                                                              \
-      mb_size,                                                                               \
-      dim,                                                                                   \
-      mb_start == 0);
-
-template <typename scalar_t>
-void causal_conv1d_fwd_kernel_impl(
-    scalar_t* __restrict__ out,
-    const scalar_t* __restrict__ input,
-    const scalar_t* __restrict__ weight,
-    const scalar_t* __restrict__ bias,
-    scalar_t* __restrict__ conv_states,
-    const int32_t* __restrict__ conv_indices,
-    const bool* __restrict__ has_initial_state,
-    bool silu_activation,
-    int64_t batch,
-    int64_t dim,
-    int64_t seqlen,
-    int64_t width,
-    int64_t num_seq_blocks,
-    int64_t conv_state_slot_stride) {
-  // handle 32 x 64 per block
-  constexpr int64_t BLOCK_M = block_size_m();
-  constexpr int64_t BLOCK_N = block_size_n() * 2;
-  const int64_t NB = div_up(dim, BLOCK_N);
-
-  const int64_t num_blocks_per_seq = div_up(seqlen, BLOCK_M);
-  const bool has_conv_states = conv_states != nullptr;
-  const bool has_conv_indices = conv_indices != nullptr;
-
-  // parallel on [batch, seq, NB]
-  AT_DISPATCH_BOOL2(bias != nullptr, has_bias, silu_activation, has_silu, [&] {
-    at::parallel_for(0, num_seq_blocks * NB, 0, [&](int64_t begin, int64_t end) {
-      int64_t mb{0}, nb{0};
-      data_index_init(begin, mb, num_seq_blocks, nb, NB);
-
-      for (int64_t i = begin; i < end; ++i) {
-        int64_t bs = mb / num_blocks_per_seq;
-
-        int64_t mb_start = (mb % num_blocks_per_seq) * BLOCK_M;
-        int64_t mb_size = std::min(seqlen - mb_start, BLOCK_M);
-        int64_t nb_start = nb * BLOCK_N;
-        int64_t nb_size = std::min(dim - nb_start, BLOCK_N);
-
-        const bool has_initial_states_value = has_conv_states ? has_initial_state[bs] : false;
-        int32_t conv_state_index = has_conv_indices ? conv_indices[bs] : bs;
-
-        switch (width << 4 | nb_size >> 4) {
-          case 0x42:
-            LAUNCH_TINYGEMM_KERNEL(4, 32);
-            break;
-          case 0x44:
-            LAUNCH_TINYGEMM_KERNEL(4, 64);
-            break;
-          default:
-            TORCH_CHECK(false, "Unexpected block size, ", width, " x ", nb_size);
-        }
-
-        // move to the next index
-        data_index_step(mb, num_seq_blocks, nb, NB);
-      }
-    });
-  });
-
-  // update conv_states if necessary
-  if (has_conv_states) {
-    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]);
-      }
-    });
-  }
-}
-
-#define LAUNCH_TINYGEMM_VARLEN_KERNEL(K, NB_SIZE)                   \
-  tinygemm_kernel<scalar_t, K, NB_SIZE, has_bias, has_silu>::apply( \
-      input + batch_offset * dim + mb_start * dim + nb_start,       \
-      weight + nb_start * width,                                    \
-      out + batch_offset * dim + mb_start * dim + nb_start,         \
-      has_bias ? bias + nb_start : nullptr,                         \
-      nullptr,                                                      \
-      false,                                                        \
-      mb_size,                                                      \
-      dim,                                                          \
-      mb_start == 0);
-
-// TODO: add `has_initial_state` support for varlen kernel
-template <typename scalar_t>
-void causal_conv1d_fwd_varlen_kernel_impl(
-    scalar_t* __restrict__ out,
-    const scalar_t* __restrict__ input,
-    const scalar_t* __restrict__ weight,
-    const scalar_t* __restrict__ bias,
-    scalar_t* __restrict__ conv_states,
-    const int32_t* __restrict__ query_start_loc,
-    const int32_t* __restrict__ conv_indices,
-    const bool* __restrict__ has_initial_state,
-    const int32_t* __restrict__ block_indices,
-    bool silu_activation,
-    int64_t batch,
-    int64_t dim,
-    int64_t width,
-    int64_t num_seq_blocks,
-    int64_t conv_state_slot_stride) {
-  // handle 32 x 64 per block
-  constexpr int64_t BLOCK_M = block_size_m();
-  constexpr int64_t BLOCK_N = block_size_n() * 2;
-  const int64_t NB = div_up(dim, BLOCK_N);
-
-  const bool has_conv_states = conv_states != nullptr;
-  const bool has_conv_indices = conv_indices != nullptr;
-
-  // parallel on [batch, seq, NB]
-  AT_DISPATCH_BOOL2(bias != nullptr, has_bias, silu_activation, has_silu, [&] {
-    at::parallel_for(0, num_seq_blocks * NB, 0, [&](int64_t begin, int64_t end) {
-      int64_t mb{0}, nb{0};
-      data_index_init(begin, mb, num_seq_blocks, nb, NB);
-
-      for (int64_t i = begin; i < end; ++i) {
-        int32_t bs = block_indices[mb * 2 + 0];
-        int32_t batch_offset = query_start_loc[bs];
-        int32_t seqlen = query_start_loc[bs + 1] - query_start_loc[bs];
-
-        int64_t mb_start = block_indices[mb * 2 + 1] * BLOCK_M;
-        int64_t mb_size = std::min(seqlen - mb_start, BLOCK_M);
-        int64_t nb_start = nb * BLOCK_N;
-        int64_t nb_size = std::min(dim - nb_start, BLOCK_N);
-
-        switch (width << 4 | nb_size >> 4) {
-          case 0x42:
-            LAUNCH_TINYGEMM_VARLEN_KERNEL(4, 32);
-            break;
-          case 0x44:
-            LAUNCH_TINYGEMM_VARLEN_KERNEL(4, 64);
-            break;
-          default:
-            TORCH_CHECK(false, "Unexpected block size, ", width, " x ", nb_size);
-        }
-
-        // move to the next index
-        data_index_step(mb, num_seq_blocks, nb, NB);
-      }
-    });
-  });
-
-  // update conv_states if necessary
-  if (has_conv_states) {
-    at::parallel_for(0, batch, 0, [&](int64_t begin, int64_t end) {
-      for (int64_t bs = begin; bs < end; ++bs) {
-        int32_t conv_state_index = has_conv_indices ? conv_indices[bs] : bs;
-        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,
-            input + batch_offset * dim,
-            width,
-            dim,
-            seqlen,
-            /* has_initial_state */ false);
-      }
-    });
-  }
-}
-
-template <typename scalar_t>
-void causal_conv1d_update_kernel_impl(
-    scalar_t* __restrict__ out,
-    const scalar_t* __restrict__ input,
-    scalar_t* __restrict__ conv_states,
-    const scalar_t* __restrict__ weight,
-    const scalar_t* __restrict__ bias,
-    const int32_t* __restrict__ conv_indices,
-    bool silu_activation,
-    int64_t batch,
-    int64_t dim,
-    int64_t seqlen,
-    int64_t width,
-    int64_t conv_state_slot_stride) {
-  // handle 32 x 64 per block
-  constexpr int64_t BLOCK_M = block_size_m();
-  constexpr int64_t BLOCK_N = block_size_n() * 2;
-  const int64_t NB = div_up(dim, BLOCK_N);
-
-  const bool has_conv_states = conv_states != nullptr;
-  const bool has_conv_indices = conv_indices != nullptr;
-
-  // parallel on [batch, NB]
-  AT_DISPATCH_BOOL2(bias != nullptr, has_bias, silu_activation, has_silu, [&] {
-    at::parallel_for(0, batch * NB, 0, [&](int64_t begin, int64_t end) {
-      int64_t bs{0}, nb{0};
-      data_index_init(begin, bs, batch, nb, NB);
-
-      for (int64_t i = begin; i < end; ++i) {
-        int64_t mb_start = 0;
-        int64_t mb_size = 1;
-        int64_t nb_start = nb * BLOCK_N;
-        int64_t nb_size = std::min(dim - nb_start, BLOCK_N);
-
-        const bool has_initial_states_value = true;
-        int32_t conv_state_index = has_conv_indices ? conv_indices[bs] : bs;
-
-        switch (width << 4 | nb_size >> 4) {
-          case 0x42:
-            LAUNCH_TINYGEMM_KERNEL(4, 32);
-            break;
-          case 0x44:
-            LAUNCH_TINYGEMM_KERNEL(4, 64);
-            break;
-          default:
-            TORCH_CHECK(false, "Unexpected block size, ", width, " x ", nb_size);
-        }
-
-        // move to the next index
-        data_index_step(bs, batch, nb, NB);
-      }
-    });
-  });
-
-#define CONV_STATE_INDEXR(w) conv_states + conv_state_index*conv_state_slot_stride + (w) * dim
-
-  // update conv_states
-  at::parallel_for(0, batch, 0, [&](int64_t begin, int64_t end) {
-    for (int64_t bs = begin; bs < end; ++bs) {
-      // update old states, range [1, width - 1)
-      int32_t conv_state_index = has_conv_indices ? conv_indices[bs] : bs;
-      for (int64_t w = 1; w < width - 1; ++w) {
-        std::memcpy(CONV_STATE_INDEXR(w - 1), CONV_STATE_INDEXR(w), dim * sizeof(scalar_t));
-      }
-      // copy new states
-      std::memcpy(CONV_STATE_INDEXR(width - 2), input + bs * dim, dim * sizeof(scalar_t));
-    }
-  });
-}
-
-}  // anonymous namespace
-
-// from [dim, width] or [N, K]
-// to [N/BLOCK_N, K/2, BLOCK_N, 2]
-at::Tensor causal_conv1d_weight_pack(const at::Tensor& weight) {
-  CHECK_INPUT(weight);
-
-  int64_t dim = weight.size(0);
-  int64_t width = weight.size(1);
-  constexpr int64_t BLOCK_N = block_size_n();
-  TORCH_CHECK(width == 4, "causal_conv1d_weight_pack: support only width of 4");
-  TORCH_CHECK(dim % BLOCK_N == 0, "causal_conv1d_weight_pack: invalid dim size ", dim);
-
-  const int64_t N = dim, K2 = width >> 1;
-  const int64_t NB = div_up(N, BLOCK_N);
-
-  auto packed_weight = at::empty_like(weight);
-  AT_DISPATCH_REDUCED_FLOATING_TYPES(weight.scalar_type(), "causal_conv1d_fwd_kernel_impl", [&] {
-    // cast to float32 as vnni size is 2
-    const float* w_data = reinterpret_cast<float*>(weight.data_ptr<scalar_t>());
-    float* packed_data = reinterpret_cast<float*>(packed_weight.data_ptr<scalar_t>());
-
-    at::parallel_for(0, NB * K2 * BLOCK_N, 0, [&](int64_t begin, int64_t end) {
-      int64_t nb{0}, k2{0}, n{0};
-      data_index_init(begin, nb, NB, k2, K2, n, BLOCK_N);
-
-      // TODO: optimize this if we need to online prepacking.
-      for (int64_t i = begin; i < end; ++i) {
-        packed_data[i] = w_data[nb * BLOCK_N * K2 + n * K2 + k2];
-
-        // move to the next index
-        data_index_step(nb, NB, k2, K2, n, BLOCK_N);
-      }
-    });
-  });
-  return packed_weight;
-}
-
-#define CHECK_OPTIONAL_SHAPE_DTYPE(OPT, SIZE, DTYPE) \
-  if (OPT.has_value()) {                             \
-    const auto tensor = OPT.value();                 \
-    CHECK_CONTIGUOUS(tensor);                        \
-    CHECK_EQ(tensor.size(0), SIZE);                  \
-    CHECK_EQ(tensor.scalar_type(), DTYPE);           \
-  }
-
-template <int BLOCK_M>
-int64_t get_block_count(const std::optional<at::Tensor>& offsets, int64_t batch, int64_t seqlen) {
-  if (offsets.has_value()) {
-    const int32_t* offsets_data = offsets.value().data_ptr<int32_t>();
-    int32_t num_seq_blocks = 0;
-    for (int64_t row = 0; row < batch; ++row) {
-      num_seq_blocks += div_up(offsets_data[row + 1] - offsets_data[row], BLOCK_M);
-    }
-    return num_seq_blocks;
-  }
-  return batch * div_up(seqlen, int64_t(BLOCK_M));
-}
-
-template <int BLOCK_M>
-at::Tensor get_block_indices(const std::optional<at::Tensor>& offsets, int64_t num_seq_blocks) {
-  if (!offsets.has_value()) {
-    return at::Tensor();
-  }
-
-  const at::Tensor& offsets_ = offsets.value();
-  at::Tensor indices = at::empty({num_seq_blocks, 2}, offsets_.options());
-
-  int64_t batch = offsets_.size(0) - 1;
-
-  const int32_t* offsets_data = offsets_.data_ptr<int32_t>();
-  int32_t* indices_data = indices.data_ptr<int32_t>();
-
-  int64_t idx = 0;
-  for (int32_t row = 0; row < batch; ++row) {
-    int32_t blocks = div_up(offsets_data[row + 1] - offsets_data[row], BLOCK_M);
-
-    for (int32_t col = 0; col < blocks; ++col) {
-      indices_data[idx * 2 + 0] = row;
-      indices_data[idx * 2 + 1] = col;
-      idx++;
-    }
-  }
-  return indices;
-}
-
-// API aligned with GPUs
-//
-//   x: (batch, dim, seqlen) or (dim, cu_seq_len) for varlen
-//   weight: (dim, width)
-//   bias: (dim,)
-//   query_start_loc: (batch + 1) int32
-//   cache_indices: (batch)  int32
-//   has_initial_state: (batch) bool
-//   conv_states: (..., dim, width - 1) itype
-//   activation: either None or "silu" or "swish"
-//   pad_slot_id: int
-//
-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>& conv_state_indices,
-    const std::optional<at::Tensor>& has_initial_state,
-    bool silu_activation,
-    int64_t pad_slot_id,
-    bool is_vnni) {
-  CHECK_CONTIGUOUS(weight);
-  auto packed_w = is_vnni ? weight : causal_conv1d_weight_pack(weight);
-
-  const bool is_var_seqlen = query_start_loc.has_value();
-  const int64_t input_ndim = is_var_seqlen ? 2 : 3;
-  TORCH_CHECK(x.dim() == input_ndim, "causal_conv1d_fwd_cpu: expect x to be ", input_ndim, "D tensor.");
-  TORCH_CHECK(x.stride(-2) == 1 && x.stride(-1) == x.size(-2), "causal_conv1d_fwd_cpu: expect x to be transposed.");
-
-  const int64_t batch = is_var_seqlen ? query_start_loc.value().size(0) - 1 : x.size(0);
-  const int64_t dim = x.size(-2);
-  const int64_t seqlen = x.size(-1);
-  const int64_t width = weight.size(-1);
-
-  const auto scalar_type = x.scalar_type();
-  CHECK_EQ(weight.scalar_type(), scalar_type);
-  CHECK_OPTIONAL_SHAPE_DTYPE(bias, dim, scalar_type);
-  CHECK_OPTIONAL_SHAPE_DTYPE(query_start_loc, batch + 1, at::kInt);
-  CHECK_OPTIONAL_SHAPE_DTYPE(conv_state_indices, batch, at::kInt);
-  CHECK_OPTIONAL_SHAPE_DTYPE(has_initial_state, batch, at::kBool);
-
-  if (conv_states.has_value()) {
-    auto& conv_states_val = conv_states.value();
-    int64_t padded_batch = conv_states_val.size(0);
-    CHECK_EQ(conv_states_val.scalar_type(), scalar_type);
-    CHECK_GE(padded_batch, batch);
-    CHECK_EQ(conv_states_val.size(1), dim);
-    CHECK_EQ(conv_states_val.size(2), width - 1);
-
-    // adjust `conv_states` to be contiguous on `dim`
-    // should happen only once
-    if (conv_states_val.stride(-2) != 1) {
-      auto conv_states_copy = conv_states_val.clone();
-      conv_states_val.as_strided_({padded_batch, dim, width - 1}, {(width - 1) * dim, 1, dim});
-      conv_states_val.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);
-
-  // block size for sequence blocks, 32
-  constexpr int64_t BLOCK_M = block_size_m();
-
-  // total number of sequence blocks
-  int64_t num_seq_blocks = get_block_count<BLOCK_M>(query_start_loc, batch, seqlen);
-
-  at::Tensor out = at::empty_like(x);
-  AT_DISPATCH_REDUCED_FLOATING_TYPES(scalar_type, "causal_conv1d_fwd_kernel_impl", [&] {
-    if (is_var_seqlen) {
-      // record seq blocks in Coordinate format, aka [num_seq_blocks, 2]
-      at::Tensor block_indices = get_block_indices<BLOCK_M>(query_start_loc, num_seq_blocks);
-
-      causal_conv1d_fwd_varlen_kernel_impl(
-          out.data_ptr<scalar_t>(),
-          x.data_ptr<scalar_t>(),
-          packed_w.data_ptr<scalar_t>(),
-          conditional_data_ptr<scalar_t>(bias),
-          conditional_data_ptr<scalar_t>(conv_states),
-          conditional_data_ptr<int32_t>(query_start_loc),
-          conditional_data_ptr<int32_t>(conv_state_indices),
-          conditional_data_ptr<bool>(has_initial_state),
-          block_indices.data_ptr<int32_t>(),
-          silu_activation,
-          batch,
-          dim,
-          width,
-          num_seq_blocks,
-          conv_state_slot_stride);
-    } else {
-      causal_conv1d_fwd_kernel_impl<scalar_t>(
-          out.data_ptr<scalar_t>(),
-          x.data_ptr<scalar_t>(),
-          packed_w.data_ptr<scalar_t>(),
-          conditional_data_ptr<scalar_t>(bias),
-          conditional_data_ptr<scalar_t>(conv_states),
-          conditional_data_ptr<int32_t>(conv_state_indices),
-          conditional_data_ptr<bool>(has_initial_state),
-          silu_activation,
-          batch,
-          dim,
-          seqlen,
-          width,
-          num_seq_blocks,
-          conv_state_slot_stride);
-    }
-  });
-  return out;
-}
-
-// API aligned with GPUs
-//
-//   x: (batch, dim) or (batch, dim, seqlen)
-//   conv_state: (..., dim, state_len), where state_len >= width - 1
-//   weight: (dim, width)
-//   bias: (dim,)
-//   cache_seqlens: (batch,), dtype int32.
-//   conv_state_indices: (batch,), dtype int32
-//   pad_slot_id: int
-//   out: (batch, dim) or (batch, dim, seqlen)
-//
-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) {
-  CHECK_CONTIGUOUS(x);
-  CHECK_CONTIGUOUS(weight);
-  auto packed_w = is_vnni ? weight : causal_conv1d_weight_pack(weight);
-
-  // TODO: add multi-token prediction support
-  TORCH_CHECK(x.dim() == 2, "causal_conv1d_update_cpu: expect x to be 2D tensor.");
-  TORCH_CHECK(!cache_seqlens.has_value(), "causal_conv1d_update_cpu: don't support cache_seqlens.");
-
-  int64_t batch = x.size(0);
-  int64_t dim = x.size(1);
-  int64_t seqlen = 1;
-  int64_t width = weight.size(-1);
-
-  const auto scalar_type = x.scalar_type();
-  CHECK_EQ(weight.scalar_type(), scalar_type);
-  CHECK_OPTIONAL_SHAPE_DTYPE(bias, dim, scalar_type);
-  CHECK_OPTIONAL_SHAPE_DTYPE(conv_state_indices, batch, at::kInt);
-
-  CHECK_EQ(conv_states.scalar_type(), scalar_type);
-  CHECK_EQ(conv_states.size(1), dim);
-  CHECK_EQ(conv_states.size(2), width - 1);
-
-  // adjust `conv_states` to be contiguous on `dim`
-  if (conv_states.stride(-2) != 1) {
-    int64_t num_cache_lines = conv_states.size(0);
-    auto conv_states_copy = conv_states.clone();
-    conv_states.as_strided_({num_cache_lines, dim, width - 1}, {(width - 1) * dim, 1, dim});
-    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>(
-        out.data_ptr<scalar_t>(),
-        x.data_ptr<scalar_t>(),
-        conv_states.data_ptr<scalar_t>(),
-        packed_w.data_ptr<scalar_t>(),
-        conditional_data_ptr<scalar_t>(bias),
-        conditional_data_ptr<int32_t>(conv_state_indices),
-        silu_activation,
-        batch,
-        dim,
-        seqlen,
-        width,
-        conv_state_slot_stride);
-  });
-  return out;
-}

csrc/cpu/sgl-kernels/gemm.cpp

@@ -1,12 +1,11 @@
 // Adapted from
 // https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
 
-// clang-format off
-
-#include "gemm.h"
-
 #include "common.h"
 #include "vec.h"
+#include "gemm.h"
+
+// clang-format off
 
 namespace {
 
@@ -27,13 +26,13 @@ inline void s8s8_compensation(int8_t* __restrict__ packed, int K) {
   const __m512i off = _mm512_set1_epi8(static_cast<char>(0x80));
   for (int k = 0; k < K / 4; ++k) {
     for (int col = 0; col < COLS; ++col) {
-      __m512i vb = _mm512_loadu_si512((const __m512i*)(packed + k * BLOCK_N * 4 + col * 64));
+      __m512i vb = _mm512_loadu_si512((const __m512i *)(packed + k * BLOCK_N * 4 + col * 64));
       vcomp[col] = _mm512_dpbusd_epi32(vcomp[col], off, vb);
     }
   }
 
   for (int col = 0; col < COLS; ++col) {
-    _mm512_storeu_si512((__m512i*)(packed + offset + col * 64), vcomp[col]);
+    _mm512_storeu_si512((__m512i *)(packed + offset + col * 64), vcomp[col]);
   }
 #else
   TORCH_CHECK(false, "s8s8_compensation not implemented!");
@@ -70,43 +69,6 @@ inline void pack_vnni<int8_t>(int8_t* __restrict__ packed, const int8_t* __restr
   s8s8_compensation<BLOCK_N>(packed, K);
 }
 
-// uint8_t: mxfp4 or int4
-// pack to vnni2 format as they are computed with bfloat16
-//
-// from [N, K'/2, 2] to [K'/2, N, 2], view 2x int4 as unit8:
-// from [N,    K   ] to [K,    N   ] where K = K'/2
-//
-template <>
-inline void pack_vnni<uint8_t>(uint8_t* __restrict__ packed, const uint8_t* __restrict__ weight, int N, int K) {
-  constexpr int BLOCK_N = block_size_n();
-
-  uint8_t unpacked[2 * BLOCK_N];
-
-  // 32-way pack (align with BLOCK_N), faster for avx512 unpacking
-  //
-  // for a range of (64):
-  //   {0, 1, 2, ..., 63}
-  //
-  // original format:
-  //   { 1|0,  3|2, ..., 63|62}
-  //
-  // packed format:
-  //   {32|0, 31|1, ..., 63|31}
-  //
-  for (int k = 0; k < K; ++k) {
-    // unpack first
-    for (int n = 0; n < N; ++n) {
-      uint8_t value = weight[n * K + k];
-      unpacked[n * 2 + 0] = value & 0xF;  // lower 4 bits
-      unpacked[n * 2 + 1] = value >> 4;   // higher 4 bits
-    }
-    // re-pack to 32-way
-    for (int n = 0; n < N; ++n) {
-      packed[k * N + n] = (unpacked[n + BLOCK_N] << 4) | unpacked[n];
-    }
-  }
-}
-
 template <typename scalar_t>
 inline void copy_stub(scalar_t* __restrict__ out, const float* __restrict__ input, int64_t size) {
   using bVec = at::vec::Vectorized<scalar_t>;
@@ -114,7 +76,7 @@ inline void copy_stub(scalar_t* __restrict__ out, const float* __restrict__ inpu
   constexpr int kVecSize = bVec::size();
 
   int64_t d;
-#pragma GCC unroll 4
+  #pragma GCC unroll 4
   for (d = 0; d <= size - kVecSize; d += kVecSize) {
     fVec data0 = fVec::loadu(input + d);
     fVec data1 = fVec::loadu(input + d + fVec::size());
@@ -127,34 +89,13 @@ inline void copy_stub(scalar_t* __restrict__ out, const float* __restrict__ inpu
 }
 
 template <typename scalar_t>
-inline void copy_stub(float* __restrict__ out, const scalar_t* __restrict__ input, int64_t size) {
+inline void copy_add_stub(scalar_t* __restrict__ out, const float* __restrict__ input, const float* __restrict__ bias, int64_t size) {
   using bVec = at::vec::Vectorized<scalar_t>;
   using fVec = at::vec::Vectorized<float>;
   constexpr int kVecSize = bVec::size();
 
   int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    fVec data0, data1;
-    bVec b_vec = bVec::loadu(input + d);
-    std::tie(data0, data1) = at::vec::convert_to_float(b_vec);
-    data0.store(out + d);
-    data1.store(out + d + fVec::size());
-  }
-  for (; d < size; ++d) {
-    out[d] = static_cast<float>(input[d]);
-  }
-}
-
-template <typename scalar_t>
-inline void copy_add_stub(
-    scalar_t* __restrict__ out, const float* __restrict__ input, const float* __restrict__ bias, int64_t size) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = bVec::size();
-
-  int64_t d;
-#pragma GCC unroll 4
+  #pragma GCC unroll 4
   for (d = 0; d <= size - kVecSize; d += kVecSize) {
     fVec data0 = fVec::loadu(input + d) + fVec::loadu(bias + d);
     fVec data1 = fVec::loadu(input + d + fVec::size()) + fVec::loadu(bias + d + fVec::size());
@@ -166,51 +107,11 @@ inline void copy_add_stub(
   }
 }
 
-template <typename scalar_t, bool has_bias>
-inline void scalar_sigmoid_and_mul(
-    scalar_t* __restrict__ out,
-    const float* __restrict__ input,
-    const float* __restrict__ bias,
-    const scalar_t* __restrict__ mul,
-    int SIZE) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  // scalar sigmoid
-  const fVec one = fVec(1.f);
-  fVec X;
-  if constexpr (has_bias) {
-    assert(bias != nullptr);
-    X = fVec(input[0] + bias[0]);
-  } else {
-    X = fVec(input[0]);
-  }
-  X = one / (one + X.neg().exp_u20());
-
-  // vec mul
-  constexpr int kVecSize = bVec::size();
-  for (int d = 0; d < SIZE; d += kVecSize) {
-    bVec m_bvec = bVec::loadu(mul + d);
-    fVec m_fvec0, m_fvec1;
-    std::tie(m_fvec0, m_fvec1) = at::vec::convert_to_float(m_bvec);
-    m_fvec0 = m_fvec0 * X;
-    m_fvec1 = m_fvec1 * X;
-
-    bVec out_vec = convert_from_float_ext<scalar_t>(m_fvec0, m_fvec1);
-    out_vec.store(out + d);
-  }
-}
-
 template <typename scalar_t, bool has_bias, int BLOCK_M, int BLOCK_N>
 struct tinygemm_kernel_nn {
   static inline void apply(
-      const scalar_t* __restrict__ A,
-      const scalar_t* __restrict__ B,
-      scalar_t* __restrict__ C,
-      const float* __restrict__ bias,
-      int64_t K,
-      int64_t lda,
-      int64_t ldb,
-      int64_t ldc) {
+      const scalar_t* __restrict__ A, const scalar_t* __restrict__ B, scalar_t* __restrict__ C,
+      const float* __restrict__ bias, int64_t K, int64_t lda, int64_t ldb, int64_t ldc) {
     TORCH_CHECK(false, "tinygemm_kernel_nn: scalar path not implemented!");
   }
 };
@@ -219,14 +120,9 @@ struct tinygemm_kernel_nn {
 template <bool has_bias, int BLOCK_M, int BLOCK_N>
 struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
   static inline void apply(
-      const at::BFloat16* __restrict__ A,
-      const at::BFloat16* __restrict__ B,
-      at::BFloat16* __restrict__ C,
-      const float* __restrict__ bias,
-      int64_t K,
-      int64_t lda,
-      int64_t ldb,
-      int64_t ldc) {
+      const at::BFloat16* __restrict__ A, const at::BFloat16* __restrict__ B, at::BFloat16* __restrict__ C,
+      const float* __restrict__ bias, int64_t K, int64_t lda, int64_t ldb, int64_t ldc) {
+
     constexpr int ROWS = BLOCK_M;
     constexpr int COLS = BLOCK_N / 16;
 
@@ -249,7 +145,7 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
 
     const int64_t K2 = K >> 1;
     const int64_t lda2 = lda >> 1;
-    const int64_t ldb2 = ldb;  // ldb * 2 >> 1;
+    const int64_t ldb2 = ldb; // ldb * 2 >> 1;
     const float* a_ptr = reinterpret_cast<const float*>(A);
     const float* b_ptr = reinterpret_cast<const float*>(B);
 
@@ -284,7 +180,9 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
               (__m512i)(_mm512_cvtne2ps_pbh(vc[row * COLS + col + 1], vc[row * COLS + col])));
         }
       } else {
-        _mm256_storeu_si256(reinterpret_cast<__m256i*>(C + row * ldc + col * 16), (__m256i)(_mm512_cvtneps_pbh(vc[i])));
+        _mm256_storeu_si256(
+            reinterpret_cast<__m256i*>(C + row * ldc + col * 16),
+            (__m256i)(_mm512_cvtneps_pbh(vc[i])));
       }
     };
     Unroll<ROWS * COLS>{}(storec);
@@ -292,33 +190,22 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
 };
 #endif
 
-#define LAUNCH_TINYGEMM_KERNEL_NN(MB_SIZE, NB_SIZE)                \
-  tinygemm_kernel_nn<scalar_t, has_bias, MB_SIZE, NB_SIZE>::apply( \
-      A + mb_start * lda,                                          \
-      B + nb_start * 2,                                            \
-      C + mb_start * ldc + nb_start,                               \
-      has_bias ? bias + nb_start : nullptr,                        \
-      K,                                                           \
-      lda,                                                         \
-      ldb,                                                         \
-      ldc);
+#define LAUNCH_TINYGEMM_KERNEL_NN(MB_SIZE, NB_SIZE)                          \
+    tinygemm_kernel_nn<scalar_t, has_bias, MB_SIZE, NB_SIZE>::apply(         \
+        A + mb_start * lda, B + nb_start * 2, C + mb_start * ldc + nb_start, \
+        has_bias ? bias + nb_start : nullptr, K, lda, ldb, ldc);
 
 template <typename scalar_t, bool has_bias>
 struct brgemm {
   static inline void apply(
-      const scalar_t* __restrict__ A,
-      const scalar_t* __restrict__ B,
-      scalar_t* __restrict__ C,
-      float* __restrict__ Ctmp,
-      const float* __restrict__ bias,
-      int64_t M,
-      int64_t N,
-      int64_t K,
-      int64_t lda,
-      int64_t ldb,
-      int64_t ldc) {
+      const scalar_t* __restrict__ A, const scalar_t* __restrict__ B, scalar_t* __restrict__ C,
+      float* __restrict__ Ctmp, const float* __restrict__ bias,
+      int64_t M, int64_t N, int64_t K, int64_t lda, int64_t ldb, int64_t ldc) {
+
     constexpr int BLOCK_N = block_size_n();
-    at::native::cpublas::brgemm(M, N, K, lda, ldb, BLOCK_N, /* add_C */ false, A, B, Ctmp);
+    at::native::cpublas::brgemm(
+        M, N, K, lda, ldb, BLOCK_N, /* add_C */false,
+        A, B, Ctmp);
 
     // copy from Ctmp to C
     for (int64_t m = 0; m < M; ++m) {
@@ -329,21 +216,6 @@ struct brgemm {
       }
     }
   }
-  static inline void apply(
-      const float* __restrict__ A,
-      const float* __restrict__ B,
-      scalar_t* __restrict__ C,
-      float* __restrict__ Ctmp,
-      const float* __restrict__ bias,
-      int64_t M,
-      int64_t N,
-      int64_t K,
-      int64_t lda,
-      int64_t ldb,
-      int64_t ldc) {
-    constexpr int BLOCK_N = block_size_n();
-    at::native::cpublas::brgemm(M, N, K, lda, ldb, BLOCK_N, /* add_C */ false, A, B, Ctmp);
-  }
 };
 
 template <typename scalar_t, bool has_bias>
@@ -360,12 +232,15 @@ void tinygemm_kernel(
     int64_t ldb,
     int64_t ldc,
     bool brg) {
+
   if (brg) {
-    brgemm<scalar_t, has_bias>::apply(A, B, C, Ctmp, bias, M, N, K, lda, ldb, ldc);
+    brgemm<scalar_t, has_bias>::apply(
+        A, B, C, Ctmp, bias,
+        M, N, K, lda, ldb, ldc);
     return;
   }
 
-  // pattern: 1-4-16, N = 16, 32, 48, 64
+  // pattern: 1-4-16
   constexpr int64_t BLOCK_M = 4;
   constexpr int64_t BLOCK_N = 64;
   const int64_t MB = div_up(M, BLOCK_M);
@@ -377,88 +252,25 @@ void tinygemm_kernel(
       int64_t nb_start = nb * BLOCK_N;
       int64_t nb_size = std::min(BLOCK_N, N - nb_start);
 
-      switch (mb_size << 4 | nb_size >> 4) {
+      switch(mb_size << 4 | nb_size >> 4) {
         // mb_size = 1
-        case 0x11:
-          LAUNCH_TINYGEMM_KERNEL_NN(1, 16);
-          break;
-        case 0x12:
-          LAUNCH_TINYGEMM_KERNEL_NN(1, 32);
-          break;
-        case 0x13:
-          LAUNCH_TINYGEMM_KERNEL_NN(1, 48);
-          break;
-        case 0x14:
-          LAUNCH_TINYGEMM_KERNEL_NN(1, 64);
-          break;
+        case 0x12: LAUNCH_TINYGEMM_KERNEL_NN(1, 32); break;
+        case 0x14: LAUNCH_TINYGEMM_KERNEL_NN(1, 64); break;
         // mb_size = 2
-        case 0x21:
-          LAUNCH_TINYGEMM_KERNEL_NN(2, 16);
-          break;
-        case 0x22:
-          LAUNCH_TINYGEMM_KERNEL_NN(2, 32);
-          break;
-        case 0x23:
-          LAUNCH_TINYGEMM_KERNEL_NN(2, 48);
-          break;
-        case 0x24:
-          LAUNCH_TINYGEMM_KERNEL_NN(2, 64);
-          break;
+        case 0x22: LAUNCH_TINYGEMM_KERNEL_NN(2, 32); break;
+        case 0x24: LAUNCH_TINYGEMM_KERNEL_NN(2, 64); break;
         // mb_size = 3
-        case 0x31:
-          LAUNCH_TINYGEMM_KERNEL_NN(3, 16);
-          break;
-        case 0x32:
-          LAUNCH_TINYGEMM_KERNEL_NN(3, 32);
-          break;
-        case 0x33:
-          LAUNCH_TINYGEMM_KERNEL_NN(3, 48);
-          break;
-        case 0x34:
-          LAUNCH_TINYGEMM_KERNEL_NN(3, 64);
-          break;
+        case 0x32: LAUNCH_TINYGEMM_KERNEL_NN(3, 32); break;
+        case 0x34: LAUNCH_TINYGEMM_KERNEL_NN(3, 64); break;
         // mb_size = 4
-        case 0x41:
-          LAUNCH_TINYGEMM_KERNEL_NN(4, 16);
-          break;
-        case 0x42:
-          LAUNCH_TINYGEMM_KERNEL_NN(4, 32);
-          break;
-        case 0x43:
-          LAUNCH_TINYGEMM_KERNEL_NN(4, 48);
-          break;
-        case 0x44:
-          LAUNCH_TINYGEMM_KERNEL_NN(4, 64);
-          break;
-        default:
-          TORCH_CHECK(false, "Unexpected block size, ", mb_size, " x ", nb_size);
+        case 0x42: LAUNCH_TINYGEMM_KERNEL_NN(4, 32); break;
+        case 0x44: LAUNCH_TINYGEMM_KERNEL_NN(4, 64); break;
+        default: TORCH_CHECK(false, "Unexpected block size, ", mb_size, "x", nb_size);
       }
     }
   }
 }
 
-template <typename scalar_t, bool has_bias>
-void tinygemm_kernel(
-    const float* __restrict__ A,
-    const float* __restrict__ B,
-    scalar_t* __restrict__ C,
-    float* __restrict__ Ctmp,
-    const float* __restrict__ bias,
-    int64_t M,
-    int64_t N,
-    int64_t K,
-    int64_t lda,
-    int64_t ldb,
-    int64_t ldc,
-    bool brg) {
-  TORCH_CHECK(brg, "Expected to use fp32 brgemm for small N GEMM");
-  if (brg) {
-    brgemm<scalar_t, has_bias>::apply(A, B, C, Ctmp, bias, M, N, K, lda, ldb, ldc);
-    return;
-  }
-  // TODO : add intrinsic path
-}
-
 template <typename scalar_t>
 void weight_packed_linear_kernel_impl(
     scalar_t* __restrict__ out,
@@ -470,20 +282,29 @@ void weight_packed_linear_kernel_impl(
     int64_t K,
     int64_t mat1_strideM,
     int64_t out_strideM) {
+
   constexpr int64_t BLOCK_M = block_size_m();
   constexpr int64_t BLOCK_N = block_size_n();
   const int64_t MB = div_up(M, BLOCK_M);
   const int64_t NB = div_up(N, BLOCK_N);
 
-  const bool use_brgemm = can_use_brgemm<scalar_t>(M);
+  // use avx512-bf16 when a) M is small; b) dtype is bfloat16, otherwise use amx
+  const bool use_brgemm = (M > 4) || (!std::is_same_v<scalar_t, at::BFloat16>);
+
+  // l2 cache block for n
+  int64_t cache_blocks_nb = get_cache_blocks<scalar_t>(BLOCK_N, K);
 
   // parallel on [MB, NB]
   AT_DISPATCH_BOOL(bias != nullptr, has_bias, [&] {
-    parallel_2d(MB, NB, [&](int64_t mb0, int64_t mb1, int64_t nb0, int64_t nb1) {
+    parallel_2d(MB, NB, [&](int64_t begin_mb, int64_t end_mb, int64_t begin_nb, int64_t end_nb) {
+
       // for brgemm, use float32 for accumulate
       alignas(64) float Ctmp[BLOCK_M * BLOCK_N];
 
-      loop_2d<scalar_t>(mb0, mb1, nb0, nb1, BLOCK_N * K, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
+      for (int64_t nbb = begin_nb; nbb < end_nb; nbb += cache_blocks_nb) {
+      for (int64_t mb = begin_mb; mb < end_mb; ++mb) {
+      for (int64_t nb = nbb; nb < std::min(nbb + cache_blocks_nb, end_nb); ++nb) {
+
         int64_t mb_start = mb * BLOCK_M;
         int64_t mb_size = std::min(M - mb_start, BLOCK_M);
         int64_t nb_start = nb * BLOCK_N;
@@ -502,7 +323,7 @@ void weight_packed_linear_kernel_impl(
             /* ldb */ nb_size,
             /* ldc */ out_strideM,
             /* brg */ use_brgemm);
-      });
+      }}}
 
       if (use_brgemm) {
         at::native::cpublas::brgemm_release();
@@ -511,113 +332,20 @@ void weight_packed_linear_kernel_impl(
   });
 }
 
-template <typename scalar_t>
-void weight_packed_linear_kernel_impl(
-    scalar_t* __restrict__ out,
-    const scalar_t* __restrict__ mat1,
-    const float* __restrict__ mat2,
-    const float* __restrict__ bias,
-    const scalar_t* __restrict__ post_mul_mat,
-    int64_t M,
-    int64_t N,
-    int64_t K,
-    int64_t mat1_strideM,
-    int64_t out_strideM) {
-  constexpr int64_t BLOCK_M = block_size_m();
-  constexpr int64_t BLOCK_N = block_size_n();
-  const int64_t MB = div_up(M, BLOCK_M);
-  const int64_t NB = div_up(N, BLOCK_N);
-
-  const bool use_brgemm = true;  // TODO: add intrinsic path
-  // parallel on [MB, NB]
-  AT_DISPATCH_BOOL(bias != nullptr, has_bias, [&] {
-    parallel_2d(MB, NB, [&](int64_t mb0, int64_t mb1, int64_t nb0, int64_t nb1) {
-      // for brgemm, use float32 for accumulate
-      alignas(64) float Atmp[BLOCK_M * K];
-      alignas(64) float Ctmp[BLOCK_M * BLOCK_N];
-
-      loop_2d<float>(mb0, mb1, nb0, nb1, BLOCK_N * K, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
-        int64_t mb_start = mb * BLOCK_M;
-        int64_t mb_size = std::min(M - mb_start, BLOCK_M);
-        int64_t nb_start = nb * BLOCK_N;
-        int64_t nb_size = std::min(N - nb_start, BLOCK_N);
-        for (int64_t m = 0; m < mb_size; ++m) {
-          copy_stub<scalar_t>(Atmp + m * K, mat1 + mb_start * mat1_strideM + m * K, K);
-        }
-        tinygemm_kernel<scalar_t, has_bias>(
-            /*   A */ Atmp,
-            /*   B */ mat2 + nb_start * K /* nb * BLOCK_N * K */,
-            /*   C */ out + mb_start * out_strideM + nb_start,
-            /* Ctmp*/ Ctmp,
-            /* bias*/ bias + nb_start,
-            /*   M */ mb_size,
-            /*   N */ nb_size,
-            /*   K */ K,
-            /* lda */ mat1_strideM,
-            /* ldb */ nb_size,
-            /* ldc */ out_strideM,
-            /* brg */ use_brgemm);
-
-        if (post_mul_mat != nullptr) {
-          for (int64_t m = 0; m < mb_size; ++m) {
-            scalar_sigmoid_and_mul<scalar_t, has_bias>(
-                out + mb_start * out_strideM + nb_start + m * out_strideM,
-                Ctmp + m * BLOCK_N,
-                bias + nb_start,
-                post_mul_mat + mb_start * out_strideM + m * out_strideM,
-                out_strideM);
-          }
-        } else {
-          for (int64_t m = 0; m < mb_size; ++m) {
-            if constexpr (has_bias) {
-              copy_add_stub(
-                  out + mb_start * out_strideM + nb_start + m * out_strideM, Ctmp + m * BLOCK_N, bias + nb_start, N);
-            } else {
-              copy_stub(out + mb_start * out_strideM + nb_start + m * out_strideM, Ctmp + m * BLOCK_N, N);
-            }
-          }
-        }
-      });
-
-      if (use_brgemm) {
-        at::native::cpublas::brgemm_release();
-      }
-    });
-  });
-}
-
-}  // anonymous namespace
+} // anonymous namespace
 
 // tinygemm interface
 template <typename scalar_t>
-void tinygemm_kernel(
-    const scalar_t* __restrict__ A,
-    const scalar_t* __restrict__ B,
-    scalar_t* __restrict__ C,
-    float* __restrict__ Ctmp,
-    int64_t M,
-    int64_t N,
-    int64_t K,
-    int64_t lda,
-    int64_t ldb,
-    int64_t ldc,
-    bool brg) {
+void tinygemm_kernel(const scalar_t* __restrict__ A, const scalar_t* __restrict__ B, scalar_t* __restrict__ C,
+    float* __restrict__ Ctmp, int64_t M, int64_t N, int64_t K, int64_t lda, int64_t ldb, int64_t ldc, bool brg) {
   tinygemm_kernel<scalar_t, false>(A, B, C, Ctmp, nullptr, M, N, K, lda, ldb, ldc, brg);
 }
 
-#define INSTANTIATE_TINYGEMM_TEMPLATE(TYPE) \
-  template void tinygemm_kernel<TYPE>(      \
-      const TYPE* __restrict__ A,           \
-      const TYPE* __restrict__ B,           \
-      TYPE* __restrict__ C,                 \
-      float* __restrict__ Ctmp,             \
-      int64_t M,                            \
-      int64_t N,                            \
-      int64_t K,                            \
-      int64_t lda,                          \
-      int64_t ldb,                          \
-      int64_t ldc,                          \
-      bool brg)
+#define INSTANTIATE_TINYGEMM_TEMPLATE(TYPE)                                             \
+    template void tinygemm_kernel<TYPE>(                                                \
+        const TYPE* __restrict__ A, const TYPE* __restrict__ B, TYPE* __restrict__ C,   \
+        float* __restrict__ Ctmp, int64_t M, int64_t N, int64_t K, int64_t lda,         \
+        int64_t ldb, int64_t ldc, bool brg)
 
 INSTANTIATE_TINYGEMM_TEMPLATE(at::BFloat16);
 INSTANTIATE_TINYGEMM_TEMPLATE(at::Half);
@@ -631,23 +359,14 @@ at::Tensor convert_weight_packed(at::Tensor& weight) {
 
   const int64_t ndim = weight.ndimension();
   TORCH_CHECK(ndim == 2 || ndim == 3, "expect weight to be 2d or 3d, got ", ndim, "d tensor.");
-
-  if (ndim == 2 && weight.size(0) < TILE_N) {
-    // for 2D weight and small OC shape, we use fma linear path, which needs transpose not pack
-    return weight.to(at::kFloat).t().contiguous();
-  }
-
   const auto st = weight.scalar_type();
   const int64_t E = ndim == 3 ? weight.size(0) : 1;
   const int64_t OC = ndim == 3 ? weight.size(1) : weight.size(0);
   const int64_t IC = ndim == 3 ? weight.size(2) : weight.size(1);
 
-  // mxfp4 or int4 are packed with uint8
-  const int64_t actual_IC = st == at::kByte ? IC * 2 : IC;
-
   // we handle 2 TILE_N at a time.
   TORCH_CHECK(OC % TILE_N == 0, "invalid weight out features ", OC);
-  TORCH_CHECK(actual_IC % TILE_K == 0, "invalid weight input features ", actual_IC);
+  TORCH_CHECK(IC % TILE_K == 0, "invalid weight input features ", IC);
 
   constexpr int64_t BLOCK_N = block_size_n();
   const int64_t NB = div_up(OC, BLOCK_N);
@@ -656,14 +375,12 @@ at::Tensor convert_weight_packed(at::Tensor& weight) {
   auto packed_weight = at::empty({}, weight.options());
   const int64_t stride = OC * IC;
 
-  // Note: for `kByte` (uint8), it represents either `mxfp4` or `int4`.
-  TORCH_CHECK(
-      st == at::kBFloat16 || st == at::kHalf || st == at::kChar || st == at::kFloat8_e4m3fn || st == at::kByte,
-      "expect weight to be bfloat16, float16, int8, fp8_e4m3 or uint8(mxfp4 or int4).");
+  TORCH_CHECK(st == at::kBFloat16 || st == at::kHalf || st == at::kChar || st == at::kFloat8_e4m3fn,
+      "expect weight to be bfloat16, float16, int8 or fp8_e4m3.");
 
   CPU_DISPATCH_PACKED_TYPES(st, [&] {
     // adjust most inner dimension size
-    const int packed_row_size = get_row_size<packed_t>(actual_IC);
+    const int packed_row_size = get_row_size<packed_t>(IC);
     auto sizes = weight.sizes().vec();
     sizes[ndim - 1] = packed_row_size;
     packed_weight.resize_(sizes);
@@ -682,7 +399,10 @@ at::Tensor convert_weight_packed(at::Tensor& weight) {
         int64_t n = nb * BLOCK_N;
         int64_t n_size = std::min(BLOCK_N, OC - n);
         pack_vnni<packed_t>(
-            packed_data + e * OC * packed_row_size + n * packed_row_size, w_data + e * stride + n * IC, n_size, IC);
+            packed_data + e * OC * packed_row_size + n * packed_row_size,
+            w_data + e * stride + n * IC,
+            n_size,
+            IC);
 
         // move to the next index
         data_index_step(e, E, nb, NB);
@@ -692,71 +412,33 @@ at::Tensor convert_weight_packed(at::Tensor& weight) {
   return packed_weight;
 }
 
-at::Tensor convert_scale_packed(at::Tensor& scale) {
-  CHECK_INPUT(scale);
-
-  const int64_t ndim = scale.ndimension();
-  TORCH_CHECK(ndim == 2 || ndim == 3, "expect scale to be 2d or 3d, got ", ndim, "d tensor.");
-  const auto st = scale.scalar_type();
-  const int64_t E = ndim == 3 ? scale.size(0) : 1;
-  const int64_t N = ndim == 3 ? scale.size(1) : scale.size(0);
-  // number of groups, e.g. K/32
-  const int64_t G = ndim == 3 ? scale.size(2) : scale.size(1);
-
-  constexpr int64_t BLOCK_N = block_size_n();
-  TORCH_CHECK(N % BLOCK_N == 0, "invalid weight out features ", N);
-  const int64_t NB = N / BLOCK_N;
-
-  auto packed_scale = at::empty_like(scale);
-  TORCH_CHECK(st == at::kByte, "expect scale to be uint8.");
-
-  const uint8_t* s_data = scale.data_ptr<uint8_t>();
-  uint8_t* packed_data = packed_scale.data_ptr<uint8_t>();
-
-  // parallel on src {E, NB, BLOCK_N, G}, dst {E, NB, G, BLOCK_N}
-  at::parallel_for(0, E * NB * BLOCK_N * G, 0, [&](int64_t begin, int64_t end) {
-    int64_t e{0}, nb{0}, n{0}, g{0};
-    data_index_init(begin, e, E, nb, NB, n, BLOCK_N, g, G);
-
-    for (int64_t i = begin; i < end; ++i) {
-      packed_data[e * N * G + nb * G * BLOCK_N + g * BLOCK_N + n] = s_data[i];
-      // move to the next index
-      data_index_step(e, E, nb, NB, n, BLOCK_N, g, G);
-    }
-  });
-  return packed_scale;
-}
-
 // mat1 : [M, K]
-// mat2 : [N, K] ([K, N] if use_fma_gemm)
+// mat2 : [N, K]
 // bias : [N]
 // out  : [M, N]
 //
-at::Tensor
-weight_packed_linear(at::Tensor& mat1, at::Tensor& mat2, const std::optional<at::Tensor>& bias, bool is_vnni) {
-  auto packed_w = is_vnni ? mat2 : convert_weight_packed(mat2);
-  bool use_fma_gemm = false;
-  if (packed_w.scalar_type() == at::kFloat) {
-    use_fma_gemm = true;
-  }
+at::Tensor weight_packed_linear(at::Tensor& mat1, at::Tensor& mat2,
+    const std::optional<at::Tensor>& bias, bool is_vnni) {
+  RECORD_FUNCTION(
+    "sgl-kernel::weight_packed_linear", std::vector<c10::IValue>({mat1, mat2, bias}));
 
-  int64_t M = mat1.size(0);
-  int64_t K = mat1.size(1);
-  int64_t N = use_fma_gemm ? mat2.size(1) : mat2.size(0);
+  auto packed_w = is_vnni ? mat2 : convert_weight_packed(mat2);
 
   CHECK_LAST_DIM_CONTIGUOUS_INPUT(mat1);
   CHECK_INPUT(mat2);
+
+  int64_t M = mat1.size(0);
+  int64_t N = mat2.size(0);
+  int64_t K = mat2.size(1);
+  CHECK_EQ(mat1.size(1), K);
   CHECK_DIM(2, mat1);
   CHECK_DIM(2, mat2);
-  if (!use_fma_gemm) {
-    CHECK_EQ(mat1.size(1), K);
-  }
 
-  auto dispatch_type = mat1.scalar_type();
   auto out = at::empty({M, N}, mat1.options());
+
   // strides
+  int64_t mat1_strideM = mat1.stride(0);
   int64_t out_strideM = out.stride(0);
-  int64_t mat1_strideM = mat1.stride(0);
 
   const bool has_bias = bias.has_value();
   const float* bias_data = nullptr;
@@ -765,83 +447,12 @@ weight_packed_linear(at::Tensor& mat1, at::Tensor& mat2, const std::optional<at:
     bias_data = bias.value().data_ptr<float>();
   }
 
-  AT_DISPATCH_REDUCED_FLOATING_TYPES(dispatch_type, "weight_packed_linear_kernel_impl", [&] {
-    if (use_fma_gemm) {
-      weight_packed_linear_kernel_impl<scalar_t>(
-          out.data_ptr<scalar_t>(),
-          mat1.data_ptr<scalar_t>(),
-          packed_w.data_ptr<float>(),
-          bias_data,
-          nullptr,
-          M,
-          N,
-          K,
-          mat1_strideM,
-          out_strideM);
-    } else {
-      weight_packed_linear_kernel_impl<scalar_t>(
-          out.data_ptr<scalar_t>(),
-          mat1.data_ptr<scalar_t>(),
-          packed_w.data_ptr<scalar_t>(),
-          bias_data,
-          M,
-          N,
-          K,
-          mat1_strideM,
-          out_strideM);
-    }
-  });
-
-  return out;
-}
-
-// mat1         : [M, K]
-// mat2         : [K, 1]
-// post_mul_mat : [M, K]
-// bias         : [N]
-// out          : [M, N]
-//
-at::Tensor fused_linear_sigmoid_mul(
-    at::Tensor& mat1,
-    at::Tensor& mat2,
-    const std::optional<at::Tensor>& bias,
-    bool is_vnni,
-    const at::Tensor& post_mul_mat) {
-  auto packed_w = is_vnni ? mat2 : convert_weight_packed(mat2);
-  TORCH_CHECK(packed_w.scalar_type() == at::kFloat, "fused_linear_sigmoid_mul requires packed float weight")
-
-  int64_t M = mat1.size(0);
-  int64_t K = mat1.size(1);
-  int64_t N = mat2.size(1);
-
-  CHECK_LAST_DIM_CONTIGUOUS_INPUT(mat1);
-  CHECK_INPUT(mat2);
-  CHECK_DIM(2, mat1);
-  CHECK_DIM(2, mat2);
-
-  int64_t out_strideM = post_mul_mat.size(1);
-  int64_t mat1_strideM = mat1.stride(0);
-  auto dispatch_type = mat1.scalar_type();
-  auto out = at::empty({M, out_strideM}, mat1.options());
-
-  TORCH_CHECK(
-      N == 1 && out_strideM % 32 == 0,
-      "post_mul_mat tensor size(1) should be 32 dividable, and the mat2 OC=1 (Mx1 as linear output shape)")
-
-  const bool has_bias = bias.has_value();
-  const float* bias_data = nullptr;
-  if (has_bias) {
-    CHECK_EQ(bias.value().size(0), N);
-    bias_data = bias.value().data_ptr<float>();
-  }
-
-  AT_DISPATCH_REDUCED_FLOATING_TYPES(dispatch_type, "fused_linear_sigmoid_mul", [&] {
+  AT_DISPATCH_REDUCED_FLOATING_TYPES(mat1.scalar_type(), "weight_packed_linear_kernel_impl", [&] {
     weight_packed_linear_kernel_impl<scalar_t>(
         out.data_ptr<scalar_t>(),
         mat1.data_ptr<scalar_t>(),
-        packed_w.data_ptr<float>(),
+        packed_w.data_ptr<scalar_t>(),
         bias_data,
-        post_mul_mat.data_ptr<scalar_t>(),
         M,
         N,
         K,

csrc/cpu/sgl-kernels/gemm.h

@@ -1,12 +1,8 @@
-// Adapted from
-// https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
-
-// clang-format off
-
 #pragma once
+
 #include <ATen/native/CPUBlas.h>
 
-#include "common.h"
+// clang-format off
 
 // amx-bf16
 #define TILE_M 16
@@ -14,42 +10,20 @@
 #define TILE_K 32
 
 // block size for AMX gemm
-constexpr int block_size_m() {
-  return 2 * TILE_M;
-}
-constexpr int block_size_n() {
-  return 2 * TILE_N;
-}
+constexpr int block_size_m() { return 2 * TILE_M; }
+constexpr int block_size_n() { return 2 * TILE_N; }
 
 // define threshold using brgemm (intel AMX)
-template <typename T>
-inline bool can_use_brgemm(int M);
-template <>
-inline bool can_use_brgemm<at::BFloat16>(int M) {
-  return M > 4;
-}
-template <>
-inline bool can_use_brgemm<at::Half>(int M) {
-  return true;
-}
-// this requires PyTorch 2.7 or above
-template <>
-inline bool can_use_brgemm<int8_t>(int M) {
-  return M > 4;
-}
-
-template <>
-inline bool can_use_brgemm<uint8_t>(int M) {
-  return M > 4;
-}
-
-template <>
-inline bool can_use_brgemm<at::Float8_e4m3fn>(int M) {
-  return M > 4;
-}
+template <typename T> inline bool can_use_brgemm(int M);
+template <> inline bool can_use_brgemm<at::BFloat16>(int M) { return M > 4; }
+template <> inline bool can_use_brgemm<at::Half>(int M) { return true; }
+template <> inline bool can_use_brgemm<int8_t>(int M) { return M > 4; }
+template <> inline bool can_use_brgemm<uint8_t>(int M) { return M > 4; }
+template <> inline bool can_use_brgemm<at::Float8_e4m3fn>(int M) { return M > 4; }
+template <> inline bool can_use_brgemm<at::quint4x2>(int M) { return M > 4; }
 
 // work around compiler internal error
-#define BLOCK_K 128  // 4 * TILE_K
+#define BLOCK_K 128 // 4 * TILE_K
 
 // adjust leading dimension size for K
 template <typename T>
@@ -62,57 +36,18 @@ inline int64_t get_row_size<int8_t>(int64_t K) {
   return K + sizeof(int32_t);
 }
 
-// uint8: mxfp4 or int4
-template <>
-inline int64_t get_row_size<uint8_t>(int64_t K) {
-  return K >> 1;
-}
-
 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 };
-
-constexpr bool operator==(CPUQuantMethod a, int64_t b) {
-  return static_cast<int64_t>(a) == b;
-}
-
-constexpr bool operator==(int64_t a, CPUQuantMethod b) {
-  return a == static_cast<int64_t>(b);
-}
-
-enum class CPUQuantAlgo : int64_t { AWQ = 0, GPTQ = 1 };
-
-constexpr bool operator==(CPUQuantAlgo a, int64_t b) {
-  return static_cast<int64_t>(a) == b;
-}
-
-constexpr bool operator==(int64_t a, CPUQuantAlgo b) {
-  return a == static_cast<int64_t>(b);
-}
-
 inline int64_t get_4bit_block_k_size(int64_t group_size) {
   return group_size > 128 ? 128 : group_size;
 }
 
-// pack weight to vnni format
+// pack weight into 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
+// pack weight to vnni format for int4 (adapted from sglang)
 std::tuple<at::Tensor, at::Tensor, at::Tensor>
 convert_weight_packed_scale_zp(at::Tensor qweight, at::Tensor qzeros, at::Tensor scales);
 
@@ -142,9 +77,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 +88,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 +103,36 @@ 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);
+
+// moe implementations for int4 w4a16
+template <typename scalar_t>
+void fused_experts_int4_w4a16_kernel_impl(
+    scalar_t* __restrict__ output,
+    scalar_t* __restrict__ ic0,
+    scalar_t* __restrict__ ic1,
+    scalar_t* __restrict__ ic2,
+    scalar_t* __restrict__ A_tmp,
+    scalar_t* __restrict__ B_tmp,
+    float* __restrict__ C_tmp,
+    const scalar_t* __restrict__ input,
+    const at::quint4x2* __restrict__ packed_w1,
+    const at::quint4x2* __restrict__ packed_w2,
+    const uint8_t* __restrict__ w1z,
+    const uint8_t* __restrict__ w2z,
+    const scalar_t* __restrict__ w1s,
+    const scalar_t* __restrict__ w2s,
+    int group_size,
+    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);
 
 // shared expert implementation for int8 w8a8
 template <typename scalar_t>
@@ -195,37 +153,6 @@ void shared_expert_int8_kernel_impl(
     int64_t N,
     int64_t K);
 
-template <typename scalar_t>
-void fused_experts_int4_w4a8_kernel_impl(
-    scalar_t* __restrict__ output,
-    scalar_t* __restrict__ ic0,
-    scalar_t* __restrict__ ic1,
-    scalar_t* __restrict__ ic2,
-    uint8_t* __restrict__ A_tmp,
-    uint8_t* __restrict__ Aq_tmp,
-    float* __restrict__ As_tmp,
-    int32_t* __restrict__ Azp_tmp,
-    float* __restrict__ C_tmp,
-    int8_t* __restrict__ dqB_tmp,
-    const scalar_t* __restrict__ input,
-    const uint8_t* __restrict__ packed_w1,
-    const uint8_t* __restrict__ packed_w2,
-    const int8_t* __restrict__ w1z,
-    const int8_t* __restrict__ w2z,
-    const float* __restrict__ w1s,
-    const float* __restrict__ w2s,
-    int group_size,
-    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);
-
 template <typename scalar_t>
 void shared_expert_fp8_kernel_impl(
     scalar_t* __restrict__ output,
@@ -277,7 +204,6 @@ void tinygemm_kernel(
     int64_t ldc,
     bool brg);
 
-// block quantization
 template <typename scalar_t>
 void tinygemm_kernel(
     const scalar_t* __restrict__ A,
@@ -285,7 +211,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,
@@ -294,46 +219,33 @@ void tinygemm_kernel(
     int64_t ldb,
     int64_t ldc,
     bool brg,
-    int64_t block_size_K,
-    bool do_unpack = true);
+    int64_t block_size_K);
 
-// per tensor quantization
 template <typename scalar_t>
 void tinygemm_kernel(
     const scalar_t* __restrict__ A,
-    const at::Float8_e4m3fn* __restrict__ B,
+    const at::quint4x2* __restrict__ B,
     scalar_t* __restrict__ C,
+    const uint8_t* __restrict__ Bz,
+    const scalar_t* __restrict__ Bs,
     scalar_t* __restrict__ Btmp,
     float* __restrict__ Ctmp,
-    float scale,
     int64_t M,
     int64_t N,
     int64_t K,
+    int group_size,
     int64_t lda,
     int64_t ldb,
     int64_t ldc,
+    int64_t strideBz,
+    int64_t strideBs,
     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);
+// int4 scaled GEMM (adapted from sglang)
+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);
 
+// int4 tinygemm kernel interface(adapted from sglang)
 template <typename scalar_t>
 void tinygemm_kernel(
     scalar_t* C,
@@ -354,21 +266,34 @@ void tinygemm_kernel(
     bool store_out,
     bool use_brgemm);
 
-// 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 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);
+// TODO: debug print, remove me later
+inline void print_16x32i(const __m512i x) {
+  int32_t a[16];
+  _mm512_storeu_si512((__m512i *)a, x);
+
+  for (int i = 0; i < 16; i++){
+    std::cout << a[i] << " ";
+  }
+  std::cout << std::endl;
+}
+
+inline void print_16x32(const __m512 x) {
+  float a[16];
+  _mm512_storeu_ps((__m512 *)a, x);
+
+  for (int i = 0; i < 16; i++){
+    std::cout << a[i] << " ";
+  }
+  std::cout << std::endl;
+}
+
+
+inline void print_32x8u(const __m256i x) {
+  uint8_t a[32];
+  _mm256_storeu_si256((__m256i *)a, x);
+
+  for (int i = 0; i < 32; ++i) {
+    std::cout << int32_t(a[i]) << " ";
+  }
+  std::cout << std::endl;
+}

csrc/cpu/sgl-kernels/gemm_int4.cpp

@@ -1,10 +1,10 @@
-// Adapted from
-// https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
+// SPDX-License-Identifier: Apache-2.0
+// Adapted from sgl-project/sglang
+// https://github.com/sgl-project/sglang/pull/8226
 
-// clang-format off
-
-#include <torch/all.h>
+#include <ATen/ATen.h>
 
+#include "common.h"
 #include "gemm.h"
 #include "vec.h"
 
@@ -24,22 +24,19 @@ struct ActDtype<false> {
   using type = uint8_t;
 };
 
-#if defined(CPU_CAPABILITY_AVX512)
 struct alignas(32) m256i_wrapper {
   __m256i data;
 };
 
-inline std::array<m256i_wrapper, 2> load_zps_4vnni(const int8_t* __restrict__ zps) {
-  // broadcast 01234567 to
-  // 01234567012345670123456701234567
-  __m256i vzps_low = _mm256_set1_epi64x(*reinterpret_cast<const long*>(zps));
-  __m256i vzps_high = _mm256_set1_epi64x(*reinterpret_cast<const long*>(zps + 8));
-  // shuffle from
-  // 01234567012345670123456701234567
-  // to
-  // 00001111222233334444555566667777
+#if defined(CPU_CAPABILITY_AVX512)
+inline std::array<m256i_wrapper, 2> load_zps_4vnni(
+    const int8_t* __restrict__ zps) {
+  __m256i vzps_low = _mm256_set1_epi64x(*reinterpret_cast<const int64_t*>(zps));
+  __m256i vzps_high =
+      _mm256_set1_epi64x(*reinterpret_cast<const int64_t*>(zps + 8));
   __m256i shuffle_mask =
-      _mm256_set_epi8(7, 7, 7, 7, 6, 6, 6, 6, 5, 5, 5, 5, 4, 4, 4, 4, 3, 3, 3, 3, 2, 2, 2, 2, 1, 1, 1, 1, 0, 0, 0, 0);
+      _mm256_set_epi8(7, 7, 7, 7, 6, 6, 6, 6, 5, 5, 5, 5, 4, 4, 4, 4, 3, 3, 3,
+                      3, 2, 2, 2, 2, 1, 1, 1, 1, 0, 0, 0, 0);
   vzps_low = _mm256_shuffle_epi8(vzps_low, shuffle_mask);
   vzps_high = _mm256_shuffle_epi8(vzps_high, shuffle_mask);
   m256i_wrapper vzps_low_wp, vzps_high_wp;
@@ -48,7 +45,8 @@ inline std::array<m256i_wrapper, 2> load_zps_4vnni(const int8_t* __restrict__ zp
   return {vzps_low_wp, vzps_high_wp};
 }
 
-inline std::array<m256i_wrapper, 2> load_uint4_as_int8(const uint8_t* __restrict__ qB) {
+inline std::array<m256i_wrapper, 2> load_uint4_as_int8(
+    const uint8_t* __restrict__ qB) {
   __m256i packed = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(qB));
   const __m256i low_mask = _mm256_set1_epi8(0x0f);
   __m256i high = _mm256_srli_epi16(packed, 4);
@@ -60,42 +58,37 @@ inline std::array<m256i_wrapper, 2> load_uint4_as_int8(const uint8_t* __restrict
   return {low_wp, high_wp};
 }
 
-template <int64_t N, int64_t ldb>
-void _dequant_weight_zp_only(const uint8_t* __restrict__ B, int8_t* dqB, const int8_t* __restrict__ qzeros, int64_t K) {
-  // unpack weight int8 -> two int4
-  // subtract zero point
-  // B shape = [K, ldb] = [K, N / 2], actual shape = [K / 4, N / 2, 4]
-  // dqB shape = [K, N], actual shape = [K / 4, N, 4]
-#pragma GCC unroll 2
+template <int N, int ldb>
+void _dequant_weight_zp_only(const uint8_t* __restrict__ B, int8_t* dqB,
+                             const int8_t* __restrict__ qzeros, int64_t K) {
+  #pragma GCC unroll 2
   for (int n = 0; n < N; n += 16) {
     auto [zps_low_wp, zps_high_wp] = load_zps_4vnni(&qzeros[n]);
     auto zps_low = zps_low_wp.data;
     auto zps_high = zps_high_wp.data;
     for (int k = 0; k < K; k += 4) {
-      auto [vb_low_wp, vb_high_wp] = load_uint4_as_int8(B + ldb * k + n / 2 * 4);
+      auto [vb_low_wp, vb_high_wp] =
+          load_uint4_as_int8(B + ldb * k + n / 2 * 4);
       auto vb_low = vb_low_wp.data;
       auto vb_high = vb_high_wp.data;
       vb_high = _mm256_sub_epi8(vb_high, zps_high);
       vb_low = _mm256_sub_epi8(vb_low, zps_low);
-      // store vb to B
-      _mm256_storeu_si256(reinterpret_cast<__m256i_u*>(dqB + N * k + n * 4), vb_low);
-      _mm256_storeu_si256(reinterpret_cast<__m256i_u*>(dqB + N * k + (n + 8) * 4), vb_high);
+      _mm256_storeu_si256(reinterpret_cast<__m256i_u*>(dqB + N * k + n * 4),
+                          vb_low);
+      _mm256_storeu_si256(
+          reinterpret_cast<__m256i_u*>(dqB + N * k + (n + 8) * 4), vb_high);
     }
   }
 }
 
-template <bool accum, int64_t N, bool sym_quant_act>
-void _dequant_and_store(
-    float* __restrict__ output,
-    const int32_t* __restrict__ input,
-    const float* __restrict__ scale_a,
-    const int32_t* __restrict__ zp_a,
-    const float* __restrict__ scale_b,
-    const int32_t* __restrict__ comp_b,
-    int M,
-    int ldi,
-    int ldo,
-    int ldsa = 1) {
+template <bool sym_quant_act, int N, bool accum>
+void _dequant_and_store(float* __restrict__ output,
+                        const int32_t* __restrict__ input,
+                        const float* __restrict__ scale_a,
+                        const int32_t* __restrict__ zp_a,
+                        const float* __restrict__ scale_b,
+                        const int32_t* __restrict__ comp_b, int M, int ldi,
+                        int ldo, int ldsa = 1) {
   for (int m = 0; m < M; ++m) {
     float a_scale = *(scale_a + m * ldsa);
     __m512 va_scale = _mm512_set1_ps(a_scale);
@@ -106,7 +99,7 @@ void _dequant_and_store(
       va_zp = _mm512_set1_epi32(a_zp);
     }
     int n = 0;
-#pragma GCC unroll 2
+  #pragma GCC unroll 2
     for (; n < N; n += 16) {
       __m512i vc = _mm512_loadu_si512(input + m * ldi + n);
       if constexpr (!sym_quant_act) {
@@ -129,7 +122,8 @@ void _dequant_and_store(
       if constexpr (sym_quant_act) {
         dq_val = (float)input[m * ldi + n] * a_scale * scale_b[n];
       } else {
-        dq_val = (float)(input[m * ldi + n] - a_zp * comp_b[n]) * a_scale * scale_b[n];
+        dq_val = (float)(input[m * ldi + n] - a_zp * comp_b[n]) * a_scale *
+                 scale_b[n];
       }
       if constexpr (accum) {
         output[m * ldo + n] += dq_val;
@@ -141,10 +135,9 @@ void _dequant_and_store(
 }
 
 #else
-template <int64_t N, int64_t ldb>
-void _dequant_weight_zp_only(const uint8_t* B, int8_t* dqB, const int8_t* qzeros, int64_t K) {
-  // B shape = [K, N / 2]
-  // dqB shape = [K, N]
+template <int N, int ldb>
+void _dequant_weight_zp_only(const uint8_t* B, int8_t* dqB,
+                             const int8_t* qzeros, int64_t K) {
   for (int k = 0; k < K; ++k) {
     for (int n = 0; n < N / 2; ++n) {
       int32_t b = (int32_t)B[k * ldb + n];
@@ -165,31 +158,20 @@ inline __m512i combine_m256i(std::array<m256i_wrapper, 2> two_256) {
   return combine_m256i(two_256[0].data, two_256[1].data);
 }
 
-// negate elements in a according to b's sign
 static inline __m512i _mm512_sign_epi8(__m512i a, __m512i b) {
   __m512i zero = _mm512_setzero_si512();
   __mmask64 blt0 = _mm512_movepi8_mask(b);
   return _mm512_mask_sub_epi8(a, blt0, zero, a);
 }
 
-template <int64_t M, int64_t N, int64_t ldb, bool sym_quant_act>
-void _dequant_gemm_accum_small_M(
-    float* __restrict__ C,
-    const uint8_t* A,
-    const float* scales_a,
-    const int32_t* qzeros_a,
-    const uint8_t* B,
-    const float* scales_b,
-    const int8_t* qzeros_b,
-    int64_t K,
-    int64_t lda,
-    int64_t ldc) {
-  // if sym_quant_act is true, A pointer type is passed in as uint8_t* but actually int8_t*.
-
+template <bool sym_quant_act, int M, int N, int ldb>
+void _dequant_gemm_accum_small_M(float* __restrict__ C, const uint8_t* A,
+                                 const float* scales_a, const int32_t* qzeros_a,
+                                 const uint8_t* B, const float* scales_b,
+                                 const int8_t* qzeros_b, int64_t K, int64_t lda,
+                                 int64_t ldc) {
   constexpr int COLS = N / 16;
-  // Computing compensation is faster than loading it for small M
-  // because it's memory bound.
-  __m512i ones = _mm512_set1_epi8(1);  // used for computing compensation
+  __m512i ones = _mm512_set1_epi8(1);
   __m512i va;
   __m512i vb[COLS];
   __m512i vc[M * COLS];
@@ -197,7 +179,6 @@ void _dequant_gemm_accum_small_M(
   __m512i vzps[COLS];
   __m512i vcompensate[COLS];
 
-  // Load scales and zps
   Unroll<COLS>{}([&](auto i) {
     vscales[i] = _mm512_loadu_ps(scales_b + i * 16);
     vzps[i] = combine_m256i(load_zps_4vnni(qzeros_b + i * 16));
@@ -233,25 +214,25 @@ void _dequant_gemm_accum_small_M(
     }
   };
 
-  // Accumulate along k
   constexpr const int unroll = 4;
   int k = 0;
   for (; k < K / 4 / unroll; k++) {
-    Unroll<unroll>{}([&](auto i) { Unroll<M * COLS>{}(compute, 4 * (k * unroll + i)); });
+    Unroll<unroll>{}(
+        [&](auto i) { Unroll<M * COLS>{}(compute, 4 * (k * unroll + i)); });
   }
   k *= 4 * unroll;
   for (; k < K; k += 4) {
     Unroll<M * COLS>{}(compute, k);
   }
 
-  // Store to C
   auto store = [&](auto i) {
     constexpr const int row = i / COLS;
     constexpr const int col = i % COLS;
-    // compute (qC - compensate * zp_a) * scale_a * scale_b
     __m512 vc_float;
     if constexpr (!sym_quant_act) {
-      vc[i] = _mm512_sub_epi32(vc[i], _mm512_mullo_epi32(vcompensate[col], _mm512_set1_epi32(*(qzeros_a + row))));
+      vc[i] = _mm512_sub_epi32(
+          vc[i], _mm512_mullo_epi32(vcompensate[col],
+                                    _mm512_set1_epi32(*(qzeros_a + row))));
     }
     vc_float = _mm512_cvtepi32_ps(vc[i]);
     vc_float = _mm512_mul_ps(vc_float, _mm512_set1_ps(*(scales_a + row)));
@@ -264,28 +245,17 @@ void _dequant_gemm_accum_small_M(
   Unroll<M * COLS>{}(store);
 }
 
-#define CALL_DEQUANT_GEMM_ACCUM_SMALL_M(M) \
-  _dequant_gemm_accum_small_M<M, N, ldb, sym_quant_act>(C, A, scales_a, qzeros_a, B, scales_b, qzeros_b, K, lda, ldc);
+  #define CALL_DEQUANT_GEMM_ACCUM_SMALL_M(M)               \
+    _dequant_gemm_accum_small_M<sym_quant_act, M, N, ldb>( \
+        C, A, scales_a, qzeros_a, B, scales_b, qzeros_b, K, lda, ldc);
 #endif
 
-template <int64_t N, int64_t ldb, bool sym_quant_act>
-void _dequant_gemm_accum(
-    float* C,
-    const uint8_t* A,
-    const float* scales_a,
-    const int32_t* qzeros_a,
-    const uint8_t* B,
-    const float* scales_b,
-    const int8_t* qzeros_b,
-    const int32_t* compensation,
-    int8_t* dqB,
-    int64_t M,
-    int64_t K,
-    int64_t lda,
-    int64_t ldc,
-    bool use_brgemm) {
-  // Compute GEMM int8 * int8 -> int32
-  // dequant result to float by applying scales/qzeros
+template <bool sym_quant_act, int N, int ldb>
+void _dequant_gemm_accum(float* C, const uint8_t* A, const float* scales_a,
+                         const int32_t* qzeros_a, const uint8_t* B,
+                         const float* scales_b, const int8_t* qzeros_b,
+                         const int32_t* compensation, int8_t* dqB, int64_t M,
+                         int64_t K, int64_t lda, int64_t ldc, bool use_brgemm) {
 #if defined(CPU_CAPABILITY_AVX512)
   if (!use_brgemm) {
     switch (M) {
@@ -312,12 +282,14 @@ void _dequant_gemm_accum(
   Tin* A_ptr = (Tin*)A;
   if (use_brgemm) {
     int32_t C_i32[M * N];
-    at::native::cpublas::brgemm(
-        M, N, K, lda, N /*ldb*/, N /*ldc*/, false /* add_C */, A_ptr, dqB, C_i32, true /* is_vnni */);
+    at::native::cpublas::brgemm(M, N, K, lda, N /*ldb*/, N /*ldc*/,
+                                false /* add_C */, A_ptr, dqB, C_i32,
+                                true /* is_vnni */);
     _mm_prefetch(B + N * K / 2, _MM_HINT_T0);
     _mm_prefetch(A + K, _MM_HINT_T0);
-    _dequant_and_store<true, N, sym_quant_act>(
-        C, C_i32, scales_a, qzeros_a, scales_b, compensation, M, N /*ldi*/, ldc, 1 /*ldsa*/);
+    _dequant_and_store<sym_quant_act, N, true>(C, C_i32, scales_a, qzeros_a,
+                                               scales_b, compensation, M,
+                                               N /*ldi*/, ldc, 1 /*ldsa*/);
   } else
 #endif
   {
@@ -325,13 +297,13 @@ void _dequant_gemm_accum(
   }
 }
 
-template <int64_t N>
+template <int N>
 inline void copy_bias(const float* bias_ptr, float* y_buf, int64_t m) {
   if (bias_ptr) {
     for (int i = 0; i < m; ++i) {
       int j = 0;
 #if defined(CPU_CAPABILITY_AVX512)
-#pragma GCC unroll 2
+  #pragma GCC unroll 2
       for (; j < N; j += 16) {
         __m512 bias_vec = _mm512_loadu_ps(bias_ptr + j);
         _mm512_storeu_ps(y_buf + i * N + j, bias_vec);
@@ -341,11 +313,11 @@ inline void copy_bias(const float* bias_ptr, float* y_buf, int64_t m) {
         y_buf[i * N + j] = bias_ptr[j];
       }
     }
-  } else {  // initialize to zero
+  } else {
     for (int i = 0; i < m; ++i) {
       int j = 0;
 #if defined(CPU_CAPABILITY_AVX512)
-#pragma GCC unroll 2
+  #pragma GCC unroll 2
       for (; j < N; j += 16) {
         __m512 zero_vec = _mm512_setzero_ps();
         _mm512_storeu_ps(y_buf + i * N + j, zero_vec);
@@ -358,13 +330,14 @@ inline void copy_bias(const float* bias_ptr, float* y_buf, int64_t m) {
   }
 }
 
-template <typename out_dtype, int64_t N>
-inline void store_out(const float* y_buf, out_dtype* c_ptr, int64_t m, /* int64_t n, */ int64_t lda) {
+template <int N, typename out_dtype>
+inline void store_out(const float* y_buf, out_dtype* c_ptr, int64_t m,
+                      int64_t lda) {
   for (int i = 0; i < m; ++i) {
     int j = 0;
     if constexpr (std::is_same<out_dtype, float>::value) {
 #if defined(CPU_CAPABILITY_AVX512)
-#pragma GCC unroll 2
+  #pragma GCC unroll 2
       for (; j < N; j += 16) {
         __m512 y_vec = _mm512_loadu_ps(y_buf + i * N + j);
         _mm512_storeu_ps(c_ptr + i * lda + j, y_vec);
@@ -375,11 +348,12 @@ inline void store_out(const float* y_buf, out_dtype* c_ptr, int64_t m, /* int64_
       }
     } else if constexpr (std::is_same<out_dtype, at::BFloat16>::value) {
 #if defined(CPU_CAPABILITY_AVX512)
-#pragma GCC unroll 2
+  #pragma GCC unroll 2
       for (; j < N; j += 16) {
         __m512 y_vec = _mm512_loadu_ps(y_buf + i * N + j);
         __m256i y_bf16_vec = at::vec::cvtfp32_bf16(y_vec);
-        _mm256_storeu_si256(reinterpret_cast<__m256i*>(c_ptr + i * lda + j), y_bf16_vec);
+        _mm256_storeu_si256(reinterpret_cast<__m256i*>(c_ptr + i * lda + j),
+                            y_bf16_vec);
       }
 #endif
       for (; j < N; ++j) {
@@ -387,11 +361,12 @@ inline void store_out(const float* y_buf, out_dtype* c_ptr, int64_t m, /* int64_
       }
     } else if constexpr (std::is_same<out_dtype, at::Half>::value) {
 #if defined(CPU_CAPABILITY_AVX512)
-#pragma GCC unroll 2
+  #pragma GCC unroll 2
       for (; j < N; j += 16) {
         __m512 y_vec = _mm512_loadu_ps(y_buf + i * N + j);
         __m256i y_fp16_vec = at::vec::cvtfp32_fp16(y_vec);
-        _mm256_storeu_si256(reinterpret_cast<__m256i*>(c_ptr + i * lda + j), y_fp16_vec);
+        _mm256_storeu_si256(reinterpret_cast<__m256i*>(c_ptr + i * lda + j),
+                            y_fp16_vec);
       }
 #endif
       for (; j < N; ++j) {
@@ -417,25 +392,16 @@ void fill_val_stub(int32_t* __restrict__ output, int32_t value, int64_t size) {
   }
 }
 
-template <typename act_dtype, typename out_dtype, bool sym_quant_act>
+template <bool sym_quant_act, typename act_dtype, typename out_dtype>
 void _da8w4_linear_impl(
-    act_dtype* __restrict__ input,
-    const float* __restrict__ input_scales,
+    act_dtype* __restrict__ input, const float* __restrict__ input_scales,
     const int32_t* __restrict__ input_qzeros,
-    const uint8_t* __restrict__ weight,
-    const float* __restrict__ weight_scales,
-    const int8_t* __restrict__ weight_qzeros,
-    const float* __restrict__ bias,
-    out_dtype* __restrict__ output,
-    float* __restrict__ output_temp,
-    int8_t* __restrict__ dequant_weight_temp,
-    int64_t M,
-    int64_t N,
-    int64_t K,
+    const uint8_t* __restrict__ weight, const float* __restrict__ weight_scales,
+    const int8_t* __restrict__ weight_qzeros, const float* __restrict__ bias,
+    out_dtype* __restrict__ output, float* __restrict__ output_temp,
+    int8_t* __restrict__ dequant_weight_temp, int64_t M, int64_t N, int64_t K,
     int64_t num_groups) {
-  // weight + compensation shape = [Nc, Kc, BLOCK_N * _block_k / 2 + BLOCK_N*sizeof(int32_t)]
-  // scales/qzeros shape = [Nc, G, BLOCK_N]
-  const bool use_brgemm = can_use_brgemm<int8_t>(M);
+  const bool use_brgemm = can_use_brgemm<act_dtype>(M);
   int64_t block_m = [&]() -> long {
     if (M <= 48) {
       return M;
@@ -466,24 +432,30 @@ void _da8w4_linear_impl(
       int64_t mc_end = parallel_on_M ? mc + 1 : Mc;
 
       for (int mci = mc; mci < mc_end; ++mci) {
-        int64_t m_size = mci * block_m + block_m > M ? M - mci * block_m : block_m;
-        // copy bias to y_buf if bias is not None
+        int64_t m_size =
+            mci * block_m + block_m > M ? M - mci * block_m : block_m;
         auto bias_data = bias ? bias + nc * BLOCK_N : nullptr;
         copy_bias<BLOCK_N>(bias_data, C_tmp, m_size);
         for (int kci = 0; kci < Kc; ++kci) {
           int32_t* compensation_ptr =
               sym_quant_act
                   ? nullptr
-                  : (int32_t*)(void*)(weight + (nc * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))) +
-                                      _block_k * BLOCK_N / 2) /*Bcomp*/;
-          _dequant_gemm_accum<BLOCK_N, BLOCK_N / 2, sym_quant_act>(
+                  : (int32_t*)(void*)(weight +
+                                      (nc * Kc + kci) *
+                                          (BLOCK_N *
+                                           (_block_k / 2 + sizeof(int32_t))) +
+                                      _block_k * BLOCK_N / 2);
+          _dequant_gemm_accum<sym_quant_act, BLOCK_N, BLOCK_N / 2>(
               /*C*/ C_tmp,
               /*A*/ (uint8_t*)input + mci * block_m * K + kci * _block_k,
               /*scales_a*/ input_scales + mci * block_m,
               /*qzeros_a*/ input_qzeros + mci * block_m,
-              /*B*/ weight + (nc * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))),
-              /*scales_b*/ weight_scales + nc * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N,
-              /*qzeros_b*/ weight_qzeros + nc * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N,
+              /*B*/ weight + (nc * Kc + kci) *
+                                 (BLOCK_N * (_block_k / 2 + sizeof(int32_t))),
+              /*scales_b*/ weight_scales + nc * BLOCK_N * num_groups +
+                  kci / block_per_group * BLOCK_N,
+              /*qzeros_b*/ weight_qzeros + nc * BLOCK_N * num_groups +
+                  kci / block_per_group * BLOCK_N,
               /*Bcomp*/ compensation_ptr,
               /*dqB_tmp*/ dqB_tmp,
               /*M*/ m_size,
@@ -492,8 +464,8 @@ void _da8w4_linear_impl(
               /*ldc*/ BLOCK_N,
               /*use_brgemm*/ use_brgemm);
         }
-        // store y_buf to output with dtype conversion
-        store_out<out_dtype, BLOCK_N>(C_tmp, output + mci * block_m * N + nc * BLOCK_N, m_size, N /*lda*/);
+        store_out<BLOCK_N>(C_tmp, output + mci * block_m * N + nc * BLOCK_N,
+                           m_size, N /*lda*/);
       }
     }
     if (use_brgemm) {
@@ -504,16 +476,15 @@ void _da8w4_linear_impl(
 
 }  // anonymous namespace
 
-/*
-return: packed_weight, packed_scales, packed_qzeros
-*/
-std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_int4_weight_packed_with_compensation(
-    const at::Tensor& weight, const at::Tensor& scales, const at::Tensor& qzeros) {
-  // weight shape = [N, K]
-  // scales shape = [N, G]
-  // qzeros shape = [N, G]
-  TORCH_CHECK(weight.dim() == 2, "DA8W4 CPU: Weight should be a 2D tensor for packing");
-  TORCH_CHECK(weight.size(1) % 2 == 0, "DA8W4 CPU: Weight should have even number of columns for packing");
+std::tuple<at::Tensor, at::Tensor, at::Tensor>
+convert_int4_weight_packed_with_compensation(const at::Tensor& weight,
+                                             const at::Tensor& scales,
+                                             const at::Tensor& qzeros) {
+  TORCH_CHECK(weight.dim() == 2,
+              "DA8W4 CPU: Weight should be a 2D tensor for packing");
+  TORCH_CHECK(
+      weight.size(1) % 2 == 0,
+      "DA8W4 CPU: Weight should have even number of columns for packing");
 
   auto new_scales = scales;
   auto new_qzeros = qzeros;
@@ -534,21 +505,20 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_int4_weight_packed_with_c
   int64_t Nc = N / block_n;
   int64_t Kc = K / _block_k;
 
-  // Reorder weight to [N/block_n, K/_block_k, _block_k, block_n]
-  // Reorder scales/qzeros to [N/block_n, G, block_n]
-  // weight + compensation shape = [Nc, Kc, block_n * _block_k / 2 + block_n*sizeof(int32_t)]
-  // scales/qzeros shape = [Nc, G, block_n]
   auto weight_view = weight.view({Nc, block_n, Kc, _block_k});
   at::Tensor weight_reordered = weight_view.permute({0, 2, 3, 1}).contiguous();
   at::Tensor blocked_weight;
-  at::Tensor blocked_scales = new_scales.view({Nc, block_n, G}).permute({0, 2, 1}).contiguous();
-  at::Tensor blocked_qzeros = new_qzeros.view({Nc, block_n, G}).permute({0, 2, 1}).contiguous();
-  // Compensation = Σ(k)(W[k][n] - ZP[n]) for each block.
-  auto weight_sub_qzero = weight.view({Nc, block_n, G, -1}).to(at::kInt) - new_qzeros.view({Nc, block_n, G, -1});
+  at::Tensor blocked_scales =
+      new_scales.view({Nc, block_n, G}).permute({0, 2, 1}).contiguous();
+  at::Tensor blocked_qzeros =
+      new_qzeros.view({Nc, block_n, G}).permute({0, 2, 1}).contiguous();
+  auto weight_sub_qzero = weight.view({Nc, block_n, G, -1}).to(at::kInt) -
+                          new_qzeros.view({Nc, block_n, G, -1});
   weight_sub_qzero = weight_sub_qzero.view({Nc, block_n, Kc, _block_k});
   at::Tensor compensation = weight_sub_qzero.sum(-1);
   compensation = compensation.permute({0, 2, 1}).contiguous().to(at::kInt);
-  int64_t buffer_size_nbytes = _block_k * block_n / 2 + block_n * sizeof(int32_t);
+  int64_t buffer_size_nbytes =
+      _block_k * block_n / 2 + block_n * sizeof(int32_t);
   blocked_weight = at::empty({Nc, Kc, buffer_size_nbytes}, weight.options());
 
   auto weight_ptr = weight_reordered.data_ptr<uint8_t>();
@@ -558,25 +528,24 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_int4_weight_packed_with_c
   at::parallel_for(0, num_blocks, 1, [&](int64_t begin, int64_t end) {
     for (const auto i : c10::irange(begin, end)) {
       auto in_ptr = weight_ptr + i * _block_k * block_n;
-      auto out_ptr = blocked_weight_ptr + i * block_n * (_block_k / 2 + sizeof(int32_t));
+      auto out_ptr =
+          blocked_weight_ptr + i * block_n * (_block_k / 2 + sizeof(int32_t));
       int32_t* comp_in_prt = compensation_ptr + i * block_n;
-      int32_t* comp_out_prt = (int32_t*)(void*)(blocked_weight_ptr + i * block_n * (_block_k / 2 + sizeof(int32_t)) +
-                                                _block_k * block_n / 2);
-      // Reorder weight block to VNNI4 and pack two lanes along N
-      // N=16 viewed as two lanes: a0, ...a7, b0, ...b7
-      // pack two lanes: [a0, b0], ..., [a7, b7]
-      // plain shape = [_block_k, block_n]
-      // packed shape = [_block_k / 4, block_n / 2, 4] viewed as [_block_k, block_n / 2]
+      int32_t* comp_out_prt =
+          (int32_t*)(void*)(blocked_weight_ptr +
+                            i * block_n * (_block_k / 2 + sizeof(int32_t)) +
+                            _block_k * block_n / 2);
       constexpr int n_group_size = 8;
       constexpr int vnni_size = 4;
-      constexpr int n_group = block_n / n_group_size;  // 4
+      constexpr int n_group = block_n / n_group_size;
       for (int nb = 0; nb < n_group; nb += 2) {
         for (int k = 0; k < _block_k; k += vnni_size) {
           for (int ni = 0; ni < n_group_size; ++ni) {
             for (int ki = 0; ki < vnni_size; ++ki) {
               int src_idx_1 = nb * n_group_size + ni + (k + ki) * block_n;
               int src_idx_2 = (nb + 1) * n_group_size + ni + (k + ki) * block_n;
-              int dst_idx = (nb / 2 * n_group_size + ni) * vnni_size + k * block_n / 2 + ki;
+              int dst_idx = (nb / 2 * n_group_size + ni) * vnni_size +
+                            k * block_n / 2 + ki;
               uint8_t src_1 = *(in_ptr + src_idx_1);
               uint8_t src_2 = *(in_ptr + src_idx_2);
               uint8_t dst = (src_1 & 0x0f) | ((src_2 & 0x0f) << 4);
@@ -585,128 +554,39 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_int4_weight_packed_with_c
           }
         }
       }
-      // compensation [block_n]
       for (int nb = 0; nb < block_n; nb++) {
         *(comp_out_prt + nb) = *(comp_in_prt + nb);
       }
     }
   });
 
-  return std::make_tuple(std::move(blocked_weight), std::move(blocked_scales), std::move(blocked_qzeros));
+  return std::make_tuple(std::move(blocked_weight), std::move(blocked_scales),
+                         std::move(blocked_qzeros));
 }
 
-std::tuple<at::Tensor, at::Tensor> unpack_4bit_to_32bit_signed(const at::Tensor& qweight, const at::Tensor& qzeros) {
-  TORCH_CHECK(qweight.scalar_type() == at::kInt, "qweight must be int32");
-  TORCH_CHECK(qzeros.scalar_type() == at::kInt, "qzeros must be int32");
-  const auto W0 = qweight.size(0);
-  const auto W1 = qweight.size(1);
-  const auto Z0 = qzeros.size(0);
-  const auto Z1 = qzeros.size(1);
+std::tuple<at::Tensor, at::Tensor> autoawq_to_int4pack(at::Tensor qweight,
+                                                       at::Tensor qzeros) {
+  auto bitshifts = at::tensor({0, 4, 1, 5, 2, 6, 3, 7}, at::kInt) * 4;
+  auto qweight_unsq = qweight.unsqueeze(-1);
+  auto unpacked = at::bitwise_right_shift(qweight_unsq, bitshifts) & 0xF;
+  auto qweight_final = unpacked.flatten(-2).transpose(-1, -2).to(at::kByte);
 
-  // unpacked_weights: (W0 * 8, W1), int8
-  auto unpacked_weights = at::zeros({W0 * 8, W1}, at::TensorOptions().dtype(at::kChar));
-  // unpacked_zeros: (Z0, Z1 * 8), int8
-  auto unpacked_zeros = at::zeros({Z0, Z1 * 8}, at::TensorOptions().dtype(at::kChar));
+  auto qzeros_unsq = qzeros.unsqueeze(-1);
+  auto qzeros_unpacked = at::bitwise_right_shift(qzeros_unsq, bitshifts) & 0xF;
+  auto qzeros_final = qzeros_unpacked.flatten(-2).to(at::kByte);
 
-  const int32_t* qw_ptr = qweight.data_ptr<int32_t>();
-  const int32_t* qz_ptr = qzeros.data_ptr<int32_t>();
-  int8_t* uw_ptr = unpacked_weights.data_ptr<int8_t>();
-  int8_t* uz_ptr = unpacked_zeros.data_ptr<int8_t>();
-
-  // ---- unpack qweight ----
-  for (int64_t row = 0; row < W0 * 8; ++row) {
-    const int i = row & 7;         // row % 8
-    const int src_row = row >> 3;  // row // 8
-    const int shift = 4 * i;
-    for (int64_t col = 0; col < W1; ++col) {
-      int32_t v = qw_ptr[src_row * W1 + col];
-      uw_ptr[row * W1 + col] = static_cast<int8_t>((v >> shift) & 0xF);
-    }
-  }
-  // ---- unpack qzeros ----
-  for (int64_t col = 0; col < Z1 * 8; ++col) {
-    const int i = col & 7;
-    const int src_col = col >> 3;
-    const int shift = 4 * i;
-
-    for (int64_t row = 0; row < Z0; ++row) {
-      int32_t v = qz_ptr[row * Z1 + src_col];
-      uz_ptr[row * (Z1 * 8) + col] = static_cast<int8_t>((v >> shift) & 0xF);
-    }
-  }
-
-  return std::make_tuple(unpacked_weights, unpacked_zeros + 1);
-}
-
-std::tuple<at::Tensor, at::Tensor>
-autogptq_to_int4pack(const at::Tensor& qweight_tensor, const at::Tensor& qzeros_tensor) {
-  TORCH_CHECK(qweight_tensor.scalar_type() == at::kInt, "qweight_tensor must be int32");
-  TORCH_CHECK(qzeros_tensor.scalar_type() == at::kInt, "qzeros_tensor must be int32");
-  TORCH_CHECK(qweight_tensor.is_cpu(), "CPU only implementation");
-  if (qweight_tensor.dim() == 3) {
-    const int64_t B = qweight_tensor.size(0);
-    std::vector<at::Tensor> qweight_list;
-    std::vector<at::Tensor> qzeros_list;
-    qweight_list.reserve(B);
-    qzeros_list.reserve(B);
-    for (int64_t i = 0; i < B; ++i) {
-      auto outputs = unpack_4bit_to_32bit_signed(qweight_tensor[i], qzeros_tensor[i]);
-      at::Tensor unpacked_qweight = std::get<0>(outputs);
-      at::Tensor unpacked_qzeros = std::get<1>(outputs);
-      qweight_list.push_back(unpacked_qweight.transpose(0, 1).contiguous().to(at::kByte));
-      qzeros_list.push_back(unpacked_qzeros.contiguous().to(at::kByte));
-    }
-    return std::make_tuple(at::stack(qweight_list).detach(), at::stack(qzeros_list).detach());
-  }
-  auto outputs = unpack_4bit_to_32bit_signed(qweight_tensor, qzeros_tensor);
-  at::Tensor unpacked_qweight = std::get<0>(outputs);
-  at::Tensor unpacked_qzeros = std::get<1>(outputs);
-  at::Tensor return_qweight = unpacked_qweight.transpose(0, 1).contiguous().to(at::kByte);
-  at::Tensor return_qzeros = unpacked_qzeros.contiguous().to(at::kByte);
-  return std::make_tuple(return_qweight, return_qzeros);
-}
-
-std::tuple<at::Tensor, at::Tensor> int4pack(at::Tensor qweight, at::Tensor qzeros, int64_t quant_method_4bit) {
-  if (quant_method_4bit == CPUQuantAlgo::AWQ) {
-    // autoawq unpacking
-    qweight = qweight.contiguous();
-    qzeros = qzeros.contiguous();
-    // bitshifts: [0, 4, 1, 5, 2, 6, 3, 7] * 4
-    auto bitshifts = at::tensor({0, 4, 1, 5, 2, 6, 3, 7}, at::kInt) * 4;
-    auto qweight_unsq = qweight.unsqueeze(-1);  // [..., K, N/8, 1]
-    auto unpacked = (at::bitwise_right_shift(qweight_unsq, bitshifts) & 0xF).contiguous();
-    auto qweight_final = unpacked.flatten(-2).transpose(-1, -2).to(at::kByte).clone();
-    auto qzeros_unsq = qzeros.unsqueeze(-1);
-    auto qzeros_unpacked = (at::bitwise_right_shift(qzeros_unsq, bitshifts) & 0xF).contiguous();
-    auto qzeros_final = qzeros_unpacked.flatten(-2).to(at::kByte).clone();
-    return std::make_tuple(qweight_final, qzeros_final);
-  } else if (quant_method_4bit == CPUQuantAlgo::GPTQ) {
-    // autogptq unpacking
-    auto outputs = autogptq_to_int4pack(qweight, qzeros);
-    at::Tensor unpacked_qweight = std::get<0>(outputs);
-    at::Tensor unpacked_qzeros = std::get<1>(outputs);
-    return std::make_tuple(unpacked_qweight, unpacked_qzeros);
-  } else {
-    TORCH_CHECK(false, "CPU int4 pack only support AWQ or GPTQ...");
-  }
+  return std::make_tuple(qweight_final, qzeros_final);
 }
 
 std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_weight_packed_scale_zp(
-    at::Tensor qweight,  // awq: (*, K, N / 8)  ||  gptq: (*, K / 8, N) , int32
-    at::Tensor qzeros,   // awq: (*, K / group_size, N / 8) ||  gptq: (*, K / group_size, N / 8) , int32
-    at::Tensor scales,   // awq: (*, K / group_size, N) ||  gptq: (*, K / group_size, N) , bfloat16
-    int64_t quant_method_4bit) {
-  at::Tensor _qweight;
-  at::Tensor _qzeros;
-
-  auto res = int4pack(qweight, qzeros, quant_method_4bit);
-  _qweight = std::get<0>(res);
-  _qzeros = std::get<1>(res);
-
+    at::Tensor qweight, at::Tensor qzeros, at::Tensor scales) {
+  auto res = autoawq_to_int4pack(qweight, qzeros);
+  auto _qweight = std::get<0>(res);
+  auto _qzeros = std::get<1>(res);
   auto _scales = scales;
-  _qzeros = _qzeros.transpose(-2, -1).contiguous();  // .T
+  _qzeros = _qzeros.transpose(-2, -1).contiguous();
   _scales = _scales.transpose(-2, -1).contiguous();
-  if (_qweight.dim() == 3) {  // Dim=3 for MOE packing, TODO: refine a unified loop
+  if (_qweight.dim() == 3) {
     int64_t E = _qweight.size(0);
     int64_t K = _qweight.size(2);
     int64_t G = _scales.size(2);
@@ -715,12 +595,17 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_weight_packed_scale_zp(
     int64_t block_n = block_size_n();
     int64_t Nc = _qweight.size(1) / block_n;
     int64_t Kc = K / _block_k;
-    int64_t buffer_size_nbytes = _block_k * block_n / 2 + block_n * sizeof(int32_t);
-    auto blocked_weight = at::empty({E, Nc, Kc, buffer_size_nbytes}, _qweight.options());
-    auto blocked_scales = at::empty({E, Nc, G, block_n}, _scales.options()).to(at::kFloat);
-    auto blocked_qzeros = at::empty({E, Nc, G, block_n}, _qzeros.options()).to(at::kChar);
+    int64_t buffer_size_nbytes =
+        _block_k * block_n / 2 + block_n * sizeof(int32_t);
+    auto blocked_weight =
+        at::empty({E, Nc, Kc, buffer_size_nbytes}, _qweight.options());
+    auto blocked_scales =
+        at::empty({E, Nc, G, block_n}, _scales.options()).to(at::kFloat);
+    auto blocked_qzeros =
+        at::empty({E, Nc, G, block_n}, _qzeros.options()).to(at::kChar);
     for (int i = 0; i < _qweight.size(0); i++) {
-      auto res_ = convert_int4_weight_packed_with_compensation(_qweight[i], _scales[i], _qzeros[i]);
+      auto res_ = convert_int4_weight_packed_with_compensation(
+          _qweight[i], _scales[i], _qzeros[i]);
       blocked_weight[i] = std::get<0>(res_);
       blocked_scales[i] = std::get<1>(res_);
       blocked_qzeros[i] = std::get<2>(res_);
@@ -729,7 +614,8 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_weight_packed_scale_zp(
     _scales = blocked_scales;
     _qzeros = blocked_qzeros;
   } else {
-    auto res_ = convert_int4_weight_packed_with_compensation(_qweight, _scales, _qzeros);
+    auto res_ = convert_int4_weight_packed_with_compensation(_qweight, _scales,
+                                                             _qzeros);
     _qweight = std::get<0>(res_);
     _scales = std::get<1>(res_);
     _qzeros = std::get<2>(res_);
@@ -738,93 +624,89 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> convert_weight_packed_scale_zp(
   return std::make_tuple(_qweight, _qzeros, _scales);
 }
 
-at::Tensor int4_scaled_mm_cpu_with_quant(
-    const at::Tensor& input,
-    const at::Tensor& weight,
-    const at::Tensor& weight_scales,
-    const at::Tensor& weight_qzeros,
-    const std::optional<at::Tensor>& bias,
-    at::ScalarType output_dtype) {
+at::Tensor int4_scaled_mm_cpu_with_quant(const at::Tensor& input,
+                                         const at::Tensor& weight,
+                                         const at::Tensor& weight_scales,
+                                         const at::Tensor& weight_qzeros,
+                                         const std::optional<at::Tensor>& bias,
+                                         at::ScalarType output_dtype) {
+  RECORD_FUNCTION("vllm::int4_scaled_mm_cpu_with_quant",
+                  std::vector<c10::IValue>({input, weight}));
+
   int64_t M_a = input.size(0);
   int64_t K_a = input.size(1);
   int64_t lda = input.stride(0);
 
   const auto st = input.scalar_type();
   TORCH_CHECK(
-      st == at::kBFloat16 || st == at::kHalf, "int4_scaled_mm_cpu_with_quant: expect A to be bfloat16 or half.");
+      st == at::kBFloat16 || st == at::kHalf,
+      "int4_scaled_mm_cpu_with_quant: expect A to be bfloat16 or half.");
 
-  constexpr bool sym_quant_act = false;  // TODO: add sym quant path
+  constexpr bool sym_quant_act = false;
   using Tin = typename ActDtype<sym_quant_act>::type;
-  int64_t act_buffer_size = /* act quant */ M_a * K_a +
-                            /* act scale */ M_a * sizeof(float) +
-                            /* act zp */ M_a * sizeof(int32_t);
-  auto act_buffer = at::empty({act_buffer_size}, input.options().dtype(at::kByte));
-  // asym path, activation quants into uint8_t
+  int64_t act_buffer_size =
+      M_a * K_a + M_a * sizeof(float) + M_a * sizeof(int32_t);
+  auto act_buffer =
+      at::empty({act_buffer_size}, input.options().dtype(at::kByte));
   auto Aq_data = act_buffer.data_ptr<uint8_t>();
   auto As_data = reinterpret_cast<float*>(Aq_data + M_a * K_a);
   auto Azp_data = reinterpret_cast<int32_t*>(As_data + M_a);
-  fill_val_stub(Azp_data, 128, M_a);  // sym_a s8s8 is unified to u8s8 with compensation (128)
+  fill_val_stub(Azp_data, 128, M_a);
 
   auto out_sizes = input.sizes().vec();
   int64_t N = weight_scales.size(0) * weight_scales.size(-1);
   out_sizes.back() = N;
   auto output = at::empty(out_sizes, input.options());
-  // weight + compensation shape = [Nc, Kc, BLOCK_N * _block_k / 2 + BLOCK_N*sizeof(int32_t)]
-  // scales/qzeros shape = [Nc, G, BLOCK_N]
   int64_t Nc = weight.size(0);
   int64_t Kc = weight.size(1);
   int64_t _block_k = K_a / Kc;
   TORCH_CHECK(N == Nc * BLOCK_N, "DA8W4: weight and input shapes mismatch");
-  // scales/qzeros shape = [Nc, G, BLOCK_N]
   int64_t num_groups = weight_scales.size(1);
 
   const uint8_t* b_ptr = weight.data_ptr<uint8_t>();
   const float* b_scales_ptr = weight_scales.data_ptr<float>();
   const int8_t* b_qzeros_ptr = weight_qzeros.data_ptr<int8_t>();
-  const float* bias_ptr = bias.has_value() ? bias.value().data_ptr<float>() : nullptr;
+  const float* bias_ptr =
+      bias.has_value() ? bias.value().data_ptr<float>() : nullptr;
   int num_threads = at::get_num_threads();
-  int64_t temp_buffer_size = /* output temp */ num_threads * BLOCK_M * BLOCK_N * sizeof(float) +
-                             /*  weight dequant temp */ num_threads * _block_k * BLOCK_N;
-  auto c_temp_buffer = at::empty({temp_buffer_size}, input.options().dtype(at::kChar));
+  int64_t temp_buffer_size = num_threads * BLOCK_M * BLOCK_N * sizeof(float) +
+                             num_threads * _block_k * BLOCK_N;
+  auto c_temp_buffer =
+      at::empty({temp_buffer_size}, input.options().dtype(at::kChar));
   float* c_temp_ptr = (float*)((void*)(c_temp_buffer.data_ptr<int8_t>()));
-  int8_t* dqB_temp_ptr = (int8_t*)((void*)(c_temp_ptr + num_threads * BLOCK_M * BLOCK_N));
+  int8_t* dqB_temp_ptr =
+      (int8_t*)((void*)(c_temp_ptr + num_threads * BLOCK_M * BLOCK_N));
 
-#define LAUNCH_DA8W4_LINEAR_WITH_QUANT_IMPL(sym_quant_act)                                                 \
-  AT_DISPATCH_FLOATING_TYPES_AND2(                                                                         \
-      at::ScalarType::BFloat16, at::ScalarType::Half, output_dtype, "int4_scaled_mm_cpu_with_quant", [&] { \
-        const scalar_t* __restrict__ A_data = input.data_ptr<scalar_t>();                                  \
-        scalar_t* __restrict__ c_ptr = output.data_ptr<scalar_t>();                                        \
-        at::parallel_for(0, M_a, 0, [&](int64_t begin, int64_t end) {                                      \
-          for (int64_t m = begin; m < end; ++m) {                                                          \
-            quantize_row_int8<scalar_t>(Aq_data + m * K_a, As_data[m], A_data + m * lda, K_a);             \
-          }                                                                                                \
-        });                                                                                                \
-        _da8w4_linear_impl<Tin, scalar_t, sym_quant_act>(                                                  \
-            Aq_data,                                                                                       \
-            As_data,                                                                                       \
-            Azp_data,                                                                                      \
-            b_ptr,                                                                                         \
-            b_scales_ptr,                                                                                  \
-            b_qzeros_ptr,                                                                                  \
-            bias_ptr,                                                                                      \
-            c_ptr,                                                                                         \
-            c_temp_ptr,                                                                                    \
-            dqB_temp_ptr,                                                                                  \
-            M_a,                                                                                           \
-            N,                                                                                             \
-            K_a,                                                                                           \
-            num_groups);                                                                                   \
+#define LAUNCH_DA8W4_LINEAR_WITH_QUANT_IMPL(sym_quant_act)                 \
+  AT_DISPATCH_FLOATING_TYPES_AND2(                                         \
+      at::ScalarType::BFloat16, at::ScalarType::Half, output_dtype,        \
+      "int4_scaled_mm_cpu", [&] {                                          \
+        const scalar_t* __restrict__ A_data = input.data_ptr<scalar_t>();  \
+        scalar_t* __restrict__ c_ptr = output.data_ptr<scalar_t>();        \
+        at::parallel_for(0, M_a, 0, [&](int64_t begin, int64_t end) {      \
+          for (int64_t m = begin; m < end; ++m) {                          \
+            quantize_row_int8<scalar_t>(Aq_data + m * K_a, As_data[m],     \
+                                        A_data + m * lda, K_a);            \
+          }                                                                \
+        });                                                                \
+        _da8w4_linear_impl<sym_quant_act, Tin, scalar_t>(                  \
+            Aq_data, As_data, Azp_data, b_ptr, b_scales_ptr, b_qzeros_ptr, \
+            bias_ptr, c_ptr, c_temp_ptr, dqB_temp_ptr, M_a, N, K_a,        \
+            num_groups);                                                   \
       });
 
   LAUNCH_DA8W4_LINEAR_WITH_QUANT_IMPL(sym_quant_act);
 
   return output;
 }
+
+namespace {
+
 template <typename scalar_t>
-inline void copy_stub(scalar_t* __restrict__ out, const float* __restrict__ input, int64_t size) {
+inline void copy_stub(scalar_t* __restrict__ out,
+                      const float* __restrict__ input, int64_t size) {
   using Vec = at::vec::Vectorized<scalar_t>;
   using fVec = at::vec::Vectorized<float>;
-// no remainder
 #pragma GCC unroll 4
   for (int64_t d = 0; d < size; d += Vec::size()) {
     fVec x0 = fVec::loadu(input + d);
@@ -834,61 +716,40 @@ inline void copy_stub(scalar_t* __restrict__ out, const float* __restrict__ inpu
   }
 }
 
+}  // anonymous namespace
+
 template <typename scalar_t>
-void tinygemm_kernel(
-    scalar_t* C,
-    float* C_temp,
-    const uint8_t* A,
-    const float* scales_a,
-    const int32_t* qzeros_a,
-    const uint8_t* B,
-    const float* scales_b,
-    const int8_t* qzeros_b,
-    const int32_t* compensation,
-    int8_t* dqB_tmp,
-    int64_t M,
-    int64_t K,
-    int64_t lda,
-    int64_t ldc_f,
-    int64_t ldc_s,
-    bool store_out,
-    bool use_brgemm) {
-  // TODO: add sym quant act, now only asym
-  _dequant_gemm_accum<BLOCK_N, BLOCK_N / 2, false>(
-      C_temp, A, scales_a, qzeros_a, B, scales_b, qzeros_b, compensation, dqB_tmp, M, K, lda, ldc_f, use_brgemm);
+void tinygemm_kernel(scalar_t* C, float* C_temp, const uint8_t* A,
+                     const float* scales_a, const int32_t* qzeros_a,
+                     const uint8_t* B, const float* scales_b,
+                     const int8_t* qzeros_b, const int32_t* compensation,
+                     int8_t* dqB_tmp, int64_t M, int64_t K, int64_t lda,
+                     int64_t ldc_f, int64_t ldc_s, bool store_out,
+                     bool use_brgemm) {
+  _dequant_gemm_accum<false, BLOCK_N, BLOCK_N / 2>(
+      C_temp, A, scales_a, qzeros_a, B, scales_b, qzeros_b, compensation,
+      dqB_tmp, M, K, lda, ldc_f, use_brgemm);
   if (store_out) {
-    // copy from Ctmp to C
     for (int64_t m = 0; m < M; ++m) {
       copy_stub<scalar_t>(C + m * ldc_s, C_temp + m * ldc_f, BLOCK_N);
     }
   }
 }
 
-#define INSTANTIATE_TINYGEMM_TEMPLATE(TYPE) \
-  template void tinygemm_kernel<TYPE>(      \
-      TYPE * C,                             \
-      float* C_temp,                        \
-      const uint8_t* A,                     \
-      const float* scales_a,                \
-      const int32_t* qzeros_a,              \
-      const uint8_t* B,                     \
-      const float* scales_b,                \
-      const int8_t* qzeros_b,               \
-      const int32_t* compensation,          \
-      int8_t* dqB_tmp,                      \
-      int64_t M,                            \
-      int64_t K,                            \
-      int64_t lda,                          \
-      int64_t ldc_f,                        \
-      int64_t ldc_s,                        \
-      bool store_out,                       \
-      bool use_brgemm)
+#define INSTANTIATE_TINYGEMM_TEMPLATE(TYPE)                                 \
+  template void tinygemm_kernel<TYPE>(                                      \
+      TYPE * C, float* C_temp, const uint8_t* A, const float* scales_a,     \
+      const int32_t* qzeros_a, const uint8_t* B, const float* scales_b,     \
+      const int8_t* qzeros_b, const int32_t* compensation, int8_t* dqB_tmp, \
+      int64_t M, int64_t K, int64_t lda, int64_t ldc_f, int64_t ldc_s,      \
+      bool store_out, bool use_brgemm)
 
 INSTANTIATE_TINYGEMM_TEMPLATE(at::BFloat16);
 INSTANTIATE_TINYGEMM_TEMPLATE(at::Half);
 
-// int4 gemm dispatch api register
-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) {
-  return int4_scaled_mm_cpu_with_quant(x, w, w_scales, w_zeros, bias, x.scalar_type());
+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) {
+  return int4_scaled_mm_cpu_with_quant(x, w, w_scales, w_zeros, bias,
+                                       x.scalar_type());
 }

csrc/cpu/sgl-kernels/gemm_int8.cpp

@@ -1,83 +1,20 @@
 // Adapted from
 // https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
 
+#include "common.h"
+#include "vec.h"
+#include "gemm.h"
+
 // clang-format off
 
-#include "common.h"
-#include "gemm.h"
-#include "vec.h"
-
 namespace {
 
-template <typename scalar_t, bool has_bias, int BLOCK_N>
-struct scale_C {
-  static inline void apply(
-      scalar_t* __restrict__ C,
-      const int32_t* __restrict__ Ctmp,
-      const int32_t* __restrict__ Bcomp,
-      const float* __restrict__ bias,
-      float As,
-      const float* __restrict__ Bs) {
-    TORCH_CHECK(false, "scale_C: scalar path not implemented!");
-  }
-};
-
-#if defined(CPU_CAPABILITY_AVX512)
-template <bool has_bias, int BLOCK_N>
-struct scale_C<at::BFloat16, has_bias, BLOCK_N> {
-  static inline void apply(
-      at::BFloat16* __restrict__ C,
-      const int32_t* __restrict__ Ctmp,
-      const int32_t* __restrict__ Bcomp,
-      const float* __restrict__ bias,
-      float As,
-      const float* __restrict__ Bs) {
-    constexpr int COLS = BLOCK_N / 16;
-    static_assert(COLS % 2 == 0);
-
-    __m512 vc[COLS];
-    __m512 vd0 = _mm512_set1_ps(As);
-
-    auto compute = [&](auto col) {
-      __m512 vd1 = _mm512_loadu_ps(Bs + col * 16);
-      __m512i vcomp = _mm512_loadu_si512(Bcomp + col * 16);
-      __m512i vc32 = _mm512_loadu_si512(Ctmp + col * 16);
-      vc[col] = _mm512_cvtepi32_ps(_mm512_sub_epi32(vc32, vcomp));
-      if constexpr (has_bias) {
-        __m512 vbias = _mm512_loadu_ps(bias + col * 16);
-        vc[col] = _mm512_fmadd_ps(_mm512_mul_ps(vc[col], vd0), vd1, vbias);
-      } else {
-        vc[col] = _mm512_mul_ps(_mm512_mul_ps(vc[col], vd0), vd1);
-      }
-    };
-    Unroll<COLS>{}(compute);
-
-    auto storec = [&](auto col) {
-      // for COLS = 2, 4 use 512bit store
-      if constexpr (col % 2 == 0) {
-        _mm512_storeu_si512(
-            reinterpret_cast<__m512i*>((C + col * 16)), (__m512i)(_mm512_cvtne2ps_pbh(vc[col + 1], vc[col + 0])));
-      }
-    };
-    Unroll<COLS>{}(storec);
-  }
-};
-#endif
-
 template <typename scalar_t, bool has_bias, int BLOCK_M, int BLOCK_N>
 struct tinygemm_kernel_nn {
   static inline void apply(
-      const uint8_t* __restrict__ A,
-      const int8_t* __restrict__ B,
-      scalar_t* __restrict__ C,
-      const float* __restrict__ As,
-      const float* __restrict__ Bs,
-      const int32_t* __restrict__ Bcomp,
-      const float* __restrict__ bias,
-      int64_t K,
-      int64_t lda,
-      int64_t ldb,
-      int64_t ldc) {
+      const uint8_t* __restrict__ A, const int8_t* __restrict__ B, scalar_t* __restrict__ C,
+      const float* __restrict__ As, const float* __restrict__ Bs, const int32_t* __restrict__ Bcomp,
+      const float* __restrict__ bias, int64_t K, int64_t lda, int64_t ldb, int64_t ldc) {
     TORCH_CHECK(false, "tinygemm_kernel_nn: scalar path not implemented!");
   }
 };
@@ -86,17 +23,10 @@ struct tinygemm_kernel_nn {
 template <bool has_bias, int BLOCK_M, int BLOCK_N>
 struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
   static inline void apply(
-      const uint8_t* __restrict__ A,
-      const int8_t* __restrict__ B,
-      at::BFloat16* __restrict__ C,
-      const float* __restrict__ As,
-      const float* __restrict__ Bs,
-      const int32_t* __restrict__ Bcomp,
-      const float* __restrict__ bias,
-      int64_t K,
-      int64_t lda,
-      int64_t ldb,
-      int64_t ldc) {
+      const uint8_t* __restrict__ A, const int8_t* __restrict__ B, at::BFloat16* __restrict__ C,
+      const float* __restrict__ As, const float* __restrict__ Bs, const int32_t* __restrict__ Bcomp,
+      const float* __restrict__ bias, int64_t K, int64_t lda, int64_t ldb, int64_t ldc) {
+
     constexpr int ROWS = BLOCK_M;
     constexpr int COLS = BLOCK_N / 16;
     static_assert(COLS % 2 == 0);
@@ -108,10 +38,10 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
     __m512i vb[COLS];
     __m512i vc[ROWS * COLS];
     __m512i vcomp[COLS];
-    __m512 vd0;
-    __m512 vd1[COLS];
+    __m512  vd0;
+    __m512  vd1[COLS];
 
-    // oops! 4x4 spills but we use 4x2
+    // oops! 4x4 spills but luckily we use 4x2
     __m512 vbias[COLS];
 
     // [NOTE]: s8s8 igemm compensation in avx512-vnni
@@ -124,12 +54,14 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
     // 1) 128 * b is pre-computed when packing B to vnni formats
     // 2) a + 128 is fused when dynamically quantize A
     //
-    auto loadc = [&](auto i) { vc[i] = _mm512_set1_epi32(0); };
+    auto loadc = [&](auto i) {
+      vc[i] = _mm512_set1_epi32(0);
+    };
     Unroll<ROWS * COLS>{}(loadc);
 
     const int64_t K4 = K >> 2;
     const int64_t lda4 = lda >> 2;
-    const int64_t ldb4 = ldb;  // ldb * 4 >> 2;
+    const int64_t ldb4 = ldb; // ldb * 4 >> 2;
     const int32_t* a_ptr = reinterpret_cast<const int32_t*>(A);
     const int32_t* b_ptr = reinterpret_cast<const int32_t*>(B);
 
@@ -157,7 +89,7 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
       constexpr int col = i % COLS;
 
       // load a scale
-      if constexpr (col == 0) {
+      if constexpr(col == 0) {
         vd0 = _mm512_set1_ps(As[row]);
       }
       // load b scale and vcomp per 2 vectors
@@ -188,7 +120,8 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
         }
 
         _mm512_storeu_si512(
-            reinterpret_cast<__m512i*>((C + row * ldc + col * 16)), (__m512i)(_mm512_cvtne2ps_pbh(vc1, vc0)));
+            reinterpret_cast<__m512i*>((C + row * ldc + col * 16)),
+            (__m512i)(_mm512_cvtne2ps_pbh(vc1, vc0)));
       }
     };
     Unroll<ROWS * COLS>{}(storec);
@@ -196,19 +129,11 @@ struct tinygemm_kernel_nn<at::BFloat16, has_bias, BLOCK_M, BLOCK_N> {
 };
 #endif
 
-#define LAUNCH_TINYGEMM_KERNEL_NN(MB_SIZE, NB_SIZE)                \
-  tinygemm_kernel_nn<scalar_t, has_bias, MB_SIZE, NB_SIZE>::apply( \
-      A + mb_start * lda,                                          \
-      B + nb_start * 4,                                            \
-      C + mb_start * ldc + nb_start,                               \
-      As + mb_start,                                               \
-      Bs + nb_start,                                               \
-      Bcomp + nb_start,                                            \
-      has_bias ? bias + nb_start : nullptr,                        \
-      K,                                                           \
-      lda,                                                         \
-      ldb,                                                         \
-      ldc);
+#define LAUNCH_TINYGEMM_KERNEL_NN(MB_SIZE, NB_SIZE)                          \
+    tinygemm_kernel_nn<scalar_t, has_bias, MB_SIZE, NB_SIZE>::apply(         \
+        A + mb_start * lda, B + nb_start * 4, C + mb_start * ldc + nb_start, \
+        As + mb_start, Bs + nb_start, Bcomp + nb_start,                      \
+        has_bias ? bias + nb_start : nullptr, K, lda, ldb, ldc);
 
 template <typename scalar_t, bool has_bias>
 void tinygemm_kernel(
@@ -226,20 +151,10 @@ void tinygemm_kernel(
     int64_t ldb,
     int64_t ldc,
     bool brg) {
+
   // B compensation
   const int32_t* Bcomp = reinterpret_cast<const int32_t*>(B + block_size_n() * K);
 
-  if (brg) {
-    constexpr int BLOCK_N = block_size_n();
-    at::native::cpublas::brgemm(M, N, K, lda, ldb, BLOCK_N, /* add_C */ false, A, B, Ctmp);
-
-    // apply compensation and scale
-    for (int64_t m = 0; m < M; ++m) {
-      scale_C<scalar_t, has_bias, BLOCK_N>::apply(C + m * ldc, Ctmp + m * BLOCK_N, Bcomp, bias, As[m], Bs);
-    }
-    return;
-  }
-
   // pattern: 1-4-16
   constexpr int64_t BLOCK_M = 4;
   constexpr int64_t BLOCK_N = 64;
@@ -252,43 +167,26 @@ void tinygemm_kernel(
       int64_t nb_start = nb * BLOCK_N;
       int64_t nb_size = std::min(BLOCK_N, N - nb_start);
 
-      switch (mb_size << 4 | nb_size >> 4) {
+      switch(mb_size << 4 | nb_size >> 4) {
         // mb_size = 1
-        case 0x12:
-          LAUNCH_TINYGEMM_KERNEL_NN(1, 32);
-          break;
-        case 0x14:
-          LAUNCH_TINYGEMM_KERNEL_NN(1, 64);
-          break;
+        case 0x12: LAUNCH_TINYGEMM_KERNEL_NN(1, 32); break;
+        case 0x14: LAUNCH_TINYGEMM_KERNEL_NN(1, 64); break;
         // mb_size = 2
-        case 0x22:
-          LAUNCH_TINYGEMM_KERNEL_NN(2, 32);
-          break;
-        case 0x24:
-          LAUNCH_TINYGEMM_KERNEL_NN(2, 64);
-          break;
+        case 0x22: LAUNCH_TINYGEMM_KERNEL_NN(2, 32); break;
+        case 0x24: LAUNCH_TINYGEMM_KERNEL_NN(2, 64); break;
         // mb_size = 3
-        case 0x32:
-          LAUNCH_TINYGEMM_KERNEL_NN(3, 32);
-          break;
-        case 0x34:
-          LAUNCH_TINYGEMM_KERNEL_NN(3, 64);
-          break;
+        case 0x32: LAUNCH_TINYGEMM_KERNEL_NN(3, 32); break;
+        case 0x34: LAUNCH_TINYGEMM_KERNEL_NN(3, 64); break;
         // mb_size = 4
-        case 0x42:
-          LAUNCH_TINYGEMM_KERNEL_NN(4, 32);
-          break;
-        case 0x44:
-          LAUNCH_TINYGEMM_KERNEL_NN(4, 64);
-          break;
-        default:
-          TORCH_CHECK(false, "Unexpected block size, ", mb_size, "x", "nb_size");
+        case 0x42: LAUNCH_TINYGEMM_KERNEL_NN(4, 32); break;
+        case 0x44: LAUNCH_TINYGEMM_KERNEL_NN(4, 64); break;
+        default: TORCH_CHECK(false, "Unexpected block size, ", mb_size, "x", nb_size);
       }
     }
   }
 }
 
-template <typename scalar_t>
+template<typename scalar_t>
 void int8_scaled_mm_kernel_impl(
     scalar_t* __restrict__ out,
     const uint8_t* __restrict__ mat1,
@@ -299,22 +197,28 @@ void int8_scaled_mm_kernel_impl(
     int64_t M,
     int64_t N,
     int64_t K) {
+
   constexpr int64_t BLOCK_M = block_size_m();
   constexpr int64_t BLOCK_N = block_size_n();
   const int64_t MB = div_up(M, BLOCK_M);
   const int64_t NB = div_up(N, BLOCK_N);
 
-  const bool use_brgemm = can_use_brgemm<int8_t>(M);
+  // TODO: brgemm u8s8 depends on PyTorch 2.7 release.
+  const bool use_brgemm = false;
 
   // K + 4 after compensation
   const int64_t packed_row_size = get_row_size<int8_t>(K);
 
   AT_DISPATCH_BOOL(bias != nullptr, has_bias, [&] {
-    parallel_2d(MB, NB, [&](int64_t mb0, int64_t mb1, int64_t nb0, int64_t nb1) {
+    at::parallel_for(0, MB * NB, 0, [&](int64_t begin, int64_t end) {
+      int64_t mb{0}, nb{0};
+      data_index_init(begin, mb, MB, nb, NB);
+
       // for brgemm, use int32_t for accumulate
       alignas(64) int32_t Ctmp[BLOCK_M * BLOCK_N];
 
-      loop_2d<int8_t>(mb0, mb1, nb0, nb1, BLOCK_N * K, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
+      for (int i = begin; i < end; ++i) {
+        UNUSED(i);
         int mb_start = mb * BLOCK_M;
         int mb_size = std::min(M - mb_start, BLOCK_M);
         int nb_start = nb * BLOCK_N;
@@ -335,7 +239,10 @@ void int8_scaled_mm_kernel_impl(
             /* ldb */ nb_size,
             /* ldc */ N,
             /* brg */ use_brgemm);
-      });
+
+        // move to the next index
+        data_index_step(mb, MB, nb, NB);
+      }
 
       if (use_brgemm) {
         at::native::cpublas::brgemm_release();
@@ -344,47 +251,28 @@ void int8_scaled_mm_kernel_impl(
   });
 }
 
-}  // anonymous namespace
+} // anonymous namespace
 
 // tinygemm interface
 template <typename scalar_t>
-void tinygemm_kernel(
-    const uint8_t* __restrict__ A,
-    const int8_t* __restrict__ B,
-    scalar_t* __restrict__ C,
-    int32_t* __restrict__ Ctmp,
-    const float* __restrict__ As,
-    const float* __restrict__ Bs,
-    int64_t M,
-    int64_t N,
-    int64_t K,
-    int64_t lda,
-    int64_t ldb,
-    int64_t ldc,
-    bool brg) {
+void tinygemm_kernel(const uint8_t* __restrict__ A, const int8_t* __restrict__ B, scalar_t* __restrict__ C,
+    int32_t* __restrict__ Ctmp,  const float* __restrict__ As, const float* __restrict__ Bs,
+    int64_t M, int64_t N, int64_t K, int64_t lda, int64_t ldb, int64_t ldc, bool brg) {
   tinygemm_kernel<scalar_t, false>(A, B, C, Ctmp, As, Bs, nullptr, M, N, K, lda, ldb, ldc, brg);
 }
 
-#define INSTANTIATE_TINYGEMM_TEMPLATE(TYPE) \
-  template void tinygemm_kernel<TYPE>(      \
-      const uint8_t* __restrict__ A,        \
-      const int8_t* __restrict__ B,         \
-      TYPE* __restrict__ C,                 \
-      int32_t* __restrict__ Ctmp,           \
-      const float* __restrict__ As,         \
-      const float* __restrict__ Bs,         \
-      int64_t M,                            \
-      int64_t N,                            \
-      int64_t K,                            \
-      int64_t lda,                          \
-      int64_t ldb,                          \
-      int64_t ldc,                          \
-      bool brg)
+#define INSTANTIATE_TINYGEMM_TEMPLATE(TYPE)                                                     \
+    template void tinygemm_kernel<TYPE>(                                                        \
+        const uint8_t* __restrict__ A, const int8_t* __restrict__ B, TYPE* __restrict__ C,      \
+        int32_t* __restrict__ Ctmp, const float* __restrict__ As, const float* __restrict__ Bs, \
+        int64_t M, int64_t N, int64_t K, int64_t lda, int64_t ldb, int64_t ldc, bool brg)
 
 INSTANTIATE_TINYGEMM_TEMPLATE(at::BFloat16);
 INSTANTIATE_TINYGEMM_TEMPLATE(at::Half);
 
 std::tuple<at::Tensor, at::Tensor> per_token_quant_int8_cpu(at::Tensor& A) {
+  RECORD_FUNCTION("sgl-kernel::per_token_quant_int8_cpu", std::vector<c10::IValue>({A}));
+
   CHECK_LAST_DIM_CONTIGUOUS_INPUT(A);
   CHECK_DIM(2, A);
 
@@ -393,7 +281,8 @@ std::tuple<at::Tensor, at::Tensor> per_token_quant_int8_cpu(at::Tensor& A) {
   int64_t lda = A.stride(0);
 
   const auto st = A.scalar_type();
-  TORCH_CHECK(st == at::kBFloat16 || st == at::kHalf, "per_token_quant_int8: expect A to be bfloat16 or half.");
+  TORCH_CHECK(st == at::kBFloat16 || st == at::kHalf,
+      "per_token_quant_int8: expect A to be bfloat16 or half.");
 
   auto Aq = at::empty({M, K}, A.options().dtype(at::kByte));
   auto As = at::empty({M}, A.options().dtype(at::kFloat));
@@ -403,9 +292,13 @@ std::tuple<at::Tensor, at::Tensor> per_token_quant_int8_cpu(at::Tensor& A) {
     float* __restrict__ As_data = As.data_ptr<float>();
     const scalar_t* __restrict__ A_data = A.data_ptr<scalar_t>();
 
-    at::parallel_for(0, M, 0, [&](int64_t begin, int64_t end) {
+    at::parallel_for(0, M, 0, [&] (int64_t begin, int64_t end) {
       for (int64_t m = begin; m < end; ++m) {
-        quantize_row_int8<scalar_t>(Aq_data + m * K, As_data[m], A_data + m * lda, K);
+        quantize_row_int8<scalar_t>(
+            Aq_data + m * K,
+            As_data[m],
+            A_data + m * lda,
+            K);
       }
     });
   });
@@ -422,14 +315,11 @@ std::tuple<at::Tensor, at::Tensor> per_token_quant_int8_cpu(at::Tensor& A) {
 // bias    : [N]
 // out     : [M, N]
 //
-at::Tensor int8_scaled_mm_cpu(
-    at::Tensor& mat1,
-    at::Tensor& mat2,
-    at::Tensor& scales1,
-    at::Tensor& scales2,
-    const std::optional<at::Tensor>& bias,
-    at::ScalarType out_dtype,
-    bool is_vnni) {
+at::Tensor int8_scaled_mm_cpu(at::Tensor& mat1, at::Tensor& mat2,
+    at::Tensor& scales1, at::Tensor& scales2,
+    std::optional<at::Tensor>& bias, at::ScalarType out_dtype, bool is_vnni) {
+  RECORD_FUNCTION("sgl-kernel::int8_scaled_mm_cpu", std::vector<c10::IValue>({mat1, mat2, scales1, scales2, bias}));
+
   auto packed_w = is_vnni ? mat2 : convert_weight_packed(mat2);
 
   CHECK_INPUT(mat1);
@@ -450,8 +340,7 @@ at::Tensor int8_scaled_mm_cpu(
 
   TORCH_CHECK(mat1.scalar_type() == at::kByte, "int8_scaled_mm: expect mat1 to be uint8.");
   TORCH_CHECK(mat2.scalar_type() == at::kChar, "int8_scaled_mm: expect mat2 to be int8.");
-  TORCH_CHECK(
-      scales1.scalar_type() == at::kFloat && scales2.scalar_type() == at::kFloat,
+  TORCH_CHECK(scales1.scalar_type() == at::kFloat && scales2.scalar_type() == at::kFloat,
       "int8_scaled_mm: expect scales to be float32.");
 
   auto out = at::empty({M, N}, mat1.options().dtype(out_dtype));
@@ -479,13 +368,10 @@ at::Tensor int8_scaled_mm_cpu(
 }
 
 // fused `per_token_quant_int8_cpu` and `int8_scaled_mm_cpu`
-at::Tensor int8_scaled_mm_with_quant(
-    at::Tensor& mat1,
-    at::Tensor& mat2,
-    at::Tensor& scales2,
-    const std::optional<at::Tensor>& bias,
-    at::ScalarType out_dtype,
-    bool is_vnni) {
+at::Tensor int8_scaled_mm_with_quant(at::Tensor& mat1, at::Tensor& mat2, at::Tensor& scales2,
+    const std::optional<at::Tensor>& bias, at::ScalarType out_dtype, bool is_vnni) {
+  RECORD_FUNCTION("sgl-kernel::int8_scaled_mm_cpu", std::vector<c10::IValue>({mat1, mat2, scales2, bias}));
+
   auto packed_w = is_vnni ? mat2 : convert_weight_packed(mat2);
 
   CHECK_LAST_DIM_CONTIGUOUS_INPUT(mat1);
@@ -504,10 +390,14 @@ at::Tensor int8_scaled_mm_with_quant(
   CHECK_EQ(scales2.numel(), N);
 
   const auto st = mat1.scalar_type();
-  TORCH_CHECK(st == at::kBFloat16 || st == at::kHalf, "int8_scaled_mm_with_quant: expect A to be bfloat16 or half.");
-  TORCH_CHECK(st == out_dtype, "int8_scaled_mm_with_quant: expect A has same dtype with out_dtype.");
-  TORCH_CHECK(mat2.scalar_type() == at::kChar, "int8_scaled_mm_with_quant: expect mat2 to be int8.");
-  TORCH_CHECK(scales2.scalar_type() == at::kFloat, "int8_scaled_mm_with_quant: expect scales to be float32.");
+  TORCH_CHECK(st == at::kBFloat16 || st == at::kHalf,
+      "int8_scaled_mm_with_quant: expect A to be bfloat16 or half.");
+  TORCH_CHECK(st == out_dtype,
+      "int8_scaled_mm_with_quant: expect A has same dtype with out_dtype.");
+  TORCH_CHECK(mat2.scalar_type() == at::kChar,
+      "int8_scaled_mm_with_quant: expect mat2 to be int8.");
+  TORCH_CHECK(scales2.scalar_type() == at::kFloat,
+      "int8_scaled_mm_with_quant: expect scales to be float32.");
 
   const int64_t buffer_size = M * K + M * sizeof(float);
   auto buffer = at::empty({buffer_size}, mat1.options().dtype(at::kByte));
@@ -525,9 +415,13 @@ at::Tensor int8_scaled_mm_with_quant(
     float* __restrict__ As_data = (float*)((void*)(Aq_data + M * K));
     const scalar_t* __restrict__ A_data = mat1.data_ptr<scalar_t>();
 
-    at::parallel_for(0, M, 0, [&](int64_t begin, int64_t end) {
+    at::parallel_for(0, M, 0, [&] (int64_t begin, int64_t end) {
       for (int64_t m = begin; m < end; ++m) {
-        quantize_row_int8<scalar_t>(Aq_data + m * K, As_data[m], A_data + m * lda, K);
+        quantize_row_int8<scalar_t>(
+            Aq_data + m * K,
+            As_data[m],
+            A_data + m * lda,
+            K);
       }
     });
 

csrc/cpu/sgl-kernels/moe.h

@@ -1,284 +0,0 @@
-// Adapted from
-// https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
-
-// clang-format off
-
-#pragma once
-#include "vec.h"
-
-template <typename scalar_t>
-inline void fill_stub(scalar_t* __restrict__ out, scalar_t val, int64_t size) {
-  using Vec = at::vec::Vectorized<scalar_t>;
-  const Vec data_vec(val);
-  at::vec::map<scalar_t>([data_vec](Vec out) { return out = data_vec; }, out, out, size);
-}
-
-template <typename scalar_t>
-inline void copy_stub(scalar_t* __restrict__ out, const scalar_t* __restrict__ input, int64_t size) {
-  using Vec = at::vec::Vectorized<scalar_t>;
-  constexpr int kVecSize = Vec::size();
-  int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    Vec data = Vec::loadu(input + d);
-    data.store(out + d);
-  }
-  for (; d < size; ++d) {
-    out[d] = input[d];
-  }
-}
-
-template <typename scalar_t>
-inline void copy_stub(scalar_t* __restrict__ out, const float* __restrict__ input, int64_t size) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = bVec::size();
-  int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    auto [x0, x1] = load_float_vec2(input + d);
-    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]);
-  }
-}
-
-template <>
-inline void copy_stub<uint8_t>(uint8_t* __restrict__ out, const uint8_t* __restrict__ input, int64_t size) {
-  // size might be 64x + 32
-  std::memcpy(out, input, size * sizeof(uint8_t));
-}
-
-template <typename scalar_t, typename input_t>
-inline void copy_mul_stub(scalar_t* __restrict__ out, const input_t* __restrict__ input, float weight, int64_t size) {
-  static_assert(
-      std::is_same_v<input_t, float> || std::is_same_v<input_t, scalar_t>,
-      "copy_mul_stub only supports input_t == float or input_t == scalar_t");
-  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) {
-    auto [x0, x1] = load_float_vec2(input + d);
-    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);
-  }
-}
-
-// acc from [topk, K] to [K]
-template <typename scalar_t>
-inline void sum_stub(scalar_t* __restrict__ out, const scalar_t* __restrict__ input, int64_t topk, int64_t K) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = bVec::size();
-  if (topk == 1) {
-    // do copy for topk = 1
-    copy_stub(out, input, K);
-  } else {
-    // do sum for topk != 1
-    int64_t d;
-#pragma GCC unroll 4
-    for (d = 0; d <= K - kVecSize; d += kVecSize) {
-      fVec sum_fvec0 = fVec(0.f);
-      fVec sum_fvec1 = fVec(0.f);
-      for (int t = 0; t < topk; ++t) {
-        bVec x_bvec = bVec::loadu(input + t * K + d);
-        fVec x_fvec0, x_fvec1;
-        std::tie(x_fvec0, x_fvec1) = at::vec::convert_to_float(x_bvec);
-
-        sum_fvec0 += x_fvec0;
-        sum_fvec1 += x_fvec1;
-      }
-      bVec out_bvec = convert_from_float_ext<scalar_t>(sum_fvec0, sum_fvec1);
-      out_bvec.store(out + d);
-    }
-    for (; d < K; ++d) {
-      float sum_val = 0.f;
-      for (int t = 0; t < topk; ++t) {
-        sum_val += static_cast<float>(input[t * K + d]);
-      }
-      out[d] = static_cast<scalar_t>(sum_val);
-    }
-  }
-}
-
-// out = input + input2 * scale
-template <typename scalar_t, typename input_t>
-inline void add_mul_stub(
-    scalar_t* __restrict__ out,
-    const input_t* __restrict__ input,
-    const scalar_t* __restrict__ input2,
-    float scale,
-    int64_t size) {
-  static_assert(
-      std::is_same_v<input_t, float> || std::is_same_v<input_t, scalar_t>,
-      "add_mul_stub only supports input_t == float or input_t == scalar_t");
-
-  // out = input (without scale factor)
-  if (input2 == nullptr) {
-    copy_stub(out, input, size);
-    return;
-  }
-
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = bVec::size();
-  const fVec s_vec = fVec(scale);
-  int64_t d;
-#pragma GCC unroll 4
-  for (d = 0; d <= size - kVecSize; d += kVecSize) {
-    auto [x0, x1] = load_float_vec2(input + d);
-
-    bVec y_bvec = bVec::loadu(input2 + d);
-    fVec y0, y1;
-    std::tie(y0, y1) = at::vec::convert_to_float(y_bvec);
-
-    x0 = x0 + y0 * s_vec;
-    x1 = x1 + y1 * s_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] + float(input2[d]) * scale);
-  }
-}
-
-template <typename scalar_t>
-inline void silu_and_mul_stub(
-    scalar_t* __restrict__ out, const scalar_t* __restrict__ input, const scalar_t* __restrict__ input2, int64_t size) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  const fVec one = fVec(1.f);
-
-  // no remainder
-#pragma GCC unroll 4
-  for (int64_t d = 0; d < size; d += bVec::size()) {
-    bVec x = bVec::loadu(input + d);
-    fVec x0, x1;
-    std::tie(x0, x1) = at::vec::convert_to_float(x);
-    bVec y = bVec::loadu(input2 + d);
-    fVec y0, y1;
-    std::tie(y0, y1) = at::vec::convert_to_float(y);
-    x0 = x0 / (one + x0.neg().exp_u20());
-    x1 = x1 / (one + x1.neg().exp_u20());
-    x0 = x0 * y0;
-    x1 = x1 * y1;
-    bVec out_vec = convert_from_float_ext<scalar_t>(x0, x1);
-    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

@@ -1,14 +1,150 @@
 // Adapted from
 // https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
 
-// clang-format off
-
 #include "common.h"
 #include "gemm.h"
-#include "moe.h"
+#include "vec.h"
 
-template <typename scalar_t, typename packed_t, typename param_t, bool is_mxfp4>
-void fused_experts_fp_kernel_impl(
+// clang-format off
+
+namespace {
+
+template <typename scalar_t>
+inline void copy_stub(scalar_t* __restrict__ out, const scalar_t* __restrict__ input, int64_t size) {
+  using Vec = at::vec::Vectorized<scalar_t>;
+  // no remainder
+  #pragma GCC unroll 4
+  for (int64_t d = 0; d < size; d += Vec::size()) {
+    Vec data = Vec::loadu(input + d);
+    data.store(out + 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);
+  }
+}
+
+// acc from [topk, K] to [K]
+template <typename scalar_t>
+inline void sum_stub(scalar_t* __restrict__ out, const scalar_t* __restrict__ input, int64_t topk, int64_t K) {
+  using bVec = at::vec::Vectorized<scalar_t>;
+  using fVec = at::vec::Vectorized<float>;
+  constexpr int kVecSize = bVec::size();
+  if (topk == 1) {
+    // do copy for topk = 1
+    copy_stub(out, input, K);
+  } else {
+    // do sum for topk != 1
+    int64_t d;
+    #pragma GCC unroll 4
+    for (d = 0; d <= K - kVecSize; d += kVecSize) {
+      fVec sum_fvec0 = fVec(0.f);
+      fVec sum_fvec1 = fVec(0.f);
+      for (int t = 0; t < topk; ++t) {
+        bVec x_bvec = bVec::loadu(input + t * K + d);
+        fVec x_fvec0, x_fvec1;
+        std::tie(x_fvec0, x_fvec1) = at::vec::convert_to_float(x_bvec);
+
+        sum_fvec0 += x_fvec0;
+        sum_fvec1 += x_fvec1;
+      }
+      bVec out_bvec = convert_from_float_ext<scalar_t>(sum_fvec0, sum_fvec1);
+      out_bvec.store(out + d);
+    }
+    for (; d < K; ++d) {
+      float sum_val = 0.f;
+      for (int t = 0; t < topk; ++t) {
+        sum_val += static_cast<float>(input[t * K + d]);
+      }
+      out[d] = static_cast<scalar_t>(sum_val);
+    }
+  }
+}
+
+// out = input + input2 * scale
+template <typename scalar_t>
+inline void add_mul_stub(
+    scalar_t* __restrict__ out,
+    const scalar_t* __restrict__ input,
+    const scalar_t* __restrict__ input2,
+    float scale,
+    int64_t size) {
+  using bVec = at::vec::Vectorized<scalar_t>;
+  using fVec = at::vec::Vectorized<float>;
+  constexpr int kVecSize = bVec::size();
+  const fVec s_vec = fVec(scale);
+
+  int64_t d;
+#pragma GCC unroll 4
+  for (d = 0; d <= size - kVecSize; d += kVecSize) {
+    bVec x_bvec = bVec::loadu(input + d);
+    fVec x0, x1;
+    std::tie(x0, x1) = at::vec::convert_to_float(x_bvec);
+
+    bVec y_bvec = bVec::loadu(input2 + d);
+    fVec y0, y1;
+    std::tie(y0, y1) = at::vec::convert_to_float(y_bvec);
+
+    x0 = x0 + y0 * s_vec;
+    x1 = x1 + y1 * s_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] + float(input2[d]) * scale);
+  }
+}
+
+template <typename scalar_t>
+inline void silu_and_mul_stub(
+    scalar_t* __restrict__ out,
+    const scalar_t* __restrict__ input,
+    const scalar_t* __restrict__ input2,
+    int64_t size) {
+  using bVec = at::vec::Vectorized<scalar_t>;
+  using fVec = at::vec::Vectorized<float>;
+  const fVec one = fVec(1.f);
+
+  // no remainder
+#pragma GCC unroll 4
+  for (int64_t d = 0; d < size; d += bVec::size()) {
+    bVec x = bVec::loadu(input + d);
+    fVec x0, x1;
+    std::tie(x0, x1) = at::vec::convert_to_float(x);
+    bVec y = bVec::loadu(input2 + d);
+    fVec y0, y1;
+    std::tie(y0, y1) = at::vec::convert_to_float(y);
+    x0 = x0 / (one + x0.neg().exp_u20());
+    x1 = x1 / (one + x1.neg().exp_u20());
+    x0 = x0 * y0;
+    x1 = x1 * y1;
+    bVec out_vec = convert_from_float_ext<scalar_t>(x0, x1);
+    out_vec.store(out + d);
+  }
+}
+
+} // anonymous namespace
+
+template <typename scalar_t>
+void fused_experts_fp8_kernel_impl(
     scalar_t* __restrict__ output,
     scalar_t* __restrict__ ic0,
     scalar_t* __restrict__ ic1,
@@ -17,12 +153,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 +168,8 @@ 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,52 +179,39 @@ 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;
-
-  int64_t avg_M = std::max(int64_t(1), M * topk / E);
-  const bool use_brgemm = can_use_brgemm<packed_t>(avg_M);
-
-  int64_t B_tmp_size_per_thread = MAX_CACHE_BLOCK_SIZE * BLOCK_N * std::max(K, N);
+  const int64_t stride_e = 2 * N * K;
+  const int64_t stride_n = K;
 
   // here we only parallel on half of 2N to fuse silu_and_mul with gemm
-  parallel_2d(MB, NB, [&](int64_t mb0, int64_t mb1, int64_t nb0, int64_t nb1) {
+  at::parallel_for(0, MB * NB, 0, [&](int64_t begin, int64_t end) {
     // get local pointers
-    int tid = get_thread_num();
+    int tid = at::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) {
+    bool is_brgemm_used = false;
+
+    for (int64_t i = begin; i < end; ++i) {
+      int64_t mb = i / NB;
+      int64_t nb = i % NB;
+
       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;
-
-      // do unpacking for the first row or a new expert
-      int32_t pre_expert_id = mb == 0 ? -1 : expert_ids[mb - 1];
-      bool do_unpack = (mb == mb0) || (expert_id != pre_expert_id);
+      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;
 
+      // 1.a load A
+      const int32_t* A_ids = sorted_ids + mb * BLOCK_M;
       int64_t m_size = offsets[mb + 1] - offsets[mb];
 
-      if (nb_offset == 0) {
-        // 1.a load A
-        const int32_t* A_ids = sorted_ids + mb * BLOCK_M;
-        for (int64_t m = 0; m < m_size; ++m) {
-          int32_t index = A_ids[m] / topk;
-          copy_stub(A + m * K, input + index * K, K);
-        }
+      const bool use_brgemm = can_use_brgemm<at::Float8_e4m3fn>(m_size);
+      is_brgemm_used = is_brgemm_used || use_brgemm;
+
+      for (int64_t m = 0; m < m_size; ++m) {
+        int32_t index = A_ids[m] / topk;
+        copy_stub(A + m * K, input + index * K, K);
       }
 
       const int64_t offset = offsets[mb];
@@ -101,9 +219,8 @@ void fused_experts_fp_kernel_impl(
           /*   A            */ A,
           /*   B            */ B,
           /*   C            */ ic0 + offset * 2 * N + nb * BLOCK_N,
-          /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * K,
+          /*   Btmp         */ B_tmp + tid * BLOCK_N * std::max(K, N),
           /*   Ctmp         */ C_tmp + tid * 2 * BLOCK_M * BLOCK_N,
-          /*   Bbias        */ B_bias,
           /*   scale        */ Bs,
           /*   M            */ m_size,
           /*   N            */ n_size,
@@ -112,30 +229,25 @@ void fused_experts_fp_kernel_impl(
           /*   ldb          */ n_size,
           /*   ldc          */ 2 * N,
           /*   brg          */ use_brgemm,
-          /*   block_size_K */ block_size_K,
-          /*   do_unpack    */ do_unpack);
-    });
+          /*   block_size_K */ block_size_K);
+    }
 
-    if (use_brgemm) {
+    if (is_brgemm_used) {
       at::native::cpublas::brgemm_release();
     }
   });
 
   // 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,19 +256,26 @@ 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();
+  at::parallel_for(0, MB2 * NB2, 0, [&](int64_t begin, int64_t end) {
+    int tid = at::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) {
+    bool is_brgemm_used = false;
+
+    for (int64_t i = begin; i < end; ++i) {
+      int64_t mb = i / NB2;
+      int64_t nb = i % NB2;
+
       int64_t m_size = offsets[mb + 1] - offsets[mb];
       int64_t n_size = std::min(OC - nb * BLOCK_N, BLOCK_N);
 
+      const bool use_brgemm = can_use_brgemm<at::Float8_e4m3fn>(m_size);
+      is_brgemm_used = is_brgemm_used || use_brgemm;
+
       // A ptr from ic1 of [M * topk, N] in sorted order
       // so as to avoid copy A to tmp buffer again
       const scalar_t* __restrict__ A = ic1 + offsets[mb] * N;
@@ -164,22 +283,15 @@ 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;
-
-      // do unpacking for the first row or a new expert
-      int32_t pre_expert_id = mb == 0 ? -1 : expert_ids[mb - 1];
-      bool do_unpack = (mb == mb0) || (expert_id != pre_expert_id);
+      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;
 
       tinygemm_kernel<scalar_t>(
           /*   A            */ A,
           /*   B            */ B,
           /*   C            */ C,
-          /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * IC,
+          /*   Btmp         */ B_tmp + tid * BLOCK_N * std::max(K, N),
           /*   Ctmp         */ C_tmp + tid * 2 * BLOCK_M * BLOCK_N,
-          /*   Bbias        */ B_bias,
           /*   scale        */ Bs,
           /*   M            */ m_size,
           /*   N            */ n_size,
@@ -188,8 +300,7 @@ void fused_experts_fp_kernel_impl(
           /*   ldb          */ n_size,
           /*   ldc          */ BLOCK_N,
           /*   brg          */ use_brgemm,
-          /*   block_size_K */ block_size_K,
-          /*   do_unpack    */ do_unpack);
+          /*   block_size_K */ block_size_K);
 
       // 2.b copy from C to ic2 in original order
       //   and also mul topk_weights in float32
@@ -198,12 +309,13 @@ void fused_experts_fp_kernel_impl(
         float weight = topk_weights[index];
         copy_mul_stub(ic2 + index * K + nb * BLOCK_N, C + m * BLOCK_N, weight, n_size);
       }
-    });
+    }
 
-    if (use_brgemm) {
+    if (is_brgemm_used) {
       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 +325,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(
@@ -270,6 +374,7 @@ void shared_expert_fp8_kernel_impl(
     int64_t M,
     int64_t N,
     int64_t K) {
+
   constexpr int64_t BLOCK_M = block_size_m();
   constexpr int64_t BLOCK_N = block_size_n();
 
@@ -280,27 +385,22 @@ void shared_expert_fp8_kernel_impl(
   int64_t blocks_n_per_group = block_size_N / BLOCK_N;
 
   const bool use_brgemm = can_use_brgemm<at::Float8_e4m3fn>(M);
-  const bool apply_scaling_factor = fused_experts_out != nullptr;
 
-  int64_t B_tmp_size_per_thread = MAX_CACHE_BLOCK_SIZE * BLOCK_N * std::max(K, N);
+  at::parallel_for(0, MB * NB, 0, [&](int64_t begin, int64_t end) {
+    int tid = at::get_thread_num();
 
-  parallel_2d(MB, NB, [&](int64_t mb0, int64_t mb1, int64_t nb0, int64_t nb1) {
-    int tid = get_thread_num();
-
-    loop_2d<at::Float8_e4m3fn>(mb0, mb1, nb0, nb1, BLOCK_N * K, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
+    for (int64_t i = begin; i < end; ++i) {
+      int64_t mb = i / NB;
+      int64_t nb = i % NB;
       int64_t m_size = std::min(M - mb * BLOCK_M, BLOCK_M);
       int64_t n_size = std::min(2 * N - nb * BLOCK_N, BLOCK_N);
 
-      // do unpacking for the first row
-      bool do_unpack = (mb == mb0);
-
       tinygemm_kernel<scalar_t>(
           /*   A            */ input + mb * BLOCK_M * K,
           /*   B            */ packed_w1 + nb * BLOCK_N * K,
           /*   C            */ ic0 + mb * BLOCK_M * 2 * N + nb * BLOCK_N,
-          /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * K,
+          /*   Btmp         */ B_tmp + tid * BLOCK_N * std::max(K, N),
           /*   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,
@@ -309,9 +409,8 @@ void shared_expert_fp8_kernel_impl(
           /*   ldb          */ n_size,
           /*   ldc          */ 2 * N,
           /*   brg          */ use_brgemm,
-          /*   block_size_K */ block_size_K,
-          /*   do_unpack    */ do_unpack);
-    });
+          /*   block_size_K */ block_size_K);
+    }
 
     if (use_brgemm) {
       at::native::cpublas::brgemm_release();
@@ -321,7 +420,11 @@ void shared_expert_fp8_kernel_impl(
   // stage 1.5: intermediate_cache1 = silu(intermediate_cache0)
   at::parallel_for(0, M, 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);
+      silu_and_mul_stub(
+          ic1 + m * N,
+          ic0 + m * 2 * N,
+          ic0 + m * 2 * N + N,
+          N);
     }
   });
 
@@ -334,25 +437,23 @@ void shared_expert_fp8_kernel_impl(
   scale_size_K = div_up(N, block_size_K);
 
   // 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();
+  at::parallel_for(0, MB2 * NB2, 0, [&](int64_t begin, int64_t end) {
+    int tid = at::get_thread_num();
     alignas(64) scalar_t C[BLOCK_M * BLOCK_K];
 
-    loop_2d<at::Float8_e4m3fn>(mb0, mb1, nb0, nb1, BLOCK_N * IC, [&](int64_t mb, int64_t nb, int64_t nb_offset) {
+    for (int64_t i = begin; i < end; ++i) {
+      int64_t mb = i / NB2;
+      int64_t nb = i % NB2;
       int64_t m_size = std::min(M - mb * BLOCK_M, BLOCK_M);
       int64_t n_size = std::min(OC - nb * BLOCK_N, BLOCK_N);
 
-      // do unpacking for the first row
-      bool do_unpack = (mb == mb0);
-
       // 2.a gemm: C = A @ B
       tinygemm_kernel<scalar_t>(
           /*   A            */ ic1 + mb * BLOCK_M * N,
           /*   B            */ packed_w2 + nb * BLOCK_N * N,
           /*   C            */ C,
-          /*   Btmp         */ B_tmp + tid * B_tmp_size_per_thread + nb_offset * BLOCK_N * IC,
+          /*   Btmp         */ B_tmp + tid * BLOCK_N * std::max(K, N),
           /*   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,
@@ -361,18 +462,15 @@ void shared_expert_fp8_kernel_impl(
           /*   ldb          */ n_size,
           /*   ldc          */ BLOCK_N,
           /*   brg          */ use_brgemm,
-          /*   block_size_K */ block_size_K,
-          /*   do_unpack    */ do_unpack);
+          /*   block_size_K */ block_size_K);
 
       // 2.b copy from C to output and add fused_experts_out
       scalar_t* __restrict__ out = output + mb * BLOCK_M * K + nb * BLOCK_N;
-      const scalar_t* __restrict__ fused_out =
-          apply_scaling_factor ? fused_experts_out + mb * BLOCK_M * K + nb * BLOCK_N : nullptr;
+      const scalar_t* __restrict__ fused_out = fused_experts_out + mb * BLOCK_M * K + nb * BLOCK_N;
       for (int64_t m = 0; m < m_size; ++m) {
-        const scalar_t* __restrict__ fused_out_row = apply_scaling_factor ? (fused_out + m * K) : nullptr;
-        add_mul_stub(out + m * K, C + m * BLOCK_N, fused_out_row, routed_scaling_factor, n_size);
+        add_mul_stub(out + m * K, C + m * BLOCK_N, fused_out + m * K, routed_scaling_factor, n_size);
       }
-    });
+    }
   });
 
   if (use_brgemm) {

csrc/cpu/sgl-kernels/moe_int4.cpp

@@ -1,323 +0,0 @@
-// Adapted from
-// https://github.com/sgl-project/sglang/tree/main/sgl-kernel/csrc/cpu
-
-// clang-format off
-
-#include "common.h"
-#include "gemm.h"
-#include "moe.h"
-
-template <int64_t N>
-inline void copy_bias(const float* bias_ptr, float* y_buf, int64_t m, int64_t ldn) {
-  using Vec = at::vec::Vectorized<float>;
-  constexpr int kVecSize = Vec::size();
-  static_assert(N % kVecSize == 0, "copy_bias requires N to be a multiple of Vectorized<float>::size()");
-  const bool has_bias = bias_ptr != nullptr;
-  const Vec zero_vec(0.f);
-  for (int i = 0; i < m; ++i) {
-#pragma GCC unroll 2
-    for (int j = 0; j < N; j += kVecSize) {
-      Vec vec = has_bias ? Vec::loadu(bias_ptr + j) : zero_vec;
-      vec.store(y_buf + i * ldn + j);
-    }
-  }
-}
-
-template <typename scalar_t>
-void fused_experts_int4_w4a8_kernel_impl(
-    scalar_t* __restrict__ output,
-    scalar_t* __restrict__ ic0,
-    scalar_t* __restrict__ ic1,
-    scalar_t* __restrict__ ic2,
-    uint8_t* __restrict__ A_tmp,
-    uint8_t* __restrict__ Aq_tmp,
-    float* __restrict__ As_tmp,
-    int32_t* __restrict__ Azp_tmp,
-    float* __restrict__ C_tmp,
-    int8_t* __restrict__ dqB_tmp,
-    const scalar_t* __restrict__ input,
-    const uint8_t* __restrict__ packed_w1,
-    const uint8_t* __restrict__ packed_w2,
-    const int8_t* __restrict__ w1z,
-    const int8_t* __restrict__ w2z,
-    const float* __restrict__ w1s,
-    const float* __restrict__ w2s,
-    int group_size,
-    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) {
-  constexpr int64_t BLOCK_M = block_size_m();
-  constexpr int64_t BLOCK_N = block_size_n();
-  int num_threads = at::get_num_threads();
-  // int64_t buffer_size_nbytes = M * topk * N * 2
-  //                              M * topk * K * 2 +
-  //                              num_threads * BLOCK_M * K +
-  //                              num_threads * 2 * BLOCK_M * BLOCK_N * sizeof(float)  +
-  //                              M * topk * 2 * N * 2 +
-  //                              max(M * K, M * topk * N)  +
-  //                              M * topk * sizeof(float);
-
-  // intermediate_cache1 (scalar_t):     START + M * topk * N
-  // intermediate_cache2 (scalar_t):     + M * topk * K
-  // A_tmp (uint8_t):                    + num_threads * BLOCK_M * K
-  // C_tmp (float):                      + num_threads * 2 * BLOCK_M * BLOCK_N
-  // intermediate_cache0 (scalar_t):     + M * topk * 2 * N
-  // Aq_tmp (uint8_t):                   + max(M * K, M * topk * N)
-  // As_tmp (float):                     + M * topk
-  // dqB_tmp (int8_t)                    + num_threads * _block_k * BlOCK_N
-
-  // stage 0: quantize input to uint8, [M, K]
-  at::parallel_for(0, M, 0, [&](int64_t begin, int64_t end) {
-    for (int64_t m = begin; m < end; ++m) {
-      quantize_row_int8<scalar_t>(Aq_tmp + m * K, As_tmp[m], input + m * K, K);
-    }
-  });
-  int64_t _block_k = get_4bit_block_k_size(group_size);
-  auto Azp = at::ones({M * topk}).to(at::kInt).mul(128);
-  auto Azp_ptr = Azp.data_ptr<int32_t>();
-  // stage 1: intermediate_cache0 = hidden_states @ w1
-  const int64_t MB = div_up(num_tokens_post_pad, BLOCK_M);
-  const int64_t NB = div_up(N, BLOCK_N);
-
-  int64_t block_per_group = group_size / _block_k;
-  int64_t Kc = K / _block_k;
-  int64_t num_groups = K / group_size;
-
-  const int64_t stride_e = 2 * NB * Kc * (BLOCK_N * (_block_k / 2 + sizeof(int32_t)));
-  const bool sym_quant_act = false;
-  // weight + compensation shape = [E, Nc, Kc, block_n * _block_k / 2 + block_n*sizeof(int32_t)]
-  // scales/qzeros shape = [E, Nc, G, block_n]
-
-  // here we only parallel on half of 2N to fuse silu_and_mul with gemm
-  at::parallel_for(0, MB * NB, 0, [&](int64_t begin, int64_t end) {
-    // get local pointers
-    int tid = at::get_thread_num();
-    int8_t* dqB_tmp1 = dqB_tmp + tid * 2 * _block_k * BLOCK_N;
-    int8_t* dqB_tmp2 = dqB_tmp1 + _block_k * BLOCK_N;
-    alignas(64) float As[BLOCK_M];
-    uint8_t* __restrict__ A = A_tmp + tid * BLOCK_M * K;
-    float* __restrict__ C0 = C_tmp + tid * 2 * BLOCK_M * BLOCK_N;
-    float* __restrict__ C1 = C0 + BLOCK_M * BLOCK_N;
-    bool is_brgemm_used = false;
-    for (int64_t i = begin; i < end; ++i) {
-      int64_t mb = i / NB;
-      int64_t nb = i % NB;
-      int64_t nb1 = nb + NB;
-      int64_t n_size = std::min(N - nb * BLOCK_N, BLOCK_N);
-      // B shape [K, n_size] in vnni format
-      int32_t expert_id = expert_ids[mb];
-      const uint8_t* __restrict__ B = packed_w1 + expert_id * stride_e;
-      // Bz and Bs: [E, K/gs, 2N]
-      const int8_t* __restrict__ Bz = w1z + expert_id * (num_groups) * (2 * N);
-      const float* __restrict__ Bs = w1s + expert_id * (num_groups) * (2 * N);
-
-      // 1.a load A
-      const int32_t* A_ids = sorted_ids + mb * BLOCK_M;
-      int64_t m_size = offsets[mb + 1] - offsets[mb];
-      const bool use_brgemm = can_use_brgemm<int8_t>(m_size);
-      is_brgemm_used = is_brgemm_used || use_brgemm;
-      // copy to A [BLOCK_M, K]
-      for (int64_t m = 0; m < m_size; ++m) {
-        int32_t index = A_ids[m] / topk;
-        copy_stub(A + m * K, Aq_tmp + index * K, K);
-        As[m] = As_tmp[index];
-      }
-      const int64_t offset = offsets[mb];
-      copy_bias<BLOCK_N>(nullptr, C0, m_size, BLOCK_N);
-      copy_bias<BLOCK_N>(nullptr, C1, m_size, BLOCK_N);
-      for (int kci = 0; kci < Kc; ++kci) {
-        int32_t* compensation_ptr =
-            sym_quant_act ? nullptr
-                          : (int32_t*)(void*)(B + (nb * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))) +
-                                              _block_k * BLOCK_N / 2) /*Bcomp*/;
-        tinygemm_kernel<scalar_t>(
-            ic0 + offset * 2 * N + nb * BLOCK_N,
-            C0,
-            A + kci * _block_k,
-            As,
-            Azp_ptr,
-            B + (nb * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))) /*B*/,
-            Bs + nb * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N /*scales_b*/,
-            Bz + nb * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N /*qzeros_b*/,
-            compensation_ptr,
-            dqB_tmp1,
-            m_size,
-            _block_k,
-            K,
-            BLOCK_N,
-            2 * N,
-            kci == Kc - 1,
-            use_brgemm);
-      }
-
-      for (int kci = 0; kci < Kc; ++kci) {
-        int32_t* compensation_ptr =
-            sym_quant_act ? nullptr
-                          : (int32_t*)(void*)(B + (nb1 * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))) +
-                                              _block_k * BLOCK_N / 2) /*Bcomp*/;
-        tinygemm_kernel<scalar_t>(
-            ic0 + offset * 2 * N + nb1 * BLOCK_N,
-            C1,
-            A + kci * _block_k,
-            As,
-            Azp_ptr,
-            B + (nb1 * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))) /*B*/,
-            Bs + nb1 * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N /*scales_b*/,
-            Bz + nb1 * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N /*qzeros_b*/,
-            compensation_ptr,
-            dqB_tmp2,
-            m_size,
-            _block_k,
-            K,
-            BLOCK_N,
-            2 * N,
-            kci == Kc - 1,
-            use_brgemm);
-      }
-    }
-
-    if (is_brgemm_used) {
-      at::native::cpublas::brgemm_release();
-    }
-  });
-
-  // stage 1.5: intermediate_cache1 = silu(intermediate_cache0)
-  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 1.5: quantize ic1 to uint8, [M * topk, N]
-  at::parallel_for(0, M * topk, 0, [&](int64_t begin, int64_t end) {
-    for (int64_t m = begin; m < end; ++m) {
-      quantize_row_int8<scalar_t>(Aq_tmp + m * N, As_tmp[m], ic1 + m * 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
-  const int64_t IC = N;  // rename N as IC
-  const int64_t MB2 = MB;
-  const int64_t NB2 = div_up(OC, BLOCK_N);
-  const int64_t stride_oc = IC;
-  num_groups = IC / group_size;
-  Kc = IC / _block_k;
-  const int64_t stride_e2 = NB2 * Kc * (BLOCK_N * (_block_k / 2 + sizeof(int32_t)));
-  // parallel on [MB2, NB2]
-  at::parallel_for(0, MB2 * NB2, 0, [&](int64_t begin, int64_t end) {
-    int tid = at::get_thread_num();
-    int8_t* dqB_tmp1 = dqB_tmp + tid * 2 * _block_k * BLOCK_N;
-    float* __restrict__ C2 = C_tmp + tid * 2 * BLOCK_M * BLOCK_N;
-    bool is_brgemm_used = false;
-    for (int64_t i = begin; i < end; ++i) {
-      int64_t mb = i / NB2;
-      int64_t nb = i % NB2;
-
-      int64_t m_size = offsets[mb + 1] - offsets[mb];
-      int64_t n_size = std::min(OC - nb * BLOCK_N, BLOCK_N);
-      const bool use_brgemm = can_use_brgemm<int8_t>(m_size);
-      is_brgemm_used = is_brgemm_used || use_brgemm;
-      const int32_t* A_ids = sorted_ids + mb * BLOCK_M;
-
-      // B shape [IC, n_size] in vnni format
-      int32_t expert_id = expert_ids[mb];
-      const uint8_t* __restrict__ B = packed_w2 + expert_id * stride_e2;
-
-      // Bz and Bs: [E, IC/gs, OC]
-      const int8_t* __restrict__ Bz = w2z + expert_id * (num_groups)*OC;
-      const float* __restrict__ Bs = w2s + expert_id * (num_groups)*OC;
-
-      // A ptr from ic1 of [M * topk, N] in sorted order
-      // so as to avoid copy A to tmp buffer again
-      const uint8_t* __restrict__ A = Aq_tmp + offsets[mb] * IC;
-      const float* __restrict__ As = As_tmp + offsets[mb];
-      copy_bias<BLOCK_N>(nullptr, C2, m_size, BLOCK_N);
-      for (int kci = 0; kci < Kc; ++kci) {
-        int32_t* compensation_ptr =
-            sym_quant_act ? nullptr
-                          : (int32_t*)(void*)(B + (nb * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))) +
-                                              _block_k * BLOCK_N / 2) /*Bcomp*/;
-        tinygemm_kernel<scalar_t>(
-            nullptr, /*store_out is false*/
-            C2,
-            A + kci * _block_k,
-            As,
-            Azp_ptr,
-            B + (nb * Kc + kci) * (BLOCK_N * (_block_k / 2 + sizeof(int32_t))),
-            Bs + nb * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N /*scales_b*/,
-            Bz + nb * BLOCK_N * num_groups + kci / block_per_group * BLOCK_N /*zeros_b*/,
-            compensation_ptr,
-            dqB_tmp1,
-            m_size,
-            _block_k,
-            IC,
-            BLOCK_N,
-            BLOCK_N,
-            false,
-            use_brgemm);
-      }
-
-      // 2.b copy from C to ic2 in original order
-      //   and also mul topk_weights in float32
-      for (int64_t m = 0; m < m_size; ++m) {
-        int32_t index = A_ids[m];
-        float weight = topk_weights[index];
-        copy_mul_stub(ic2 + index * K + nb * BLOCK_N, C2 + m * BLOCK_N, weight, n_size);
-      }
-    }
-
-    if (is_brgemm_used) {
-      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) {
-    for (int64_t m = begin; m < end; ++m) {
-      sum_stub(output + m * K, ic2 + m * topk * K, topk, K);
-    }
-  });
-}
-
-#define INSTANTIATE_MOE_INT4_W4A8_TEMPLATE(TYPE)           \
-  template void fused_experts_int4_w4a8_kernel_impl<TYPE>( \
-      TYPE* __restrict__ output,                           \
-      TYPE* __restrict__ ic0,                              \
-      TYPE* __restrict__ ic1,                              \
-      TYPE* __restrict__ ic2,                              \
-      uint8_t* __restrict__ A_tmp,                         \
-      uint8_t* __restrict__ Aq_tmp,                        \
-      float* __restrict__ As_tmp,                          \
-      int32_t* __restrict__ Azp_tmp,                       \
-      float* __restrict__ C_tmp,                           \
-      int8_t* __restrict__ dqB_tmp,                        \
-      const TYPE* __restrict__ input,                      \
-      const uint8_t* __restrict__ packed_w1,               \
-      const uint8_t* __restrict__ packed_w2,               \
-      const int8_t* __restrict__ w1z,                      \
-      const int8_t* __restrict__ w2z,                      \
-      const float* __restrict__ w1s,                       \
-      const float* __restrict__ w2s,                       \
-      int group_size,                                      \
-      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_INT4_W4A8_TEMPLATE(at::BFloat16);
-INSTANTIATE_MOE_INT4_W4A8_TEMPLATE(at::Half);