Merge branch 'main' into wentao-skip-work-when-empty

refactor only inside nixl
Signed-off-by: yewentao256 <zhyanwentao@126.com>
2026-04-03 11:31:23 -04:00 · 2026-04-01 18:28:13 +00:00 · 2026-04-01 18:09:06 +00:00 · 2026-03-31 13:01:59 -04:00 · 2026-03-30 21:06:12 +00:00 · 2026-03-30 19:17:12 +00:00
303 changed files with 7703 additions and 21476 deletions
@@ -5,6 +5,7 @@ steps:
    depends_on: []
    device: amd_cpu
    no_plugin: true
+    soft_fail: true
    commands:
    - >
      docker build
@@ -1,9 +1,6 @@
 # For hf script, without -t option (tensor parallel size).
 # bash .buildkite/lm-eval-harness/run-lm-eval-mmlupro-vllm-baseline.sh -m meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8 -l 250 -t 8 -f 5
 model_name: "meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8"
-required_gpu_arch:
-  - gfx942
-  - gfx950
 tasks:
 - name: "mmlu_pro"
  metrics:
@@ -1,9 +1,6 @@
 # For vllm script, with -t option (tensor parallel size)
 # bash .buildkite/lm-eval-harness/run-lm-eval-gsm-vllm-baseline.sh -m RedHatAI/Qwen2.5-VL-3B-Instruct-FP8-Dynamic -l 1319 -t 1
 model_name: "RedHatAI/Qwen2.5-VL-3B-Instruct-FP8-Dynamic"
-required_gpu_arch:
-  - gfx942
-  - gfx950
 tasks:
 - name: "gsm8k"
  metrics:
@@ -1,7 +1,4 @@
 model_name: "Qwen/Qwen3-235B-A22B-Instruct-2507-FP8"
-required_gpu_arch:
-  - gfx942
-  - gfx950
 tasks:
  - name: "mmlu_pro"
    metrics:
@@ -1,6 +1,5 @@
 Qwen2.5-1.5B-Instruct.yaml
 Meta-Llama-3.2-1B-Instruct-INT8-compressed-tensors.yaml
-Meta-Llama-3-8B-Instruct-INT8-compressed-tensors-asym.yaml
 Meta-Llama-3-8B-Instruct-nonuniform-compressed-tensors.yaml
 Qwen2.5-VL-3B-Instruct-FP8-dynamic.yaml
 Qwen1.5-MoE-W4A16-compressed-tensors.yaml
@@ -13,7 +13,6 @@ import os
 from contextlib import contextmanager

 import lm_eval
-import pytest
 import yaml

 from vllm.platforms import current_platform
@@ -90,40 +89,9 @@ def launch_lm_eval(eval_config, tp_size):
    return results


-def _check_rocm_gpu_arch_requirement(eval_config):
-    """Skip the test if the model requires a ROCm GPU arch not present.
-
-    Model YAML configs can specify::
-
-        required_gpu_arch:
-          - gfx942
-          - gfx950
-
-    The check only applies on ROCm.  On other platforms (e.g. CUDA) the
-    field is ignored so that shared config files work for both NVIDIA and
-    AMD CI pipelines.
-    """
-    required_archs = eval_config.get("required_gpu_arch")
-    if not required_archs:
-        return
-
-    if not current_platform.is_rocm():
-        return
-
-    from vllm.platforms.rocm import _GCN_ARCH  # noqa: E402
-
-    if not any(arch in _GCN_ARCH for arch in required_archs):
-        pytest.skip(
-            f"Model requires GPU arch {required_archs}, "
-            f"but detected arch is '{_GCN_ARCH}'"
-        )
-
-
 def test_lm_eval_correctness_param(config_filename, tp_size):
    eval_config = yaml.safe_load(config_filename.read_text(encoding="utf-8"))

-    _check_rocm_gpu_arch_requirement(eval_config)
-
    results = launch_lm_eval(eval_config, tp_size)

    rtol = eval_config.get("rtol", DEFAULT_RTOL)
@@ -35,6 +35,23 @@ export PYTHONPATH=".."
 # Helper Functions
 ###############################################################################

+wait_for_clean_gpus() {
+  local timeout=${1:-300}
+  local start=$SECONDS
+  echo "--- Waiting for clean GPU state (timeout: ${timeout}s)"
+  while true; do
+    if grep -q clean /opt/amdgpu/etc/gpu_state; then
+      echo "GPUs state is \"clean\""
+      return
+    fi
+    if (( SECONDS - start >= timeout )); then
+      echo "Error: GPUs did not reach clean state within ${timeout}s" >&2
+      exit 1
+    fi
+    sleep 3
+  done
+}
+
 cleanup_docker() {
  # Get Docker's root directory
  docker_root=$(docker info -f '{{.DockerRootDir}}')
@@ -348,12 +365,19 @@ apply_rocm_test_overrides() {
 ###############################################################################

 # --- GPU initialization ---
+echo "--- Confirming Clean Initial State"
+wait_for_clean_gpus
+
 echo "--- ROCm info"
 rocminfo

 # --- Docker housekeeping ---
 cleanup_docker

+echo "--- Resetting GPUs"
+echo "reset" > /opt/amdgpu/etc/gpu_state
+wait_for_clean_gpus
+
 # --- Pull test image ---
 echo "--- Pulling container"
 image_name="rocm/vllm-ci:${BUILDKITE_COMMIT}"
@@ -751,7 +751,6 @@ steps:
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
  agent_pool: mi250_1
-  optional: true
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - csrc/
@@ -2036,6 +2035,7 @@ steps:
  timeout_in_minutes: 38
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
  agent_pool: mi325_1
+  optional: true
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - csrc/
@@ -2165,15 +2165,7 @@ steps:
  - vllm/platforms/rocm.py
  - tests/quantization
  commands:
-
-  # temporary install here since we need nightly, will move to requirements/test.in
-  # after torchao 0.12 release, and pin a working version of torchao nightly here
-
-  # since torchao nightly is only compatible with torch nightly currently
-  # https://github.com/pytorch/ao/issues/2919, we'll have to skip new torchao tests for now
-  # we can only upgrade after this is resolved
-  # TODO(jerryzh168): resolve the above comment
-  - uv pip install --system torchao==0.17.0
+  - uv pip install --system torchao==0.14.1
  - uv pip install --system conch-triton-kernels
  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py

@@ -2698,24 +2690,6 @@ steps:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt


- label: LM Eval Small Models (MI325) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
-  agent_pool: mi325_1
-  working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
-  source_file_dependencies:
-  - csrc/
-  - vllm/model_executor/layers/quantization
-  - vllm/model_executor/models/
-  - vllm/model_executor/model_loader/
-  - vllm/v1/attention/backends/
-  - vllm/v1/attention/selector.py
-  - vllm/_aiter_ops.py
-  - vllm/platforms/rocm.py
-  commands:
-  - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-small-rocm.txt
-
-
 - label: LM Eval Small Models (B200-MI325) # TBD
  timeout_in_minutes: 180
  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
@@ -2932,10 +2906,10 @@ steps:
  - bash .buildkite/scripts/scheduled_integration_test/qwen3_next_mtp_async_eplb.sh 0.8 1319 8040

 ##### .buildkite/test_areas/compile.yaml #####
-# Slowly setting up the tests so that it is also easier for the
+# Slowly setting up the tests so that it is also easier for the 
 # CI team to review and upstream to the pipelinev2.
 # The following tests are important for vLLM IR Ops refactoring,
-# which affects fusion passes on ROCm. So we have to
+# which affects fusion passes on ROCm. So we have to 
 # enable them as as soon as possible.

 ## TODO: Enable the test in this group
@@ -3014,7 +2988,7 @@ steps:

 ## There are no ops on ROCm for these tests.
 ## The test still passes but the logs are not useful.
-## fused ops just call torch.ops.symm_mem which
+## fused ops just call torch.ops.symm_mem which 
 ## exists in ROCm even though they don't work
 # - label: AsyncTP Correctness Tests (2xH100-2xMI325)
 # - label: Fusion E2E TP2 Quick (H100-MI325)
@@ -3346,7 +3320,7 @@ steps:
  - vllm/_aiter_ops.py
  - vllm/platforms/rocm.py
  commands:
-  - uv pip install --system torchao==0.17.0
+  - uv pip install --system torchao==0.14.1
  - uv pip install --system conch-triton-kernels
  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py

@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: Basic Correctness
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/basic_correctness/test_basic_correctness
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
  timeout_in_minutes: 20
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/benchmarks/
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: Platform Tests (CUDA)
  timeout_in_minutes: 15
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/cuda
@@ -224,6 +224,20 @@ steps:
  commands:
    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_multi_node_assignment.py && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"

+- label: MessageQueue TCP Multi-Node (2 GPUs)
+  timeout_in_minutes: 10
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 1
+  num_nodes: 2
+  no_plugin: true
+  optional: true
+  source_file_dependencies:
+  - vllm/distributed/device_communicators/shm_broadcast.py
+  - vllm/distributed/parallel_state.py
+  - tests/distributed/test_mq_tcp_multinode.py
+  commands:
+    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 1 $IMAGE_TAG "torchrun --nnodes 2 --nproc-per-node=1 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_mq_tcp_multinode.py" "torchrun --nnodes 2 --nproc-per-node=1 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_mq_tcp_multinode.py"
+
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: Engine
  timeout_in_minutes: 15
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/engine
@@ -26,7 +25,6 @@ steps:

 - label: e2e Scheduling (1 GPU)
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/
    - tests/v1/e2e/general/
@@ -61,7 +61,6 @@ steps:

 - label: Entrypoints Integration (API Server openai - Part 3)
  timeout_in_minutes: 50
-  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -106,7 +105,6 @@ steps:

 - label: OpenAI API Correctness
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
  - csrc/
  - vllm/entrypoints/openai/
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: EPLB Algorithm
  timeout_in_minutes: 15
-  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/distributed/eplb
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: vLLM IR Tests
  timeout_in_minutes: 10
-  device: h200_18gb
  working_dir: "/vllm-workspace/"
  source_file_dependencies:
    - vllm/ir
@@ -18,9 +17,10 @@ steps:
  source_file_dependencies:
  - csrc/
  - tests/kernels/core
+  - tests/kernels/test_top_k_per_row.py
  - tests/kernels/test_concat_mla_q.py
  commands:
-    - pytest -v -s kernels/core kernels/test_concat_mla_q.py
+    - pytest -v -s kernels/core kernels/test_top_k_per_row.py kernels/test_concat_mla_q.py

 - label: Kernels Attention Test %N
  timeout_in_minutes: 35
@@ -106,7 +106,6 @@ steps:
  - vllm/v1/attention/backends/mla/flashinfer_mla.py
  - vllm/v1/attention/selector.py
  - vllm/platforms/cuda.py
-  - tests/kernels/test_top_k_per_row.py
  commands:
    - nvidia-smi
    - python3 examples/basic/offline_inference/chat.py
@@ -117,7 +116,6 @@ steps:
    - pytest -v -s tests/kernels/attention/test_flashinfer_trtllm_attention.py
    - pytest -v -s tests/kernels/attention/test_cutlass_mla_decode.py
    - pytest -v -s tests/kernels/attention/test_flashinfer_mla_decode.py
-    - pytest -v -s tests/kernels/test_top_k_per_row.py
    # Quantization
    - pytest -v -s tests/kernels/quantization/test_cutlass_scaled_mm.py -k 'fp8'
    - pytest -v -s tests/kernels/quantization/test_nvfp4_quant.py
@@ -181,21 +179,3 @@ steps:
    - pytest -v -s kernels/moe/test_flashinfer_moe.py
    - pytest -v -s kernels/moe/test_nvfp4_moe.py
    - pytest -v -s kernels/moe/test_ocp_mx_moe.py
-
-
- label: Kernels FusedMoE Layer Test (2 H100s)
-  timeout_in_minutes: 90
-  device: h100
-  num_devices: 2
-  optional: true
-  commands:
-    - pytest -v -s kernels/moe/test_moe_layer.py
-
-
- label: Kernels FusedMoE Layer Test (2 B200s)
-  timeout_in_minutes: 90
-  device: b200
-  num_devices: 2
-  optional: true
-  commands:
-    - pytest -v -s kernels/moe/test_moe_layer.py
@@ -19,7 +19,6 @@ steps:

 - label: V1 Sample + Logits
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
    - vllm/
    - tests/v1/sample
@@ -87,7 +86,6 @@ steps:

 - label: Regression
  timeout_in_minutes: 20
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/test_regression
@@ -176,7 +174,6 @@ steps:
  - tests/renderers
  - tests/standalone_tests/lazy_imports.py
  - tests/tokenizers_
-  - tests/reasoning
  - tests/tool_parsers
  - tests/transformers_utils
  - tests/config
@@ -190,7 +187,6 @@ 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 --ignore=reasoning/test_gemma4_reasoning_parser.py
  - pytest -v -s tool_parsers
  - pytest -v -s transformers_utils
  - pytest -v -s config
@@ -78,6 +78,7 @@ steps:
    - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py -k "not ray"
    - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py

+# These require fix https://github.com/vllm-project/vllm/pull/36280
 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
@@ -100,13 +101,11 @@ steps:
  - vllm/v1/worker/gpu/
  - vllm/v1/worker/gpu_worker.py
  - tests/v1/spec_decode/test_max_len.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_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"
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
  timeout_in_minutes: 45
-  device: h200_18gb
  torch_nightly: true
  source_file_dependencies:
  - vllm/
@@ -38,7 +38,7 @@ steps:
    # Install fast path packages for testing against transformers
    # Note: also needed to run plamo2 model in vLLM
    - 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'
+    - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.5.2'
    # Shard hybrid language model tests
    - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
  parallelism: 2
@@ -53,7 +53,7 @@ steps:
    # Install fast path packages for testing against transformers
    # Note: also needed to run plamo2 model in vLLM
    - 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'
+    - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.5.2'
    - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'
  mirror:
    amd:
@@ -67,7 +67,6 @@ steps:

 - label: Language Models Test (PPL)
  timeout_in_minutes: 110
-  device: h200_18gb
  optional: true
  source_file_dependencies:
  - vllm/
@@ -91,7 +90,6 @@ steps:

 - label: Language Models Test (MTEB)
  timeout_in_minutes: 110
-  device: h200_18gb
  optional: true
  source_file_dependencies:
  - vllm/
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
  timeout_in_minutes: 45
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -20,7 +19,6 @@ steps:

 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
  timeout_in_minutes: 45
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -79,7 +77,6 @@ steps:

 - label: Multi-Modal Processor # 44min
  timeout_in_minutes: 60
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -134,7 +131,6 @@ steps:

 - label: Multi-Modal Models (Extended Pooling)
  optional: true
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal/pooling
@@ -49,7 +49,6 @@ steps:

 - label: PyTorch Fullgraph
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
  - vllm/
  - tests/compile
@@ -61,7 +60,6 @@ steps:
  # if this test fails, it means the nightly torch version is not compatible with some
  # of the dependencies. Please check the error message and add the package to whitelist
  # in /vllm/tools/pre_commit/generate_nightly_torch_test.py
-  device: h200_18gb
  soft_fail: true
  source_file_dependencies:
  - requirements/nightly_torch_test.txt
@@ -1,5 +1,5 @@
 group: Quantization
-depends_on:
+depends_on: 
  - image-build
 steps:
 - label: Quantization
@@ -16,7 +16,7 @@ steps:
  # https://github.com/pytorch/ao/issues/2919, we'll have to skip new torchao tests for now
  # we can only upgrade after this is resolved
  # TODO(jerryzh168): resolve the above comment
-  - uv pip install --system torchao==0.17.0 --index-url https://download.pytorch.org/whl/cu130
+  - uv pip install --system torchao==0.14.1 --index-url https://download.pytorch.org/whl/cu129
  - uv pip install --system conch-triton-kernels
  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py

@@ -7,7 +7,6 @@ steps:
  # If this fails, it means the PR introduces a dependency that
  # conflicts with Ray's dependency constraints.
  # See https://github.com/vllm-project/vllm/issues/33599
-  device: h200_18gb
  soft_fail: true
  timeout_in_minutes: 10
  source_file_dependencies:
@@ -4,7 +4,6 @@ depends_on:
 steps:
 - label: Spec Decode Eagle
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
    - vllm/v1/worker/gpu/spec_decode/
@@ -14,7 +13,6 @@ steps:

 - label: Spec Decode Speculators + MTP
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
    - vllm/v1/worker/gpu/spec_decode/
@@ -25,7 +23,6 @@ steps:

 - label: Spec Decode Ngram + Suffix
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
    - vllm/v1/worker/gpu/spec_decode/
@@ -35,7 +32,6 @@ steps:

 - label: Spec Decode Draft Model
  timeout_in_minutes: 30
-  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
    - vllm/v1/worker/gpu/spec_decode/
@@ -39,7 +39,7 @@ repos:
  rev: 0.11.1
  hooks:
    - id: pip-compile
-      args: [requirements/test.in, -c, requirements/common.txt, -o, requirements/test.txt, --index-strategy, unsafe-best-match, --torch-backend, cu130, --python-platform, x86_64-manylinux_2_28, --python-version, "3.12"]
+      args: [requirements/test.in, -c, requirements/common.txt, -o, requirements/test.txt, --index-strategy, unsafe-best-match, --torch-backend, cu129, --python-platform, x86_64-manylinux_2_28, --python-version, "3.12"]
      files: ^requirements/test\.(in|txt)$
    - id: pip-compile
      alias: pip-compile-rocm
@@ -56,8 +56,8 @@ endif()
 # requirements.txt files and should be kept consistent.  The ROCm torch
 # versions are derived from docker/Dockerfile.rocm
 #
-set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
-set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
+set(TORCH_SUPPORTED_VERSION_CUDA "2.10.0")
+set(TORCH_SUPPORTED_VERSION_ROCM "2.10.0")

 #
 # Try to find python package with an executable that exactly matches
@@ -225,8 +225,8 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
  # Certain HIP functions are marked as [[nodiscard]], yet vllm ignores the result which generates
  # a lot of warnings that always mask real issues. Suppressing until this is properly addressed.
  #
-  set(CMAKE_${VLLM_GPU_LANG}_FLAGS "${CMAKE_${VLLM_GPU_LANG}_FLAGS} -Wno-unused-result -Wno-unused-value")
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-unused-result -Wno-unused-value")
+  set(CMAKE_${VLLM_GPU_LANG}_FLAGS "${CMAKE_${VLLM_GPU_LANG}_FLAGS} -Wno-unused-result")
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-unused-result")
 endif()

 #
@@ -299,7 +299,6 @@ set(VLLM_EXT_SRC
  "csrc/quantization/w8a8/int8/scaled_quant.cu"
  "csrc/quantization/w8a8/fp8/common.cu"
  "csrc/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
-  "csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
  "csrc/quantization/gguf/gguf_kernel.cu"
  "csrc/quantization/activation_kernels.cu"
  "csrc/cuda_utils_kernels.cu"
@@ -341,7 +340,8 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")

  list(APPEND VLLM_EXT_SRC
    "csrc/quantization/awq/gemm_kernels.cu"
-    "csrc/cutlass_extensions/common.cpp")
+    "csrc/cutlass_extensions/common.cpp"
+    "csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu")

  set_gencode_flags_for_srcs(
    SRCS "${VLLM_EXT_SRC}"
@@ -23,54 +23,47 @@ For events, please visit [vllm.ai/events](https://vllm.ai/events) to join us.

 vLLM is a fast and easy-to-use library for LLM inference and serving.

-Originally developed in the [Sky Computing Lab](https://sky.cs.berkeley.edu) at UC Berkeley, vLLM has grown into one of the most active open-source AI projects built and maintained by a diverse community of many dozens of academic institutions and companies from over 2000 contributors.
+Originally developed in the [Sky Computing Lab](https://sky.cs.berkeley.edu) at UC Berkeley, vLLM has evolved into a community-driven project with contributions from both academia and industry.

 vLLM is fast with:

 - State-of-the-art serving throughput
 - Efficient management of attention key and value memory with [**PagedAttention**](https://blog.vllm.ai/2023/06/20/vllm.html)
- Continuous batching of incoming requests, chunked prefill, prefix caching
- Fast and flexible model execution with piecewise and full CUDA/HIP graphs
- Quantization: FP8, MXFP8/MXFP4, NVFP4, INT8, INT4, GPTQ/AWQ, GGUF, compressed-tensors, ModelOpt, TorchAO, and [more](https://docs.vllm.ai/en/latest/features/quantization/index.html)
- Optimized attention kernels including FlashAttention, FlashInfer, TRTLLM-GEN, FlashMLA, and Triton
- Optimized GEMM/MoE kernels for various precisions using CUTLASS, TRTLLM-GEN, CuTeDSL
- Speculative decoding including n-gram, suffix, EAGLE, DFlash
- Automatic kernel generation and graph-level transformations using torch.compile
- Disaggregated prefill, decode, and encode
+- Continuous batching of incoming requests
+- Fast model execution with CUDA/HIP graph
+- Quantizations: [GPTQ](https://arxiv.org/abs/2210.17323), [AWQ](https://arxiv.org/abs/2306.00978), [AutoRound](https://arxiv.org/abs/2309.05516), INT4, INT8, and FP8
+- Optimized CUDA kernels, including integration with FlashAttention and FlashInfer
+- Speculative decoding
+- Chunked prefill

 vLLM is flexible and easy to use with:

 - Seamless integration with popular Hugging Face models
 - High-throughput serving with various decoding algorithms, including *parallel sampling*, *beam search*, and more
- Tensor, pipeline, data, expert, and context parallelism for distributed inference
+- Tensor, pipeline, data and expert parallelism support for distributed inference
 - Streaming outputs
- Generation of structured outputs using xgrammar or guidance
- Tool calling and reasoning parsers
- OpenAI-compatible API server, plus Anthropic Messages API and gRPC support
- Efficient multi-LoRA support for dense and MoE layers
- Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
+- OpenAI-compatible API server
+- Support for NVIDIA GPUs, AMD CPUs and GPUs, Intel CPUs and GPUs, PowerPC CPUs, Arm CPUs, and TPU. Additionally, support for diverse hardware plugins such as Intel Gaudi, IBM Spyre and Huawei Ascend.
+- Prefix caching support
+- Multi-LoRA support

-vLLM seamlessly supports 200+ model architectures on HuggingFace, including:
+vLLM seamlessly supports most popular open-source models on HuggingFace, including:

- Decoder-only LLMs (e.g., Llama, Qwen, Gemma)
- Mixture-of-Expert LLMs (e.g., Mixtral, DeepSeek-V3, Qwen-MoE, GPT-OSS)
- Hybrid attention and state-space models (e.g., Mamba, Qwen3.5)
- Multi-modal models (e.g., LLaVA, Qwen-VL, Pixtral)
- Embedding and retrieval models (e.g., E5-Mistral, GTE, ColBERT)
- Reward and classification models (e.g., Qwen-Math)
+- Transformer-like LLMs (e.g., Llama)
+- Mixture-of-Expert LLMs (e.g., Mixtral, Deepseek-V2 and V3)
+- Embedding Models (e.g., E5-Mistral)
+- Multi-modal LLMs (e.g., LLaVA)

 Find the full list of supported models [here](https://docs.vllm.ai/en/latest/models/supported_models.html).

 ## Getting Started

-Install vLLM with [`uv`](https://docs.astral.sh/uv/) (recommended) or `pip`:
+Install vLLM with `pip` or [from source](https://docs.vllm.ai/en/latest/getting_started/installation/gpu/index.html#build-wheel-from-source):

 ```bash
-uv pip install vllm
+pip install vllm
 ```

-Or [build from source](https://docs.vllm.ai/en/latest/getting_started/installation/gpu/index.html#build-wheel-from-source) for development.
-
 Visit our [documentation](https://docs.vllm.ai/en/latest/) to learn more.

 - [Installation](https://docs.vllm.ai/en/latest/getting_started/installation.html)
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG f5bc33cfc02c744d24a2e9d50e6db656de40611c
+          GIT_TAG c0ec424fd8a546d0cbbf4bf050bbcfe837c55afb
          GIT_PROGRESS TRUE
          # Don't share the vllm-flash-attn build between build types
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
@@ -87,30 +87,18 @@ endforeach()
 #
 add_custom_target(_vllm_fa4_cutedsl_C)

-# Install flash_attn/cute directory (needed for FA4).
-# When using a local source dir (VLLM_FLASH_ATTN_SRC_DIR), create a symlink
-# so edits to cute-dsl Python files take effect immediately without rebuilding.
-# Otherwise, copy files and transform flash_attn.cute imports to
-# vllm.vllm_flash_attn.cute to match our package structure.
-if(VLLM_FLASH_ATTN_SRC_DIR)
-  install(CODE "
-    set(LINK_TARGET \"${vllm-flash-attn_SOURCE_DIR}/flash_attn/cute\")
-    set(LINK_NAME \"\${CMAKE_INSTALL_PREFIX}/vllm/vllm_flash_attn/cute\")
-    file(MAKE_DIRECTORY \"\${CMAKE_INSTALL_PREFIX}/vllm/vllm_flash_attn\")
-    file(REMOVE_RECURSE \"\${LINK_NAME}\")
-    file(CREATE_LINK \"\${LINK_TARGET}\" \"\${LINK_NAME}\" SYMBOLIC)
-  " COMPONENT _vllm_fa4_cutedsl_C)
-else()
-  install(CODE "
-    file(GLOB_RECURSE CUTE_PY_FILES \"${vllm-flash-attn_SOURCE_DIR}/flash_attn/cute/*.py\")
-    foreach(SRC_FILE \${CUTE_PY_FILES})
-      file(RELATIVE_PATH REL_PATH \"${vllm-flash-attn_SOURCE_DIR}/flash_attn/cute\" \${SRC_FILE})
-      set(DST_FILE \"\${CMAKE_INSTALL_PREFIX}/vllm/vllm_flash_attn/cute/\${REL_PATH}\")
-      get_filename_component(DST_DIR \${DST_FILE} DIRECTORY)
-      file(MAKE_DIRECTORY \${DST_DIR})
-      file(READ \${SRC_FILE} FILE_CONTENTS)
-      string(REPLACE \"flash_attn.cute\" \"vllm.vllm_flash_attn.cute\" FILE_CONTENTS \"\${FILE_CONTENTS}\")
-      file(WRITE \${DST_FILE} \"\${FILE_CONTENTS}\")
-    endforeach()
-  " COMPONENT _vllm_fa4_cutedsl_C)
-endif()
+# Copy flash_attn/cute directory (needed for FA4) and transform imports
+# The cute directory uses flash_attn.cute imports internally, which we replace
+# with vllm.vllm_flash_attn.cute to match our package structure.
+install(CODE "
+  file(GLOB_RECURSE CUTE_PY_FILES \"${vllm-flash-attn_SOURCE_DIR}/flash_attn/cute/*.py\")
+  foreach(SRC_FILE \${CUTE_PY_FILES})
+    file(RELATIVE_PATH REL_PATH \"${vllm-flash-attn_SOURCE_DIR}/flash_attn/cute\" \${SRC_FILE})
+    set(DST_FILE \"\${CMAKE_INSTALL_PREFIX}/vllm/vllm_flash_attn/cute/\${REL_PATH}\")
+    get_filename_component(DST_DIR \${DST_FILE} DIRECTORY)
+    file(MAKE_DIRECTORY \${DST_DIR})
+    file(READ \${SRC_FILE} FILE_CONTENTS)
+    string(REPLACE \"flash_attn.cute\" \"vllm.vllm_flash_attn.cute\" FILE_CONTENTS \"\${FILE_CONTENTS}\")
+    file(WRITE \${DST_FILE} \"\${FILE_CONTENTS}\")
+  endforeach()
+" COMPONENT _vllm_fa4_cutedsl_C)
@@ -91,9 +91,9 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,

  if (n == 0) return;

-  int64_t* src_data = src_ptrs.mutable_data_ptr<int64_t>();
-  int64_t* dst_data = dst_ptrs.mutable_data_ptr<int64_t>();
-  int64_t* size_data = sizes.mutable_data_ptr<int64_t>();
+  const int64_t* src_data = src_ptrs.data_ptr<int64_t>();
+  const int64_t* dst_data = dst_ptrs.data_ptr<int64_t>();
+  const int64_t* size_data = sizes.data_ptr<int64_t>();

  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

@@ -107,24 +107,15 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
  CUmemcpyAttributes attr = {};
  attr.srcAccessOrder = CU_MEMCPY_SRC_ACCESS_ORDER_STREAM;
  size_t attrs_idx = 0;
-  #if defined(CUDA_VERSION) && CUDA_VERSION >= 13000
-  CUresult result = cuMemcpyBatchAsync(
-      reinterpret_cast<CUdeviceptr*>(dst_data),
-      reinterpret_cast<CUdeviceptr*>(src_data),
-      reinterpret_cast<size_t*>(size_data), static_cast<size_t>(n), &attr,
-      &attrs_idx, 1, static_cast<CUstream>(stream));
-  TORCH_CHECK(result == CUDA_SUCCESS, "cuMemcpyBatchAsync failed with error ",
-              result);
-  #else
  size_t fail_idx = 0;
  CUresult result = cuMemcpyBatchAsync(
-      reinterpret_cast<CUdeviceptr*>(dst_data),
-      reinterpret_cast<CUdeviceptr*>(src_data),
-      reinterpret_cast<size_t*>(size_data), static_cast<size_t>(n), &attr,
-      &attrs_idx, 1, &fail_idx, static_cast<CUstream>(stream));
+      reinterpret_cast<CUdeviceptr*>(const_cast<int64_t*>(dst_data)),
+      reinterpret_cast<CUdeviceptr*>(const_cast<int64_t*>(src_data)),
+      reinterpret_cast<size_t*>(const_cast<int64_t*>(size_data)),
+      static_cast<size_t>(n), &attr, &attrs_idx, 1, &fail_idx,
+      static_cast<CUstream>(stream));
  TORCH_CHECK(result == CUDA_SUCCESS, "cuMemcpyBatchAsync failed at index ",
              fail_idx, " with error ", result);
-  #endif
 #else
  // Fallback for CUDA < 12.8 and ROCm: individual async copies.
  // cudaMemcpyDefault lets the driver infer direction from pointer types.
@@ -53,7 +53,7 @@ class TileGemm82 {
                         const int64_t ldb, const int64_t ldc,
                         const int32_t block_size, const int32_t dynamic_k_size,
                         const bool accum_c) {
-    static_assert(0 < M && M <= 8);
+    static_assert(0 < M <= 8);
    using load_vec_t = typename VecTypeTrait<kv_cache_t>::vec_t;

    kv_cache_t* __restrict__ curr_b_0 = b_tile;
@@ -68,7 +68,7 @@ class TileGemm161 {
                         const int64_t ldb, const int64_t ldc,
                         const int32_t block_size, const int32_t dynamic_k_size,
                         const bool accum_c) {
-    static_assert(0 < M && M <= 16);
+    static_assert(0 < M <= 16);
    using load_vec_t = typename VecTypeTrait<kv_cache_t>::vec_t;

    kv_cache_t* __restrict__ curr_b_0 = b_tile;
@@ -39,7 +39,7 @@ class TileGemm82 {

  template <int32_t M>
  static void gemm_micro(DEFINE_CPU_MICRO_GEMM_PARAMS) {
-    static_assert(0 < M && M <= 8);
+    static_assert(0 < M <= 8);
    using load_vec_t = typename cpu_utils::VecTypeTrait<scalar_t>::vec_t;

    scalar_t* __restrict__ curr_b_0 = b_ptr;
@@ -55,8 +55,7 @@ struct Counter {

 inline int64_t get_available_l2_size() {
  static int64_t size = []() {
-    auto caps = at::cpu::get_cpu_capabilities();
-    const uint32_t l2_cache_size = caps.at("l2_cache_size").toInt();
+    const uint32_t l2_cache_size = at::cpu::L2_cache_size();
    return l2_cache_size >> 1;  // use 50% of L2 cache
  }();
  return size;
@@ -114,9 +114,9 @@ void top_k_per_row_decode(const torch::Tensor& logits, int64_t next_n,
                          int64_t numRows, int64_t stride0, int64_t stride1,
                          int64_t topK);

-void persistent_topk(const torch::Tensor& logits, const torch::Tensor& lengths,
-                     torch::Tensor& output, torch::Tensor& workspace, int64_t k,
-                     int64_t max_seq_len);
+void large_context_topk(const torch::Tensor& score, torch::Tensor& indices,
+                        const torch::Tensor& lengths,
+                        std::optional<torch::Tensor> row_starts_opt);

 void rms_norm_static_fp8_quant(torch::Tensor& out, torch::Tensor& input,
                               torch::Tensor& weight, torch::Tensor& scale,
@@ -143,11 +143,13 @@ void rms_norm_per_block_quant(torch::Tensor& out, torch::Tensor const& input,
                              std::optional<torch::Tensor> residual,
                              int64_t group_size, bool is_scale_transposed);

+#ifndef USE_ROCM
 void silu_and_mul_per_block_quant(torch::Tensor& out,
                                  torch::Tensor const& input,
                                  torch::Tensor& scales, int64_t group_size,
                                  std::optional<torch::Tensor> scale_ub,
                                  bool is_scale_transposed);
+#endif

 void rotary_embedding(torch::Tensor& positions, torch::Tensor& query,
                      std::optional<torch::Tensor> key, int64_t head_size,
@@ -6,7 +6,7 @@

 #include "libtorch_stable/quantization/vectorization.cuh"
 // TODO(luka/varun):refactor common.cuh to use this file instead
-#include "../w8a8/fp8/common.cuh"
+#include "quantization/w8a8/fp8/common.cuh"

 namespace vllm {

@@ -1,7 +1,7 @@
 #pragma once

 #include "libtorch_stable/quantization/vectorization.cuh"
-#include "../../utils.cuh"
+#include "quantization/utils.cuh"

 #include <cmath>

@@ -1,154 +1,373 @@
-// Persistent TopK kernel for DeepSeek V3 sparse attention indexer.
-// See persistent_topk.cuh for kernel implementation.
+// Portions of this file are adapted from SGLang PR:
+// https://github.com/sgl-project/sglang/pull/11194
+// and
+// https://github.com/sgl-project/sglang/pull/17747

-#include <torch/all.h>
-#include <ATen/cuda/CUDAContext.h>
-#include <cuda_runtime.h>
-#include <algorithm>
+#include "cuda_compat.h"
+#include "dispatch_utils.h"
+
+#include <torch/cuda.h>
+#include <c10/cuda/CUDAGuard.h>

 #ifndef USE_ROCM
-  #include "persistent_topk.cuh"
-#endif
-
-void persistent_topk(const torch::Tensor& logits, const torch::Tensor& lengths,
-                     torch::Tensor& output, torch::Tensor& workspace, int64_t k,
-                     int64_t max_seq_len) {
-#ifndef USE_ROCM
-  TORCH_CHECK(logits.is_cuda(), "logits must be CUDA tensor");
-  TORCH_CHECK(lengths.is_cuda(), "lengths must be CUDA tensor");
-  TORCH_CHECK(output.is_cuda(), "output must be CUDA tensor");
-  TORCH_CHECK(logits.dtype() == torch::kFloat32, "Only float32 supported");
-  TORCH_CHECK(lengths.dtype() == torch::kInt32, "lengths must be int32");
-  TORCH_CHECK(output.dtype() == torch::kInt32, "output must be int32");
-  TORCH_CHECK(logits.dim() == 2, "logits must be 2D");
-  TORCH_CHECK(lengths.dim() == 1, "lengths must be 1D");
-  TORCH_CHECK(output.dim() == 2, "output must be 2D");
-
-  const int64_t num_rows = logits.size(0);
-  const int64_t stride = logits.size(1);
-
-  TORCH_CHECK(lengths.size(0) == num_rows, "lengths size mismatch");
-  TORCH_CHECK(output.size(0) == num_rows && output.size(1) == k,
-              "output size mismatch");
-  namespace P = vllm::persistent;
-
-  TORCH_CHECK(k == P::TopK, "k must be 2048");
-  TORCH_CHECK(k <= stride, "k out of range");
-
-  cudaStream_t stream = at::cuda::getCurrentCUDAStream();
-
-  static int num_sms = 0;
-  static int max_smem_per_block = 0;
-  if (num_sms == 0) {
-    int device;
-    cudaGetDevice(&device);
-    cudaDeviceGetAttribute(&num_sms, cudaDevAttrMultiProcessorCount, device);
-    cudaDeviceGetAttribute(&max_smem_per_block,
-                           cudaDevAttrMaxSharedMemoryPerBlockOptin, device);
-  }
-
-  if (num_rows > 32 && max_smem_per_block >= 128 * 1024) {
-    cudaError_t status = vllm::FilteredTopKRaggedTransform<float, int32_t>(
-        logits.data_ptr<float>(), output.data_ptr<int32_t>(),
-        lengths.data_ptr<int32_t>(), static_cast<uint32_t>(num_rows),
-        static_cast<uint32_t>(k), static_cast<uint32_t>(stride), stream);
-    TORCH_CHECK(status == cudaSuccess,
-                "FilteredTopK failed: ", cudaGetErrorString(status));
-  } else {
-    TORCH_CHECK(workspace.is_cuda(), "workspace must be CUDA tensor");
-    TORCH_CHECK(workspace.dtype() == torch::kUInt8, "workspace must be uint8");
-
-    // Smem cap: smaller smem → more CTAs/group → more per-row parallelism for
-    // large path. Empirically tuned.
-    int effective_max_smem;
-    if (num_rows <= 4) {
-      effective_max_smem =
-          std::min(max_smem_per_block, static_cast<int>(P::kSmemMedium));
-    } else if (num_rows <= 8) {
-      constexpr int kSmemCapMedium = 48 * 1024;
-      effective_max_smem = std::min(max_smem_per_block, kSmemCapMedium);
-    } else {
-      effective_max_smem = max_smem_per_block;
-    }
-
-    size_t available_for_ordered =
-        static_cast<size_t>(effective_max_smem) - P::kFixedSmemLarge;
-    uint32_t max_chunk_elements =
-        static_cast<uint32_t>(available_for_ordered / sizeof(uint32_t));
-
-    uint32_t vec_size = 1;
-    if (stride % 4 == 0)
-      vec_size = 4;
-    else if (stride % 2 == 0)
-      vec_size = 2;
-
-    max_chunk_elements = (max_chunk_elements / vec_size) * vec_size;
-    uint32_t min_chunk = vec_size * P::kThreadsPerBlock;
-    if (max_chunk_elements < min_chunk) max_chunk_elements = min_chunk;
-
-    uint32_t ctas_per_group =
-        (static_cast<uint32_t>(stride) + max_chunk_elements - 1) /
-        max_chunk_elements;
-    uint32_t chunk_size =
-        (static_cast<uint32_t>(stride) + ctas_per_group - 1) / ctas_per_group;
-    chunk_size = ((chunk_size + vec_size - 1) / vec_size) * vec_size;
-    if (chunk_size > max_chunk_elements) chunk_size = max_chunk_elements;
-
-    size_t smem_size = P::kFixedSmemLarge + chunk_size * sizeof(uint32_t);
-    if (smem_size < P::kSmemMedium) smem_size = P::kSmemMedium;
-
-    int occupancy = 1;
-    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
-        &occupancy, P::persistent_topk_kernel<4>, P::kThreadsPerBlock,
-        smem_size);
-    if (occupancy < 1) occupancy = 1;
-
-    uint32_t max_resident_ctas = static_cast<uint32_t>(num_sms) * occupancy;
-    uint32_t num_groups = std::min(max_resident_ctas / ctas_per_group,
-                                   static_cast<uint32_t>(num_rows));
-    if (num_groups == 0) num_groups = 1;
-    uint32_t total_ctas = num_groups * ctas_per_group;
-
-    size_t state_bytes = num_groups * sizeof(P::RadixRowState);
-    TORCH_CHECK(workspace.size(0) >= static_cast<int64_t>(state_bytes),
-                "workspace too small, need ", state_bytes, " bytes");
-
-    P::PersistentTopKParams params;
-    params.input = logits.data_ptr<float>();
-    params.output = output.data_ptr<int32_t>();
-    params.lengths = lengths.data_ptr<int32_t>();
-    params.num_rows = static_cast<uint32_t>(num_rows);
-    params.stride = static_cast<uint32_t>(stride);
-    params.chunk_size = chunk_size;
-    params.row_states =
-        reinterpret_cast<P::RadixRowState*>(workspace.data_ptr<uint8_t>());
-    params.ctas_per_group = ctas_per_group;
-    params.max_seq_len = static_cast<uint32_t>(max_seq_len);
-
-  #define LAUNCH_PERSISTENT(VS)                                               \
-    do {                                                                      \
-      auto kernel = &P::persistent_topk_kernel<VS>;                           \
-      cudaError_t err = cudaFuncSetAttribute(                                 \
-          kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size);    \
-      TORCH_CHECK(err == cudaSuccess,                                         \
-                  "Failed to set smem: ", cudaGetErrorString(err));           \
-      kernel<<<total_ctas, P::kThreadsPerBlock, smem_size, stream>>>(params); \
-    } while (0)
-
-    if (vec_size == 4) {
-      LAUNCH_PERSISTENT(4);
-    } else if (vec_size == 2) {
-      LAUNCH_PERSISTENT(2);
-    } else {
-      LAUNCH_PERSISTENT(1);
-    }
-  #undef LAUNCH_PERSISTENT
-  }
-
-  cudaError_t err = cudaGetLastError();
-  TORCH_CHECK(err == cudaSuccess,
-              "persistent_topk failed: ", cudaGetErrorString(err));
+  #include <cub/cub.cuh>
 #else
-  TORCH_CHECK(false, "persistent_topk is not supported on ROCm");
+  #include <hipcub/hipcub.hpp>
 #endif
+
+namespace vllm {
+
+constexpr int TopK = 2048;              // DeepSeek V3 sparse attention top-k
+constexpr int kThreadsPerBlock = 1024;  // Threads per block
+
+// Shared memory budget
+#if defined(USE_ROCM)
+constexpr size_t kSmem = 48 * 1024;  // ROCm default: 48KB
+#else
+// Reduced from 128KB to 32KB to improve occupancy.
+// Each radix pass needs at most ~TopK candidates in the threshold bin,
+// so 4K entries per round (2 rounds = 8K entries = 32KB) is sufficient.
+constexpr size_t kSmem = 8 * 1024 * sizeof(uint32_t);  // 32KB (bytes)
+#endif
+
+struct FastTopKParams {
+  const float* __restrict__ input;         // [batch, seq_len] Logits
+  const int32_t* __restrict__ row_starts;  // [batch] Offset into each row
+                                           // (optional)
+  int32_t* __restrict__ indices;           // [batch, TopK] Output top-k indices
+  int32_t* __restrict__ lengths;           // [batch] Sequence lengths per row
+  int64_t input_stride;                    // Stride between rows
+};
+
+__device__ __forceinline__ auto convert_to_uint32_v2(float x) -> uint32_t {
+  uint32_t bits = __float_as_uint(x);
+  return (bits & 0x80000000u) ? ~bits : (bits | 0x80000000u);
 }
+
+__device__ __forceinline__ auto convert_to_uint8(float x) -> uint8_t {
+  __half h = __float2half_rn(x);
+  uint16_t bits = __half_as_ushort(h);
+  uint16_t key = (bits & 0x8000) ? static_cast<uint16_t>(~bits)
+                                 : static_cast<uint16_t>(bits | 0x8000);
+  return static_cast<uint8_t>(key >> 8);
+}
+
+__device__ void naive_topk_cuda(const float* __restrict__ logits,
+                                int32_t* __restrict__ output_indices,
+                                int32_t seq_len) {
+  const int thread_id = threadIdx.x;
+  for (int i = thread_id; i < TopK; i += kThreadsPerBlock) {
+    output_indices[i] = (i < seq_len) ? i : -1;
+  }
+}
+
+// Adapted from:
+// https://github.com/sgl-project/sglang/blob/v0.5.8/sgl-kernel/csrc/elementwise/topk.cu#L87
+// by: DarkSharpness
+// which at the same time is an optimized topk kernel copied from tilelang
+// kernel
+__device__ void fast_topk_cuda_tl(
+    const float* __restrict__ logits,  // Input logits [seq_len]
+    int* __restrict__ output_indices,  // Output top-k indices [TopK]
+    int logits_offset,                 // Starting offset in logits array
+    int seq_len)                       // Number of valid logits to process
+{
+  constexpr int RADIX = 256;
+  constexpr int MAX_BUFFERED_ITEMS = kSmem / (2 * sizeof(int));
+
+  alignas(128) __shared__ int shared_histogram[2][RADIX + 128];
+  alignas(128) __shared__ int shared_output_count;
+  alignas(128) __shared__ int shared_threshold_bin;
+  alignas(128) __shared__ int shared_buffered_count[2];
+
+  extern __shared__ int buffered_indices[][MAX_BUFFERED_ITEMS];
+
+  const int thread_id = threadIdx.x;
+  int remaining_k = TopK;
+
+  // Pass 0: Build coarse 8-bit histogram using FP16 high bits
+  if (thread_id < RADIX + 1) {
+    shared_histogram[0][thread_id] = 0;
+  }
+  __syncthreads();
+
+  for (int idx = thread_id; idx < seq_len; idx += kThreadsPerBlock) {
+    const auto bin = convert_to_uint8(logits[idx + logits_offset]);
+    ::atomicAdd(&shared_histogram[0][bin], 1);
+  }
+  __syncthreads();
+
+  // Helper: Compute cumulative sum (suffix sum) over histogram using ping-pong
+  // buffers
+  auto compute_cumulative_sum = [&]() {
+    static_assert(1 << 8 == RADIX,
+                  "Radix must be 256 for 8 unrolled iterations");
+#pragma unroll 8
+    for (int i = 0; i < 8; ++i) {
+      if (C10_LIKELY(thread_id < RADIX)) {
+        const int stride = 1 << i;
+        const int src_buffer = i & 1;
+        const int dst_buffer = src_buffer ^ 1;
+
+        int value = shared_histogram[src_buffer][thread_id];
+        if (thread_id < RADIX - stride) {
+          value += shared_histogram[src_buffer][thread_id + stride];
+        }
+        shared_histogram[dst_buffer][thread_id] = value;
+      }
+      __syncthreads();
+    }
+  };
+
+  compute_cumulative_sum();
+
+  // Find threshold bin where cumsum crosses remaining_k
+  if (thread_id < RADIX && shared_histogram[0][thread_id] > remaining_k &&
+      shared_histogram[0][thread_id + 1] <= remaining_k) {
+    shared_threshold_bin = thread_id;
+    shared_buffered_count[0] = 0;
+    shared_output_count = 0;
+  }
+  __syncthreads();
+
+  const int threshold_bin = shared_threshold_bin;
+  remaining_k -= shared_histogram[0][threshold_bin + 1];
+
+  // Early exit if threshold bin perfectly matches remaining_k
+  if (remaining_k == 0) {
+    for (int idx = thread_id; idx < seq_len; idx += kThreadsPerBlock) {
+      const int bin = convert_to_uint8(logits[idx + logits_offset]);
+      if (bin > threshold_bin) {
+        const int output_pos = ::atomicAdd(&shared_output_count, 1);
+        output_indices[output_pos] = idx;
+      }
+    }
+    __syncthreads();
+    return;
+  }
+
+  // Prepare for refinement passes: Process threshold bin
+  __syncthreads();
+  if (thread_id < RADIX + 1) {
+    shared_histogram[0][thread_id] = 0;
+  }
+  __syncthreads();
+
+  // Scan all elements and:
+  // 1. Write indices > threshold_bin to output
+  // 2. Buffer indices == threshold_bin for refinement
+  // 3. Build histogram for next refinement pass (fused optimization)
+  for (int idx = thread_id; idx < seq_len; idx += kThreadsPerBlock) {
+    const float logit_value = logits[idx + logits_offset];
+    const int bin = convert_to_uint8(logit_value);
+
+    if (bin > threshold_bin) {
+      // in top-k, write to output
+      const int output_pos = ::atomicAdd(&shared_output_count, 1);
+      output_indices[output_pos] = idx;
+    } else if (bin == threshold_bin) {
+      // Candidate for top-k, needs refinement
+      const int buffer_pos = ::atomicAdd(&shared_buffered_count[0], 1);
+      if (C10_LIKELY(buffer_pos < MAX_BUFFERED_ITEMS)) {
+        buffered_indices[0][buffer_pos] = idx;
+        // Fused: Build histogram for next pass
+        const uint32_t fp32_bits = convert_to_uint32_v2(logit_value);
+        const int next_bin = (fp32_bits >> 24) & 0xFF;
+        ::atomicAdd(&shared_histogram[0][next_bin], 1);
+      }
+    }
+  }
+  __syncthreads();
+
+  // ============================================================================
+  // Passes 1-4: Refine using 8-bit passes over FP32 bits
+  // ============================================================================
+  // FP32 bits [31:0] split into 4 bytes processed MSB-first:
+  // Pass 1: bits [31:24], Pass 2: bits [23:16], Pass 3: bits [15:8], Pass 4:
+  // bits [7:0]
+#pragma unroll 4
+  for (int pass = 0; pass < 4; ++pass) {
+    __shared__ int shared_final_k;  // For final pass: remaining slots to fill
+    const int src_buffer = pass % 2;
+    const int dst_buffer = src_buffer ^ 1;
+
+    // Clamp buffered count to prevent overflow
+    const int raw_buffered = shared_buffered_count[src_buffer];
+    const int num_buffered =
+        (raw_buffered < MAX_BUFFERED_ITEMS) ? raw_buffered : MAX_BUFFERED_ITEMS;
+
+    compute_cumulative_sum();
+
+    // Find threshold bin for this pass
+    if (thread_id < RADIX && shared_histogram[0][thread_id] > remaining_k &&
+        shared_histogram[0][thread_id + 1] <= remaining_k) {
+      shared_threshold_bin = thread_id;
+      shared_buffered_count[dst_buffer] = 0;
+      shared_final_k = remaining_k - shared_histogram[0][thread_id + 1];
+    }
+    __syncthreads();
+
+    const int threshold_bin = shared_threshold_bin;
+    remaining_k -= shared_histogram[0][threshold_bin + 1];
+
+    // Bit offset for this pass: 24, 16, 8, 0
+    const int bit_offset = 24 - pass * 8;
+
+    // Early exit if threshold bin perfectly matches
+    if (remaining_k == 0) {
+      for (int i = thread_id; i < num_buffered; i += kThreadsPerBlock) {
+        const int idx = buffered_indices[src_buffer][i];
+        const uint32_t fp32_bits =
+            convert_to_uint32_v2(logits[idx + logits_offset]);
+        const int bin = (fp32_bits >> bit_offset) & 0xFF;
+        if (bin > threshold_bin) {
+          const int output_pos = ::atomicAdd(&shared_output_count, 1);
+          output_indices[output_pos] = idx;
+        }
+      }
+      __syncthreads();
+      break;
+    }
+
+    // Continue refinement
+    __syncthreads();
+    if (thread_id < RADIX + 1) {
+      shared_histogram[0][thread_id] = 0;
+    }
+    __syncthreads();
+
+    for (int i = thread_id; i < num_buffered; i += kThreadsPerBlock) {
+      const int idx = buffered_indices[src_buffer][i];
+      const float logit_value = logits[idx + logits_offset];
+      const uint32_t fp32_bits = convert_to_uint32_v2(logit_value);
+      const int bin = (fp32_bits >> bit_offset) & 0xFF;
+
+      if (bin > threshold_bin) {
+        // Definitely in top-k
+        const int output_pos = ::atomicAdd(&shared_output_count, 1);
+        output_indices[output_pos] = idx;
+      } else if (bin == threshold_bin) {
+        if (pass == 3) {
+          // Final pass (bits [7:0]): No more refinement possible
+          // Fill remaining slots in reverse order to maintain descending order
+          const int slot = ::atomicAdd(&shared_final_k, -1);
+          if (slot > 0) {
+            output_indices[TopK - slot] = idx;
+          }
+        } else {
+          // Buffer for next pass and build next histogram
+          const int buffer_pos =
+              ::atomicAdd(&shared_buffered_count[dst_buffer], 1);
+          if (C10_LIKELY(buffer_pos < MAX_BUFFERED_ITEMS)) {
+            buffered_indices[dst_buffer][buffer_pos] = idx;
+            // Fused: Build histogram for next pass
+            const int next_bit_offset = bit_offset - 8;
+            const int next_bin = (fp32_bits >> next_bit_offset) & 0xFF;
+            ::atomicAdd(&shared_histogram[0][next_bin], 1);
+          }
+        }
+      }
+    }
+    __syncthreads();
+  }
+}
+
+__global__ __launch_bounds__(kThreadsPerBlock) void topk_kernel(
+    const FastTopKParams params) {
+  const auto& [input, row_starts, indices, lengths, input_stride] = params;
+  const uint64_t batch_idx = blockIdx.x;
+  const int logits_offset = row_starts == nullptr ? 0 : row_starts[batch_idx];
+  const int seq_len = lengths[batch_idx];
+  int* output_indices = indices + batch_idx * TopK;
+  const float* logits = input + batch_idx * input_stride;
+
+  if (seq_len <= TopK) {
+    // Shortcut: All elements are in top-k
+    return naive_topk_cuda(logits, output_indices, seq_len);
+  } else {
+    return fast_topk_cuda_tl(logits, output_indices, logits_offset, seq_len);
+  }
+}
+
+FastTopKParams get_params(
+    const at::Tensor& score, const at::Tensor& lengths,
+    std::optional<at::Tensor> row_starts_opt = std::nullopt,
+    std::optional<at::Tensor> indices_opt = std::nullopt) {
+  const int64_t batch_size = score.size(0);
+
+  TORCH_CHECK(score.dim() == 2 && score.stride(1) == 1,
+              "score must be 2D with contiguous rows");
+  TORCH_CHECK(lengths.dim() == 1 && lengths.is_contiguous() &&
+                  lengths.size(0) == batch_size,
+              "lengths must be 1D contiguous with size matching batch");
+
+  const int32_t* row_starts_ptr = nullptr;
+  if (row_starts_opt.has_value()) {
+    const auto& row_starts = *row_starts_opt;
+    TORCH_CHECK(row_starts.dim() == 1 && row_starts.size(0) == batch_size,
+                "row_starts must be 1D with size matching batch");
+    row_starts_ptr = row_starts.data_ptr<int32_t>();
+  }
+
+  int32_t* indices_ptr = nullptr;
+  if (indices_opt.has_value()) {
+    const auto& indices = *indices_opt;
+    TORCH_CHECK(indices.dim() == 2 && indices.is_contiguous() &&
+                    indices.size(0) == batch_size && indices.size(1) == TopK,
+                "indices must be 2D contiguous [batch, TopK]");
+    indices_ptr = indices.data_ptr<int32_t>();
+  }
+
+  return FastTopKParams{
+      .input = score.data_ptr<float>(),
+      .row_starts = row_starts_ptr,
+      .indices = indices_ptr,
+      .lengths = lengths.data_ptr<int32_t>(),
+      .input_stride = score.stride(0),
+  };
+}
+
+template <auto* kernel_func, size_t smem_bytes>
+void setup_kernel_smem_once() {
+  static const cudaError_t result = []() -> cudaError_t {
+#ifdef USE_ROCM
+    auto func_ptr = reinterpret_cast<const void*>(kernel_func);
+#else
+    auto func_ptr = kernel_func;
+#endif
+    return cudaFuncSetAttribute(
+        func_ptr, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_bytes);
+  }();
+
+  TORCH_CHECK(
+      result == cudaSuccess,
+      "Failed to set kernel shared memory limit: ", cudaGetErrorString(result));
+}
+
+}  // namespace vllm
+
+void large_context_topk(
+    const torch::Tensor& logits, torch::Tensor& indices,
+    const torch::Tensor& seq_lens,
+    std::optional<torch::Tensor> row_starts = std::nullopt) {
+  TORCH_CHECK(logits.is_cuda(), "logits must be a CUDA tensor");
+  TORCH_CHECK(indices.is_cuda(), "indices must be a CUDA tensor");
+  TORCH_CHECK(seq_lens.is_cuda(), "seq_lens must be a CUDA tensor");
+  if (row_starts.has_value()) {
+    TORCH_CHECK(row_starts->is_cuda(), "row_starts must be a CUDA tensor");
+  }
+
+  const auto params = vllm::get_params(logits, seq_lens, row_starts, indices);
+  const int64_t batch_size = logits.size(0);
+
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+  const dim3 grid(static_cast<uint32_t>(batch_size));
+  const dim3 block(vllm::kThreadsPerBlock);
+
+  vllm::setup_kernel_smem_once<vllm::topk_kernel, vllm::kSmem>();
+  vllm::topk_kernel<<<grid, block, vllm::kSmem, stream>>>(params);
+
+  const cudaError_t result = cudaGetLastError();
+  TORCH_CHECK(result == cudaSuccess,
+              "large_context_topk kernel failed: ", cudaGetErrorString(result));
+}
@@ -110,18 +110,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      "silu_and_mul_quant(Tensor! result, Tensor input, Tensor scale) -> ()");
  ops.impl("silu_and_mul_quant", torch::kCUDA, &silu_and_mul_quant);

-  // Fused SiLU+Mul + per-block quantization
-  ops.def(
-      "silu_and_mul_per_block_quant("
-      "Tensor! out, "
-      "Tensor input, "
-      "Tensor! scales, "
-      "int group_size, "
-      "Tensor? scale_ub=None, "
-      "bool is_scale_transposed=False) -> ()");
-  ops.impl("silu_and_mul_per_block_quant", torch::kCUDA,
-           &silu_and_mul_per_block_quant);
-
  ops.def("mul_and_silu(Tensor! out, Tensor input) -> ()");
  ops.impl("mul_and_silu", torch::kCUDA, &mul_and_silu);

@@ -197,9 +185,10 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
  ops.impl("top_k_per_row_decode", torch::kCUDA, &top_k_per_row_decode);

  ops.def(
-      "persistent_topk(Tensor logits, Tensor lengths, Tensor! output, "
-      "Tensor workspace, int k, int max_seq_len) -> ()");
-  ops.impl("persistent_topk", torch::kCUDA, &persistent_topk);
+      "large_context_topk(Tensor score, Tensor indices, Tensor lengths, "
+      "Tensor? "
+      "row_starts_opt) -> ()");
+  ops.impl("large_context_topk", torch::kCUDA, &large_context_topk);

  // Layernorm-quant
  // Apply Root Mean Square (RMS) Normalization to the input tensor.
@@ -244,6 +233,17 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {

  // Quantization ops
 #ifndef USE_ROCM
+  // Fused SiLU+Mul + per-block quantization
+  ops.def(
+      "silu_and_mul_per_block_quant("
+      "Tensor! out, "
+      "Tensor input, "
+      "Tensor! scales, "
+      "int group_size, "
+      "Tensor? scale_ub=None, "
+      "bool is_scale_transposed=False) -> ()");
+  ops.impl("silu_and_mul_per_block_quant", torch::kCUDA,
+           &silu_and_mul_per_block_quant);
  // DeepSeek V3 fused A GEMM (SM 9.0+, bf16 only, 1-16 tokens).
  ops.def(
      "dsv3_fused_a_gemm(Tensor! output, Tensor mat_a, Tensor mat_b) -> ()");
@@ -22,7 +22,7 @@
 #   docker buildx bake -f docker/docker-bake.hcl -f docker/versions.json
 # =============================================================================

-ARG CUDA_VERSION=13.0.0
+ARG CUDA_VERSION=12.9.1
 ARG PYTHON_VERSION=3.12
 ARG UBUNTU_VERSION=22.04

@@ -37,7 +37,7 @@ ARG UBUNTU_VERSION=22.04
 # compatibility with other Linux OSes. The main reason for this is that the
 # glibc version is baked into the distro, and binaries built with one glibc
 # version are not backwards compatible with OSes that use an earlier version.
-ARG BUILD_BASE_IMAGE=nvidia/cuda:${CUDA_VERSION}-devel-ubuntu22.04
+ARG BUILD_BASE_IMAGE=nvidia/cuda:${CUDA_VERSION}-devel-ubuntu20.04
 # Using cuda base image with minimal dependencies necessary for JIT compilation (FlashInfer, DeepGEMM, EP kernels)
 ARG FINAL_BASE_IMAGE=nvidia/cuda:${CUDA_VERSION}-base-ubuntu${UBUNTU_VERSION}

@@ -546,21 +546,17 @@ RUN apt-get update -y \
 # Install CUDA development tools for runtime JIT compilation
 # (FlashInfer, DeepGEMM, EP kernels all require compilation at runtime)
 RUN CUDA_VERSION_DASH=$(echo $CUDA_VERSION | cut -d. -f1,2 | tr '.' '-') && \
-    CUDA_VERSION_SHORT=$(echo $CUDA_VERSION | cut -d. -f1,2) && \
    apt-get update -y && \
-    apt-get install -y --no-install-recommends --allow-change-held-packages \
+    apt-get install -y --no-install-recommends \
        cuda-nvcc-${CUDA_VERSION_DASH} \
        cuda-cudart-${CUDA_VERSION_DASH} \
        cuda-nvrtc-${CUDA_VERSION_DASH} \
        cuda-cuobjdump-${CUDA_VERSION_DASH} \
        libcurand-dev-${CUDA_VERSION_DASH} \
-        libcublas-${CUDA_VERSION_DASH} && \
-    # Fixes nccl_allocator requiring nccl.h at runtime
-    # https://github.com/vllm-project/vllm/blob/1336a1ea244fa8bfd7e72751cabbdb5b68a0c11a/vllm/distributed/device_communicators/pynccl_allocator.py#L22
-    # NCCL packages don't use the cuda-MAJOR-MINOR naming convention,
-    # so we pin the version to match our CUDA version
-    NCCL_VER=$(apt-cache madison libnccl-dev | grep "+cuda${CUDA_VERSION_SHORT}" | head -1 | awk -F'|' '{gsub(/^ +| +$/, "", $2); print $2}') && \
-    apt-get install -y --no-install-recommends --allow-change-held-packages libnccl-dev=${NCCL_VER} libnccl2=${NCCL_VER} && \
+        libcublas-${CUDA_VERSION_DASH} \
+        # Fixes nccl_allocator requiring nccl.h at runtime
+        # https://github.com/vllm-project/vllm/blob/1336a1ea244fa8bfd7e72751cabbdb5b68a0c11a/vllm/distributed/device_communicators/pynccl_allocator.py#L22
+        libnccl-dev && \
    rm -rf /var/lib/apt/lists/*

 # Install uv for faster pip installs
@@ -826,7 +822,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
        uv pip install --system -r /tmp/kv_connectors.txt --no-build || ( \
            # if the above fails, install from source
            apt-get update -y && \
-            apt-get install -y --no-install-recommends --allow-change-held-packages ${BUILD_PKGS} && \
+            apt-get install -y --no-install-recommends ${BUILD_PKGS} && \
            uv pip install --system -r /tmp/kv_connectors.txt --no-build-isolation && \
            apt-get purge -y ${BUILD_PKGS} && \
            # clean up -dev packages, keep runtime libraries
@@ -140,7 +140,7 @@ RUN \
    esac; \
    }; \
    remove_packages_not_supported_on_aarch64 && \
-    sed -i 's/^torch==.*/torch==2.11.0/g' requirements/cpu-test.in && \
+    sed -i 's/^torch==.*/torch==2.10.0/g' requirements/cpu-test.in && \
    sed -i 's/torchaudio.*/torchaudio/g' requirements/cpu-test.in && \
    sed -i 's/torchvision.*/torchvision/g' requirements/cpu-test.in && \
    uv pip compile requirements/cpu-test.in -o requirements/cpu-test.txt --index-strategy unsafe-best-match --torch-backend cpu
@@ -390,21 +390,20 @@ ENV MIOPEN_DEBUG_CONV_GEMM=0
 RUN mkdir src && mv vllm src/vllm

 # This is a workaround to ensure pytest exits with the correct status code in CI tests.
-RUN printf '%s\n' \
-    'import os' \
-    '' \
-    '_exit_code = 1' \
-    '' \
-    'def pytest_sessionfinish(session, exitstatus):' \
-    '    global _exit_code' \
-    '    _exit_code = int(exitstatus)' \
-    '' \
-    'def pytest_unconfigure(config):' \
-    '    import sys' \
-    '    sys.stdout.flush()' \
-    '    sys.stderr.flush()' \
-    '    os._exit(_exit_code)' \
-    > /vllm-workspace/conftest.py
+RUN cat << 'EOF' > /vllm-workspace/conftest.py
+import os
+
+_exit_code = 1
+
+def pytest_sessionfinish(session, exitstatus):
+    global _exit_code
+    _exit_code = int(exitstatus)
+
+def pytest_unconfigure(config):
+    sys.stdout.flush()
+    sys.stderr.flush()
+    os._exit(_exit_code)
+EOF

 # -----------------------
 # Final vLLM image
@@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.9.0"
 ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git"
 ARG FA_BRANCH="0e60e394"
 ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git"
-ARG AITER_BRANCH="v0.1.12"
+ARG AITER_BRANCH="v0.1.10.post2"
 ARG AITER_REPO="https://github.com/ROCm/aiter.git"
 ARG MORI_BRANCH="2d02c6a9"
 ARG MORI_REPO="https://github.com/ROCm/mori.git"
@@ -112,14 +112,10 @@ FROM base AS build_triton
 ARG TRITON_BRANCH
 ARG TRITON_REPO
 RUN git clone ${TRITON_REPO}
-# Cherry picking the following
-# https://github.com/triton-lang/triton/pull/8991
-# https://github.com/triton-lang/triton/pull/9541
 RUN cd triton \
    && git checkout ${TRITON_BRANCH} \
    && git config --global user.email "you@example.com" && git config --global user.name "Your Name" \
    && git cherry-pick 555d04f \
-    && git cherry-pick dd998b6 \
    && if [ ! -f setup.py ]; then cd python; fi \
    && python3 setup.py bdist_wheel --dist-dir=dist \
    && mkdir -p /app/install && cp dist/*.whl /app/install
@@ -93,13 +93,13 @@ RUN curl https://sh.rustup.rs -sSf | sh -s -- -y && \

 FROM python-install AS torch-vision
 # Install torchvision
-ARG TORCH_VISION_VERSION=v0.26.0
+ARG TORCH_VISION_VERSION=v0.25.0
 WORKDIR /tmp
 RUN --mount=type=cache,target=/root/.cache/uv \
    git clone https://github.com/pytorch/vision.git && \
    cd vision && \
    git checkout $TORCH_VISION_VERSION && \
-    uv pip install torch==2.11.0 --index-url https://download.pytorch.org/whl/cpu && \
+    uv pip install torch==2.10.0 --index-url https://download.pytorch.org/whl/cpu && \
    python setup.py bdist_wheel

 FROM python-install AS hf-xet-builder
@@ -253,7 +253,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     NUMBA_WHL_FILE=$(ls /tmp/numba-wheels/*.whl) && \
     OPENCV_WHL_FILE=$(ls /tmp/opencv-wheels/*.whl) && \
     OUTLINES_CORE_WHL_FILE=$(ls /tmp/outlines-core/dist/*.whl) && \
-     uv pip install -v \
+     uv pip install -v \    
        $ARROW_WHL_FILE  \
        $VISION_WHL_FILE \
        $HF_XET_WHL_FILE \
@@ -2,7 +2,7 @@
  "_comment": "Auto-generated from Dockerfile ARGs. Do not edit manually. Run: python tools/generate_versions_json.py",
  "variable": {
    "CUDA_VERSION": {
-      "default": "13.0.0"
+      "default": "12.9.1"
    },
    "PYTHON_VERSION": {
      "default": "3.12"
@@ -11,10 +11,10 @@
      "default": "22.04"
    },
    "BUILD_BASE_IMAGE": {
-      "default": "nvidia/cuda:13.0.0-devel-ubuntu22.04"
+      "default": "nvidia/cuda:12.9.1-devel-ubuntu20.04"
    },
    "FINAL_BASE_IMAGE": {
-      "default": "nvidia/cuda:13.0.0-base-ubuntu22.04"
+      "default": "nvidia/cuda:12.9.1-base-ubuntu22.04"
    },
    "GET_PIP_URL": {
      "default": "https://bootstrap.pypa.io/get-pip.py"
@@ -140,80 +140,6 @@ Data parallelism replicates the entire model across multiple GPU sets and proces
 Data parallelism can be combined with the other parallelism strategies and is set by `data_parallel_size=N`.
 Note that MoE layers will be sharded according to the product of the tensor parallel size and data parallel size.

-### NUMA Binding for Multi-Socket GPU Nodes
-
-On multi-socket GPU servers, GPU worker processes can lose performance if their
-CPU execution and memory allocation drift away from the NUMA node nearest to the
-GPU. vLLM can pin each worker with `numactl` before the Python subprocess starts,
-so the interpreter, imports, and early allocator state are created with the
-desired NUMA policy from the beginning.
-
-Use `--numa-bind` to enable the feature. By default, vLLM auto-detects the
-GPU-to-NUMA mapping and uses `--cpunodebind=<node> --membind=<node>` for each
-worker. When you need a custom CPU policy, add `--numa-bind-cpus` and vLLM will
-switch to `--physcpubind=<cpu-list> --membind=<node>`.
-
-These `--numa-bind*` options only apply to GPU execution processes. They do not
-configure the CPU backend's separate thread-affinity controls. Automatic
-GPU-to-NUMA detection is currently implemented for CUDA/NVML-based platforms;
-other GPU backends must provide explicit binding lists if they use these
-options.
-
-`--numa-bind-nodes` takes one non-negative NUMA node index per visible GPU, in
-the same order as the GPU indices.
-`--numa-bind-cpus` takes one `numactl` CPU list per visible GPU, in the same
-order as the GPU indices. Each CPU list must use
-`numactl --physcpubind` syntax such as `0-3`, `0,2,4-7`, or `16-31,48-63`.
-
-```bash
-# Auto-detect NUMA nodes for visible GPUs
-vllm serve meta-llama/Llama-3.1-8B-Instruct \
-  --tensor-parallel-size 4 \
-  --numa-bind
-
-# Explicit NUMA-node mapping
-vllm serve meta-llama/Llama-3.1-8B-Instruct \
-  --tensor-parallel-size 4 \
-  --numa-bind \
-  --numa-bind-nodes 0 0 1 1
-
-# Explicit CPU pinning, useful for PCT or other high-frequency core layouts
-vllm serve meta-llama/Llama-3.1-8B-Instruct \
-  --tensor-parallel-size 4 \
-  --numa-bind \
-  --numa-bind-nodes 0 0 1 1 \
-  --numa-bind-cpus 0-3 4-7 48-51 52-55
-```
-
-Notes:
-
- CLI usage forces multiprocessing to use the `spawn` method automatically. If you enable NUMA binding through the Python API, also set `VLLM_WORKER_MULTIPROC_METHOD=spawn`.
- Automatic detection relies on NVML and NUMA support from the host. If it cannot determine the mapping reliably, pass `--numa-bind-nodes` explicitly.
- Explicit `--numa-bind-nodes` and `--numa-bind-cpus` values must be valid `numactl` inputs. vLLM does a small amount of validation, but the effective binding semantics are still determined by `numactl`.
- The current implementation binds GPU execution processes such as `EngineCore` and multiprocessing workers. It does not apply NUMA binding to frontend API server processes or the DP coordinator.
- In containerized environments, NUMA policy syscalls may require extra permissions, such as `--cap-add SYS_NICE` when running via `docker run`.
-
-### CPU Backend Thread Affinity
-
-The CPU backend uses a different mechanism from `--numa-bind`. CPU execution is
-configured through CPU-specific environment variables such as
-`VLLM_CPU_OMP_THREADS_BIND`, `VLLM_CPU_NUM_OF_RESERVED_CPU`, and
-`CPU_VISIBLE_MEMORY_NODES`, rather than the GPU-oriented `--numa-bind*` CLI
-options.
-
-By default, `VLLM_CPU_OMP_THREADS_BIND=auto` derives OpenMP placement from the
-available CPU and NUMA topology for each CPU worker. To override the automatic
-policy, set `VLLM_CPU_OMP_THREADS_BIND` explicitly using the CPU list format
-documented for the CPU backend, or use `nobind` to disable this behavior.
-
-For the current CPU backend setup and tuning guidance, see:
-
- [Related runtime environment variables](../getting_started/installation/cpu.md#related-runtime-environment-variables)
- [How to decide `VLLM_CPU_OMP_THREADS_BIND`](../getting_started/installation/cpu.md#how-to-decide-vllm_cpu_omp_threads_bind)
-
-The GPU-only `--numa-bind`, `--numa-bind-nodes`, and `--numa-bind-cpus` options
-do not configure CPU worker affinity.
-
 ### Batch-level DP for Multi-Modal Encoders

 By default, TP is used to shard the weights of multi-modal encoders just like for language decoders,
@@ -57,8 +57,8 @@ Modular kernels are supported by the following `FusedMoEMethodBase` classes.

 - [`ModelOptFp8MoEMethod`][vllm.model_executor.layers.quantization.modelopt.ModelOptFp8MoEMethod]
 - [`Fp8MoEMethod`][vllm.model_executor.layers.quantization.fp8.Fp8MoEMethod]
- [`CompressedTensorsW4A4Nvfp4MoEMethod`][vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe.compressed_tensors_moe_w4a4_nvfp4.CompressedTensorsW4A4Nvfp4MoEMethod]
- [`CompressedTensorsW8A8Fp8MoEMethod`][vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe.compressed_tensors_moe_w8a8_fp8.CompressedTensorsW8A8Fp8MoEMethod]
+- [`CompressedTensorsW4A4Nvfp4MoEMethod`][vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe.CompressedTensorsW4A4Nvfp4MoEMethod]
+- [`CompressedTensorsW8A8Fp8MoEMethod`][vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe.CompressedTensorsW8A8Fp8MoEMethod]
 - [`Mxfp4MoEMethod`][vllm.model_executor.layers.quantization.mxfp4.Mxfp4MoEMethod]
 - [`UnquantizedFusedMoEMethod`][vllm.model_executor.layers.fused_moe.layer.UnquantizedFusedMoEMethod]

@@ -1,167 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-MkDocs hook to automatically convert inline code references to API doc links.
-
-For example, `WeightTransferConfig` becomes
-[`WeightTransferConfig`][vllm.config.WeightTransferConfig]
-
-This works with the `autorefs` plugin to create clickable cross-references
-to API documentation pages generated by `mkdocstrings`.
-
-The hook builds an index of all documented public Python names (classes and
-functions with docstrings) from the vllm package at startup using AST parsing,
-then substitutes matching inline code spans on each page. Names without
-docstrings are excluded because mkdocstrings will not generate a page for them.
-"""
-
-import ast
-import logging
-from pathlib import Path
-
-import regex as re
-from mkdocs.config.defaults import MkDocsConfig
-from mkdocs.structure.files import Files
-from mkdocs.structure.pages import Page
-
-logger = logging.getLogger("mkdocs")
-
-ROOT_DIR = Path(__file__).parent.parent.parent.parent.resolve()
-VLLM_DIR = ROOT_DIR / "vllm"
-
-# Maps short name -> qualified name (e.g. "ModelConfig" -> "vllm.config.ModelConfig")
-_name_index: dict[str, str] = {}
-
-# Fenced code block pattern (``` or ~~~, with optional language specifier).
-_FENCED_BLOCK = re.compile(
-    r"(?:^|\n)(?P<fence>`{3,}|~{3,})[^\n]*\n.*?(?:\n(?P=fence))", re.DOTALL
-)
-
-# Inline code that is NOT already part of a markdown link.
-# Matches `Name` but not [`Name`] and not [`Name`][...] or [`Name`](...).
-_INLINE_CODE = re.compile(
-    r"(?<!\[)"  # not preceded by [
-    r"`(?P<name>[A-Za-z0-9_]*)`"  # `UpperCamelCase` or `UPPER_SNAKE`
-    r"(?!\])"  # not followed by ]
-)
-
-
-def _has_docstring(node: ast.AST) -> bool:
-    """Check if a class or function node has a docstring."""
-    if not isinstance(node, ast.ClassDef | ast.FunctionDef | ast.AsyncFunctionDef):
-        return False
-    return ast.get_docstring(node, clean=False) is not None
-
-
-def _module_path(filepath: Path) -> str:
-    """Convert a filesystem path to a dotted module path."""
-    rel = filepath.relative_to(ROOT_DIR)
-    parts = list(rel.with_suffix("").parts)
-    if parts[-1] == "__init__":
-        parts = parts[:-1]
-    return ".".join(parts)
-
-
-def _index_file(filepath: Path) -> dict[str, str]:
-    """Extract documented public names from a Python file using AST parsing.
-
-    Only classes and functions with docstrings are included, since
-    mkdocstrings won't generate a page for undocumented symbols.
-    """
-    names: dict[str, str] = {}
-    try:
-        source = filepath.read_text(encoding="utf-8")
-        tree = ast.parse(source, filename=str(filepath))
-    except (SyntaxError, UnicodeDecodeError):
-        return names
-
-    module = _module_path(filepath)
-
-    for node in ast.iter_child_nodes(tree):
-        if (
-            # Class definitions (with docstring)
-            isinstance(node, ast.ClassDef)
-            and not node.name.startswith("_")
-            and _has_docstring(node)
-        ) or (
-            # Function definitions (with docstring, only uppercase/CamelCase)
-            isinstance(node, ast.FunctionDef | ast.AsyncFunctionDef)
-            and not node.name.startswith("_")
-            and node.name[0].isupper()
-            and _has_docstring(node)
-        ):
-            names[node.name] = f"{module}.{node.name}"
-
-    return names
-
-
-def _build_index() -> dict[str, str]:
-    """Walk the vllm package and build a name -> qualified path index."""
-    index: dict[str, str] = {}
-    # Track conflicts: if multiple modules define the same name,
-    # prefer shallower modules (more likely to be the public API).
-    depth: dict[str, int] = {}
-
-    for filepath in sorted(VLLM_DIR.rglob("*.py")):
-        # Skip internal/private modules
-        if any(part.startswith("_") and part != "__init__" for part in filepath.parts):
-            continue
-        # Skip third-party vendored code
-        rel = filepath.relative_to(VLLM_DIR)
-        if rel.parts and rel.parts[0] in ("third_party", "vllm_flash_attn"):
-            continue
-
-        module_depth = len(filepath.relative_to(ROOT_DIR).parts)
-        file_names = _index_file(filepath)
-
-        for name, qualified in file_names.items():
-            if name not in index or module_depth < depth[name]:
-                index[name] = qualified
-                depth[name] = module_depth
-
-    return index
-
-
-def on_startup(*, command: str, dirty: bool) -> None:
-    """Build the name index once at startup."""
-    global _name_index
-    _name_index = _build_index()
-    logger.info("autoref_code: indexed %d names from vllm/", len(_name_index))
-
-
-def on_page_markdown(
-    markdown: str, *, page: Page, config: MkDocsConfig, files: Files
-) -> str:
-    """Replace inline code references with autoref links."""
-    if not _name_index:
-        return markdown
-
-    # Skip API reference pages to avoid circular/redundant links.
-    if page.file.src_path.startswith("api/"):
-        return markdown
-
-    # Step 1: Mask fenced code blocks so we don't touch code inside them.
-    masks: list[str] = []
-
-    def _mask_block(match: re.Match) -> str:
-        masks.append(match.group(0))
-        return f"\ue000CODEBLOCK{len(masks) - 1}\ue000"
-
-    masked = _FENCED_BLOCK.sub(_mask_block, markdown)
-
-    # Step 2: Replace inline code references.
-    def _replace(match: re.Match) -> str:
-        name = match.group("name")
-        qualified = _name_index.get(name)
-        if qualified is None:
-            return match.group(0)
-        logger.debug("autoref_code: linking `%s` to [%s]", name, qualified)
-        return f"[`{name}`][{qualified}]"
-
-    result = _INLINE_CODE.sub(_replace, masked)
-
-    # Step 3: Restore masked code blocks.
-    result = re.sub(
-        r"\ue000CODEBLOCK(\d+)\ue000", lambda m: masks[int(m.group(1))], result
-    )
-    return result
@@ -59,7 +59,7 @@ class PydanticMagicMock(MagicMock):
    """`MagicMock` that's able to generate pydantic-core schemas."""

    def __init__(self, *args, **kwargs):
-        name = kwargs.get("name")
+        name = kwargs.pop("name", None)
        super().__init__(*args, **kwargs)
        self.__spec__ = ModuleSpec(name, None)

@@ -85,8 +85,7 @@ def auto_mock(module_name: str, attr: str, max_mocks: int = 100):
            logger.info("Mocking %s for argparse doc generation", e.name)
            sys.modules[e.name] = PydanticMagicMock(name=e.name)
        except Exception:
-            logger.exception("Failed to import %s.%s", module_name, attr)
-            raise
+            logger.exception("Failed to import %s.%s: %s", module_name, attr)

    raise ImportError(
        f"Failed to import {module_name}.{attr} after mocking {max_mocks} imports"
@@ -457,7 +457,6 @@ th {
 | `PanguEmbeddedForCausalLM` | openPangu-Embedded-7B | `FreedomIntelligence/openPangu-Embedded-7B-V1.1` | ✅︎ | ✅︎ |
 | `PanguProMoEV2ForCausalLM` | openpangu-pro-moe-v2 | | ✅︎ | ✅︎ |
 | `PanguUltraMoEForCausalLM` | openpangu-ultra-moe-718b-model | `FreedomIntelligence/openPangu-Ultra-MoE-718B-V1.1` | ✅︎ | ✅︎ |
-| `Param2MoEForCausalLM` | param2moe | `bharatgenai/Param2-17B-A2.4B-Thinking`, etc. | ✅︎ | ✅︎ |
 | `PhiForCausalLM` | Phi | `microsoft/phi-1_5`, `microsoft/phi-2`, etc. | ✅︎ | ✅︎ |
 | `Phi3ForCausalLM` | Phi-4, Phi-3 | `microsoft/Phi-4-mini-instruct`, `microsoft/Phi-4`, `microsoft/Phi-3-mini-4k-instruct`, `microsoft/Phi-3-mini-128k-instruct`, `microsoft/Phi-3-medium-128k-instruct`, etc. | ✅︎ | ✅︎ |
 | `PhiMoEForCausalLM` | Phi-3.5-MoE | `microsoft/Phi-3.5-MoE-instruct`, etc. | ✅︎ | ✅︎ |
@@ -601,7 +600,6 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `PaliGemmaForConditionalGeneration` | PaliGemma, PaliGemma 2 | T + I<sup>E</sup> | `google/paligemma-3b-pt-224`, `google/paligemma-3b-mix-224`, `google/paligemma2-3b-ft-docci-448`, etc. | ✅︎ | ✅︎ |
 | `Phi3VForCausalLM` | Phi-3-Vision, Phi-3.5-Vision | T + I<sup>E+</sup> | `microsoft/Phi-3-vision-128k-instruct`, `microsoft/Phi-3.5-vision-instruct`, etc. | | ✅︎ |
 | `Phi4MMForCausalLM` | Phi-4-multimodal | T + I<sup>+</sup> / T + A<sup>+</sup> / I<sup>+</sup> + A<sup>+</sup> | `microsoft/Phi-4-multimodal-instruct`, etc. | ✅︎ | ✅︎ |
-| `Phi4ForCausalLMV` | Phi-4-reasoning-vision | T + I<sup>+</sup> | `microsoft/Phi-4-reasoning-vision-15B`, etc. | | ✅︎ |
 | `PixtralForConditionalGeneration` | Ministral 3 (Mistral format), Mistral 3 (Mistral format), Mistral Large 3 (Mistral format), Pixtral (Mistral format) | T + I<sup>+</sup> | `mistralai/Ministral-3-3B-Instruct-2512`, `mistralai/Mistral-Small-3.1-24B-Instruct-2503`, `mistralai/Mistral-Large-3-675B-Instruct-2512` `mistralai/Pixtral-12B-2409` etc. | ✅︎ | ✅︎ |
 | `QwenVLForConditionalGeneration`<sup>^</sup> | Qwen-VL | T + I<sup>E+</sup> | `Qwen/Qwen-VL`, `Qwen/Qwen-VL-Chat`, etc. | ✅︎ | ✅︎ |
 | `Qwen2AudioForConditionalGeneration` | Qwen2-Audio | T + A<sup>+</sup> | `Qwen/Qwen2-Audio-7B-Instruct` | | ✅︎ |
@@ -1741,27 +1741,6 @@ def run_phi4mm(questions: list[str], modality: str) -> ModelRequestData:
    )


-# Phi-4-reasoning-vision
-def run_phi4siglip(questions: list[str], modality: str) -> ModelRequestData:
-    assert modality == "image"
-    model_name = "microsoft/Phi-4-reasoning-vision-15B"
-    prompts = [
-        f"<|user|>\n<image>\n{question}<|end|>\n<|assistant|>\n"
-        for question in questions
-    ]
-    engine_args = EngineArgs(
-        model=model_name,
-        trust_remote_code=True,
-        max_model_len=8192,
-        max_num_seqs=2,
-        limit_mm_per_prompt={modality: 1},
-    )
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompts=prompts,
-    )
-
-
 # Pixtral HF-format
 def run_pixtral_hf(questions: list[str], modality: str) -> ModelRequestData:
    assert modality == "image"
@@ -2243,7 +2222,6 @@ model_example_map = {
    "paligemma2": run_paligemma2,
    "phi3_v": run_phi3v,
    "phi4_mm": run_phi4mm,
-    "phi4_siglip": run_phi4siglip,
    "pixtral_hf": run_pixtral_hf,
    "qwen_vl": run_qwen_vl,
    "qwen2_vl": run_qwen2_vl,
@@ -957,24 +957,6 @@ def load_phi4mm(question: str, image_urls: list[str]) -> ModelRequestData:
    )


-def load_phi4siglip(question: str, image_urls: list[str]) -> ModelRequestData:
-    model_name = "microsoft/Phi-4-reasoning-vision-15B"
-    placeholders = "\n".join("<image>" for _ in image_urls)
-    prompt = f"<|user|>\n{placeholders}\n{question}<|end|>\n<|assistant|>\n"
-    engine_args = EngineArgs(
-        model=model_name,
-        trust_remote_code=True,
-        max_model_len=8192,
-        max_num_seqs=2,
-        limit_mm_per_prompt={"image": len(image_urls)},
-    )
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompt=prompt,
-        image_data=[fetch_image(url) for url in image_urls],
-    )
-
-
 def load_qwen_vl_chat(question: str, image_urls: list[str]) -> ModelRequestData:
    model_name = "Qwen/Qwen-VL-Chat"
    engine_args = EngineArgs(
@@ -1473,7 +1455,6 @@ model_example_map = {
    "paddleocr_vl": load_paddleocr_vl,
    "phi3_v": load_phi3v,
    "phi4_mm": load_phi4mm,
-    "phi4_siglip": load_phi4siglip,
    "pixtral_hf": load_pixtral_hf,
    "qwen_vl_chat": load_qwen_vl_chat,
    "qwen2_vl": load_qwen2_vl,
@@ -54,7 +54,6 @@ hooks:
  - docs/mkdocs/hooks/generate_argparse.py
  - docs/mkdocs/hooks/generate_metrics.py
  - docs/mkdocs/hooks/url_schemes.py
-  - docs/mkdocs/hooks/autoref_code.py

 plugins:
  - meta
@@ -6,7 +6,7 @@ requires = [
    "packaging>=24.2",
    "setuptools>=77.0.3,<81.0.0",
    "setuptools-scm>=8.0",
-    "torch == 2.11.0",
+    "torch == 2.10.0",
    "wheel",
    "jinja2",
 ]
@@ -4,7 +4,7 @@ ninja
 packaging>=24.2
 setuptools>=77.0.3,<81.0.0
 setuptools-scm>=8
-torch==2.11.0
+torch==2.10.0
 wheel
 jinja2>=3.1.6
 regex
@@ -7,7 +7,7 @@ requests >= 2.26.0
 tqdm
 blake3
 py-cpuinfo
-transformers >= 5.5.0
+transformers >= 4.56.0, < 5
 tokenizers >= 0.21.1  # Required for fast incremental detokenization.
 protobuf >= 5.29.6, !=6.30.*, !=6.31.*, !=6.32.*, !=6.33.0.*, !=6.33.1.*, !=6.33.2.*, !=6.33.3.*, !=6.33.4.* # Required by LlamaTokenizer, gRPC. CVE-2026-0994
 fastapi[standard] >= 0.115.0 # Required by FastAPI's form models in the OpenAI API server's audio transcriptions endpoint.
@@ -31,13 +31,13 @@ partial-json-parser # used for parsing partial JSON outputs
 pyzmq >= 25.0.0
 msgspec
 gguf >= 0.17.0
-mistral_common[image] >= 1.11.0
+mistral_common[image] >= 1.10.0
 opencv-python-headless >= 4.13.0    # required for video IO
 pyyaml
 six>=1.16.0; python_version > '3.11' # transitive dependency of pandas that needs to be the latest version for python 3.12
 setuptools>=77.0.3,<81.0.0; python_version > '3.11' # Setuptools is used by triton, we need to ensure a modern version is installed for 3.12+ so that it does not try to import distutils, which was removed in 3.12
 einops # Required for Qwen2-VL.
-compressed-tensors >= 0.15.0 # required for compressed-tensors
+compressed-tensors == 0.14.0.1 # required for compressed-tensors
 depyf==0.20.0 # required for profiling and debugging with compilation config
 cloudpickle # allows pickling lambda functions in model_executor/models/registry.py
 watchfiles # required for http server to monitor the updates of TLS files
@@ -1,11 +1,10 @@
--extra-index-url https://download.pytorch.org/whl/cpu
 cmake>=3.26.1
 ninja
 packaging>=24.2
 setuptools==77.0.3 # this version can reuse CMake build dir
 setuptools-scm>=8
-torch==2.11.0+cpu; platform_machine == "x86_64" or platform_machine == "s390x" or platform_machine == "aarch64"
-torch==2.11.0; platform_system == "Darwin" or platform_machine == "ppc64le"  or platform_machine == "riscv64"
+torch==2.10.0+cpu; platform_machine == "x86_64" or platform_machine == "s390x"
+torch==2.10.0; platform_machine == "aarch64" or platform_system == "Darwin" or platform_machine == "ppc64le"
 wheel
 jinja2>=3.1.6
 regex
@@ -1,4 +1,3 @@
--extra-index-url https://download.pytorch.org/whl/cpu
 # Common dependencies
 -r common.txt

@@ -7,8 +6,8 @@ setuptools==77.0.3 # this version can reuse CMake build dir
 numba == 0.61.2; platform_machine != "s390x" # Required for N-gram speculative decoding

 # Dependencies for CPUs
-torch==2.11.0+cpu; platform_machine == "x86_64" or platform_machine == "s390x" or platform_machine == "aarch64"
-torch==2.11.0; platform_system == "Darwin" or platform_machine == "ppc64le" or platform_machine == "riscv64"
+torch==2.10.0+cpu; platform_machine == "x86_64" or platform_machine == "s390x"
+torch==2.10.0; platform_machine == "aarch64" or platform_system == "Darwin" or platform_machine == "ppc64le" or platform_machine == "riscv64"

 # required for the image processor of minicpm-o-2_6, this must be updated alongside torch
 torchaudio; platform_machine != "s390x" and platform_machine != "riscv64"
@@ -4,10 +4,10 @@
 numba == 0.61.2 # Required for N-gram speculative decoding

 # Dependencies for NVIDIA GPUs
-torch==2.11.0
-torchaudio==2.11.0
+torch==2.10.0
+torchaudio==2.10.0
 # These must be updated alongside torch
-torchvision==0.26.0 # Required for phi3v processor. See https://github.com/pytorch/vision?tab=readme-ov-file#installation for corresponding version
+torchvision==0.25.0 # Required for phi3v processor. See https://github.com/pytorch/vision?tab=readme-ov-file#installation for corresponding version
 # FlashInfer should be updated together with the Dockerfile
 flashinfer-python==0.6.7
 flashinfer-cubin==0.6.7
@@ -1,3 +1,3 @@
 lmcache >= 0.3.9
-nixl[cu13] >= 0.7.1, < 0.10.0 # Required for disaggregated prefill
+nixl >= 0.7.1, < 0.10.0 # Required for disaggregated prefill
 mooncake-transfer-engine >= 0.3.8
@@ -23,14 +23,14 @@ jiwer # required for audio tests
 timm # required for internvl test
 transformers_stream_generator # required for qwen-vl test
 matplotlib # required for qwen-vl test
-mistral_common[image,audio] >= 1.11.0 # required for voxtral test
+mistral_common[image,audio] >= 1.9.1 # required for voxtral test
 num2words # required for smolvlm test
 opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.11 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.5.0
-tokenizers==0.22.2
+transformers==4.57.5
+tokenizers==0.22.0
 schemathesis>=3.39.15 # Required for openai schema test.
 # quantization
 bitsandbytes>=0.49.2
@@ -1,11 +1,10 @@
 # Common dependencies
 -r common.txt

-
 --extra-index-url https://download.pytorch.org/whl/rocm7.1
-torch==2.11.0
-torchvision==0.26.0
-torchaudio==2.11.0
+torch==2.10.0
+torchvision==0.25.0
+torchaudio==2.10.0
 triton==3.6.0
 cmake>=3.26.1,<4
 packaging>=24.2
@@ -1,5 +1,3 @@
-r common.txt
-
 # testing
 pytest
 tensorizer==2.10.1
@@ -31,15 +29,15 @@ tblib # for pickling test exceptions
 timm>=1.0.17 # required for internvl and gemma3n-mm test
 transformers_stream_generator # required for qwen-vl test
 matplotlib # required for qwen-vl test
-mistral_common[image,audio]>=1.11.0 # required for voxtral test
+mistral_common[image,audio]>=1.10.0 # required for voxtral test
 num2words # required for smolvlm test
 open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_common.py
 opencv-python-headless>=4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.11 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.5.0
-tokenizers==0.22.2
+transformers==4.57.5
+tokenizers==0.22.0
 schemathesis>=3.39.15 # Required for openai schema test
 # quantization
 bitsandbytes==0.49.2
@@ -82,4 +80,4 @@ plotly # required for perf comparison html report
 rapidfuzz
 torchgeo==0.7.0
 multiprocess==0.70.16
-huggingface-hub==1.9.2
+huggingface-hub==0.36.2
@@ -15,7 +15,6 @@ aiohappyeyeballs==2.6.1
 aiohttp==3.13.3
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   aiohttp-cors
    #   fsspec
    #   gpt-oss
@@ -39,31 +38,20 @@ annotated-doc==0.0.4
    #   typer
 annotated-types==0.7.0
    # via pydantic
-anthropic==0.89.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 antlr4-python3-runtime==4.9.3
    # via
    #   hydra-core
    #   omegaconf
-anyio==4.13.0
+anyio==4.6.2.post1
    # via
-    #   anthropic
    #   httpx
-    #   mcp
-    #   openai
-    #   sse-starlette
    #   starlette
-    #   watchfiles
 arctic-inference==0.1.1
    # via -r requirements/rocm-test.in
 argcomplete==3.6.3
    # via datamodel-code-generator
 arrow==1.4.0
    # via isoduration
-astor==0.8.1
-    # via depyf
 attrs==26.1.0
    # via
    #   aiohttp
@@ -95,8 +83,6 @@ bitsandbytes==0.49.2
    #   lightning
 black==26.3.1
    # via datamodel-code-generator
-blake3==1.0.8
-    # via -r requirements/common.txt
 blobfile==3.0.0
    # via -r requirements/rocm-test.in
 bm25s==0.2.13
@@ -113,10 +99,6 @@ bounded-pool-executor==0.0.3
    # via pqdm
 buildkite-test-collector==0.1.9
    # via -r requirements/rocm-test.in
-cachetools==7.0.5
-    # via -r requirements/common.txt
-cbor2==5.9.0
-    # via -r requirements/common.txt
 certifi==2026.2.25
    # via
    #   fiona
@@ -150,7 +132,6 @@ click==8.3.1
    #   nltk
    #   rasterio
    #   ray
-    #   rich-toolkit
    #   schemathesis
    #   typer
    #   uvicorn
@@ -161,8 +142,6 @@ cligj==0.7.2
    # via
    #   fiona
    #   rasterio
-cloudpickle==3.1.2
-    # via -r requirements/common.txt
 colorama==0.4.6
    # via
    #   perceptron
@@ -172,10 +151,6 @@ colorful==0.5.8
    # via ray
 colorlog==6.10.1
    # via optuna
-compressed-tensors==0.15.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 contourpy==1.3.3
    # via matplotlib
 coverage==7.13.5
@@ -207,42 +182,24 @@ decorator==5.2.1
    # via librosa
 decord==0.6.0
    # via -r requirements/rocm-test.in
-depyf==0.20.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 diffusers==0.37.0
    # via terratorch
 dill==0.3.8
    # via
    #   datasets
-    #   depyf
    #   evaluate
    #   lm-eval
    #   multiprocess
-diskcache==5.6.3
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 distlib==0.4.0
    # via virtualenv
-distro==1.9.0
-    # via
-    #   anthropic
-    #   openai
-dnspython==2.8.0
-    # via email-validator
 docker==7.1.0
    # via gpt-oss
 docopt==0.6.2
    # via num2words
 docstring-parser==0.17.0
-    # via
-    #   anthropic
-    #   jsonargparse
+    # via jsonargparse
 einops==0.8.2
    # via
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   encodec
    #   terratorch
@@ -251,10 +208,6 @@ einops==0.8.2
    #   vocos
 einx==0.4.2
    # via vector-quantize-pytorch
-email-validator==2.3.0
-    # via
-    #   fastapi
-    #   pydantic
 encodec==0.1.1
    # via vocos
 et-xmlfile==2.0.0
@@ -264,15 +217,7 @@ evaluate==0.4.6
 fastapi==0.135.2
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   gpt-oss
-    #   model-hosting-container-standards
-fastapi-cli==0.0.24
-    # via fastapi
-fastapi-cloud-cli==0.15.1
-    # via fastapi-cli
-fastar==0.9.0
-    # via fastapi-cloud-cli
 fastparquet==2026.3.0
    # via genai-perf
 fastsafetensors==0.2.2
@@ -280,7 +225,6 @@ fastsafetensors==0.2.2
 filelock==3.25.2
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   blobfile
    #   datasets
    #   diffusers
@@ -288,6 +232,7 @@ filelock==3.25.2
    #   python-discovery
    #   ray
    #   torch
+    #   transformers
    #   virtualenv
 fiona==1.10.1
    # via torchgeo
@@ -319,10 +264,6 @@ genson==1.3.0
    # via datamodel-code-generator
 geopandas==1.1.3
    # via terratorch
-gguf==0.18.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 gitdb==4.0.12
    # via gitpython
 gitpython==3.1.46
@@ -349,10 +290,7 @@ google-crc32c==1.8.0
 google-resumable-media==2.8.0
    # via google-cloud-storage
 googleapis-common-protos==1.73.0
-    # via
-    #   google-api-core
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
+    # via google-api-core
 gpt-oss==0.0.8
    # via -r requirements/rocm-test.in
 graphql-core==3.2.8
@@ -364,7 +302,6 @@ grpcio==1.78.0
    #   -c requirements/rocm.txt
    #   -r requirements/rocm-test.in
    #   grpcio-reflection
-    #   opentelemetry-exporter-otlp-proto-grpc
    #   ray
    #   tensorboard
 grpcio-reflection==1.78.0
@@ -381,7 +318,7 @@ h5py==3.16.0
    # via terratorch
 harfile==0.4.0
    # via schemathesis
-hf-xet==1.4.3
+hf-xet==1.4.2
    # via huggingface-hub
 hiredis==3.3.1
    # via tensorizer
@@ -391,24 +328,13 @@ html2text==2025.4.15
    # via gpt-oss
 httpcore==1.0.9
    # via httpx
-httptools==0.7.1
-    # via uvicorn
 httpx==0.27.2
    # via
    #   -r requirements/rocm-test.in
-    #   anthropic
    #   diffusers
-    #   fastapi
-    #   fastapi-cloud-cli
-    #   huggingface-hub
-    #   mcp
-    #   model-hosting-container-standards
-    #   openai
    #   perceptron
    #   schemathesis
-httpx-sse==0.4.3
-    # via mcp
-huggingface-hub==1.9.2
+huggingface-hub==0.36.2
    # via
    #   -r requirements/rocm-test.in
    #   accelerate
@@ -444,13 +370,10 @@ hypothesis-jsonschema==0.23.1
 idna==3.11
    # via
    #   anyio
-    #   email-validator
    #   httpx
    #   jsonschema
    #   requests
    #   yarl
-ijson==3.5.0
-    # via -r requirements/common.txt
 imagehash==4.3.2
    # via -r requirements/rocm-test.in
 imageio==2.37.3
@@ -467,8 +390,6 @@ iniconfig==2.3.0
    # via pytest
 instanttensor==0.1.6
    # via -r requirements/rocm-test.in
-interegular==0.3.3
-    # via lm-format-enforcer
 isodate==0.7.2
    # via azure-storage-blob
 isoduration==20.11.0
@@ -478,21 +399,15 @@ isort==8.0.1
 jinja2==3.1.6
    # via
    #   datamodel-code-generator
-    #   fastapi
    #   genai-perf
    #   lm-eval
    #   torch
-jiter==0.13.0
-    # via
-    #   anthropic
-    #   openai
 jiwer==4.0.0
    # via -r requirements/rocm-test.in
 jmespath==1.1.0
    # via
    #   boto3
    #   botocore
-    #   model-hosting-container-standards
 joblib==1.5.3
    # via
    #   librosa
@@ -511,7 +426,6 @@ jsonpointer==3.1.0
 jsonschema==4.26.0
    # via
    #   hypothesis-jsonschema
-    #   mcp
    #   mistral-common
    #   ray
    #   schemathesis
@@ -529,10 +443,6 @@ kornia==0.8.2
    # via torchgeo
 kornia-rs==0.1.10
    # via kornia
-lark==1.2.2
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 lazy-loader==0.4
    # via
    #   librosa
@@ -556,24 +466,14 @@ lightning-utilities==0.15.3
    #   lightning
    #   pytorch-lightning
    #   torchmetrics
-llguidance==1.3.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 llvmlite==0.44.0
    # via numba
 lm-eval==0.4.11
    # via -r requirements/rocm-test.in
-lm-format-enforcer==0.11.3
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 logistro==2.0.1
    # via
    #   choreographer
    #   kaleido
-loguru==0.7.3
-    # via compressed-tensors
 lxml==6.0.2
    # via
    #   blobfile
@@ -600,19 +500,12 @@ mbstrdecoder==1.1.4
    #   dataproperty
    #   pytablewriter
    #   typepy
-mcp==1.27.0
-    # via -r requirements/common.txt
 mdurl==0.1.2
    # via markdown-it-py
-mistral-common==1.11.0
+mistral-common==1.10.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
-model-hosting-container-standards==0.1.14
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 more-itertools==10.8.0
    # via
    #   inflect
@@ -629,8 +522,6 @@ msgpack==1.1.2
    # via
    #   librosa
    #   ray
-msgspec==0.20.0
-    # via -r requirements/common.txt
 mteb==2.11.5
    # via -r requirements/rocm-test.in
 multidict==6.7.1
@@ -650,8 +541,6 @@ networkx==3.6.1
    # via
    #   scikit-image
    #   torch
-ninja==1.13.0
-    # via -r requirements/common.txt
 nltk==3.9.3
    # via rouge-score
 num2words==0.5.14
@@ -666,7 +555,6 @@ numkong==7.1.1
    # via albucore
 numpy==2.2.6
    # via
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   accelerate
    #   albucore
@@ -684,7 +572,6 @@ numpy==2.2.6
    #   fastparquet
    #   genai-perf
    #   geopandas
-    #   gguf
    #   h5py
    #   imagehash
    #   imageio
@@ -733,60 +620,15 @@ numpy==2.2.6
    #   tritonclient
    #   vocos
    #   xarray
-    #   xgrammar
-nvidia-cublas-cu12==12.8.4.1
-    # via
-    #   nvidia-cudnn-cu12
-    #   nvidia-cusolver-cu12
-    #   torch
-nvidia-cuda-cupti-cu12==12.8.90
-    # via torch
-nvidia-cuda-nvrtc-cu12==12.8.93
-    # via torch
-nvidia-cuda-runtime-cu12==12.8.90
-    # via torch
-nvidia-cudnn-cu12==9.10.2.21
-    # via torch
-nvidia-cufft-cu12==11.3.3.83
-    # via torch
-nvidia-cufile-cu12==1.13.1.3
-    # via torch
-nvidia-curand-cu12==10.3.9.90
-    # via torch
-nvidia-cusolver-cu12==11.7.3.90
-    # via torch
-nvidia-cusparse-cu12==12.5.8.93
-    # via
-    #   nvidia-cusolver-cu12
-    #   torch
-nvidia-cusparselt-cu12==0.7.1
-    # via torch
-nvidia-nccl-cu12==2.27.5
-    # via torch
-nvidia-nvjitlink-cu12==12.8.93
-    # via
-    #   nvidia-cufft-cu12
-    #   nvidia-cusolver-cu12
-    #   nvidia-cusparse-cu12
-    #   torch
-nvidia-nvshmem-cu12==3.4.5
-    # via torch
-nvidia-nvtx-cu12==12.8.90
-    # via torch
 omegaconf==2.3.0
    # via
    #   hydra-core
    #   lightning
 open-clip-torch==2.32.0
    # via -r requirements/rocm-test.in
-openai==2.30.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 openai-harmony==0.0.8
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   gpt-oss
 opencensus==0.11.4
    # via ray
@@ -795,7 +637,6 @@ opencensus-context==0.1.3
 opencv-python-headless==4.13.0.92
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   albumentations
    #   mistral-common
@@ -804,59 +645,26 @@ openpyxl==3.1.5
 opentelemetry-api==1.40.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
    #   opentelemetry-exporter-prometheus
    #   opentelemetry-sdk
    #   opentelemetry-semantic-conventions
-opentelemetry-exporter-otlp==1.40.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
-opentelemetry-exporter-otlp-proto-common==1.40.0
-    # via
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
-opentelemetry-exporter-otlp-proto-grpc==1.40.0
-    # via opentelemetry-exporter-otlp
-opentelemetry-exporter-otlp-proto-http==1.40.0
-    # via opentelemetry-exporter-otlp
 opentelemetry-exporter-prometheus==0.61b0
    # via ray
 opentelemetry-proto==1.40.0
-    # via
-    #   opentelemetry-exporter-otlp-proto-common
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
-    #   ray
+    # via ray
 opentelemetry-sdk==1.40.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
    #   opentelemetry-exporter-prometheus
-    #   opentelemetry-semantic-conventions-ai
    #   ray
 opentelemetry-semantic-conventions==0.61b0
-    # via
-    #   opentelemetry-sdk
-    #   opentelemetry-semantic-conventions-ai
-opentelemetry-semantic-conventions-ai==0.5.1
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
+    # via opentelemetry-sdk
 optuna==3.6.1
    # via genai-perf
 orjson==3.11.7
    # via
    #   genai-perf
    #   kaleido
-outlines-core==0.2.11
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 packaging==26.0
    # via
    #   -c requirements/rocm.txt
@@ -874,7 +682,6 @@ packaging==26.0
    #   lazy-loader
    #   lightning
    #   lightning-utilities
-    #   lm-format-enforcer
    #   matplotlib
    #   optuna
    #   peft
@@ -906,8 +713,6 @@ pandas==3.0.1
    #   tacoreader
    #   torchgeo
    #   xarray
-partial-json-parser==0.2.1.1.post7
-    # via -r requirements/common.txt
 pathspec==1.0.4
    # via black
 pathvalidate==3.3.1
@@ -922,7 +727,6 @@ perf-analyzer==0.1.0
    # via genai-perf
 pillow==12.1.1
    # via
-    #   -r requirements/common.txt
    #   diffusers
    #   genai-perf
    #   imagehash
@@ -964,14 +768,8 @@ pqdm==0.2.0
 prometheus-client==0.24.1
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   opentelemetry-exporter-prometheus
-    #   prometheus-fastapi-instrumentator
    #   ray
-prometheus-fastapi-instrumentator==7.1.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 propcache==0.4.1
    # via
    #   aiohttp
@@ -981,7 +779,6 @@ proto-plus==1.27.1
 protobuf==6.33.6
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   google-api-core
    #   googleapis-common-protos
    #   grpcio-reflection
@@ -994,14 +791,11 @@ protobuf==6.33.6
    #   wandb
 psutil==7.2.2
    # via
-    #   -r requirements/common.txt
    #   accelerate
    #   peft
    #   tensorizer
 py==1.11.0
    # via pytest-forked
-py-cpuinfo==9.0.0
-    # via -r requirements/common.txt
 py-spy==0.4.1
    # via ray
 pyarrow==23.0.1
@@ -1014,8 +808,6 @@ pyasn1==0.6.3
    # via pyasn1-modules
 pyasn1-modules==0.4.2
    # via google-auth
-pybase64==1.4.3
-    # via -r requirements/common.txt
 pycocotools==2.0.11
    # via terratorch
 pycountry==26.2.16
@@ -1027,44 +819,26 @@ pycryptodomex==3.23.0
 pydantic==2.12.5
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   albumentations
-    #   anthropic
-    #   compressed-tensors
    #   datamodel-code-generator
    #   fastapi
-    #   fastapi-cloud-cli
    #   gpt-oss
    #   lightly
-    #   lm-format-enforcer
-    #   mcp
    #   mistral-common
-    #   model-hosting-container-standards
    #   mteb
-    #   openai
    #   openai-harmony
    #   pydantic-extra-types
-    #   pydantic-settings
    #   ray
    #   wandb
-    #   xgrammar
 pydantic-core==2.41.5
    # via pydantic
 pydantic-extra-types==2.11.1
-    # via
-    #   fastapi
-    #   mistral-common
-pydantic-settings==2.13.1
-    # via
-    #   fastapi
-    #   mcp
+    # via mistral-common
 pygments==2.19.2
    # via rich
 pyjwt==2.12.1
-    # via
-    #   mcp
-    #   msal
+    # via msal
 pyogrio==0.12.1
    # via geopandas
 pyparsing==3.3.2
@@ -1124,16 +898,6 @@ python-dateutil==2.9.0.post0
    #   typepy
 python-discovery==1.2.0
    # via virtualenv
-python-dotenv==1.2.2
-    # via
-    #   pydantic-settings
-    #   uvicorn
-python-json-logger==4.1.0
-    # via -r requirements/common.txt
-python-multipart==0.0.22
-    # via
-    #   fastapi
-    #   mcp
 python-rapidjson==1.23
    # via tritonclient
 pytokens==0.4.1
@@ -1150,17 +914,14 @@ pywavelets==1.9.0
    # via imagehash
 pyyaml==6.0.3
    # via
-    #   -r requirements/common.txt
    #   accelerate
    #   albumentations
    #   datamodel-code-generator
    #   datasets
    #   genai-perf
-    #   gguf
    #   huggingface-hub
    #   jsonargparse
    #   lightning
-    #   lm-format-enforcer
    #   omegaconf
    #   optuna
    #   peft
@@ -1170,13 +931,8 @@ pyyaml==6.0.3
    #   schemathesis
    #   timm
    #   transformers
-    #   uvicorn
    #   vocos
    #   wandb
-pyzmq==27.1.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 rapidfuzz==3.12.1
    # via
    #   -r requirements/rocm-test.in
@@ -1196,7 +952,6 @@ referencing==0.37.0
    #   jsonschema-specifications
 regex==2026.2.28
    # via
-    #   -r requirements/common.txt
    #   diffusers
    #   nltk
    #   open-clip-torch
@@ -1206,23 +961,21 @@ regex==2026.2.28
 requests==2.32.5
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   azure-core
    #   buildkite-test-collector
    #   datasets
    #   diffusers
    #   docker
    #   evaluate
-    #   gguf
    #   google-api-core
    #   google-cloud-storage
    #   gpt-oss
+    #   huggingface-hub
    #   lightly
    #   lm-eval
    #   mistral-common
    #   msal
    #   mteb
-    #   opentelemetry-exporter-otlp-proto-http
    #   pooch
    #   ray
    #   responses
@@ -1230,6 +983,7 @@ requests==2.32.5
    #   starlette-testclient
    #   tacoreader
    #   tiktoken
+    #   transformers
    #   wandb
 resampy==0.4.3
    # via -r requirements/rocm-test.in
@@ -1245,15 +999,8 @@ rich==14.3.3
    #   lightning
    #   mteb
    #   perceptron
-    #   rich-toolkit
    #   terratorch
    #   typer
-rich-toolkit==0.19.7
-    # via
-    #   fastapi-cli
-    #   fastapi-cloud-cli
-rignore==0.7.6
-    # via fastapi-cloud-cli
 rioxarray==0.22.0
    # via terratorch
 rouge-score==0.1.2
@@ -1323,20 +1070,12 @@ sentence-transformers==5.3.0
    # via
    #   -r requirements/rocm-test.in
    #   mteb
-sentencepiece==0.2.1
-    # via -r requirements/common.txt
 sentry-sdk==2.55.0
-    # via
-    #   fastapi-cloud-cli
-    #   wandb
-setproctitle==1.3.7
-    # via -r requirements/common.txt
+    # via wandb
 setuptools==79.0.1
    # via
    #   -c requirements/common.txt
    #   -c requirements/rocm.txt
-    #   -r requirements/common.txt
-    #   model-hosting-container-standards
    #   pytablewriter
    #   tensorboard
    #   torch
@@ -1353,7 +1092,6 @@ simplejson==3.20.2
 six==1.17.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   junit-xml
    #   lightly
    #   opencensus
@@ -1366,9 +1104,8 @@ smmap==5.0.3
    # via gitdb
 sniffio==1.3.1
    # via
-    #   anthropic
+    #   anyio
    #   httpx
-    #   openai
 sortedcontainers==2.4.0
    # via hypothesis
 soundfile==0.13.1
@@ -1387,16 +1124,10 @@ sqlalchemy==2.0.48
    #   optuna
 sqlitedict==2.1.0
    # via lm-eval
-sse-starlette==3.3.4
-    # via mcp
 starlette==0.52.1
    # via
    #   fastapi
-    #   mcp
-    #   model-hosting-container-standards
-    #   prometheus-fastapi-instrumentator
    #   schemathesis
-    #   sse-starlette
    #   starlette-testclient
 starlette-testclient==0.4.1
    # via schemathesis
@@ -1406,8 +1137,6 @@ stringzilla==4.6.0
    # via albucore
 structlog==25.5.0
    # via gpt-oss
-supervisor==4.3.0
-    # via model-hosting-container-standards
 sympy==1.14.0
    # via
    #   einx
@@ -1451,7 +1180,6 @@ tifffile==2026.3.3
 tiktoken==0.12.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   gpt-oss
    #   lm-eval
    #   mistral-common
@@ -1463,10 +1191,9 @@ timm==1.0.17
    #   segmentation-models-pytorch
    #   terratorch
    #   torchgeo
-tokenizers==0.22.2
+tokenizers==0.22.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   transformers
 tomli==2.4.0
@@ -1485,10 +1212,8 @@ torchmetrics==1.9.0
    #   torchgeo
 tqdm==4.67.3
    # via
-    #   -r requirements/common.txt
    #   datasets
    #   evaluate
-    #   gguf
    #   huggingface-hub
    #   lightly
    #   lightning
@@ -1496,7 +1221,6 @@ tqdm==4.67.3
    #   mteb
    #   nltk
    #   open-clip-torch
-    #   openai
    #   optuna
    #   peft
    #   pqdm
@@ -1506,17 +1230,14 @@ tqdm==4.67.3
    #   tacoreader
    #   terratorch
    #   transformers
-transformers==5.5.0
+transformers==4.57.5
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
-    #   compressed-tensors
    #   genai-perf
    #   peft
    #   sentence-transformers
    #   transformers-stream-generator
-    #   xgrammar
 transformers-stream-generator==0.0.5
    # via -r requirements/rocm-test.in
 tritonclient==2.66.0
@@ -1530,23 +1251,16 @@ typepy==1.3.4
    #   tabledata
 typer==0.24.1
    # via
-    #   fastapi-cli
-    #   fastapi-cloud-cli
    #   fastsafetensors
-    #   huggingface-hub
    #   perceptron
-    #   transformers
 typeshed-client==2.9.0
    # via jsonargparse
 typing-extensions==4.15.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   aiosignal
    #   albumentations
    #   alembic
-    #   anthropic
-    #   anyio
    #   azure-core
    #   azure-identity
    #   azure-storage-blob
@@ -1558,13 +1272,9 @@ typing-extensions==4.15.0
    #   lightning
    #   lightning-utilities
    #   lm-eval
-    #   mcp
    #   mistral-common
    #   mteb
-    #   openai
    #   opentelemetry-api
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
    #   opentelemetry-sdk
    #   opentelemetry-semantic-conventions
    #   pqdm
@@ -1573,7 +1283,6 @@ typing-extensions==4.15.0
    #   pydantic-extra-types
    #   pytorch-lightning
    #   referencing
-    #   rich-toolkit
    #   sentence-transformers
    #   sqlalchemy
    #   starlette
@@ -1583,13 +1292,10 @@ typing-extensions==4.15.0
    #   typeshed-client
    #   typing-inspection
    #   wandb
-    #   xgrammar
 typing-inspection==0.4.2
    # via
    #   fastapi
-    #   mcp
    #   pydantic
-    #   pydantic-settings
 tzdata==2025.3
    # via arrow
 uri-template==1.3.0
@@ -1605,14 +1311,7 @@ urllib3==2.6.3
    #   sentry-sdk
    #   tritonclient
 uvicorn==0.42.0
-    # via
-    #   fastapi
-    #   fastapi-cli
-    #   fastapi-cloud-cli
-    #   gpt-oss
-    #   mcp
-uvloop==0.22.1
-    # via uvicorn
+    # via gpt-oss
 vector-quantize-pytorch==1.28.0
    # via -r requirements/rocm-test.in
 virtualenv==21.2.0
@@ -1621,16 +1320,10 @@ vocos==0.1.0
    # via -r requirements/rocm-test.in
 wandb==0.25.1
    # via terratorch
-watchfiles==1.1.1
-    # via
-    #   -r requirements/common.txt
-    #   uvicorn
 wcwidth==0.6.0
    # via ftfy
 webcolors==25.10.0
    # via jsonschema
-websockets==16.0
-    # via uvicorn
 werkzeug==3.1.6
    # via
    #   schemathesis
@@ -1641,10 +1334,6 @@ wrapt==2.1.2
    # via smart-open
 xarray==2026.2.0
    # via rioxarray
-xgrammar==0.1.33
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 xxhash==3.6.0
    # via
    #   datasets
@@ -1665,4 +1354,20 @@ zstandard==0.25.0
 # triton
 # cuda-bindings
 # cuda-pathfinder
+# cuda-toolkit
 # cupy-cuda12x
+# nvidia-cublas
+# nvidia-cuda-cupti
+# nvidia-cuda-nvrtc
+# nvidia-cuda-runtime
+# nvidia-cudnn-cu13
+# nvidia-cufft
+# nvidia-cufile
+# nvidia-curand
+# nvidia-cusolver
+# nvidia-cusparse
+# nvidia-cusparselt-cu13
+# nvidia-nccl-cu13
+# nvidia-nvjitlink
+# nvidia-nvshmem-cu13
+# nvidia-nvtx
@@ -27,20 +27,20 @@ soundfile # required for audio tests
 jiwer # required for audio tests
 tblib # for pickling test exceptions
 timm >=1.0.17 # required for internvl and gemma3n-mm test
-torch==2.11.0
-torchaudio==2.11.0
-torchvision==0.26.0
+torch==2.10.0
+torchaudio==2.10.0
+torchvision==0.25.0
 transformers_stream_generator # required for qwen-vl test
 matplotlib # required for qwen-vl test
-mistral_common[image,audio] >= 1.11.0 # required for voxtral test
+mistral_common[image,audio] >= 1.9.1 # required for voxtral test
 num2words # required for smolvlm test
 open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_common.py
 opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.11 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.5.0
-tokenizers==0.22.2
+transformers==4.57.5
+tokenizers==0.22.0
 schemathesis>=3.39.15 # Required for openai schema test.
 # quantization
 bitsandbytes==0.49.2
@@ -1,5 +1,5 @@
 # This file was autogenerated by uv via the following command:
-#    uv pip compile requirements/test.in -c requirements/common.txt -o requirements/test.txt --index-strategy unsafe-best-match --torch-backend cu130 --python-platform x86_64-manylinux_2_28 --python-version 3.12
+#    uv pip compile requirements/test.in -c requirements/common.txt -o requirements/test.txt --index-strategy unsafe-best-match --torch-backend cu129 --python-platform x86_64-manylinux_2_28 --python-version 3.12
 absl-py==2.1.0
    # via
    #   rouge-score
@@ -165,12 +165,10 @@ cryptography==46.0.5
    #   azure-storage-blob
    #   msal
    #   pyjwt
-cuda-bindings==13.0.3
+cuda-bindings==12.9.4
    # via torch
 cuda-pathfinder==1.3.3
    # via cuda-bindings
-cuda-toolkit==13.0.2
-    # via torch
 cupy-cuda12x==13.6.0
    # via ray
 cycler==0.12.1
@@ -246,6 +244,7 @@ filelock==3.16.1
    #   huggingface-hub
    #   ray
    #   torch
+    #   transformers
    #   virtualenv
 fiona==1.10.1
    # via torchgeo
@@ -328,7 +327,7 @@ h5py==3.13.0
    # via terratorch
 harfile==0.3.0
    # via schemathesis
-hf-xet==1.4.3
+hf-xet==1.1.7
    # via huggingface-hub
 hiredis==3.0.0
    # via tensorizer
@@ -342,10 +341,9 @@ httpx==0.27.2
    # via
    #   -r requirements/test.in
    #   diffusers
-    #   huggingface-hub
    #   perceptron
    #   schemathesis
-huggingface-hub==1.9.2
+huggingface-hub==0.36.2
    # via
    #   accelerate
    #   datasets
@@ -510,7 +508,7 @@ mbstrdecoder==1.1.3
    #   typepy
 mdurl==0.1.2
    # via markdown-it-py
-mistral-common==1.11.0
+mistral-common==1.10.0
    # via
    #   -c requirements/common.txt
    #   -r requirements/test.in
@@ -617,45 +615,45 @@ numpy==2.2.6
    #   tritonclient
    #   vocos
    #   xarray
-nvidia-cublas==13.1.0.3
+nvidia-cublas-cu12==12.9.1.4
    # via
-    #   cuda-toolkit
-    #   nvidia-cudnn-cu13
-    #   nvidia-cusolver
-nvidia-cuda-cupti==13.0.85
-    # via cuda-toolkit
-nvidia-cuda-nvrtc==13.0.88
-    # via cuda-toolkit
-nvidia-cuda-runtime==13.0.96
-    # via cuda-toolkit
-nvidia-cudnn-cu13==9.19.0.56
+    #   nvidia-cudnn-cu12
+    #   nvidia-cusolver-cu12
+    #   torch
+nvidia-cuda-cupti-cu12==12.9.79
    # via torch
-nvidia-cufft==12.0.0.61
-    # via cuda-toolkit
-nvidia-cufile==1.15.1.6
-    # via cuda-toolkit
-nvidia-curand==10.4.0.35
-    # via cuda-toolkit
-nvidia-cusolver==12.0.4.66
-    # via cuda-toolkit
-nvidia-cusparse==12.6.3.3
+nvidia-cuda-nvrtc-cu12==12.9.86
+    # via torch
+nvidia-cuda-runtime-cu12==12.9.79
+    # via torch
+nvidia-cudnn-cu12==9.10.2.21
+    # via torch
+nvidia-cufft-cu12==11.4.1.4
+    # via torch
+nvidia-cufile-cu12==1.14.1.1
+    # via torch
+nvidia-curand-cu12==10.3.10.19
+    # via torch
+nvidia-cusolver-cu12==11.7.5.82
+    # via torch
+nvidia-cusparse-cu12==12.5.10.65
    # via
-    #   cuda-toolkit
-    #   nvidia-cusolver
-nvidia-cusparselt-cu13==0.8.0
+    #   nvidia-cusolver-cu12
+    #   torch
+nvidia-cusparselt-cu12==0.7.1
    # via torch
-nvidia-nccl-cu13==2.28.9
+nvidia-nccl-cu12==2.27.5
    # via torch
-nvidia-nvjitlink==13.0.88
+nvidia-nvjitlink-cu12==12.9.86
    # via
-    #   cuda-toolkit
-    #   nvidia-cufft
-    #   nvidia-cusolver
-    #   nvidia-cusparse
-nvidia-nvshmem-cu13==3.4.5
+    #   nvidia-cufft-cu12
+    #   nvidia-cusolver-cu12
+    #   nvidia-cusparse-cu12
+    #   torch
+nvidia-nvshmem-cu12==3.4.5
+    # via torch
+nvidia-nvtx-cu12==12.9.79
    # via torch
-nvidia-nvtx==13.0.85
-    # via cuda-toolkit
 omegaconf==2.3.0
    # via
    #   hydra-core
@@ -979,7 +977,7 @@ referencing==0.35.1
    # via
    #   jsonschema
    #   jsonschema-specifications
-regex==2026.4.4
+regex==2024.9.11
    # via
    #   diffusers
    #   nltk
@@ -999,6 +997,7 @@ requests==2.32.3
    #   google-api-core
    #   google-cloud-storage
    #   gpt-oss
+    #   huggingface-hub
    #   lightly
    #   lm-eval
    #   mistral-common
@@ -1011,6 +1010,7 @@ requests==2.32.3
    #   starlette-testclient
    #   tacoreader
    #   tiktoken
+    #   transformers
    #   wandb
 resampy==0.4.3
    # via -r requirements/test.in
@@ -1211,7 +1211,7 @@ timm==1.0.17
    #   segmentation-models-pytorch
    #   terratorch
    #   torchgeo
-tokenizers==0.22.2
+tokenizers==0.22.0
    # via
    #   -c requirements/common.txt
    #   -r requirements/test.in
@@ -1220,7 +1220,7 @@ tomli==2.2.1
    # via schemathesis
 tomli-w==1.2.0
    # via schemathesis
-torch==2.11.0+cu130
+torch==2.10.0+cu129
    # via
    #   -r requirements/test.in
    #   accelerate
@@ -1240,12 +1240,13 @@ torch==2.11.0+cu130
    #   tensorizer
    #   terratorch
    #   timm
+    #   torchaudio
    #   torchgeo
    #   torchmetrics
    #   torchvision
    #   vector-quantize-pytorch
    #   vocos
-torchaudio==2.11.0+cu130
+torchaudio==2.10.0+cu129
    # via
    #   -r requirements/test.in
    #   encodec
@@ -1258,7 +1259,7 @@ torchmetrics==1.7.4
    #   pytorch-lightning
    #   terratorch
    #   torchgeo
-torchvision==0.26.0+cu130
+torchvision==0.25.0+cu129
    # via
    #   -r requirements/test.in
    #   lightly
@@ -1287,7 +1288,7 @@ tqdm==4.67.3
    #   tacoreader
    #   terratorch
    #   transformers
-transformers==5.5.0
+transformers==4.57.5
    # via
    #   -c requirements/common.txt
    #   -r requirements/test.in
@@ -1309,9 +1310,7 @@ typepy==1.3.2
 typer==0.15.2
    # via
    #   fastsafetensors
-    #   huggingface-hub
    #   perceptron
-    #   transformers
 types-python-dateutil==2.9.0.20241206
    # via arrow
 typeshed-client==2.8.2
@@ -1,5 +1,5 @@
 # This file was autogenerated by uv via the following command:
-#    uv pip compile requirements/xpu-test.in -o requirements/xpu-test.txt -c requirements/xpu.txt --python-version 3.12 --index-strategy unsafe-best-match --python-platform x86_64-manylinux_2_28
+#    uv pip compile requirements/xpu-test.in -o requirements/xpu-test.txt -c requirements/xpu.txt --python-version 3.12 --index-strategy unsafe-best-match
 absl-py==2.4.0
    # via
    #   -r requirements/xpu-test.in
@@ -19,9 +19,7 @@ aiosignal==1.4.0
 albumentations==1.4.6
    # via -r requirements/xpu-test.in
 annotated-doc==0.0.4
-    # via
-    #   fastapi
-    #   typer
+    # via fastapi
 annotated-types==0.7.0
    # via pydantic
 anyio==4.13.0
@@ -66,7 +64,6 @@ click==8.3.1
    #   jiwer
    #   nltk
    #   schemathesis
-    #   typer
    #   uvicorn
 colorama==0.4.6
    # via sacrebleu
@@ -93,7 +90,7 @@ docker==7.1.0
    # via gpt-oss
 docopt==0.6.2
    # via num2words
-dpcpp-cpp-rt==2025.3.2
+dpcpp-cpp-rt==2025.3.1
    # via
    #   onemkl-sycl-blas
    #   onemkl-sycl-dft
@@ -115,6 +112,7 @@ filelock==3.25.2
    #   huggingface-hub
    #   modelscope
    #   torch
+    #   transformers
 frozenlist==1.8.0
    # via
    #   aiohttp
@@ -137,7 +135,7 @@ harfile==0.4.0
    # via schemathesis
 hf-transfer==0.1.9
    # via -r requirements/xpu-test.in
-hf-xet==1.4.3
+hf-xet==1.4.2
    # via huggingface-hub
 html2text==2025.4.15
    # via gpt-oss
@@ -146,9 +144,8 @@ httpcore==1.0.9
 httpx==0.28.1
    # via
    #   datasets
-    #   huggingface-hub
    #   schemathesis
-huggingface-hub==1.9.2
+huggingface-hub==0.36.2
    # via
    #   accelerate
    #   datasets
@@ -174,27 +171,27 @@ idna==3.11
    #   yarl
 imageio==2.37.3
    # via scikit-image
-impi-rt==2021.17.2
+impi-rt==2021.17.0
    # via
    #   oneccl
    #   torch
 iniconfig==2.3.0
    # via pytest
-intel-cmplr-lib-rt==2025.3.2
+intel-cmplr-lib-rt==2025.3.1
    # via
    #   intel-sycl-rt
    #   torch
-intel-cmplr-lib-ur==2025.3.2
+intel-cmplr-lib-ur==2025.3.1
    # via
    #   intel-openmp
    #   intel-sycl-rt
    #   torch
-intel-cmplr-lic-rt==2025.3.2
+intel-cmplr-lic-rt==2025.3.1
    # via
    #   intel-opencl-rt
    #   intel-sycl-rt
    #   torch
-intel-opencl-rt==2025.3.2
+intel-opencl-rt==2025.3.1
    # via
    #   dpcpp-cpp-rt
    #   onemkl-sycl-blas
@@ -203,14 +200,14 @@ intel-opencl-rt==2025.3.2
    #   onemkl-sycl-rng
    #   onemkl-sycl-sparse
    #   torch
-intel-openmp==2025.3.2
+intel-openmp==2025.3.1
    # via
    #   dpcpp-cpp-rt
    #   mkl
    #   torch
-intel-pti==0.16.0
+intel-pti==0.15.0
    # via torch
-intel-sycl-rt==2025.3.2
+intel-sycl-rt==2025.3.1
    # via
    #   dpcpp-cpp-rt
    #   oneccl
@@ -268,11 +265,11 @@ mbstrdecoder==1.1.4
    #   typepy
 mdurl==0.1.2
    # via markdown-it-py
-mistral-common==1.11.0
+mistral-common==1.10.0
    # via
    #   -c requirements/common.txt
    #   -r requirements/xpu-test.in
-mkl==2025.3.1
+mkl==2025.3.0
    # via
    #   onemkl-sycl-blas
    #   onemkl-sycl-dft
@@ -337,28 +334,28 @@ numpy==2.2.6
    #   tifffile
    #   torchvision
    #   transformers
-oneccl==2021.17.2
+oneccl==2021.17.1
    # via
    #   oneccl-devel
    #   torch
-oneccl-devel==2021.17.2
+oneccl-devel==2021.17.1
    # via torch
-onemkl-license==2025.3.1
+onemkl-license==2025.3.0
    # via
    #   mkl
    #   torch
-onemkl-sycl-blas==2025.3.1
+onemkl-sycl-blas==2025.3.0
    # via
    #   onemkl-sycl-lapack
    #   onemkl-sycl-sparse
    #   torch
-onemkl-sycl-dft==2025.3.1
+onemkl-sycl-dft==2025.3.0
    # via torch
-onemkl-sycl-lapack==2025.3.1
+onemkl-sycl-lapack==2025.3.0
    # via torch
-onemkl-sycl-rng==2025.3.1
+onemkl-sycl-rng==2025.3.0
    # via torch
-onemkl-sycl-sparse==2025.3.1
+onemkl-sycl-sparse==2025.3.0
    # via torch
 openai-harmony==0.0.8
    # via
@@ -518,6 +515,7 @@ requests==2.33.1
    #   docker
    #   evaluate
    #   gpt-oss
+    #   huggingface-hub
    #   lm-eval
    #   mistral-common
    #   modelscope
@@ -526,11 +524,11 @@ requests==2.33.1
    #   schemathesis
    #   starlette-testclient
    #   tiktoken
+    #   transformers
 rich==14.3.3
    # via
    #   mteb
    #   schemathesis
-    #   typer
 rouge-score==0.1.2
    # via lm-eval
 rpds-py==0.30.0
@@ -574,8 +572,6 @@ setuptools==80.10.2
    #   modelscope
    #   pytablewriter
    #   torch
-shellingham==1.5.4
-    # via typer
 six==1.17.0
    # via
    #   -c requirements/common.txt
@@ -610,7 +606,7 @@ tabledata==1.3.4
    # via pytablewriter
 tabulate==0.10.0
    # via sacrebleu
-tbb==2022.3.1
+tbb==2022.3.0
    # via
    #   intel-opencl-rt
    #   mkl
@@ -647,7 +643,7 @@ tokenizers==0.22.2
    # via
    #   -c requirements/common.txt
    #   transformers
-torch==2.11.0+xpu
+torch==2.10.0+xpu
    # via
    #   -c requirements/xpu.txt
    #   accelerate
@@ -655,7 +651,7 @@ torch==2.11.0+xpu
    #   sentence-transformers
    #   timm
    #   torchvision
-torchvision==0.26.0+xpu
+torchvision==0.25.0+xpu
    # via timm
 tqdm==4.67.3
    # via
@@ -669,21 +665,17 @@ tqdm==4.67.3
    #   pqdm
    #   sentence-transformers
    #   transformers
-transformers==5.5.0
+transformers==4.57.6
    # via
    #   -c requirements/common.txt
    #   sentence-transformers
-triton-xpu==3.7.0
+triton-xpu==3.6.0
    # via torch
 typepy==1.3.4
    # via
    #   dataproperty
    #   pytablewriter
    #   tabledata
-typer==0.24.1
-    # via
-    #   huggingface-hub
-    #   transformers
 typing-extensions==4.15.0
    # via
    #   -c requirements/common.txt
@@ -712,7 +704,7 @@ typing-inspection==0.4.2
    # via
    #   fastapi
    #   pydantic
-umf==1.0.3
+umf==1.0.2
    # via
    #   intel-cmplr-lib-ur
    #   torch
@@ -11,7 +11,7 @@ jinja2>=3.1.6
 datasets # for benchmark scripts
 numba == 0.61.2 # Required for N-gram speculative decoding
 --extra-index-url=https://download.pytorch.org/whl/xpu
-torch==2.11.0+xpu
+torch==2.10.0+xpu
 torchaudio
 torchvision

@@ -1013,7 +1013,6 @@ package_data = {
        "model_executor/layers/quantization/utils/configs/*.json",
        "entrypoints/serve/instrumentator/static/*.js",
        "entrypoints/serve/instrumentator/static/*.css",
-        "distributed/kv_transfer/kv_connector/v1/hf3fs/utils/*.cpp",
    ]
 }

@@ -1061,6 +1060,8 @@ setup(
        ],  # Required for audio processing
        "video": [],  # Kept for backwards compatibility
        "flashinfer": [],  # Kept for backwards compatibility
+        # Optional deps for AMD FP4 quantization support
+        "petit-kernel": ["petit-kernel"],
        # Optional deps for Helion kernel development
        # NOTE: When updating helion version, also update CI files:
        #   - .buildkite/test_areas/kernels.yaml
@@ -216,14 +216,12 @@ def test_splitting_ops_dynamic():
        compilation_config=CompilationConfig(
            mode=CompilationMode.VLLM_COMPILE,
            use_inductor_graph_partition=True,
-            splitting_ops=["vllm::unified_attention_with_output"],
+            splitting_ops=["vllm::unified_attention"],
        )
    )
    # with inductor partition we use splitting_ops directly for
    # partition rules
-    assert config.compilation_config.splitting_ops == [
-        "vllm::unified_attention_with_output"
-    ]
+    assert config.compilation_config.splitting_ops == ["vllm::unified_attention"]

    # When attn_fusion pass enabled.
    config = VllmConfig(
@@ -283,7 +281,7 @@ def test_moe_splitting_ops_deepep_ht_inductor_partition():
            mode=CompilationMode.VLLM_COMPILE,
            use_inductor_graph_partition=True,
            splitting_ops=[
-                "vllm::unified_attention_with_output",
+                "vllm::unified_attention",
                "vllm::moe_forward",
                "vllm::moe_forward_shared",
            ],
@@ -291,7 +289,7 @@ def test_moe_splitting_ops_deepep_ht_inductor_partition():
    )
    splitting_ops = config.compilation_config.splitting_ops
    assert splitting_ops == [
-        "vllm::unified_attention_with_output",
+        "vllm::unified_attention",
        "vllm::moe_forward",
        "vllm::moe_forward_shared",
    ]
@@ -0,0 +1,119 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Multi-node integration test for MessageQueue TCP fallback.
+
+Verifies that when writer and readers span separate nodes (Docker containers
+with isolated /dev/shm), `create_from_process_group` correctly detects
+cross-node ranks via `in_the_same_node_as()` and falls back to ZMQ TCP
+transport — and that data actually arrives.
+"""
+
+import numpy as np
+import torch.distributed as dist
+
+from vllm.distributed.device_communicators.shm_broadcast import MessageQueue
+from vllm.distributed.parallel_state import in_the_same_node_as
+
+
+def main():
+    dist.init_process_group(backend="gloo")
+
+    rank = dist.get_rank()
+    world_size = dist.get_world_size()
+    assert world_size >= 2, (
+        f"Need at least 2 ranks across nodes, got world_size={world_size}"
+    )
+
+    # Verify that in_the_same_node_as detects cross-node correctly
+    status = in_the_same_node_as(dist.group.WORLD, source_rank=0)
+    local_count = sum(status)
+    print(
+        f"[Rank {rank}] in_the_same_node_as(source=0): {status}  "
+        f"(local={local_count}/{world_size})"
+    )
+    # With 2 Docker containers (1 proc each), rank 0 and rank 1
+    # should be on different nodes.
+    assert local_count < world_size, (
+        f"Expected cross-node ranks but all {world_size} ranks appear local."
+    )
+
+    # Create MessageQueue
+    writer_rank = 0
+    mq = MessageQueue.create_from_process_group(
+        dist.group.WORLD,
+        max_chunk_bytes=1024 * 1024,  # 1 MiB
+        max_chunks=10,
+        writer_rank=writer_rank,
+    )
+
+    # Verify the transport path selection
+    if rank == writer_rank:
+        print(
+            f"[Rank {rank}] Writer: n_local_reader={mq.n_local_reader}, "
+            f"n_remote_reader={mq.n_remote_reader}"
+        )
+        assert mq.n_remote_reader > 0, (
+            "Writer should have at least 1 remote (TCP) reader in a multi-node setup."
+        )
+    else:
+        if status[rank]:
+            assert mq._is_local_reader, (
+                f"Rank {rank} is on the same node as writer but is not a local reader."
+            )
+            print(f"[Rank {rank}] Reader: local (shared memory)")
+        else:
+            assert mq._is_remote_reader, (
+                f"Rank {rank} is on a different node but is not a remote (TCP) reader."
+            )
+            print(f"[Rank {rank}] Reader: remote (TCP)")
+
+    # Test data transfer: simple objects
+    dist.barrier()
+    if rank == writer_rank:
+        mq.enqueue("hello_from_node0")
+    else:
+        msg = mq.dequeue(timeout=10)
+        assert msg == "hello_from_node0"
+    dist.barrier()
+    print(f"[Rank {rank}] Simple object test passed")
+
+    # Test data transfer: numpy arrays
+    np.random.seed(42)
+    arrays = [
+        np.random.randint(0, 100, size=np.random.randint(100, 5000)) for _ in range(100)
+    ]
+
+    dist.barrier()
+    if rank == writer_rank:
+        for arr in arrays:
+            mq.enqueue(arr)
+    else:
+        for i, expected in enumerate(arrays):
+            received = mq.dequeue(timeout=10)
+            assert np.array_equal(expected, received), (
+                f"Array mismatch at index {i}: "
+                f"expected shape {expected.shape}, got shape {received.shape}"
+            )
+    dist.barrier()
+    print(f"[Rank {rank}] Numpy array test passed")
+
+    # Test data transfer: large payload (> max_chunk_bytes)
+    dist.barrier()
+    big_array = np.zeros(200_000, dtype=np.int64)  # ~1.6 MiB > 1 MiB chunk
+    if rank == writer_rank:
+        mq.enqueue(big_array)
+    else:
+        received = mq.dequeue(timeout=10)
+        assert np.array_equal(big_array, received)
+    dist.barrier()
+    print(f"[Rank {rank}] Large payload test passed")
+
+    # Done -- cleanup
+    dist.barrier()
+    print(f"[Rank {rank}] All MessageQueue TCP multi-node tests passed!")
+    dist.destroy_process_group()
+
+
+if __name__ == "__main__":
+    main()
@@ -525,29 +525,6 @@ def test_human_readable_model_len():
            parser.parse_args(["--max-model-len", invalid])


-def test_numa_bind_args():
-    parser = EngineArgs.add_cli_args(FlexibleArgumentParser())
-    args = parser.parse_args(
-        [
-            "--numa-bind",
-            "--numa-bind-nodes",
-            "0",
-            "0",
-            "1",
-            "1",
-            "--numa-bind-cpus",
-            "0-3",
-            "4-7",
-            "8-11",
-            "12-15",
-        ]
-    )
-    engine_args = EngineArgs.from_cli_args(args=args)
-    assert engine_args.numa_bind is True
-    assert engine_args.numa_bind_nodes == [0, 0, 1, 1]
-    assert engine_args.numa_bind_cpus == ["0-3", "4-7", "8-11", "12-15"]
-
-
 def test_ir_op_priority():
    from vllm.config.kernel import IrOpPriorityConfig, KernelConfig

@@ -628,31 +628,6 @@ def _identity_increment(event):
    return event


-def _mock_parser_with_reasoning(serving, delta_sequence: list[DeltaMessage]):
-    """Set up serving.parser so that it returns a mock parser instance
-    with a reasoning parser that returns the given delta_sequence.
-
-    The mock has reasoning_parser set (truthy) but tool_parser as None,
-    so the parser's parse_delta enters the reasoning-only branch.
-    """
-    call_count = 0
-
-    def mock_parse_delta(**kwargs):
-        nonlocal call_count
-        if call_count >= len(delta_sequence):
-            return None
-        result = delta_sequence[call_count]
-        call_count += 1
-        return result
-
-    mock_parser_instance = MagicMock()
-    mock_parser_instance.reasoning_parser = MagicMock()  # truthy
-    mock_parser_instance.tool_parser = None
-    mock_parser_instance.parse_delta = mock_parse_delta
-    mock_parser_instance.is_reasoning_end = MagicMock(return_value=False)
-    serving.parser = MagicMock(return_value=mock_parser_instance)
-
-
 class TestStreamingReasoningToContentTransition:
    """Tests for _process_simple_streaming_events reasoning-to-content
    transition, specifically the fix for mixed deltas that carry both
@@ -671,13 +646,27 @@ class TestStreamingReasoningToContentTransition:
        monkeypatch.setattr(envs, "VLLM_USE_EXPERIMENTAL_PARSER_CONTEXT", False)
        serving = _make_serving_instance_with_reasoning()

-        # Sequence of DeltaMessages the mock orchestrator will return
+        # Sequence of DeltaMessages the mock reasoning parser will return
        delta_sequence = [
            DeltaMessage(reasoning="thinking..."),
            DeltaMessage(reasoning=" end", content="hello"),  # mixed delta
            DeltaMessage(content=" world"),
        ]
-        _mock_parser_with_reasoning(serving, delta_sequence)
+        call_count = 0
+
+        def mock_extract_reasoning_streaming(**kwargs):
+            nonlocal call_count
+            result = delta_sequence[call_count]
+            call_count += 1
+            return result
+
+        # Mock the reasoning parser on the serving instance
+        mock_parser = MagicMock()
+        mock_parser.extract_reasoning_streaming = mock_extract_reasoning_streaming
+        mock_parser.extract_tool_calls_streaming = mock_extract_reasoning_streaming
+        serving.parser = MagicMock()
+        serving.parser.reasoning_parser_cls = MagicMock(return_value=mock_parser)
+        serving.parser.tool_parser_cls = MagicMock(return_value=mock_parser)
        # Create contexts for each streaming chunk
        contexts = [
            _make_simple_context_with_output("chunk1", [10]),
@@ -745,7 +734,20 @@ class TestStreamingReasoningToContentTransition:
            DeltaMessage(reasoning="thinking"),
            DeltaMessage(content="answer"),
        ]
-        _mock_parser_with_reasoning(serving, delta_sequence)
+        call_count = 0
+
+        def mock_extract_reasoning_streaming(**kwargs):
+            nonlocal call_count
+            result = delta_sequence[call_count]
+            call_count += 1
+            return result
+
+        mock_parser = MagicMock()
+        mock_parser.extract_reasoning_streaming = mock_extract_reasoning_streaming
+        mock_parser.extract_tool_calls_streaming = mock_extract_reasoning_streaming
+        serving.parser = MagicMock()
+        serving.parser.reasoning_parser_cls = MagicMock(return_value=mock_parser)
+        serving.parser.tool_parser_cls = MagicMock(return_value=mock_parser)

        contexts = [
            _make_simple_context_with_output("chunk1", [10]),
@@ -807,7 +809,20 @@ class TestStreamingReasoningToContentTransition:
            DeltaMessage(reasoning="step 1"),
            DeltaMessage(reasoning=" step 2"),
        ]
-        _mock_parser_with_reasoning(serving, delta_sequence)
+        call_count = 0
+
+        def mock_extract_reasoning_streaming(**kwargs):
+            nonlocal call_count
+            result = delta_sequence[call_count]
+            call_count += 1
+            return result
+
+        mock_parser = MagicMock()
+        mock_parser.extract_reasoning_streaming = mock_extract_reasoning_streaming
+        mock_parser.extract_tool_calls_streaming = mock_extract_reasoning_streaming
+        serving.parser = MagicMock()
+        serving.parser.reasoning_parser_cls = MagicMock(return_value=mock_parser)
+        serving.parser.tool_parser_cls = MagicMock(return_value=mock_parser)

        contexts = [
            _make_simple_context_with_output("chunk1", [10]),
@@ -3,4 +3,4 @@
 model_name: openai/gpt-oss-20b
 metric_threshold: 0.568
 reasoning_effort: low
-server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN --tensor-parallel-size 2"
+server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN"
@@ -3,6 +3,6 @@
 model_name: amd/gpt-oss-20b-w-mxfp4-a-bf16
 metric_threshold: 0.568
 reasoning_effort: low
-server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN --moe-backend aiter --tokenizer openai/gpt-oss-20b --tensor-parallel-size 2"
+server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN --moe-backend aiter"
 env:
-  VLLM_ROCM_USE_AITER: "1"
+  VLLM_ROCM_USE_AITER: "1"
@@ -3,4 +3,4 @@
 model_name: amd/gpt-oss-20b-w-mxfp4-a-bf16
 metric_threshold: 0.568
 reasoning_effort: low
-server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN --moe-backend triton --tokenizer openai/gpt-oss-20b --tensor-parallel-size 2"
+server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN --moe-backend triton"
@@ -3,6 +3,6 @@
 model_name: amd/gpt-oss-20b-MoE-Quant-W-MXFP4-A-FP8-KV-FP8
 metric_threshold: 0.568
 reasoning_effort: low
-server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN --tensor-parallel-size 2"
+server_args: "--attention-backend ROCM_AITER_UNIFIED_ATTN"
 env:
  VLLM_ROCM_USE_AITER: "1"
@@ -16,7 +16,7 @@ from vllm.model_executor.layers.quantization.utils.int8_utils import (
 from vllm.platforms import current_platform

 DTYPES = [torch.float16, torch.bfloat16]
-QUANT_DTYPES = [current_platform.fp8_dtype(), torch.int8]
+QUANT_DTYPES = [torch.float8_e4m3fn, torch.int8]
 VEC_HIDDEN_SIZES = [1024, 1025, 1027, 1029]
 NUM_TOKENS_HIDDEN_SIZES = [
    *[(1, i) for i in [64, *VEC_HIDDEN_SIZES, 2048, 5120]],
@@ -28,7 +28,9 @@ SCALE_UBS = [False]
 GROUP_SIZES = [64, 128]
 IS_SCALE_TRANSPOSED = [False, True]
 SEEDS = [0]
-CUDA_DEVICES = [i for i in range(1 if torch.accelerator.device_count() == 1 else 2)]
+CUDA_DEVICES = [
+    f"cuda:{i}" for i in range(1 if torch.accelerator.device_count() == 1 else 2)
+]


 def ref_silu_and_mul_per_block_quant(
@@ -58,7 +60,7 @@ def ref_silu_and_mul_per_block_quant(
@pytest.mark.parametrize("group_size", GROUP_SIZES)
@pytest.mark.parametrize("is_scale_transposed", IS_SCALE_TRANSPOSED)
@pytest.mark.parametrize("seed", SEEDS)
-@pytest.mark.parametrize("device_idx", CUDA_DEVICES)
+@pytest.mark.parametrize("device", CUDA_DEVICES)
@torch.inference_mode()
 def test_silu_and_mul_per_block_quant(
    default_vllm_config,
@@ -70,11 +72,9 @@ def test_silu_and_mul_per_block_quant(
    group_size: int,
    is_scale_transposed: bool,
    seed: int,
-    device_idx: str,
+    device: str,
 ) -> None:
    """Test SiLU+Mul+Block Quantization kernel correctness."""
-    torch.accelerator.set_device_index(device_idx)
-    device = f"cuda:{device_idx}"
    torch.random.manual_seed(seed)
    torch.set_default_device(device)

@@ -147,7 +147,7 @@ def test_silu_block_quant_shapes(
    out, scales = ops.silu_and_mul_per_block_quant(
        x,
        group_size=group_size,
-        quant_dtype=current_platform.fp8_dtype(),
+        quant_dtype=torch.float8_e4m3fn,
        is_scale_transposed=False,
    )
    assert out.shape == (num_tokens, hidden_size)
@@ -157,7 +157,7 @@ def test_silu_block_quant_shapes(
    out, scales = ops.silu_and_mul_per_block_quant(
        x,
        group_size=group_size,
-        quant_dtype=current_platform.fp8_dtype(),
+        quant_dtype=torch.float8_e4m3fn,
        is_scale_transposed=True,
    )
    assert out.shape == (num_tokens, hidden_size)
@@ -177,12 +177,12 @@ def test_silu_block_quant_edge_cases(
    out, scales = ops.silu_and_mul_per_block_quant(
        x,
        group_size=128,
-        quant_dtype=current_platform.fp8_dtype(),
+        quant_dtype=torch.float8_e4m3fn,
        is_scale_transposed=False,
    )

    assert out.shape == (batch_size, hidden_size)
-    assert out.dtype == current_platform.fp8_dtype()
+    assert out.dtype == torch.float8_e4m3fn
    assert scales.dtype == torch.float32
    assert not torch.isnan(out.float()).any()
    assert not torch.isnan(scales).any()
@@ -6,21 +6,13 @@ import torch

 from tests.kernels.quant_utils import FP8_DTYPE
 from tests.kernels.utils import opcheck
-from vllm.model_executor.layers.layernorm import GemmaRMSNorm, RMSNorm
-from vllm.platforms import current_platform
+from vllm.model_executor.layers.layernorm import RMSNorm
 from vllm.utils.torch_utils import set_random_seed

-if current_platform.is_rocm():
-    from vllm.platforms.rocm import on_gfx90a
-
-    on_mi250 = on_gfx90a()
-else:
-    on_mi250 = False
-
 DTYPES = [torch.half, torch.bfloat16, torch.float]
 NUM_TOKENS = [7, 83, 4096]  # Arbitrary values for testing
 HIDDEN_SIZES = [8, 768, 769, 5120, 5125, 8192]  # Arbitrary values for testing
-ADD_RESIDUAL = [False, True] if not on_mi250 else [True]
+ADD_RESIDUAL = [False, True]
 SEEDS = [0]
 CUDA_DEVICES = [
    f"cuda:{i}" for i in range(1 if torch.accelerator.device_count() == 1 else 2)
@@ -162,31 +154,3 @@ def test_fused_rms_norm_quant(
        atol=1e-3,
        rtol=1e-3,
    )
-
-
-@torch.inference_mode()
-def test_gemma_rms_norm_mixed_input_weight_dtype(default_vllm_config) -> None:
-    if not torch.cuda.is_available():
-        pytest.skip("CUDA required")
-
-    device = CUDA_DEVICES[0]
-    torch.set_default_device(device)
-
-    num_tokens, hidden_size = 32, 1024
-    x = torch.randn(num_tokens, hidden_size, dtype=torch.bfloat16, device=device)
-    layer = GemmaRMSNorm(hidden_size, eps=1e-6).to(device=device)
-    layer.weight.data.normal_(mean=0.0, std=0.1)
-
-    # Gemma uses fp32 weight parameter while activations can be bf16.
-    assert layer.weight.dtype == torch.float32
-    out = layer(x)
-
-    x_fp32 = x.float()
-    weight_fp32 = layer.weight.data.float() + 1.0
-    variance = x_fp32.pow(2).mean(dim=-1, keepdim=True)
-    ref = (x_fp32 * torch.rsqrt(variance + layer.variance_epsilon) * weight_fp32).to(
-        x.dtype
-    )
-
-    assert out.dtype == x.dtype
-    torch.testing.assert_close(out, ref, atol=1e-2, rtol=1e-2)
@@ -67,7 +67,6 @@ class TestMakeFxHop:
    def setup_method(self):
        helion_kernel_side_table.reset_table()

-    @pytest.mark.skip(reason="SymInt proxy tracking issue with PyTorch 2.11+")
    def test_make_fx_symbolic(self):
        def raw_add_scale(
            x: torch.Tensor, y: torch.Tensor, scale: float
@@ -129,7 +128,6 @@ class TestMakeFxHop:
            for out_s, in_s in zip(val.shape, input_shape):
                assert out_s == in_s

-    @pytest.mark.skip(reason="SymInt proxy tracking issue with PyTorch 2.11+")
    def test_pattern_matcher_replaces_with_helion_hop(self):
        def raw_silu_mul(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
            M, N = x.size()
@@ -1,14 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import pytest
-
-
-def pytest_addoption(parser):
-    parser.addoption(
-        "--subtests", action="store", type=str, default=None, help="subtest ids"
-    )
-
-
-@pytest.fixture
-def subtests(request):
-    return request.config.getoption("--subtests")
@@ -11,11 +11,7 @@ from torch.multiprocessing import spawn  # pyright: ignore[reportPrivateImportUs
 from typing_extensions import ParamSpec

 from vllm.config import VllmConfig, set_current_vllm_config
-from vllm.distributed import (
-    cleanup_dist_env_and_memory,
-    init_distributed_environment,
-    initialize_model_parallel,
-)
+from vllm.distributed import init_distributed_environment, initialize_model_parallel
 from vllm.utils.network_utils import get_open_port

 ## Parallel Processes Utils
@@ -40,17 +36,10 @@ def _set_vllm_config(

    temp_file = tempfile.mkstemp()[1]

-    # When DP is enabled, processes are organized as:
-    #  rank = dp_rank * tp_pp_world_size + tp_pp_rank
-    tp_pp_world_size = vllm_config.parallel_config.world_size
-    vllm_config.parallel_config.data_parallel_rank = rank // tp_pp_world_size
-    tp_pp_rank = rank % tp_pp_world_size
-    vllm_config.parallel_config.rank = tp_pp_rank
-
    with set_current_vllm_config(vllm_config):
        init_distributed_environment(
-            world_size=tp_pp_world_size,
-            rank=tp_pp_rank,
+            world_size=world_size,
+            rank=rank,
            distributed_init_method=f"file://{temp_file}",
            local_rank=local_rank,
            backend="nccl",
@@ -70,11 +59,11 @@ def _worker_parallel_launch(
    world_local_size: int,
    node_rank: int,
    init_method: str,
-    worker: Callable[..., None],
+    worker: Callable[Concatenate[ProcessGroupInfo, VllmConfig | None, Any, P], None],
    vllm_config: VllmConfig | None,
    env_dict: dict | None,
-    worker_kwargs: dict[str, Any],
-    *args: Any,
+    *args: P.args,
+    **kwargs: P.kwargs,
 ) -> None:
    rank = node_rank * world_local_size + local_rank
    torch.accelerator.set_device_index(local_rank)
@@ -109,17 +98,14 @@ def _worker_parallel_launch(
            vllm_config,
            cpu_group,
            *args,
-            **worker_kwargs,
+            **kwargs,
        )
    except Exception as ex:
        print(ex)
        traceback.print_exc()
        raise
    finally:
-        if vllm_config is not None:
-            cleanup_dist_env_and_memory()
-        else:
-            torch.distributed.destroy_process_group()
+        torch.distributed.destroy_process_group()


 def parallel_launch_with_config(
@@ -130,6 +116,7 @@ def parallel_launch_with_config(
    *args: P.args,
    **kwargs: P.kwargs,
 ) -> None:
+    assert not kwargs
    spawn(
        _worker_parallel_launch,
        args=(
@@ -140,7 +127,6 @@ def parallel_launch_with_config(
            worker,
            vllm_config,
            env_dict,
-            kwargs,
        )
        + args,
        nprocs=world_size,
@@ -17,7 +17,7 @@ from flashinfer import fp4_quantize
 from torch.nn import functional as F

 from vllm.model_executor.layers.activation import SiluAndMul
-from vllm.model_executor.layers.fused_moe.experts.flashinfer_cutedsl_batched_moe import (  # noqa: E501
+from vllm.model_executor.layers.fused_moe.experts.flashinfer_cutedsl_moe import (
    flashinfer_cutedsl_moe_masked,
 )
 from vllm.utils.flashinfer import (
@@ -23,12 +23,16 @@ from triton_kernels.numerics_details.mxfp import downcast_to_mxfp, upcast_from_m
 from triton_kernels.tensor import FP4, convert_layout, wrap_torch_tensor
 from triton_kernels.tensor_details import layout
 from triton_kernels.testing import assert_close
+from triton_kernels.topk import topk as topk_fn

 from vllm.model_executor.layers.fused_moe.config import mxfp4_w4a16_moe_quant_config
 from vllm.model_executor.layers.fused_moe.gpt_oss_triton_kernels_moe import (
+    legacy_routing,
+    make_routing_data,
    triton_kernel_moe_forward,
 )
 from vllm.utils.math_utils import round_up
+from vllm.utils.torch_utils import set_random_seed

 from .utils import shuffle_weight

@@ -93,18 +97,10 @@ def init_compute_data(M, K, N, E, a_dtype: str, w_dtype: str, num_warps: int):
    if w_dtype != "mx4":
        pytest.skip("NYI")
    else:  # quantize to mx4
-        # Padding alignment depends on the platform.  On CDNA4 the scale
-        # swizzle requires SCALE_K % 8 == 0 (K % 256) and
-        # SCALE_N % 32 == 0 (2*N % 512), matching the production
-        # alignment in mxfp4_round_up_hidden_size_and_intermediate_size.
-        # On CUDA (Hopper) the scale layout pads internally, so the
-        # original 64/128 alignment is sufficient.
-        if current_platform.is_rocm():
-            k_align, n2_align = 256, 512
-        else:
-            k_align, n2_align = 64, 128
-        w1_bottom_pad = round_up(w1_tri.shape[1], k_align) - w1_tri.shape[1]
-        w1_right_pad = round_up(w1_tri.shape[2], n2_align) - w1_tri.shape[2]
+        # careful on the padding here, the activation padding need to be
+        # multiple of 64, the actual engine is not implemented
+        w1_bottom_pad = round_up(w1_tri.shape[1], 64) - w1_tri.shape[1]
+        w1_right_pad = round_up(w1_tri.shape[2], 128) - w1_tri.shape[2]

        w2_bottom_pad = w1_right_pad // 2
        w2_right_pad = w1_bottom_pad
@@ -371,3 +367,52 @@ def test_unit_shuffle():
    )

    assert_close(ref=out_ref, tri=out)
+
+
+@pytest.mark.parametrize("num_tokens", [2, 8, 64])
+@pytest.mark.parametrize("num_experts", [32, 128])
+@pytest.mark.parametrize("topk", [1, 4])
+@pytest.mark.parametrize("renormalize", [True, False])
+@pytest.mark.parametrize("dtype", [torch.bfloat16])
+def test_legacy_routing(
+    num_tokens: int, num_experts: int, topk: int, renormalize: bool, dtype: torch.dtype
+):
+    set_random_seed(0)
+    gating_output = torch.randn(num_tokens, num_experts, device="cuda", dtype=dtype)
+
+    sm_first = not renormalize
+    logits = gating_output
+    if sm_first:
+        logits = torch.softmax(logits, dim=-1)
+    topk_result = topk_fn(logits, topk, apply_softmax=not sm_first)
+    # topk_fn returns SparseMatrix on NVIDIA, plain tuple on ROCm.
+    if isinstance(topk_result, tuple):
+        topk_weights, topk_ids_raw, bitmatrix = topk_result
+        from triton_kernels.routing import routing_from_bitmatrix
+
+        routing_data_ref, gather_indx_ref, scatter_indx_ref = routing_from_bitmatrix(
+            bitmatrix, topk_weights, topk_ids_raw, num_experts, topk
+        )
+    else:
+        topk_ids = topk_result.indx.to(torch.long)
+        topk_weights = topk_result.vals
+        routing_data_ref, gather_indx_ref, scatter_indx_ref = make_routing_data(
+            topk_ids, topk_weights, num_experts
+        )
+
+    routing_data, gather_indx, scatter_indx = legacy_routing(
+        gating_output, topk, sm_first=sm_first
+    )
+
+    assert_close(
+        ref=gather_indx_ref.src_indx, tri=gather_indx.src_indx, maxtol=0, rmstol=0
+    )
+    assert_close(
+        ref=gather_indx_ref.dst_indx, tri=gather_indx.dst_indx, maxtol=0, rmstol=0
+    )
+    assert_close(
+        ref=scatter_indx_ref.src_indx, tri=scatter_indx.src_indx, maxtol=0, rmstol=0
+    )
+    assert_close(
+        ref=scatter_indx_ref.dst_indx, tri=scatter_indx.dst_indx, maxtol=0, rmstol=0
+    )
@@ -248,7 +248,7 @@ def make_quantized_test_activations(
    return a, a_q, a_scale


-def moe_quantize_weights_2d(
+def moe_quantize_weights(
    w: torch.Tensor,
    w_s: torch.Tensor | None,
    quant_dtype: torch.dtype | str | None,
@@ -293,40 +293,6 @@ def moe_quantize_weights_2d(
    return w, w_s, w_gs


-def moe_quantize_weights(
-    w: torch.Tensor,
-    w_s: torch.Tensor | None,
-    quant_dtype: torch.dtype | str | None,
-    per_token_quant: bool,
-    block_shape: list[int] | None,
-) -> tuple[torch.Tensor, torch.Tensor | None, torch.Tensor | None]:
-    assert w.dim() == 3
-    e, rows, cols = w.shape
-    w_l = [None] * e
-    w_s_l = [None] * e
-    w_gs_l = [None] * e
-    for idx in range(e):
-        w_l[idx], w_s_l[idx], w_gs_l[idx] = moe_quantize_weights_2d(
-            w[idx], None, quant_dtype, per_token_quant, block_shape
-        )
-
-    w = torch.stack(w_l)
-    w_s = torch.stack(w_s_l)
-    w_gs = torch.stack(w_gs_l) if e > 0 and w_gs_l[0] is not None else None
-
-    if w_s.ndim == 2:
-        assert w_s.shape[-1] == 1
-        w_s = w_s.view(-1, 1, 1)
-
-    if block_shape is not None:
-        block_n, block_k = block_shape
-        n_tiles = (rows + block_n - 1) // block_n
-        k_tiles = (cols + block_k - 1) // block_k
-        assert w_s.shape == (e, n_tiles, k_tiles)
-
-    return w, w_s, w_gs
-
-
 def make_test_weight(
    e: int,
    rows: int,
@@ -337,11 +303,30 @@ def make_test_weight(
    per_out_ch_quant: bool = False,
 ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor | None, torch.Tensor | None]:
    w_16 = torch.randn((e, rows, cols), device="cuda", dtype=in_dtype) / 15
+    w_gs = None

    if quant_dtype is not None:
-        w, w_s, w_gs = moe_quantize_weights(
-            w_16, None, quant_dtype, per_out_ch_quant, block_shape
-        )
+        w_l = [None] * e
+        w_s_l = [None] * e
+        w_gs_l = [None] * e
+        for idx in range(e):
+            w_l[idx], w_s_l[idx], w_gs_l[idx] = moe_quantize_weights(
+                w_16[idx], None, quant_dtype, per_out_ch_quant, block_shape
+            )
+
+        w = torch.stack(w_l)
+        w_s = torch.stack(w_s_l)
+        if e > 0 and w_gs_l[0] is not None:
+            w_gs = torch.stack(w_gs_l)
+        if w_s.ndim == 2:
+            assert w_s.shape[-1] == 1
+            w_s = w_s.view(-1, 1, 1)
+
+        if block_shape is not None:
+            block_n, block_k = block_shape
+            n_tiles = (rows + block_n - 1) // block_n
+            k_tiles = (cols + block_k - 1) // block_k
+            assert w_s.shape == (e, n_tiles, k_tiles)
    else:
        w = w_16
        w_s = None
@@ -469,6 +454,7 @@ def fused_moe(
    )


+# CustomOp?
 class BaselineMM(torch.nn.Module):
    def __init__(
        self,
@@ -476,22 +462,13 @@ class BaselineMM(torch.nn.Module):
        out_dtype: torch.dtype,
    ):
        super().__init__()
-        self.b = torch.nn.Parameter(b.to(dtype=torch.float32))
+        self.b = b.to(dtype=torch.float32)
        self.out_dtype = out_dtype

    def forward(self, a: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor | None]:
        return torch.mm(a.to(dtype=torch.float32), self.b).to(self.out_dtype), None


-class BaselineSiluAndMul(torch.nn.Module):
-    def __init__(self):
-        super().__init__()
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        d = x.shape[-1] // 2
-        return torch.nn.functional.silu(x[..., :d]) * x[..., d:]
-
-
 class TestMLP(torch.nn.Module):
    def __init__(
        self,
@@ -502,7 +479,7 @@ class TestMLP(torch.nn.Module):
        super().__init__()
        self.gate_up_proj = BaselineMM(w1, out_dtype)
        self.down_proj = BaselineMM(w2, out_dtype)
-        self.act_fn = BaselineSiluAndMul()
+        self.act_fn = SiluAndMul()

    def forward(self, x):
        x, _ = self.gate_up_proj(x)
@@ -587,24 +564,35 @@ class RealMLP(torch.nn.Module):
        return x


-def make_shared_experts_with_weights(
+def make_shared_experts(
    N: int,
    K: int,
-    in_dtype: torch.dtype,
-    w1: torch.Tensor,
-    w2: torch.Tensor,
-    w1_s: torch.Tensor | None = None,
-    w2_s: torch.Tensor | None = None,
+    in_dtype: torch.dtype = torch.bfloat16,
    quant_dtype: torch.dtype | str | None = None,
 ) -> torch.nn.Module:
+    from vllm.model_executor.layers.quantization.fp8 import Fp8Config
+
+    (_, w1, w1_s, _), (_, w2, w2_s, _) = make_test_weights(
+        1,
+        N,
+        K,
+        in_dtype=in_dtype,
+        quant_dtype=quant_dtype,
+    )
    old_dtype = torch.get_default_dtype()
    try:
        torch.set_default_dtype(in_dtype)
        if quant_dtype == torch.float8_e4m3fn:
-            from vllm.model_executor.layers.quantization.fp8 import Fp8Config
-
+            w1 = w1[0].transpose(0, 1)
+            w2 = w2[0].transpose(0, 1)
+            w1_s = w1_s[0].transpose(0, 1) if w1_s is not None else None
+            w2_s = w2_s[0].transpose(0, 1) if w2_s is not None else None
            quant_config = Fp8Config(True)
        else:
+            w1 = w1[0]
+            w2 = w2[0]
+            w1_s = None
+            w2_s = None
            quant_config = None

        return RealMLP(K, N, w1, w2, "silu", quant_config, w1_s=w1_s, w2_s=w2_s)
@@ -626,22 +614,3 @@ def modular_triton_fused_moe(
        TritonExperts(moe_config, quant_config),
        inplace=False,
    )
-
-
-def make_shared_experts(
-    N: int,
-    K: int,
-    in_dtype: torch.dtype = torch.bfloat16,
-    quant_dtype: torch.dtype | str | None = None,
-) -> torch.nn.Module:
-    (_, w1, w1_s, _), (_, w2, w2_s, _) = make_test_weights(
-        1,
-        N,
-        K,
-        in_dtype=in_dtype,
-        quant_dtype=quant_dtype,
-    )
-
-    return make_shared_experts_with_weights(
-        N, K, in_dtype, w1, w2, w1_s=w1_s, w2_s=w2_s, quant_dtype=quant_dtype
-    )
@@ -4,9 +4,12 @@
 Tests that triton_kernel_moe_forward correctly applies expert_map
 remapping when expert parallelism (EP) is enabled.

-Both EP and non-EP paths use topk + make_routing_data. When expert_map
-is provided, global expert IDs are remapped to local IDs before building
-routing structures.
+Previously, legacy_routing was always used and it produced routing data
+with global expert IDs that didn't correspond to local weight indices,
+causing illegal memory access with EP.  The fix splits routing: when
+expert_map is provided, topk selection is performed first, expert_map is
+applied to remap global→local IDs, and make_routing_data builds routing
+structures from the local IDs.
 """

 from unittest.mock import MagicMock, patch
@@ -21,15 +24,21 @@ class TestTritonMoeForwardExpertMap:

    @pytest.mark.parametrize("expert_map_present", [False, True])
    def test_routing_path_selection(self, expert_map_present):
-        """Verify that both EP and non-EP paths use topk + make_routing_data,
-        and that expert_map remapping is applied when present."""
+        """Verify that the EP-aware routing path is taken when expert_map
+        is present, and the legacy_routing path is taken otherwise."""

        device = "cuda" if torch.cuda.is_available() else "cpu"
+        # This is a structural test: we mock the routing functions to
+        # verify the correct path is exercised.
        mock_expert_map = (
            torch.tensor([0, -1, 1, -1], device=device) if expert_map_present else None
        )

        with (
+            patch(
+                "vllm.model_executor.layers.fused_moe."
+                "gpt_oss_triton_kernels_moe.legacy_routing"
+            ) as mock_legacy,
            patch("triton_kernels.topk.topk") as mock_topk,
            patch(
                "vllm.model_executor.layers.fused_moe."
@@ -44,19 +53,27 @@ class TestTritonMoeForwardExpertMap:
                triton_kernel_moe_forward,
            )

+            # Set up return values
            mock_routing_data = MagicMock()
            mock_gather = MagicMock()
            mock_scatter = MagicMock()

-            sparse_result = MagicMock()
-            sparse_result.indx = torch.tensor([[0, 2]], dtype=torch.int32)
-            sparse_result.vals = torch.tensor([[0.6, 0.4]])
-            mock_topk.return_value = sparse_result
-            mock_make_routing.return_value = (
-                mock_routing_data,
-                mock_gather,
-                mock_scatter,
-            )
+            if expert_map_present:
+                sparse_result = MagicMock()
+                sparse_result.indx = torch.tensor([[0, 2]], dtype=torch.int32)
+                sparse_result.vals = torch.tensor([[0.6, 0.4]])
+                mock_topk.return_value = sparse_result
+                mock_make_routing.return_value = (
+                    mock_routing_data,
+                    mock_gather,
+                    mock_scatter,
+                )
+            else:
+                mock_legacy.return_value = (
+                    mock_routing_data,
+                    mock_gather,
+                    mock_scatter,
+                )

            mock_fused_experts.return_value = torch.zeros((1, 8), device=device)

@@ -75,14 +92,20 @@ class TestTritonMoeForwardExpertMap:
                expert_map=mock_expert_map,
            )

-            # Both paths use topk + make_routing_data
-            mock_topk.assert_called_once()
-            mock_make_routing.assert_called_once()
-
            if expert_map_present:
+                # EP path: should use topk + make_routing_data, NOT
+                # legacy_routing
+                mock_topk.assert_called_once()
+                mock_make_routing.assert_called_once()
+                mock_legacy.assert_not_called()
                # expert_map should be None in the fused_experts call
                # (already applied)
                call_kwargs = mock_fused_experts.call_args
                assert call_kwargs[1].get("expert_map") is None or (
                    len(call_kwargs[0]) > 0
                )
+            else:
+                # Non-EP path: should use legacy_routing
+                mock_legacy.assert_called_once()
+                mock_topk.assert_not_called()
+                mock_make_routing.assert_not_called()
@@ -122,39 +122,6 @@ def compare_top_k_results(
    return True


-def validate_topk_against_reference(
-    logits: torch.Tensor,
-    cuda_indices: torch.Tensor,
-    row_starts: torch.Tensor,
-    row_ends: torch.Tensor,
-    top_k: int,
-    kernel_name: str,
-) -> None:
-    """
-    Validate CUDA top-k results against PyTorch reference implementation.
-
-    Args:
-        logits: Input logits tensor
-        cuda_indices: CUDA kernel output indices
-        row_starts: Row start positions
-        row_ends: Row end positions
-        top_k: Number of top elements to select
-        kernel_name: Name of the kernel being tested (for error messages)
-    """
-    num_rows = cuda_indices.shape[0]
-    torch_indices = torch.empty((num_rows, top_k), dtype=torch.int32, device="cuda")
-
-    for i in range(num_rows):
-        row_end = int(row_ends[i])
-        k_i = min(top_k, row_end)
-        idx = logits[i, :row_end].topk(k_i, dim=-1)[1]
-        torch_indices[i, :k_i] = idx
-
-    assert compare_top_k_results(
-        logits, cuda_indices, torch_indices, row_starts, row_ends, top_k
-    ), f"{kernel_name} results don't match torch.topk"
-
-
@pytest.mark.parametrize("num_rows", NUM_ROWS)
@pytest.mark.parametrize("top_k", TOP_K_VALUES)
@pytest.mark.parametrize("clean_logits", [True, False])
@@ -311,540 +278,111 @@ def test_top_k_per_row_decode_large_vocab_size(clean_logits: bool) -> None:


@pytest.mark.skipif(not current_platform.is_cuda(), reason="This test requires CUDA")
-@pytest.mark.parametrize(
-    "seq_len_range,test_id",
-    [
-        pytest.param((4000, 8000), "short_sequences", id="short"),
-        pytest.param((8000, 32000), "medium_sequences", id="medium"),
-        pytest.param((32000, 163840), "long_sequences", id="long"),
-    ],
-)
@pytest.mark.parametrize("clean_logits", [True, False])
-@pytest.mark.parametrize("top_k", [2048])
-@pytest.mark.parametrize("next_n", [1, 4])
@torch.inference_mode()
-def test_deepseek_persistent_topk(
-    seq_len_range: tuple[int, int],
-    test_id: str,
-    clean_logits: bool,
-    top_k: int,
-    next_n: int,
-) -> None:
-    """
-    Test persistent_topk with varying sequence lengths and speculative decoding.
-    Supports speculative decoding with next_n > 1.
-    """
-    set_random_seed(42 if test_id == "short_sequences" else 43)
+def test_deepseek_hybrid_topk(clean_logits: bool) -> None:
    torch.set_default_device("cuda:0")

-    batch_size = 4
-    num_rows = batch_size * next_n
+    top_k = 2048

-    seq_lens = torch.randint(
-        seq_len_range[0],
-        seq_len_range[1],
-        (batch_size,),
-        dtype=torch.int32,
-        device="cuda",
+    # Test case 1: Short sequences (< 8192)
+    batch_size_short = 4
+    next_n = 1
+    num_rows_short = batch_size_short * next_n
+
+    # Create sequences with max length < 8192
+    seq_lens_short = torch.randint(
+        4000, 8000, (batch_size_short,), dtype=torch.int32, device="cuda"
    )

-    # Compute row boundaries for speculative decoding
-    row_starts = torch.zeros(num_rows, dtype=torch.int32, device="cuda")
-    row_indices = torch.arange(num_rows, device="cuda") // next_n
-    next_n_offset = torch.arange(num_rows, device="cuda") % next_n
-    row_ends = seq_lens[row_indices] - next_n + next_n_offset + 1
-
-    logits = create_random_logits(
-        row_starts, row_ends, torch.float32, 42, clean_logits, "random"
+    row_starts_short = torch.zeros(num_rows_short, dtype=torch.int32, device="cuda")
+    row_indices_short = torch.arange(num_rows_short, device="cuda") // next_n
+    next_n_offset_short = torch.arange(num_rows_short, device="cuda") % next_n
+    row_ends_short = (
+        seq_lens_short[row_indices_short] - next_n + next_n_offset_short + 1
    )

-    indices = torch.empty((num_rows, top_k), dtype=torch.int32, device="cuda")
+    logits_short = create_random_logits(
+        row_starts_short, row_ends_short, torch.float32, 42, clean_logits, "random"
+    )

+    indices_vllm = torch.empty(
+        (num_rows_short, top_k), dtype=torch.int32, device="cuda"
+    )
+
+    # Use vllm's kernel for short sequences
+    torch.ops._C.top_k_per_row_decode(
+        logits_short,
+        next_n,
+        seq_lens_short,
+        indices_vllm,
+        num_rows_short,
+        logits_short.stride(0),
+        logits_short.stride(1),
+        top_k,
+    )
+
+    # Test case 2: Long sequences (>= 8192) - should use large_context_topk kernel
+    batch_size_long = 4
+    num_rows_long = batch_size_long * next_n
+
+    # Create sequences with max length >= 8192
+    seq_lens_long = torch.randint(
+        8192, 16384, (batch_size_long,), dtype=torch.int32, device="cuda"
+    )
+
+    row_starts_long = torch.zeros(num_rows_long, dtype=torch.int32, device="cuda")
+    row_indices_long = torch.arange(num_rows_long, device="cuda") // next_n
+    next_n_offset_long = torch.arange(num_rows_long, device="cuda") % next_n
+    row_ends_long = seq_lens_long[row_indices_long] - next_n + next_n_offset_long + 1
+
+    logits_long = create_random_logits(
+        row_starts_long, row_ends_long, torch.float32, 43, clean_logits, "random"
+    )
+
+    indices = torch.empty((num_rows_long, top_k), dtype=torch.int32, device="cuda")
+
+    # Use large_context_topk kernel for long sequences
    if next_n == 1:
-        lengths = seq_lens
+        lengths = seq_lens_long
    else:
-        offsets = torch.arange(next_n, device=logits.device, dtype=torch.int32)
-        lengths = (seq_lens.unsqueeze(1) - next_n + 1 + offsets).flatten()
+        offsets = torch.arange(next_n, device=logits_long.device, dtype=torch.int32)
+        lengths = (seq_lens_long.unsqueeze(1) - next_n + 1 + offsets).flatten()

-    workspace = torch.empty(1024 * 1024, dtype=torch.uint8, device="cuda")
-    max_seq_len = int(seq_lens.max().item())
-    torch.ops._C.persistent_topk(
-        logits, lengths, indices, workspace, top_k, max_seq_len
+    torch.ops._C.large_context_topk(
+        logits_long,
+        indices,
+        lengths,
+        None,
    )

-    validate_topk_against_reference(
-        logits, indices, row_starts, row_ends, top_k, f"persistent_topk ({test_id})"
+    torch_indices_short = torch.empty(
+        (num_rows_short, top_k), dtype=torch.int32, device="cuda"
    )
+    for i in range(num_rows_short):
+        row_end = int(row_ends_short[i])
+        k_i = min(top_k, row_end)
+        idx = logits_short[i, :row_end].topk(k_i, dim=-1)[1]
+        torch_indices_short[i, :k_i] = idx

+    assert compare_top_k_results(
+        logits_short,
+        indices_vllm,
+        torch_indices_short,
+        row_starts_short,
+        row_ends_short,
+        top_k,
+    ), "top_k_per_row_decode kernel (short sequences) doesn't match torch.topk"

-def run_large_context_topk_test(
-    batch_size: int,
-    seq_lens: list[int],
-    top_k: int,
-    data_type: str = "random",
-    seed: int = 42,
-) -> None:
-    """
-    Helper to run persistent_topk kernel test with given parameters.
-
-    Args:
-        batch_size: Number of rows/sequences
-        seq_lens: List of sequence lengths (one per row)
-        top_k: Number of top elements to select
-        data_type: Type of test data to generate
-        seed: Random seed for reproducibility
-    """
-    torch.set_default_device("cuda:0")
-    set_random_seed(seed)
-
-    # Create test data
-    num_rows = batch_size
-    max_len = max(seq_lens)
-    lengths = torch.tensor(seq_lens, dtype=torch.int32, device="cuda")
-
-    if data_type == "random":
-        logits = torch.randn(num_rows, max_len, dtype=torch.float32, device="cuda")
-    elif data_type == "sorted_asc":
-        # Each row gets its own ascending sequence based on its length
-        logits = torch.empty(num_rows, max_len, dtype=torch.float32, device="cuda")
-        for i, length in enumerate(seq_lens):
-            logits[i, :length] = torch.arange(
-                length, dtype=torch.float32, device="cuda"
-            )
-            if length < max_len:
-                logits[i, length:] = float("-inf")
-    elif data_type == "sorted_desc":
-        # Each row gets its own descending sequence based on its length
-        logits = torch.empty(num_rows, max_len, dtype=torch.float32, device="cuda")
-        for i, length in enumerate(seq_lens):
-            logits[i, :length] = torch.arange(
-                length, 0, -1, dtype=torch.float32, device="cuda"
-            )
-            if length < max_len:
-                logits[i, length:] = float("-inf")
-    elif data_type == "all_same":
-        logits = torch.ones(num_rows, max_len, dtype=torch.float32, device="cuda")
-        for i, length in enumerate(seq_lens):
-            if length < max_len:
-                logits[i, length:] = float("-inf")
-    elif data_type == "many_ties":
-        # Only 10 unique values, many duplicates
-        logits = torch.randint(0, 10, (num_rows, max_len), device="cuda").float() / 10.0
-        for i, length in enumerate(seq_lens):
-            if length < max_len:
-                logits[i, length:] = float("-inf")
-    elif data_type == "small_differences":
-        # Very small differences to test float precision
-        base = torch.randn(num_rows, max_len, dtype=torch.float32, device="cuda")
-        noise = (
-            torch.randn(num_rows, max_len, dtype=torch.float32, device="cuda") * 1e-6
-        )
-        logits = base + noise
-        for i, length in enumerate(seq_lens):
-            if length < max_len:
-                logits[i, length:] = float("-inf")
-    else:
-        raise ValueError(f"Unknown data_type: {data_type}")
-
-    # Create output tensor
-    indices = torch.empty((num_rows, top_k), dtype=torch.int32, device="cuda")
-
-    workspace = torch.empty(1024 * 1024, dtype=torch.uint8, device="cuda")
-    max_seq_len = max(seq_lens)
-    torch.ops._C.persistent_topk(
-        logits, lengths, indices, workspace, top_k, max_seq_len
+    torch_indices_long = torch.empty(
+        (num_rows_long, top_k), dtype=torch.int32, device="cuda"
    )
+    for i in range(num_rows_long):
+        row_end = int(row_ends_long[i])
+        k_i = min(top_k, row_end)
+        idx = logits_long[i, :row_end].topk(k_i, dim=-1)[1]
+        torch_indices_long[i, :k_i] = idx

-    torch.accelerator.synchronize()
-
-    torch_indices = torch.empty((num_rows, top_k), dtype=torch.int32, device="cuda")
-    for i in range(num_rows):
-        length = seq_lens[i]
-        k_i = min(top_k, length)
-        if k_i > 0:
-            idx = logits[i, :length].topk(k_i, dim=-1)[1]
-            torch_indices[i, :k_i] = idx
-            if k_i < top_k:
-                torch_indices[i, k_i:] = -1
-        else:
-            torch_indices[i, :] = -1
-
-    # Compare results
-    for i in range(num_rows):
-        length = seq_lens[i]
-        k_i = min(top_k, length)
-
-        if k_i == 0:
-            continue
-
-        cuda_row = indices[i, :k_i].cpu()
-        torch_row = torch_indices[i, :k_i].cpu()
-
-        # Filter out -1 padding values from cuda_row
-        valid_mask = cuda_row >= 0
-        cuda_row = cuda_row[valid_mask]
-
-        # Compare sets (order may differ for ties)
-        cuda_set = set(cuda_row.tolist())
-        torch_set = set(torch_row.tolist())
-
-        if cuda_set == torch_set:
-            continue
-
-        # If sets differ, check if it's due to equal values (ties)
-        cuda_vals = logits[i, cuda_row].cpu()
-        torch_vals = logits[i, torch_row].cpu()
-
-        # Check that min CUDA value >= max of values NOT in top-k
-        if k_i < length:
-            non_topk_indices = torch.tensor(
-                list(set(range(length)) - cuda_set), dtype=torch.int32
-            )
-            if len(non_topk_indices) > 0:
-                non_topk_vals = logits[i, non_topk_indices].cpu()
-                min_cuda_val = cuda_vals.min()
-                max_non_topk = non_topk_vals.max()
-
-                # Allow small tolerance for floating point errors
-                assert min_cuda_val >= max_non_topk - 1e-4, (
-                    f"Row {i}: CUDA top-k contains values smaller than non-top-k. "
-                    f"Min CUDA: {min_cuda_val}, Max non-top-k: {max_non_topk}, "
-                    f"Length: {length}, k: {k_i}, CUDA indices: {sorted(cuda_set)[:10]}..., "  # noqa: E501
-                    f"Expected indices: {sorted(torch_set)[:10]}..."
-                )
-
-        # For ties, verify the values are close
-        assert torch.allclose(
-            cuda_vals.sort(descending=True)[0],
-            torch_vals.sort(descending=True)[0],
-            rtol=1e-4,
-            atol=1e-4,
-        ), f"""Row {i}: Top-k values don't match.
-            CUDA: {cuda_vals.sort(descending=True)[0][:10]},
-            Torch: {torch_vals.sort(descending=True)[0][:10]}"""
-
-
-@pytest.mark.skipif(not current_platform.is_cuda(), reason="This test requires CUDA")
-@pytest.mark.parametrize(
-    "test_config",
-    [
-        # ==================== CATEGORY: Sequence Length Edge Cases ====================
-        pytest.param(
-            {"seq_lens": [1, 10, 100, 2048], "top_k": 2048, "data_type": "random"},
-            id="seq_len_edge_very_small_to_medium",
-        ),
-        pytest.param(
-            {
-                "seq_lens": [2049, 2100, 2500, 3000],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="seq_len_edge_above_k",
-        ),
-        pytest.param(
-            {"seq_lens": [8000, 16384, 20000], "top_k": 2048, "data_type": "random"},
-            id="algo_transition_filtered_radix",
-        ),
-        # ==================== CATEGORY: Data Distributions ====================
-        pytest.param(
-            {"seq_lens": [5000, 10000], "top_k": 2048, "data_type": "sorted_asc"},
-            id="data_sorted_ascending",
-        ),
-        pytest.param(
-            {"seq_lens": [5000, 10000], "top_k": 2048, "data_type": "sorted_desc"},
-            id="data_sorted_descending",
-        ),
-        pytest.param(
-            {"seq_lens": [5000, 10000], "top_k": 2048, "data_type": "all_same"},
-            id="data_all_same",
-        ),
-        pytest.param(
-            {"seq_lens": [5000, 10000], "top_k": 2048, "data_type": "many_ties"},
-            id="data_many_ties",
-        ),
-        pytest.param(
-            {
-                "seq_lens": [5000, 10000],
-                "top_k": 2048,
-                "data_type": "small_differences",
-            },
-            id="data_float_precision",
-        ),
-        # ==================== CATEGORY: Alignment / Vectorization ====================
-        pytest.param(
-            {
-                "seq_lens": [2055, 2056, 2057, 2063],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="align_vec_boundaries_low",
-        ),
-        pytest.param(
-            {
-                "seq_lens": [4095, 4096, 4097, 4102],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="align_4k_boundary",
-        ),
-        pytest.param(
-            {
-                "seq_lens": [8191, 8192, 8193, 8198],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="align_8k_boundary",
-        ),
-        pytest.param(
-            {
-                "seq_lens": [16383, 16384, 16385, 16390],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="align_16k_boundary",
-        ),
-    ],
-)
-@torch.inference_mode()
-def test_persistent_topk_correctness(test_config: dict) -> None:
-    """
-    Comprehensive correctness tests covering:
-    - Sequence length edge cases (trivial, boundary, varied)
-    - Very small sequences (< 100 elements)
-    - Mixed sequence lengths in same batch
-    - Data distributions (sorted, ties, precision)
-    - Memory alignment / vectorization boundaries
-    """
-    run_large_context_topk_test(
-        batch_size=len(test_config["seq_lens"]),
-        seq_lens=test_config["seq_lens"],
-        top_k=test_config["top_k"],
-        data_type=test_config.get("data_type", "random"),
-    )
-
-
-@pytest.mark.skipif(not current_platform.is_cuda(), reason="This test requires CUDA")
-@pytest.mark.parametrize(
-    "test_config",
-    [
-        # ==================== CATEGORY: Batch Size Scalability ====================
-        pytest.param(
-            {"batch_size": 1, "seq_len": 5000, "top_k": 2048},
-            id="batch_1",
-        ),
-        pytest.param(
-            {"batch_size": 4, "seq_len": 5000, "top_k": 2048},
-            id="batch_4",
-        ),
-        pytest.param(
-            {"batch_size": 32, "seq_len": 5000, "top_k": 2048},
-            id="batch_32",
-        ),
-        pytest.param(
-            {"batch_size": 256, "seq_len": 5000, "top_k": 2048},
-            id="batch_256",
-        ),
-        # ==================== CATEGORY: Single-CTA vs Multi-CTA ====================
-        pytest.param(
-            {"batch_size": 2, "seq_len": 4096, "top_k": 2048},
-            id="single_cta_4k",
-        ),
-        pytest.param(
-            {"batch_size": 2, "seq_len": 8192, "top_k": 2048},
-            id="single_cta_8k",
-        ),
-        pytest.param(
-            {"batch_size": 2, "seq_len": 163840, "top_k": 2048},
-            id="multi_cta_163840_dsv3_max",
-        ),
-        # ==================== CATEGORY: Extreme Cases ====================
-        pytest.param(
-            {"batch_size": 512, "seq_len": 5000, "top_k": 2048},
-            id="extreme_large_batch",
-        ),
-        pytest.param(
-            {"batch_size": 2, "seq_len": 163840, "top_k": 2048},
-            id="extreme_dsv3_max_context",
-        ),
-    ],
-)
-@torch.inference_mode()
-def test_persistent_topk_algorithm_paths(test_config: dict) -> None:
-    """
-    Test different algorithm execution paths (capped at 163840 for DeepSeek V3.2):
-    - Batch size scalability (1, 4, 32, 256)
-    - Single-CTA vs Multi-CTA execution
-    - Extreme configurations (large batch, max context length)
-    """
-    run_large_context_topk_test(
-        batch_size=test_config["batch_size"],
-        seq_lens=[test_config["seq_len"]] * test_config["batch_size"],
-        top_k=test_config["top_k"],
-    )
-
-
-@pytest.mark.skipif(not current_platform.is_cuda(), reason="This test requires CUDA")
-@torch.inference_mode()
-def test_persistent_topk_stress() -> None:
-    """
-    Stress test with random configurations to catch edge cases.
-    Capped at 163840 (DeepSeek V3.2 max context) for realistic testing.
-    """
-    torch.set_default_device("cuda:0")
-    top_k = 2048
-
-    for seed in range(3):
-        set_random_seed(seed)
-
-        # Random batch size (limited for speed)
-        batch_size = torch.randint(1, 32, (1,)).item()
-
-        # Random sequence lengths capped at DeepSeek V3.2 max context
-        seq_lens = torch.randint(100, 163840, (batch_size,)).tolist()
-
-        run_large_context_topk_test(
-            batch_size=batch_size,
-            seq_lens=seq_lens,
-            top_k=top_k,
-            seed=seed,
-        )
-
-
-@pytest.mark.skipif(not current_platform.is_cuda(), reason="This test requires CUDA")
-@pytest.mark.parametrize(
-    "test_config",
-    [
-        # Mixed batch: rows spanning all four paths (trivial, decode, medium, large)
-        pytest.param(
-            {
-                "seq_lens": [2000, 6000, 30000, 80000],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="mixed_all_paths",
-        ),
-        # All decode/medium rows (typical decode scenario)
-        pytest.param(
-            {
-                "seq_lens": [2048, 4096, 8192, 16000],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="all_decode_medium",
-        ),
-        # All large rows
-        pytest.param(
-            {
-                "seq_lens": [70000, 100000, 163840],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="all_large",
-        ),
-        # Boundary around LARGE_THRESHOLD (32K)
-        pytest.param(
-            {
-                "seq_lens": [32767, 32768, 32769, 32772],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="large_threshold_boundary",
-        ),
-        # Single row medium
-        pytest.param(
-            {
-                "seq_lens": [5000],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="single_row_medium",
-        ),
-        # Single row large
-        pytest.param(
-            {
-                "seq_lens": [100000],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="single_row_large",
-        ),
-        # Trivial rows mixed with medium and large
-        pytest.param(
-            {
-                "seq_lens": [100, 2048, 10000, 80000],
-                "top_k": 2048,
-                "data_type": "random",
-            },
-            id="trivial_medium_large_mix",
-        ),
-    ],
-)
-@torch.inference_mode()
-def test_persistent_topk(test_config: dict) -> None:
-    """
-    Tests specific to the persistent_topk kernel:
-    - Mixed medium/large rows in the same batch (dynamic per-row dispatch)
-    - Boundary around LARGE_THRESHOLD (32K)
-    - Trivial + medium + large rows in a single batch
-    """
-    run_large_context_topk_test(
-        batch_size=len(test_config["seq_lens"]),
-        seq_lens=test_config["seq_lens"],
-        top_k=test_config["top_k"],
-        data_type=test_config.get("data_type", "random"),
-    )
-
-
-@pytest.mark.skipif(not current_platform.is_cuda(), reason="This test requires CUDA")
-@torch.inference_mode()
-def test_persistent_topk_padded_stride() -> None:
-    """
-    Test persistent_topk with padded logits (large stride, small seq_len)
-    to simulate the e2e CUDAGraph scenario where fp8_paged_mqa_logits
-    returns [B, max_model_len] with max_model_len=163840.
-    """
-    set_random_seed(42)
-    torch.set_default_device("cuda:0")
-
-    top_k = 2048
-    batch_size = 4
-    padded_stride = 163840  # DeepSeek-V3.2 max_model_len
-    actual_seq_lens = [3000, 5000, 8000, 12000]
-
-    # Create padded logits tensor (like fp8_paged_mqa_logits output)
-    logits = torch.full(
-        (batch_size, padded_stride),
-        float("-inf"),
-        dtype=torch.float32,
-        device="cuda",
-    )
-    for i, sl in enumerate(actual_seq_lens):
-        logits[i, :sl] = torch.randn(sl, dtype=torch.float32, device="cuda")
-
-    lengths = torch.tensor(actual_seq_lens, dtype=torch.int32, device="cuda")
-    indices = torch.empty((batch_size, top_k), dtype=torch.int32, device="cuda")
-    workspace = torch.empty(1024 * 1024, dtype=torch.uint8, device="cuda")
-
-    torch.ops._C.persistent_topk(
-        logits, lengths, indices, workspace, top_k, max(actual_seq_lens)
-    )
-    torch.accelerator.synchronize()
-
-    # Validate against torch.topk
-    for i in range(batch_size):
-        sl = actual_seq_lens[i]
-        k_i = min(top_k, sl)
-        expected = logits[i, :sl].topk(k_i, dim=-1)[1].cpu()
-        actual = indices[i, :k_i].cpu()
-
-        expected_set = set(expected.tolist())
-        actual_set = set(actual.tolist())
-
-        if expected_set != actual_set:
-            # Allow ties
-            expected_vals = logits[i, expected].cpu().sort(descending=True)[0]
-            actual_vals = logits[i, actual].cpu().sort(descending=True)[0]
-            assert torch.allclose(expected_vals, actual_vals, rtol=1e-4, atol=1e-4), (
-                f"Row {i}: persistent_topk with padded stride doesn't match. "
-                f"seq_len={sl}, stride={padded_stride}"
-            )
+    assert compare_top_k_results(
+        logits_long, indices, torch_indices_long, row_starts_long, row_ends_long, top_k
+    ), "large_context_topk kernel (long sequences) doesn't match torch.topk"
@@ -109,14 +109,6 @@ def _load_hf_model(model_name: str, hf_spec: dict, device: torch.device):
        **extra,
    ).to(device)
    model.eval()
-
-    # Transformers 5.0 weight materialization can clear non-persistent
-    # buffers (e.g. rotary inv_freq) that were registered with
-    # persistent=False.  Re-compute them so the model produces valid output.
-    for mod in model.modules():
-        if hasattr(mod, "_compute_inv_freq") and hasattr(mod, "inv_freq"):
-            mod.inv_freq = mod._compute_inv_freq(device=device)
-
    return model


@@ -8,13 +8,7 @@ import pytest
 from ...utils import EmbedModelInfo

 MODELS = [
-    EmbedModelInfo(
-        "nomic-ai/nomic-embed-text-v1",
-        # Fixme:
-        #  Update nomic-embed code to support the latest
-        #  HF version and remove revision set.
-        revision="720244025c1a7e15661a174c63cce63c8218e52b",
-    ),
+    EmbedModelInfo("nomic-ai/nomic-embed-text-v1"),
    # EmbedModelInfo("nomic-ai/nomic-embed-text-v1.5"),
    # EmbedModelInfo("nomic-ai/CodeRankEmbed"),
    EmbedModelInfo("nomic-ai/nomic-embed-text-v2-moe"),
@@ -30,10 +24,7 @@ max_model_len = int(original_max_position_embeddings * factor)
@pytest.mark.parametrize("model_info", MODELS)
 def test_default(model_info, vllm_runner):
    with vllm_runner(
-        model_info.name,
-        revision=model_info.revision,
-        runner="pooling",
-        max_model_len=None,
+        model_info.name, runner="pooling", max_model_len=None
    ) as vllm_model:
        model_config = vllm_model.llm.llm_engine.model_config
        if model_info.name == "nomic-ai/nomic-embed-text-v2-moe":
@@ -48,10 +39,7 @@ def test_default(model_info, vllm_runner):
 def test_set_max_model_len_legal(model_info, vllm_runner):
    # set max_model_len <= 512
    with vllm_runner(
-        model_info.name,
-        revision=model_info.revision,
-        runner="pooling",
-        max_model_len=256,
+        model_info.name, runner="pooling", max_model_len=256
    ) as vllm_model:
        model_config = vllm_model.llm.llm_engine.model_config
        assert model_config.max_model_len == 256
@@ -61,19 +49,11 @@ def test_set_max_model_len_legal(model_info, vllm_runner):
        # For nomic-embed-text-v2-moe the length is set to 512
        # by sentence_bert_config.json.
        with pytest.raises(ValueError):
-            with vllm_runner(
-                model_info.name,
-                revision=model_info.revision,
-                runner="pooling",
-                max_model_len=1024,
-            ):
+            with vllm_runner(model_info.name, runner="pooling", max_model_len=1024):
                pass
    else:
        with vllm_runner(
-            model_info.name,
-            revision=model_info.revision,
-            runner="pooling",
-            max_model_len=1024,
+            model_info.name, runner="pooling", max_model_len=1024
        ) as vllm_model:
            model_config = vllm_model.llm.llm_engine.model_config
            assert model_config.max_model_len == 1024
@@ -83,12 +63,7 @@ def test_set_max_model_len_legal(model_info, vllm_runner):
 def test_set_max_model_len_illegal(model_info, vllm_runner):
    # set max_model_len > 2048
    with pytest.raises(ValueError):
-        with vllm_runner(
-            model_info.name,
-            revision=model_info.revision,
-            runner="pooling",
-            max_model_len=4096,
-        ):
+        with vllm_runner(model_info.name, runner="pooling", max_model_len=4096):
            pass

    # set max_model_len > 2048 by hf_overrides
@@ -96,7 +71,6 @@ def test_set_max_model_len_illegal(model_info, vllm_runner):
    with pytest.raises(ValueError):
        with vllm_runner(
            model_info.name,
-            revision=model_info.revision,
            runner="pooling",
            max_model_len=None,
            hf_overrides=hf_overrides,
@@ -117,11 +91,7 @@ def test_use_rope_scaling_legal(model_info, vllm_runner):
    }

    with vllm_runner(
-        model_info.name,
-        revision=model_info.revision,
-        runner="pooling",
-        max_model_len=None,
-        hf_overrides=hf_overrides,
+        model_info.name, runner="pooling", max_model_len=None, hf_overrides=hf_overrides
    ):
        pass

@@ -140,7 +110,6 @@ def test_use_rope_scaling_illegal(model_info, vllm_runner):
    with pytest.raises(ValueError):
        with vllm_runner(
            model_info.name,
-            revision=model_info.revision,
            runner="pooling",
            max_model_len=max_model_len + 1,
            hf_overrides=hf_overrides,
@@ -160,7 +129,6 @@ def test_use_rope_scaling_illegal(model_info, vllm_runner):
    with pytest.raises(ValueError):
        with vllm_runner(
            model_info.name,
-            revision=model_info.revision,
            runner="pooling",
            max_model_len=None,
            hf_overrides=hf_overrides,
@@ -151,7 +151,6 @@ def mteb_test_embed_models(

    with vllm_runner(
        model_info.name,
-        revision=model_info.revision,
        runner="pooling",
        max_model_len=model_info.max_model_len,
        **vllm_extra_kwargs,
@@ -202,7 +201,6 @@ def mteb_test_embed_models(
    if model_info.mteb_score is None:
        with hf_runner(
            model_info.name,
-            revision=model_info.revision,
            is_sentence_transformer=True,
            dtype=ci_envs.VLLM_CI_HF_DTYPE or model_info.hf_dtype,
        ) as hf_model:
@@ -241,7 +241,6 @@ def mteb_test_rerank_models(

    with vllm_runner(
        model_info.name,
-        revision=model_info.revision,
        runner="pooling",
        max_model_len=None,
        max_num_seqs=8,
@@ -287,9 +286,7 @@ def mteb_test_rerank_models(
    # Accelerate mteb test by setting
    # SentenceTransformers mteb score to a constant
    if model_info.mteb_score is None:
-        with hf_runner(
-            model_info.name, revision=model_info.revision, dtype=model_info.hf_dtype
-        ) as hf_model:
+        with hf_runner(model_info.name, dtype=model_info.hf_dtype) as hf_model:
            hf_model.chat_template = chat_template
            st_main_score = run_mteb_rerank(
                hf_model,
@@ -12,10 +12,6 @@ MODELS = [
    EmbedModelInfo(
        "nomic-ai/nomic-embed-text-v1",
        architecture="NomicBertModel",
-        # Fixme:
-        #  Update nomic-embed code to support the latest
-        #  HF version and remove revision set.
-        revision="720244025c1a7e15661a174c63cce63c8218e52b",
        mteb_score=0.737568559,
        enable_test=True,
        seq_pooling_type="MEAN",
@@ -89,33 +89,22 @@ def test_models(example_prompts, model_name) -> None:

 EAGER = [True, False]

-SM_100_NVFP4_BACKENDS = [
-    "flashinfer-cudnn",
-    "flashinfer-trtllm",
-    "flashinfer-cutlass",
-]
-

+@pytest.mark.skipif(
+    not current_platform.has_device_capability(100),
+    reason="modelopt_fp4 is not supported on this GPU type.",
+)
@pytest.mark.parametrize("model", ["nvidia/Llama-3.1-8B-Instruct-NVFP4"])
@pytest.mark.parametrize("eager", EAGER)
@pytest.mark.parametrize(
    "backend",
    [
-        "emulation",
        "flashinfer-cudnn",
        "flashinfer-trtllm",  # the small seq_len ensures trtllm_8x4_layout backend is used
        "flashinfer-cutlass",
    ],
 )
 def test_nvfp4(vllm_runner, model, eager, backend, monkeypatch):
-    if (
-        not current_platform.has_device_capability(100)
-        and backend in SM_100_NVFP4_BACKENDS
-    ):
-        pytest.skip(
-            f"The backend {backend} is not supported with current_platform.has_device_capability(100) == False"
-        )
-
    monkeypatch.setenv("VLLM_NVFP4_GEMM_BACKEND", backend)
    with vllm_runner(model, enforce_eager=eager) as llm:
        output = llm.generate_greedy(["1 2 3 4 5"], max_tokens=2)
@@ -461,10 +461,6 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
        trust_remote_code=True,
        is_available_online=False,
    ),
-    "Param2MoEForCausalLM": _HfExamplesInfo(
-        "bharatgenai/Param2-17B-A2.4B-Thinking",
-        trust_remote_code=True,
-    ),
    "PersimmonForCausalLM": _HfExamplesInfo("adept/persimmon-8b-chat"),
    "PhiForCausalLM": _HfExamplesInfo("microsoft/phi-2"),
    "Phi3ForCausalLM": _HfExamplesInfo("microsoft/Phi-3-mini-4k-instruct"),
@@ -1250,12 +1246,6 @@ _SPECULATIVE_DECODING_EXAMPLE_MODELS = {
        use_original_num_layers=True,
        max_model_len=10240,
    ),
-    "Eagle3MiniMaxM2ForCausalLM": _HfExamplesInfo(
-        "MiniMaxAI/MiniMax-M2",
-        trust_remote_code=True,
-        speculative_model="yuhuili/EAGLE3-LLaMA3.1-Instruct-8B",
-        tokenizer="MiniMaxAI/MiniMax-M2",
-    ),
    "EagleMistralLarge3ForCausalLM": _HfExamplesInfo(
        "mistralai/Mistral-Large-3-675B-Instruct-2512",
        speculative_model="mistralai/Mistral-Large-3-675B-Instruct-2512-Eagle",
@@ -375,7 +375,6 @@ def softmax(data):
@dataclass
 class ModelInfo:
    name: str
-    revision: str | None = None
    architecture: str = ""
    dtype: str = "auto"
    max_model_len: int | None = None
@@ -366,6 +366,9 @@ def test_compressed_tensors_kv_cache_fp8_per_attn_head(vllm_runner):
        assert output


+@pytest.mark.skipif(
+    not current_platform.is_cuda(), reason="This test is skipped on non-CUDA platform."
+)
@pytest.mark.parametrize(
    "args",
    [
@@ -395,7 +398,7 @@ def test_compressed_tensors_nvfp4(vllm_runner, args):
            assert qkv_proj.scheme.group_size == 16

        llm.apply_model(check_model)
-        output = llm.generate_greedy(["Hello my name is"], max_tokens=4)
+        output = llm.generate_greedy("Hello my name is", max_tokens=4)
        print(output)
        assert output

@@ -1,179 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Tests online quantization."""
-
-import pytest
-import torch
-
-from tests.quantization.utils import (
-    _test_online_quant_peak_mem_impl,
-    is_quant_method_supported,
-)
-from vllm.model_executor.layers.linear import UnquantizedLinearMethod
-from vllm.model_executor.layers.quantization.online.fp8 import (
-    Fp8PerBlockOnlineLinearMethod,
-    Fp8PerBlockOnlineMoEMethod,
-    Fp8PerTensorOnlineLinearMethod,
-    Fp8PerTensorOnlineMoEMethod,
-)
-from vllm.platforms import current_platform
-
-
-@pytest.mark.skipif(
-    not is_quant_method_supported("fp8"),
-    reason="FP8 is not supported on this GPU type.",
-)
-@pytest.mark.parametrize(
-    "quant_scheme,online_quant_args,expected_linear_cls,expected_moe_cls",
-    [
-        # simple case - quantization='fp8_per_tensor'
-        (
-            "fp8_per_tensor",
-            None,
-            Fp8PerTensorOnlineLinearMethod,
-            Fp8PerTensorOnlineMoEMethod,
-        ),
-        # simple case - quantization='fp8_per_block'
-        (
-            "fp8_per_block",
-            None,
-            Fp8PerBlockOnlineLinearMethod,
-            Fp8PerBlockOnlineMoEMethod,
-        ),
-        # quantization='online with linear_scheme_override and
-        # moe_scheme_override
-        (
-            "online",
-            {
-                "linear_scheme_override": "fp8_per_block",
-                "moe_scheme_override": "fp8_per_tensor",
-            },
-            Fp8PerBlockOnlineLinearMethod,
-            Fp8PerTensorOnlineMoEMethod,
-        ),
-        # ignore with direct layer name
-        (
-            "fp8_per_tensor",
-            # qkv_proj is fused from q_proj/k_proj/v_proj, so currently the
-            # ignore regex must match the unfused shard names
-            # TODO(future PR): also make 're:.*qkv_proj.*' work
-            {"ignore": ["model.layers.1.self_attn.o_proj", "re:.*[qkv]_proj"]},
-            Fp8PerTensorOnlineLinearMethod,
-            Fp8PerTensorOnlineMoEMethod,
-        ),
-    ],
-)
-@pytest.mark.parametrize(
-    "use_rocm_aiter", [True, False] if current_platform.is_rocm() else [False]
-)
-def test_online_quantization(
-    vllm_runner,
-    quant_scheme: str,
-    online_quant_args: dict | None,
-    expected_linear_cls,
-    expected_moe_cls,
-    use_rocm_aiter: bool,
-    monkeypatch,
-) -> None:
-    """
-    Tests that online quantization frontend configuration works -
-    selecting quant schemes, overriding quant schemes by type, ignoring
-    layers.
-
-    Does not test performance, peak memory usage, etc.
-    """
-
-    if use_rocm_aiter:
-        monkeypatch.setenv("VLLM_ROCM_USE_AITER", "1")
-
-    # `LLM.apply_model` requires pickling a function.
-    monkeypatch.setenv("VLLM_ALLOW_INSECURE_SERIALIZATION", "1")
-
-    # a tiny model with both dense and MoE layers
-    model_name = "ibm-granite/granite-3.0-1b-a400m-base"
-
-    runner_kwargs = dict(
-        quantization=quant_scheme,
-        enforce_eager=True,
-    )
-    if online_quant_args is not None:
-        runner_kwargs["quantization_config"] = online_quant_args
-
-    with vllm_runner(
-        model_name,
-        **runner_kwargs,
-    ) as llm:
-
-        def check_model(model):
-            # checks further down in the test case are hardcoded for this
-            # model
-            assert model_name == "ibm-granite/granite-3.0-1b-a400m-base"
-
-            o_proj = model.model.layers[0].self_attn.o_proj
-            moe = model.model.layers[0].block_sparse_moe.experts
-
-            # o_proj and moe in layer 0 are always quantized (never ignored)
-            # because of how we craft the test case inputs
-            assert isinstance(o_proj.quant_method, expected_linear_cls)
-            if moe is not None:
-                assert isinstance(moe.quant_method, expected_moe_cls)
-
-            if current_platform.is_cuda():
-                assert o_proj.weight.dtype == torch.float8_e4m3fn
-            elif current_platform.is_rocm():
-                assert o_proj.weight.dtype == current_platform.fp8_dtype()
-            else:
-                pytest.skip("Only runs on CUDA and ROCm.")
-
-            # Verify ignored layers are unquantized.
-            if isinstance(online_quant_args, dict) and "ignore" in online_quant_args:
-                # only .*1.self_attn_o_proj is skipped
-                for layer_idx in range(len(model.model.layers)):
-                    o_proj = model.model.layers[layer_idx].self_attn.o_proj
-                    if layer_idx == 1:
-                        assert isinstance(o_proj.quant_method, UnquantizedLinearMethod)
-                    else:
-                        assert isinstance(o_proj.quant_method, expected_linear_cls)
-
-                # every .*self_attn.qkv_proj is skipped
-                for layer_idx in range(len(model.model.layers)):
-                    qkv_proj = model.model.layers[layer_idx].self_attn.qkv_proj
-                    assert isinstance(qkv_proj.quant_method, UnquantizedLinearMethod)
-
-        llm.apply_model(check_model)
-
-        outputs = llm.generate_greedy(["Hello my name is"], max_tokens=4)
-        print(outputs[0][1])
-
-
-@pytest.mark.skipif(
-    not is_quant_method_supported("fp8"),
-    reason="FP8 is not supported on this GPU type.",
-)
-def test_online_quant_peak_mem(
-    vllm_runner,
-    caplog_mp_spawn,
-    monkeypatch,
-) -> None:
-    _test_online_quant_peak_mem_impl(
-        "fp8_per_tensor", vllm_runner, caplog_mp_spawn, monkeypatch
-    )
-
-
-@pytest.mark.skipif(
-    not is_quant_method_supported("fp8"),
-    reason="FP8 is not supported on this GPU type.",
-)
-def test_online_quant_load_format_dummy(
-    vllm_runner,
-    monkeypatch,
-    caplog,
-) -> None:
-    with vllm_runner(
-        "ibm-granite/granite-3.0-1b-a400m-base",
-        quantization="fp8_per_tensor",
-        enforce_eager=True,
-        load_format="dummy",
-    ) as llm:
-        outputs = llm.generate_greedy(["The future of AI is"], max_tokens=4)
-        print(outputs[0][1])
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Wentao YeandGitHub	beefff2174	Merge branch 'main' into wentao-skip-work-when-empty	2026-04-03 11:31:23 -04:00
yewentao256	163266d0b2	refactor only inside nixl Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-04-01 18:28:13 +00:00
yewentao256	a2fd28a7e1	Merge branch 'main' into wentao-skip-work-when-empty Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-04-01 18:09:06 +00:00
Wentao YeandGitHub	7a80ac928f	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-31 13:01:59 -04:00
yewentao256	0ed11013b4	check empty using kv connector's method Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-03-30 21:06:12 +00:00
yewentao256	6d568b995a	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-30 19:17:12 +00:00
Wentao YeandGitHub	dfe9decbcb	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-28 10:28:43 -04:00
Wentao YeandGitHub	1324e6ff67	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-27 15:52:55 -04:00
Wentao YeandGitHub	c1aba6d7ae	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-27 09:42:19 -04:00
Wentao Ye GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	200bef28c9	Update vllm/v1/worker/gpu/kv_connector.py Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>	2026-03-26 18:04:12 -04:00
yewentao256	fd9820bbf9	skip kv connector empty work Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-03-26 21:50:28 +00:00