Merge branch 'main' into patches

Signed-off-by: MOVZX <movzx@yahoo.com>

.agents/skills/model-microbenchmark/SKILL.md

@@ -0,0 +1,72 @@
+---
+name: model-microbenchmark
+description: Time a vLLM model's forward pass on exact batch shapes (decode, prefill, chunked prefill, prefix hit, spec-decode verify, or mixes) with `vllm bench forward`, single-GPU or DP/TP/EP across nodes. Use for per-step model latency, layout/parallelism comparisons, and checking a model-level optimization. Not for single kernels (use kernel-microbenchmark) or end-to-end serving (use `vllm bench serve`).
+---
+
+# Model Microbenchmark (`vllm bench forward`)
+
+Builds the batch without the scheduler, runs it through MRV2's real
+`execute_model`, and times the forward as a one-shot CUDA graph (prefill
+included). Sampling, the speculator, and CPU step overhead are excluded, so
+expect results somewhat below a serving step.
+
+## Usage
+
+```bash
+vllm bench forward --model deepseek-ai/DeepSeek-V4-Flash \
+  --case decode=64*1@8192 --case prefill=8192 --case mix=4096,32*1@8192
+
+# DP/TP/EP: one process per GPU (multi-node: add --nnodes/--node-rank/--master-addr)
+torchrun --nproc-per-node=4 -m vllm.entrypoints.cli.main bench forward \
+  --model ... -dp 4 --enable-expert-parallel --case decode=128*1@100000
+```
+
+Requests: `[N*]Q[@K][c][:D]`, comma-separated, per DP rank.
+
+| Form | Meaning |
+| --- | --- |
+| `8192` | fresh prefill, 8192 tokens |
+| `32*1@8192` | 32 decodes, KV length 8192 (KV length **includes** the query) |
+| `4@4096:3` | spec-decode verify, 3 draft tokens |
+| `2048@10240` | prefill after an 8192-token prefix-cache hit |
+| `2048@10240c` | chunked prefill: the prefix came from earlier chunks |
+
+Hit vs chunk only differs on models with SWA bounded replay (DSV4.1): a hit
+also recomputes the last window (128 tokens) of the prefix.
+
+Or `--shapes file.json|yaml`:
+`{"cases": [{"name": ..., "requests": [{"query_len", "kv_len", "num_draft_tokens", "chunked", "count"}]}]}`.
+
+Key flags (plus all `EngineArgs`):
+
+- `--kv-len-jitter 0.5` (default): prefixes drawn within ±50%, same on every
+  DP rank, capped at `max_model_len`. Use `0` for exact lengths.
+- `--routing auto|real|balanced`: `auto` = balanced with dummy weights.
+- `--iters 20 --samples 5`, `--output-json out.json` (per-rank latency stats,
+  `samples_us`, graph-vs-eager check).
+- Defaults: `--load-format dummy`, prefix caching off, no torch.compile.
+
+Read `median_us` (each replay is already the max over ranks).
+
+## Caveats
+
+- **Dummy weights lie for data-dependent kernels.** Dummy gates route every
+  token to the same experts (hence balanced routing). Cross-check suspicious
+  kernels with real weights.
+- **torch.compile models (e.g. Qwen3/3.5) look slow**: no compile here. Use
+  them for correctness only; trust perf for models under `vllm/models/`.
+- `sample spread > 0.5%` warns of unstable timing. DSV4.1 DEP prefill is
+  inherently ~0.5–1% (MegaMoE cross-rank sync, host engram reads).
+- `graph_vs_eager` may be nonzero for small batches where models branch on
+  `is_current_stream_capturing()`; that's expected, not a bug.
+- KV cache is zero-initialized (`--kv-init zeros`); random KV was within 3%.
+- Uniform DP only: every rank gets the same case; no idle ranks.
+- CUDA only: it hangs on ROCm (untested cause: likely the in-graph timing
+  events).
+- Large prefills/prefix hits need room: raise `--max-num-batched-tokens` /
+  `--max-model-len` / `--gpu-memory-utilization` if build fails on KV blocks.
+- Startup cost is JIT/autotune, not the bench: install FlashInfer jit-cache
+  and cubin; DeepGEMM warmup is skipped (each case's eager pass compiles its
+  own shapes).
+- To compare against serving decode: `vllm bench latency` with 65- vs
+  1-token outputs, divided by 64, with a FULL decode graph.

.agents/skills/transformers-upgrade/SKILL.md

@@ -121,6 +121,9 @@ For each deleted config:
 Deleted processors follow the same steps in
 `vllm/transformers_utils/processors/__init__.py`.
 
+Before merging, rebase and grep again for the old field names: PRs that land
+in between still read them.
+
 Done when every line of the inventory output is deleted, or kept with a
 reason that meets the bar in step 2. The PR also carries the config diff
 results and evals for any changed values.
@@ -128,13 +131,15 @@ results and evals for any changed values.
 ## Upstreaming hazards
 
 Upstream configs are `@strict` dataclasses. Checkpoint keys upstream does not
-declare survive as plain attributes. Defaults that only vLLM sets, and values
+declare survive as plain attributes, but only when the checkpoint sets them.
-vLLM derives, vanish.
+Defaults that only vLLM sets, and values vLLM derives, vanish.
 
 | Hazard | Signal | Fix |
 | --- | --- | --- |
 | Upstream consumes legacy keys | `__post_init__` pops keys the model reads (`deepseek_v4` `compress_ratios`) | Read the upstream replacement (`layer_types`, `mlp_layer_types`, per-layer `rope_parameters`) |
 | Renamed or dropped fields | Diff shows `<MISSING>` for a field the model reads | Read the upstream name; derive dropped flags from what implies them (`enable_moe_block` from `num_experts`) |
+| Undeclared field without a default | `config.x` raises `AttributeError` for configs that omit `x`, e.g. minimal test configs (A.X-K1 `scoring_func`). Dumps of real checkpoints miss it | Upstream ignoring `x` usually means vLLM can stop reading it; if vLLM truly needs it, read `getattr(config, "x", <old default>)` |
+| Top-level config stops forwarding | The vendored top-level config forwarded reads to `text_config`; upstream does not. `getattr(config, "x", default)` silently takes the default, e.g. on MTP draft configs from `SpeculativeConfig.hf_config_override`, which the dumps do not build | Read from `get_text_config()`; grep `getattr` reads of every field, including speculative-decoding paths |
 | Type check on the config | `ctx.get_hf_config(SomeConfig)` relies on the vendored base class | Accept the upstream class too, e.g. a `Cosmos3ProcessingInfo` overriding `get_hf_config` |
 | Remote code takes over | Checkpoint `config.json` has `auto_map`; with `--trust-remote-code`, the hub class replaces the native one once vLLM stops registering its own | Check the remote class loads; if not, register the upstream class for that `model_type` |
 | Read-only upstream fields | Properties or derived fields break subclasses and `hf_overrides` (`nemotron_h` `hybrid_override_pattern`) | Override the source field (`layers_block_type`) instead |
@@ -159,10 +164,10 @@ Version gates from the inventory need a decision:
 - **Patches citing a Transformers issue or PR:**
     - Delete the patch when its fix is in the floor release.
     - For an issue, find the PR that closed it and check that PR. A closed
-      issue is not a fixed one: #43329 was closed while the bug was still in
+      issue is not a fixed one: issues get closed while the bug is still in
-      5.16.1.
+      the release.
     - Re-run the patched behaviour before deleting. A fix in the release does
-      not always cover vLLM's path: #47924 made Emu3's processor add BOS, but
+      not always cover vLLM's path: Emu3's processor gained BOS upstream, but
       vLLM tokenizes without the processor.
     - Links that only explain a design choice (e.g. a float32 RoPE note) are
       not patches; leave them.
@@ -176,3 +181,11 @@ Version gates from the inventory need a decision:
 
 Add `tests/models/test_initialization.py -k "<archs>"` where a GPU is
 available, and a model eval for any model whose resolved config changed.
+
+Models in `vllm/models/` are high priority: whenever their config or model
+code changes, run GSM8k (`tests/evals/gsm8k/`) and GPQA Diamond
+(`tests/evals/gpt_oss/`) on `main` and the branch, even when the dumps show
+no diff, and put both scores in the PR. FlashInfer autotuning picks kernel
+tactics afresh on every server start, which moved DeepSeek-V4-Flash GSM8k by
+about 1 point; pass `--kernel-config '{"enable_flashinfer_autotune": false}'`
+to both servers so greedy outputs are comparable.

.buildkite/hardware_tests/ascend_npu.yaml

@@ -56,7 +56,6 @@ steps:
       - "!vllm/distributed/device_communicators/shm_object_storage.py"
       - "!vllm/distributed/device_communicators/symm_mem.py"
       - "!vllm/distributed/device_communicators/mnnvl_compat.py"
-      - "!vllm/distributed/device_communicators/ray_communicator.py"
       - "!vllm/third_party/flashmla/"
       - "!vllm/third_party/pynvml.py"
       - "!vllm/_custom_ops.py"

.buildkite/hardware_tests/cpu.yaml

@@ -277,7 +277,7 @@ steps:
     # `uv pip install "vllm[zen]"`, which forces uv to re-resolve vLLM's full
     # dependency closure and chokes on its transitive triton-cpu URL dependency.
     # Keep this pin in sync with the `zen` extra in setup.py.
-    ZENTORCH_VERSION: "2.13.0.1"
+    ZENTORCH_VERSION: "2.14.0.1"
   retry:
     automatic:
       - exit_status: 1  # Transient Docker/BuildKit failure

.buildkite/hardware_tests/intel_xpu_ci/test-intel.yaml

@@ -15,7 +15,7 @@ steps:
           limit: 2
         - exit_status: -10  # Agent was lost
           limit: 2
-  - label: "XPU example Test"
+  - label: ":intel: (BMG) Example Test"
     key: xpu-example-test
     depends_on:
       - image-build-xpu
@@ -37,7 +37,7 @@ steps:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh example'
-  - label: "XPU V1 test"
+  - label: ":intel: (BMG) V1 test"
     key: xpu-v1-test
     depends_on:
       - image-build-xpu
@@ -60,7 +60,7 @@ steps:
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh v1'
     parallelism: 2
-  - label: "XPU W8A8 FP8 Linear Examples"
+  - label: ":intel: (BMG) W8A8 FP8 Linear Examples"
     key: xpu-w8a8-fp8-linear-examples
     depends_on:
       - image-build-xpu
@@ -83,7 +83,7 @@ steps:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh w8a8-fp8-linear'
-  - label: "XPU server test"
+  - label: ":intel: (BMG) Server test"
     key: xpu-server-test
     depends_on:
       - image-build-xpu
@@ -105,7 +105,7 @@ steps:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh server'
-  - label: "XPU quantization test"
+  - label: ":intel: (BMG) Quantization test"
     key: xpu-quantization-test
     depends_on:
       - image-build-xpu
@@ -127,7 +127,7 @@ steps:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh quantization'
-  - label: "XPU compressed tensors FP8 test"
+  - label: ":intel: (BMG) Compressed tensors FP8 test"
     key: xpu-compressed-tensors-fp8-test
     depends_on:
       - image-build-xpu
@@ -150,7 +150,7 @@ steps:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh compressed-tensors-fp8'
-  - label: "XPU Graph"
+  - label: ":intel: (BMG) Graph"
     key: xpu-graph
     depends_on:
       - image-build-xpu

.buildkite/intel_jobs/basic_correctness_intel.yaml

@@ -1,8 +1,8 @@
-group: Basic Correctness
+group: Basic Correctness Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: XPU Sleep Mode
+- label: ":intel: (BMG) Sleep Mode"
   key: xpu-sleep-mode
   timeout_in_minutes: 45
   device: intel_gpu

.buildkite/intel_jobs/benchmarks_intel.yaml

@@ -1,8 +1,8 @@
-group: Benchmarks
+group: Benchmarks Intel
 depends_on: 
   - image-build-xpu
 steps:
-- label: Benchmarks CLI Test
+- label: ":intel: (BMG) Benchmarks CLI"
   key: benchmarks-cli-test
   timeout_in_minutes: 40
   device: intel_gpu

.buildkite/intel_jobs/distributed_intel.yaml

@@ -0,0 +1,33 @@
+group: Distributed Intel
+depends_on:
+  - image-build-xpu
+steps:
+- label: ":intel: (BMG) Distributed Torchrun + Shutdown"
+  key: distributed-torchrun-shutdown-tests-2-gpus
+  timeout_in_minutes: 50
+  device: intel_gpu
+  agent_tags:
+    label: production
+    gpu: 2+
+    mem: 24+
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/distributed/
+  - "!vllm/distributed/kv_transfer/"
+  - vllm/engine/
+  - vllm/executor/
+  - vllm/worker/worker_base.py
+  - vllm/v1/engine/
+  - vllm/v1/worker/
+  - tests/distributed/
+  - tests/v1/shutdown
+  - tests/v1/worker/test_worker_memory_snapshot.py
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'cd tests && pytest -v -s v1/shutdown'

.buildkite/intel_jobs/engine_intel.yaml

@@ -2,7 +2,7 @@ group: Engine Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: Engine
+- label: ":intel: (BMG) Engine"
   key: engine
   timeout_in_minutes: 40
   device: intel_gpu
@@ -40,7 +40,7 @@ steps:
       'cd tests &&
       pytest -v -s engine/test_arg_utils.py test_sequence.py test_logger.py test_vllm_port.py jit_monitor/test_hooks.py'
 
-- label: Engine (1 GPU)
+- label: ":intel: (BMG) Engine (1 GPU)"
   key: engine-1-gpu
   timeout_in_minutes: 30
   device: intel_gpu
@@ -65,7 +65,7 @@ steps:
       pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py --ignore v1/engine/test_core_engine_actor_manager.py &&
       pytest -v -s test_config.py'
 
-- label: V1 e2e (2 GPUs)
+- label: ":intel: (BMG) V1 e2e (2 GPUs)"
   key: v1-e2e-2-gpus
   timeout_in_minutes: 30
   device: intel_gpu

.buildkite/intel_jobs/entrypoints_intel.yaml

@@ -1,4 +1,4 @@
-group: Entrypoints
+group: Entrypoints Intel
 depends_on:
   - image-build-xpu
 steps:

.buildkite/intel_jobs/expert_parallelism_intel.yaml

@@ -1,8 +1,8 @@
-group: Expert Parallelism
+group: Expert Parallelism Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: EPLB Algorithm
+- label: ":intel: (BMG) EPLB Algorithm"
   key: eplb-algorithm
   timeout_in_minutes: 45
   device: intel_gpu

.buildkite/intel_jobs/kernels_intel.yaml

@@ -2,7 +2,7 @@ group: Kernels Intel
 depends_on: 
   - image-build-xpu
 steps:
-- label: vLLM IR Tests
+- label: ":intel: (BMG) vLLM IR Tests"
   key: vllm-ir-tests
   timeout_in_minutes: 30
   device: intel_gpu
@@ -26,7 +26,7 @@ steps:
       pytest -v -s ir &&
       pytest -v -s kernels/ir'
 
-- label: UVA Kernel Tests
+- label: ":intel: (BMG) UVA Kernel Tests"
   key: uva-kernel-tests
   timeout_in_minutes: 20
   device: intel_gpu

.buildkite/intel_jobs/lm_eval_intel.yaml

@@ -1,8 +1,8 @@
-group: LM Eval
+group: LM Eval Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: "XPU GPQA Eval (GPT-OSS)"
+- label: ":intel: (BMG) GPQA Eval (GPT-OSS)"
   key: xpu-gpqa-eval-gpt-oss
   timeout_in_minutes: 60
   device: intel_gpu
@@ -25,7 +25,7 @@ steps:
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
       pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-xpu.txt'
-- label: LM Eval Large Models(2 GPUs)
+- label: ":intel: (BMG) LM Eval Large Models(2 GPUs)"
   key: lm-eval-large-models-2-gpus
   timeout_in_minutes: 45
   device: intel_gpu

.buildkite/intel_jobs/lora_intel.yaml

@@ -2,7 +2,7 @@ group: LoRA Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: LoRA Runtime + Utils
+- label: ":intel: (BMG) LoRA Runtime + Utils"
   key: lora-runtime-utils
   timeout_in_minutes: 45
   device: intel_gpu
@@ -35,7 +35,7 @@ steps:
       pytest -v -s lora/test_utils.py &&
       pytest -v -s lora/test_worker.py'
 
-- label: LoRA Fused/MoE Kernels
+- label: ":intel: (BMG) LoRA Fused/MoE Kernels"
   key: lora-fused-moe-kernels
   timeout_in_minutes: 45
   device: intel_gpu
@@ -60,7 +60,7 @@ steps:
       pytest -v -s lora/test_fused_moe_lora_kernel.py && 
       pytest -v -s lora/test_moe_lora_align_sum.py --deselect="tests/lora/test_moe_lora_align_sum.py::test_moe_lora_align_block_size_mixed_base_and_lora[1]"'
 
-- label: LoRA Punica Kernels
+- label: ":intel: (BMG) LoRA Punica Kernels"
   key: lora-punica-kernels
   timeout_in_minutes: 45
   device: intel_gpu
@@ -87,7 +87,7 @@ steps:
       pytest -v -s lora/test_punica_ops.py::test_kernels_hidden_size &&
       pytest -v -s lora/test_punica_ops.py::test_add_lora_fused_moe_early_exit'
 
-- label: LoRA Punica FP8/XPU Ops
+- label: ":intel: (BMG) LoRA Punica FP8/XPU Ops"
   key: lora-punica-fp8-xpu-ops
   timeout_in_minutes: 60
   device: intel_gpu
@@ -112,7 +112,7 @@ steps:
       pytest -v -s lora/test_punica_ops_fp8.py &&
       pytest -v -s lora/test_punica_xpu_ops.py'
 
-- label: LoRA Models
+- label: ":intel: (BMG) LoRA Models"
   key: lora-models
   timeout_in_minutes: 45
   device: intel_gpu
@@ -139,7 +139,7 @@ steps:
       pytest -v -s lora/test_chatglm3_tp.py &&
       pytest -v -s lora/test_llama_tp.py::test_llama_lora'
 
-- label: LoRA Multimodal
+- label: ":intel: (BMG) LoRA Multimodal"
   key: lora-multimodal
   timeout_in_minutes: 45
   device: intel_gpu
@@ -163,7 +163,7 @@ steps:
       export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_whisper.py'
 
-- label: LoRA Add + Multimodal
+- label: ":intel: (BMG) LoRA Add + Multimodal"
   key: lora-add-multimodal
   timeout_in_minutes: 45
   device: intel_gpu

.buildkite/intel_jobs/misc_intel.yaml

@@ -2,7 +2,7 @@ group: Miscellaneous Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: XPU Batch Invariance
+- label: ":intel: (BMG) Batch Invariance"
   key: xpu-batch-invariance
   timeout_in_minutes: 60
   num_devices: 2
@@ -30,7 +30,7 @@ steps:
       pytest -v -s v1/determinism/test_xpu_batch_invariant_ut.py &&
       bash ../.buildkite/scripts/hardware_ci/run-xpu-batch-invariance.sh'
 
-- label: V1 Core + KV + Metrics
+- label: ":intel: (BMG) V1 Core + KV + Metrics"
   key: v1-core-kv-metrics-intel
   timeout_in_minutes: 45
   device: intel_gpu
@@ -64,7 +64,7 @@ steps:
       VLLM_BATCH_INVARIANT=1 pytest -v -s -m "not cpu_test" v1/ec_connector/unit
       --ignore=v1/ec_connector/unit/test_ec_cpu_connector.py'
 
-- label: EC CPU Connector (XPU)
+- label: ":intel: (BMG) EC CPU Connector (XPU)"
   key: ec-cpu-connector-xpu
   timeout_in_minutes: 45
   device: intel_gpu
@@ -96,7 +96,7 @@ steps:
       cd tests &&
       pytest -v -s v1/ec_connector/unit/test_ec_cpu_connector.py'
 
-- label: V1 Sample + Logits
+- label: ":intel: (BMG) V1 Sample + Logits"
   key: v1-sample-logits-intel
   timeout_in_minutes: 90
   device: intel_gpu
@@ -141,7 +141,7 @@ steps:
         limit: 1
   parallelism: 4
 
-- label: Basic Models Tests (Initialization)
+- label: ":intel: (BMG) Basic Models Tests (Initialization)"
   key: basic-models-tests-initialization
   timeout_in_minutes: 60
   device: intel_gpu
@@ -166,7 +166,7 @@ steps:
       cd tests &&
       pytest -v -s models/test_initialization.py::test_can_initialize_large_subset[Eagle3MiniMaxM2ForCausalLM]'
 
-- label: XPU CPU Offload
+- label: ":intel: (BMG) CPU Offload"
   key: xpu-cpu-offload
   timeout_in_minutes: 60
   device: intel_gpu
@@ -194,7 +194,7 @@ steps:
       pytest -v -s v1/kv_offload &&
       pytest -v -s v1/kv_connector/unit/test_offloading_connector.py'
 
-- label: NixlConnector PD accuracy (4 GPUs)
+- label: ":intel: (BMG) NixlConnector PD accuracy (4 GPUs)"
   key: nixlconnector-pd-accuracy-4-gpus
   timeout_in_minutes: 60
   num_devices: 4
@@ -223,7 +223,7 @@ steps:
       PREFILLER_TP_SIZE=1 DECODER_TP_SIZE=2 bash v1/kv_connector/nixl_integration/run_xpu_disagg_accuracy_test.sh &&
       PREFILLER_TP_SIZE=2 DECODER_TP_SIZE=2 bash v1/kv_connector/nixl_integration/run_xpu_disagg_accuracy_test.sh'
 
-- label: Regression
+- label: ":intel: (BMG) Regression"
   key: regression
   timeout_in_minutes: 50
   device: intel_gpu
@@ -256,7 +256,7 @@ steps:
       'cd tests &&
       pytest -v -s test_regression.py'
 
-- label: Metrics, Tracing (2 GPUs)
+- label: ":intel: (BMG) Metrics, Tracing (2 GPUs)"
   key: metrics-tracing-2-gpus
   timeout_in_minutes: 45
   num_devices: 2
@@ -291,7 +291,7 @@ steps:
     'cd tests &&
     pytest -v -s v1/tracing'
 
-- label: Async Engine, Inputs, Utils, Worker
+- label: ":intel: (BMG) Async Engine, Inputs, Utils, Worker"
   key: async-engine-inputs-utils-worker
   timeout_in_minutes: 55
   device: intel_gpu
@@ -330,7 +330,7 @@ steps:
     pytest -v -s -m "not cpu_test" ./multimodal &&
     pytest -v -s utils_ --ignore=utils_/test_mem_utils.py'
 
-- label: Detokenizer Stop Strings (XPU)
+- label: ":intel: (BMG) Detokenizer Stop Strings (XPU)"
   key: detokenizer-stop-strings-xpu
   timeout_in_minutes: 45
   device: intel_gpu
@@ -367,7 +367,7 @@ steps:
       'cd tests &&
       pytest -v -s detokenizer/test_stop_strings.py'
 
-- label: Fusion Unit Tests
+- label: ":intel: (BMG) Fusion Unit Tests"
   key: fusion-unit-tests
   timeout_in_minutes: 30
   device: intel_gpu

.buildkite/intel_jobs/model_executor_intel.yaml

@@ -2,7 +2,7 @@ group: Model Executor Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: Model Executor (Intel)
+- label: ":intel: (BMG) Model Executor"
   key: model-executor-intel
   timeout_in_minutes: 45
   device: intel_gpu
@@ -28,9 +28,9 @@ steps:
       export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       export PYTHONFAULTHANDLER=1 &&
       cd tests &&
-      pytest -v -s model_executor -m "not slow_test" --ignore="model_executor/layers/test_rocm_unquantized_gemm.py" --ignore="model_executor/model_loader/tensorizer_loader/test_tensorizer.py" --deselect="tests/model_executor/model_loader/test_reload.py::test_kv_scale_reload"'
+      pytest -v -s model_executor -m "not slow_test" --ignore="model_executor/layers/test_rocm_unquantized_gemm.py" --ignore="model_executor/model_loader/tensorizer_loader/test_tensorizer.py" --ignore="model_executor/model_loader/test_reload.py"'
 
-- label: Tensorizer Model Loader Tests (XPU)
+- label: ":intel: (BMG) Tensorizer Model Loader Tests (XPU)"
   key: tensorizer-model-loader-tests-xpu
   timeout_in_minutes: 45
   device: intel_gpu

.buildkite/intel_jobs/models_distributed_intel.yaml

@@ -1,8 +1,8 @@
-group: Models - Distributed
+group: Models - Distributed Intel
 depends_on: 
   - image-build-xpu
 steps:
-- label: Distributed Model Tests (2 GPUs)
+- label: ":intel: (BMG) Distributed Model Tests"
   key: distributed-model-tests-2-gpus
   timeout_in_minutes: 65
   device: intel_gpu

.buildkite/intel_jobs/models_multimodal_intel.yaml

@@ -1,8 +1,8 @@
-group: Models - Multimodal
+group: Models - Multimodal Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: "Multi-Modal Models (Standard) 1: qwen2"
+- label: ":intel: (BMG) Multi-Modal Models (Standard) 1: qwen2"
   key: multi-modal-models-standard-1-qwen2
   timeout_in_minutes: 70
   device: intel_gpu
@@ -26,7 +26,7 @@ steps:
       pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2" &&
       pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model'
 
-- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
+- label: ":intel: (BMG) Multi-Modal Models (Standard) 2: qwen3 + gemma"
   key: multi-modal-models-standard-2-qwen3-gemma
   timeout_in_minutes: 90
   device: intel_gpu
@@ -51,7 +51,7 @@ steps:
       pytest -v -s models/multimodal/generation/test_mm_prefix_lm.py -m core_model &&
       pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model'
 
-- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
+- label: ":intel: (BMG) Multi-Modal Models (Standard) 3: llava + qwen2_vl"
   key: multi-modal-models-standard-3-llava-qwen2-vl
   timeout_in_minutes: 65
   device: intel_gpu
@@ -76,7 +76,7 @@ steps:
       pytest -v -s models/multimodal/generation/test_gemma_ragged_audio.py -m core_model
       pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model'
 
-- label: "Multi-Modal Models (Standard) 4: other + whisper"
+- label: ":intel: (BMG) Multi-Modal Models (Standard) 4: other + whisper"
   key: multi-modal-models-standard-4-other-whisper
   timeout_in_minutes: 45
   device: intel_gpu
@@ -99,7 +99,7 @@ steps:
       'cd tests &&
       pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_gemma_ragged_audio.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing'
 
-- label: Multi-Modal Processor # 44min
+- label: ":intel: (BMG) Multi-Modal Processor" # 44min
   key: multi-modal-processor
   timeout_in_minutes: 60
   device: intel_gpu
@@ -126,7 +126,7 @@ steps:
       --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB'
   parallelism: 4
 
-- label: Multi-Modal Models (Extended Generation 1)
+- label: ":intel: (BMG) Multi-Modal Models (Extended Generation 1)"
   key: multi-modal-models-extended-generation-1
   timeout_in_minutes: 60
   device: intel_gpu

.buildkite/intel_jobs/quantization_intel.yaml

@@ -1,8 +1,8 @@
-group: Quantization
+group: Quantization Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: Quantization
+- label: ":intel: (BMG) Quantization"
   key: quantization
   timeout_in_minutes: 30
   env:

.buildkite/intel_jobs/samplers_intel.yaml

@@ -2,7 +2,7 @@ group: Samplers Intel
 depends_on:
   - image-build-xpu
 steps:
-- label: Samplers Test (XPU Kernel)
+- label: ":intel: (BMG) Samplers Test (XPU Kernel)"
   key: samplers-test-xpu-kernel-intel
   timeout_in_minutes: 40
   device: intel_gpu
@@ -29,7 +29,7 @@ steps:
       VLLM_XPU_USE_SAMPLER_KERNEL=1 pytest -v -s samplers
       --ignore=samplers/test_beam_search.py'
 
-- label: Samplers Beam Search Test (XPU Kernel)
+- label: ":intel: (BMG) Samplers Beam Search Test (XPU Kernel)"
   key: samplers-beam-search-test-xpu-kernel-intel
   timeout_in_minutes: 40
   device: intel_gpu

.buildkite/intel_jobs/test-intel.yaml

@@ -14,7 +14,7 @@ steps:
           limit: 2
         - exit_status: -10  # Agent was lost
           limit: 2
-  - label: "XPU example Test"
+  - label: ":intel: (BMG) Example Test"
     key: xpu-example-test
     depends_on:
       - image-build-xpu
@@ -30,7 +30,7 @@ steps:
       REPO: "vllm-ci-test-repo"
     source_file_dependencies:
       - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml 
+      - .buildkite/intel_jobs/test-intel.yaml
     commands:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
@@ -48,7 +48,7 @@ steps:
         python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192 &&
         python3 examples/basic/offline_inference/generate.py --model INCModel/Qwen3-30B-A3B-Instruct-2507-MXFP4-LLMC -tp 2 --max-model-len 8192
         '
-  - label: "XPU W8A8 FP8 Linear Examples"
+  - label: ":intel: (BMG) W8A8 FP8 Linear Examples"
     key: xpu-w8a8-fp8-linear-examples
     depends_on:
       - image-build-xpu
@@ -76,7 +76,7 @@ steps:
         python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model meta-llama/Llama-3.2-1B-Instruct --quantization fp8_per_tensor --enforce-eager --max-model-len 4096 &&
         python3 examples/basic/offline_inference/generate.py --linear-backend torch --model meta-llama/Llama-3.2-1B-Instruct --quantization fp8_per_tensor --enforce-eager --max-model-len 4096
         '
-  - label: "XPU V1 test"
+  - label: ":intel: (BMG) V1 test"
     key: xpu-v1-test
     depends_on:
       - image-build-xpu
@@ -92,7 +92,7 @@ steps:
       REPO: "vllm-ci-test-repo"
     source_file_dependencies:
       - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml 
+      - .buildkite/intel_jobs/test-intel.yaml
     commands:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
@@ -105,9 +105,9 @@ steps:
         pytest -v -s v1/test_serial_utils.py &&
         pytest -v -s v1/e2e/general/test_correctness_sliding_window.py &&
         pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py --ignore=v1/spec_decode/test_speculators_correctness.py --deselect=tests/v1/spec_decode/test_mtp.py::test_glm_mtp_defers_lm_head &&
-        pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py --ignore=v1/kv_connector/unit/test_nixl_connector_hma.py --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --deselect=tests/v1/kv_connector/unit/test_nixl_connector.py::test_abort_timeout_on_prefiller[ray]'
+        pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --deselect=tests/v1/kv_connector/unit/test_nixl_connector.py::test_abort_timeout_on_prefiller[ray]'
     parallelism: 2
-  - label: "XPU Qwen3-30B + V1 KV Connector"
+  - label: ":intel: (BMG) Qwen3-30B + V1 KV Connector"
     key: xpu-qwen3-30b-v1-kv-connector
     depends_on:
       - image-build-xpu
@@ -133,7 +133,7 @@ steps:
         pytest -v -s v1/kv_connector/unit/test_example_connector.py &&
         cd .. &&
         VLLM_XPU_FUSED_MOE_USE_REF=1 python3 examples/basic/offline_inference/generate.py --model Qwen/Qwen3-30B-A3B-Instruct-2507-FP8 --enforce-eager -tp 2 --max-model-len 8192'
-  - label: "XPU server test"
+  - label: ":intel: (BMG) Server test"
     key: xpu-server-test
     depends_on:
       - image-build-xpu
@@ -149,14 +149,14 @@ steps:
       REPO: "vllm-ci-test-repo"
     source_file_dependencies:
       - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml 
+      - .buildkite/intel_jobs/test-intel.yaml
     commands:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'cd tests &&
         pytest -v -s entrypoints/multimodal/openai/chat_completion/test_audio_in_video.py &&
         pytest -v -s benchmarks/test_serve_cli.py'
-  - label: "XPU quantization test"
+  - label: ":intel: (BMG) Quantization test"
     key: xpu-quantization-test
     depends_on:
       - image-build-xpu
@@ -172,14 +172,14 @@ steps:
       REPO: "vllm-ci-test-repo"
     source_file_dependencies:
       - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml 
+      - .buildkite/intel_jobs/test-intel.yaml
     commands:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'cd tests &&
         pytest -v -s quantization/test_auto_round.py &&
         pytest -v -s quantization/test_online.py --deselect=tests/quantization/test_online.py::test_online_quantization_loads_real_weights --deselect="tests/quantization/test_online.py::test_online_quantization[False-requantization_mxfp8_ptcp_fp8]"'
-  - label: "XPU compressed tensors FP8 test"
+  - label: ":intel: (BMG) Compressed tensors FP8 test"
     key: xpu-compressed-tensors-fp8-test
     depends_on:
       - image-build-xpu
@@ -204,7 +204,7 @@ steps:
         'cd tests &&
         pytest -v -s quantization/test_compressed_tensors.py::test_compressed_tensors_fp8 &&
         pytest -v -s quantization/test_compressed_tensors.py::test_compressed_tensors_fp8_block_enabled'
-  - label: "XPU Graph"
+  - label: ":intel: (BMG) Graph"
     key: xpu-graph
     depends_on:
       - image-build-xpu

.buildkite/scripts/crcr-report.sh

@@ -227,7 +227,12 @@ SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
 # five-minute lifetime and the loop above can run to the build-age deadline.
 # Build 91827 waited 6h48m and then took a 401 on all 360 callbacks from a
 # token that had expired hours earlier.
-OIDC_TOKEN="$(buildkite-agent oidc request-token --skip-redaction --audience "${AUDIENCE}" 2>/dev/null)"
+# Posting is paced under the relay's rate limit, which takes several minutes
+# for a full nightly, so ask for a longer lifetime than the default. Fall back
+# to the default if this agent rejects --lifetime.
+OIDC_TOKEN="$(buildkite-agent oidc request-token --skip-redaction --audience "${AUDIENCE}" \
+    --lifetime 1800 2>/dev/null \
+    || buildkite-agent oidc request-token --skip-redaction --audience "${AUDIENCE}" 2>/dev/null)"
 if [[ -z "${OIDC_TOKEN}" ]]; then
     echo "could not mint a Buildkite OIDC token after the wait -- skipping report"
     exit 0

.buildkite/scripts/crcr_report.py

@@ -24,7 +24,9 @@ from __future__ import annotations
 
 import argparse
 import json
+import os
 import sys
+import time
 import urllib.error
 import urllib.request
 
@@ -43,6 +45,14 @@ _CONCLUSION = {
 }
 
 
+# The relay rate-limits callbacks per repo over a sliding 60 s window
+# (RATE_LIMIT_PER_MIN in pytorch/test-infra), and a nightly has far more jobs
+# than that. Pace below the limit, and back off on a 429 for about a minute in
+# total so the last retry lands in a new window.
+_POST_INTERVAL_S = float(os.environ.get("CRCR_POST_INTERVAL_S", "0.65"))
+_RATE_LIMIT_BACKOFF_S = (2, 4, 8, 16, 32)
+
+
 def post(url: str, token: str, body: dict, timeout: int = 30) -> tuple[int, str]:
     req = urllib.request.Request(
         url,
@@ -113,6 +123,12 @@ def main() -> int:
             ok += 1
             continue
         status, detail = post(args.callback_url, args.oidc_token, body)
+        for delay in _RATE_LIMIT_BACKOFF_S:
+            if status != 429:
+                break
+            time.sleep(delay)
+            status, detail = post(args.callback_url, args.oidc_token, body)
+        time.sleep(_POST_INTERVAL_S)
         if 200 <= status < 300:
             ok += 1
         else:

.buildkite/scripts/rocm/aiter-nightly-overlay.sh

@@ -0,0 +1,112 @@
+#!/bin/bash
+# Install the AITER nightly wheel over this build's ci_base.
+#
+# Runs in CI when AITER_NIGHTLY=1, once the stock ci_base it builds on is
+# published. In order:
+#   1. Select the newest nightly wheel that fits the stock ci_base
+#      (select_aiter_nightly_wheel.py, beside this script, run inside that image)
+#   2. Build a new image: pip install that wheel over the stock ci_base
+#   3. Push the new image as this build's ci_base-build tag, replacing the stock
+#      one, so later steps and native jobs use it; record its digest
+#
+# Buildkite meta-data, a key-value store shared by every job in this build:
+#   reads   rocm-ci-base-image       the stock ci_base
+#   r/w     aiter-nightly-wheel-url  the wheel step 1 chose; a retry reads it
+#                                    back and installs that same wheel
+#   writes  aiter-nightly-ci-base    the new image as tag@digest; the vLLM
+#                                    image build uses it as its ci_base
+#
+# Exit status:
+#   10  no wheel fits this ci_base
+#   11  the wheel does not install
+#   Anything else is unexpected and passed through as is.
+#   Buildkite retries this step once on exit status 1 or 2, which it treats as
+#   transient Docker/registry failures (amd_infra_retry in the template).
+#   10 and 11 sit outside that range: they are real AITER failures, and a retry
+#   would only fail the same way.
+
+set -euo pipefail
+
+SELECT_WHEEL="$(dirname "${BASH_SOURCE[0]}")/select_aiter_nightly_wheel.py"
+BUILD_TAG="rocm/vllm-dev:ci_base-build-${BUILDKITE_BUILD_ID:?BUILDKITE_BUILD_ID is required}"
+
+# Scratch space for the build context and captured errors. The trap deletes it
+# however the script exits: success, fail, or set -e.
+scratch="$(mktemp -d)"
+trap 'rm -rf "${scratch}"' EXIT
+errors="${scratch}/errors"
+
+fail() {
+    local status="$1" title="$2" detail="$3"
+    echo "+++ :x: AITER nightly: ${title}"
+    echo "${detail}" >&2
+    # shellcheck disable=SC2016 # the backticks are a literal markdown fence
+    printf ':x: **AITER nightly: %s**\n\n```\n%s\n```\n' "${title}" "${detail}" \
+        | buildkite-agent annotate --style error --context aiter-nightly
+    exit "${status}"
+}
+
+# Set by ensure-ci-base-amd; this step's depends_on guarantees it is there.
+ci_base="$(buildkite-agent meta-data get rocm-ci-base-image)"
+# Set below on the first attempt. Meta-data is per build, so a retry reuses
+# that wheel, not a newer nightly.
+pinned="$(buildkite-agent meta-data get aiter-nightly-wheel-url --default "")"
+
+echo "--- :docker: Pulling ${ci_base}"
+docker pull "${ci_base}"
+
+echo "--- :mag: Selecting the AITER nightly for ${ci_base}"
+set +e
+out="$(docker run --rm -i --entrypoint python3 "${ci_base}" \
+    -P - ${pinned:+--wheel-url "${pinned}"} \
+    < "${SELECT_WHEEL}" 2> "${errors}")"
+status=$?
+set -e
+cat "${errors}" >&2
+case "${status}" in
+    0) ;;
+    10) fail 10 "no wheel for this image" "${out}" ;;
+    *) fail "${status}" "could not select a wheel" "$(tail -n 20 "${errors}")" ;;
+esac
+# The first line is the wheel's URL; the rest, the annotation naming it.
+wheel_url="$(head -n 1 <<< "${out}")"
+note="$(tail -n +2 <<< "${out}")"
+buildkite-agent meta-data set aiter-nightly-wheel-url "${wheel_url}"
+
+cat > "${scratch}/Dockerfile" <<'EOF'
+ARG CI_BASE_IMAGE
+FROM ${CI_BASE_IMAGE}
+ARG AITER_WHEEL_URL
+RUN python3 -m pip install --no-cache-dir "$AITER_WHEEL_URL"
+EOF
+
+echo "--- :docker: Installing ${wheel_url##*/}"
+if ! docker buildx build \
+    --builder default \
+    --provenance=false \
+    --progress plain \
+    --build-arg "CI_BASE_IMAGE=${ci_base}" \
+    --build-arg "AITER_WHEEL_URL=${wheel_url}" \
+    -t "${BUILD_TAG}" \
+    "${scratch}"; then
+    fail 11 "the wheel does not install" \
+        "Installing ${wheel_url##*/} over ${ci_base} failed; see this step's log."
+fi
+
+docker push "${BUILD_TAG}"
+# Handed off as tag@digest: build-test-image.sh rejects a ci_base that is not
+# digest-pinned.
+digest="$(docker buildx imagetools inspect "${BUILD_TAG}" | awk '$1 == "Digest:" { print $2; exit }')"
+if [[ ! "${digest}" =~ ^sha256:[0-9a-f]{64}$ ]]; then
+    echo "Could not resolve the pushed digest of ${BUILD_TAG}" >&2
+    exit 1
+fi
+# The vLLM image build uses this as its ci_base.
+buildkite-agent meta-data set aiter-nightly-ci-base "${BUILD_TAG}@${digest}"
+# The annotation starts with :warning: when the wheel is stale.
+style=info
+if [[ "${note}" == :warning:* ]]; then
+    style=warning
+fi
+echo "${note}" | buildkite-agent annotate --style "${style}" --context aiter-nightly
+echo "--- :white_check_mark: ${BUILD_TAG}@${digest}"

.buildkite/scripts/rocm/select_aiter_nightly_wheel.py

@@ -0,0 +1,260 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Select the AITER nightly wheel for a vLLM ROCm image.
+
+aiter-nightly-overlay.sh runs this inside the stock ci_base, so it reads that
+image's ROCm, torch and Python. It prints the wheel's URL, then the Buildkite
+annotation naming it.
+"""
+
+import argparse
+import datetime
+import importlib.metadata
+import os
+import sys
+import urllib.request
+from dataclasses import dataclass
+from html.parser import HTMLParser
+from pathlib import Path
+from urllib.parse import unquote, urldefrag, urljoin
+
+import regex as re
+
+INDEX_URL = "https://rocm.frameworks-nightlies.amd.com/whl-multi-arch/amd-aiter/"
+AITER_COMMIT_URL = "https://github.com/ROCm/aiter/commit/"
+STALE_AFTER = datetime.timedelta(days=1)
+
+# Example:
+#   amd_aiter-0.1.25+rocm7.2.3.torch2.12.ae984f5.d20261008-cp312-cp312-linux_x86_64.whl
+# is AITER main at commit ae984f5, built 2026-10-08, for ROCm 7.2.3, torch 2.12
+# and Python 3.12. 0.1.25 is a guess at the next release, not a real one.
+WHEEL_NAME = re.compile(
+    r"^amd_aiter-(?P<base_version>[^+-]+)\+(?P<local>[^-]+)"
+    r"-(?P<python>cp\d+)-(?P=python)-linux_x86_64\.whl$"
+)
+ROCM_FIELD = re.compile(r"(?:^|\.)rocm(\d+(?:\.\d+)*)(?:a\d+)?(?=\.|$)")
+TORCH_FIELD = re.compile(r"(?:^|\.)torch(\d+\.\d+)(?=\.|$)")
+COMMIT_AND_DATE_FIELD = re.compile(r"(?:^|\.)([0-9a-f]{7,40})\.d(\d{8})(?=\.|$)")
+
+
+@dataclass(frozen=True)
+class Wheel:
+    """An AITER nightly wheel, read from its filename (see WHEEL_NAME)."""
+
+    filename: str
+    url: str
+    version: str
+    base_version: str
+    rocm: str
+    torch: str
+    commit: str
+    build_date: datetime.date
+    python: str
+
+
+class _LinkParser(HTMLParser):
+    """Collects every link on the index page, which lists one per wheel."""
+
+    def __init__(self):
+        super().__init__()
+        self.hrefs: list[str] = []
+
+    def handle_starttag(self, tag, attrs):
+        if tag == "a":
+            self.hrefs.extend(
+                value for name, value in attrs if name == "href" and value
+            )
+
+
+def wheel_from_url(url: str) -> Wheel | None:
+    """The AITER nightly wheel an absolute URL points at, if it is one."""
+    filename = unquote(urldefrag(url).url.rsplit("/", 1)[-1])
+    match = WHEEL_NAME.match(filename)
+    if not match:
+        return None
+    local = match["local"]
+    rocm, torch, build = (
+        field.search(local)
+        for field in (ROCM_FIELD, TORCH_FIELD, COMMIT_AND_DATE_FIELD)
+    )
+    if not (rocm and torch and build):
+        return None
+    return Wheel(
+        filename=filename,
+        url=url,
+        version=f"{match['base_version']}+{local}",
+        base_version=match["base_version"],
+        rocm=rocm[1],
+        torch=torch[1],
+        commit=build[1],
+        build_date=datetime.datetime.strptime(build[2], "%Y%m%d").date(),
+        python=match["python"],
+    )
+
+
+def parse_index(html: str, index_url: str) -> list[Wheel]:
+    """Every AITER nightly wheel linked from the index page, a plain HTML list
+    of links, one per file (the format pip reads)."""
+    parser = _LinkParser()
+    parser.feed(html)
+    wheels = (wheel_from_url(urljoin(index_url, href)) for href in parser.hrefs)
+    return [wheel for wheel in wheels if wheel]
+
+
+def rocm_version() -> str:
+    """This image's ROCm version; ROCm 10+ ships as the `rocm` pip SDK."""
+    try:
+        return importlib.metadata.version("rocm")
+    except importlib.metadata.PackageNotFoundError:
+        # No pip SDK: fall back to older ROCm releases, which install to ROCM_PATH.
+        rocm_path = Path(os.environ.get("ROCM_PATH", "/opt/rocm"))
+        return (rocm_path / ".info" / "version").read_text()
+
+
+def rocm_release(version: str) -> str:
+    """`7.2.3-49` or `10.0.0a20261001` -> `7.2.3` or `10.0.0`, as wheel names
+    spell it."""
+    match = re.match(r"\d+(?:\.\d+)*", version.strip())
+    if not match:
+        raise ValueError(f"Unrecognised ROCm version: {version!r}")
+    return match.group(0)
+
+
+def torch_release(version: str) -> str:
+    """`2.12.0a0+git6bbd260` -> `2.12`, as wheel names spell it."""
+    match = re.match(r"\d+\.\d+", version)
+    if not match:
+        raise ValueError(f"Unrecognised torch version: {version!r}")
+    return match.group(0)
+
+
+def python_tag(version_info: tuple[int, int]) -> str:
+    """(3, 12) -> `cp312`, as wheel names spell it."""
+    return f"cp{version_info[0]}{version_info[1]}"
+
+
+def _base_version_key(wheel: Wheel) -> tuple[int, ...]:
+    """`0.1.25` -> (0, 1, 25), so versions sort as numbers: 0.1.10 > 0.1.9."""
+    return tuple(int(part) for part in re.findall(r"\d+", wheel.base_version))
+
+
+def _label(rocm: str, torch: str, python: str) -> str:
+    """`rocm7.2.3/torch2.12/cp312`, for messages."""
+    return f"rocm{rocm}/torch{torch}/{python}"
+
+
+def select_wheel(wheels: list[Wheel], rocm: str, torch: str, python: str) -> Wheel:
+    """Newest build for this ROCm, torch and Python, by date: version labels lag.
+
+    Two builds with the same date and version: whichever the index lists first.
+    """
+    target = (rocm, torch, python)
+    matching = [w for w in wheels if (w.rocm, w.torch, w.python) == target]
+    if not matching:
+        available = sorted({_label(w.rocm, w.torch, w.python) for w in wheels})
+        raise LookupError(
+            f"No AITER nightly for {_label(*target)}. "
+            f"Available: {', '.join(available) or 'none'}"
+        )
+    return max(matching, key=lambda w: (w.build_date, _base_version_key(w)))
+
+
+def fetch(url: str) -> str:
+    """GET url as text. A failure exits 1, so Buildkite retries the step."""
+    with urllib.request.urlopen(url, timeout=60) as response:
+        return response.read().decode()
+
+
+def annotation(wheel: Wheel, today: datetime.date) -> str:
+    """The Buildkite annotation naming the wheel this build tests.
+
+    Starts with :warning: if the wheel is stale; the overlay script styles it so.
+    """
+    age = today - wheel.build_date
+    stale = (
+        f":warning: **AITER nightly is {age.days} days old**: "
+        f"no newer build has been published for "
+        f"{_label(wheel.rocm, wheel.torch, wheel.python)}.\n\n"
+        if age > STALE_AFTER
+        else ""
+    )
+    return stale + (
+        f":crescent_moon: **AITER nightly**: main at "
+        f"[`{wheel.commit}`]({AITER_COMMIT_URL}{wheel.commit}), "
+        f"built {wheel.build_date.isoformat()} for torch {wheel.torch} "
+        f"(`{wheel.version}`)"
+    )
+
+
+class NoWheel(Exception):
+    """No wheel fits this image's ROCm, torch and Python. Exit 10."""
+
+
+def _image_wheel(wheel_url: str) -> Wheel:
+    """The wheel at wheel_url, if it fits this image's ROCm, torch and Python."""
+    wheel = wheel_from_url(wheel_url)
+    if not wheel:
+        raise NoWheel(f"Not an AITER nightly wheel: {wheel_url}")
+    return _for_image([wheel])
+
+
+def _image() -> tuple[str, str, str]:
+    """This image's ROCm, torch and Python, as wheel names spell them."""
+    try:
+        rocm = rocm_release(rocm_version())
+    except (OSError, ValueError) as error:
+        raise NoWheel(f"Cannot read this image's ROCm version: {error}") from error
+    try:
+        torch = torch_release(importlib.metadata.version("torch"))
+    except (importlib.metadata.PackageNotFoundError, ValueError) as error:
+        raise NoWheel(f"Cannot read this image's torch version: {error}") from error
+    return rocm, torch, python_tag(sys.version_info[:2])
+
+
+def _for_image(wheels: list[Wheel]) -> Wheel:
+    """The newest of wheels for the image this runs in."""
+    try:
+        return select_wheel(wheels, *_image())
+    except LookupError as error:
+        raise NoWheel(str(error)) from error
+
+
+def select(index_url: str, wheel_url: str = "") -> Wheel:
+    """The wheel to install: wheel_url on a retry, else the newest.
+
+    However old the newest is, it is tested; the annotation shows its build date.
+    """
+    if wheel_url:
+        return _image_wheel(wheel_url)
+    return _for_image(parse_index(fetch(index_url), index_url))
+
+
+def _today() -> datetime.date:
+    """Today in UTC; a function so tests can set the date."""
+    return datetime.datetime.now(datetime.UTC).date()
+
+
+# Not 1 or 2, which Buildkite retries; see aiter-nightly-overlay.sh.
+EXIT_NO_WHEEL = 10
+
+
+def main(argv=None) -> int:
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--index-url", default=INDEX_URL)
+    parser.add_argument(
+        "--wheel-url", default="", help="a retry's earlier choice; skips the index"
+    )
+    args = parser.parse_args(argv)
+    try:
+        wheel = select(args.index_url, args.wheel_url)
+    except NoWheel as error:
+        print(error)
+        return EXIT_NO_WHEEL
+    print(wheel.url)
+    print(annotation(wheel, _today()))
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

.buildkite/test-amd.yaml

@@ -273,7 +273,7 @@ steps:
 #------------------------------------------------------------  mi250 · v1  -------------------------------------------------------------#
 
 - label: ":amd: (MI250) Batch Invariance" # TBD
-  timeout_in_minutes: 35
+  timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
   dind: false
   agent_pool: mi250_1
@@ -282,14 +282,18 @@ steps:
   source_file_dependencies:
   - vllm/v1/attention
   - vllm/model_executor/layers
+  - vllm/model_executor/kernels/linear
+  - vllm/model_executor/determinism
   - tests/v1/determinism/
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
+  - export VLLM_GPU_MEMORY_UTILIZATION=0.8
+  - export VLLM_TEST_CLEAN_GPU_MEMORY=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pip install pytest-timeout pytest-forked
   - pytest -v -s v1/determinism/test_batch_invariance.py
-  - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
+  - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
 
 - label: ":amd: (MI250) CUDAGraph" # TBD
   timeout_in_minutes: 55
@@ -825,9 +829,7 @@ steps:
   - examples/features/data_parallel/data_parallel_offline.py
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
-  - csrc/custom_quickreduce.cu
+  - csrc/rocm/quickreduce/
-  - csrc/ops.h
-  - csrc/torch_bindings.cpp
   - vllm/model_executor/layers/
   - vllm/entrypoints/llm.py
   - vllm/config/parallel.py
@@ -938,7 +940,6 @@ steps:
   - tests/basic_correctness/models/test_basic_correctness.py
   - vllm/platforms/rocm.py
   commands:
-  - export VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1
   - pytest -v -s distributed/test_ray_v2_executor.py
   - pytest -v -s distributed/test_ray_v2_executor_e2e.py
   - pytest -v -s distributed/test_pipeline_parallel.py -k "ray"
@@ -1510,7 +1511,7 @@ steps:
   parallelism: 5
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - csrc/moe/
   - csrc/rocm/
   - tests/kernels/moe
@@ -1526,7 +1527,32 @@ steps:
   commands:
   - pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --ignore=kernels/moe/test_modular_kernel_combinations.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
   - pytest -v -s kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-  - VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_MOE=1 pytest -v -s kernels/moe/test_modular_kernel_combinations.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_MOE=1 pytest -v -s kernels/moe/test_modular_kernel_combinations.py -k "not multigpu" --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+
+- label: ":amd: (MI300) MoE Kernels Multi-GPU" # TBD
+  timeout_in_minutes: 30
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  dind: false
+  agent_pool: mi300_2
+  num_gpus: 2
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/moe/
+  - csrc/rocm/
+  - tests/kernels/moe/test_modular_kernel_combinations.py
+  - tests/kernels/moe/modular_kernel_tools/
+  - tests/kernels/utils.py
+  - vllm/model_executor/layers/fused_moe/
+  - vllm/distributed/device_communicators/
+  - vllm/distributed/parallel_state.py
+  - vllm/envs.py
+  - vllm/config
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  commands:
+  - VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_MOE=1 pytest -v -s kernels/moe/test_modular_kernel_combinations.py -k "multigpu"
 
 - label: ":amd: (MI300) vLLM IR" # TBD
   timeout_in_minutes: 50
@@ -1556,6 +1582,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - csrc/attention/
+  - csrc/core/vec_dtypes/
   - vllm/v1/attention
   - vllm/model_executor/layers/attention
   - tests/kernels/attention
@@ -1581,10 +1608,10 @@ steps:
   - vllm/platforms/rocm.py
   - tests/kernels/core/test_minimax_reduce_rms.py
   - tests/utils.py
-  - csrc/libtorch_stable/minimax_reduce_rms_kernel.cu
+  - csrc/minimax_reduce_rms_kernel.cu
-  - csrc/libtorch_stable/minimax_reduce_rms_kernel.h
+  - csrc/minimax_reduce_rms_kernel.h
-  - csrc/libtorch_stable/ops.h
+  - csrc/ops.h
-  - csrc/libtorch_stable/torch_bindings.cpp
+  - csrc/torch_bindings.cpp
   - vllm/model_executor/layers/mamba/linear/minimax_linear_attn.py
   optional: true
   commands:
@@ -1606,9 +1633,9 @@ steps:
   - vllm/models/kimi_k3/nvidia/ops/third_party/kda/
   - vllm/models/kimi_k3/amd/
   - CMakeLists.txt
-  - csrc/libtorch_stable/kimi_k3/fused_kda_decode_kernel_rocm.cu
+  - csrc/kimi_k3/fused_kda_decode_kernel_rocm.cu
-  - csrc/libtorch_stable/ops.h
+  - csrc/ops.h
-  - csrc/libtorch_stable/torch_bindings.cpp
+  - csrc/torch_bindings.cpp
   - tests/models/kimi_k3/test_amd_attn_res.py
   - tests/models/kimi_k3/test_amd_kda_decode.py
   - tests/models/kimi_k3/test_kda.py
@@ -1629,7 +1656,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/models/kimi_k3/
-  - csrc/libtorch_stable/kimi_k3/
+  - csrc/kimi_k3/
   - tests/models/kimi_k3/
   - vllm/platforms/rocm.py
   commands:
@@ -1986,6 +2013,19 @@ steps:
   commands:
   - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'
 
+- label: ":amd: (MI300) Multimodal Models (PPL)"
+  timeout_in_minutes: 90
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  dind: false
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/models/multimodal/generation_ppl_test/
+  commands:
+  - pytest -v -s models/multimodal/generation_ppl_test/
+
 - label: ":amd: (MI300) Multimodal Models (Standard) 3: llava + qwen2_vl" # TBD
   timeout_in_minutes: 55
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -2046,7 +2086,6 @@ steps:
   - vllm/model_executor/model_loader/
   optional: true
   commands:
-  - unset VLLM_USE_V2_MODEL_RUNNER
   - pytest -v -s models/quantization
 
 #--------------------------------------------------  mi300 · models / transformers  ---------------------------------------------------#
@@ -2353,7 +2392,6 @@ 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
-  - unset VLLM_USE_V2_MODEL_RUNNER
   - uv pip install --system torchao==0.18.0
   - uv pip install --system conch-triton-kernels
   - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py --ignore quantization/test_rocm_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
@@ -2445,6 +2483,56 @@ steps:
 
 #------------------------------------------------------------  mi300 · v1  -------------------------------------------------------------#
 
+- label: ":amd: (MI300) Batch Invariance" # TBD
+  timeout_in_minutes: 60
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  dind: false
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/attention
+  - vllm/model_executor/layers
+  - vllm/lora/ops/triton_ops
+  - vllm/model_executor/kernels/linear
+  - vllm/model_executor/determinism
+  - tests/v1/determinism/
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pip install pytest-timeout pytest-forked
+  - pytest -v -s v1/determinism/test_batch_invariance.py
+  - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
+  - pytest -v -s v1/determinism/test_lora_shrink_batch_invariant.py
+  - pytest -v -s v1/determinism/test_cutlass_batch_invariance.py
+  - pytest -v -s v1/determinism/test_matmul_batch_invariant.py
+  - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
+  - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_ATTN
+
+- label: ":amd: (MI300) Batch Invariance (Distributed)" # TBD
+  timeout_in_minutes: 45
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  dind: false
+  agent_pool: mi300_4
+  num_gpus: 4
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/distributed/device_communicators/
+  - vllm/model_executor/determinism
+  - vllm/model_executor/layers
+  - vllm/v1/attention
+  - tests/v1/determinism/
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - export VLLM_TEST_CLEAN_GPU_MEMORY=1
+  - VLLM_TEST_TP_SIZE=4 pytest -v -s v1/determinism/test_batch_invariance.py -k "bs1_vs_bsN or without_batch_invariance_should_fail"
+  - VLLM_TEST_TP_SIZE=4 VLLM_TEST_DISABLE_CUSTOM_AR=1 pytest -v -s v1/determinism/test_batch_invariance.py -k bs1_vs_bsN
+  - VLLM_TP_SIZE=4 VLLM_NEEDLE_TRIALS=2 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_ATTN
+  - VLLM_TEST_MODEL=allenai/OLMoE-1B-7B-0924 VLLM_TEST_SERVER_ARGS="--data-parallel-size 4 --enable-expert-parallel --max-model-len 4096" pytest -v -s v1/determinism/test_online_batch_invariance.py
+
 - label: ":amd: (MI300) Spec Decode Speculators + MTP" # TBD
   timeout_in_minutes: 75
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -2466,6 +2554,87 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/speculators/
     - pytest -v -s v1/e2e/spec_decode/mtp/
 
+- label: ":amd: (MI300) Spec Decode DeepSeek MTP Parallel Load"
+  timeout_in_minutes: 45
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  dind: false
+  agent_pool: mi300_2
+  num_gpus: 2
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/spec_decode/llm_base_proposer.py
+  - vllm/v1/spec_decode/eagle.py
+  - vllm/v1/worker/gpu/spec_decode/eagle/
+  - vllm/model_executor/models/deepseek_mtp.py
+  - vllm/model_executor/models/deepseek_v2.py
+  - tests/v1/e2e/spec_decode/test_mtp_parallel_load.py
+  - vllm/distributed/
+  - vllm/model_executor/model_loader/
+  - vllm/model_executor/layers/
+  - vllm/v1/attention/
+  - vllm/v1/engine/
+  - vllm/v1/spec_decode/
+  - vllm/v1/worker/gpu/spec_decode/
+  - vllm/v1/worker/gpu_model_runner.py
+  - vllm/platforms/rocm.py
+  commands:
+  - pytest -v -s v1/e2e/spec_decode/test_mtp_parallel_load.py
+
+- label: ":amd: (MI300) Spec Decode AL DFlash Nightly"
+  timeout_in_minutes: 120
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  dind: false
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/spec_decode/
+  - vllm/v1/worker/gpu/spec_decode/
+  - vllm/model_executor/models/qwen3_dflash.py
+  - vllm/model_executor/models/qwen3_dflash2.py
+  - vllm/model_executor/models/qwen3_5.py
+  - vllm/model_executor/models/laguna_dflash.py
+  - tests/v1/e2e/spec_decode/
+  - vllm/v1/attention/
+  - vllm/v1/sample/
+  - vllm/model_executor/model_loader/
+  - vllm/model_executor/layers/
+  - vllm/model_executor/models/qwen2.py
+  - vllm/model_executor/models/qwen3.py
+  - vllm/model_executor/models/laguna.py
+  - vllm/platforms/rocm.py
+  commands:
+  - pytest -v -s v1/e2e/spec_decode/acceptance_rates/dflash/ -k "not dflash2"
+
+- label: ":amd: (MI300) Spec Decode AL DSpark Nightly"
+  timeout_in_minutes: 90
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  dind: false
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/spec_decode/
+  - vllm/v1/worker/gpu/spec_decode/
+  - vllm/model_executor/models/qwen3_dspark.py
+  - vllm/model_executor/models/gemma4_dspark.py
+  - tests/v1/e2e/spec_decode/
+  - vllm/v1/attention/
+  - vllm/v1/sample/
+  - vllm/model_executor/model_loader/
+  - vllm/model_executor/layers/
+  - vllm/model_executor/models/qwen2.py
+  - vllm/model_executor/models/qwen3.py
+  - vllm/model_executor/models/qwen3_dflash.py
+  - vllm/model_executor/models/gemma4.py
+  - vllm/model_executor/models/gemma4_mm.py
+  - vllm/model_executor/models/gemma4_unified.py
+  - vllm/model_executor/models/gemma4_mtp.py
+  - vllm/platforms/rocm.py
+  commands:
+  - pytest -v -s v1/e2e/spec_decode/acceptance_rates/dspark/
+
 - label: ":amd: (MI300) V1 Spec Decode" # TBD
   timeout_in_minutes: 50
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -2499,7 +2668,6 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/v1/spec_decode/
-  - vllm/model_executor/models/mlp_speculator.py
   - tests/v1/spec_decode/test_acceptance_length.py
   - vllm/platforms/rocm.py
   commands:
@@ -2634,8 +2802,11 @@ steps:
   - vllm/_aiter_ops.py
   - vllm/envs.py
   - vllm/platforms/rocm.py
+  - vllm/models/deepseek_v32/amd/rocm.py
   commands:
-  - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - pytest -v -s v1/attention --ignore=v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+
+  - VLLM_ROCM_USE_AITER=1 pytest -v -s v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
 
 - label: ":amd: (MI300) Metrics, Tracing" # TBD
   timeout_in_minutes: 30
@@ -3444,7 +3615,7 @@ steps:
   - export MODEL_NAMES='moonshotai/Kimi-Linear-48B-A3B-Instruct'
   - export PREFILL_BLOCK_SIZE='2048'
   - export DECODE_BLOCK_SIZE='2048'
-  - export VLLM_SERVE_EXTRA_ARGS='--trust-remote-code,--mamba-cache-mode align,--max-model-len 32768'
+  - export VLLM_SERVE_EXTRA_ARGS='--trust-remote-code,--max-model-len 32768'
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
   - bash v1/kv_connector/nixl_integration/run_accuracy_test.sh
 
@@ -3457,9 +3628,7 @@ steps:
   optional: true
   working_dir: "/vllm-workspace/"
   source_file_dependencies:
-  - csrc/custom_quickreduce.cu
+  - csrc/rocm/quickreduce/
-  - csrc/ops.h
-  - csrc/torch_bindings.cpp
   - vllm/distributed/
   - vllm/model_executor/layers/
   - vllm/entrypoints/llm.py
@@ -3487,8 +3656,11 @@ steps:
   - examples/rl/rlhf_async_new_apis.py
   - tests/distributed/test_weight_transfer.py
   - tests/distributed/test_packed_tensor.py
+  - tests/v1/attention/test_rocm_mla_dcp_context.py
+  - vllm/v1/attention/ops/rocm_aiter_mla_prefill.py
   commands:
   - pytest -v -s tests/distributed/test_context_parallel.py
+  - pytest -v -s tests/v1/attention/test_rocm_mla_dcp_context.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
   - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
   - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=allgather_reducescatter --disable-nccl-for-dp-synchronization
@@ -3978,7 +4150,7 @@ steps:
   optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - csrc/libtorch_stable/gdn/
+  - csrc/gdn/
   - tests/evals/qwen4_exp/
   - vllm/model_executor/layers/mamba/
   - vllm/models/qwen4_exp/
@@ -4266,6 +4438,7 @@ steps:
   - vllm/v1/spec_decode/
   - vllm/v1/worker/gpu/spec_decode/
   - vllm/v1/core/kv_cache_utils.py
+  - vllm/v1/core/kv_cache_planning.py
   - vllm/model_executor/layers/layernorm.py
   - csrc/
   - vllm/_aiter_ops.py
@@ -4299,6 +4472,7 @@ steps:
   - vllm/v1/spec_decode/
   - vllm/v1/worker/gpu/spec_decode/
   - vllm/v1/core/kv_cache_utils.py
+  - vllm/v1/core/kv_cache_planning.py
   - vllm/model_executor/layers/layernorm.py
   - csrc/
   - vllm/_aiter_ops.py
@@ -4332,6 +4506,7 @@ steps:
   - vllm/v1/spec_decode/
   - vllm/v1/worker/gpu/spec_decode/
   - vllm/v1/core/kv_cache_utils.py
+  - vllm/v1/core/kv_cache_planning.py
   - vllm/model_executor/layers/layernorm.py
   - csrc/
   - vllm/_aiter_ops.py
@@ -4365,6 +4540,7 @@ steps:
   - vllm/v1/spec_decode/
   - vllm/v1/worker/gpu/spec_decode/
   - vllm/v1/core/kv_cache_utils.py
+  - vllm/v1/core/kv_cache_planning.py
   - vllm/model_executor/layers/layernorm.py
   - csrc/
   - vllm/_aiter_ops.py
@@ -4398,6 +4574,7 @@ steps:
   - vllm/v1/spec_decode/
   - vllm/v1/worker/gpu/spec_decode/
   - vllm/v1/core/kv_cache_utils.py
+  - vllm/v1/core/kv_cache_planning.py
   - vllm/model_executor/layers/layernorm.py
   - csrc/
   - vllm/_aiter_ops.py
@@ -4549,10 +4726,10 @@ steps:
   - vllm/platforms/rocm.py
   - tests/kernels/core/test_minimax_reduce_rms.py
   - tests/utils.py
-  - csrc/libtorch_stable/minimax_reduce_rms_kernel.cu
+  - csrc/minimax_reduce_rms_kernel.cu
-  - csrc/libtorch_stable/minimax_reduce_rms_kernel.h
+  - csrc/minimax_reduce_rms_kernel.h
-  - csrc/libtorch_stable/ops.h
+  - csrc/ops.h
-  - csrc/libtorch_stable/torch_bindings.cpp
+  - csrc/torch_bindings.cpp
   - vllm/model_executor/layers/mamba/linear/minimax_linear_attn.py
   optional: true
   commands:
@@ -4567,7 +4744,7 @@ steps:
   source_file_dependencies:
   - csrc/quantization/fp4/
   - csrc/attention/mla/
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - vllm/model_executor/layers/fused_moe/cutlass_moe.py
   - vllm/v1/attention/backends/triton_attn.py
   - vllm/v1/attention/backends/rocm_attn.py
@@ -4619,7 +4796,7 @@ steps:
   optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
+  - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
   - vllm/models/deepseek_v4/common/ops/
   - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
   - vllm/platforms/rocm.py
@@ -4635,6 +4812,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - csrc/attention/
+  - csrc/core/vec_dtypes/
   - vllm/v1/attention
   - vllm/model_executor/layers/attention
   - tests/kernels/attention
@@ -4654,7 +4832,7 @@ steps:
   parallelism: 5
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - csrc/moe/
   - csrc/rocm/
   - tests/kernels/moe
@@ -4704,7 +4882,7 @@ steps:
   - vllm/utils/torch_utils.py
   - vllm/platforms/
   - vllm/_aiter_ops.py
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - vllm/distributed/device_communicators/
   - vllm/config
   commands:
@@ -4740,6 +4918,23 @@ steps:
   - amd-smi
   - VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_MOE=1 VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS=0 pytest -v -s kernels/moe/test_mori_moe.py
 
+- label: ":amd: (MI355) DeepSeek-V4.1 Mono Decode Numerics"
+  timeout_in_minutes: 40
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+  dind: false
+  agent_pool: mi355_4
+  num_gpus: 4
+  optional: true
+  working_dir: "/vllm-workspace/"
+  source_file_dependencies:
+  - docker/Dockerfile.rocm_base
+  - vllm/models/deepseek_v41/amd/
+  - tests/models/test_dsv41_mono_numerics.py
+  - tests/models/test_dsv41_mono_layers.py
+  commands:
+  - pytest -v -s tests/models/test_dsv41_mono_numerics.py
+  - pytest -v -s tests/models/test_dsv41_mono_layers.py
+
 - label: ":amd: (MI355) Quantization Kernels Shard %N"
   timeout_in_minutes: 120
   parallelism: 6
@@ -4816,6 +5011,32 @@ steps:
     - pytest -v -s kernels/moe/test_triton_moe_no_act_mul.py
     - pytest -v -s kernels/moe/test_triton_moe_ptpc_fp8.py
 
+- label: ":amd: (MI355 DPX) FP4 and INT4 MoE Kernels"
+  timeout_in_minutes: 40
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+  dind: false
+  agent_pool: mi355_dpx
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - csrc/moe/
+  - csrc/rocm/
+  - vllm/model_executor/layers/fused_moe/
+  - vllm/model_executor/layers/quantization/
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  - vllm/envs.py
+  - vllm/config/
+  - tests/kernels/moe/test_rocm_aiter_moe.py
+  - tests/kernels/moe/test_ocp_mx_moe.py
+  - tests/kernels/moe/test_flydsl_moe.py
+  - tests/kernels/moe/utils.py
+  - tests/kernels/utils.py
+  commands:
+  - pytest -v -s kernels/moe/test_rocm_aiter_moe.py -k 'mi350_mxfp4 or quant_scheme_support_matches_gfx950'
+  - pytest -v -s kernels/moe/test_ocp_mx_moe.py::test_rocm_mxfp4_moe_oracle -k AITER_MXFP4
+  - pytest -v -s kernels/moe/test_flydsl_moe.py
+
 - label: ":amd: (MI355) DeepEP FP8 MoE Kernels" # TBD
   timeout_in_minutes: 55
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -4874,6 +5095,23 @@ steps:
   - VLLM_ROCM_USE_AITER=1 pytest -v -s models/test_hyv4_rocm.py -m 'not distributed'
   - pytest -v -s models/test_vision.py::test_simple_mrope_vision_model_spatial_merge
 
+- label: ":amd: (MI355 DPX) Qwen4 Exp"
+  timeout_in_minutes: 40
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+  dind: false
+  agent_pool: mi355_dpx
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/models/qwen4_exp/
+  - vllm/config/compilation.py
+  - vllm/v1/attention/backends/short_conv_attn.py
+  - tests/models/qwen4_exp/
+  commands:
+  - pytest -v -s models/qwen4_exp/test_qsa_amd.py models/qwen4_exp/test_qsa_prepare_amd.py
+  # PLE and QSA import conflicting HC registrations in a shared process.
+  - pytest -v -s models/qwen4_exp/test_ple.py
+
 - label: ":amd: (MI355 DPX) GLM5Next"
   timeout_in_minutes: 35
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -5163,7 +5401,6 @@ steps:
   - vllm/model_executor/model_loader/
   optional: true
   commands:
-  - unset VLLM_USE_V2_MODEL_RUNNER
   - pytest -v -s models/quantization
 
 #-------------------------------------------------------  mi355 · quantization  --------------------------------------------------------#
@@ -5186,8 +5423,7 @@ steps:
   - tests/rocm/test_moe_weight_replay.py
   optional: true
   commands:
-  - unset VLLM_USE_V2_MODEL_RUNNER
+  - uv pip install --system torchao==0.18.0
-  - uv pip install --system torchao==0.17.0
   - uv pip install --system conch-triton-kernels
   - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py --ignore quantization/test_rocm_moe.py --ignore quantization/test_gfx950_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
   - if [ "$$BUILDKITE_PARALLEL_JOB" = "0" ]; then pytest -v -s rocm/test_moe_weight_replay.py; fi
@@ -5216,6 +5452,37 @@ steps:
 
 #------------------------------------------------------------  mi355 · v1  -------------------------------------------------------------#
 
+- label: ":amd: (MI355) Batch Invariance" # TBD
+  timeout_in_minutes: 60
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+  dind: false
+  agent_pool: mi355_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/attention
+  - vllm/model_executor/layers
+  - vllm/lora/ops/triton_ops
+  - vllm/model_executor/kernels/linear
+  - vllm/model_executor/determinism
+  - tests/v1/determinism/
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pip install pytest-timeout pytest-forked
+  - pytest -v -s v1/determinism/test_batch_invariance.py
+  - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
+  - pytest -v -s v1/determinism/test_lora_shrink_batch_invariant.py
+  - pytest -v -s v1/determinism/test_batch_invariance_vlm.py
+  - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
+  - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_ATTN
+  - pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
+  - pytest -v -s v1/determinism/test_matmul_batch_invariant.py
+  - pytest -v -s v1/determinism/test_cutlass_batch_invariance.py
+  - pytest -v -s v1/determinism/test_online_batch_invariance.py
+  - pytest -v -s v1/determinism/test_mx_linear_batch_invariant.py
+
 - label: ":amd: (MI355) Spec Decode AL DFlash2 Nightly"
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -5646,10 +5913,11 @@ steps:
   - vllm/_aiter_ops.py
   - vllm/envs.py
   - vllm/platforms/rocm.py
+  - vllm/models/deepseek_v32/amd/rocm.py
   optional: true
   commands:
-  - pytest -v -s v1/attention --ignore=v1/attention/test_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - pytest -v -s v1/attention --ignore=v1/attention/test_mla_backends.py --ignore=v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-  - VLLM_ROCM_USE_AITER=1 pytest -v -s v1/attention/test_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - VLLM_ROCM_USE_AITER=1 pytest -v -s v1/attention/test_mla_backends.py v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
 
 - label: ":amd: (MI355 DPX) CUDAGraph"
   timeout_in_minutes: 55
@@ -5725,7 +5993,6 @@ steps:
   - vllm/model_executor/models/gemma4_mtp.py
   - vllm/model_executor/models/llama.py
   - vllm/model_executor/models/llama_eagle3.py
-  - vllm/model_executor/models/medusa.py
   - vllm/model_executor/models/registry.py
   - tests/evals/gsm8k/
   - tests/utils.py

.buildkite/test_areas/attention.yaml

@@ -2,8 +2,9 @@ group: Attention
 depends_on: 
   - image-build
 steps:
-- label: ":nvidia: (H100) V1 Attention Shard %N"
+- label: ":nvidia: (H100) V1 Attention"
   key: v1-attention-h100-mi300
+  automatic_shard: true
   timeout_in_minutes: 85
   device: h100
   source_file_dependencies:
@@ -12,14 +13,18 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+    - pytest -v -s v1/attention
-  parallelism: 2
   mirror:
     amd:
       label: ":amd: (MI300) V1 Attention Shard %N"
       dind: false
       device: mi300_1
       timeout_in_minutes: 125
+      # automatic_shard covers NVIDIA steps only; the mirror keeps its
+      # fixed pytest shards.
+      parallelism: 2
+      commands:
+      - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -30,9 +35,14 @@ steps:
       - vllm/_aiter_ops.py
       - vllm/envs.py
       - vllm/platforms/rocm.py
+      - vllm/models/deepseek_v32/amd/rocm.py
+      commands:
+      - pytest -v -s v1/attention --ignore=v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+      - VLLM_ROCM_USE_AITER=1 pytest -v -s v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
 
-- label: ":nvidia: (B200) V1 Attention Shard %N"
+- label: ":nvidia: (B200) V1 Attention"
   key: v1-attention-b200
+  automatic_shard: true
   timeout_in_minutes: 80
   device: b200-k8s
   source_file_dependencies:
@@ -41,14 +51,16 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+    - pytest -v -s v1/attention
-  parallelism: 2
   mirror:
     amd:
       label: ":amd: (MI355 DPX) V1 Attention Shard %N"
       dind: false
       device: mi355_dpx
       timeout_in_minutes: 125
+      # automatic_shard covers NVIDIA steps only; the mirror keeps its
+      # fixed pytest shards.
+      parallelism: 2
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -59,6 +71,21 @@ steps:
       - vllm/_aiter_ops.py
       - vllm/envs.py
       - vllm/platforms/rocm.py
+      - vllm/models/deepseek_v32/amd/rocm.py
       commands:
-      - pytest -v -s v1/attention --ignore=v1/attention/test_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+      - pytest -v -s v1/attention --ignore=v1/attention/test_mla_backends.py --ignore=v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-      - VLLM_ROCM_USE_AITER=1 pytest -v -s v1/attention/test_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+      - VLLM_ROCM_USE_AITER=1 pytest -v -s v1/attention/test_mla_backends.py v1/attention/test_sparse_mla_backends.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+
+- label: ":nvidia: (DGX) Spark FlashInfer NVFP4 XQA"
+  key: flashinfer-nvfp4-xqa-dgx-spark
+  timeout_in_minutes: 30
+  device: dgx-spark
+  optional: true
+  num_devices: 1
+  depends_on:
+  - arm64-image-build
+  source_file_dependencies:
+  - vllm/v1/attention/backends/flashinfer.py
+  - tests/v1/attention/test_attention_backends.py
+  commands:
+  - pytest -v -s v1/attention/test_attention_backends.py -k test_flashinfer_xqa_nvfp4_decode_correctness

.buildkite/test_areas/disaggregated.yaml

@@ -426,7 +426,7 @@ steps:
     # Keep TP=2 prefill and TP=1 decode above both hybrid page-size minima.
     PREFILL_BLOCK_SIZE: "2048"
     DECODE_BLOCK_SIZE: "2048"
-    VLLM_SERVE_EXTRA_ARGS: "--trust-remote-code,--mamba-cache-mode align,--max-model-len 32768"
+    VLLM_SERVE_EXTRA_ARGS: "--trust-remote-code,--max-model-len 32768"
   commands:
     - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - bash v1/kv_connector/nixl_integration/run_accuracy_test.sh

.buildkite/test_areas/disaggregated_mooncake.yaml

@@ -4,6 +4,8 @@ depends_on:
 steps:
 - label: ":nvidia: (B200) Distributed MooncakeConnector PD accuracy"
   key: distributed-mooncakeconnector-pd-accuracy-4-gpus
+  # Nightly only: 4-GPU B200 pods wait hours behind smaller jobs on per-PR runs.
+  optional: true
   timeout_in_minutes: 10
   working_dir: "/vllm-workspace/tests"
   # Runs on b200-k8s because it is the only NVIDIA pool that can carry mooncake
@@ -22,6 +24,7 @@ steps:
   mirror:
     amd:
       label: ":amd: (MI355) Distributed MooncakeConnector HIP PD accuracy"
+      optional: true
       dind: false
       device: mi355_4
       num_devices: 4

.buildkite/test_areas/distributed.yaml

@@ -77,6 +77,7 @@ steps:
 
 - label: ":nvidia: (L4) Distributed DP Basic"
   key: distributed-dp-tests-2-gpus
+  automatic_shard: true
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   device: l4
@@ -300,6 +301,7 @@ steps:
 
 - label: ":nvidia: (L4) Distributed DP Extended"
   key: distributed-dp-tests-4-gpus
+  automatic_shard: true
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   device: l4
@@ -508,9 +510,7 @@ steps:
       - examples/features/data_parallel/data_parallel_offline.py
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.py
-      - csrc/custom_quickreduce.cu
+      - csrc/rocm/quickreduce/
-      - csrc/ops.h
-      - csrc/torch_bindings.cpp
       - vllm/model_executor/layers/
       - vllm/entrypoints/llm.py
       - vllm/config/parallel.py
@@ -522,13 +522,17 @@ steps:
       - tests/distributed/test_rocm_aiter_custom_ar.py
       - tests/distributed/test_rocm_quick_reduce.py
       - tests/distributed/test_quick_all_reduce.py
+      - tests/distributed/test_aiter_rms_fallback.py
       - tests/v1/e2e/general/test_rocm_aiter_custom_ar.py
       - tests/utils.py
       - examples/rl/rlhf_async_new_apis.py
       - tests/distributed/test_weight_transfer.py
       - tests/distributed/test_packed_tensor.py
+      - tests/v1/attention/test_rocm_mla_dcp_context.py
+      - vllm/v1/attention/ops/rocm_aiter_mla_prefill.py
       commands:
       - pytest -v -s tests/distributed/test_context_parallel.py
+      - pytest -v -s tests/v1/attention/test_rocm_mla_dcp_context.py
       - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
       - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
       - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=allgather_reducescatter --disable-nccl-for-dp-synchronization
@@ -539,6 +543,7 @@ steps:
       - pytest -v -s tests/v1/e2e/general/test_rocm_aiter_custom_ar.py
       - pytest -v -s tests/distributed/test_rocm_quick_reduce.py
       - pytest -v -s tests/distributed/test_quick_all_reduce.py
+      - pytest -v -s tests/distributed/test_aiter_rms_fallback.py
 
 - label: ":nvidia: (B200) Distributed"
   key: distributed-tests-2xb200
@@ -636,6 +641,7 @@ steps:
 
 - label: ":nvidia: (L4) Pipeline + Context Parallelism"
   key: pipeline-context-parallelism-4-gpus
+  automatic_shard: true
   timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   device: l4
@@ -677,6 +683,7 @@ steps:
 
 - label: ":nvidia: (L4) RayExecutorV2"
   key: rayexecutorv2-4-gpus
+  automatic_shard: true
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   device: l4
@@ -690,7 +697,6 @@ steps:
   - tests/distributed/test_pipeline_parallel.py
   - tests/basic_correctness/models/test_basic_correctness.py
   commands:
-  - export VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1
   - export NCCL_CUMEM_HOST_ENABLE=0
   - pytest -v -s distributed/test_ray_v2_executor.py
   - pytest -v -s distributed/test_ray_v2_executor_e2e.py
@@ -715,7 +721,6 @@ steps:
       - tests/basic_correctness/models/test_basic_correctness.py
       - vllm/platforms/rocm.py
       commands:
-      - export VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1
       - pytest -v -s distributed/test_ray_v2_executor.py
       - pytest -v -s distributed/test_ray_v2_executor_e2e.py
       - pytest -v -s distributed/test_pipeline_parallel.py -k "ray"

.buildkite/test_areas/engine.yaml

@@ -141,7 +141,7 @@ steps:
     - tests/v1/e2e/general/
   commands:
     - pytest -v -s v1/e2e/general/test_kv_sharing_fast_prefill.py
-    - pytest -v -s v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2 v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2_async
+    - pytest -v -s v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2 v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2_async v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2_async_flashinfer
   mirror:
     amd:
       label: ":amd: (MI355 DPX) E2E Core Large Memory"
@@ -164,6 +164,7 @@ steps:
   device: h100
   num_devices: 2
   key: replayssm-e2e
+  automatic_shard: true
   timeout_in_minutes: 90
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:

.buildkite/test_areas/entrypoints.yaml

@@ -31,6 +31,7 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/entrypoints
+  - vllm/utils/flashinfer.py
   - tests/entrypoints/unit_tests
   - tests/entrypoints/launchers
   commands:
@@ -120,6 +121,7 @@ steps:
 - label: ":nvidia: (H200 MIG 35GB) Entrypoints Integration (OpenAI API chat_completion)"
   device: h200_35gb
   key: entrypoints-integration-api-server-openai-chat_completion
+  automatic_shard: true
   timeout_in_minutes: 83
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -141,6 +143,7 @@ steps:
 - label: ":nvidia: (H200 MIG 35GB) Entrypoints Integration (API Server Generate)"
   device: h200_35gb
   key: entrypoints-integration-api-server-generate
+  automatic_shard: true
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -213,6 +216,7 @@ steps:
 - label: ":nvidia: (H200 MIG 35GB) Entrypoints Integration (Multimodal)"
   device: h200_35gb
   key: entrypoints-integration-multimodal
+  automatic_shard: true
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:

.buildkite/test_areas/expert_parallelism.yaml

@@ -58,6 +58,8 @@ steps:
 
 - label: ":nvidia: (H100) Elastic EP Scaling"
   key: elastic-ep-scaling-test
+  # Nightly only: 4-GPU H100 pods wait behind smaller jobs on per-PR runs.
+  optional: true
   timeout_in_minutes: 40
   device: h100
   working_dir: "/vllm-workspace/tests"

.buildkite/test_areas/kernels.yaml

@@ -30,9 +30,10 @@ steps:
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.py
 
-- label: ":nvidia: (H200 MIG 35GB) Core Operation Kernels Shard %N"
+- label: ":nvidia: (H200 MIG 35GB) Core Operation Kernels"
   device: h200_35gb
   key: kernels-core-operation-test
+  automatic_shard: true
   timeout_in_minutes: 120
   source_file_dependencies:
   - csrc/
@@ -40,10 +41,10 @@ steps:
   - tests/kernels/test_concat_mla_q.py
   - tests/kernels/test_fused_qk_norm_rope_gate.py
   commands:
-    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py
-  parallelism: 3
   mirror:
     amd:
+      parallelism: 3
       label: ":amd: (MI300) Core Operation Kernels Shard %N"
       dind: false
       device: mi300_1
@@ -72,10 +73,10 @@ steps:
   num_devices: 2
   device: h100
   source_file_dependencies:
-  - csrc/libtorch_stable/minimax_reduce_rms_kernel.cu
+  - csrc/minimax_reduce_rms_kernel.cu
-  - csrc/libtorch_stable/minimax_reduce_rms_kernel.h
+  - csrc/minimax_reduce_rms_kernel.h
-  - csrc/libtorch_stable/ops.h
+  - csrc/ops.h
-  - csrc/libtorch_stable/torch_bindings.cpp
+  - csrc/torch_bindings.cpp
   - vllm/model_executor/layers/minimax_rms_norm/
   - vllm/model_executor/layers/mamba/linear/minimax_linear_attn.py
   - tests/kernels/core/test_minimax_reduce_rms.py
@@ -122,7 +123,7 @@ steps:
       depends_on:
       - image-build-amd
       source_file_dependencies:
-      - csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
+      - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
       - vllm/models/deepseek_v4/common/ops/
       - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
       - vllm/platforms/rocm.py
@@ -230,8 +231,8 @@ steps:
 # Catch-all for test files at the tests/kernels root. This job collects
 # the whole root so new files are wired by default.
 # Files with dedicated jobs elsewhere in this file are excluded via --ignore
-# (test_bf16x3_router_gemm_cutedsl and test_ll_bf16_gemm run in their own jobs /
+# (test_bf16x3_router_gemm_cutedsl runs in Kernels (B200); test_ll_bf16_gemm and
-# Kernels (B200)).
+# test_ll_fp32w_gemm run in Gate Linear CuTeDSL Kernels (B200)).
 - label: ":nvidia: (B200) Miscellaneous Kernels"
   key: kernels-root-misc-test-b200
   timeout_in_minutes: 30
@@ -260,6 +261,7 @@ steps:
         --ignore=kernels/test_bf16x3_router_gemm_cutedsl.py
         --ignore=kernels/test_bf16x3_router_gemm_rocm.py
         --ignore=kernels/test_ll_bf16_gemm.py
+        --ignore=kernels/test_ll_fp32w_gemm.py
         --ignore=kernels/test_shuffle_rows.py
     # BROKEN on main, pending kernel fixes (B200):
     #   test_shuffle_rows.py (1: test_shuffle_rows_edge_cases)
@@ -316,6 +318,7 @@ steps:
   timeout_in_minutes: 30
   source_file_dependencies:
   - csrc/attention/
+  - csrc/core/vec_dtypes/
   - vllm/v1/attention
     # TODO: remove this dependency (https://github.com/vllm-project/vllm/issues/32267)
   - vllm/model_executor/layers/attention
@@ -335,6 +338,7 @@ steps:
       - image-build-amd
       source_file_dependencies:
       - csrc/attention/
+      - csrc/core/vec_dtypes/
       - vllm/v1/attention
       - vllm/model_executor/layers/attention
       - tests/kernels/attention
@@ -407,10 +411,14 @@ steps:
       depends_on:
       - image-build-amd
       source_file_dependencies:
+      - vllm/models/deepseek_v41/amd/rocm.py
+      - vllm/models/deepseek_v41/sparse_mla.py
       - vllm/v1/attention/backends/mla/rocm_aiter_mla.py
+      - vllm/v1/attention/ops/rocm_aiter_mla_sparse.py
       - vllm/v1/attention/selector.py
       - vllm/platforms/rocm.py
       - vllm/_aiter_ops.py
+      - tests/kernels/attention/test_rocm_triton_attn_dsv4.py
       - tests/kernels/attention/test_rocm_aiter_mla_decode.py
       - tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py
       - tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py
@@ -421,6 +429,7 @@ steps:
       - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode.py
       - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py
       - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_op_registration.py
+      - pytest -v -s tests/kernels/attention/test_rocm_triton_attn_dsv4.py::test_dsv41_adaptive_mla_swa_metadata_graph_replay
 
 - label: ":nvidia: (L4) Quantization Kernels Shard %N"
   key: kernels-quantization-test
@@ -467,7 +476,7 @@ steps:
   num_devices: 1
   timeout_in_minutes: 50
   source_file_dependencies:
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - csrc/moe/
   - tests/kernels/moe
   - vllm/model_executor/layers/fused_moe/
@@ -485,7 +494,7 @@ steps:
       device: mi300_1
       timeout_in_minutes: 55
       source_file_dependencies:
-      - csrc/quantization/cutlass_w8a8/moe/
+      - csrc/quantization/w8a8/cutlass/moe/
       - csrc/moe/
       - tests/kernels/moe
       - vllm/model_executor/layers/fused_moe/
@@ -509,8 +518,10 @@ steps:
   source_file_dependencies:
   - csrc/mamba/
   - tests/kernels/mamba
+  - vllm/model_executor/layers/mamba/checkpoint.py
   - vllm/model_executor/layers/mamba/gdn
   - vllm/model_executor/layers/mamba/ops
+  - vllm/v1/attention/backends/gdn_attn.py
   commands:
     - pytest -v -s kernels/mamba
   mirror:
@@ -576,11 +587,10 @@ steps:
   # optional: true
   source_file_dependencies:
   - csrc/quantization/fp4/
-  - csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu
+  - csrc/nvfp4_kv_cache_kernels.cu
-  - csrc/libtorch_stable/quantization/fp4/
+  - csrc/quantization/w8a8/cutlass/
-  - csrc/libtorch_stable/quantization/w8a8/cutlass/
   - csrc/attention/mla/
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - vllm/model_executor/layers/fusion/
   - vllm/model_executor/layers/fused_moe/experts/cutlass_moe.py
   - vllm/model_executor/layers/fused_moe/experts/flashinfer_cutlass_moe.py
@@ -594,9 +604,6 @@ steps:
   - vllm/v1/attention/backends/mla/flashinfer_mla.py
   - vllm/v1/attention/selector.py
   - vllm/platforms/cuda.py
-  - vllm/model_executor/kernels/linear/cute_dsl/ll_bf16.py
-  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_dotprod.py
-  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py
   - vllm/cute_utils/
   - vllm/model_executor/layers/mamba/ops/gdn_chunk_cutedsl/
   - vllm/model_executor/layers/fused_moe/router/bf16x3_router_gemm_cutedsl.py
@@ -606,9 +613,9 @@ steps:
   - vllm/config/attention.py
   - vllm/v1/attention/backends/registry.py
   - vllm/_custom_ops.py
-  - csrc/libtorch_stable/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu
+  - csrc/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu
-  - csrc/libtorch_stable/ops.h
+  - csrc/ops.h
-  - csrc/libtorch_stable/torch_bindings.cpp
+  - csrc/torch_bindings.cpp
   - tests/kernels/attention/test_flashinfer_mla_decode.py
   - tests/kernels/attention/test_minimax_m3_msa_cutlass_sparse_decode.py
   - tests/kernels/attention/test_cache.py
@@ -616,7 +623,6 @@ steps:
   - tests/kernels/test_fused_minimax_m3_qknorm_rope_kv_insert.py
   - tests/kernels/mamba/test_gdn_prefill_cutedsl.py
   - tests/kernels/test_bf16x3_router_gemm_cutedsl.py
-  - tests/kernels/test_ll_bf16_gemm.py
   - tests/kernels/test_top_k_per_row.py
   - tests/kernels/quantization/test_silu_mul_nvfp4_quant.py
   - tests/kernels/quantization/test_block_fp8.py
@@ -637,7 +643,6 @@ steps:
           pytest -v -s tests/kernels/moe/test_nvfp4_moe.py
           pytest -v -s tests/kernels/moe/test_ocp_mx_moe.py
           pytest -v -s tests/kernels/test_bf16x3_router_gemm_cutedsl.py
-          pytest -v -s tests/kernels/test_ll_bf16_gemm.py
           ;;
         1)
           python3 examples/basic/offline_inference/chat.py
@@ -694,6 +699,28 @@ steps:
       - amd-smi
       - pytest -v -s kernels/quantization/test_nvfp4_emulation.py -k test_nvfp4_dense_emulation --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
 
+- label: ":nvidia: (B200) Gate Linear CuTeDSL Kernels"
+  key: kernels-gate-linear-cutedsl-b200
+  timeout_in_minutes: 15
+  working_dir: "/vllm-workspace/"
+  device: b200-k8s
+  source_file_dependencies:
+  - vllm/cute_utils/
+  - vllm/platforms/cuda.py
+  - vllm/model_executor/layers/fused_moe/router/gate_linear.py
+  - vllm/model_executor/kernels/linear/cute_dsl/_ll_router_common.py
+  - vllm/model_executor/kernels/linear/cute_dsl/ll_bf16.py
+  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_dotprod.py
+  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py
+  - tests/kernels/test_ll_bf16_gemm.py
+  - vllm/model_executor/kernels/linear/cute_dsl/ll_fp32w.py
+  - vllm/model_executor/kernels/linear/cute_dsl/_ll_fp32w_dotprod.py
+  - tests/kernels/test_ll_fp32w_gemm.py
+  commands:
+    - nvidia-smi
+    - pytest -v -s tests/kernels/test_ll_bf16_gemm.py
+    - pytest -v -s tests/kernels/test_ll_fp32w_gemm.py
+
 - label: ":nvidia: (DGX) Spark b12x Linear Kernels Nightly"
   key: b12x-linear-kernels-dgx-spark-nightly
   timeout_in_minutes: 60
@@ -726,7 +753,7 @@ steps:
   - tests/quantization/reference_mxfp4.py
   - tests/quantization/test_auto_round.py
   commands:
-  - uv pip install --system b12x==1.3.0
+  - uv pip install --system b12x==1.5.1
   - pytest -v -s model_executor/kernels/test_b12x_linear.py
       model_executor/test_b12x_warmup.py
       kernels/quantization/test_block_fp8.py
@@ -853,7 +880,7 @@ steps:
   device: h100
   num_devices: 2
   source_file_dependencies:
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - csrc/moe/
   - tests/kernels/moe
   - vllm/model_executor/layers/fused_moe/
@@ -896,7 +923,7 @@ steps:
   device: b200-k8s
   num_devices: 2
   source_file_dependencies:
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - csrc/moe/
   - tests/kernels/moe
   - vllm/model_executor/layers/fused_moe/
@@ -1009,7 +1036,7 @@ steps:
   source_file_dependencies:
   - csrc/quantization/fp4/
   - csrc/attention/mla/
-  - csrc/quantization/cutlass_w8a8/moe/
+  - csrc/quantization/w8a8/cutlass/moe/
   - vllm/model_executor/layers/fused_moe/cutlass_moe.py
   - vllm/v1/attention/backends/triton_attn.py
   - vllm/v1/attention/backends/rocm_attn.py

.buildkite/test_areas/lm_eval.yaml

@@ -105,6 +105,7 @@ steps:
 
 - label: ":nvidia: (H100) LM Eval Large Models"
   key: lm-eval-large-models-4xh100
+  automatic_shard: true
   device: h100
   optional: true
   num_devices: 4
@@ -139,6 +140,7 @@ steps:
 
 - label: ":nvidia: (B200) LM Eval Small Models"
   key: lm-eval-small-models-1xb200
+  automatic_shard: true
   timeout_in_minutes: 50
   device: b200-k8s
   optional: true
@@ -230,12 +232,13 @@ steps:
   - vllm/v1/attention/backends/mla/
   - vllm/model_executor/layers/attention/sparse_mla_attention.py
   - vllm/v1/worker/gpu/
-  - csrc/libtorch_stable/hisparse_kernels.cu
+  - csrc/hisparse_kernels.cu
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-hisparse.txt
 
 - label: ":nvidia: (B200) LM Eval PCP"
   key: lm-eval-pcp-4xb200
+  automatic_shard: true
   timeout_in_minutes: 360
   device: b200-k8s
   num_devices: 4
@@ -371,6 +374,7 @@ steps:
       - vllm/tokenizers/
       - vllm/v1/attention/
       - vllm/v1/core/kv_cache_utils.py
+      - vllm/v1/core/kv_cache_planning.py
       - vllm/v1/spec_decode/metrics.py
       - vllm/v1/worker/
       - vllm/_aiter_ops.py
@@ -383,6 +387,7 @@ steps:
 
 - label: ":nvidia: (B200) LM Eval Large Models EP"
   key: lm-eval-large-models-ep-2xb200
+  automatic_shard: true
   timeout_in_minutes: 60
   device: b200-k8s
   optional: true
@@ -516,6 +521,7 @@ steps:
       - vllm/v1/spec_decode/
       - vllm/v1/worker/gpu/spec_decode/
       - vllm/v1/core/kv_cache_utils.py
+      - vllm/v1/core/kv_cache_planning.py
       - vllm/model_executor/layers/layernorm.py
       - csrc/
       - vllm/_aiter_ops.py
@@ -559,6 +565,7 @@ steps:
       - vllm/v1/spec_decode/
       - vllm/v1/worker/gpu/spec_decode/
       - vllm/v1/core/kv_cache_utils.py
+      - vllm/v1/core/kv_cache_planning.py
       - vllm/model_executor/layers/layernorm.py
       - csrc/
       - vllm/_aiter_ops.py
@@ -602,6 +609,7 @@ steps:
       - vllm/v1/spec_decode/
       - vllm/v1/worker/gpu/spec_decode/
       - vllm/v1/core/kv_cache_utils.py
+      - vllm/v1/core/kv_cache_planning.py
       - vllm/model_executor/layers/layernorm.py
       - csrc/
       - vllm/_aiter_ops.py
@@ -645,6 +653,7 @@ steps:
       - vllm/v1/spec_decode/
       - vllm/v1/worker/gpu/spec_decode/
       - vllm/v1/core/kv_cache_utils.py
+      - vllm/v1/core/kv_cache_planning.py
       - vllm/model_executor/layers/layernorm.py
       - csrc/
       - vllm/_aiter_ops.py
@@ -688,6 +697,7 @@ steps:
       - vllm/v1/spec_decode/
       - vllm/v1/worker/gpu/spec_decode/
       - vllm/v1/core/kv_cache_utils.py
+      - vllm/v1/core/kv_cache_planning.py
       - vllm/model_executor/layers/layernorm.py
       - csrc/
       - vllm/_aiter_ops.py
@@ -1185,7 +1195,7 @@ steps:
   optional: true
   num_devices: 4
   source_file_dependencies:
-  - csrc/libtorch_stable/gdn/
+  - csrc/gdn/
   - tests/evals/qwen4_exp/
   - vllm/model_executor/layers/mamba/
   - vllm/models/qwen4_exp/
@@ -1203,7 +1213,7 @@ steps:
   optional: true
   num_devices: 4
   source_file_dependencies:
-  - csrc/libtorch_stable/gdn/
+  - csrc/gdn/
   - tests/evals/qwen4_exp/
   - vllm/model_executor/layers/mamba/
   - vllm/models/qwen4_exp/
@@ -1224,7 +1234,7 @@ steps:
       depends_on:
       - image-build-amd
       source_file_dependencies:
-      - csrc/libtorch_stable/gdn/
+      - csrc/gdn/
       - tests/evals/qwen4_exp/
       - vllm/model_executor/layers/mamba/
       - vllm/models/qwen4_exp/
@@ -1335,6 +1345,8 @@ steps:
 
 - label: ":nvidia: (H100) LM Eval KV-Offload Large"
   key: kv-offload-large
+  # Nightly only: 4-GPU H100 pods wait behind smaller jobs on per-PR runs.
+  optional: true
   timeout_in_minutes: 40
   device: h100
   num_devices: 4

.buildkite/test_areas/misc.yaml

@@ -446,6 +446,17 @@ steps:
       - setup.py
       - vllm/platforms/rocm.py
 
+- label: ":nvidia: (H200 MIG 18GB) Python 3.14 Smoke Nightly"
+  key: python-314-smoke-nightly
+  depends_on: ~
+  optional: true
+  timeout_in_minutes: 30
+  device: h200_18gb
+  source_file_dependencies:
+  - tests/standalone_tests/python_only_compile.sh
+  commands:
+  - PYTHON_ONLY_COMPILE_PYTHON=3.14 bash standalone_tests/python_only_compile.sh
+
 - label: ":nvidia: (H200 MIG 35GB) Async Engine, Inputs, Utils, Worker"
   device: h200_35gb
   key: async-engine-inputs-utils-worker
@@ -644,21 +655,45 @@ steps:
   source_file_dependencies:
     - vllm/v1/attention
     - vllm/model_executor/layers
+    - vllm/model_executor/kernels/linear
+    - vllm/model_executor/determinism
     - tests/v1/determinism/
   commands:
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
     - pip install pytest-timeout pytest-forked
     - pytest -v -s v1/determinism/test_batch_invariance.py
     - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
+  mirror:
+    amd:
+      label: ":amd: (MI250) Batch Invariance"
+      dind: false
+      device: mi250_1
+      timeout_in_minutes: 60
+      env:
+        VLLM_GPU_MEMORY_UTILIZATION: "0.8"
+        VLLM_TEST_CLEAN_GPU_MEMORY: "1"
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/v1/attention
+      - vllm/model_executor/layers
+      - vllm/model_executor/kernels/linear
+      - vllm/model_executor/determinism
+      - tests/v1/determinism/
+      - vllm/_aiter_ops.py
+      - vllm/platforms/rocm.py
 
 - label: ":nvidia: (H100) Batch Invariance"
   key: batch-invariance-h100
+  automatic_shard: true
   timeout_in_minutes: 60
   device: h100
   source_file_dependencies:
     - vllm/v1/attention
     - vllm/model_executor/layers
     - vllm/lora/ops/triton_ops
+    - vllm/model_executor/kernels/linear
+    - vllm/model_executor/determinism
     - tests/v1/determinism/
   commands:
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
@@ -670,15 +705,18 @@ steps:
     - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k FLASH_ATTN
   mirror:
     amd:
-      label: ":amd: (MI250) Batch Invariance"
+      label: ":amd: (MI300) Batch Invariance"
       dind: false
-      device: mi250_1
+      device: mi300_1
-      timeout_in_minutes: 35
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       source_file_dependencies:
       - vllm/v1/attention
       - vllm/model_executor/layers
+      - vllm/lora/ops/triton_ops
+      - vllm/model_executor/kernels/linear
+      - vllm/model_executor/determinism
       - tests/v1/determinism/
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.py
@@ -687,15 +725,23 @@ steps:
       - pip install pytest-timeout pytest-forked
       - pytest -v -s v1/determinism/test_batch_invariance.py
       - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
+      - pytest -v -s v1/determinism/test_lora_shrink_batch_invariant.py
+      - pytest -v -s v1/determinism/test_cutlass_batch_invariance.py
+      - pytest -v -s v1/determinism/test_matmul_batch_invariant.py
+      - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
+      - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_ATTN
 
 - label: ":nvidia: (B200) Batch Invariance"
   key: batch-invariance-b200
+  automatic_shard: true
   timeout_in_minutes: 45
   device: b200-k8s
   source_file_dependencies:
     - vllm/v1/attention
     - vllm/model_executor/layers
     - vllm/lora/ops/triton_ops
+    - vllm/model_executor/kernels/linear
+    - vllm/model_executor/determinism
     - tests/v1/determinism/
   commands:
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
@@ -710,6 +756,80 @@ steps:
     - pytest -v -s v1/determinism/test_matmul_batch_invariant.py
     - pytest -v -s v1/determinism/test_cutlass_batch_invariance.py
     - pytest -v -s v1/determinism/test_online_batch_invariance.py
+  mirror:
+    amd:
+      label: ":amd: (MI355) Batch Invariance"
+      dind: false
+      device: mi355_1
+      timeout_in_minutes: 60
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/v1/attention
+      - vllm/model_executor/layers
+      - vllm/lora/ops/triton_ops
+      - vllm/model_executor/kernels/linear
+      - vllm/model_executor/determinism
+      - tests/v1/determinism/
+      - vllm/_aiter_ops.py
+      - vllm/platforms/rocm.py
+      commands:
+      - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+      - pip install pytest-timeout pytest-forked
+      - pytest -v -s v1/determinism/test_batch_invariance.py
+      - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
+      - pytest -v -s v1/determinism/test_lora_shrink_batch_invariant.py
+      - pytest -v -s v1/determinism/test_batch_invariance_vlm.py
+      - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_MLA
+      - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_ATTN
+      - pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
+      - pytest -v -s v1/determinism/test_matmul_batch_invariant.py
+      - pytest -v -s v1/determinism/test_cutlass_batch_invariance.py
+      - pytest -v -s v1/determinism/test_online_batch_invariance.py
+      - pytest -v -s v1/determinism/test_mx_linear_batch_invariant.py
+
+- label: ":nvidia: (H100) Batch Invariance (Distributed)"
+  key: batch-invariance-distributed-4xh100
+  automatic_shard: true
+  timeout_in_minutes: 45
+  device: h100
+  optional: true
+  num_devices: 4
+  source_file_dependencies:
+    - vllm/distributed/device_communicators/
+    - vllm/model_executor/determinism
+    - vllm/model_executor/layers
+    - vllm/v1/attention
+    - tests/v1/determinism/
+  commands:
+    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+    - export VLLM_TEST_CLEAN_GPU_MEMORY=1
+    - VLLM_TEST_TP_SIZE=4 pytest -v -s v1/determinism/test_batch_invariance.py -k "bs1_vs_bsN or without_batch_invariance_should_fail"
+    - VLLM_TEST_TP_SIZE=4 VLLM_TEST_DISABLE_CUSTOM_AR=1 pytest -v -s v1/determinism/test_batch_invariance.py -k bs1_vs_bsN
+    - VLLM_TP_SIZE=4 VLLM_NEEDLE_TRIALS=2 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k FLASH_ATTN
+    - VLLM_TEST_MODEL=allenai/OLMoE-1B-7B-0924 VLLM_TEST_SERVER_ARGS="--data-parallel-size 4 --enable-expert-parallel --max-model-len 4096" pytest -v -s v1/determinism/test_online_batch_invariance.py
+  mirror:
+    amd:
+      label: ":amd: (MI300) Batch Invariance (Distributed)"
+      dind: false
+      device: mi300_4
+      timeout_in_minutes: 45
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/distributed/device_communicators/
+      - vllm/model_executor/determinism
+      - vllm/model_executor/layers
+      - vllm/v1/attention
+      - tests/v1/determinism/
+      - vllm/platforms/rocm.py
+      commands:
+      - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+      - export VLLM_TEST_CLEAN_GPU_MEMORY=1
+      - VLLM_TEST_TP_SIZE=4 pytest -v -s v1/determinism/test_batch_invariance.py -k "bs1_vs_bsN or without_batch_invariance_should_fail"
+      - VLLM_TEST_TP_SIZE=4 VLLM_TEST_DISABLE_CUSTOM_AR=1 pytest -v -s v1/determinism/test_batch_invariance.py -k bs1_vs_bsN
+      - VLLM_TP_SIZE=4 VLLM_NEEDLE_TRIALS=2 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle -k TRITON_ATTN
+      - VLLM_TEST_MODEL=allenai/OLMoE-1B-7B-0924 VLLM_TEST_SERVER_ARGS="--data-parallel-size 4 --enable-expert-parallel --max-model-len 4096" pytest -v -s v1/determinism/test_online_batch_invariance.py
 
 - label: ":nvidia: (H200 MIG 35GB) Acceptance Length (Large Models)"
   device: h200_35gb
@@ -721,7 +841,6 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/v1/spec_decode/
-  - vllm/model_executor/models/mlp_speculator.py
   - tests/v1/spec_decode/test_acceptance_length.py
   commands:
     - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
@@ -737,6 +856,5 @@ steps:
       working_dir: /vllm-workspace/tests
       source_file_dependencies:
       - vllm/v1/spec_decode/
-      - vllm/model_executor/models/mlp_speculator.py
       - tests/v1/spec_decode/test_acceptance_length.py
       - vllm/platforms/rocm.py

.buildkite/test_areas/model_executor.yaml

@@ -5,6 +5,7 @@ steps:
 - label: ":nvidia: (H200 MIG 35GB) Model Executor"
   device: h200_35gb
   key: model-executor
+  automatic_shard: true
   timeout_in_minutes: 60
   source_file_dependencies:
   - vllm/engine/arg_utils.py

.buildkite/test_areas/models_basic.yaml

@@ -58,6 +58,7 @@ steps:
 - label: ":nvidia: (H200 MIG 35GB) Basic Models (Other)"
   device: h200_35gb
   key: basic-models-tests-other
+  automatic_shard: true
   timeout_in_minutes: 38
   source_file_dependencies:
   - vllm/
@@ -82,17 +83,20 @@ steps:
       - image-build-amd
       source_file_dependencies:
       - vllm/models/deepseek_v4/
+      - vllm/models/deepseek_v41/
       - tests/models/test_deepseek_v4_vl_rocm.py
       - tests/models/test_hyv4_rocm.py
       - tests/models/transformers/test_layer_registry_e2e.py
       - tests/models/test_deepseek_v41_replay_start.py
       - tests/models/test_deepseek_v41_decoder_replay_layers.py
       - tests/models/test_deepseek_v41_replay_batch.py
+      - tests/models/test_deepseek_v41_amd_replay_forward.py
       - tests/models/test_deepseek_v4_dspark_rocm.py
       - tests/models/test_deepseek_v4_rocm_wo_a.py
       - tests/models/test_deepseek_v4_rocm_compressor_gemm_fusion.py
+      - tests/models/test_dsv41_mono_decode.py
       commands:
-      - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py models/test_deepseek_v4_vl_rocm.py models/transformers/test_layer_registry_e2e.py models/test_deepseek_v41_replay_start.py models/test_deepseek_v41_decoder_replay_layers.py models/test_deepseek_v41_replay_batch.py models/test_deepseek_v4_dspark_rocm.py models/test_deepseek_v4_rocm_wo_a.py models/test_deepseek_v4_rocm_compressor_gemm_fusion.py
+      - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py models/test_deepseek_v4_vl_rocm.py models/transformers/test_layer_registry_e2e.py models/test_deepseek_v41_replay_start.py models/test_deepseek_v41_decoder_replay_layers.py models/test_deepseek_v41_replay_batch.py models/test_deepseek_v41_amd_replay_forward.py models/test_deepseek_v4_dspark_rocm.py models/test_deepseek_v4_rocm_wo_a.py models/test_deepseek_v4_rocm_compressor_gemm_fusion.py models/test_dsv41_mono_decode.py
       - VLLM_ROCM_USE_AITER=1 pytest -v -s models/test_hyv4_rocm.py -m 'not distributed'
 
 - label: ":nvidia: (B200) Inkling"
@@ -131,7 +135,7 @@ steps:
   device: b200-k8s
   source_file_dependencies:
   - vllm/models/kimi_k3/
-  - csrc/libtorch_stable/kimi_k3/
+  - csrc/kimi_k3/
   - tests/models/kimi_k3/
   - tests/kernels/attention/test_kimi_k3_mla_fused_epilogue.py
   commands:
@@ -155,8 +159,11 @@ steps:
       # cuda.bindings.driver), so it is deliberately left out here.
       - pytest -v -s models/kimi_k3
 
-- label: ":nvidia: (4xB200) Kimi K3 Prefix Cache Shard %N"
+- label: ":nvidia: (4xB200) Kimi K3 Prefix Cache"
   key: kimi-k3-prefix-cache-4xb200
+  # Nightly only: 4-GPU B200 pods wait hours behind smaller jobs on per-PR runs.
+  optional: true
+  automatic_shard: true
   timeout_in_minutes: 75
   device: b200-k8s
   num_devices: 4
@@ -165,6 +172,8 @@ steps:
   - vllm/v1/core/
   - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
   - vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
+  - vllm/distributed/kv_transfer/kv_connector/v1/mooncake/store/
+  - vllm/distributed/mooncake_store.py
   - vllm/v1/simple_kv_offload/
   - vllm/v1/spec_decode/
   - vllm/model_executor/layers/mamba/
@@ -172,8 +181,7 @@ steps:
   - .buildkite/scripts/install-kv-connectors.sh
   commands:
   - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
-  - pytest -v -s models/kimi_k3/test_prefix_cache.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - pytest -v -s models/kimi_k3/test_prefix_cache.py
-  parallelism: 2
 
 - label: ":nvidia: (H200) GLM5Next"
   key: glm5next-unit-tests
@@ -202,12 +210,12 @@ steps:
   device: h200_35gb
   source_file_dependencies:
   - vllm/models/qwen4_exp/
-  - vllm/transformers_utils/configs/qwen4_exp.py
   - vllm/config/compilation.py
   - vllm/v1/attention/backends/short_conv_attn.py
   - tests/models/qwen4_exp/
   commands:
     - pytest -v -s models/qwen4_exp/test_hc_ops.py models/qwen4_exp/test_qsa_prepare.py models/qwen4_exp/test_qsa_reference.py
+    - pytest -v -s models/qwen4_exp/test_ple.py
   mirror:
     amd:
       label: ":amd: (MI355 DPX) Qwen4 Exp"
@@ -228,7 +236,6 @@ steps:
   device: cpu-small
   source_file_dependencies:
   - vllm/models/qwen4_exp/
-  - vllm/transformers_utils/configs/qwen4_exp.py
   - vllm/config/compilation.py
   - vllm/v1/attention/backends/short_conv_attn.py
   - tests/models/qwen4_exp/
@@ -247,7 +254,10 @@ steps:
   - tests/models/test_vision.py
   - tests/models/test_adapters.py
   - tests/models/test_qwen3_5_mtp_config.py
+  - tests/models/test_alice_ai.py
+  - tests/models/test_alice_ai_config.py
   - tests/models/transformers/fusers/
+  - tests/models/transformers/test_dynamic_rope.py
   - tests/model_executor/layers/test_activation.py
   - tests/models/test_language_model_cache_is_weak.py
   - tests/models/deepseek_v32/test_sequence_parallel.py
@@ -256,4 +266,22 @@ steps:
   - tests/models/transformers/test_layer_registry.py
   device: cpu-small
   commands:
-    - pytest -v -s models/test_utils.py models/test_vision.py models/test_adapters.py models/test_qwen3_5_mtp_config.py models/transformers/fusers/ model_executor/layers/test_activation.py models/test_language_model_cache_is_weak.py models/deepseek_v32/test_sequence_parallel.py models/test_deepseek_v4_fi_moe_ep.py models/test_dspark_mla.py models/transformers/test_layer_registry.py
+    - pytest -v -s models/test_utils.py models/test_vision.py models/test_adapters.py models/test_qwen3_5_mtp_config.py models/transformers/fusers/ models/transformers/test_dynamic_rope.py model_executor/layers/test_activation.py models/test_language_model_cache_is_weak.py models/deepseek_v32/test_sequence_parallel.py models/test_deepseek_v4_fi_moe_ep.py models/test_dspark_mla.py models/transformers/test_layer_registry.py models/test_alice_ai.py models/test_alice_ai_config.py
+
+- label: ':amd: (MI355) DeepSeek-V4.1 Mono Decode Numerics'
+  key: dsv41-mono-decode-numerics-amd
+  dind: false
+  device: mi355_4
+  num_devices: 4
+  timeout_in_minutes: 40
+  working_dir: /vllm-workspace/
+  depends_on:
+  - image-build-amd
+  source_file_dependencies:
+  - docker/Dockerfile.rocm_base
+  - vllm/models/deepseek_v41/amd/
+  - tests/models/test_dsv41_mono_numerics.py
+  - tests/models/test_dsv41_mono_layers.py
+  commands:
+  - pytest -v -s tests/models/test_dsv41_mono_numerics.py
+  - pytest -v -s tests/models/test_dsv41_mono_layers.py

.buildkite/test_areas/models_distributed.yaml

@@ -2,6 +2,23 @@ group: Models - Distributed
 depends_on: 
   - image-build
 steps:
+- label: ":nvidia: (L4) Qwen4 Exp SP"
+  key: qwen4-exp-sp-tests-4-gpus
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  device: l4
+  num_devices: 4
+  source_file_dependencies:
+  - vllm/config/parallel.py
+  - vllm/models/common/ops/sequence_parallel.py
+  - vllm/models/qwen4_exp/
+  - vllm/model_executor/layers/vocab_parallel_embedding.py
+  - vllm/model_executor/models/qwen3_next.py
+  - tests/models/qwen4_exp/test_ple.py
+  - tests/models/qwen4_exp/test_tpsp.py
+  commands:
+  - pytest -v -s models/qwen4_exp/test_tpsp.py
+
 - label: ":nvidia: (L4) Distributed Models Shard %N"
   key: distributed-model-tests-2-gpus
   timeout_in_minutes: 35

.buildkite/test_areas/models_language.yaml

@@ -199,20 +199,23 @@ steps:
       - vllm/platforms/rocm.py
       - tests/models/language/generation_ppl_test
 
-- label: ":nvidia: (H200 MIG 35GB) Language Models (Extended Pooling) Shard %N"
+- label: ":nvidia: (H200 MIG 35GB) Language Models (Extended Pooling)"
   device: h200_35gb
   key: language-models-test-extended-pooling
+  automatic_shard: true
   timeout_in_minutes: 120
-  parallelism: 4
   optional: true
   source_file_dependencies:
   - vllm/
   - "!vllm/distributed/kv_transfer/"
   - tests/models/language/pooling
   commands:
-    - pytest -v -s models/language/pooling -m 'not core_model' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+    - pytest -v -s models/language/pooling -m 'not core_model'
   mirror:
     amd:
+      parallelism: 4
+      commands:
+        - "pytest -v -s models/language/pooling -m 'not core_model' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
       label: ":amd: (MI300) Language Models (Extended Pooling) Shard %N"
       dind: false
       device: mi300_1

.buildkite/test_areas/models_multimodal.yaml

@@ -24,6 +24,7 @@ steps:
 
 - label: ":nvidia: (H200 MIG 18GB) Multimodal Models (Standard) 2: qwen3 + gemma"
   key: multi-modal-models-standard-2-qwen3-gemma
+  automatic_shard: true
   timeout_in_minutes: 75
   device: h200_18gb
   source_file_dependencies:

.buildkite/test_areas/quantization.yaml

@@ -2,13 +2,11 @@ group: Quantization
 depends_on:
   - image-build
 steps:
-- label: ":nvidia: (H200 MIG 35GB) Quantization Shard %N"
+- label: ":nvidia: (H200 MIG 35GB) Quantization"
   device: h200_35gb
   key: quantization
+  automatic_shard: true
   timeout_in_minutes: 40
-  parallelism: 4
-  env:
-    VLLM_USE_V2_MODEL_RUNNER: "0"
   source_file_dependencies:
   - csrc/
   - vllm/model_executor/layers/quantization
@@ -16,13 +14,14 @@ steps:
   - tests/quantization
   commands:
   # Pin torchao to the version compatible with the CI PyTorch and CUDA stack.
-  - uv pip install --system torchao==0.17.0 --index-url https://download.pytorch.org/whl/cu130
+  - uv pip install --system torchao==0.18.0 --index-url https://download.pytorch.org/whl/cu130
   - uv pip install --system conch-triton-kernels
   # The SM90-only checkpoint currently contains a removed weight_chan_scale
   # parameter. It was not exercised by the previous L4 job.
-  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py --ignore quantization/test_rocm_moe.py -k 'not test_compressed_tensors_w4a8_fp8' --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py --ignore quantization/test_rocm_moe.py -k 'not test_compressed_tensors_w4a8_fp8'
   mirror:
     amd:
+      parallelism: 4
       label: ":amd: (MI300) Quantization Shard %N"
       dind: false
       device: mi300_1
@@ -41,8 +40,6 @@ steps:
       - tests/rocm/test_moe_weight_replay.py
       - tests/rocm/test_moe_padding.py
       commands:
-      # Use the ROCm default instead of the NVIDIA runner override.
-      - unset VLLM_USE_V2_MODEL_RUNNER
       - uv pip install --system torchao==0.17.0
       - uv pip install --system conch-triton-kernels
       - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py --ignore quantization/test_rocm_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
@@ -130,8 +127,6 @@ steps:
   key: quantized-models-test
   automatic_shard: true
   timeout_in_minutes: 65
-  env:
-    VLLM_USE_V2_MODEL_RUNNER: "0"
   source_file_dependencies:
   - vllm/model_executor/layers/quantization
   - tests/models/quantization
@@ -153,6 +148,4 @@ steps:
       - vllm/platforms/rocm.py
       - vllm/model_executor/model_loader/
       commands:
-      # Use the ROCm default instead of the NVIDIA runner override.
-      - unset VLLM_USE_V2_MODEL_RUNNER
       - pytest -v -s models/quantization

.buildkite/test_areas/spec_decode.yaml

@@ -594,7 +594,6 @@ steps:
       - vllm/model_executor/models/gemma4_mtp.py
       - vllm/model_executor/models/llama.py
       - vllm/model_executor/models/llama_eagle3.py
-      - vllm/model_executor/models/medusa.py
       - vllm/model_executor/models/registry.py
       - tests/evals/gsm8k/
       - tests/utils.py

.claude/skills/model-microbenchmark

@@ -0,0 +1 @@
+../../.agents/skills/model-microbenchmark

.github/CODEOWNERS

@@ -3,13 +3,14 @@
 
 # This lists cover the "core" components of vLLM that require careful review
 /vllm/compilation @zou3519 @youkaichao @ProExpertProg @BoyuanFeng
+/vllm/compilation/breakable_cudagraph.py @ZJY0516
 /vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery @xuechendi @ivanium @wzhao18
 /vllm/distributed/kv_transfer/kv_connector/v1/hisparse @MatthewBonanni @LucasWilkinson
 /vllm/lora @jeejeelee
 /vllm/model_executor/layers/attention @LucasWilkinson @MatthewBonanni
 /vllm/model_executor/layers/fused_moe @mgoin @pavanimajety @zyongye @wzhao18
 /vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety @zyongye
-/vllm/model_executor/layers/mamba @tdoublep @tomeras91
+/vllm/model_executor/layers/mamba @tdoublep @tomeras91 @ZJY0516
 /vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py @tdoublep @ZJY0516 @vadiklyutiy
 /vllm/model_executor/model_loader @22quinn
 /vllm/model_executor/determinism/ @yewentao256
@@ -26,7 +27,7 @@
 # Any change to the VllmConfig changes can have a large user-facing impact,
 # so spam a lot of people
 /vllm/config @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @yewentao256 @ProExpertProg
-/vllm/config/cache.py @heheda12345 @ivanium
+/vllm/config/cache.py @heheda12345 @ivanium @ZJY0516
 
 # Config utils
 /vllm/config/utils.py @hmellor
@@ -34,23 +35,48 @@
 /vllm/utils/argparse_utils.py
 
 # Entrypoints
-/vllm/entrypoints/anthropic @mgoin @DarkLight1337
+/vllm/entrypoints/anthropic @chaunceyjiang @mgoin @DarkLight1337
+/vllm/entrypoints/cohere @DarkLight1337
 /vllm/entrypoints/cli @hmellor @mgoin @DarkLight1337
 /vllm/entrypoints/mcp @heheda12345
-/vllm/entrypoints/openai @aarnphm @chaunceyjiang @DarkLight1337
+/vllm/entrypoints/openai/chat_completion @chaunceyjiang @DarkLight1337
-/vllm/entrypoints/speech_to_text/realtime @njhill
+/vllm/entrypoints/openai/completion @chaunceyjiang @DarkLight1337
+/vllm/entrypoints/openai/models @chaunceyjiang @DarkLight1337
+/vllm/entrypoints/openai/parser @aarnphm @chaunceyjiang @sfeng33
+/vllm/entrypoints/openai/responses @aarnphm @chaunceyjiang @sfeng33
+/vllm/entrypoints/pooling @noooop @DarkLight1337
+/vllm/entrypoints/rl @aoshen02 @njhill
+/vllm/entrypoints/scale_out @aoshen02 @noooop @DarkLight1337
 /vllm/entrypoints/speech_to_text @NickLucche
-/vllm/entrypoints/pooling @noooop
+/vllm/entrypoints/speech_to_text/realtime @njhill @NickLucche
-/vllm/entrypoints/serve/sagemaker @DarkLight1337
 /vllm/entrypoints/serve @njhill
-/vllm/entrypoints/*.py @njhill
+/vllm/entrypoints/serve/sagemaker @DarkLight1337
-/vllm/entrypoints/chat_utils.py @DarkLight1337
 /vllm/entrypoints/common @DarkLight1337
 /vllm/entrypoints/llm.py @DarkLight1337
 
+/tests/entrypoints @robertgshaw2-redhat @aarnphm @NickLucche @AndreasKaratzas
+/tests/entrypoints/anthropic @mgoin @DarkLight1337
+/tests/entrypoints/cohere @DarkLight1337
+/tests/entrypoints/cli @hmellor @mgoin @DarkLight1337
+/tests/entrypoints/openai/chat_completion @chaunceyjiang @DarkLight1337
+/tests/entrypoints/openai/completion @chaunceyjiang @DarkLight1337
+/tests/entrypoints/openai/models @chaunceyjiang @DarkLight1337
+/tests/entrypoints/openai/parser @aarnphm @chaunceyjiang @sfeng33
+/tests/entrypoints/openai/responses @aarnphm @chaunceyjiang @sfeng33
+/tests/entrypoints/pooling @noooop @DarkLight1337
+/tests/entrypoints/rl @aoshen02 @njhill
+/tests/entrypoints/scale_out @noooop @DarkLight1337
+/tests/entrypoints/speech_to_text @NickLucche
+/tests/entrypoints/speech_to_text/realtime @njhill @NickLucche
+/tests/entrypoints/serve @njhill
+/tests/entrypoints/serve/sagemaker @DarkLight1337
+/tests/entrypoints/llm @DarkLight1337
+/tests/entrypoints/llm/test_struct_output_generate.py @mgoin @aarnphm @arpera
+
 # Rust Frontend
 /rust/ @BugenZhao @njhill
 /rust/src/bench @esmeetu
+/rust/src/server/src/routes/inference @aoshen02 @BugenZhao @njhill
 /tools/build_rust.sh @BugenZhao @njhill
 /rust-toolchain.toml @BugenZhao @njhill
 /.buildkite/test_areas/rust* @BugenZhao @njhill
@@ -64,6 +90,12 @@
 /vllm/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
 /vllm/parser @aarnphm @chaunceyjiang @sfeng33 @bbrowning
 
+/tests/tokenizers_ @DarkLight1337 @njhill
+/tests/renderers @DarkLight1337 @njhill
+/tests/reasoning @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/tests/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/tests/tool_use @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+
 # vLLM V1
 /vllm/v1/attention @LucasWilkinson @MatthewBonanni
 /vllm/v1/attention/backend.py @WoosukKwon @zhuohan123 @youkaichao @alexm-redhat @njhill
@@ -71,17 +103,18 @@
 /vllm/v1/attention/backends/flashinfer.py @mgoin @pavanimajety @vadiklyutiy @wzhao18
 /vllm/v1/attention/backends/triton_attn.py @tdoublep
 /vllm/v1/attention/backends/gdn_attn.py @ZJY0516 @vadiklyutiy
-/vllm/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery @ivanium
+/vllm/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery @ivanium @ZJY0516
 /vllm/v1/sample @22quinn @houseroad @njhill
 /vllm/v1/spec_decode @benchislett @luccafong @MatthewBonanni
 /vllm/v1/structured_output @mgoin @aarnphm @benchislett @arpera
-/vllm/v1/kv_cache_interface.py @heheda12345 @ivanium
+/vllm/v1/kv_cache_interface.py @heheda12345 @ivanium @ZJY0516
 /vllm/v1/kv_offload @ApostaC @orozery
 /vllm/v1/simple_kv_offload @ivanium
 /vllm/v1/hisparse @MatthewBonanni @LucasWilkinson
 /vllm/v1/engine @njhill
 /vllm/v1/executor @njhill
 /vllm/v1/worker @njhill
+/vllm/v1/worker/mamba_utils.py @njhill @ZJY0516
 /vllm/v1/worker/kv_connector_model_runner_mixin.py @orozery @NickLucche @ivanium
 /vllm/v1/worker/encoder_cudagraph.py @shen-shanshan
 /vllm/v1/worker/encoder_cudagraph_defs.py @shen-shanshan
@@ -89,6 +122,7 @@
 # Model runner V2
 /vllm/v1/worker/gpu @WoosukKwon @njhill @yewentao256
 /vllm/v1/worker/gpu/kv_connector.py @orozery
+/vllm/v1/worker/gpu/model_states/mamba_hybrid.py @WoosukKwon @njhill @yewentao256 @ZJY0516
 
 # CI & building
 /.buildkite @Harry-Chen @khluu
@@ -101,27 +135,22 @@
 /tests/distributed/test_multi_node_assignment.py @youkaichao
 /tests/distributed/test_pipeline_parallel.py @youkaichao
 /tests/distributed/test_same_node.py @youkaichao
-/tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche @AndreasKaratzas
+/tests/evals @mgoin @vadiklyutiy @AndreasKaratzas @ZJY0516
-/tests/evals @mgoin @vadiklyutiy @AndreasKaratzas
 /tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye @AndreasKaratzas
 /tests/kernels/ir @ProExpertProg @tjtanaa
 /tests/models @DarkLight1337 @ywang96 @AndreasKaratzas
 /tests/multimodal @DarkLight1337 @ywang96 @NickLucche @shen-shanshan
 /tests/multimodal/media @Isotr0py
 /tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye @AndreasKaratzas
-/tests/entrypoints/llm/test_struct_output_generate.py @mgoin @aarnphm @arpera
 /tests/v1/structured_output @mgoin @aarnphm @arpera
-/tests/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery @ivanium
+/tests/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery @ivanium @ZJY0516
 /tests/lora @jeejeelee
-/tests/models/language/generation/test_hybrid.py @tdoublep @tomeras91
+/tests/models/language/generation/test_hybrid.py @tdoublep @tomeras91 @ZJY0516
 /tests/v1/kv_connector/nixl_integration @NickLucche
 /tests/v1/kv_connector @ApostaC @orozery @ivanium @wzhao18
 /tests/v1/kv_offload @ApostaC @orozery
 /tests/v1/simple_kv_offload @ivanium
 /tests/v1/determinism @yewentao256
-/tests/reasoning @aarnphm @chaunceyjiang @sfeng33 @bbrowning
-/tests/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
-/tests/tool_use @aarnphm @chaunceyjiang @sfeng33 @bbrowning
 
 # Transformers modeling backend
 /vllm/model_executor/models/transformers @hmellor
@@ -223,15 +252,12 @@ mkdocs.yaml @hmellor
 /examples/pooling @noooop
 /docs/models/pooling_models @noooop
 /tests/models/*/pooling* @noooop
-/tests/entrypoints/pooling @noooop
 /vllm/config/pooler.py @noooop
 /vllm/model_executor/layers/pooler @noooop
 
 # Reinforcement learning
 /vllm/distributed/weight_transfer @aoshen02
-/vllm/entrypoints/rl @njhill @aoshen02
 /tests/distributed/test_weight_transfer.py @aoshen02
-/tests/entrypoints/rl @DarkLight1337 @robertgshaw2-redhat @njhill @aarnphm @NickLucche @AndreasKaratzas @aoshen02
 /tests/v1/worker/test_gpu_worker_weight_transfer.py @aoshen02
 /examples/rl @aoshen02
 /docs/training/rlhf.md @aoshen02

.github/mergify.yml

@@ -494,8 +494,7 @@ pull_request_rules:
       - files~=^tests/v1/e2e/spec_decode/
       - files~=^examples/features/speculative_decoding/
       - files~=^vllm/model_executor/models/.*eagle.*\.py
-      - files=vllm/model_executor/models/mlp_speculator.py
+      - files~=^vllm/transformers_utils/configs/eagle\.py
-      - files~=^vllm/transformers_utils/configs/(eagle|medusa|mlp_speculator)\.py
   actions:
     label:
       add:

.github/workflows/issue_autolabel.yml

@@ -279,7 +279,6 @@ jobs:
               "speculative-decoding": {
                 keywords: [
                   { term: "eagle",  searchIn: "title" },
-                  { term: "medusa", searchIn: "title" },
                 ],
                 substrings: [
                   { term: "speculative decoding", searchIn: "title" },

.github/workflows/run-ci-command.yml

@@ -18,22 +18,8 @@ jobs:
     if: >-
       github.repository_owner == 'vllm-project' &&
       github.event.issue.pull_request &&
-      (github.event.comment.body == '/ci run' ||
+      (startsWith(github.event.comment.body, '/ci') ||
-      github.event.comment.body == '/ci run all' ||
+      startsWith(github.event.comment.body, '/amd-ci'))
-      github.event.comment.body == '/ci run nightly' ||
-      github.event.comment.body == '/ci run --allow-stale' ||
-      github.event.comment.body == '/ci run all --allow-stale' ||
-      github.event.comment.body == '/ci run nightly --allow-stale' ||
-      github.event.comment.body == '/ci retry' ||
-      github.event.comment.body == '/ci cancel' ||
-      github.event.comment.body == '/amd-ci run' ||
-      github.event.comment.body == '/amd-ci run all' ||
-      github.event.comment.body == '/amd-ci run nightly' ||
-      github.event.comment.body == '/amd-ci run --allow-stale' ||
-      github.event.comment.body == '/amd-ci run all --allow-stale' ||
-      github.event.comment.body == '/amd-ci run nightly --allow-stale' ||
-      github.event.comment.body == '/amd-ci retry' ||
-      github.event.comment.body == '/amd-ci cancel')
     runs-on: [self-hosted, linux, x64, vllm-runners]
     steps:
       - uses: actions/checkout@9c091bb21b7c1c1d1991bb908d89e4e9dddfe3e0 # v7.0.0
@@ -53,4 +39,16 @@ jobs:
           BUILDKITE_ORGANIZATION: vllm
           BUILDKITE_PIPELINE: ci
           CI_TRUSTED_USERS: ${{ vars.CI_TRUSTED_USERS }}
+          GH_BRANCH_UPDATE_TOKEN: ${{ secrets.CI_BRANCH_UPDATE_TOKEN }}
           GH_TOKEN: ${{ github.token }}
+      - name: Report failure on the PR
+        if: failure()
+        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
+        with:
+          script: |
+            const runUrl = `${context.serverUrl}/${context.repo.owner}/${context.repo.repo}/actions/runs/${context.runId}`;
+            await github.rest.issues.createComment({
+              ...context.repo,
+              issue_number: context.issue.number,
+              body: `❌ @${context.payload.comment.user.login}, your CI command failed with an unexpected error, so CI may not have started. See the [workflow run](${runUrl}) for details.`,
+            });

.github/workflows/scripts/run_ci_command.py

@@ -58,9 +58,12 @@ AMD_CI_COMMANDS = frozenset(
 RETRY_COMMANDS = frozenset({COMMAND_RETRY_FAILED, COMMAND_RETRY_AMD_FAILED})
 CANCEL_COMMANDS = frozenset({COMMAND_CANCEL_CI, COMMAND_CANCEL_AMD_CI})
 ALL_CI_COMMANDS = UPSTREAM_CI_COMMANDS | AMD_CI_COMMANDS
+CI_COMMAND_PREFIXES = frozenset({"/ci", "/amd-ci"})
 CI_AUTHORIZED_COMMENT_MARKER = "<!-- vllm-ci-authorized -->"
 READY_LABELS = {"ready", "ready-run-all-tests"}
 TRUSTED_PERMISSIONS = {"admin", "maintain", "write"}
+BRANCH_UPDATE_POLL_SECONDS = 3
+BRANCH_UPDATE_POLL_ATTEMPTS = 40
 ACTIVE_BUILD_STATES = {
     "blocked",
     "creating",
@@ -258,6 +261,13 @@ class GitHubClient:
             raise ApiError(None, "GitHub returned an invalid commits-behind count.")
         return behind
 
+    def update_branch(self, number: int, expected_head_sha: str) -> None:
+        self._request(
+            self._repo_path(f"/pulls/{number}/update-branch"),
+            body={"expected_head_sha": expected_head_sha},
+            method="PUT",
+        )
+
     def list_pulls_for_commit(self, commit: str) -> list[dict[str, Any]]:
         commit = urllib.parse.quote(commit, safe="")
         response = self._request(self._repo_path(f"/commits/{commit}/pulls"))
@@ -461,11 +471,17 @@ class BuildkiteClient:
 
 
 def parse_command(body: str) -> str | None:
+    body = body.strip()
     if body in ALL_CI_COMMANDS:
         return body
     return None
 
 
+def is_ci_command_attempt(body: str) -> bool:
+    words = body.split(maxsplit=1)
+    return bool(words) and words[0].casefold() in CI_COMMAND_PREFIXES
+
+
 def pipeline_for_command(
     command: str,
     *,
@@ -771,10 +787,11 @@ def notify_authorized(
         (
             f"✅ @{author}, CI is now available for this PR.\n\n"
             "- `/ci run` starts upstream CI; `/amd-ci run` starts AMD CI only.\n"
-            "- Your branch must contain every commit currently on its upstream "
+            "- Run commands test your branch with every commit currently on its "
-            "target branch. Merge or rebase onto the latest target branch, then "
+            "upstream target branch. If it is behind, the bot merges the target "
-            "rerun the command. Append `--allow-stale` to a run command to test "
+            "branch into it first (pull before pushing again). Append "
-            "an outdated branch at your own risk.\n"
+            "`--allow-stale` to a run command to test an outdated branch as is, "
+            "at your own risk.\n"
             "- `/ci retry` retries failed jobs in the CI build for the current "
             "PR head. If the current head has no CI build, it starts a new CI "
             "build for the current head containing only jobs that failed in "
@@ -822,10 +839,65 @@ def resolve_workflow_run_pr(
     return find_matching_pr(github.list_pulls_for_commit(head_sha))
 
 
+def merge_base_into_pr(
+    github: GitHubClient,
+    branch_updater: GitHubClient,
+    pr: Mapping[str, Any],
+    command: str,
+    lag: str,
+) -> tuple[dict[str, Any], str]:
+    base_ref = pr["base"]["ref"]
+
+    def merge_failed(reason: str) -> CiPreparationError:
+        return CiPreparationError(
+            f"{lag} The bot could not merge upstream `{base_ref}` into this "
+            f"branch ({reason}). No new CI build was started. Either merge or "
+            f"rebase onto the latest `{base_ref}` yourself and rerun `{command}`, "
+            f"or use `{command}{ALLOW_STALE_SUFFIX}` to test this branch as is."
+        )
+
+    try:
+        branch_updater.update_branch(pr["number"], pr["head"]["sha"])
+        for _ in range(BRANCH_UPDATE_POLL_ATTEMPTS):
+            time.sleep(BRANCH_UPDATE_POLL_SECONDS)
+            current_pr = github.get_pr(pr["number"])
+            if current_pr["head"]["sha"] == pr["head"]["sha"]:
+                continue
+            try:
+                behind = github.get_commits_behind_base(
+                    base_ref, current_pr["head"]["sha"]
+                )
+            except ApiError as error:
+                # A fork's new head is briefly unknown to the base repository.
+                if error.status == 404:
+                    continue
+                raise
+            break
+        else:
+            raise merge_failed("the merge commit did not appear in time")
+        if (
+            current_pr["state"] != "open"
+            or current_pr["base"]["ref"] != base_ref
+            or behind
+        ):
+            raise merge_failed("the PR changed during the merge")
+    except ApiError as error:
+        raise merge_failed(str(error)) from error
+
+    note = (
+        f"ℹ️ {lag} Merged upstream `{base_ref}` into this branch as "
+        f"`{current_pr['head']['sha'][:12]}` before starting CI. "
+        "Pull the branch before pushing more changes."
+    )
+    print(note)
+    return current_pr, note
+
+
 def prepare_pr_for_ci(
     github: GitHubClient,
     pr: Mapping[str, Any],
     command: str,
+    branch_updater: GitHubClient | None = None,
 ) -> tuple[dict[str, Any], str]:
     base_ref = pr["base"]["ref"]
     no_action = "No new CI build was started."
@@ -861,6 +933,8 @@ def prepare_pr_for_ci(
             print(warning)
         return current_pr, warning
 
+    if behind and branch_updater is not None:
+        return merge_base_into_pr(github, branch_updater, current_pr, command, lag)
     if behind:
         raise CiPreparationError(
             f"{lag} Your branch must contain every commit currently on upstream "
@@ -880,6 +954,8 @@ def handle_run_ci(
     command: str,
     github: GitHubClient,
     pr: Mapping[str, Any],
+    branch_updater: GitHubClient | None = None,
+    allow_repeat_run: bool = False,
 ) -> str:
     ci_name = ci_name_for_command(command)
     duplicate_builds = buildkite.list_builds(
@@ -906,7 +982,19 @@ def handle_run_ci(
             f"{ci_name} is already running for this commit: {active_build['web_url']}"
         )
 
-    current_pr, warning = prepare_pr_for_ci(github, pr, command)
+    tested_build = select_latest_build(
+        [build for build in current_builds if build.get("state") != "canceled"],
+        pr["number"],
+    )
+    if tested_build and not allow_repeat_run:
+        raise CiPreparationError(
+            f"{ci_name} already ran on this commit: {tested_build['web_url']}. "
+            "No new CI build was started. Push new changes to run CI again, use "
+            f"`{retry_command_for_command(command)}` to rerun failed jobs, or ask "
+            "a reviewer with write access to rerun it."
+        )
+
+    current_pr, warning = prepare_pr_for_ci(github, pr, command, branch_updater)
 
     build = buildkite.create_build(
         create_build_payload(
@@ -1086,11 +1174,34 @@ def handle_cancel_ci(
     )
 
 
+def reply_unrecognized_command(
+    event: Mapping[str, Any],
+    github: GitHubClient,
+) -> None:
+    if "pull_request" not in event["issue"]:
+        return
+    issue_number = event["issue"]["number"]
+    comment_id = event["comment"]["id"]
+    actor = event["comment"]["user"]["login"]
+    if is_already_handled(github, issue_number, comment_id):
+        print(f"Comment {comment_id} was already handled.")
+        return
+    add_reaction_safely(github, comment_id, "-1")
+    commands = ", ".join(f"`{command}`" for command in sorted(ALL_CI_COMMANDS))
+    github.add_comment(
+        issue_number,
+        f"❌ @{actor}, that is not a recognized CI command, so no CI was "
+        "started. Post a new comment containing only one of: "
+        f"{commands}.\n\n{command_comment_marker(comment_id)}",
+    )
+
+
 def run(
     event: Mapping[str, Any],
     github: GitHubClient,
     buildkite: BuildkiteClient,
     trusted_users_value: str = "",
+    branch_updater: GitHubClient | None = None,
 ) -> None:
     command = parse_command(event["comment"]["body"])
     if not command or "pull_request" not in event["issue"]:
@@ -1153,6 +1264,8 @@ def run(
                 command=command,
                 github=github,
                 pr=pr,
+                branch_updater=branch_updater,
+                allow_repeat_run=is_trusted_permission(permission),
             )
         elif command in RETRY_COMMANDS:
             message = handle_retry_failed(
@@ -1217,8 +1330,11 @@ def main() -> None:
         )
         return
 
-    command = parse_command(event["comment"]["body"])
+    body = event["comment"]["body"]
+    command = parse_command(body)
     if not command:
+        if is_ci_command_attempt(body):
+            reply_unrecognized_command(event, github)
         return
 
     pipeline = pipeline_for_command(
@@ -1231,11 +1347,15 @@ def main() -> None:
         os.environ.get("BUILDKITE_ORGANIZATION", "vllm"),
         pipeline,
     )
+    branch_update_token = os.environ.get("GH_BRANCH_UPDATE_TOKEN", "")
     run(
         event,
         github,
         buildkite,
         os.environ.get("CI_TRUSTED_USERS", ""),
+        GitHubClient(branch_update_token, os.environ["GITHUB_REPOSITORY"])
+        if branch_update_token
+        else None,
     )
 
 

.github/workflows/scripts/test_run_ci_command.py

@@ -32,10 +32,12 @@ from run_ci_command import (
     has_trusted_approval,
     is_active_build,
     is_build_for_pr,
+    is_ci_command_attempt,
     notify_authorized,
     parse_command,
     parse_trusted_users,
     pipeline_for_command,
+    reply_unrecognized_command,
     resolve_workflow_run_pr,
     run,
     select_latest_build,
@@ -142,6 +144,23 @@ class FakeGitHub:
         self.comments.append(body)
 
 
+class FakeBranchUpdater:
+    def __init__(self, github: FakeGitHub, error: ApiError | None = None) -> None:
+        self.calls: list[tuple[int, str]] = []
+        self.error = error
+        self.github = github
+
+    def update_branch(self, number: int, expected_head_sha: str) -> None:
+        self.calls.append((number, expected_head_sha))
+        if self.error:
+            raise self.error
+        self.github.prs[number] = {
+            **self.github.prs[number],
+            "head": {**self.github.prs[number]["head"], "sha": "merge-commit-sha"},
+        }
+        self.github.behind = 0
+
+
 class FakeBuildkite:
     def __init__(
         self,
@@ -340,14 +359,46 @@ class RunCiCommandTest(unittest.TestCase):
         self.assertIsNone(parse_command("/ci run please"))
         self.assertIsNone(parse_command("/ci run all please"))
         self.assertIsNone(parse_command("/ci cancel please"))
-        self.assertIsNone(parse_command(" /ci run"))
         self.assertIsNone(parse_command("/amd-ci run please"))
-        self.assertIsNone(parse_command("/amd-ci retry "))
         self.assertIsNone(parse_command("/AMD-CI run"))
         self.assertIsNone(parse_command("/amdci run"))
         self.assertIsNone(parse_command("/ci run --allow-stale all"))
         self.assertIsNone(parse_command("/amd-ci run --allow-stale please"))
 
+    def test_surrounding_whitespace_is_ignored(self) -> None:
+        self.assertEqual(parse_command("/ci run\r\n"), COMMAND_RUN_CI)
+        self.assertEqual(parse_command(" /ci run "), COMMAND_RUN_CI)
+        self.assertEqual(parse_command("/amd-ci retry\n"), COMMAND_RETRY_AMD_FAILED)
+
+    def test_ci_command_attempts_are_detected(self) -> None:
+        for body in ("/ci", "/ci rn", "/CI run", "/amd-ci run please", " /ci\nrun"):
+            with self.subTest(body=body):
+                self.assertTrue(is_ci_command_attempt(body))
+        for body in ("", "LGTM", "/cifoo", "/amdci run", "please /ci run"):
+            with self.subTest(body=body):
+                self.assertFalse(is_ci_command_attempt(body))
+
+    def test_unrecognized_command_gets_feedback(self) -> None:
+        github = FakeGitHub()
+
+        reply_unrecognized_command(make_event("/ci rn", actor="author"), github)
+
+        self.assertEqual(github.reactions, ["-1"])
+        self.assertEqual(len(github.comments), 1)
+        self.assertIn(
+            "@author, that is not a recognized CI command", github.comments[0]
+        )
+        self.assertIn(f"`{COMMAND_RUN_CI}`", github.comments[0])
+        self.assertIn("<!-- vllm-ci-command:99 -->", github.comments[0])
+
+    def test_unrecognized_command_feedback_is_not_repeated(self) -> None:
+        github = FakeGitHub(comments=["❌ earlier\n\n<!-- vllm-ci-command:99 -->"])
+
+        reply_unrecognized_command(make_event("/ci rn"), github)
+
+        self.assertEqual(len(github.comments), 1)
+        self.assertEqual(github.reactions, [])
+
     def test_commands_select_only_their_configured_pipeline(self) -> None:
         cases = (
             (COMMAND_RUN_CI, "ci"),
@@ -582,6 +633,82 @@ class RunCiCommandTest(unittest.TestCase):
                         )
                         self.assertIn(command, github.comments[0])
 
+    def test_run_merges_base_into_outdated_branch_before_building(self) -> None:
+        github = FakeGitHub(behind=5)
+        updater = FakeBranchUpdater(github)
+        buildkite = FakeBuildkite([[], []])
+
+        with patch("run_ci_command.time.sleep"):
+            run(make_event(COMMAND_RUN_CI), github, buildkite, "", updater)
+
+        self.assertEqual(updater.calls, [(42, "0123456789abcdef")])
+        self.assertEqual(buildkite.created_builds[0]["commit"], "merge-commit-sha")
+        self.assertTrue(github.comments[0].startswith("✅ "))
+        self.assertIn("5 commits behind", github.comments[0])
+        self.assertIn("Merged upstream `main`", github.comments[0])
+
+    def test_only_committers_can_rerun_a_commit_that_already_ran_ci(self) -> None:
+        finished_build = {
+            "created_at": "2026-07-28T01:00:00Z",
+            "pull_request": {"id": 42},
+            "state": "failed",
+            "web_url": "https://buildkite.example/builds/122",
+        }
+        for actor, permission, command, allowed in (
+            ("author", "read", COMMAND_RUN_CI, False),
+            ("author", "read", f"{COMMAND_RUN_CI} --allow-stale", False),
+            ("reviewer", "write", COMMAND_RUN_CI, True),
+        ):
+            with self.subTest(actor=actor, command=command):
+                github = FakeGitHub(
+                    behind=5,
+                    permission=permission,
+                    pr=make_pr(labels=[{"name": "ready"}]),
+                )
+                updater = FakeBranchUpdater(github)
+                buildkite = FakeBuildkite([[], [finished_build]])
+
+                with patch("run_ci_command.time.sleep"):
+                    run(make_event(command, actor), github, buildkite, "", updater)
+
+                self.assertEqual(bool(buildkite.created_builds), allowed)
+                if not allowed:
+                    self.assertEqual(updater.calls, [])
+                    self.assertIn("CI already ran on this commit", github.comments[0])
+                    self.assertIn("`/ci retry`", github.comments[0])
+
+    def test_run_waits_for_fork_merge_commit_to_reach_base_repo(self) -> None:
+        github = FakeGitHub(behind=5)
+        updater = FakeBranchUpdater(github)
+        buildkite = FakeBuildkite([[], []])
+
+        with (
+            patch("run_ci_command.time.sleep"),
+            patch.object(
+                github,
+                "get_commits_behind_base",
+                side_effect=[5, ApiError(404, "Not Found"), 0],
+            ),
+        ):
+            run(make_event(COMMAND_RUN_CI), github, buildkite, "", updater)
+
+        self.assertEqual(buildkite.created_builds[0]["commit"], "merge-commit-sha")
+        self.assertTrue(github.comments[0].startswith("✅ "))
+
+    def test_run_reports_failed_merge_without_building(self) -> None:
+        github = FakeGitHub(behind=5)
+        updater = FakeBranchUpdater(github, error=ApiError(422, "merge conflict"))
+        buildkite = FakeBuildkite([[], []])
+
+        run(make_event(COMMAND_RUN_CI), github, buildkite, "", updater)
+
+        self.assertEqual(buildkite.created_builds, [])
+        self.assertTrue(github.comments[0].startswith("❌ "))
+        self.assertIn("could not merge upstream `main`", github.comments[0])
+        self.assertIn("merge conflict", github.comments[0])
+        self.assertIn("No new CI build was started.", github.comments[0])
+        self.assertIn("`/ci run --allow-stale`", github.comments[0])
+
     def test_allow_stale_preserves_author_permissions_and_run_configuration(
         self,
     ) -> None:

.gitignore

@@ -237,9 +237,6 @@ benchmarks/**/*.json
 actionlint
 shellcheck*/
 
-# Ignore moe/marlin_moe gen code
-csrc/libtorch_stable/moe/marlin_moe_wna16/kernel_*
-
 # Ignore ep_kernels_workspace folder
 ep_kernels_workspace/
 

.pre-commit-config.yaml

@@ -21,7 +21,7 @@ repos:
   rev: v21.1.2
   hooks:
   - id: clang-format
-    exclude: 'csrc/libtorch_stable/moe/topk_softmax_kernels.cu|vllm/third_party/.*'
+    exclude: 'csrc/moe/topk_softmax_kernels.cu|vllm/third_party/.*'
     types_or: [c++, cuda]
     args: [--style=file, --verbose]
 - repo: https://github.com/DavidAnson/markdownlint-cli2
@@ -205,7 +205,7 @@ repos:
     language: python
     types_or: [python, pyi]
     require_serial: true
-    additional_dependencies: ["mypy==1.20.2", regex, types-cachetools, types-setuptools, types-PyYAML, types-requests, types-torch, pydantic]
+    additional_dependencies: ["mypy==1.20.2", regex, types-cachetools, types-setuptools, types-PyYAML, types-requests, types-torch, "pydantic==2.13.5"]
   - id: png-lint
     name: Lint PNG exports from excalidraw
     entry: tools/pre_commit/png-lint.sh

CMakeLists.txt

@@ -68,8 +68,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.13.0")
+set(TORCH_SUPPORTED_VERSION_CUDA "2.14.0")
-set(TORCH_SUPPORTED_VERSION_ROCM "2.13.0")
+set(TORCH_SUPPORTED_VERSION_ROCM "2.14.0")
 # TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
 # version check would always warn. Only treat it as a nightly build when the
 # value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
@@ -363,8 +363,8 @@ endif()
 
 if(VLLM_GPU_LANG STREQUAL "HIP")
   set(VLLM_EXT_SRC
-    "csrc/torch_bindings.cpp"
+    "csrc/rocm/quickreduce/torch_bindings.cpp"
-    "csrc/custom_quickreduce.cu")
+    "csrc/rocm/quickreduce/custom_quickreduce.cu")
 
   message(STATUS "Enabling C extension.")
   define_extension_target(
@@ -374,16 +374,8 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
     SOURCES ${VLLM_EXT_SRC}
     COMPILE_FLAGS ${VLLM_GPU_FLAGS}
     ARCHITECTURES ${VLLM_GPU_ARCHES}
-    INCLUDE_DIRECTORIES ${CUTLASS_INCLUDE_DIR}
-    INCLUDE_DIRECTORIES ${CUTLASS_TOOLS_UTIL_INCLUDE_DIR}
     USE_SABI 3
     WITH_SOABI)
-
-  # If CUTLASS is compiled on NVCC >= 12.5, it by default uses
-  # cudaGetDriverEntryPointByVersion as a wrapper to avoid directly calling the
-  # driver API. This causes problems when linking with earlier versions of CUDA.
-  # Setting this variable sidesteps the issue by calling the driver directly.
-  target_compile_definitions(_C PRIVATE CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
 endif() # _C HIP endif
 
 if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
@@ -419,36 +411,35 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
 
   set(VLLM_STABLE_EXT_SRC
-    "csrc/libtorch_stable/torch_bindings.cpp"
+    "csrc/torch_bindings.cpp"
-    "csrc/libtorch_stable/cuda_view.cu"
+    "csrc/cuda_view.cu"
-    "csrc/libtorch_stable/cuda_utils_kernels.cu"
+    "csrc/cuda_utils_kernels.cu"
-    "csrc/libtorch_stable/activation_kernels.cu"
+    "csrc/activation_kernels.cu"
-    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
+    "csrc/ngram_embedding_kernels.cu"
-    "csrc/libtorch_stable/vocab_parallel_embedding_kernels.cu"
+    "csrc/vocab_parallel_embedding_kernels.cu"
-    "csrc/libtorch_stable/quantization/activation_kernels.cu"
+    "csrc/quantization/activation_kernels.cu"
-    "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
+    "csrc/quantization/w8a8/int8/scaled_quant.cu"
-    "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
+    "csrc/quantization/w8a8/fp8/common.cu"
-    "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
+    "csrc/quantization/w8a8/fp8/per_token_group_quant.cu"
-    "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
+    "csrc/quantization/gptq/q_gemm.cu"
-    "csrc/libtorch_stable/quantization/gptq/q_gemm.cu"
+    "csrc/pos_encoding_kernels.cu"
-    "csrc/libtorch_stable/pos_encoding_kernels.cu"
+    "csrc/fused_qknorm_rope_kernel.cu"
-    "csrc/libtorch_stable/fused_qknorm_rope_kernel.cu"
+    "csrc/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu"
-    "csrc/libtorch_stable/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu"
+    "csrc/layernorm_kernels.cu"
-    "csrc/libtorch_stable/layernorm_kernels.cu"
+    "csrc/layernorm_quant_kernels.cu"
-    "csrc/libtorch_stable/layernorm_quant_kernels.cu"
+    "csrc/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
-    "csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
+    "csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
-    "csrc/libtorch_stable/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
+    "csrc/attention/merge_attn_states.cu"
-    "csrc/libtorch_stable/attention/merge_attn_states.cu"
+    "csrc/sampler.cu"
-    "csrc/libtorch_stable/sampler.cu"
+    "csrc/topk.cu"
-    "csrc/libtorch_stable/topk.cu"
+    "csrc/mamba/selective_scan_fwd.cu"
-    "csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
+    "csrc/cache_kernels.cu"
-    "csrc/libtorch_stable/cache_kernels.cu"
+    "csrc/cache_kernels_fused.cu"
-    "csrc/libtorch_stable/cache_kernels_fused.cu"
+    "csrc/custom_all_gather_reduce_scatter.cu"
-    "csrc/libtorch_stable/custom_all_gather_reduce_scatter.cu"
+    "csrc/custom_all_gather_reduce_scatter_ops.cpp"
-    "csrc/libtorch_stable/custom_all_gather_reduce_scatter_ops.cpp"
+    "csrc/custom_all_reduce.cu"
-    "csrc/libtorch_stable/custom_all_reduce.cu"
+    "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu"
-    "csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu"
+    "csrc/fused_kimi_k3_mla_key_concat_kv_cache_kernel.cu")
-    "csrc/libtorch_stable/fused_kimi_k3_mla_key_concat_kv_cache_kernel.cu")
 
   if(VLLM_GPU_LANG STREQUAL "CUDA" AND
      DEFINED CMAKE_CUDA_COMPILER_VERSION AND
@@ -470,10 +461,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
 
   if(VLLM_GPU_LANG STREQUAL "CUDA")
     list(APPEND VLLM_STABLE_EXT_SRC
-      "csrc/libtorch_stable/hisparse_kernels.cu"
+      "csrc/hisparse_kernels.cu"
-      "csrc/libtorch_stable/attention/dcp_utils/dcp_direct_a2a_lse_reduce.cu"
+      "csrc/attention/dcp_utils/dcp_direct_a2a_lse_reduce.cu"
-      "csrc/libtorch_stable/attention/dcp_utils/dcp_direct_kv_gather.cu"
+      "csrc/attention/dcp_utils/dcp_direct_kv_gather.cu"
-      "csrc/libtorch_stable/attention/dcp_utils/dcp_direct_q_gather.cu")
+      "csrc/attention/dcp_utils/dcp_direct_q_gather.cu")
 
     SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")
 
@@ -508,12 +499,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     FetchContent_MakeAvailable(cutlass)
 
     list(APPEND VLLM_STABLE_EXT_SRC
-      "csrc/libtorch_stable/cutlass_extensions/common.cpp"
+      "csrc/quantization/w8a8/cutlass/scaled_mm_entry.cu"
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
+      "csrc/quantization/fp4/nvfp4_quant_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
+      "csrc/quantization/fp4/nvfp4_scaled_mm_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_entry.cu"
+      "csrc/quantization/awq/gemm_kernels.cu"
-      "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu"
+      "csrc/minimax_reduce_rms_kernel.cu")
-      "csrc/libtorch_stable/minimax_reduce_rms_kernel.cu")
 
     #
     # Machete kernels
@@ -527,7 +517,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       # preselected input type pairs and schedules.
       # Generate sources:
       set(MACHETE_GEN_SCRIPT
-        ${CMAKE_CURRENT_SOURCE_DIR}/csrc/libtorch_stable/quantization/machete/generate.py)
+        ${CMAKE_CURRENT_SOURCE_DIR}/csrc/quantization/machete/generate.py)
       file(MD5 ${MACHETE_GEN_SCRIPT} MACHETE_GEN_SCRIPT_HASH)
 
       message(STATUS "Machete generation script hash: ${MACHETE_GEN_SCRIPT_HASH}")
@@ -560,7 +550,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       endif()
 
       # Add machete generated sources
-      file(GLOB MACHETE_GEN_SOURCES "csrc/libtorch_stable/quantization/machete/generated/*.cu")
+      file(GLOB MACHETE_GEN_SOURCES "csrc/quantization/machete/generated/*.cu")
       list(APPEND VLLM_STABLE_EXT_SRC ${MACHETE_GEN_SOURCES})
 
       # forward compatible
@@ -569,7 +559,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
         CUDA_ARCHS "${MACHETE_ARCHS}")
 
       list(APPEND VLLM_STABLE_EXT_SRC
-        csrc/libtorch_stable/quantization/machete/machete_pytorch.cu)
+        csrc/quantization/machete/machete_pytorch.cu)
       message(STATUS "Building Machete kernels for archs: ${MACHETE_ARCHS}")
     else()
       if (NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0
@@ -587,7 +577,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     # The fused MiniMax-M3 KV insert quantizes NVFP4 with SM100-family-only
     # cvt instructions, so its SM10x/11x code is built for the family targets.
     set(MINIMAX_M3_FUSED_SRC
-      "csrc/libtorch_stable/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu")
+      "csrc/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu")
     list(REMOVE_ITEM VLLM_STABLE_EXT_SRC "${MINIMAX_M3_FUSED_SRC}")
     set_gencode_flags_for_srcs(
       SRCS "${VLLM_STABLE_EXT_SRC}"
@@ -611,9 +601,9 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
 
     if(COOPERATIVE_TOPK_ARCHS)
       list(APPEND VLLM_STABLE_EXT_SRC
-        "csrc/libtorch_stable/cooperative_topk.cu")
+        "csrc/cooperative_topk.cu")
       set_gencode_flags_for_srcs(
-        SRCS "csrc/libtorch_stable/cooperative_topk.cu"
+        SRCS "csrc/cooperative_topk.cu"
         CUDA_ARCHS "${COOPERATIVE_TOPK_ARCHS}")
     endif()
 
@@ -656,7 +646,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       # preselected input type pairs and schedules.
       # Generate sources:
       set(MARLIN_GEN_SCRIPT
-        ${CMAKE_CURRENT_SOURCE_DIR}/csrc/libtorch_stable/quantization/marlin/generate_kernels.py)
+        ${CMAKE_CURRENT_SOURCE_DIR}/csrc/quantization/marlin/generate_kernels.py)
       file(MD5 ${MARLIN_GEN_SCRIPT} MARLIN_GEN_SCRIPT_HASH)
       list(JOIN CUDA_ARCHS "," CUDA_ARCHS_STR)
       set(MARLIN_GEN_SCRIPT_HASH_AND_ARCH "${MARLIN_GEN_SCRIPT_HASH}(ARCH:${CUDA_ARCHS_STR})")
@@ -691,7 +681,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       endif()
 
       if (MARLIN_ARCHS)
-        file(GLOB MARLIN_TEMPLATE_KERNEL_SRC "csrc/libtorch_stable/quantization/marlin/sm80_kernel_*_float16.cu")
+        file(GLOB MARLIN_TEMPLATE_KERNEL_SRC "csrc/quantization/marlin/sm80_kernel_*_float16.cu")
         set_gencode_flags_for_srcs(
           SRCS "${MARLIN_TEMPLATE_KERNEL_SRC}"
           CUDA_ARCHS "${MARLIN_ARCHS}")
@@ -701,7 +691,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
         endif()
         list(APPEND VLLM_STABLE_EXT_SRC ${MARLIN_TEMPLATE_KERNEL_SRC})
 
-        file(GLOB MARLIN_TEMPLATE_BF16_KERNEL_SRC "csrc/libtorch_stable/quantization/marlin/sm80_kernel_*_bfloat16.cu")
+        file(GLOB MARLIN_TEMPLATE_BF16_KERNEL_SRC "csrc/quantization/marlin/sm80_kernel_*_bfloat16.cu")
         set_gencode_flags_for_srcs(
           SRCS "${MARLIN_TEMPLATE_BF16_KERNEL_SRC}"
           CUDA_ARCHS "${MARLIN_BF16_ARCHS}")
@@ -713,7 +703,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       endif()
 
       if (MARLIN_SM75_ARCHS)
-        file(GLOB MARLIN_TEMPLATE_SM75_KERNEL_SRC "csrc/libtorch_stable/quantization/marlin/sm75_kernel_*.cu")
+        file(GLOB MARLIN_TEMPLATE_SM75_KERNEL_SRC "csrc/quantization/marlin/sm75_kernel_*.cu")
         set_gencode_flags_for_srcs(
           SRCS "${MARLIN_TEMPLATE_SM75_KERNEL_SRC}"
           CUDA_ARCHS "${MARLIN_SM75_ARCHS}")
@@ -725,7 +715,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       endif()
 
       if (MARLIN_FP8_ARCHS)
-        file(GLOB MARLIN_TEMPLATE_FP8_KERNEL_SRC "csrc/libtorch_stable/quantization/marlin/sm89_kernel_*.cu")
+        file(GLOB MARLIN_TEMPLATE_FP8_KERNEL_SRC "csrc/quantization/marlin/sm89_kernel_*.cu")
         set_gencode_flags_for_srcs(
           SRCS "${MARLIN_TEMPLATE_FP8_KERNEL_SRC}"
           CUDA_ARCHS "${MARLIN_FP8_ARCHS}")
@@ -737,10 +727,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       endif()
 
       set(MARLIN_SRCS
-         "csrc/libtorch_stable/quantization/marlin/marlin.cu"
+         "csrc/quantization/marlin/marlin.cu"
-         "csrc/libtorch_stable/quantization/marlin/marlin_int4_fp8_preprocess.cu"
+         "csrc/quantization/marlin/marlin_int4_fp8_preprocess.cu"
-         "csrc/libtorch_stable/quantization/marlin/gptq_marlin_repack.cu"
+         "csrc/quantization/marlin/gptq_marlin_repack.cu"
-         "csrc/libtorch_stable/quantization/marlin/awq_marlin_repack.cu")
+         "csrc/quantization/marlin/awq_marlin_repack.cu")
       set_gencode_flags_for_srcs(
         SRCS "${MARLIN_SRCS}"
         CUDA_ARCHS "${MARLIN_OTHER_ARCHS}")
@@ -763,7 +753,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
-    set(DSV3_FUSED_A_GEMM_SRCS "csrc/libtorch_stable/dsv3_fused_a_gemm.cu")
+    set(DSV3_FUSED_A_GEMM_SRCS "csrc/dsv3_fused_a_gemm.cu")
     set_gencode_flags_for_srcs(
       SRCS "${DSV3_FUSED_A_GEMM_SRCS}"
       CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
@@ -778,8 +768,8 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   cuda_archs_sm90plus(FP32_ROUTER_GEMM_ARCHS "${CUDA_ARCHS}")
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND FP32_ROUTER_GEMM_ARCHS)
     set(FP32_ROUTER_GEMM_SRCS
-      "csrc/libtorch_stable/fp32_router_gemm_entry.cu"
+      "csrc/fp32_router_gemm_entry.cu"
-      "csrc/libtorch_stable/fp32_router_gemm.cu")
+      "csrc/fp32_router_gemm.cu")
     set_gencode_flags_for_srcs(
       SRCS "${FP32_ROUTER_GEMM_SRCS}"
       CUDA_ARCHS "${FP32_ROUTER_GEMM_ARCHS}")
@@ -799,11 +789,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   cuda_archs_loose_intersection(SCALED_MM_ARCHS "9.0a;" "${CUDA_ARCHS}")
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND SCALED_MM_ARCHS)
     set(SCALED_MM_SM90_SRCS
-       "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c3x_sm90.cu"
+       "csrc/quantization/w8a8/cutlass/scaled_mm_c3x_sm90.cu"
-       "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_fp8.cu"
+       "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_fp8.cu"
-       "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_int8.cu"
+       "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_int8.cu"
-       "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_azp_sm90_int8.cu"
+       "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_azp_sm90_int8.cu"
-       "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm90_fp8.cu")
+       "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm90_fp8.cu")
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_SM90_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
@@ -834,9 +824,9 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
     set(SCALED_MM_SM120_SRCS
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c3x_sm120.cu"
+      "csrc/quantization/w8a8/cutlass/scaled_mm_c3x_sm120.cu"
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8.cu"
+      "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8.cu"
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm120_fp8.cu"
+      "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm120_fp8.cu"
     )
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_SM120_SRCS}"
@@ -868,9 +858,9 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
     set(SCALED_MM_SM100_SRCS
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c3x_sm100.cu"
+      "csrc/quantization/w8a8/cutlass/scaled_mm_c3x_sm100.cu"
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm100_fp8.cu"
+      "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm100_fp8.cu"
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm100_fp8.cu"
+      "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm100_fp8.cu"
     )
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_SM100_SRCS}"
@@ -901,7 +891,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # subtract out the archs that are already built for 3x
   list(REMOVE_ITEM SCALED_MM_2X_ARCHS ${SCALED_MM_3X_ARCHS})
   if (SCALED_MM_2X_ARCHS)
-    set(SCALED_MM_C2X_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c2x.cu")
+    set(SCALED_MM_C2X_SRCS "csrc/quantization/w8a8/cutlass/scaled_mm_c2x.cu")
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_C2X_SRCS}"
       CUDA_ARCHS "${SCALED_MM_2X_ARCHS}")
@@ -927,7 +917,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # if it's possible to compile MoE kernels that use its output.
   cuda_archs_loose_intersection(SCALED_MM_ARCHS "9.0a" "${CUDA_ARCHS}")
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.3 AND SCALED_MM_ARCHS)
-    set(CUTLASS_MOE_SM90_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm90.cu")
+    set(CUTLASS_MOE_SM90_SRCS "csrc/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm90.cu")
     set_gencode_flags_for_srcs(
       SRCS "${CUTLASS_MOE_SM90_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
@@ -951,7 +941,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
-    set(CUTLASS_MOE_SM100_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm100.cu")
+    set(CUTLASS_MOE_SM100_SRCS "csrc/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm100.cu")
     set_gencode_flags_for_srcs(
       SRCS "${CUTLASS_MOE_SM100_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
@@ -976,7 +966,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     cuda_archs_loose_intersection(CUTLASS_MOE_DATA_ARCHS "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.3 AND CUTLASS_MOE_DATA_ARCHS)
-    set(CUTLASS_MOE_DATA_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/moe_data.cu")
+    set(CUTLASS_MOE_DATA_SRCS "csrc/quantization/w8a8/cutlass/moe/moe_data.cu")
     set_gencode_flags_for_srcs(
       SRCS "${CUTLASS_MOE_DATA_SRCS}"
       CUDA_ARCHS "${CUTLASS_MOE_DATA_ARCHS}")
@@ -1008,12 +998,12 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_SM120_ARCHS)
     set(FP4_SM120_SRCS
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_quant_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
+      "csrc/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
+      "csrc/quantization/fp4/nvfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
+      "csrc/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
+      "csrc/nvfp4_kv_cache_kernels.cu")
     set_gencode_flags_for_srcs(
       SRCS "${FP4_SM120_SRCS}"
       CUDA_ARCHS "${FP4_SM120_ARCHS}")
@@ -1034,15 +1024,15 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_SM100_ARCHS)
     set(FP4_SM100_SRCS
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_quant_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
+      "csrc/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
+      "csrc/quantization/fp4/nvfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
+      "csrc/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu"
+      "csrc/quantization/fp4/mxfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu"
+      "csrc/quantization/fp4/mxfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu"
+      "csrc/nvfp4_kv_cache_kernels.cu"
-      "csrc/libtorch_stable/nvfp4_ds_mla_cache_kernels.cu")
+      "csrc/nvfp4_ds_mla_cache_kernels.cu")
     if(NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9)
       message(STATUS
         "Building mxfp4_experts_quant unsupported stubs because CUDA compiler version is not >= 12.9 (found ${CMAKE_CUDA_COMPILER_VERSION}).")
@@ -1066,9 +1056,9 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   cuda_archs_loose_intersection(W4A8_ARCHS "9.0a" "${CUDA_ARCHS}")
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND W4A8_ARCHS)
     set(W4A8_SRCS
-       "csrc/libtorch_stable/quantization/cutlass_w4a8/w4a8_mm_entry.cu"
+       "csrc/quantization/cutlass_w4a8/w4a8_mm_entry.cu"
-       "csrc/libtorch_stable/quantization/cutlass_w4a8/w4a8_grouped_mm_entry.cu"
+       "csrc/quantization/cutlass_w4a8/w4a8_grouped_mm_entry.cu"
-       "csrc/libtorch_stable/quantization/cutlass_w4a8/w4a8_utils.cu"
+       "csrc/quantization/cutlass_w4a8/w4a8_utils.cu"
        )
 
     set_gencode_flags_for_srcs(
@@ -1101,7 +1091,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
     set(CUTLASS_MLA_SRCS
-      "csrc/libtorch_stable/attention/mla/sm100_cutlass_mla_kernel.cu")
+      "csrc/attention/mla/sm100_cutlass_mla_kernel.cu")
     set_gencode_flags_for_srcs(
       SRCS "${CUTLASS_MLA_SRCS}"
       CUDA_ARCHS "${MLA_ARCHS}")
@@ -1125,7 +1115,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(FUSED_KDA_DECODE_ARCHS)
     set(FUSED_KDA_DECODE_SRC
-      "csrc/libtorch_stable/kimi_k3/fused_kda_decode_kernel.cu")
+      "csrc/kimi_k3/fused_kda_decode_kernel.cu")
     set_gencode_flags_for_srcs(
       SRCS "${FUSED_KDA_DECODE_SRC}"
       CUDA_ARCHS "${FUSED_KDA_DECODE_ARCHS}")
@@ -1137,7 +1127,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(FUSED_GDN_DECODE_ARCHS)
     set(FUSED_GDN_DECODE_SRC
-      "csrc/libtorch_stable/gdn/fused_gdn_decode_kernel.cu")
+      "csrc/gdn/fused_gdn_decode_kernel.cu")
     set_gencode_flags_for_srcs(
       SRCS "${FUSED_GDN_DECODE_SRC}"
       CUDA_ARCHS "${FUSED_GDN_DECODE_ARCHS}")
@@ -1154,7 +1144,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(KIMI_K3_ATTN_RES_ARCHS)
     set(KIMI_K3_ATTN_RES_SRC
-      "csrc/libtorch_stable/kimi_k3/attn_res_kernel.cu")
+      "csrc/kimi_k3/attn_res_kernel.cu")
     set_gencode_flags_for_srcs(
       SRCS "${KIMI_K3_ATTN_RES_SRC}"
       CUDA_ARCHS "${KIMI_K3_ATTN_RES_ARCHS}")
@@ -1169,7 +1159,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # Hadacore kernels
   cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
   if(HADACORE_ARCHS)
-    set(HADACORE_SRCS "csrc/libtorch_stable/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
+    set(HADACORE_SRCS "csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
     set_gencode_flags_for_srcs(
       SRCS "${HADACORE_SRCS}"
       CUDA_ARCHS "${HADACORE_ARCHS}")
@@ -1193,7 +1183,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     list(FILTER FUSED_KDA_DECODE_HIP_ARCHS INCLUDE REGEX "gfx942|gfx950")
     if(FUSED_KDA_DECODE_HIP_ARCHS)
       set(FUSED_KDA_DECODE_HIP_SRC
-        "csrc/libtorch_stable/kimi_k3/fused_kda_decode_kernel_rocm.cu")
+        "csrc/kimi_k3/fused_kda_decode_kernel_rocm.cu")
       list(APPEND VLLM_STABLE_EXT_SRC "${FUSED_KDA_DECODE_HIP_SRC}")
       message(STATUS
         "Building fused KDA decode for HIP archs: ${FUSED_KDA_DECODE_HIP_ARCHS}")
@@ -1203,7 +1193,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     list(FILTER FUSED_KDA_CHUNK_HIP_ARCHS INCLUDE REGEX "gfx950")
     if(FUSED_KDA_CHUNK_HIP_ARCHS)
       set(FUSED_KDA_CHUNK_HIP_SRC
-        "csrc/libtorch_stable/kimi_k3/fused_kda_chunk_kernel_rocm.cu")
+        "csrc/kimi_k3/fused_kda_chunk_kernel_rocm.cu")
       list(APPEND VLLM_STABLE_EXT_SRC "${FUSED_KDA_CHUNK_HIP_SRC}")
       message(STATUS
         "Building fused KDA chunk for HIP archs: ${FUSED_KDA_CHUNK_HIP_ARCHS}")
@@ -1289,21 +1279,22 @@ endif()
 #
 
 set(VLLM_MOE_EXT_SRC
-  "csrc/libtorch_stable/moe/torch_bindings.cpp"
+  "csrc/moe/torch_bindings.cpp"
-  "csrc/libtorch_stable/moe/moe_align_sum_kernels.cu"
+  "csrc/moe/moe_align_sum_kernels.cu"
-  "csrc/libtorch_stable/moe/topk_softmax_kernels.cu"
+  "csrc/moe/topk_softmax_kernels.cu"
-  "csrc/libtorch_stable/moe/topk_softplus_sqrt_kernels.cu")
+  "csrc/moe/topk_softplus_sqrt_kernels.cu")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   list(APPEND VLLM_MOE_EXT_SRC
-    "csrc/libtorch_stable/moe/moe_wna16.cu"
+    "csrc/moe/moe_wna16.cu"
-    "csrc/libtorch_stable/moe/grouped_topk_kernels.cu")
+    "csrc/moe/grouped_topk_kernels.cu"
+    "csrc/moe/fused_globalize_align_block_size.cu")
 endif()
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   set(MOE_PERMUTE_SRC
-      "csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.cu"
+      "csrc/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.cu"
-      "csrc/libtorch_stable/moe/moe_permute_unpermute_op.cu")
+      "csrc/moe/moe_permute_unpermute_op.cu")
 
   list(APPEND VLLM_MOE_EXT_SRC "${MOE_PERMUTE_SRC}")
 endif()
@@ -1314,7 +1305,7 @@ set_gencode_flags_for_srcs(
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   set(VLLM_MOE_WNA16_SRC
-    "csrc/libtorch_stable/moe/moe_wna16.cu")
+    "csrc/moe/moe_wna16.cu")
 
   set_gencode_flags_for_srcs(
     SRCS "${VLLM_MOE_WNA16_SRC}"
@@ -1349,7 +1340,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     # preselected input type pairs and schedules.
     # Generate sources:
     set(MOE_MARLIN_GEN_SCRIPT
-      ${CMAKE_CURRENT_SOURCE_DIR}/csrc/libtorch_stable/moe/marlin_moe_wna16/generate_kernels.py)
+      ${CMAKE_CURRENT_SOURCE_DIR}/csrc/quantization/marlin/generate_kernels.py)
     file(MD5 ${MOE_MARLIN_GEN_SCRIPT} MOE_MARLIN_GEN_SCRIPT_HASH)
     list(JOIN CUDA_ARCHS "," CUDA_ARCHS_STR)
     set(MOE_MARLIN_GEN_SCRIPT_HASH_AND_ARCH "${MOE_MARLIN_GEN_SCRIPT_HASH}(ARCH:${CUDA_ARCHS_STR})")
@@ -1363,6 +1354,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
         COMMAND ${CMAKE_COMMAND} -E env
         PYTHONPATH=$ENV{PYTHONPATH}
           ${Python_EXECUTABLE} ${MOE_MARLIN_GEN_SCRIPT} ${CUDA_ARCHS_STR}
+          ${CMAKE_CURRENT_SOURCE_DIR}/csrc/moe/marlin_moe_wna16
         RESULT_VARIABLE moe_marlin_generation_result
         OUTPUT_VARIABLE moe_marlin_generation_output
         OUTPUT_FILE ${CMAKE_CURRENT_BINARY_DIR}/moe_marlin_generation.log
@@ -1384,7 +1376,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
 
     if (MARLIN_MOE_ARCHS)
-      file(GLOB MARLIN_MOE_SRC "csrc/libtorch_stable/moe/marlin_moe_wna16/sm80_kernel_*.cu")
+      file(GLOB MARLIN_MOE_SRC "csrc/moe/marlin_moe_wna16/sm80_kernel_*.cu")
       set_gencode_flags_for_srcs(
         SRCS "${MARLIN_MOE_SRC}"
         CUDA_ARCHS "${MARLIN_MOE_ARCHS}")
@@ -1396,7 +1388,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
 
     if (MARLIN_MOE_SM75_ARCHS)
-      file(GLOB MARLIN_MOE_SM75_SRC "csrc/libtorch_stable/moe/marlin_moe_wna16/sm75_kernel_*.cu")
+      file(GLOB MARLIN_MOE_SM75_SRC "csrc/moe/marlin_moe_wna16/sm75_kernel_*.cu")
       set_gencode_flags_for_srcs(
         SRCS "${MARLIN_MOE_SM75_SRC}"
         CUDA_ARCHS "${MARLIN_MOE_SM75_ARCHS}")
@@ -1408,7 +1400,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
 
     if (MARLIN_MOE_FP8_ARCHS)
-      file(GLOB MARLIN_MOE_FP8_SRC "csrc/libtorch_stable/moe/marlin_moe_wna16/sm89_kernel_*.cu")
+      file(GLOB MARLIN_MOE_FP8_SRC "csrc/moe/marlin_moe_wna16/sm89_kernel_*.cu")
       set_gencode_flags_for_srcs(
         SRCS "${MARLIN_MOE_FP8_SRC}"
         CUDA_ARCHS "${MARLIN_MOE_FP8_ARCHS}")
@@ -1419,7 +1411,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
       list(APPEND VLLM_MOE_EXT_SRC ${MARLIN_MOE_FP8_SRC})
     endif()
 
-    set(MARLIN_MOE_OTHER_SRC "csrc/libtorch_stable/moe/marlin_moe_wna16/ops.cu")
+    set(MARLIN_MOE_OTHER_SRC "csrc/moe/marlin_moe_wna16/ops.cu")
     set_gencode_flags_for_srcs(
       SRCS "${MARLIN_MOE_OTHER_SRC}"
       CUDA_ARCHS "${MARLIN_MOE_OTHER_ARCHS}")

README.md

@@ -67,6 +67,7 @@ Install vLLM with [`uv`](https://docs.astral.sh/uv/) (recommended) or `pip`:
 
 ```bash
 uv pip install vllm
+vllm download-kernels  # precompiled kernels for faster startup
 ```
 
 Or [build from source](https://docs.vllm.ai/en/latest/getting_started/installation/gpu/index.html#build-wheel-from-source) for development.

benchmarks/benchmark_topk_topp.py

@@ -1,23 +1,29 @@
 #!/usr/bin/env python3
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Benchmark comparing Triton vs PyTorch sort-based top-k/top-p implementations.
+"""Benchmark top-k/top-p filtering with Triton, FlashInfer, and optional PyTorch.
 
-Compares:
+All implementations return masked logits. Timings include a fresh-input copy.
-- apply_top_k_top_p_triton (Triton binary search)
+FlashInfer timings include softmax and mask reconstruction when its APIs require them.
-- apply_top_k_top_p (PyTorch sort-based)
+L2 is flushed before each measured iteration, outside the timed region.
 
 Scenarios:
 - top_k only (whole batch, partial batch)
 - top_p only (whole batch, partial batch)
 - mix of top_k and top_p
+- optional sparse inputs via --num-finite-logits
 """
 
 import argparse
 import gc
+import json
 from dataclasses import dataclass
+from pathlib import Path
 
+import flashinfer.sampling as fi_sampling
+import pandas as pd
 import torch
+from flashinfer.testing import bench_gpu_time_with_cupti
 
 from vllm.v1.sample.ops.topk_topp_sampler import apply_top_k_top_p_pytorch
 from vllm.v1.sample.ops.topk_topp_triton import (
@@ -65,14 +71,18 @@ def calculate_ops_pct(
 
 
 def create_logits(
-    batch_size: int, vocab_size: int, device: str = "cuda"
+    batch_size: int,
+    vocab_size: int,
+    device: str = "cuda",
+    num_finite_logits: int | None = None,
 ) -> torch.Tensor:
     """Create random logits mimicking a realistic LLM distribution.
 
     Uses a Zipf-like probability distribution (rank^-1.1) converted to logits
     via log, then randomly permuted per row. This produces a peaked distribution
     where a small number of tokens capture most probability mass, similar to
-    real model outputs.
+    real model outputs. When num_finite_logits is set, retain that many ranks
+    before permuting their positions independently in each row.
     """
     # Create Zipf-like probabilities: p(rank) ~ rank^(-alpha)
     ranks = torch.arange(1, vocab_size + 1, dtype=torch.float32, device=device)
@@ -81,6 +91,8 @@ def create_logits(
 
     # Convert to logits (log-probabilities, unnormalized is fine)
     base_logits = probs.log()
+    if num_finite_logits is not None:
+        base_logits[num_finite_logits:] = -float("inf")
 
     # Broadcast to batch and randomly permute each row
     logits = base_logits.unsqueeze(0).expand(batch_size, -1).clone()
@@ -90,6 +102,26 @@ def create_logits(
     return logits
 
 
+def apply_top_k_top_p_flashinfer_mask(logits, k, p):
+    """Keep top-k in logits; adapt probability-based top-p back to logits."""
+    if k is not None:
+        logits = fi_sampling.top_k_mask_logits(logits, k)
+    if p is not None:
+        # Top-p renormalization requires softmax, even after logits-based top-k.
+        probs = logits.softmax(dim=-1)
+        probs = fi_sampling.top_p_renorm_probs(probs, p, is_deterministic=True)
+        # Return the selected original logits, not normalized probabilities.
+        logits = logits.masked_fill(probs == 0, -float("inf"))
+    return logits
+
+
+IMPLEMENTATIONS = {
+    "triton": apply_top_k_top_p_triton,
+    "flashinfer-mask": apply_top_k_top_p_flashinfer_mask,
+    "pytorch": apply_top_k_top_p_pytorch,
+}
+
+
 def measure_memory() -> tuple[int, int]:
     """Return (allocated, reserved) memory in bytes."""
     torch.accelerator.synchronize()
@@ -107,53 +139,6 @@ def reset_memory_stats():
     gc.collect()
 
 
-def benchmark_function(
-    func,
-    logits: torch.Tensor,
-    k: torch.Tensor | None,
-    p: torch.Tensor | None,
-    warmup_iters: int = 5,
-    benchmark_iters: int = 20,
-) -> tuple[float, int]:
-    """Benchmark a function and return (avg_time_ms, peak_memory_bytes).
-
-    Returns average time in milliseconds and peak memory usage.
-    """
-    # Warmup
-    for _ in range(warmup_iters):
-        logits_copy = logits.clone()
-        func(logits_copy, k, p)
-    torch.accelerator.synchronize()
-
-    # Reset memory stats before benchmark
-    reset_memory_stats()
-
-    # Benchmark
-    start_events = [
-        torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)
-    ]
-    end_events = [torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)]
-
-    for i in range(benchmark_iters):
-        logits_copy = logits.clone()
-        start_events[i].record()
-        func(logits_copy, k, p)
-        end_events[i].record()
-
-    torch.accelerator.synchronize()
-
-    # Calculate timing
-    times = [
-        start_events[i].elapsed_time(end_events[i]) for i in range(benchmark_iters)
-    ]
-    avg_time = sum(times) / len(times)
-
-    # Get peak memory
-    _, peak_memory = measure_memory()
-
-    return avg_time, peak_memory
-
-
 def create_benchmark_configs(
     batch_sizes: list[int],
     vocab_sizes: list[int],
@@ -286,118 +271,103 @@ def format_memory(bytes_val: int) -> str:
 
 def run_benchmark(
     configs: list[BenchmarkConfig],
+    implementations: list[str],
     warmup_iters: int = 5,
     benchmark_iters: int = 20,
     verbose: bool = True,
+    num_finite_logits: int | None = None,
+    check: bool = False,
+    output: Path | None = None,
 ):
-    """Run all benchmarks and print results."""
+    """Run selected implementations with identical inputs and masked-logit outputs."""
     results = []
 
-    print("=" * 100)
+    for index, config in enumerate(configs):
-    print("Top-k/Top-p Benchmark: Triton vs PyTorch Sort-based")
+        logits = create_logits(
-    print("=" * 100)
+            config.batch_size,
-    print()
+            config.vocab_size,
-
+            num_finite_logits=num_finite_logits,
-    for config in configs:
-        if verbose:
-            print(f"Running: {config.description}")
-
-        # Create fresh logits for this config
-        logits = create_logits(config.batch_size, config.vocab_size)
-
-        # Benchmark Triton
-        reset_memory_stats()
-        triton_time, triton_mem = benchmark_function(
-            apply_top_k_top_p_triton,
-            logits,
-            config.k_values,
-            config.p_values,
-            warmup_iters,
-            benchmark_iters,
         )
-
+        checks = {}
-        # Benchmark PyTorch
+        if check:
-        reset_memory_stats()
+            checks = check_implementations(logits, config, implementations)
-        pytorch_time, pytorch_mem = benchmark_function(
+        k, p = config.k_values, config.p_values
-            apply_top_k_top_p_pytorch,
+        timings = {}
-            logits,
+        for name in implementations:
-            config.k_values,
+            func = IMPLEMENTATIONS[name]
-            config.p_values,
+            reset_memory_stats()
-            warmup_iters,
+            times = bench_gpu_time_with_cupti(
-            benchmark_iters,
+                # note: measured runtime includes logits.clone()
-        )
+                lambda func=func, logits=logits, k=k, p=p: func(logits.clone(), k, p),
-
+                dry_run_iters=warmup_iters,
-        speedup = pytorch_time / triton_time if triton_time > 0 else float("inf")
+                repeat_iters=benchmark_iters,
-        mem_ratio = pytorch_mem / triton_mem if triton_mem > 0 else float("inf")
+                use_cuda_graph=False,  # sampling is outside cuda graph
-
+                cold_l2_cache=True,
-        result = {
+            )
-            "config": config,
+            elapsed = sum(times) / len(times)
-            "triton_time_ms": triton_time,
+            _, memory = measure_memory()
-            "pytorch_time_ms": pytorch_time,
+            timings[name] = dict(time_ms=elapsed, peak_memory=memory)
-            "triton_mem": triton_mem,
+        row = create_result_row(config, timings, checks)
-            "pytorch_mem": pytorch_mem,
+        results.append(row)
-            "speedup": speedup,
+        if output is not None:
-            "mem_ratio": mem_ratio,
+            with output.open("a") as file:
-        }
+                file.write(json.dumps(row) + "\n")
-        results.append(result)
-
         if verbose:
-            print(f"  Triton:  {triton_time:.3f} ms, {format_memory(triton_mem)}")
+            print(f"Completed {index + 1}/{len(configs)}: {config.name}", flush=True)
-            print(f"  PyTorch: {pytorch_time:.3f} ms, {format_memory(pytorch_mem)}")
-            print(f"  Speedup: {speedup:.2f}x, Memory ratio: {mem_ratio:.2f}x")
-            print()
-
-        # Clean up
         del logits
         reset_memory_stats()
 
     return results
 
 
+def check_implementations(
+    logits: torch.Tensor,
+    config: BenchmarkConfig,
+    implementations: list[str],
+) -> dict:
+    """Compare filtered distributions against PyTorch outside the timed region."""
+    k, p = config.k_values, config.p_values
+    reference = apply_top_k_top_p_pytorch(logits.clone(), k, p).softmax(dim=-1)
+    checks = {}
+    for name in implementations:
+        actual = IMPLEMENTATIONS[name](logits.clone(), k, p).softmax(dim=-1)
+        total_variation = (actual - reference).abs().sum(-1).max().item() / 2
+        checks[name] = {
+            "close": torch.allclose(actual, reference, rtol=1e-5, atol=1e-6),
+            "max_total_variation": total_variation,
+        }
+        if not checks[name]["close"]:
+            print(f"Distribution difference: {name} {config.name} {checks[name]}")
+        del actual
+    return checks
+
+
 def print_summary_table(results: list[dict]):
-    """Print a summary table of results."""
+    """Print the same result rows saved to the output file."""
-    print()
-    print("=" * 130)
-    print("SUMMARY TABLE")
-    print("=" * 130)
     print()
+    print(pd.DataFrame(results).to_string(index=False, float_format="{:.3g}".format))
 
-    # Header
-    header = (
-        f"{'Scenario':<40} {'Batch':>6} {'Vocab':>7} {'Ops%':>6} "
-        f"{'Triton (ms)':>12} {'PyTorch (ms)':>13} {'Speedup':>8} "
-        f"{'Tri Mem':>10} {'Pyt Mem':>10}"
-    )
-    print(header)
-    print("-" * 130)
 
-    # Group by scenario type
+def create_result_row(config, timings, checks) -> dict:
-    current_vocab = None
+    """Build a shared row for the summary table and JSONL output."""
-    for result in results:
+    row = {
-        config = result["config"]
+        "Scenario": config.name.split("_b")[0],
-
+        "Batch": config.batch_size,
-        # Add separator between vocab sizes
+        "Vocab": config.vocab_size,
-        if current_vocab != config.vocab_size:
+        "Ops%": f"{config.ops_pct:.0f}%",
-            if current_vocab is not None:
+    }
-                print("-" * 130)
+    row.update({f"{name} (ms)": data["time_ms"] for name, data in timings.items()})
-            current_vocab = config.vocab_size
+    for name, data in timings.items():
-
+        row[f"{name} mem"] = format_memory(data["peak_memory"])
-        scenario = config.name.split("_b")[0]  # Extract scenario name
+    for name, data in checks.items():
-        print(
+        row[f"{name} check"] = data["close"]
-            f"{scenario:<40} {config.batch_size:>6} {config.vocab_size:>7} "
+        row[f"{name} max TV"] = data["max_total_variation"]
-            f"{config.ops_pct:>5.0f}% "
+    return row
-            f"{result['triton_time_ms']:>12.3f} {result['pytorch_time_ms']:>13.3f} "
-            f"{result['speedup']:>7.2f}x "
-            f"{format_memory(result['triton_mem']):>10} "
-            f"{format_memory(result['pytorch_mem']):>10}"
-        )
-
-    print("=" * 130)
 
 
 def main():
     parser = argparse.ArgumentParser(
-        description="Benchmark Triton vs PyTorch sort-based top-k/top-p implementations"
+        description="Benchmark Triton and FlashInfer top-k/top-p filtering"
     )
     parser.add_argument(
         "--batch-sizes",
@@ -431,29 +401,60 @@ def main():
         help="Only print summary table",
     )
 
-    args = parser.parse_args()
+    parser.add_argument(
+        "--num-finite-logits",
+        type=int,
+        help="Finite logits per row; remaining entries are -inf (default: dense)",
+    )
 
-    # Print configuration
+    parser.add_argument(
-    print(f"Batch sizes: {args.batch_sizes}")
+        "--implementations",
-    print(f"Vocab sizes: {args.vocab_sizes}")
+        nargs="+",
-    print(f"Warmup iterations: {args.warmup_iters}")
+        choices=list(IMPLEMENTATIONS),
-    print(f"Benchmark iterations: {args.benchmark_iters}")
+        default=["triton", "flashinfer-mask"],
-    print()
+        help="Implementations to compare (PyTorch is an optional reference)",
+    )
+
+    parser.add_argument(
+        "--check",
+        action="store_true",
+        help="Check filtered distributions against PyTorch outside the timed region",
+    )
+
+    parser.add_argument(
+        "--scenarios",
+        nargs="+",
+        choices=[
+            "topk_whole",
+            "topk_partial",
+            "topp_whole",
+            "topp_partial",
+            "topk_topp_whole",
+            "mixed_partial",
+        ],
+        help="Run only the selected scenarios (default: all)",
+    )
+    parser.add_argument("--output", type=Path, help="Save incremental JSONL results")
+
+    args = parser.parse_args()
+    if args.output is not None:
+        args.output.parent.mkdir(parents=True, exist_ok=True)
+        args.output.open("x").close()
+    if args.num_finite_logits is not None and not (
+        0 <= args.num_finite_logits <= min(args.vocab_sizes)
+    ):
+        parser.error("--num-finite-logits must be between 0 and each vocabulary size")
+    torch.manual_seed(0)
 
     # Check CUDA
     if not torch.cuda.is_available():
         print("ERROR: CUDA is not available. This benchmark requires a GPU.")
         return
 
-    device_name = torch.cuda.get_device_name(0)
-    print(f"GPU: {device_name}")
-    print()
-
     # Create configs
-    configs = create_benchmark_configs(
+    configs = create_benchmark_configs(args.batch_sizes, args.vocab_sizes)
-        args.batch_sizes,
+    if args.scenarios is not None:
-        args.vocab_sizes,
+        configs = [c for c in configs if c.name.split("_b")[0] in args.scenarios]
-    )
 
     # Run benchmarks
     results = run_benchmark(
@@ -461,6 +462,10 @@ def main():
         warmup_iters=args.warmup_iters,
         benchmark_iters=args.benchmark_iters,
         verbose=not args.quiet,
+        num_finite_logits=args.num_finite_logits,
+        implementations=args.implementations,
+        check=args.check,
+        output=args.output,
     )
 
     # Print summary

benchmarks/kernels/benchmark_sparse_indexer_topk_prefill.py

@@ -0,0 +1,119 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Compare in-tree vs AITER sparse-indexer prefill top-k at GLM-5.3-Flash shapes.
+
+Example:
+    .venv/bin/python benchmarks/kernels/benchmark_sparse_indexer_topk_prefill.py \
+        --num-rows 1024 --num-cols 131072 --topk 512
+
+"""
+
+import argparse
+import os
+import time
+
+import torch
+
+from vllm import _custom_ops  # noqa: F401
+from vllm._aiter_ops import rocm_aiter_ops
+
+
+def _make_inputs(num_rows, num_cols, causal, device):
+    logits = torch.randn(num_rows, num_cols, dtype=torch.float32, device=device)
+    row_starts = torch.zeros(num_rows, dtype=torch.int32, device=device)
+    if causal:
+        ramp = torch.arange(1, num_rows + 1, device=device, dtype=torch.float32)
+        row_ends = (ramp * (num_cols / num_rows)).ceil().to(torch.int32)
+    else:
+        row_ends = torch.full((num_rows,), num_cols, dtype=torch.int32, device=device)
+    return logits, row_starts, row_ends
+
+
+def _time(fn, iters, warmup):
+    for _ in range(warmup):
+        fn()
+    torch.accelerator.synchronize()
+    start = time.perf_counter()
+    for _ in range(iters):
+        fn()
+    torch.accelerator.synchronize()
+    return (time.perf_counter() - start) / iters * 1e3
+
+
+def _selected_set(indices, row_starts):
+    valid = indices >= 0
+    return [set(r[m].tolist()) for r, m in zip(indices.cpu(), valid.cpu())]
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--num-rows", type=int, default=1024)
+    parser.add_argument("--num-cols", type=int, default=131072)
+    parser.add_argument("--topk", type=int, default=512)
+    parser.add_argument("--causal", action="store_true")
+    parser.add_argument("--iters", type=int, default=20)
+    parser.add_argument("--warmup", type=int, default=5)
+    args = parser.parse_args()
+
+    device = torch.device("cuda")
+    logits, row_starts, row_ends = _make_inputs(
+        args.num_rows, args.num_cols, args.causal, device
+    )
+    print(
+        f"rows={args.num_rows} cols={args.num_cols} k={args.topk} "
+        f"causal={args.causal} logits={logits.numel() * 4 / 2**20:.0f} MiB "
+        f"path={os.environ.get('TOPK_FORCE_PATH', 'auto')}"
+    )
+
+    native = torch.empty(args.num_rows, args.topk, dtype=torch.int32, device=device)
+
+    def run_native():
+        torch.ops._C.top_k_per_row_prefill(
+            logits,
+            row_starts,
+            row_ends,
+            native,
+            args.num_rows,
+            logits.stride(0),
+            logits.stride(1),
+            args.topk,
+        )
+
+    native_ms = _time(run_native, args.iters, args.warmup)
+    print(f"in-tree  topKPerRowPrefill : {native_ms:8.3f} ms")
+
+    if not rocm_aiter_ops.is_indexer_top_k_enabled():
+        print(
+            "AITER path unavailable: requires ROCm gfx950 with AITER installed "
+            "and VLLM_ROCM_USE_AITER=1"
+        )
+        return
+
+    from vllm.v1.attention.ops.rocm_aiter_mla_sparse import (
+        _launch_aiter_top_k_per_row_prefill,
+    )
+
+    aiter_out = torch.empty(args.num_rows, args.topk, dtype=torch.int32, device=device)
+
+    def run_aiter():
+        _launch_aiter_top_k_per_row_prefill(
+            logits, row_starts, row_ends, aiter_out, args.topk
+        )
+
+    try:
+        aiter_ms = _time(run_aiter, args.iters, args.warmup)
+    except Exception as exc:
+        print(f"AITER path failed: {type(exc).__name__}: {exc}")
+        return
+    print(f"AITER    top_k_per_row_prefill: {aiter_ms:8.3f} ms")
+    print(f"speedup  : {native_ms / aiter_ms:.2f}x")
+
+    ref = _selected_set(native, row_starts)
+    got = _selected_set(aiter_out, row_starts)
+    mismatched = sum(1 for a, b in zip(ref, got) if a != b)
+    worst = max((len(a ^ b) for a, b in zip(ref, got)), default=0)
+    print(f"rows differing: {mismatched}/{args.num_rows} (worst sym-diff {worst})")
+
+
+if __name__ == "__main__":
+    main()

cmake/cpu_extension.cmake

@@ -497,7 +497,7 @@ set(VLLM_EXT_SRC
     "csrc/cpu/mla_decode.cpp"
     "csrc/cpu/pos_encoding.cpp"
     "csrc/cpu/mamba_cpu.cpp"
-    "csrc/moe/dynamic_4bit_int_moe_cpu.cpp"
+    "csrc/cpu/dynamic_4bit_int_moe_cpu.cpp"
     "csrc/cpu/cpu_fused_moe.cpp"
     "csrc/cpu/cpu_attn.cpp"
     "csrc/cpu/cpu_isa.cpp"
@@ -599,7 +599,7 @@ if (ENABLE_X86_ISA)
         "csrc/cpu/layernorm.cpp"
         "csrc/cpu/mla_decode.cpp"
         "csrc/cpu/pos_encoding.cpp"
-        "csrc/moe/dynamic_4bit_int_moe_cpu.cpp") 
+        "csrc/cpu/dynamic_4bit_int_moe_cpu.cpp") 
 
     set(VLLM_EXT_SRC_AVX2
         "csrc/cpu/sgl-kernels/fla.cpp"
@@ -617,7 +617,7 @@ if (ENABLE_X86_ISA)
         "csrc/cpu/layernorm.cpp"
         "csrc/cpu/mla_decode.cpp"
         "csrc/cpu/pos_encoding.cpp"
-        "csrc/moe/dynamic_4bit_int_moe_cpu.cpp") 
+        "csrc/cpu/dynamic_4bit_int_moe_cpu.cpp") 
 
     message(STATUS "CPU extension (AVX512F + BF16 + VNNI + AMX) source files: ${VLLM_EXT_SRC_AVX512} ${VLLM_EXT_SRC_SGL}")
     message(STATUS "CPU extension (AVX512F) source files: ${VLLM_EXT_SRC_AVX512}")

cmake/external_projects/deepgemm.cmake

@@ -29,12 +29,11 @@ if(DEEPGEMM_SRC_DIR)
 else()
   # Keep in sync with tools/install_deepgemm.sh
   set(_DEEPGEMM_UPSTREAM_REPO "https://github.com/vllm-project/DeepGEMM.git")
-  # Pinned to the fork's dev branch: upstream 2.8.0 plus the SM120
+  # Pinned to the fork's dev branch: upstream's 26/09/30 release (2.8.1)
-  # port, the SM90 paged-MQA kv_block=32/next_n=4 port, configurable SwiGLU
+  # plus the SM120 port, the SM90 paged-MQA kv_block=32/next_n=4 port,
-  # alpha/beta, and SiTU for FP8/FP4 Mega MoE. Also includes the CUDA 12.x
+  # configurable SwiGLU alpha/beta, SiTU for FP8/FP4 Mega MoE, and NVFP4
-  # layout header fix from vllm-project/DeepGEMM#12.
+  # Mega MoE. TORCH_LIBRARY build with a CPython 3.10 stable-ABI floor.
-  # TORCH_LIBRARY build with a CPython 3.10 stable-ABI floor.
+  set(_DEEPGEMM_UPSTREAM_TAG "6ce1c5ec6627a6bf11dbcc698b0b1e19c9248432")
-  set(_DEEPGEMM_UPSTREAM_TAG "1e1842a833699298f7afc02eefb2ed168fff6938")
 
   set(_deepgemm_fc_root "${FETCHCONTENT_BASE_DIR}")
   if(NOT _deepgemm_fc_root)
@@ -146,6 +145,7 @@ if(DEEPGEMM_ARCHS)
   #
   install(FILES
     "${deepgemm_SOURCE_DIR}/deep_gemm/__init__.py"
+    "${deepgemm_SOURCE_DIR}/deep_gemm/locality_domain.py"
     DESTINATION vllm/third_party/deep_gemm
     COMPONENT _deep_gemm_C)
 

cmake/external_projects/fmha_sm100.cmake

@@ -17,7 +17,7 @@ else()
   FetchContent_Declare(
     fmha_sm100
     GIT_REPOSITORY https://github.com/vllm-project/MSA.git
-    GIT_TAG be4a52aea420b6bef73be9af1a649859964790cb
+    GIT_TAG 968061ab16cb14991a5c7d4c0e8eabf531be0dcf
     GIT_PROGRESS TRUE
     CONFIGURE_COMMAND ""
     BUILD_COMMAND ""
@@ -36,11 +36,14 @@ set(FMHA_SM100_PY_ROOT "${fmha_sm100_SOURCE_DIR}/python/fmha_sm100")
 
 install(FILES
   "${FMHA_SM100_PY_ROOT}/__init__.py"
+  "${FMHA_SM100_PY_ROOT}/_jit_cache.py"
   "${FMHA_SM100_PY_ROOT}/api.py"
   "${FMHA_SM100_PY_ROOT}/bench_utils.py"
   "${FMHA_SM100_PY_ROOT}/jit.py"
+  "${FMHA_SM100_PY_ROOT}/msa_warmup.py"
   "${FMHA_SM100_PY_ROOT}/nvfp4_kv.py"
   "${FMHA_SM100_PY_ROOT}/q8kv4_decode_adapter.py"
+  "${FMHA_SM100_PY_ROOT}/q8kv4_prefill_adapter.py"
   "${FMHA_SM100_PY_ROOT}/sparse.py"
   "${FMHA_SM100_PY_ROOT}/sparse_fmha_adapter.py"
   DESTINATION vllm/third_party/fmha_sm100
@@ -65,6 +68,12 @@ install(DIRECTORY "${FMHA_SM100_PY_ROOT}/kvouter/"
   PATTERN "__pycache__" EXCLUDE
   PATTERN "*.pyc" EXCLUDE)
 
+install(DIRECTORY "${FMHA_SM100_PY_ROOT}/prefill_q8kv4/"
+  DESTINATION vllm/third_party/fmha_sm100/prefill_q8kv4
+  COMPONENT fmha_sm100
+  PATTERN "__pycache__" EXCLUDE
+  PATTERN "*.pyc" EXCLUDE)
+
 install(DIRECTORY "${FMHA_SM100_PY_ROOT}/cute/"
   DESTINATION vllm/third_party/fmha_sm100/cute
   COMPONENT fmha_sm100

cmake/external_projects/vllm_flash_attn.cmake

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

cmake/hipify.py

@@ -65,7 +65,7 @@ if __name__ == "__main__":
         project_directory=args.project_dir,
         output_directory=args.output_dir,
         # Hipify resolves quoted includes next to the including file first; vLLM
-        # uses paths relative to csrc/ (e.g. "libtorch_stable/torch_utils.h"
+        # uses paths relative to csrc/ (e.g. "torch_utils.h"
         # from quantization/w8a8/fp8/*.cu). Without an include root here, those
         # headers are never found and are not hipified or rewritten in dependents.
         header_include_dirs=["."],

csrc/activation_kernels.cu

@@ -10,7 +10,7 @@
 #include <cmath>
 #include <type_traits>
 
-#include "../cuda_compat.h"
+#include "core/cuda_compat.h"
 #include "async_util.cuh"
 #include "cuda_vec_utils.cuh"
 #include "dispatch_utils.h"

csrc/attention/merge_attn_states.cu

@@ -6,9 +6,8 @@
 #include "../dispatch_utils.h"
 #include <torch/headeronly/core/ScalarType.h>
 
-#include "../../attention/attention_dtypes.h"
+#include "core/vec_dtypes/vec_dtypes.h"
-#include "attention_utils.cuh"
+#include "quantization/w8a8/fp8/common.cuh"
-#include "../../quantization/w8a8/fp8/common.cuh"
 
 namespace vllm {
 

csrc/attention/mla/sm100_cutlass_mla_kernel.cu

@@ -18,7 +18,7 @@ limitations under the License.
  * Taken from SGLANG PR https://github.com/sgl-project/sglang/pull/6929
  * by Alcanderian JieXin Liang
  */
-#include "libtorch_stable/torch_utils.h"
+#include "torch_utils.h"
 
 #include <torch/csrc/stable/library.h>
 

csrc/cache.h

@@ -1,93 +0,0 @@
-#pragma once
-
-#include <torch/all.h>
-#include <c10/util/Optional.h>
-
-#include <map>
-#include <vector>
-
-void swap_blocks(torch::Tensor& src, torch::Tensor& dst,
-                 int64_t block_size_in_bytes,
-                 const torch::Tensor& block_mapping);
-
-void swap_blocks_batch(const torch::Tensor& src_ptrs,
-                       const torch::Tensor& dst_ptrs,
-                       const torch::Tensor& sizes,
-                       bool is_src_access_order_any);
-
-void reshape_and_cache(torch::Tensor& key, torch::Tensor& value,
-                       torch::Tensor& key_cache, torch::Tensor& value_cache,
-                       torch::Tensor& slot_mapping,
-                       const std::string& kv_cache_dtype,
-                       torch::Tensor& k_scale, torch::Tensor& v_scale);
-
-void reshape_and_cache_flash(torch::Tensor& key, torch::Tensor& value,
-                             torch::Tensor& key_cache,
-                             torch::Tensor& value_cache,
-                             torch::Tensor& slot_mapping,
-                             const std::string& kv_cache_dtype,
-                             torch::Tensor& k_scale, torch::Tensor& v_scale);
-
-void concat_and_cache_mla(torch::Tensor& kv_c, torch::Tensor& k_pe,
-                          torch::Tensor& kv_cache, torch::Tensor& slot_mapping,
-                          const std::string& kv_cache_dtype,
-                          torch::Tensor& scale);
-
-// NOTE: k_pe and kv_c order is flipped compared to concat_and_cache_mla
-void concat_and_cache_mla_rope_fused(
-    torch::Tensor& positions, torch::Tensor& q_pe, torch::Tensor& k_pe,
-    torch::Tensor& kv_c, torch::Tensor& rope_cos_sin_cache, bool rope_is_neox,
-    torch::Tensor& kv_cache_slot_mapping, torch::Tensor& kv_cache,
-    const std::string& kv_cache_dtype, torch::Tensor& kv_cache_quant_scale);
-
-// Just for unittest
-void convert_fp8(torch::Tensor& dst_cache, torch::Tensor& src_cache,
-                 const double scale, const std::string& kv_cache_dtype);
-
-void gather_and_maybe_dequant_cache(
-    torch::Tensor const& src_cache,     // [NUM_BLOCKS, BLOCK_SIZE, ENTRIES...]
-    torch::Tensor const& dst,           // [TOT_TOKENS, ENTRIES...]
-    torch::Tensor const& block_table,   // [BATCH, BLOCK_INDICES]
-    torch::Tensor const& cu_seq_lens,   // [BATCH+1]
-    torch::Tensor const& token_to_seq,  // [MAX_TOKEN_ACROSS_CHUNKS]
-    int64_t num_tokens, const std::string& kv_cache_dtype,
-    torch::Tensor const& scale,
-    std::optional<torch::Tensor> seq_starts = std::nullopt);
-
-// TODO(hc): cp_gather_cache need support scaled kvcahe in the future.
-void cp_gather_cache(
-    torch::Tensor const& src_cache,    // [NUM_BLOCKS, BLOCK_SIZE, ENTRIES...]
-    torch::Tensor const& dst,          // [TOT_TOKENS, ENTRIES...]
-    torch::Tensor const& block_table,  // [BATCH, BLOCK_INDICES]
-    torch::Tensor const& cu_seq_lens,  // [BATCH+1]
-    int64_t batch_size, std::optional<torch::Tensor> seq_starts = std::nullopt);
-
-// Gather and upconvert FP8 KV cache to BF16 workspace
-void cp_gather_and_upconvert_fp8_kv_cache(
-    torch::Tensor const& src_cache,         // [NUM_BLOCKS, BLOCK_SIZE, 656]
-    torch::Tensor const& dst,               // [TOT_TOKENS, 576]
-    torch::Tensor const& block_table,       // [BATCH, BLOCK_INDICES]
-    torch::Tensor const& workspace_starts,  // [BATCH]
-    int64_t batch_size, std::optional<torch::Tensor> seq_starts = std::nullopt);
-
-// Indexer K quantization and cache function
-void indexer_k_quant_and_cache(
-    torch::Tensor& k,             // [num_tokens, head_dim]
-    torch::Tensor& kv_cache,      // [num_blocks, block_size, cache_stride]
-    torch::Tensor& slot_mapping,  // [num_tokens]
-    int64_t quant_block_size,     // quantization block size
-    const std::string& scale_fmt);
-
-// Concatenate query nope and rope for MLA/DSA attention
-void concat_mla_q(
-    torch::Tensor& ql_nope,  // [num_tokens, num_heads, nope_dim]
-    torch::Tensor& q_pe,     // [num_tokens, num_heads, rope_dim]
-    torch::Tensor& q_out);   // [num_tokens, num_heads, nope_dim + rope_dim]
-
-// Extract function to gather quantized K cache
-void cp_gather_indexer_k_quant_cache(
-    const torch::Tensor& kv_cache,  // [num_blocks, block_size, cache_stride]
-    torch::Tensor& dst_k,           // [num_tokens, head_dim]
-    torch::Tensor& dst_scale,  // [num_tokens, head_dim / quant_block_size * 4]
-    const torch::Tensor& block_table,   // [batch_size, num_blocks]
-    const torch::Tensor& cu_seq_lens);  // [batch_size + 1]

csrc/cache_kernels.cu

@@ -1,8 +1,8 @@
 #include "torch_utils.h"
 #include "dispatch_utils.h"
 
-#include "../cuda_utils.h"
+#include "core/utils.hpp"
-#include "../cuda_compat.h"
+#include "core/cuda_compat.h"
 
 #include "quantization/vectorization_utils.cuh"
 #include "concat_mla_q.cuh"
@@ -12,9 +12,9 @@
 #endif
 
 #ifdef USE_ROCM
-  #include "../quantization/w8a8/fp8/amd/quant_utils.cuh"
+  #include "quantization/w8a8/fp8/amd/quant_utils.cuh"
 #else
-  #include "../quantization/w8a8/fp8/nvidia/quant_utils.cuh"
+  #include "quantization/w8a8/fp8/nvidia/quant_utils.cuh"
 #endif
 
 #include <algorithm>
@@ -630,7 +630,8 @@ __global__ void indexer_k_quant_and_cache_kernel(
     const int cache_block_size,                // cache block size
     const int64_t cache_block_stride,  // stride for each block in kv_cache
 
-    const bool use_ue8m0  // use ue8m0 scale format
+    const bool use_ue8m0,    // use ue8m0 scale format
+    const bool use_mx_scale  // store the ue8m0 scale as packed MX bytes
 ) {
   constexpr int VEC_SIZE = 4;
   const int64_t token_idx = blockIdx.x;
@@ -679,7 +680,13 @@ __global__ void indexer_k_quant_and_cache_kernel(
     const int64_t dst_scale_idx =
         block_idx * cache_block_stride + cache_block_size * head_dim +
         (block_offset * head_dim + head_dim_idx) * 4 / quant_block_size;
-    reinterpret_cast<float*>(kv_cache)[dst_scale_idx / 4] = scale;
+    if (use_mx_scale) {
+      // One UE8M0 byte per 32 elements, all equal to the shared scale's.
+      reinterpret_cast<uint32_t*>(kv_cache)[dst_scale_idx / 4] =
+          ((__float_as_uint(scale) >> 23) & 0xFF) * 0x01010101u;
+    } else {
+      reinterpret_cast<float*>(kv_cache)[dst_scale_idx / 4] = scale;
+    }
   }
 }
 
@@ -712,8 +719,7 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
   }
   __syncthreads();
 
-  for (int iter = 0; iter < cuda_utils::ceil_div(batch_size, int(blockDim.x));
+  for (int iter = 0; iter < div_ceil(batch_size, int(blockDim.x)); iter++) {
-       iter++) {
     int tid = iter * blockDim.x + threadIdx.x;
     if (tid < batch_size) {
       const int seq_start = cu_seq_lens[tid];
@@ -1229,7 +1235,7 @@ __device__ __forceinline__ bool map_gather_page_task(
     const int32_t source_begin = seq_starts == nullptr ? 0 : seq_starts[req_id];
     const int32_t first_block = source_begin / block_size;
     const int32_t num_pages =
-        cuda_utils::ceil_div(source_begin + seq_len, block_size) - first_block;
+        div_ceil(source_begin + seq_len, block_size) - first_block;
     if (relative_page < num_pages) {
       page.req_id = req_id;
       page.logical_block = first_block + relative_page;
@@ -1408,8 +1414,7 @@ void gather_and_maybe_dequant_cache(
   const int64_t dst_entry_stride = dst.stride(0);
   const int32_t page_threads = num_tokens >= (1 << 20) ? 128 : 256;
   const int32_t required_blocks =
-      cuda_utils::ceil_div(static_cast<int32_t>(num_tokens), block_size) +
+      div_ceil(static_cast<int32_t>(num_tokens), block_size) + 2 * num_reqs;
-      2 * num_reqs;
   const dim3 grid(required_blocks);
   const dim3 block(page_threads);
 
@@ -1685,7 +1690,7 @@ void cp_gather_cache(
 
   const int32_t num_reqs = static_cast<int32_t>(batch_size);
   const int32_t required_blocks =
-      cuda_utils::ceil_div(total_tokens, block_size) + 2 * num_reqs;
+      div_ceil(total_tokens, block_size) + 2 * num_reqs;
   const dim3 grid(required_blocks);
   const dim3 block(256);
 
@@ -1828,8 +1833,8 @@ void cp_gather_and_upconvert_fp8_kv_cache(
   const int32_t* seq_starts_ptr =
       seq_starts.has_value() ? seq_starts.value().const_data_ptr<int32_t>()
                              : nullptr;
-  const int required_blocks = cuda_utils::ceil_div(total_tokens, block_size) +
+  const int required_blocks =
-                              2 * static_cast<int32_t>(batch_size);
+      div_ceil(total_tokens, block_size) + 2 * static_cast<int32_t>(batch_size);
   vllm::cp_gather_and_upconvert_fp8_kv_cache_page<<<
       required_blocks, block_size_threads, 0, stream>>>(
       src_ptr, reinterpret_cast<__nv_bfloat16*>(dst.data_ptr()),
@@ -1914,7 +1919,7 @@ void cp_gather_and_upconvert_nvfp4_kv_cache(
           reinterpret_cast<KV_T*>(k.data_ptr()),                              \
           reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),                    \
           slot_mapping.const_data_ptr<int64_t>(), head_dim, quant_block_size, \
-          cache_block_size, cache_block_stride, use_ue8m0);
+          cache_block_size, cache_block_stride, use_ue8m0, use_mx_scale);
 
 void indexer_k_quant_and_cache(
     torch::stable::Tensor& k,         // [num_tokens, head_dim]
@@ -1926,8 +1931,13 @@ void indexer_k_quant_and_cache(
   int head_dim = k.size(1);
   int cache_block_size = kv_cache.size(1);
   int64_t cache_block_stride = kv_cache.stride(0);
-  bool use_ue8m0 = scale_fmt == "ue8m0";
+  // "mxfp8" stores the ue8m0 scale as four packed UE8M0 bytes, the MX scales
+  // DeepGEMM's SM100 MQA logits kernels read.
+  bool use_mx_scale = scale_fmt == "mxfp8";
+  bool use_ue8m0 = use_mx_scale || scale_fmt == "ue8m0";
 
+  STD_TORCH_CHECK(!use_mx_scale || (head_dim == 128 && quant_block_size == 128),
+                  "mxfp8 scales require head_dim == quant_block_size == 128");
   STD_TORCH_CHECK(k.device() == kv_cache.device(),
                   "k and kv_cache must be on the same device");
   STD_TORCH_CHECK(k.device() == slot_mapping.device(),

csrc/cache_kernels_fused.cu

@@ -1,13 +1,13 @@
 #include "torch_utils.h"
 #include "dispatch_utils.h"
 
-#include "../cuda_compat.h"
+#include "core/cuda_compat.h"
 
-#include "../quantization/w8a8/fp8/common.cuh"
+#include "quantization/w8a8/fp8/common.cuh"
 #ifdef USE_ROCM
-  #include "../quantization/w8a8/fp8/amd/quant_utils.cuh"
+  #include "quantization/w8a8/fp8/amd/quant_utils.cuh"
 #else
-  #include "../quantization/w8a8/fp8/nvidia/quant_utils.cuh"
+  #include "quantization/w8a8/fp8/nvidia/quant_utils.cuh"
 #endif
 
 #ifdef USE_ROCM

csrc/core/batch_invariant.hpp

@@ -1,18 +0,0 @@
-#pragma once
-#include <cstdlib>
-#include <string>
-
-namespace vllm {
-
-// vllm_is_batch_invariant(); returns true
-// if env VLLM_BATCH_INVARIANT=1
-inline bool vllm_is_batch_invariant() {
-  static bool cached = []() {
-    std::string env_key = "VLLM_BATCH_INVARIANT";
-    const char* val = std::getenv(env_key.c_str());
-    return (val && std::atoi(val) != 0) ? 1 : 0;
-  }();
-  return cached;
-}
-
-}  // namespace vllm

csrc/core/cuda_compat.h

@@ -54,12 +54,6 @@ struct Utils {
     __shfl_xor(var, lane_mask, width)
 #endif
 
-#ifndef USE_ROCM
-  #define VLLM_SHFL_SYNC(var, src_lane) __shfl_sync(uint32_t(-1), var, src_lane)
-#else
-  #define VLLM_SHFL_SYNC(var, src_lane) __shfl(var, src_lane)
-#endif
-
 #ifndef USE_ROCM
   #define VLLM_SHFL_DOWN_SYNC(var, lane_delta) \
     __shfl_down_sync(uint32_t(-1), var, lane_delta)

csrc/core/exception.hpp

@@ -1,3 +0,0 @@
-#pragma once
-
-#define VLLM_IMPLIES(p, q) (!(p) || (q))

csrc/core/registration.h

@@ -12,11 +12,6 @@
 // could be a macro instead of a literal token.
 #define TORCH_LIBRARY_EXPAND(NAME, MODULE) TORCH_LIBRARY(NAME, MODULE)
 
-// A version of the TORCH_LIBRARY_IMPL macro that expands the NAME, i.e. so NAME
-// could be a macro instead of a literal token.
-#define TORCH_LIBRARY_IMPL_EXPAND(NAME, DEVICE, MODULE) \
-  TORCH_LIBRARY_IMPL(NAME, DEVICE, MODULE)
-
 // REGISTER_EXTENSION allows the shared library to be loaded and initialized
 // via python's import statement.
 #define REGISTER_EXTENSION(NAME)                                               \

csrc/core/utils.hpp

@@ -1,7 +1,8 @@
 #pragma once
 
 #include <climits>
-#include <iostream>
+#include <cstdint>
+#include <cstdlib>
 
 inline constexpr uint32_t next_pow_2(uint32_t const num) {
   if (num <= 1) return num;
@@ -26,3 +27,16 @@ template <typename T>
 inline constexpr T round_to_next_multiple_of(T a, T b) {
   return a % b == 0 ? a : ((a / b) + 1) * b;
 }
+
+namespace vllm {
+
+// Returns true if env VLLM_BATCH_INVARIANT=1
+inline bool vllm_is_batch_invariant() {
+  static bool cached = []() {
+    const char* val = std::getenv("VLLM_BATCH_INVARIANT");
+    return val && std::atoi(val) != 0;
+  }();
+  return cached;
+}
+
+}  // namespace vllm

csrc/core/vec_dtypes/dtype_bfloat16.cuh

@@ -20,7 +20,7 @@
  */
 #pragma once
 
-#include "attention_generic.cuh"
+#include "dtype_generic.cuh"
 #include "dtype_float32.cuh"
 
 #ifndef USE_ROCM

csrc/core/vec_dtypes/dtype_float16.cuh

@@ -20,7 +20,7 @@
  */
 #pragma once
 
-#include "attention_generic.cuh"
+#include "dtype_generic.cuh"
 #include "dtype_float32.cuh"
 
 #ifdef USE_ROCM

csrc/core/vec_dtypes/dtype_float32.cuh

@@ -20,7 +20,7 @@
  */
 #pragma once
 
-#include "attention_generic.cuh"
+#include "dtype_generic.cuh"
 
 #include <stdint.h>
 

csrc/core/vec_dtypes/dtype_fp8.cuh

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

csrc/core/vec_dtypes/vec_dtypes.h

@@ -1,6 +1,6 @@
 #pragma once
 
-#include "attention_generic.cuh"
+#include "dtype_generic.cuh"
 #include "dtype_float16.cuh"
 #include "dtype_float32.cuh"
 #include "dtype_bfloat16.cuh"

csrc/cpu/cpu_attn_impl.hpp

@@ -14,7 +14,7 @@
 namespace cpu_attention {
 enum class ISA { AMX, VEC, VEC16, NEON, VXE, RVV, VSX, AMX_FP8 };
 
-// Mirrors csrc/attention/dtype_fp8.cuh Fp8KVCacheDataType exactly.
+// Mirrors csrc/core/vec_dtypes/dtype_fp8.cuh Fp8KVCacheDataType exactly.
 enum class Fp8KVCacheDataType {
   kAuto = 0,
   kFp8E4M3 = 1,

csrc/cpu/ops.h

@@ -0,0 +1,46 @@
+#pragma once
+
+#include <optional>
+#include <string>
+#include <torch/library.h>
+#include <tuple>
+#include <vector>
+
+void rms_norm(torch::Tensor& out, torch::Tensor& input,
+              std::optional<torch::Tensor> weight, double epsilon);
+
+void fused_add_rms_norm(torch::Tensor& input, torch::Tensor& residual,
+                        std::optional<torch::Tensor> weight, double epsilon);
+
+void rotary_embedding(torch::Tensor& positions, torch::Tensor& query,
+                      std::optional<torch::Tensor> key, int64_t head_size,
+                      torch::Tensor& cos_sin_cache, bool is_neox,
+                      int64_t rope_dim_offset, bool inverse);
+
+void silu_and_mul(torch::Tensor& out, torch::Tensor& input);
+
+void gelu_and_mul(torch::Tensor& out, torch::Tensor& input);
+
+void gelu_tanh_and_mul(torch::Tensor& out, torch::Tensor& input);
+
+void gelu_tanh(torch::Tensor& out, torch::Tensor& input);
+
+void gelu_new(torch::Tensor& out, torch::Tensor& input);
+
+void gelu_fast(torch::Tensor& out, torch::Tensor& input);
+
+void gelu_quick(torch::Tensor& out, torch::Tensor& input);
+
+void static_scaled_int8_quant(torch::Tensor& out, torch::Tensor const& input,
+                              torch::Tensor const& scale,
+                              std::optional<torch::Tensor> const& azp);
+
+void dynamic_scaled_int8_quant(torch::Tensor& out, torch::Tensor const& input,
+                               torch::Tensor& scales,
+                               std::optional<torch::Tensor> const& azp);
+
+torch::Tensor dynamic_4bit_int_moe_cpu(
+    torch::Tensor x, torch::Tensor topk_ids, torch::Tensor topk_weights,
+    torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t hidden_size,
+    int64_t intermediate_size, int64_t group_size,
+    bool apply_router_weight_on_input, int64_t activation_kind);

csrc/cpu/torch_bindings.cpp

@@ -1,4 +1,4 @@
-#include "cache.h"
+#include <torch/all.h>
 #include "ops.h"
 #include "core/registration.h"
 

csrc/cuda_utils.h

@@ -1,21 +1,8 @@
 #pragma once
 
+#include <cstdint>
 #include <stdio.h>
 
-#if defined(__HIPCC__)
-  #define HOST_DEVICE_INLINE __host__ __device__
-  #define DEVICE_INLINE __device__
-  #define HOST_INLINE __host__
-#elif defined(__CUDACC__) || defined(_NVHPC_CUDA)
-  #define HOST_DEVICE_INLINE __host__ __device__ __forceinline__
-  #define DEVICE_INLINE __device__ __forceinline__
-  #define HOST_INLINE __host__ __forceinline__
-#else
-  #define HOST_DEVICE_INLINE inline
-  #define DEVICE_INLINE inline
-  #define HOST_INLINE inline
-#endif
-
 #define CUDA_CHECK(cmd)                                             \
   do {                                                              \
     cudaError_t e = cmd;                                            \
@@ -29,13 +16,3 @@
 int64_t get_device_attribute(int64_t attribute, int64_t device_id);
 
 int64_t get_max_shared_memory_per_block_device_attribute(int64_t device_id);
-
-namespace cuda_utils {
-
-template <typename T>
-HOST_DEVICE_INLINE constexpr std::enable_if_t<std::is_integral_v<T>, T>
-ceil_div(T a, T b) {
-  return (a + b - 1) / b;
-}
-
-};  // namespace cuda_utils