[Bugfix][Frontend] Stop beam search on generation config EOS tokens (#62072)

Signed-off-by: mottopanikeiku <176798723+mottopanikeiku@users.noreply.github.com>
Co-authored-by: mottopanikeiku <176798723+mottopanikeiku@users.noreply.github.com>

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

@@ -0,0 +1,115 @@
+---
+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 |
+
+Verify cases require `--speculative-config` with `num_speculative_tokens`
+at least as large as the case's draft count; insufficient capacity is rejected.
+
+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).
+
+## Kernel breakdown
+
+Add `--profile-dir DIR` to see where the time goes. After timing, each case is
+replayed 3 more times under `torch.profiler`; every rank writes
+`DIR/<case>.rank<r>.kernels.json` (GPU time per kernel per forward, sorted)
+and `DIR/<case>.rank<r>.trace.json.gz` (open in Perfetto). Rank 0's top
+kernels are printed and included in `--output-json`.
+
+```bash
+python -m vllm.benchmarks.forward.profile DIR/decode.rank0.kernels.json
+# Before/after a change, same case: per-kernel deltas, largest first.
+python -m vllm.benchmarks.forward.profile base/decode.rank0.kernels.json new/decode.rank0.kernels.json
+```
+
+- Kernels are grouped by name without template or function arguments; the
+  JSON keeps one full name per group in `example`.
+- Only the timed forward is included: the in-graph KV restore and rank
+  rendezvous are excluded (`pre_forward_us`). `span_us` (first to last
+  forward kernel) should be close to `median_us`, a bit higher from profiler
+  overhead; `busy_us` can exceed it when streams overlap.
+- Compare the same rank across runs: under EP, ranks differ (e.g. collective
+  wait time shows up inside communication kernels).
+
+For Nsight Systems (NVTX, NCCL, CPU side, GPU metrics), skip `--profile-dir`:
+the two can't share CUPTI. Each case's 3 profiled replays run between
+`cudaProfilerStart`/`Stop` in an NVTX range `forward/<case>`, so a capture range
+records only them, one report per case (and rank):
+
+```bash
+nsys profile --capture-range=cudaProfilerApi --capture-range-end=repeat \
+  --cuda-graph-trace=node -o prof vllm bench forward ...       # prof.1, prof.2, ...
+torchrun --nproc-per-node=4 --no-python nsys profile ... -o prof_rank%q{RANK} \
+  python -m vllm.entrypoints.cli.main bench forward ...
+nsys stats --report cuda_gpu_kern_sum --format csv -o kern prof.1.nsys-rep
+```
+
+Keep `--capture-range-end=repeat`: nsys's default stops after the first case
+and leaves the bench hanging. `nsys stats` also counts the KV restore copies
+that the JSON summary excludes.
+
+## 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/model-recipes/SKILL.md

@@ -5,9 +5,12 @@ description: Find vLLM deployment recipes and configs for LLMs and Hugging Face
 
 # Model Recipes
 
-Model serving recipes are maintained in the
-[models/ directory of vllm-project/recipes](https://github.com/vllm-project/recipes/tree/main/models).
-The published catalog is at
+Model serving recipes are maintained in this repository's
+[recipes/models/ directory](../../../recipes/models). Omni (image, video,
+audio) recipes stay in vllm-project/recipes for now (moving to vllm-omni),
+and Ascend NPU content is inline in these recipes (moving to vllm-ascend
+overlays). The published catalog, rendered by
+[vllm-project/recipes](https://github.com/vllm-project/recipes), is at
 [recipes.vllm.ai](https://recipes.vllm.ai).
 
 ## Usage notes
@@ -146,6 +149,9 @@ uploads, untested hardware claims, or non-standard flags that only work for one
 setup. Before opening a pull request, confirm that the user ran the
 configuration and can share the results.
 
-Follow the recipes repository's
+Edit the YAML under [recipes/models/](../../../recipes/models) in this
+repository, following the recipes repository's
 [add-recipe skill](https://github.com/vllm-project/recipes/blob/main/.agents/skills/add-recipe/SKILL.md)
-in a checkout of `vllm-project/recipes` or the user's intended fork.
+for the schema. Validate and preview it from a checkout of
+`vllm-project/recipes` pointed at this directory, as
+[recipes/README.md](../../../recipes/README.md) describes.

.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
-vLLM derives, vanish.
+declare survive as plain attributes, but only when the checkpoint sets them.
+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
-      5.16.1.
+      issue is not a fixed one: issues get closed while the bug is still in
+      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/check-torch-abi.py

@@ -17,7 +17,6 @@ ALLOWED_UNSTABLE_LIBRARIES: tuple[str, ...] = (
     # _C_stable_libtorch, but these two modules still bind libtorch directly.
     "_C.abi3.so",
     "_rocm_C.abi3.so",
-    "third_party/deep_gemm/_C*.so",
 )
 
 

.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" --deselect="tests/model_executor/test_enabled_custom_ops.py::test_topk_gating_enabled_by_cdna"'
 
-- 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/release-pipeline.yaml

@@ -703,385 +703,6 @@ steps:
           DOCKER_BUILDKIT: "1"
           DOCKERHUB_USERNAME: "vllmbot"
 
-  # =============================================================================
-  # ROCm Release Pipeline (x86_64 only)
-  # =============================================================================
-  #
-  # vLLM version is determined by the Buildkite checkout (like CUDA pipeline).
-  # To build a specific version, trigger the build from that branch/tag.
-  #
-  # Environment variables for ROCm builds (set via Buildkite UI or schedule):
-  #
-  # Note: ROCm version is determined by BASE_IMAGE in docker/Dockerfile.rocm_base
-  #
-  # =============================================================================
-
-  - group: "Build ROCm Wheel / Image "
-    key: "build-rocm-wheel-image"
-    depends_on: ~
-    steps:
-      # ROCm Job 1: Build ROCm Base Wheels (with S3 caching)
-      - label: ":rocm: Build ROCm Base Image & Wheels"
-        id: build-rocm-base-wheels
-        depends_on: ~
-        agents:
-          queue: cpu_queue_release
-        commands:
-          - |
-            set -euo pipefail
-
-            # Generate cache key
-            CACHE_KEY=$$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
-            ECR_CACHE_TAG="public.ecr.aws/q9t5s3a7/vllm-release-repo:$${CACHE_KEY}-rocm-base"
-
-            echo "========================================"
-            echo "ROCm Base Build Configuration"
-            echo "========================================"
-            echo "  CACHE_KEY: $${CACHE_KEY}"
-            echo "  ECR_CACHE_TAG: $${ECR_CACHE_TAG}"
-            echo "========================================"
-            
-            # Login to ECR
-            aws ecr-public get-login-password --region us-east-1 | \
-              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-            
-            IMAGE_EXISTS=false
-            WHEELS_EXIST=false
-            
-            # Check ECR for Docker image
-
-            if docker manifest inspect "$${ECR_CACHE_TAG}" > /dev/null 2>&1; then
-              IMAGE_EXISTS=true
-              echo "ECR image cache HIT"
-            fi
-            
-            # Check S3 for wheels
-            WHEEL_CACHE_STATUS=$(.buildkite/scripts/cache-rocm-base-wheels.sh check)
-            if [ "$${WHEEL_CACHE_STATUS}" = "hit" ]; then
-              WHEELS_EXIST=true
-              echo "S3 wheels cache HIT"
-            fi
-
-            
-            # Scenario 1: Both cached (best case)
-            if [ "$${IMAGE_EXISTS}" = "true" ] && [ "$${WHEELS_EXIST}" = "true" ]; then
-              echo ""
-              echo "FULL CACHE HIT - Reusing both image and wheels"
-              echo ""
-
-              # Download wheels
-              .buildkite/scripts/cache-rocm-base-wheels.sh download
-              
-              # Save ECR tag for downstream jobs
-              buildkite-agent meta-data set "rocm-base-image-tag" "$${ECR_CACHE_TAG}"
-              
-            # Scenario 2: Full rebuild needed
-            else
-              echo ""
-              echo " CACHE MISS - Building from scratch..."
-              echo ""
-              
-              # Build full base image and push to ECR
-              DOCKER_BUILDKIT=1 docker buildx build \
-                --file docker/Dockerfile.rocm_base \
-                --tag "$${ECR_CACHE_TAG}" \
-                --build-arg USE_SCCACHE=1 \
-                --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-                --build-arg SCCACHE_REGION_NAME=us-west-2 \
-                --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-                --push \
-                .
-              
-              # Build wheel extraction stage
-              DOCKER_BUILDKIT=1 docker buildx build \
-                --file docker/Dockerfile.rocm_base \
-                --tag rocm-base-debs:$${BUILDKITE_BUILD_NUMBER} \
-                --target debs_wheel_release \
-                --build-arg USE_SCCACHE=1 \
-                --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-                --build-arg SCCACHE_REGION_NAME=us-west-2 \
-                --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-                --load \
-                .
-              
-              # Extract and upload wheels
-              mkdir -p artifacts/rocm-base-wheels
-              cid=$(docker create rocm-base-debs:$${BUILDKITE_BUILD_NUMBER})
-              docker cp $${cid}:/app/debs/. artifacts/rocm-base-wheels/
-              docker rm $${cid}
-              
-              .buildkite/scripts/cache-rocm-base-wheels.sh upload
-
-              # Cache base docker image to ECR
-              docker push "$${ECR_CACHE_TAG}"
-              
-              buildkite-agent meta-data set "rocm-base-image-tag" "$${ECR_CACHE_TAG}"
-              
-              echo ""
-              echo " Build complete - Image and wheels cached"
-            fi
-
-        artifact_paths:
-          - "artifacts/rocm-base-wheels/*.whl"
-        env:
-          DOCKER_BUILDKIT: "1"
-          S3_BUCKET: "vllm-wheels"
-
-      # ROCm Job 2: Build vLLM ROCm Wheel
-      - label: ":python: Build vLLM ROCm Wheel - x86_64"
-        id: build-rocm-vllm-wheel
-        depends_on:
-          - step: build-rocm-base-wheels
-            allow_failure: false
-        agents:
-          queue: cpu_queue_release
-        timeout_in_minutes: 180
-        commands:
-          # Download artifacts and prepare Docker image
-          - |
-            set -euo pipefail
-
-            # Ensure git tags are up-to-date (Buildkite's default fetch doesn't update tags)
-            # This fixes version detection when tags are moved/force-pushed
-            echo "Fetching latest tags from origin..."
-            git fetch --tags --force origin
-            
-            # Log tag information for debugging version detection
-            echo "========================================"
-            echo "Git Tag Verification"
-            echo "========================================"
-            echo "Current HEAD: $(git rev-parse HEAD)"
-            echo "git describe --tags: $(git describe --tags 2>/dev/null || echo 'No tags found')"
-            echo ""
-            echo "Recent tags (pointing to commits near HEAD):"
-            git tag -l --sort=-creatordate | head -5
-            echo "setuptools_scm version detection:"
-            pip install -q setuptools_scm 2>/dev/null || true
-            python3 -c "import setuptools_scm; print('  Detected version:', setuptools_scm.get_version())" 2>/dev/null || echo "  (setuptools_scm not available in this environment)"
-            echo "========================================"
-
-            # Download wheel artifacts from current build
-            echo "Downloading wheel artifacts from current build"
-            buildkite-agent artifact download "artifacts/rocm-base-wheels/*.whl" .
-
-            # Get ECR image tag from metadata (set by build-rocm-base-wheels)
-            ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rocm-base-image-tag 2>/dev/null || echo '')"
-            if [ -z "$${ECR_IMAGE_TAG}" ]; then
-              echo "ERROR: rocm-base-image-tag metadata not found"
-              echo "This should have been set by the build-rocm-base-wheels job"
-              exit 1
-            fi
-            
-            echo "Pulling base Docker image from ECR: $${ECR_IMAGE_TAG}"
-            
-            # Login to ECR
-            aws ecr-public get-login-password --region us-east-1 | \
-              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-            
-            # Pull base Docker image from ECR
-            docker pull "$${ECR_IMAGE_TAG}"
-            
-            echo "Loaded base image: $${ECR_IMAGE_TAG}"
-            
-            # Prepare base wheels for Docker build context
-            mkdir -p docker/context/base-wheels
-            touch docker/context/base-wheels/.keep
-            cp artifacts/rocm-base-wheels/*.whl docker/context/base-wheels/
-            echo "Base wheels for vLLM build:"
-            ls -lh docker/context/base-wheels/
-
-            echo "========================================"
-            echo "Building vLLM wheel with:"
-            echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
-            echo "  BUILDKITE_BRANCH: $${BUILDKITE_BRANCH}"
-            echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
-            echo "========================================"
-
-            # Build vLLM wheel using local checkout (REMOTE_VLLM=0)
-            DOCKER_BUILDKIT=1 docker build \
-              --file docker/Dockerfile.rocm \
-              --target export_vllm_wheel_release \
-              --output type=local,dest=rocm-dist \
-              --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
-              --build-arg REMOTE_VLLM=0 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-              --build-arg SCCACHE_REGION_NAME=us-west-2 \
-              --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-              .
-            echo "Built vLLM wheel:"
-            ls -lh rocm-dist/*.whl
-            # Copy wheel to artifacts directory
-            mkdir -p artifacts/rocm-vllm-wheel
-            cp rocm-dist/*.whl artifacts/rocm-vllm-wheel/
-            echo "Final vLLM wheel:"
-            ls -lh artifacts/rocm-vllm-wheel/
-        artifact_paths:
-          - "artifacts/rocm-vllm-wheel/*.whl"
-        env:
-          DOCKER_BUILDKIT: "1"
-          S3_BUCKET: "vllm-wheels"
-
-      # ROCm Job 3: Upload Wheels to S3
-      - label: ":s3: Upload ROCm Wheels to S3"
-        id: upload-rocm-wheels
-        depends_on:
-          - step: build-rocm-vllm-wheel
-            allow_failure: false
-        agents:
-          queue: cpu_queue_release
-        timeout_in_minutes: 60
-        commands:
-          # Download all wheel artifacts and run upload
-          - |
-            set -euo pipefail
-
-            # Download artifacts from current build
-            echo "Downloading artifacts from current build"
-            buildkite-agent artifact download "artifacts/rocm-base-wheels/*.whl" .
-            buildkite-agent artifact download "artifacts/rocm-vllm-wheel/*.whl" .
-
-            # # Run upload script
-            bash .buildkite/scripts/upload-rocm-wheels.sh
-        env:
-          DOCKER_BUILDKIT: "1"
-          S3_BUCKET: "vllm-wheels"
-
-      # ROCm Job 4: Annotate ROCm Wheel Release
-      - label: ":memo: Annotate ROCm wheel release"
-        id: annotate-rocm-release
-        depends_on:
-          - upload-rocm-wheels
-        agents:
-          queue: cpu_queue_release
-        commands:
-          - "bash .buildkite/scripts/annotate-rocm-release.sh"
-        env:
-          S3_BUCKET: "vllm-wheels"
-
-      # ROCm Job 5: Generate Root Index for ROCm Wheels (for release only)
-      # This is the job to create https://wheels.vllm.ai/rocm/ index allowing
-      # users to install with `uv pip install vllm --extra-index-url https://wheels.vllm.ai/rocm/`
-      - block: "Generate Root Index for ROCm Wheels for Release"
-        key: block-generate-root-index-rocm-wheels
-        depends_on: upload-rocm-wheels
-
-      - label: ":package: Generate Root Index for ROCm Wheels for Release"
-        depends_on: block-generate-root-index-rocm-wheels
-        id: generate-root-index-rocm-wheels
-        agents:
-          queue: cpu_queue_release
-        commands:
-          - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
-        env:
-          S3_BUCKET: "vllm-wheels"
-          VARIANT: "rocm723"
-
-      # ROCm Job 6: Build ROCm Release Docker Image
-      - label: ":docker: Build release image - x86_64 - ROCm"
-        id: build-rocm-release-image
-        depends_on:
-          - step: block-build-release-images
-            allow_failure: true
-          - step: build-rocm-base-wheels
-            allow_failure: false
-        agents:
-          queue: cpu_queue_release
-        timeout_in_minutes: 60
-        commands:
-          - |
-            set -euo pipefail
-            
-            # Login to ECR
-            aws ecr-public get-login-password --region us-east-1 | \
-              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-            
-            # Get ECR image tag from metadata (set by build-rocm-base-wheels)
-            ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rocm-base-image-tag 2>/dev/null || echo '')"
-            if [ -z "$${ECR_IMAGE_TAG}" ]; then
-              echo "ERROR: rocm-base-image-tag metadata not found"
-              echo "This should have been set by the build-rocm-base-wheels job"
-              exit 1
-            fi
-            
-            echo "Pulling base Docker image from ECR: $${ECR_IMAGE_TAG}"
-            
-            # Pull base Docker image from ECR
-            docker pull "$${ECR_IMAGE_TAG}"
-            
-            echo "Loaded base image: $${ECR_IMAGE_TAG}"
-            
-            # Pass the base image ECR tag to downstream steps (nightly publish)
-            buildkite-agent meta-data set "rocm-base-ecr-tag" "$${ECR_IMAGE_TAG}"
-            
-            echo "========================================"
-            echo "Building vLLM ROCm release image with:"
-            echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
-            echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
-            echo "========================================"
-            
-            # Build vLLM ROCm release image using cached base
-            DOCKER_BUILDKIT=1 docker build \
-              --build-arg max_jobs=16 \
-              --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-              --build-arg SCCACHE_REGION_NAME=us-west-2 \
-              --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-              --build-arg INSTALL_LMCACHE=true \
-              --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile.rocm .
-            
-            # Push to ECR
-            docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm
-
-            echo ""
-            echo " Successfully built and pushed ROCm release image"
-            echo "   Image: public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm"
-            echo ""
-        env:
-          DOCKER_BUILDKIT: "1"
-          S3_BUCKET: "vllm-wheels"
-
-      - label: "Publish nightly XPU image to DockerHub"
-        depends_on:
-          - create-manifest-xpu
-        if: build.env("NIGHTLY") == "1"
-        agents:
-          queue: small_cpu_queue_release
-        commands:
-          - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
-          - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
-        plugins:
-          - docker-login#v3.0.0:
-              username: vllmbot
-              password-env: DOCKERHUB_TOKEN
-        env:
-          DOCKER_BUILDKIT: "1"
-          DOCKERHUB_USERNAME: "vllmbot"
-
-      - label: "Publish nightly ROCm image to DockerHub"
-        depends_on:
-          - build-rocm-release-image
-        if: build.env("NIGHTLY") == "1"
-        agents:
-          queue: small_cpu_queue_release
-        commands:
-          - "bash .buildkite/scripts/push-nightly-builds-rocm.sh"
-          # Clean up old nightly builds (keep only last 14)
-          - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-rocm"
-          - "bash .buildkite/scripts/cleanup-nightly-builds.sh base-nightly- vllm/vllm-openai-rocm"
-        plugins:
-          - docker-login#v3.0.0:
-              username: vllmbot
-              password-env: DOCKERHUB_TOKEN
-        env:
-          DOCKER_BUILDKIT: "1"
-          DOCKERHUB_USERNAME: "vllmbot"
-
   - trigger: "perf-eval"
     label: "Trigger perf-eval (CUDA)"
     key: trigger-perf-eval-cuda
@@ -1103,7 +724,7 @@ steps:
     key: trigger-perf-eval-rocm
     if: build.env("NIGHTLY") == "1" || build.env("RUN_PERF_EVAL") == "1"
     depends_on:
-      - build-rocm-release-image
+      - build-rocm100-release-image
     async: true
     soft_fail: true
     build:
@@ -1112,138 +733,250 @@ steps:
       env:
         NIGHTLY: 1 # run all perf-eval workloads
         VLLM_COMMIT: "$BUILDKITE_COMMIT"
-        VLLM_IMAGE_ROCM: "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-rocm"
+        VLLM_IMAGE_ROCM: "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-rocm100"
 
   # =============================================================================
-  # ROCk Release Pipeline (x86_64 only)
+  # ROCm Release Pipelines (x86_64 only)
   # =============================================================================
   #
-  # Same shape as the ROCm pipeline above (Jobs 1 and 6 only), but built from
-  # docker/Dockerfile.rock_base and docker/Dockerfile.rock, which install the
-  # ROCm SDK from TheRock wheels instead of the ROCm apt repos.
+  # One group per ROCm stack, identified by its base Dockerfile (ROCM_BASE_DOCKERFILE).
+  # The scripts derive the ROCm version, wheel variant, cache keys and image tags from
+  # it (.buildkite/scripts/rocm/stack.sh), so both groups share one set of step
+  # definitions, anchored on the ROCm 10.0 group.
   #
-  # Note: ROCm SDK version is determined by ROCM_SDK_VERSION in
-  # docker/Dockerfile.rock_base
+  # - ROCm 10.0 (TheRock, docker/Dockerfile.rocm{,_base}) is the default: it owns the
+  #   :nightly / :latest images and the root wheel index.
+  # - ROCm 7.2 (docker/Dockerfile.rocm_72{,_base}) is kept for a transition period as
+  #   rocm72 images and wheels.
   #
+  # vLLM version is determined by the Buildkite checkout (like the CUDA pipeline).
   # =============================================================================
 
-  - group: "Build ROCk Image"
-    key: "build-rock-image"
+  - group: "Build ROCm 10.0 Wheel / Image"
+    key: "build-rocm100"
     depends_on: ~
     steps:
-      # ROCk Job 1: Build ROCk Base Image (with ECR caching)
-      - label: ":rocm: Build ROCk Base Image"
-        id: build-rock-base-image
+      - label: ":rocm: Build ROCm 10.0 Base Image & Wheels"
+        id: build-rocm100-base
         depends_on: ~
-        agents:
+        agents: &rocm_build_agents
           queue: cpu_queue_release
         timeout_in_minutes: 360
-        commands:
+        env: &rocm100_env
+          ROCM_BASE_DOCKERFILE: "docker/Dockerfile.rocm_base"
+          ROCM_DEFAULT_ARG: "--default"
+          DOCKER_BUILDKIT: "1"
+          S3_BUCKET: "vllm-wheels"
+          DOCKERHUB_USERNAME: "vllmbot"
+        commands: &rocm_base_commands
           - |
             set -euo pipefail
-
-            # The SDK version, torch wheels and component branches all live in
-            # the Dockerfile, so its hash is the cache key.
-            CACHE_KEY=$$(sha256sum docker/Dockerfile.rock_base | cut -c1-16)
-            ECR_CACHE_TAG="public.ecr.aws/q9t5s3a7/vllm-release-repo:$${CACHE_KEY}-rock-base"
-
-            echo "========================================"
-            echo "ROCk Base Build Configuration"
-            echo "========================================"
-            echo "  CACHE_KEY: $${CACHE_KEY}"
-            echo "  ECR_CACHE_TAG: $${ECR_CACHE_TAG}"
-            echo "========================================"
-
-            # Login to ECR
+            source .buildkite/scripts/rocm/stack.sh "$${ROCM_BASE_DOCKERFILE}"
+            SCCACHE_ARGS=(--build-arg USE_SCCACHE=1 --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache
+              --build-arg SCCACHE_REGION_NAME=us-west-2 --build-arg SCCACHE_S3_NO_CREDENTIALS=0)
             aws ecr-public get-login-password --region us-east-1 | \
               docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
 
-            if docker manifest inspect "$${ECR_CACHE_TAG}" > /dev/null 2>&1; then
-              echo "ECR image cache HIT"
+            # Base image, cached in ECR by the base Dockerfile's hash
+            if docker manifest inspect "$${ROCM_STACK_ECR_BASE}" > /dev/null 2>&1; then
+              echo "ECR image cache HIT: $${ROCM_STACK_ECR_BASE}"
             else
-              echo "CACHE MISS - Building from scratch..."
-
-              DOCKER_BUILDKIT=1 docker buildx build \
-                --file docker/Dockerfile.rock_base \
-                --tag "$${ECR_CACHE_TAG}" \
-                --build-arg USE_SCCACHE=1 \
-                --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-                --build-arg SCCACHE_REGION_NAME=us-west-2 \
-                --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-                --push \
-                .
+              docker buildx build --file "$${ROCM_BASE_DOCKERFILE}" --tag "$${ROCM_STACK_ECR_BASE}" \
+                "$${SCCACHE_ARGS[@]}" --push .
             fi
 
-            # Save ECR tag for downstream jobs
-            buildkite-agent meta-data set "rock-base-image-tag" "$${ECR_CACHE_TAG}"
-        env:
-          DOCKER_BUILDKIT: "1"
+            # Base wheels, cached in S3 (see cache-rocm-base-wheels.sh)
+            CACHE=.buildkite/scripts/cache-rocm-base-wheels.sh
+            if [ "$$($${CACHE} "$${ROCM_BASE_DOCKERFILE}" check)" = "hit" ]; then
+              $${CACHE} "$${ROCM_BASE_DOCKERFILE}" download
+            else
+              DEBS="rocm-base-debs:$${BUILDKITE_BUILD_NUMBER}-$${ROCM_STACK_VARIANT}"
+              docker buildx build --file "$${ROCM_BASE_DOCKERFILE}" --tag "$${DEBS}" \
+                --target debs_wheel_release "$${SCCACHE_ARGS[@]}" --load .
+              mkdir -p "$${ROCM_STACK_BASE_WHEELS_DIR}"
+              cid=$$(docker create "$${DEBS}")
+              docker cp "$${cid}":/app/debs/. "$${ROCM_STACK_BASE_WHEELS_DIR}/"
+              docker rm "$${cid}"
+              $${CACHE} "$${ROCM_BASE_DOCKERFILE}" upload
+            fi
+            buildkite-agent artifact upload "$${ROCM_STACK_BASE_WHEELS_DIR}/*.whl"
 
-      # ROCk Job 2: Build ROCk Release Docker Image
-      - label: ":docker: Build release image - x86_64 - ROCk"
-        id: build-rock-release-image
+      - label: ":python: Build vLLM ROCm 10.0 Wheel - x86_64"
+        id: build-rocm100-vllm-wheel
+        depends_on: build-rocm100-base
+        agents: *rocm_build_agents
+        timeout_in_minutes: 180
+        env: *rocm100_env
+        commands: &rocm_vllm_wheel_commands
+          - |
+            set -euo pipefail
+            source .buildkite/scripts/rocm/stack.sh "$${ROCM_BASE_DOCKERFILE}"
+            # Buildkite's default fetch doesn't update moved tags; version detection needs them
+            git fetch --tags --force origin
+            buildkite-agent artifact download "$${ROCM_STACK_BASE_WHEELS_DIR}/*.whl" .
+            aws ecr-public get-login-password --region us-east-1 | \
+              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+            docker pull "$${ROCM_STACK_ECR_BASE}"
+            mkdir -p docker/context/base-wheels
+            touch docker/context/base-wheels/.keep
+            cp "$${ROCM_STACK_BASE_WHEELS_DIR}"/*.whl docker/context/base-wheels/
+            docker build --file "$${ROCM_STACK_DOCKERFILE}" --target export_vllm_wheel_release \
+              --output type=local,dest=rocm-dist \
+              --build-arg BASE_IMAGE="$${ROCM_STACK_ECR_BASE}" --build-arg REMOTE_VLLM=0 \
+              --build-arg GIT_REPO_CHECK=1 --build-arg USE_SCCACHE=1 \
+              --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
+              --build-arg SCCACHE_REGION_NAME=us-west-2 --build-arg SCCACHE_S3_NO_CREDENTIALS=0 .
+            mkdir -p "$${ROCM_STACK_VLLM_WHEEL_DIR}"
+            cp rocm-dist/*.whl "$${ROCM_STACK_VLLM_WHEEL_DIR}/"
+            buildkite-agent artifact upload "$${ROCM_STACK_VLLM_WHEEL_DIR}/*.whl"
+
+      - label: ":s3: Upload ROCm 10.0 Wheels to S3"
+        id: upload-rocm100-wheels
+        depends_on: build-rocm100-vllm-wheel
+        agents: *rocm_build_agents
+        timeout_in_minutes: 60
+        env: *rocm100_env
+        commands: &rocm_upload_commands
+          - |
+            set -euo pipefail
+            source .buildkite/scripts/rocm/stack.sh "$${ROCM_BASE_DOCKERFILE}"
+            buildkite-agent artifact download "$${ROCM_STACK_BASE_WHEELS_DIR}/*.whl" .
+            buildkite-agent artifact download "$${ROCM_STACK_VLLM_WHEEL_DIR}/*.whl" .
+            bash .buildkite/scripts/upload-rocm-wheels.sh "$${ROCM_BASE_DOCKERFILE}"
+
+      - label: ":memo: Annotate ROCm 10.0 wheel release"
+        id: annotate-rocm100-release
+        depends_on: upload-rocm100-wheels
+        agents: *rocm_build_agents
+        env: *rocm100_env
+        commands: &rocm_annotate_commands
+          - bash .buildkite/scripts/annotate-rocm-release.sh "$${ROCM_BASE_DOCKERFILE}" $${ROCM_DEFAULT_ARG}
+
+      # Generate https://wheels.vllm.ai/rocm/ (the default stack's variant), so users can
+      # `uv pip install vllm --extra-index-url https://wheels.vllm.ai/rocm/` (release only)
+      - block: "Generate Root Index for ROCm Wheels for Release"
+        key: block-generate-root-index-rocm-wheels
+        depends_on:
+          - upload-rocm100-wheels
+          - upload-rocm72-wheels
+
+      - label: ":package: Generate Root Index for ROCm Wheels for Release"
+        depends_on: block-generate-root-index-rocm-wheels
+        id: generate-root-index-rocm-wheels
+        agents: *rocm_build_agents
+        env: *rocm100_env
+        commands:
+          - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
+
+      - label: ":docker: Build release image - x86_64 - ROCm 10.0"
+        id: build-rocm100-release-image
         depends_on:
           - step: block-build-release-images
             allow_failure: true
-          - step: build-rock-base-image
+          - step: build-rocm100-base
             allow_failure: false
-        agents:
-          queue: cpu_queue_release
+        agents: *rocm_build_agents
         timeout_in_minutes: 180
-        commands:
+        env: *rocm100_env
+        commands: &rocm_release_image_commands
           - |
             set -euo pipefail
-
-            # Login to ECR
+            source .buildkite/scripts/rocm/stack.sh "$${ROCM_BASE_DOCKERFILE}"
             aws ecr-public get-login-password --region us-east-1 | \
               docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-
-            ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rock-base-image-tag 2>/dev/null || echo '')"
-            if [ -z "$${ECR_IMAGE_TAG}" ]; then
-              echo "ERROR: rock-base-image-tag metadata not found"
-              echo "This should have been set by the build-rock-base-image job"
-              exit 1
-            fi
-
-            docker pull "$${ECR_IMAGE_TAG}"
-
-            # Pass the base image ECR tag to downstream steps (nightly publish)
-            buildkite-agent meta-data set "rock-base-ecr-tag" "$${ECR_IMAGE_TAG}"
-
-            echo "========================================"
-            echo "Building vLLM ROCk release image with:"
-            echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
-            echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
-            echo "========================================"
-
-            DOCKER_BUILDKIT=1 docker build \
-              --build-arg max_jobs=16 \
-              --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
-              --build-arg USE_SCCACHE=1 \
+            docker pull "$${ROCM_STACK_ECR_BASE}"
+            docker build --file "$${ROCM_STACK_DOCKERFILE}" --target vllm-openai \
+              --tag "$${ROCM_STACK_ECR_IMAGE}" --progress plain \
+              --build-arg max_jobs=16 --build-arg BASE_IMAGE="$${ROCM_STACK_ECR_BASE}" \
+              --build-arg INSTALL_LMCACHE=true --build-arg USE_SCCACHE=1 \
               --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-              --build-arg SCCACHE_REGION_NAME=us-west-2 \
-              --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-              --build-arg INSTALL_LMCACHE=true \
-              --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rock \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile.rock .
+              --build-arg SCCACHE_REGION_NAME=us-west-2 --build-arg SCCACHE_S3_NO_CREDENTIALS=0 .
+            docker push "$${ROCM_STACK_ECR_IMAGE}"
 
-            docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rock
-        env:
+      - label: "Publish nightly ROCm 10.0 image to DockerHub"
+        id: publish-rocm100-nightly
+        depends_on: build-rocm100-release-image
+        if: build.env("NIGHTLY") == "1"
+        agents: &rocm_publish_agents
+          queue: small_cpu_queue_release
+        env: *rocm100_env
+        commands: &rocm_publish_nightly_commands
+          - bash .buildkite/scripts/push-nightly-builds-rocm.sh "$${ROCM_BASE_DOCKERFILE}" $${ROCM_DEFAULT_ARG}
+        plugins: &rocm_dockerhub_login
+          - docker-login#v3.0.0:
+              username: vllmbot
+              password-env: DOCKERHUB_TOKEN
+
+  - group: "Build ROCm 7.2 Wheel / Image"
+    key: "build-rocm72"
+    depends_on: ~
+    steps:
+      - label: ":rocm: Build ROCm 7.2 Base Image & Wheels"
+        id: build-rocm72-base
+        depends_on: ~
+        agents: *rocm_build_agents
+        timeout_in_minutes: 360
+        env: &rocm72_env
+          ROCM_BASE_DOCKERFILE: "docker/Dockerfile.rocm_72_base"
+          ROCM_DEFAULT_ARG: ""
           DOCKER_BUILDKIT: "1"
+          S3_BUCKET: "vllm-wheels"
+          DOCKERHUB_USERNAME: "vllmbot"
+        commands: *rocm_base_commands
 
-      - label: "Publish nightly ROCk image to DockerHub"
+      - label: ":python: Build vLLM ROCm 7.2 Wheel - x86_64"
+        id: build-rocm72-vllm-wheel
+        depends_on: build-rocm72-base
+        agents: *rocm_build_agents
+        timeout_in_minutes: 180
+        env: *rocm72_env
+        commands: *rocm_vllm_wheel_commands
+
+      - label: ":s3: Upload ROCm 7.2 Wheels to S3"
+        id: upload-rocm72-wheels
+        depends_on: build-rocm72-vllm-wheel
+        agents: *rocm_build_agents
+        timeout_in_minutes: 60
+        env: *rocm72_env
+        commands: *rocm_upload_commands
+
+      - label: ":memo: Annotate ROCm 7.2 wheel release"
+        id: annotate-rocm72-release
+        depends_on: upload-rocm72-wheels
+        agents: *rocm_build_agents
+        env: *rocm72_env
+        commands: *rocm_annotate_commands
+
+      - label: ":docker: Build release image - x86_64 - ROCm 7.2"
+        id: build-rocm72-release-image
         depends_on:
-          - build-rock-release-image
+          - step: block-build-release-images
+            allow_failure: true
+          - step: build-rocm72-base
+            allow_failure: false
+        agents: *rocm_build_agents
+        timeout_in_minutes: 180
+        env: *rocm72_env
+        commands: *rocm_release_image_commands
+
+      - label: "Publish nightly ROCm 7.2 image to DockerHub"
+        id: publish-rocm72-nightly
+        depends_on: build-rocm72-release-image
+        if: build.env("NIGHTLY") == "1"
+        agents: *rocm_publish_agents
+        env: *rocm72_env
+        commands: *rocm_publish_nightly_commands
+        plugins: *rocm_dockerhub_login
+
+      - label: "Publish nightly XPU image to DockerHub"
+        depends_on:
+          - create-manifest-xpu
         if: build.env("NIGHTLY") == "1"
         agents:
           queue: small_cpu_queue_release
         commands:
-          - "bash .buildkite/scripts/push-nightly-builds-rocm.sh rock rocm100"
-          # Clean up old nightly builds (keep only last 14)
-          - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly-rocm100- vllm/vllm-openai-rocm"
-          - "bash .buildkite/scripts/cleanup-nightly-builds.sh base-nightly-rocm100- vllm/vllm-openai-rocm"
+          - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
+          - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
         plugins:
           - docker-login#v3.0.0:
               username: vllmbot
@@ -1264,7 +997,8 @@ steps:
       - create-multi-arch-manifest-ubuntu2404
       - create-multi-arch-manifest-cuda-12-9-ubuntu2404
       - create-manifest-xpu
-      - build-rocm-release-image
+      - build-rocm100-release-image
+      - build-rocm72-release-image
       - input-release-version
       # Wait for CPU builds if their block steps were unblocked, so publish
       # doesn't race the in-progress CPU build. allow_failure lets publish
@@ -1289,6 +1023,7 @@ steps:
       - "cuda-13-0-ubuntu-24-04"
       - "cuda-12-9-ubuntu-24-04"
       - "rocm"
+      - "rocm72"
       - "xpu"
       - "cpu"
     plugins:

.buildkite/scripts/annotate-rocm-release.sh

@@ -2,15 +2,23 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 #
-# Generate Buildkite annotation for ROCm wheel release
+# Generate Buildkite annotation for a ROCm stack's wheel release.
+#
+# Usage: annotate-rocm-release.sh <base-dockerfile> [--default]
+#   base-dockerfile  e.g. docker/Dockerfile.rocm_base or docker/Dockerfile.rocm_72_base
+#   --default        the stack also owns the plain :latest / :v<ver> Docker Hub tags
 set -ex
 
-# Extract build configuration from Dockerfile.rocm_base (single source of truth)
-# Extract ROCm version dynamically from Dockerfile.rocm_base
-# BASE_IMAGE format: rocm/dev-ubuntu-22.04:7.0-complete -> extracts "7.0"
-ROCM_VERSION=$(grep -E '^ARG BASE_IMAGE=' docker/Dockerfile.rocm_base | sed -E 's/.*:([0-9]+\.[0-9]+).*/\1/' || echo "unknown")
-PYTHON_VERSION=$(grep '^ARG PYTHON_VERSION=' docker/Dockerfile.rocm_base | sed 's/^ARG PYTHON_VERSION=//')
-PYTORCH_ROCM_ARCH=$(grep '^ARG PYTORCH_ROCM_ARCH=' docker/Dockerfile.rocm_base | sed 's/^ARG PYTORCH_ROCM_ARCH=//')
+# shellcheck source=.buildkite/scripts/rocm/stack.sh
+source "$(dirname "${BASH_SOURCE[0]}")/rocm/stack.sh" "${1:?Usage: $0 <base-dockerfile> [--default]}"
+TAG_VARIANTS=("-${ROCM_STACK_VARIANT}")
+if [[ "${2:-}" == "--default" ]]; then
+  TAG_VARIANTS=("" "-${ROCM_STACK_VARIANT}")
+fi
+ROCM_VERSION="$ROCM_STACK_VERSION"
+VARIANT="$ROCM_STACK_VARIANT"
+PYTHON_VERSION=$(sed -nE 's/^ARG PYTHON_VERSION="?([^" ]+)"?.*/\1/p' "$ROCM_STACK_BASE_DOCKERFILE")
+PYTORCH_ROCM_ARCH=$(sed -nE 's/^ARG PYTORCH_ROCM_ARCH="?([^"]+)"?.*/\1/p' "$ROCM_STACK_BASE_DOCKERFILE")
 
 # Get release version, default to 1.0.0.dev for nightly/per-commit builds
 RELEASE_VERSION=$(buildkite-agent meta-data get release-version 2>/dev/null || echo "")
@@ -18,22 +26,32 @@ if [ -z "${RELEASE_VERSION}" ]; then
   RELEASE_VERSION="1.0.0.dev"
 fi
 
-ROCM_BASE_CACHE_KEY=$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
-
 # S3 URLs
 S3_BUCKET="${S3_BUCKET:-vllm-wheels}"
 S3_REGION="${AWS_DEFAULT_REGION:-us-west-2}"
-S3_URL="http://${S3_BUCKET}.s3-website-${S3_REGION}.amazonaws.com"
+S3_HOST="${S3_BUCKET}.s3-website-${S3_REGION}.amazonaws.com"
+S3_URL="http://${S3_HOST}"
+WHEEL_PATH="rocm/${BUILDKITE_COMMIT}/${VARIANT}-wheels/"
 
-# Format ROCm version for path (e.g., "7.1" -> "rocm710")
-ROCM_VERSION_PATH="rocm$(echo "${ROCM_VERSION}" | tr -d '.')"
-ROCM_PATH="rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}"
-buildkite-agent annotate --style 'success' --context 'rocm-release-workflow' << EOF
-## ROCm Wheel and Docker Image Releases
+DOCKER_CMDS="docker pull ${ROCM_STACK_ECR_BASE}
+docker pull ${ROCM_STACK_ECR_IMAGE}"
+for suffix in "${TAG_VARIANTS[@]}"; do
+  for tag in "latest${suffix}" "v${RELEASE_VERSION}${suffix}"; do
+    DOCKER_CMDS+="
+docker tag ${ROCM_STACK_ECR_BASE} vllm/vllm-openai-rocm:${tag}-base
+docker tag ${ROCM_STACK_ECR_IMAGE} vllm/vllm-openai-rocm:${tag}
+docker push vllm/vllm-openai-rocm:${tag}-base
+docker push vllm/vllm-openai-rocm:${tag}"
+  done
+done
+
+buildkite-agent annotate --style 'success' --context "rocm-release-workflow-${VARIANT}" << EOF
+## ROCm ${ROCM_VERSION} Wheel and Docker Image Releases
 ### Build Configuration
 | Setting | Value |
 |---------|-------|
 | **ROCm Version** | ${ROCM_VERSION} |
+| **Index Variant** | ${VARIANT} |
 | **Python Version** | ${PYTHON_VERSION} |
 | **GPU Architectures** | ${PYTORCH_ROCM_ARCH} |
 | **Branch** | \`${BUILDKITE_BRANCH}\` |
@@ -41,49 +59,20 @@ buildkite-agent annotate --style 'success' --context 'rocm-release-workflow' <<
 
 ### :package: Installation
 
-**Install from this build (by commit):**
-
 \`\`\`bash
-pip install vllm --extra-index-url ${S3_URL}/${ROCM_PATH}/ --trusted-host ${S3_BUCKET}.s3-website-${S3_REGION}.amazonaws.com
-
-# Example for ROCm ${ROCM_VERSION}:
-pip install vllm --extra-index-url ${S3_URL}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/ --trusted-host ${S3_BUCKET}.s3-website-${S3_REGION}.amazonaws.com
-\`\`\`
-
-**Install from nightly (if published):**
-
-\`\`\`bash
-pip install vllm --extra-index-url ${S3_URL}/rocm/nightly/ --trusted-host ${S3_BUCKET}.s3-website-${S3_REGION}.amazonaws.com
+# This build (by commit)
+pip install vllm --extra-index-url ${S3_URL}/rocm/${BUILDKITE_COMMIT}/${VARIANT}/ --trusted-host ${S3_HOST}
+# Nightly (if published)
+pip install vllm --extra-index-url ${S3_URL}/rocm/nightly/${VARIANT}/ --trusted-host ${S3_HOST}
 \`\`\`
 
 ### :floppy_disk: Download Wheels Directly
 
 \`\`\`bash
-# List all ROCm wheels
-aws s3 ls s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/
-# Download specific wheels
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/vllm-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/torch-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/triton-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/triton-kernels-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/torchvision-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/torchaudio-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/amdsmi-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/amd_aiter-*.whl .
-aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/flash-attn-*.whl .
+aws s3 ls s3://${S3_BUCKET}/${WHEEL_PATH}
+aws s3 cp --recursive --exclude "*" --include "*.whl" s3://${S3_BUCKET}/${WHEEL_PATH} .
 \`\`\`
 
-### :gear: Included Packages
-- **vllm**: vLLM with ROCm support
-- **torch**: PyTorch built for ROCm ${ROCM_VERSION}
-- **triton**: Triton
-- **triton-kernels**: Triton kernels
-- **torchvision**: TorchVision for ROCm PyTorch
-- **torchaudio**: Torchaudio for ROCm PyTorch
-- **amdsmi**: AMD SMI Python bindings
-- **amd_aiter**: Aiter for ROCm
-- **flash-attn**: Flash Attention for ROCm
-
 ### :warning: Notes
 - These wheels are built for **ROCm ${ROCM_VERSION}** and will NOT work with CUDA GPUs
 - Supported GPU architectures: ${PYTORCH_ROCM_ARCH}
@@ -91,23 +80,8 @@ aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/flash-
 
 ### :package: Docker Image Release
 
-To download and upload the image:
-
 \`\`\`
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm-base
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base
-docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base vllm/vllm-openai-rocm:latest-base
-docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
-docker push vllm/vllm-openai-rocm:latest-base
-docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}
-docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT} vllm/vllm-openai-rocm:latest
-docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT} vllm/vllm-openai-rocm:v${RELEASE_VERSION}
-docker push vllm/vllm-openai-rocm:latest
-docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}
+${DOCKER_CMDS}
 \`\`\`
 
 EOF

.buildkite/scripts/cache-rocm-base-wheels.sh

@@ -5,39 +5,33 @@
 # Cache helper for ROCm base wheels
 #
 # This script manages caching of pre-built ROCm base wheels (torch, triton, etc.)
-# to avoid rebuilding them when Dockerfile.rocm_base hasn't changed.
+# to avoid rebuilding them when the stack's base Dockerfile hasn't changed.
 #
-# Usage:
-#   cache-rocm-base-wheels.sh check    - Check if cache exists, outputs "hit" or "miss"
-#   cache-rocm-base-wheels.sh upload   - Upload wheels to cache
-#   cache-rocm-base-wheels.sh download - Download wheels from cache
-#   cache-rocm-base-wheels.sh key      - Output the cache key
+# Usage: cache-rocm-base-wheels.sh <base-dockerfile> <command>
+#   base-dockerfile  e.g. docker/Dockerfile.rocm_base or docker/Dockerfile.rocm_72_base
+# Commands:
+#   check    - Check if cache exists, outputs "hit" or "miss"
+#   upload   - Upload wheels to cache
+#   download - Download wheels from cache
+#   key      - Output the cache key
+#   path     - Output the full S3 cache path
 #
 # Environment variables:
 #   S3_BUCKET          - S3 bucket name (default: vllm-wheels)
 #
-# Note: ROCm version is determined by BASE_IMAGE in Dockerfile.rocm_base,
-#       so changes to ROCm version are captured by the Dockerfile hash.
+# The cache key hashes the base Dockerfile and the repo files it COPYs in, which
+# capture the ROCm version and every component pin (see rocm/stack.sh).
 
 set -euo pipefail
 
 BUCKET="${S3_BUCKET:-vllm-wheels}"
-DOCKERFILE="docker/Dockerfile.rocm_base"
 CACHE_PREFIX="rocm/cache"
 
-# Generate hash from Dockerfile content + build args
-generate_cache_key() {
-    # Include Dockerfile content
-    if [[ ! -f "$DOCKERFILE" ]]; then
-        echo "ERROR: Dockerfile not found: $DOCKERFILE" >&2
-        exit 1
-    fi
-    local dockerfile_hash=$(sha256sum "$DOCKERFILE" | cut -c1-16)
-
-    echo "${dockerfile_hash}"
-}
-
-CACHE_KEY=$(generate_cache_key)
+# shellcheck source=.buildkite/scripts/rocm/stack.sh
+source "$(dirname "${BASH_SOURCE[0]}")/rocm/stack.sh" "${1:?Usage: $0 <base-dockerfile> <command>}"
+shift
+CACHE_KEY="$ROCM_STACK_WHEEL_KEY"
+WHEELS_DIR="$ROCM_STACK_BASE_WHEELS_DIR"
 CACHE_PATH="s3://${BUCKET}/${CACHE_PREFIX}/${CACHE_KEY}/"
 
 case "${1:-}" in
@@ -67,19 +61,19 @@ case "${1:-}" in
         echo "Cache path: ${CACHE_PATH}"
         echo ""
 
-        if [[ ! -d "artifacts/rocm-base-wheels" ]]; then
-            echo "ERROR: artifacts/rocm-base-wheels directory not found" >&2
+        if [[ ! -d "${WHEELS_DIR}" ]]; then
+            echo "ERROR: ${WHEELS_DIR} directory not found" >&2
             exit 1
         fi
 
-        WHEEL_COUNT=$(find artifacts/rocm-base-wheels -maxdepth 1 -name '*.whl' 2>/dev/null | wc -l)
+        WHEEL_COUNT=$(find "${WHEELS_DIR}" -maxdepth 1 -name '*.whl' 2>/dev/null | wc -l)
         if [[ "$WHEEL_COUNT" -eq 0 ]]; then
-            echo "ERROR: No wheels found in artifacts/rocm-base-wheels/" >&2
+            echo "ERROR: No wheels found in ${WHEELS_DIR}/" >&2
             exit 1
         fi
 
         echo "Uploading $WHEEL_COUNT wheels..."
-        aws s3 cp --recursive artifacts/rocm-base-wheels/ "${CACHE_PATH}"
+        aws s3 cp --recursive "${WHEELS_DIR}/" "${CACHE_PATH}"
 
         echo ""
         echo "Cache upload complete!"
@@ -93,17 +87,17 @@ case "${1:-}" in
         echo "Cache key: ${CACHE_KEY}"
         echo "Cache path: ${CACHE_PATH}"
         echo ""
-        mkdir -p artifacts/rocm-base-wheels
+        mkdir -p "${WHEELS_DIR}"
         
         # Use sync with include/exclude to only download .whl files
-        aws s3 sync "${CACHE_PATH}" artifacts/rocm-base-wheels/ \
+        aws s3 sync "${CACHE_PATH}" "${WHEELS_DIR}/" \
             --exclude "*" \
             --include "*.whl"
         
         echo ""
         echo "Downloaded wheels:"
-        find artifacts/rocm-base-wheels -maxdepth 1 -name '*.whl' -exec ls -lh {} \;
-        WHEEL_COUNT=$(find artifacts/rocm-base-wheels -maxdepth 1 -name '*.whl' 2>/dev/null | wc -l)
+        find "${WHEELS_DIR}" -maxdepth 1 -name '*.whl' -exec ls -lh {} \;
+        WHEEL_COUNT=$(find "${WHEELS_DIR}" -maxdepth 1 -name '*.whl' 2>/dev/null | wc -l)
         echo ""
         echo "Total: $WHEEL_COUNT wheels"
         echo "========================================"
@@ -118,12 +112,12 @@ case "${1:-}" in
         ;;
 
     *)
-        echo "Usage: $0 {check|upload|download|key|path}" >&2
+        echo "Usage: $0 <base-dockerfile> {check|upload|download|key|path}" >&2
         echo "" >&2
         echo "Commands:" >&2
         echo "  check    - Check if cache exists, outputs 'hit' or 'miss'" >&2
-        echo "  upload   - Upload wheels from artifacts/rocm-base-wheels/ to cache" >&2
-        echo "  download - Download wheels from cache to artifacts/rocm-base-wheels/" >&2
+        echo "  upload   - Upload wheels from ${WHEELS_DIR}/ to cache" >&2
+        echo "  download - Download wheels from cache to ${WHEELS_DIR}/" >&2
         echo "  key      - Output the cache key" >&2
         echo "  path     - Output the full S3 cache path" >&2
         exit 1

.buildkite/scripts/ci-bake-rocm.sh

@@ -17,7 +17,7 @@ DEFAULT_REPO_SLUG="vllm-project/vllm"
 DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
 DEFAULT_CI_BASE_CONTENT_FILES=".dockerignore requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt tools/install_torchcodec_rocm.sh rust-toolchain.toml tests/vllm_test_utils"
 DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
-DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust-toolchain-input rust-toolchain build_nixl lmcache_source build_lmcache build_rocshmem build_deepep mori_base ci_base"
+DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust-toolchain-input rust-toolchain build_nixl lmcache_source build_lmcache build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="3"
 # ROCm CI forces REMOTE_VLLM=0, so content identity covers only the selected
 # local-source stages rather than unreachable remote-fetch alternatives.
@@ -555,9 +555,6 @@ configure_custom_rocm_stages() {
     CI_BASE_DOCKERFILE_STAGES="${CI_BASE_DOCKERFILE_STAGES:-${stages}}"
     ROCM_CSRC_DOCKERFILE_STAGES="${ROCM_CSRC_DOCKERFILE_STAGES:-${stages}}"
     ROCM_RUST_DOCKERFILE_STAGES="${ROCM_RUST_DOCKERFILE_STAGES:-${stages}}"
-    if [[ " ${stages} " != *" build_rocshmem "* ]]; then
-        unset ROCSHMEM_BRANCH ROCSHMEM_CACHE_KEY DEEPEP_CACHE_KEY
-    fi
 }
 
 compose_dependency_cache_key() {
@@ -1724,14 +1721,11 @@ ci_base_metadata_pairs() {
     metadata_pair "vllm.rocm.nixl_commit" "${NIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "NIXL_BRANCH")}"
     metadata_pair "vllm.rocm.ucx_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_REPO")"
     metadata_pair "vllm.rocm.ucx_commit" "${UCX_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_BRANCH")}"
-    metadata_pair "vllm.rocm.rocshmem_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "ROCSHMEM_REPO")"
-    metadata_pair "vllm.rocm.rocshmem_commit" "${ROCSHMEM_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "ROCSHMEM_BRANCH")}"
     metadata_pair "vllm.rocm.deepep_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_REPO")"
     metadata_pair "vllm.rocm.deepep_commit" "${DEEPEP_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_BRANCH")}"
     metadata_pair "vllm.rocm.deepep_nic" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_NIC")"
     metadata_pair "vllm.rocm.deepep_rocm_arch" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_ROCM_ARCH")"
     metadata_pair "vllm.rocm.nixl_cache_key" "${NIXL_CACHE_KEY:-}"
-    metadata_pair "vllm.rocm.rocshmem_cache_key" "${ROCSHMEM_CACHE_KEY:-}"
     metadata_pair "vllm.rocm.deepep_cache_key" "${DEEPEP_CACHE_KEY:-}"
 }
 
@@ -2314,7 +2308,7 @@ extract_dependency_pins() {
         return 0
     fi
 
-    for var in NIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
+    for var in NIXL_BRANCH UCX_BRANCH DEEPEP_BRANCH; do
         if [[ -n "${!var:-}" ]]; then
             echo "Using provided ${var}: ${!var}"
             continue
@@ -2336,17 +2330,14 @@ compute_dependency_cache_keys() {
     local dockerfile_rocm=""
     local nixl_branch=""
     local ucx_branch=""
-    local rocshmem_branch=""
     local deepep_branch=""
     local nixl_material=""
-    local rocshmem_material=""
     local deepep_material=""
 
     bake_dir=$(dirname "${VLLM_BAKE_FILE}")
     dockerfile_rocm="${CI_BASE_DOCKERFILE:-${bake_dir}/Dockerfile.rocm}"
     nixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "NIXL_BRANCH")
     ucx_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "UCX_BRANCH")
-    rocshmem_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "ROCSHMEM_BRANCH")
     deepep_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "DEEPEP_BRANCH")
 
     if [[ -n "${nixl_branch}" && -n "${ucx_branch}" ]]; then
@@ -2360,22 +2351,11 @@ compute_dependency_cache_keys() {
         echo "NIXL dependency cache key: ${NIXL_CACHE_KEY}"
     fi
 
-    if [[ -n "${rocshmem_branch}" ]]; then
-        rocshmem_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_rocshmem")
-        ROCSHMEM_CACHE_KEY=$(
-            compose_dependency_cache_key \
-                "${rocshmem_branch}" \
-                "${rocshmem_material}"
-        )
-        export ROCSHMEM_CACHE_KEY
-        echo "ROCShmem dependency cache key: ${ROCSHMEM_CACHE_KEY}"
-    fi
-
-    if [[ -n "${deepep_branch}" && -n "${rocshmem_branch}" ]]; then
-        deepep_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_rocshmem build_deepep")
+    if [[ -n "${deepep_branch}" ]]; then
+        deepep_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_deepep")
         DEEPEP_CACHE_KEY=$(
             compose_dependency_cache_key \
-                "${deepep_branch}-rocshmem-${rocshmem_branch}" \
+                "${deepep_branch}" \
                 "${deepep_material}"
         )
         export DEEPEP_CACHE_KEY
@@ -2415,18 +2395,11 @@ dependency_cache_ref_for_target() {
                 printf '%s\n' "${cache_repo}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH:-}"
             fi
             ;;
-        rocshmem-rocm-ci)
-            if [[ -n "${ROCSHMEM_CACHE_KEY:-}" ]]; then
-                printf '%s\n' "${cache_repo}:rocshmem-rocm-${ROCSHMEM_CACHE_KEY}"
-            elif [[ -n "${ROCSHMEM_BRANCH:-}" ]]; then
-                printf '%s\n' "${cache_repo}:rocshmem-rocm-${ROCSHMEM_BRANCH}"
-            fi
-            ;;
         deepep-rocm-ci)
             if [[ -n "${DEEPEP_CACHE_KEY:-}" ]]; then
                 printf '%s\n' "${cache_repo}:deepep-rocm-${DEEPEP_CACHE_KEY}"
             elif [[ -n "${DEEPEP_BRANCH:-}" ]]; then
-                printf '%s\n' "${cache_repo}:deepep-rocm-${DEEPEP_BRANCH}-rocshmem-${ROCSHMEM_BRANCH:-}"
+                printf '%s\n' "${cache_repo}:deepep-rocm-${DEEPEP_BRANCH}"
             fi
             ;;
     esac
@@ -2444,7 +2417,6 @@ add_dependency_cache_target() {
 resolve_ci_base_dependency_targets() {
     local mode="${ROCM_DEP_CACHE_EXPORT_MODE:-missing}"
     local nixl_ref=""
-    local rocshmem_ref=""
     local deepep_ref=""
 
     [[ "${TARGET}" == "ci-base-rocm-ci-with-deps" ]] || return 0
@@ -2452,7 +2424,7 @@ resolve_ci_base_dependency_targets() {
     case "${mode}" in
         always)
             echo "ROCM_DEP_CACHE_EXPORT_MODE=always; exporting all dependency caches serially"
-            for target in nixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
+            for target in nixl-rocm-ci deepep-rocm-ci; do
                 if [[ -n "$(dependency_cache_ref_for_target "${target}")" ]]; then
                     add_dependency_cache_target "${target}"
                 fi
@@ -2482,16 +2454,6 @@ resolve_ci_base_dependency_targets() {
         fi
     fi
 
-    if [[ "${mode}" != "always" && -n "${ROCSHMEM_CACHE_KEY:-}" ]]; then
-        rocshmem_ref=$(dependency_cache_ref_for_target "rocshmem-rocm-ci")
-        if dependency_cache_ref_exists "${rocshmem_ref}"; then
-            echo "ROCShmem dependency cache exists: ${rocshmem_ref}"
-        else
-            echo "ROCShmem dependency cache missing; will seed: ${rocshmem_ref}"
-            add_dependency_cache_target "rocshmem-rocm-ci"
-        fi
-    fi
-
     if [[ "${mode}" != "always" && -n "${DEEPEP_CACHE_KEY:-}" ]]; then
         deepep_ref=$(dependency_cache_ref_for_target "deepep-rocm-ci")
         if dependency_cache_ref_exists "${deepep_ref}"; then
@@ -2502,15 +2464,6 @@ resolve_ci_base_dependency_targets() {
         fi
     fi
 
-    # DeepEP inherits from ROCShmem. If ROCShmem is being seeded, seed DeepEP too
-    # so the pair stays consistent for future ci_base rebuilds.
-    if printf '%s\n' "${DEPENDENCY_CACHE_TARGETS[@]}" | grep -qx "rocshmem-rocm-ci" \
-        && ! printf '%s\n' "${DEPENDENCY_CACHE_TARGETS[@]}" | grep -qx "deepep-rocm-ci" \
-        && [[ -n "${DEEPEP_BRANCH:-}" ]]; then
-        echo "ROCShmem cache is missing; also seeding DeepEP cache"
-        add_dependency_cache_target "deepep-rocm-ci"
-    fi
-
     BAKE_TARGETS=("ci-base-rocm-ci")
     if [[ ${#DEPENDENCY_CACHE_TARGETS[@]} -eq 0 ]]; then
         echo "All dependency caches exist; building ci_base without dependency cache exports"

.buildkite/scripts/cleanup-nightly-builds.sh

@@ -5,14 +5,17 @@ set -ex
 # Clean up old nightly builds from DockerHub, keeping only the last 14 builds
 # This script uses DockerHub API to list and delete old tags with specified prefix
 # Tags starting with "nightly-dev" are always excluded from deletion
-# Usage: cleanup-nightly-builds.sh [TAG_PREFIX] [REPO]
+# Usage: cleanup-nightly-builds.sh [TAG_PREFIX] [REPO] [EXCLUDE_PREFIX]
 # Example: cleanup-nightly-builds.sh "nightly-"
 # Example: cleanup-nightly-builds.sh "cu130-nightly-"
 # Example: cleanup-nightly-builds.sh "nightly-" "vllm/vllm-openai-rocm"
+# Example: cleanup-nightly-builds.sh "nightly-" "vllm/vllm-openai-rocm" "nightly-rocm"
+#   (EXCLUDE_PREFIX skips tags owned by another variant sharing TAG_PREFIX)
 
 # Get tag prefix and repo from arguments
 TAG_PREFIX="${1:-nightly-}"
 REPO="${2:-vllm/vllm-openai}"
+EXCLUDE_PREFIX="${3:-}"
 
 echo "Cleaning up tags with prefix: $TAG_PREFIX in repository: $REPO"
 
@@ -57,7 +60,7 @@ get_all_tags() {
         
         # Get both last_updated timestamp and tag name, separated by |
         # Exclude mutable architecture tags and nightly-dev tags from cleanup
-        local tags=$(echo "$response" | jq -r --arg prefix "$TAG_PREFIX" '.results[] | select(.name | startswith($prefix)) | select(.name != ($prefix + "x86_64") and .name != ($prefix + "aarch64")) | select(.name | startswith("nightly-dev") | not) | "\(.last_updated)|\(.name)"')
+        local tags=$(echo "$response" | jq -r --arg prefix "$TAG_PREFIX" --arg exclude "$EXCLUDE_PREFIX" '.results[] | select(.name | startswith($prefix)) | select($exclude == "" or (.name | startswith($exclude) | not)) | select(.name != ($prefix + "x86_64") and .name != ($prefix + "aarch64")) | select(.name | startswith("nightly-dev") | not) | "\(.last_updated)|\(.name)"')
         
         if [ -z "$tags" ]; then
             break

.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/generate-nightly-index.py

@@ -12,7 +12,7 @@ from dataclasses import asdict, dataclass
 from datetime import datetime
 from pathlib import Path
 from typing import Any
-from urllib.parse import quote
+from urllib.parse import quote, unquote, urlparse
 
 import regex as re
 
@@ -51,6 +51,8 @@ class WheelFileInfo:
     platform_tag: str
     variant: str | None
     filename: str
+    # Absolute URL for files hosted elsewhere (see --external-links)
+    url: str | None = None
 
 
 def parse_from_filename(file: str) -> WheelFileInfo:
@@ -110,6 +112,29 @@ def parse_from_filename(file: str) -> WheelFileInfo:
     )
 
 
+def parse_external_url(url: str) -> WheelFileInfo:
+    """Parse an absolute wheel or sdist URL hosted on another index."""
+    filename = unquote(urlparse(url).path.rsplit("/", 1)[-1])
+    if filename.endswith(".whl"):
+        info = parse_from_filename(filename)
+    else:
+        match = re.match(r"^(?P<name>.+)-(?P<version>[^-]+)\.(tar\.gz|zip)$", filename)
+        if not match:
+            raise ValueError(f"Unsupported external file: {filename}")
+        info = WheelFileInfo(
+            package_name=match.group("name"),
+            version=match.group("version"),
+            build_tag=None,
+            python_tag="source",
+            abi_tag="none",
+            platform_tag="any",
+            variant=None,
+            filename=filename,
+        )
+    info.url = url
+    return info
+
+
 def generate_project_list(subdir_names: list[str], comment: str = "") -> str:
     """Generate project list HTML content linking to each project & variant subdirectory."""
     href_tags = []
@@ -129,14 +154,19 @@ def generate_package_index_and_metadata(
     href_tags = []
     metadata = []
     for file in sorted(wheel_files, key=lambda x: x.filename):
-        relative_path = (
-            wheel_base_dir.relative_to(index_base_dir, walk_up=True) / file.filename
-        )
-        # handle with '+' in URL, and avoid double-encoding '/' and already-encoded '%2B'
-        # NOTE: this is AWS S3 specific behavior!
-        file_path_quoted = quote(relative_path.as_posix(), safe=":%/")
+        if file.url:
+            file_path_quoted = file.url
+        else:
+            relative_path = (
+                wheel_base_dir.relative_to(index_base_dir, walk_up=True) / file.filename
+            )
+            # handle with '+' in URL, and avoid double-encoding '/' and already-encoded '%2B'
+            # NOTE: this is AWS S3 specific behavior!
+            file_path_quoted = quote(relative_path.as_posix(), safe=":%/")
         href_tags.append(f'    <a href="{file_path_quoted}">{file.filename}</a><br/>')
         file_meta = asdict(file)
+        if file_meta["url"] is None:
+            del file_meta["url"]
         file_meta["path"] = file_path_quoted
         metadata.append(file_meta)
     index_str = INDEX_HTML_TEMPLATE.format(items="\n".join(href_tags), comment=comment)
@@ -151,6 +181,7 @@ def generate_index_and_metadata(
     default_variant: str | None = None,
     alias_to_default: str | None = None,
     comment: str = "",
+    external_urls: list[str] | None = None,
 ):
     """Generate index for all wheel files.
 
@@ -160,6 +191,8 @@ def generate_index_and_metadata(
         index_base_dir (Path): Base directory to store index files.
         default_variant (str | None): The default variant name, if any.
         alias_to_default (str | None): Alias variant name for the default variant, if any.
+        external_urls (list[str] | None): Absolute URLs of files hosted on another index,
+            listed in the generated indices without being copied.
         comment (str | None): Optional comment to include in the generated HTML files.
 
     First, parse all wheel files to extract metadata.
@@ -214,6 +247,7 @@ def generate_index_and_metadata(
 
     """
     parsed_files = [parse_from_filename(f) for f in whl_files]
+    parsed_files += [parse_external_url(u) for u in external_urls or []]
 
     if not parsed_files:
         print("No wheel files found, skipping index generation.")
@@ -232,10 +266,14 @@ def generate_index_and_metadata(
             print(f"Detected ROCm variant from vllm: {rocm_variant}")
             break
 
-    # Apply ROCm variant to all wheels without a variant
+    # Apply ROCm variant to all wheels without a variant. Dependencies with their
+    # own ROCm local version (e.g. TheRock's torch-2.12.0+rocm10.0.0) belong to
+    # the same variant rather than one of their own.
     if rocm_variant:
         for file in parsed_files:
-            if file.variant is None:
+            if file.variant is None or (
+                file.variant.startswith("rocm") and file.variant != rocm_variant
+            ):
                 file.variant = rocm_variant
                 print(f"Inherited variant '{rocm_variant}' for {file.filename}")
 
@@ -366,6 +404,13 @@ if __name__ == "__main__":
         default=None,
         help="Alias variant name for the default variant",
     )
+    parser.add_argument(
+        "--external-links",
+        type=str,
+        default=None,
+        help="Optional file with one absolute URL per line for packages hosted on "
+        "another index (listed alongside, but not copied into, this index)",
+    )
     parser.add_argument(
         "--comment",
         type=str,
@@ -463,5 +508,14 @@ if __name__ == "__main__":
         default_variant=None,
         alias_to_default=args.alias_to_default,
         comment=args.comment.strip(),
+        external_urls=(
+            [
+                line.strip()
+                for line in Path(args.external_links).read_text().splitlines()
+                if line.strip() and not line.startswith("#")
+            ]
+            if args.external_links
+            else None
+        ),
     )
     print(f"Successfully generated index and metadata in {output_dir}")

.buildkite/scripts/publish-release-images.sh

@@ -11,9 +11,9 @@ set -euo pipefail
 TARGET="${1:-all}"
 case "${TARGET}" in
   cuda-13-0 | cuda-12-9 | cuda-13-0-ubuntu-24-04 | \
-    cuda-12-9-ubuntu-24-04 | rocm | xpu | cpu | all) ;;
+    cuda-12-9-ubuntu-24-04 | rocm | rocm72 | xpu | cpu | all) ;;
   *)
-    echo "Usage: $0 {cuda-13-0|cuda-12-9|cuda-13-0-ubuntu-24-04|cuda-12-9-ubuntu-24-04|rocm|xpu|cpu|all}"
+    echo "Usage: $0 {cuda-13-0|cuda-12-9|cuda-13-0-ubuntu-24-04|cuda-12-9-ubuntu-24-04|rocm|rocm72|xpu|cpu|all}"
     exit 2
     ;;
 esac
@@ -143,23 +143,35 @@ if target_enabled cuda-12-9-ubuntu-24-04; then
 fi
 
 # ---- ROCm ----
+# Each ROCm stack is identified by its base Dockerfile (see rocm/stack.sh). The
+# default stack also owns the plain :latest / :v<ver> tags.
+
+publish_rocm_stack() {
+  local base_dockerfile="$1" default="$2" suffix
+  (
+    # shellcheck source=.buildkite/scripts/rocm/stack.sh
+    source .buildkite/scripts/rocm/stack.sh "$base_dockerfile"
+    local suffixes=("-${ROCM_STACK_VARIANT}")
+    [[ "$default" == "1" ]] && suffixes=("" "-${ROCM_STACK_VARIANT}")
+    docker pull "$ROCM_STACK_ECR_IMAGE"
+    docker pull "$ROCM_STACK_ECR_BASE"
+    for suffix in "${suffixes[@]}"; do
+      for tag in "latest${suffix}" "v${RELEASE_VERSION}${suffix}"; do
+        docker tag "$ROCM_STACK_ECR_IMAGE" "vllm/vllm-openai-rocm:${tag}"
+        docker tag "$ROCM_STACK_ECR_BASE" "vllm/vllm-openai-rocm:${tag}-base"
+        docker push "vllm/vllm-openai-rocm:${tag}"
+        docker push "vllm/vllm-openai-rocm:${tag}-base"
+      done
+    done
+  )
+}
 
 if target_enabled rocm; then
-  ROCM_BASE_CACHE_KEY=$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
-  echo "ROCm base cache key: ${ROCM_BASE_CACHE_KEY}"
+  publish_rocm_stack docker/Dockerfile.rocm_base 1
+fi
 
-  docker pull "public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm"
-  docker pull "public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base"
-
-  docker tag "public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm" vllm/vllm-openai-rocm:latest
-  docker tag "public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm" "vllm/vllm-openai-rocm:v${RELEASE_VERSION}"
-  docker push vllm/vllm-openai-rocm:latest
-  docker push "vllm/vllm-openai-rocm:v${RELEASE_VERSION}"
-
-  docker tag "public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base" vllm/vllm-openai-rocm:latest-base
-  docker tag "public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base" "vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base"
-  docker push vllm/vllm-openai-rocm:latest-base
-  docker push "vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base"
+if target_enabled rocm72; then
+  publish_rocm_stack docker/Dockerfile.rocm_72_base 0
 fi
 
 # ---- XPU ----

.buildkite/scripts/push-nightly-builds-rocm.sh

@@ -2,42 +2,36 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 #
-# Push a ROCm-family nightly base image and nightly image from ECR to Docker Hub
-# under vllm/vllm-openai-rocm, as base-nightly[-<variant>], nightly[-<variant>]
-# and their -<commit> forms.
-# Run when NIGHTLY=1 after the matching build-*-release-image step has pushed to ECR.
+# Push a ROCm stack's nightly base image and nightly image from ECR to Docker Hub
+# under vllm/vllm-openai-rocm, as base-nightly-<variant>, nightly-<variant> and their
+# -<commit> forms, then prune old nightly tags. With --default, also push the plain
+# base-nightly and nightly tags for the same images.
+# Run when NIGHTLY=1 after the stack's release image has been pushed to ECR.
 #
-# Usage: push-nightly-builds-rocm.sh [ECR_TAG_SUFFIX TAG_VARIANT]
-#   Pass both arguments or neither. With none, the default ROCm apt build is
-#   published as :nightly. The TheRock build passes "rock" (the source ECR tag
-#   suffix) and "rocm714" (the Docker Hub tag flavor), giving :nightly-rocm714.
+# Usage: push-nightly-builds-rocm.sh <base-dockerfile> [--default]
+#   e.g. "docker/Dockerfile.rocm_base --default" (ROCm 10.0, also :nightly)
+#   and "docker/Dockerfile.rocm_72_base" (ROCm 7.2, :nightly-rocm72).
 #
 # Local testing (no push to Docker Hub):
-#   BASE_ECR_IMAGE=<full-base-image-reference> \
-#     BUILDKITE_COMMIT=<commit-with-image-in-ecr> DRY_RUN=1 \
-#     bash .buildkite/scripts/push-nightly-builds-rocm.sh
+#   BUILDKITE_COMMIT=<commit-with-image-in-ecr> DRY_RUN=1 \
+#     bash .buildkite/scripts/push-nightly-builds-rocm.sh docker/Dockerfile.rocm_base --default
 # Requires: AWS CLI configured (for ECR public login), Docker. For full run: Docker Hub login.
 
 set -euxo pipefail
 
 usage() {
-  echo "Usage: $0 [ECR_TAG_SUFFIX TAG_VARIANT]" >&2
+  echo "Usage: $0 <base-dockerfile> [--default]" >&2
 }
 
-DOCKERHUB_REPO="vllm/vllm-openai-rocm"
-
+PUSH_DEFAULT=0
 case "$#" in
-  0)
-    ECR_TAG_SUFFIX="rocm"
-    TAG_VARIANT=""
-    ;;
+  1) ;;
   2)
-    if [[ -z "$1" || -z "$2" ]]; then
+    if [[ "$2" != "--default" ]]; then
       usage
       exit 2
     fi
-    ECR_TAG_SUFFIX="$1"
-    TAG_VARIANT="$2"
+    PUSH_DEFAULT=1
     ;;
   *)
     usage
@@ -49,61 +43,44 @@ esac
 BUILDKITE_COMMIT="${BUILDKITE_COMMIT:?Set BUILDKITE_COMMIT to the commit SHA that has the image in ECR (e.g. from a previous release pipeline run)}"
 DRY_RUN="${DRY_RUN:-0}"
 
-ECR_REPO="public.ecr.aws/q9t5s3a7/vllm-release-repo"
+# shellcheck source=.buildkite/scripts/rocm/stack.sh
+source "$(dirname "${BASH_SOURCE[0]}")/rocm/stack.sh" "$1"
+DOCKERHUB_REPO="vllm/vllm-openai-rocm"
+VARIANT="$ROCM_STACK_VARIANT"
 
-# Get the base image ECR tag (set by the build-*-release-image pipeline step)
-BASE_METADATA_KEY="${ECR_TAG_SUFFIX}-base-ecr-tag"
-if [[ -n "${BASE_ECR_IMAGE:-}" ]]; then
-  if [[ "$DRY_RUN" != "1" ]]; then
-    echo "ERROR: BASE_ECR_IMAGE may only be used with DRY_RUN=1" >&2
-    exit 1
-  fi
-  BASE_ORIG_TAG="$BASE_ECR_IMAGE"
-elif ! BASE_ORIG_TAG="$(buildkite-agent meta-data get "$BASE_METADATA_KEY")"; then
-  echo "ERROR: Failed to read required Buildkite metadata '$BASE_METADATA_KEY'" >&2
-  exit 1
-fi
-if [[ -z "$BASE_ORIG_TAG" ]]; then
-  echo "ERROR: Required Buildkite metadata '$BASE_METADATA_KEY' is empty" >&2
-  echo "Set BASE_ECR_IMAGE to the full ECR base image reference for local testing" >&2
-  exit 1
+BASE_TAGS=("base-nightly-${VARIANT}" "base-nightly-${VARIANT}-${BUILDKITE_COMMIT}")
+TAGS=("nightly-${VARIANT}" "nightly-${VARIANT}-${BUILDKITE_COMMIT}")
+if [[ "$PUSH_DEFAULT" == "1" ]]; then
+  BASE_TAGS+=("base-nightly" "base-nightly-${BUILDKITE_COMMIT}")
+  TAGS+=("nightly" "nightly-${BUILDKITE_COMMIT}")
 fi
 
-ORIG_TAG="${BUILDKITE_COMMIT}-${ECR_TAG_SUFFIX}"
-VARIANT_SUFFIX=""
-[[ -n "$TAG_VARIANT" ]] && VARIANT_SUFFIX="-$TAG_VARIANT"
-BASE_TAG_NAME="base-nightly${VARIANT_SUFFIX}"
-TAG_NAME="nightly${VARIANT_SUFFIX}"
-BASE_TAG_NAME_COMMIT="base-nightly${VARIANT_SUFFIX}-${BUILDKITE_COMMIT}"
-TAG_NAME_COMMIT="nightly${VARIANT_SUFFIX}-${BUILDKITE_COMMIT}"
+echo "Pushing base image $ROCM_STACK_ECR_BASE as: ${BASE_TAGS[*]}"
+echo "Pushing release image $ROCM_STACK_ECR_IMAGE as: ${TAGS[*]}"
 
-echo "Pushing base image from ECR: $BASE_ORIG_TAG"
-echo "Pushing release image from ECR tag: $ORIG_TAG to Docker Hub as $DOCKERHUB_REPO:$TAG_NAME and $DOCKERHUB_REPO:$TAG_NAME_COMMIT"
-[[ "$DRY_RUN" == "1" ]] && echo "[DRY_RUN] Skipping push to Docker Hub"
-
-# Login to ECR and pull the images built by build-*-release-image
+# Login to ECR and pull the stack's images
 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-docker pull "$BASE_ORIG_TAG"
-docker pull "${ECR_REPO}:${ORIG_TAG}"
+docker pull "$ROCM_STACK_ECR_BASE"
+docker pull "$ROCM_STACK_ECR_IMAGE"
 
-# Tag for Docker Hub (base-nightly and base-nightly-<commit>, nightly and nightly-<commit>)
-docker tag "$BASE_ORIG_TAG" "$DOCKERHUB_REPO":"$BASE_TAG_NAME"
-docker tag "$BASE_ORIG_TAG" "$DOCKERHUB_REPO":"$BASE_TAG_NAME_COMMIT"
-docker tag "${ECR_REPO}:${ORIG_TAG}" "$DOCKERHUB_REPO":"$TAG_NAME"
-docker tag "${ECR_REPO}:${ORIG_TAG}" "$DOCKERHUB_REPO":"$TAG_NAME_COMMIT"
+for tag in "${BASE_TAGS[@]}"; do docker tag "$ROCM_STACK_ECR_BASE" "$DOCKERHUB_REPO:$tag"; done
+for tag in "${TAGS[@]}"; do docker tag "$ROCM_STACK_ECR_IMAGE" "$DOCKERHUB_REPO:$tag"; done
 
 if [[ "$DRY_RUN" == "1" ]]; then
-  echo "[DRY_RUN] Would push $DOCKERHUB_REPO:$BASE_TAG_NAME and $DOCKERHUB_REPO:$BASE_TAG_NAME_COMMIT"
-  echo "[DRY_RUN] Would push $DOCKERHUB_REPO:$TAG_NAME and $DOCKERHUB_REPO:$TAG_NAME_COMMIT"
   echo "[DRY_RUN] Local tags created. Exiting without push."
   exit 0
 fi
 
 # Push to Docker Hub (docker-login plugin runs before this step in CI)
-docker push "$DOCKERHUB_REPO":"$BASE_TAG_NAME"
-docker push "$DOCKERHUB_REPO":"$BASE_TAG_NAME_COMMIT"
-docker push "$DOCKERHUB_REPO":"$TAG_NAME"
-docker push "$DOCKERHUB_REPO":"$TAG_NAME_COMMIT"
+for tag in "${BASE_TAGS[@]}" "${TAGS[@]}"; do docker push "$DOCKERHUB_REPO:$tag"; done
+echo "Pushed $DOCKERHUB_REPO: ${BASE_TAGS[*]} ${TAGS[*]}"
 
-echo "Pushed $DOCKERHUB_REPO:$BASE_TAG_NAME and $DOCKERHUB_REPO:$BASE_TAG_NAME_COMMIT"
-echo "Pushed $DOCKERHUB_REPO:$TAG_NAME and $DOCKERHUB_REPO:$TAG_NAME_COMMIT"
+# Keep only the last 14 builds of each tag family. The plain nightly- families
+# exclude the variant ones sharing their prefix; commit SHAs never start with "rocm".
+CLEANUP=.buildkite/scripts/cleanup-nightly-builds.sh
+bash "$CLEANUP" "nightly-${VARIANT}-" "$DOCKERHUB_REPO"
+bash "$CLEANUP" "base-nightly-${VARIANT}-" "$DOCKERHUB_REPO"
+if [[ "$PUSH_DEFAULT" == "1" ]]; then
+  bash "$CLEANUP" "nightly-" "$DOCKERHUB_REPO" "nightly-rocm"
+  bash "$CLEANUP" "base-nightly-" "$DOCKERHUB_REPO" "base-nightly-rocm"
+fi

.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/build-config.sh

@@ -1,8 +1,8 @@
 #!/usr/bin/env bash
 # Shared Dockerfile selection for the normal AMD CI build chain.
 # CI_ROCM_DOCKERFILE_BASE and CI_ROCM_DOCKERFILE accept repository-relative paths.
-# For Rock, set CI_ROCM_DOCKERFILE_BASE=docker/Dockerfile.rock_base and
-# CI_ROCM_DOCKERFILE=docker/Dockerfile.rock.
+# For the legacy ROCm 7.2 stack, set CI_ROCM_DOCKERFILE_BASE=docker/Dockerfile.rocm_72_base
+# and CI_ROCM_DOCKERFILE=docker/Dockerfile.rocm_72.
 # Custom recipes must preserve the CI targets and artifact layout. Additional
 # COPY inputs also need the corresponding *_CONTENT_FILES overrides.
 

.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/scripts/rocm/stack.sh

@@ -0,0 +1,78 @@
+#!/usr/bin/env bash
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+#
+# Derive everything the release pipeline needs about a ROCm stack from its base
+# Dockerfile, so the stack's image is the single input to every ROCm script.
+#
+# Usage: source .buildkite/scripts/rocm/stack.sh <base-dockerfile>
+#   e.g. docker/Dockerfile.rocm_base (ROCm 10.0) or docker/Dockerfile.rocm_72_base
+#
+# Exports:
+#   ROCM_STACK_BASE_DOCKERFILE  the base Dockerfile
+#   ROCM_STACK_DOCKERFILE       the matching vLLM Dockerfile (base name minus _base)
+#   ROCM_STACK_THEROCK          1 if the ROCm SDK comes from TheRock wheels
+#   ROCM_STACK_VERSION          ROCm version, e.g. 10.0.0 or 7.2.3
+#   ROCM_STACK_VARIANT          wheel variant and Docker Hub tag flavor from major.minor,
+#                               same rule as setup.py (rocm100, rocm72)
+#   ROCM_STACK_IMAGE_KEY        hash of the base Dockerfile (base image cache key)
+#   ROCM_STACK_WHEEL_KEY        hash of the base Dockerfile plus files it COPYs from
+#                               the repo (base wheel cache key)
+#   ROCM_STACK_ECR_BASE         ECR tag of the cached base image
+#   ROCM_STACK_ECR_IMAGE        ECR tag of this commit's release image
+#   ROCM_STACK_BASE_WHEELS_DIR  local dir for base wheels (also the artifact path)
+#   ROCM_STACK_VLLM_WHEEL_DIR   local dir for the vLLM wheel (also the artifact path)
+
+rocm_stack_init() {
+    local base="$1"
+    if [[ ! -f "$base" ]]; then
+        echo "ROCm base Dockerfile not found: $base" >&2
+        return 2
+    fi
+    ROCM_STACK_BASE_DOCKERFILE="$base"
+    ROCM_STACK_DOCKERFILE="${base%_base}"
+
+    local sdk
+    sdk=$(sed -nE 's/^ARG ROCM_SDK_VERSION=([0-9.]+).*/\1/p' "$base" | head -1)
+    if [[ -n "$sdk" ]]; then
+        ROCM_STACK_THEROCK=1
+        ROCM_STACK_VERSION="$sdk"
+    else
+        ROCM_STACK_THEROCK=0
+        # BASE_IMAGE format: rocm/dev-ubuntu-22.04:7.2.3-complete -> 7.2.3
+        ROCM_STACK_VERSION=$(sed -nE 's/^ARG BASE_IMAGE=.*:([0-9]+(\.[0-9]+)+).*/\1/p' "$base" | head -1)
+    fi
+    if [[ -z "$ROCM_STACK_VERSION" ]]; then
+        echo "Could not determine the ROCm version from $base" >&2
+        return 2
+    fi
+
+    local major minor
+    IFS=. read -r major minor _ <<< "$ROCM_STACK_VERSION"
+    ROCM_STACK_VARIANT="rocm${major}${minor}"
+
+    ROCM_STACK_IMAGE_KEY=$(sha256sum "$base" | cut -c1-16)
+    local inputs
+    inputs=$(grep -E '^COPY ' "$base" | grep -v -- '--from' | awk '{print $2}' || true)
+    if [[ -z "$inputs" ]]; then
+        ROCM_STACK_WHEEL_KEY="$ROCM_STACK_IMAGE_KEY"
+    else
+        # shellcheck disable=SC2086
+        ROCM_STACK_WHEEL_KEY=$(sha256sum "$base" $inputs | sha256sum | cut -c1-16)
+    fi
+
+    local ecr="public.ecr.aws/q9t5s3a7/vllm-release-repo"
+    ROCM_STACK_ECR_BASE="${ecr}:${ROCM_STACK_IMAGE_KEY}-${ROCM_STACK_VARIANT}-base"
+    ROCM_STACK_ECR_IMAGE="${ecr}:${BUILDKITE_COMMIT:-unknown}-${ROCM_STACK_VARIANT}"
+    ROCM_STACK_BASE_WHEELS_DIR="artifacts/${ROCM_STACK_VARIANT}-base-wheels"
+    ROCM_STACK_VLLM_WHEEL_DIR="artifacts/${ROCM_STACK_VARIANT}-vllm-wheel"
+
+    export ROCM_STACK_BASE_DOCKERFILE ROCM_STACK_DOCKERFILE ROCM_STACK_THEROCK \
+        ROCM_STACK_VERSION ROCM_STACK_VARIANT ROCM_STACK_IMAGE_KEY \
+        ROCM_STACK_WHEEL_KEY ROCM_STACK_ECR_BASE ROCM_STACK_ECR_IMAGE \
+        ROCM_STACK_BASE_WHEELS_DIR ROCM_STACK_VLLM_WHEEL_DIR
+}
+
+if [[ $# -gt 0 ]]; then
+    rocm_stack_init "$1"
+fi

.buildkite/scripts/upload-rocm-wheels.sh

@@ -4,21 +4,31 @@
 #
 # Upload ROCm wheels to S3 with proper index generation
 #
+# Usage: upload-rocm-wheels.sh <base-dockerfile>
+#   base-dockerfile  the ROCm stack, e.g. docker/Dockerfile.rocm_base (see rocm/stack.sh)
+#
 # Required environment variables:
 #   AWS_ACCESS_KEY_ID, AWS_SECRET_ACCESS_KEY (or IAM role)
 #   S3_BUCKET (default: vllm-wheels)
 #
 # S3 path structure:
-#   s3://vllm-wheels/rocm/{commit}/     - All wheels for this commit
-#   s3://vllm-wheels/rocm/nightly/      - Index pointing to latest nightly
-#   s3://vllm-wheels/rocm/{version}/    - Index for release versions
+#   s3://vllm-wheels/rocm/{commit}/<variant>-wheels/ - This stack's wheels for this commit
+#   s3://vllm-wheels/rocm/{commit}/<variant>/   - Index per ROCm variant (e.g. rocm100)
+#   s3://vllm-wheels/rocm/nightly/<variant>/    - Index pointing to latest nightly
+#   s3://vllm-wheels/rocm/{version}/<variant>/  - Index for release versions
+# The top-level index.html at each prefix lists every variant present, since
+# several ROCm stacks publish into the same prefixes.
 
 set -ex
 
 # ======== Configuration ========
 BUCKET="${S3_BUCKET:-vllm-wheels}"
+# shellcheck source=.buildkite/scripts/rocm/stack.sh
+source "$(dirname "${BASH_SOURCE[0]}")/rocm/stack.sh" "${1:?Usage: $0 <base-dockerfile>}"
 ROCM_SUBPATH="rocm/${BUILDKITE_COMMIT}"
-S3_COMMIT_PREFIX="s3://$BUCKET/$ROCM_SUBPATH/"
+WHEEL_DIRS="$ROCM_STACK_BASE_WHEELS_DIR $ROCM_STACK_VLLM_WHEEL_DIR"
+WHEEL_SUBPATH="$ROCM_SUBPATH/${ROCM_STACK_VARIANT}-wheels"
+S3_WHEEL_PREFIX="s3://$BUCKET/$WHEEL_SUBPATH/"
 INDICES_OUTPUT_DIR="rocm-indices"
 
 echo "========================================"
@@ -49,9 +59,11 @@ source .buildkite/scripts/lib/manylinux.sh
 # ======== Part 1: Collect and prepare wheels ========
 
 # Collect all wheels
+rm -rf all-rocm-wheels "$INDICES_OUTPUT_DIR"
 mkdir -p all-rocm-wheels
-cp artifacts/rocm-base-wheels/*.whl all-rocm-wheels/ 2>/dev/null || true
-cp artifacts/rocm-vllm-wheel/*.whl all-rocm-wheels/ 2>/dev/null || true
+for dir in $WHEEL_DIRS; do
+    cp "$dir"/*.whl all-rocm-wheels/ 2>/dev/null || true
+done
 
 WHEEL_COUNT=$(find all-rocm-wheels -maxdepth 1 -name '*.whl' 2>/dev/null | wc -l)
 echo "Total wheels to upload: $WHEEL_COUNT"
@@ -84,9 +96,9 @@ ls -lh all-rocm-wheels/
 # ======== Part 2: Upload wheels to S3 ========
 
 echo ""
-echo "Uploading wheels to $S3_COMMIT_PREFIX"
+echo "Uploading wheels to $S3_WHEEL_PREFIX"
 for wheel in all-rocm-wheels/*.whl; do
-    aws s3 cp "$wheel" "$S3_COMMIT_PREFIX"
+    aws s3 cp "$wheel" "$S3_WHEEL_PREFIX"
 done
 
 # ======== Part 3: Generate and upload indices ========
@@ -95,28 +107,63 @@ done
 echo ""
 echo "Generating indices..."
 obj_json="rocm-objects.json"
-aws s3api list-objects-v2 --bucket "$BUCKET" --prefix "$ROCM_SUBPATH/" --delimiter / --output json > "$obj_json"
+aws s3api list-objects-v2 --bucket "$BUCKET" --prefix "$WHEEL_SUBPATH/" --delimiter / --output json > "$obj_json"
 
 mkdir -p "$INDICES_OUTPUT_DIR"
 
 # Use the existing generate-nightly-index.py
 # HACK: Replace regex module with stdlib re (same as CUDA script)
 sed -i 's/import regex as re/import re/g' .buildkite/scripts/generate-nightly-index.py
+sed -i 's/import regex as re/import re/g' tools/vllm-rocm/therock_wheels.py
+sed -i 's/import pybase64 as base64/import base64/g' tools/vllm-rocm/therock_wheels.py
 
+INDEX_ARGS=(--version "$ROCM_SUBPATH" --wheel-dir "${WHEEL_SUBPATH#rocm/}")
+# TheRock's SDK, device kernels and torchvision/torchaudio are linked, not hosted
+if [[ "$ROCM_STACK_THEROCK" == "1" ]]; then
+    $PYTHON tools/vllm-rocm/therock_wheels.py external-links \
+        --dockerfile "$ROCM_STACK_BASE_DOCKERFILE" \
+        --torch-wheel "$(ls "$ROCM_STACK_BASE_WHEELS_DIR"/torch-*.whl)" > therock-external-links.txt
+    cat therock-external-links.txt
+    INDEX_ARGS+=(--external-links therock-external-links.txt)
+fi
 $PYTHON .buildkite/scripts/generate-nightly-index.py \
-    --version "$ROCM_SUBPATH" \
+    "${INDEX_ARGS[@]}" \
     --current-objects "$obj_json" \
     --output-dir "$INDICES_OUTPUT_DIR" \
     --comment "ROCm commit $BUILDKITE_COMMIT"
 
-# Upload indices to commit directory
-echo "Uploading indices to $S3_COMMIT_PREFIX"
-aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
+# Upload this build's variant indices under an S3 prefix (e.g. rocm/nightly), then
+# rebuild that prefix's top-level project list from every variant now present.
+publish_indices() {
+    local prefix="$1"
+    local listing variants merged
+    echo "Uploading indices to s3://$BUCKET/$prefix/"
+    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/$prefix/" --exclude "index.html"
+    listing=$(aws s3api list-objects-v2 --bucket "$BUCKET" --prefix "$prefix/" \
+        --delimiter / --query 'CommonPrefixes[].Prefix' --output text)
+    variants=$(tr '\t' '\n' <<< "$listing" | sed -n "s|^$prefix/\(rocm[0-9]*\)/$|\1|p" | sort -u)
+    merged=$(mktemp)
+    {
+        echo '<!DOCTYPE html>'
+        echo '<html>'
+        echo '  <meta name="pypi:repository-version" content="1.0">'
+        echo '  <body>'
+        for variant in $variants; do
+            echo "    <a href=\"$variant/\">$variant/</a><br/>"
+        done
+        echo '  </body>'
+        echo '</html>'
+    } > "$merged"
+    echo "Variants under $prefix/: $(echo "$variants" | tr '\n' ' ')"
+    aws s3 cp "$merged" "s3://$BUCKET/$prefix/index.html"
+    rm -f "$merged"
+}
+
+publish_indices "$ROCM_SUBPATH"
 
 # Only scheduled nightly builds should update the moving nightly index.
 if [[ "${NIGHTLY:-0}" == "1" ]]; then
-    echo "Updating rocm/nightly/ index..."
-    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/rocm/nightly/"
+    publish_indices "rocm/nightly"
 fi
 
 # Extract version from vLLM wheel and update version-specific index
@@ -129,8 +176,7 @@ if [ -n "$VLLM_WHEEL" ]; then
     echo "Pure version: $PURE_VERSION"
 
     if [[ "$VERSION" != *"dev"* ]]; then
-        echo "Updating rocm/$PURE_VERSION/ index..."
-        aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/rocm/$PURE_VERSION/"
+        publish_indices "rocm/$PURE_VERSION"
     fi
 fi
 
@@ -142,14 +188,14 @@ echo "ROCm Wheel Upload Complete!"
 echo "========================================"
 echo ""
 echo "Wheels available at:"
-echo "  s3://$BUCKET/$ROCM_SUBPATH/"
+echo "  s3://$BUCKET/$WHEEL_SUBPATH/"
 echo ""
 echo "Install command (by commit):"
-echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/$ROCM_SUBPATH/"
+echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/$ROCM_SUBPATH/$ROCM_STACK_VARIANT/"
 echo ""
 if [[ "${NIGHTLY:-0}" == "1" ]]; then
     echo "Install command (nightly):"
-    echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/rocm/nightly/"
+    echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/rocm/nightly/$ROCM_STACK_VARIANT/"
 fi
 echo ""
 echo "Wheel count: $WHEEL_COUNT"

.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
-  parallelism: 2
+    - pytest -v -s v1/attention
   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
-  parallelism: 2
+    - pytest -v -s v1/attention
   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
-      - 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
+      - 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 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

@@ -14,6 +14,7 @@ steps:
   - tests/distributed
   commands:
   - pytest -v -s distributed/test_comm_ops.py
+  - pytest -v -s distributed/test_cpu_all_reduce.py
   - pytest -v -s distributed/test_shm_broadcast.py
   - pytest -v -s distributed/test_shm_buffer.py
   - pytest -v -s distributed/test_shm_storage.py
@@ -77,6 +78,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
@@ -108,6 +110,7 @@ steps:
     amd:
       label: ":amd: (MI355) Distributed DP Basic"
       dind: false
+      automatic_shard: true
       device: mi355_2
       timeout_in_minutes: 45
       depends_on:
@@ -300,6 +303,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
@@ -329,6 +333,7 @@ steps:
     amd:
       label: ':amd: (MI355) Distributed DP Extended'
       dind: false
+      automatic_shard: true
       device: mi355_4
       timeout_in_minutes: 60
       working_dir: /vllm-workspace/tests
@@ -508,9 +513,7 @@ steps:
       - examples/features/data_parallel/data_parallel_offline.py
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.py
-      - csrc/custom_quickreduce.cu
-      - csrc/ops.h
-      - csrc/torch_bindings.cpp
+      - csrc/rocm/quickreduce/
       - vllm/model_executor/layers/
       - vllm/entrypoints/llm.py
       - vllm/config/parallel.py
@@ -522,13 +525,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 +546,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
@@ -632,10 +640,11 @@ steps:
   - tests/distributed/
   - examples/features/data_parallel/data_parallel_offline.py
   commands:
-    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_multi_node_assignment.py && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"
+    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"
 
 - 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 +686,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 +700,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
@@ -700,6 +709,7 @@ steps:
     amd:
       label: ':amd: (MI355) RayExecutorV2'
       dind: false
+      automatic_shard: true
       device: mi355_4
       timeout_in_minutes: 75
       working_dir: /vllm-workspace/tests
@@ -715,7 +725,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:
@@ -237,9 +238,9 @@ steps:
     - pytest -v -s v1/e2e/test_hybrid_chunked_prefill.py
   mirror:
     amd:
-      label: ":amd: (MI300) V1 E2E Hybrid Chunked Prefill"
+      label: ":amd: (MI355 DPX) V1 E2E Hybrid Chunked Prefill"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 35
       depends_on:
       - image-build-amd

.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:
@@ -40,6 +41,7 @@ steps:
     amd:
       label: ":amd: (MI355 DPX) Entrypoints Unit"
       dind: false
+      automatic_shard: true
       device: mi355_dpx
       timeout_in_minutes: 50
       depends_on:
@@ -120,6 +122,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 +144,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:
@@ -204,6 +208,7 @@ steps:
     amd:
       label: ":amd: (MI355 DPX) Entrypoints Integration (Speech to Text)"
       dind: false
+      automatic_shard: true
       device: mi355_dpx
       timeout_in_minutes: 60
       working_dir: "/vllm-workspace/tests"
@@ -213,6 +218,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:
@@ -248,6 +254,7 @@ steps:
     amd:
       label: ":amd: (MI355 DPX) Entrypoints Integration (Pooling)"
       dind: false
+      automatic_shard: true
       device: mi355_dpx
       timeout_in_minutes: 65
       working_dir: "/vllm-workspace/tests"

.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
-  parallelism: 3
+    - 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
   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/libtorch_stable/minimax_reduce_rms_kernel.h
-  - csrc/libtorch_stable/ops.h
-  - csrc/libtorch_stable/torch_bindings.cpp
+  - csrc/minimax_reduce_rms_kernel.cu
+  - csrc/minimax_reduce_rms_kernel.h
+  - csrc/ops.h
+  - 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
@@ -137,11 +138,15 @@ steps:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
   - vllm/models/deepseek_v4/common/ops/
   - vllm/models/deepseek_v4/nvidia/
+  - vllm/model_executor/layers/fused_moe/deep_gemm_mega_moe.py
+  - vllm/model_executor/layers/fused_moe/prepare_megamoe.py
   - vllm/utils/deep_gemm.py
   - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
+  - tests/kernels/moe/test_deep_gemm_mega_moe.py
   - tests/models/test_deepseek_v4_mega_moe.py
   commands:
     - pytest -v -s kernels/test_fused_deepseek_v4_*.py
+    - pytest -v -s kernels/moe/test_deep_gemm_mega_moe.py
     - pytest -v -s models/test_deepseek_v4_mega_moe.py
   mirror:
     amd:
@@ -230,8 +235,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 /
-# Kernels (B200)).
+# (test_bf16x3_router_gemm_cutedsl runs in Kernels (B200); test_ll_bf16_gemm and
+# 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 +265,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)
@@ -284,7 +290,6 @@ steps:
           kernels/test_apply_repetition_penalties.py
           kernels/test_awq_int4_to_int8.py
           kernels/test_bf16_skinny_gemm.py
-          kernels/test_bf16x3_router_gemm_rocm.py
           kernels/test_cache_kernels.py
           kernels/test_compressor_kv_cache.py
           kernels/test_cp_gather_fp8.py
@@ -299,14 +304,12 @@ steps:
           kernels/test_fused_inv_rope_fp8_quant.py
           kernels/test_fused_minimax_m3_qknorm_rope_kv_insert.py
           kernels/test_fused_quant_activation.py
-          kernels/test_gate_linear_rocm_dispatch.py
           kernels/test_gemm_rs_ar.py
           kernels/test_kpool_decode_update_batched.py
           kernels/test_minimax_m3_amd_ops.py
           kernels/test_minimax_m3_sparse_attn_fp8_scale.py
           kernels/test_onednn.py
           kernels/test_relu2_fp8_quant.py
-          kernels/test_rocm_fp32_router_gemm.py
           --deselect=tests/kernels/test_bf16_skinny_gemm.py::test_every_dsv3_routed_shape_is_instantiated
 
 - label: ":nvidia: (L4) Attention Kernels Shard %N"
@@ -316,6 +319,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 +339,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 +412,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 +430,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 +477,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 +495,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,14 +519,17 @@ 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:
     amd:
       label: ":amd: (MI250) Mamba Kernels"
       dind: false
+      automatic_shard: true
       device: mi250_1
       timeout_in_minutes: 55
       depends_on:
@@ -532,7 +545,6 @@ steps:
   num_devices: 1
   source_file_dependencies:
   - cmake/external_projects/deepgemm.cmake
-  - tools/install_deepgemm.sh
   - tools/build_deepgemm_C.py
   - tools/setup_deepgemm_pythons.sh
   - tools/check_wheel_deepgemm.py
@@ -576,11 +588,10 @@ steps:
   # optional: true
   source_file_dependencies:
   - csrc/quantization/fp4/
-  - csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu
-  - csrc/libtorch_stable/quantization/fp4/
-  - csrc/libtorch_stable/quantization/w8a8/cutlass/
+  - csrc/nvfp4_kv_cache_kernels.cu
+  - csrc/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 +605,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 +614,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/libtorch_stable/ops.h
-  - csrc/libtorch_stable/torch_bindings.cpp
+  - csrc/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu
+  - csrc/ops.h
+  - 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 +624,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 +644,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 +700,46 @@ 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
+  mirror:
+    amd:
+      label: ":amd: (MI355) Gate Linear Router Kernels"
+      dind: false
+      device: mi355_1
+      timeout_in_minutes: 20
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/model_executor/layers/fused_moe/router/gate_linear.py
+      - vllm/model_executor/layers/fused_moe/router/rocm_fp32_router_gemm.py
+      - vllm/model_executor/layers/fused_moe/router/bf16x3_router_gemm_rocm.py
+      - vllm/platforms/rocm.py
+      - tests/kernels/test_rocm_fp32_router_gemm.py
+      - tests/kernels/test_gate_linear_rocm_dispatch.py
+      - tests/kernels/test_bf16x3_router_gemm_rocm.py
+      commands:
+      - pytest -v -s tests/kernels/test_rocm_fp32_router_gemm.py tests/kernels/test_gate_linear_rocm_dispatch.py tests/kernels/test_bf16x3_router_gemm_rocm.py
+
 - label: ":nvidia: (DGX) Spark b12x Linear Kernels Nightly"
   key: b12x-linear-kernels-dgx-spark-nightly
   timeout_in_minutes: 60
@@ -726,7 +772,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 +899,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/
@@ -892,11 +938,11 @@ steps:
 
 - label: ":nvidia: (B200) FusedMoE Layer Kernels"
   key: kernels-fusedmoe-layer-test-2-b200s
-  timeout_in_minutes: 90
+  timeout_in_minutes: 45
   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 +1055,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

@@ -15,9 +15,9 @@ steps:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
   mirror:
     amd:
-      label: ":amd: (MI300) LM Eval Small Models"
+      label: ":amd: (MI355) LM Eval Small Models"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 45
       depends_on:
       - image-build-amd
@@ -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
@@ -697,18 +707,17 @@ steps:
       - export VLLM_WORKER_MULTIPROC_METHOD=spawn
       - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi355-nemotron-3-super.txt
 
-- label: ":amd: (MI300) LM Eval Large Models"
+- label: ":amd: (MI355) LM Eval Large Models"
   key: amd-lm-eval-large-models
   dind: false
-  device: mi300_8
+  device: mi355_8
   optional: true
   timeout_in_minutes: 65
   depends_on:
   - image-build-amd
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - export PYTORCH_ROCM_ARCH=gfx942 # Limit Quark compilation to save time
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-gfx950-large.txt
   source_file_dependencies:
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
@@ -1185,7 +1194,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 +1212,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 +1233,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/
@@ -1288,9 +1297,9 @@ steps:
   - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "nemotron-h-8b or gemma-4-e4b-it"
   mirror:
     amd:
-      label: ":amd: (MI300) LM Eval KV-Offload"
+      label: ":amd: (MI355 DPX) LM Eval KV-Offload"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 55
       depends_on:
       - image-build-amd
@@ -1335,6 +1344,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
@@ -1348,9 +1359,9 @@ steps:
   - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "deepseek-v4-flash"
   mirror:
     amd:
-      label: ":amd: (MI300) LM Eval KV-Offload Large"
+      label: ":amd: (MI355) LM Eval KV-Offload Large"
       dind: false
-      device: mi300_4
+      device: mi355_4
       timeout_in_minutes: 65
       depends_on:
       - image-build-amd
@@ -1415,10 +1426,10 @@ steps:
   # Qwen2-57B has 28 attention heads and cannot run at TP8.
   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4
 
-- label: ":amd: (MI300) LM Eval Large Models ROCm Harness"
+- label: ":amd: (MI355) LM Eval Large Models ROCm Harness"
   key: amd-lm-eval-large-models-rocm-harness
   dind: false
-  device: mi300_8
+  device: mi355_8
   optional: true
   timeout_in_minutes: 40
   depends_on:

.buildkite/test_areas/misc.yaml

@@ -32,6 +32,7 @@ steps:
     amd:
       label: ":amd: (MI355 DPX) V1 Sample"
       dind: false
+      automatic_shard: true
       device: mi355_dpx
       timeout_in_minutes: 70
       depends_on:
@@ -163,6 +164,59 @@ steps:
       depends_on:
       - image-build-amd
 
+# The shared KV Connectors steps above exclude vllm/distributed/kv_transfer/ from
+# their triggers, and CUDA skips MoRIIO's mori/aiter-gated tests. Run the whole
+# MoRIIO suite on ROCm whenever the connector or its tests change.
+- label: ":amd: (MI355 DPX) MoRIIO Connector Unit Tests"
+  key: amd-moriio-connector-unit-mi355
+  dind: false
+  device: mi355_dpx
+  num_devices: 1
+  timeout_in_minutes: 30
+  working_dir: /vllm-workspace/tests
+  depends_on:
+  - image-build-amd
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/moriio/
+  - vllm/distributed/kv_transfer/kv_connector/v1/base.py
+  - vllm/distributed/kv_transfer/kv_connector/v1/ssm_conv_transfer_utils.py
+  - examples/disaggregated/disaggregated_serving/moriio_toy_proxy_server.py
+  - tests/v1/kv_connector/unit/test_moriio_
+  - tests/v1/kv_connector/unit/utils.py
+  - requirements/kv_connectors_rocm.txt
+  commands:
+  - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s -rs v1/kv_connector/unit/test_moriio_*.py
+
+# Single-node MoRIIO WRITE P/D on dummy weights: prefill and decode on separate
+# GPUs over xGMI, compared exactly against a direct serve.
+- label: ":amd: (MI355) MoRIIO WRITE P/D Dummy-Weight E2E Shard %N"
+  key: amd-moriio-write-dummy-e2e-mi355
+  dind: false
+  device: mi355_2
+  num_devices: 2
+  # ~3.5-5 min per case; pytest-shard hashes test ids, so with 20 cases the
+  # largest of 6 shards holds 5 (~22 min on MI355). Every case picks free ports
+  # for its servers and proxy and its own compile-cache dir.
+  parallelism: 6
+  timeout_in_minutes: 45
+  working_dir: /vllm-workspace/tests
+  depends_on:
+  - image-build-amd
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/moriio/
+  - vllm/distributed/kv_transfer/kv_connector/v1/base.py
+  - vllm/distributed/kv_transfer/kv_connector/v1/ssm_conv_transfer_utils.py
+  - vllm/model_executor/model_loader/dummy_loader.py
+  - examples/disaggregated/disaggregated_serving/moriio_toy_proxy_server.py
+  - tests/v1/kv_connector/moriio_integration/
+  - requirements/kv_connectors_rocm.txt
+  commands:
+  - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s -rs v1/kv_connector/moriio_integration/test_write_dummy_e2e_rocm.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+
 - label: ":nvidia: (H200 MIG 35GB) V1 Metrics + LM Eval"
   device: h200_35gb
   key: v1-metrics-lmeval
@@ -213,12 +267,14 @@ steps:
     - vllm/v1/
     - tests/v1
     - tests/watermarking
+    - examples/disaggregated/mooncake_connector/mooncake_connector_proxy.py
   device: cpu-small
   commands:
     # split the test to avoid interference
     - pytest -v -s -m 'cpu_test' v1/core
     - pytest -v -s v1/structured_output
     - pytest -v -s v1/test_serial_utils.py
+    - pytest -v -s v1/worker/test_cpu_worker_kv_budget.py
     - pytest -v -s v1/test_kv_cache_spec_registry.py
     - pytest -v -s v1/cudagraph/test_cudagraph_manager.py
     - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
@@ -446,6 +502,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 +711,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 +761,19 @@ 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
-      timeout_in_minutes: 35
+      automatic_shard: true
+      device: mi300_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
@@ -687,15 +782,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 +813,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 +898,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 +913,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/
@@ -68,10 +69,12 @@ steps:
   - tests/models/test_deepseek_v4_vl_rocm.py
   - tests/models/test_deepseek_v41_decoder_replay_layers.py
   - tests/models/test_deepseek_v41_replay_batch.py
+  - tests/models/test_deepseek_v41_pcp.py
   - tests/models/transformers/test_layer_registry_e2e.py
   - tests/models/test_deepseek_v41_replay_start.py
+  - tests/models/test_deepseek_v4_vl_stream_weights.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
+    - 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_pcp.py models/test_deepseek_v4_vl_stream_weights.py
   mirror:
     amd:
       label: ":amd: (MI355) Basic Models (Other)"
@@ -82,17 +85,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 +137,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 +161,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 +174,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 +183,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
-  parallelism: 2
+  - pytest -v -s models/kimi_k3/test_prefix_cache.py
 
 - label: ":nvidia: (H200) GLM5Next"
   key: glm5next-unit-tests
@@ -202,12 +212,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 +238,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 +256,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 +268,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,39 @@ 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
+  mirror:
+    amd:
+      label: ":amd: (MI355) Qwen4 Exp SP"
+      dind: false
+      device: mi355_4
+      depends_on:
+      - image-build-amd
+      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
+      - vllm/platforms/rocm.py
+      - tests/models/qwen4_exp/test_ple.py
+      - tests/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,23 +199,29 @@ 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
+    # Per-test timeout: a hung test fails with a stack dump instead of holding
+    # the job until Buildkite cancels it, which also blocks the full run's
+    # coverage collection. A whole shard takes 12-21 minutes.
+    - pytest -v -s models/language/pooling -m 'not core_model' --timeout=900 --timeout-method=thread
   mirror:
     amd:
-      label: ":amd: (MI300) Language Models (Extended Pooling) Shard %N"
+      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: (MI355 DPX) Language Models (Extended Pooling) Shard %N"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 95
       depends_on:
       - image-build-amd

.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:
@@ -162,9 +163,9 @@ steps:
     - pytest -v -s models/multimodal/test_mapping.py
   mirror:
     amd:
-      label: ":amd: (MI300) Multimodal Models (Extended Generation 1)"
+      label: ":amd: (MI355 DPX) Multimodal Models (Extended Generation 1)"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 100
       depends_on:
       - image-build-amd
@@ -181,9 +182,9 @@ steps:
     - pytest -v -s models/multimodal/generation_ppl_test/
   mirror:
     amd:
-      label: ":amd: (MI300) Multimodal Models (PPL)"
+      label: ":amd: (MI355 DPX) Multimodal Models (PPL)"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 90
       depends_on:
       - image-build-amd
@@ -239,9 +240,9 @@ steps:
     - pytest -v -s models/multimodal/pooling -m 'not core_model'
   mirror:
     amd:
-      label: ":amd: (MI300) Multimodal Models (Extended Pooling)"
+      label: ":amd: (MI355 DPX) Multimodal Models (Extended Pooling)"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 75
       depends_on:
       - image-build-amd

.buildkite/test_areas/pytorch.yaml

@@ -160,6 +160,7 @@ steps:
     amd:
       label: ':amd: (MI300) PyTorch Compilation Passes'
       dind: false
+      automatic_shard: true
       device: mi300_1
       timeout_in_minutes: 80
       working_dir: /vllm-workspace/tests

.buildkite/test_areas/quantization.yaml

@@ -2,27 +2,27 @@ 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
   - vllm/models/deepseek_v41/quant_config.py
+  - vllm/utils/humming.py
   - 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 +41,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 +128,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
@@ -141,6 +137,7 @@ steps:
     amd:
       label: ":amd: (MI355 DPX) Quantized Models"
       dind: false
+      automatic_shard: true
       device: mi355_dpx
       timeout_in_minutes: 80
       working_dir: "/vllm-workspace/tests"
@@ -153,6 +150,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

@@ -28,6 +28,7 @@ steps:
     amd:
       label: ":amd: (MI300) V1 Spec Decode"
       dind: false
+      automatic_shard: true
       device: mi300_1
       timeout_in_minutes: 50
       depends_on:
@@ -136,9 +137,9 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/speculators/
   mirror:
     amd:
-      label: ":amd: (MI300) Spec Decode Speculators"
+      label: ":amd: (MI355 DPX) Spec Decode Speculators"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 40
       depends_on:
       - image-build-amd
@@ -166,9 +167,9 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/mtp/gemma4/
   mirror:
     amd:
-      label: ":amd: (MI300) Spec Decode MTP Gemma4"
+      label: ":amd: (MI355 DPX) Spec Decode MTP Gemma4"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 40
       depends_on:
       - image-build-amd
@@ -197,9 +198,9 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/mtp/mimo/
   mirror:
     amd:
-      label: ":amd: (MI300) Spec Decode MTP DeepSeek + MiMo"
+      label: ":amd: (MI355 DPX) Spec Decode MTP DeepSeek + MiMo"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 40
       depends_on:
       - image-build-amd
@@ -227,9 +228,9 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/mtp/qwen3_5/
   mirror:
     amd:
-      label: ":amd: (MI300) Spec Decode MTP Qwen3.5"
+      label: ":amd: (MI355 DPX) Spec Decode MTP Qwen3.5"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 40
       depends_on:
       - image-build-amd
@@ -386,9 +387,9 @@ steps:
     - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
   mirror:
     amd:
-      label: ":amd: (MI300) Speculators Correctness Nightly"
+      label: ":amd: (MI355 DPX) Speculators Correctness Nightly"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 40
       depends_on:
       - image-build-amd
@@ -451,9 +452,9 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/acceptance_rates/dflash/ -k "not dflash2"
   mirror:
     amd:
-      label: ":amd: (MI300) Spec Decode AL DFlash Nightly"
+      label: ":amd: (MI355 DPX) Spec Decode AL DFlash Nightly"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 120
       commands:
       - pytest -v -s v1/e2e/spec_decode/acceptance_rates/dflash/ -k "not dflash2"
@@ -530,9 +531,9 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/acceptance_rates/dspark/
   mirror:
     amd:
-      label: ":amd: (MI300) Spec Decode AL DSpark Nightly"
+      label: ":amd: (MI355 DPX) Spec Decode AL DSpark Nightly"
       dind: false
-      device: mi300_1
+      device: mi355_dpx
       timeout_in_minutes: 90
       depends_on:
       - image-build-amd
@@ -594,7 +595,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/fused_moe @mgoin @pavanimajety @zyongye @wzhao18 @bnellnm
 /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/speech_to_text/realtime @njhill
+/vllm/entrypoints/openai/chat_completion @chaunceyjiang @DarkLight1337
+/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/serve/sagemaker @DarkLight1337
+/vllm/entrypoints/speech_to_text/realtime @njhill @NickLucche
 /vllm/entrypoints/serve @njhill
-/vllm/entrypoints/*.py @njhill
-/vllm/entrypoints/chat_utils.py @DarkLight1337
+/vllm/entrypoints/serve/sagemaker @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,23 @@
 /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
+/tests/evals @mgoin @vadiklyutiy @AndreasKaratzas @ZJY0516
 /tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye @AndreasKaratzas
 /tests/kernels/ir @ProExpertProg @tjtanaa
+/tests/kernels/moe @bnellnm
 /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
@@ -135,6 +165,9 @@
 .readthedocs.yaml @hmellor
 mkdocs.yaml @hmellor
 
+# Model recipes
+/recipes @esmeetu
+
 # Linting
 /tools/pre_commit/.markdownlint.yaml @hmellor
 .pre-commit-config.yaml @hmellor
@@ -223,15 +256,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|medusa|mlp_speculator)\.py
+      - files~=^vllm/transformers_utils/configs/eagle\.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/recipes.yml

@@ -0,0 +1,28 @@
+name: Recipes
+
+# Validate recipe changes against the renderer that publishes them
+# (recipes.vllm.ai). The check logic lives in vllm-project/recipes and tracks
+# its main branch, so a PR is held to the same validation that runs when
+# vllm-project/recipes picks up this directory after merge.
+
+on:
+  pull_request:
+    paths:
+      - 'recipes/**'
+      - '.github/workflows/recipes.yml'
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.ref }}
+  cancel-in-progress: true
+
+permissions:
+  contents: read
+
+jobs:
+  validate:
+    name: Recipes
+    if: github.repository_owner == 'vllm-project'
+    uses: vllm-project/recipes/.github/workflows/validate-recipes.yml@main
+    with:
+      source: vllm
+      path: recipes

.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' ||
-      github.event.comment.body == '/ci run all' ||
-      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')
+      (startsWith(github.event.comment.body, '/ci') ||
+      startsWith(github.event.comment.body, '/amd-ci'))
     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,
     *,
@@ -605,6 +621,16 @@ def blocks_new_run(command: str, build: Mapping[str, Any]) -> bool:
     return is_active_build(build)
 
 
+def list_commit_builds(
+    buildkite: BuildkiteClient, command: str, commit: str
+) -> list[dict[str, Any]]:
+    builds = buildkite.list_builds(commit)
+    if command not in AMD_CI_COMMANDS:
+        return builds
+    # bootstrap-amd.sh exits before uploading any test steps for PR webhook builds.
+    return [build for build in builds if build.get("source") != "webhook"]
+
+
 def select_latest_build(
     builds: Sequence[dict[str, Any]],
     pr_number: int,
@@ -771,10 +797,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 "
-            "target branch. Merge or rebase onto the latest target branch, then "
-            "rerun the command. Append `--allow-stale` to a run command to test "
-            "an outdated branch at your own risk.\n"
+            "- Run commands test your branch with every commit currently on its "
+            "upstream target branch. If it is behind, the bot merges the target "
+            "branch into it first (pull before pushing again). Append "
+            "`--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 +849,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 +943,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 +964,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(
@@ -892,7 +978,7 @@ def handle_run_ci(
             f"{ci_name} was already requested by this comment: {duplicate['web_url']}"
         )
 
-    current_builds = buildkite.list_builds(pr["head"]["sha"])
+    current_builds = list_commit_builds(buildkite, command, pr["head"]["sha"])
     active_build = next(
         (
             build
@@ -906,7 +992,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(
@@ -935,7 +1033,7 @@ def handle_retry_failed(
     ci_name = ci_name_for_command(command)
     run_command = run_command_for_command(command)
     retry_command = retry_command_for_command(command)
-    builds = buildkite.list_builds(pr["head"]["sha"])
+    builds = list_commit_builds(buildkite, command, pr["head"]["sha"])
     build = select_latest_build(builds, pr["number"])
     if build:
         metadata = build.get("meta_data") or {}
@@ -1086,11 +1184,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 +1274,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 +1340,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 +1357,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:
@@ -756,6 +883,42 @@ class RunCiCommandTest(unittest.TestCase):
         self.assertEqual(buildkite.created_builds, [])
         self.assertIn("AMD CI is already running", github.comments[0])
 
+    def test_amd_run_ignores_finished_webhook_build_for_author(self) -> None:
+        webhook_build = {
+            "created_at": "2026-08-18T01:00:00Z",
+            "meta_data": {},
+            "number": 122,
+            "pull_request": {"id": 42},
+            "source": "webhook",
+            "state": "passed",
+            "web_url": "https://buildkite.example/amd-ci/builds/122",
+        }
+        ui_build = {**webhook_build, "source": "ui"}
+        command_build = {
+            **webhook_build,
+            "meta_data": {"github-comment-id": "98"},
+            "source": "api",
+        }
+        for build, allowed in (
+            (webhook_build, True),
+            (ui_build, False),
+            (command_build, False),
+        ):
+            with self.subTest(source=build["source"]):
+                github = FakeGitHub(
+                    permission="read",
+                    pr=make_pr(labels=[{"name": "ready"}]),
+                )
+                buildkite = FakeBuildkite([[], [build]])
+
+                run(make_event(COMMAND_RUN_AMD_CI, "author"), github, buildkite)
+
+                self.assertEqual(bool(buildkite.created_builds), allowed)
+                if not allowed:
+                    self.assertIn(
+                        "AMD CI already ran on this commit", github.comments[0]
+                    )
+
     def test_run_all_sets_buildkite_environment(self) -> None:
         github = FakeGitHub()
         buildkite = FakeBuildkite([[], []])
@@ -1102,6 +1265,29 @@ class RunCiCommandTest(unittest.TestCase):
         )
         self.assertIn("Use `/amd-ci run`", github.comments[0])
 
+    def test_amd_ci_retry_ignores_webhook_build_for_the_current_head(self) -> None:
+        github = FakeGitHub(
+            permission="read",
+            pr=make_pr(labels=[{"name": "ready"}]),
+        )
+        webhook_build = {
+            "created_at": "2026-08-18T01:00:00Z",
+            "meta_data": {},
+            "number": 122,
+            "pull_request": {"id": 42},
+            "source": "webhook",
+            "state": "passed",
+            "web_url": "https://buildkite.example/amd-ci/builds/122",
+        }
+        buildkite = FakeBuildkite([[webhook_build]])
+
+        run(make_event(COMMAND_RETRY_AMD_FAILED, "author"), github, buildkite)
+
+        self.assertEqual(buildkite.retry_calls, [])
+        self.assertIn(
+            "No AMD CI build exists for the current PR head", github.comments[0]
+        )
+
     def test_ci_retry_creates_filtered_build_for_new_head(self) -> None:
         github = FakeGitHub(
             behind=100,

.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
@@ -277,6 +277,12 @@ repos:
     files: ^docker/(Dockerfile|versions\.json)$
     pass_filenames: false
     additional_dependencies: [dockerfile-parse]
+  - id: check-dockerignore-tracked
+    name: Check .dockerignore does not exclude tracked files
+    entry: python tools/pre_commit/check_dockerignore_tracked.py
+    language: python
+    always_run: true
+    pass_filenames: false
   - id: check-boolean-context-manager
     name: Check for boolean ops in with-statements
     entry: python tools/pre_commit/check_boolean_context_manager.py

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_ROCM "2.13.0")
+set(TORCH_SUPPORTED_VERSION_CUDA "2.14.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/custom_quickreduce.cu")
+    "csrc/rocm/quickreduce/torch_bindings.cpp"
+    "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,42 @@ 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/libtorch_stable/cuda_view.cu"
-    "csrc/libtorch_stable/cuda_utils_kernels.cu"
-    "csrc/libtorch_stable/activation_kernels.cu"
-    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
-    "csrc/libtorch_stable/vocab_parallel_embedding_kernels.cu"
-    "csrc/libtorch_stable/quantization/activation_kernels.cu"
-    "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
-    "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
-    "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
-    "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
-    "csrc/libtorch_stable/quantization/gptq/q_gemm.cu"
-    "csrc/libtorch_stable/pos_encoding_kernels.cu"
-    "csrc/libtorch_stable/fused_qknorm_rope_kernel.cu"
-    "csrc/libtorch_stable/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu"
-    "csrc/libtorch_stable/layernorm_kernels.cu"
-    "csrc/libtorch_stable/layernorm_quant_kernels.cu"
-    "csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
-    "csrc/libtorch_stable/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
-    "csrc/libtorch_stable/attention/merge_attn_states.cu"
-    "csrc/libtorch_stable/sampler.cu"
-    "csrc/libtorch_stable/topk.cu"
-    "csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
-    "csrc/libtorch_stable/cache_kernels.cu"
-    "csrc/libtorch_stable/cache_kernels_fused.cu"
-    "csrc/libtorch_stable/custom_all_gather_reduce_scatter.cu"
-    "csrc/libtorch_stable/custom_all_gather_reduce_scatter_ops.cpp"
-    "csrc/libtorch_stable/custom_all_reduce.cu"
-    "csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu"
-    "csrc/libtorch_stable/fused_kimi_k3_mla_key_concat_kv_cache_kernel.cu")
+    "csrc/torch_bindings.cpp"
+    "csrc/cuda_view.cu"
+    "csrc/cuda_utils_kernels.cu"
+    "csrc/activation_kernels.cu"
+    "csrc/muse_spark_mul_tanh_soft_cap_kernel.cu"
+    "csrc/ngram_embedding_kernels.cu"
+    "csrc/vocab_parallel_embedding_kernels.cu"
+    "csrc/quantization/activation_kernels.cu"
+    "csrc/quantization/w8a8/int8/scaled_quant.cu"
+    "csrc/quantization/w8a8/fp8/common.cu"
+    "csrc/quantization/w8a8/fp8/per_token_group_quant.cu"
+    "csrc/quantization/gptq/q_gemm.cu"
+    "csrc/pos_encoding_kernels.cu"
+    "csrc/fused_muse_spark_qknorm_rope_kv_insert_kernel.cu"
+    "csrc/fused_qknorm_rope_kernel.cu"
+    "csrc/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu"
+    "csrc/layernorm_kernels.cu"
+    "csrc/layernorm_quant_kernels.cu"
+    "csrc/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
+    "csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu"
+    "csrc/attention/merge_attn_states.cu"
+    "csrc/sampler.cu"
+    "csrc/topk.cu"
+    "csrc/mamba/selective_scan_fwd.cu"
+    "csrc/cache_kernels.cu"
+    "csrc/cache_kernels_fused.cu"
+    "csrc/custom_all_gather_reduce_scatter.cu"
+    "csrc/custom_all_gather_reduce_scatter_ops.cpp"
+    "csrc/custom_all_reduce.cu"
+    "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu"
+    "csrc/fused_kimi_k3_mla_key_concat_kv_cache_kernel.cu")
+
+  if(VLLM_GPU_LANG STREQUAL "CUDA")
+    list(APPEND VLLM_STABLE_EXT_SRC
+      "csrc/fused_qknorm_rope_cache_kernel.cu")
+  endif()
 
   if(VLLM_GPU_LANG STREQUAL "CUDA" AND
      DEFINED CMAKE_CUDA_COMPILER_VERSION AND
@@ -470,10 +468,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/libtorch_stable/attention/dcp_utils/dcp_direct_a2a_lse_reduce.cu"
-      "csrc/libtorch_stable/attention/dcp_utils/dcp_direct_kv_gather.cu"
-      "csrc/libtorch_stable/attention/dcp_utils/dcp_direct_q_gather.cu")
+      "csrc/hisparse_kernels.cu"
+      "csrc/attention/dcp_utils/dcp_direct_a2a_lse_reduce.cu"
+      "csrc/attention/dcp_utils/dcp_direct_kv_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 +506,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/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_entry.cu"
-      "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu"
-      "csrc/libtorch_stable/minimax_reduce_rms_kernel.cu")
+      "csrc/quantization/w8a8/cutlass/scaled_mm_entry.cu"
+      "csrc/quantization/fp4/nvfp4_quant_entry.cu"
+      "csrc/quantization/fp4/nvfp4_scaled_mm_entry.cu"
+      "csrc/quantization/awq/gemm_kernels.cu"
+      "csrc/minimax_reduce_rms_kernel.cu")
 
     #
     # Machete kernels
@@ -527,7 +524,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 +557,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 +566,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 +584,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 +608,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 +653,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 +688,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 +698,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 +710,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 +722,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 +734,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/libtorch_stable/quantization/marlin/marlin_int4_fp8_preprocess.cu"
-         "csrc/libtorch_stable/quantization/marlin/gptq_marlin_repack.cu"
-         "csrc/libtorch_stable/quantization/marlin/awq_marlin_repack.cu")
+         "csrc/quantization/marlin/marlin.cu"
+         "csrc/quantization/marlin/marlin_int4_fp8_preprocess.cu"
+         "csrc/quantization/marlin/gptq_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 +760,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 +775,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/libtorch_stable/fp32_router_gemm.cu")
+      "csrc/fp32_router_gemm_entry.cu"
+      "csrc/fp32_router_gemm.cu")
     set_gencode_flags_for_srcs(
       SRCS "${FP32_ROUTER_GEMM_SRCS}"
       CUDA_ARCHS "${FP32_ROUTER_GEMM_ARCHS}")
@@ -799,11 +796,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/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_fp8.cu"
-       "csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_int8.cu"
-       "csrc/libtorch_stable/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/scaled_mm_c3x_sm90.cu"
+       "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_fp8.cu"
+       "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_int8.cu"
+       "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_azp_sm90_int8.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 +831,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/libtorch_stable/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/scaled_mm_c3x_sm120.cu"
+      "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_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 +865,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/libtorch_stable/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/scaled_mm_c3x_sm100.cu"
+      "csrc/quantization/w8a8/cutlass/c3x/scaled_mm_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 +898,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 +924,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 +948,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 +973,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 +1005,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/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
+      "csrc/quantization/fp4/nvfp4_quant_kernels.cu"
+      "csrc/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_experts_quant.cu"
+      "csrc/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
+      "csrc/nvfp4_kv_cache_kernels.cu")
     set_gencode_flags_for_srcs(
       SRCS "${FP4_SM120_SRCS}"
       CUDA_ARCHS "${FP4_SM120_ARCHS}")
@@ -1034,15 +1031,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/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu"
-      "csrc/libtorch_stable/nvfp4_ds_mla_cache_kernels.cu")
+      "csrc/quantization/fp4/nvfp4_quant_kernels.cu"
+      "csrc/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_experts_quant.cu"
+      "csrc/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
+      "csrc/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
+      "csrc/quantization/fp4/mxfp4_experts_quant.cu"
+      "csrc/quantization/fp4/mxfp4_blockwise_moe_kernel.cu"
+      "csrc/nvfp4_kv_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 +1063,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/libtorch_stable/quantization/cutlass_w4a8/w4a8_grouped_mm_entry.cu"
-       "csrc/libtorch_stable/quantization/cutlass_w4a8/w4a8_utils.cu"
+       "csrc/quantization/cutlass_w4a8/w4a8_mm_entry.cu"
+       "csrc/quantization/cutlass_w4a8/w4a8_grouped_mm_entry.cu"
+       "csrc/quantization/cutlass_w4a8/w4a8_utils.cu"
        )
 
     set_gencode_flags_for_srcs(
@@ -1101,7 +1098,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}")
@@ -1122,10 +1119,15 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       "9.0a;10.0f;12.0f" "${CUDA_ARCHS}")
     cuda_archs_loose_intersection(FUSED_GDN_DECODE_ARCHS
       "8.0;8.6;8.9;9.0a;10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
+  else()
+    # The fused GDN decode kernel needs only cp.async and BF16 math. Only the
+    # family-specific targets (10.0f, 11.0f, 12.0f) need CUDA >= 13.0.
+    cuda_archs_loose_intersection(FUSED_GDN_DECODE_ARCHS
+      "8.0;8.6;8.9;9.0a" "${CUDA_ARCHS}")
   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 +1139,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 +1156,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 +1171,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 +1195,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 +1205,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 +1291,22 @@ endif()
 #
 
 set(VLLM_MOE_EXT_SRC
-  "csrc/libtorch_stable/moe/torch_bindings.cpp"
-  "csrc/libtorch_stable/moe/moe_align_sum_kernels.cu"
-  "csrc/libtorch_stable/moe/topk_softmax_kernels.cu"
-  "csrc/libtorch_stable/moe/topk_softplus_sqrt_kernels.cu")
+  "csrc/moe/torch_bindings.cpp"
+  "csrc/moe/moe_align_sum_kernels.cu"
+  "csrc/moe/topk_softmax_kernels.cu"
+  "csrc/moe/topk_softplus_sqrt_kernels.cu")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   list(APPEND VLLM_MOE_EXT_SRC
-    "csrc/libtorch_stable/moe/moe_wna16.cu"
-    "csrc/libtorch_stable/moe/grouped_topk_kernels.cu")
+    "csrc/moe/moe_wna16.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/libtorch_stable/moe/moe_permute_unpermute_op.cu")
+      "csrc/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.cu"
+      "csrc/moe/moe_permute_unpermute_op.cu")
 
   list(APPEND VLLM_MOE_EXT_SRC "${MOE_PERMUTE_SRC}")
 endif()
@@ -1314,7 +1317,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 +1352,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 +1366,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 +1388,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 +1400,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 +1412,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 +1423,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:
-- apply_top_k_top_p_triton (Triton binary search)
-- apply_top_k_top_p (PyTorch sort-based)
+All implementations return masked logits. Timings include a fresh-input copy.
+FlashInfer timings include softmax and mask reconstruction when its APIs require them.
+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)
-    print("Top-k/Top-p Benchmark: Triton vs PyTorch Sort-based")
-    print("=" * 100)
-    print()
-
-    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,
+    for index, config in enumerate(configs):
+        logits = create_logits(
+            config.batch_size,
+            config.vocab_size,
+            num_finite_logits=num_finite_logits,
         )
-
-        # Benchmark PyTorch
-        reset_memory_stats()
-        pytorch_time, pytorch_mem = benchmark_function(
-            apply_top_k_top_p_pytorch,
-            logits,
-            config.k_values,
-            config.p_values,
-            warmup_iters,
-            benchmark_iters,
-        )
-
-        speedup = pytorch_time / triton_time if triton_time > 0 else float("inf")
-        mem_ratio = pytorch_mem / triton_mem if triton_mem > 0 else float("inf")
-
-        result = {
-            "config": config,
-            "triton_time_ms": triton_time,
-            "pytorch_time_ms": pytorch_time,
-            "triton_mem": triton_mem,
-            "pytorch_mem": pytorch_mem,
-            "speedup": speedup,
-            "mem_ratio": mem_ratio,
-        }
-        results.append(result)
-
+        checks = {}
+        if check:
+            checks = check_implementations(logits, config, implementations)
+        k, p = config.k_values, config.p_values
+        timings = {}
+        for name in implementations:
+            func = IMPLEMENTATIONS[name]
+            reset_memory_stats()
+            times = bench_gpu_time_with_cupti(
+                # 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,
+                repeat_iters=benchmark_iters,
+                use_cuda_graph=False,  # sampling is outside cuda graph
+                cold_l2_cache=True,
+            )
+            elapsed = sum(times) / len(times)
+            _, memory = measure_memory()
+            timings[name] = dict(time_ms=elapsed, peak_memory=memory)
+        row = create_result_row(config, timings, checks)
+        results.append(row)
+        if output is not None:
+            with output.open("a") as file:
+                file.write(json.dumps(row) + "\n")
         if verbose:
-            print(f"  Triton:  {triton_time:.3f} ms, {format_memory(triton_mem)}")
-            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
+            print(f"Completed {index + 1}/{len(configs)}: {config.name}", flush=True)
         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()
-    print("=" * 130)
-    print("SUMMARY TABLE")
-    print("=" * 130)
+    """Print the same result rows saved to the output file."""
     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
-    current_vocab = None
-    for result in results:
-        config = result["config"]
-
-        # Add separator between vocab sizes
-        if current_vocab != config.vocab_size:
-            if current_vocab is not None:
-                print("-" * 130)
-            current_vocab = config.vocab_size
-
-        scenario = config.name.split("_b")[0]  # Extract scenario name
-        print(
-            f"{scenario:<40} {config.batch_size:>6} {config.vocab_size:>7} "
-            f"{config.ops_pct:>5.0f}% "
-            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 create_result_row(config, timings, checks) -> dict:
+    """Build a shared row for the summary table and JSONL output."""
+    row = {
+        "Scenario": config.name.split("_b")[0],
+        "Batch": config.batch_size,
+        "Vocab": config.vocab_size,
+        "Ops%": f"{config.ops_pct:.0f}%",
+    }
+    row.update({f"{name} (ms)": data["time_ms"] for name, data in timings.items()})
+    for name, data in timings.items():
+        row[f"{name} mem"] = format_memory(data["peak_memory"])
+    for name, data in checks.items():
+        row[f"{name} check"] = data["close"]
+        row[f"{name} max TV"] = data["max_total_variation"]
+    return row
 
 
 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
-    print(f"Batch sizes: {args.batch_sizes}")
-    print(f"Vocab sizes: {args.vocab_sizes}")
-    print(f"Warmup iterations: {args.warmup_iters}")
-    print(f"Benchmark iterations: {args.benchmark_iters}")
-    print()
+    parser.add_argument(
+        "--implementations",
+        nargs="+",
+        choices=list(IMPLEMENTATIONS),
+        default=["triton", "flashinfer-mask"],
+        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(
-        args.batch_sizes,
-        args.vocab_sizes,
-    )
+    configs = create_benchmark_configs(args.batch_sizes, args.vocab_sizes)
+    if args.scenarios is not None:
+        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_fused_topk.py

@@ -11,7 +11,7 @@ from vllm.utils.argparse_utils import FlexibleArgumentParser
 
 num_tokens_range = [2**i for i in range(0, 8, 2)]
 num_experts_range = [16, 32, 64, 128, 256, 512]
-topk_range = [3, 4]
+topk_range = [3, 4, 8, 10]
 configs = list(itertools.product(num_tokens_range, num_experts_range, topk_range))
 
 
@@ -39,9 +39,9 @@ def get_benchmark(scoring_func):
             x_names=["num_tokens", "num_experts", "topk"],
             x_vals=[list(_) for _ in configs],
             line_arg="provider",
-            line_vals=["torch", "vllm"],
-            line_names=["Torch", "vLLM"],
-            styles=[("blue", "-"), ("red", "-")],
+            line_vals=["torch", "vllm", "aiter_gating"],
+            line_names=["Torch", "vLLM", "AITER topk_gating"],
+            styles=[("blue", "-"), ("red", "-"), ("green", "-")],
             ylabel="us",
             plot_name=f"fused-topk-perf-{scoring_func}",
             args={},
@@ -70,6 +70,28 @@ def get_benchmark(scoring_func):
                 ),
                 quantiles=quantiles,
             )
+        elif provider == "aiter_gating":
+            from vllm._aiter_ops import rocm_aiter_ops
+
+            if scoring_func != "softmax" or not rocm_aiter_ops.is_topk_gating_enabled():
+                return float("nan"), float("nan"), float("nan")
+            topk_weights = torch.empty(
+                (num_tokens, topk), dtype=torch.float32, device="cuda"
+            )
+            topk_ids = torch.empty((num_tokens, topk), dtype=torch.int32, device="cuda")
+            token_expert_indices = torch.empty(
+                (num_tokens, topk), dtype=torch.int32, device="cuda"
+            )
+            ms, min_ms, max_ms = triton.testing.do_bench(
+                lambda: rocm_aiter_ops.topk_gating(
+                    topk_weights,
+                    topk_ids,
+                    token_expert_indices,
+                    gating_output,
+                    renormalize,
+                ),
+                quantiles=quantiles,
+            )
         else:
             ms, min_ms, max_ms = triton.testing.do_bench(
                 lambda: fused_topk(

benchmarks/kernels/benchmark_muse_spark_qknorm_rope_kv_insert.py

@@ -0,0 +1,326 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Benchmark the fused Muse-Spark QK-norm + RoPE + KV-insert kernel vs unfused.
+
+``tuned_us`` forces the cp.async schedule ``fused_qk_norm_rope`` only auto-selects
+on SM 9.0; ``unfused_us`` is what production runs.
+"""
+
+import argparse
+import itertools
+
+import torch
+
+import vllm._custom_ops as ops
+
+HEAD_DIM = 64
+DEFAULT_TOKENS = [1, 8, 64, 256, 1024, 4096, 8192, 16384, 32768, 65536]
+
+
+def make_cos_sin_cache(max_pos, rotary_dim, base, device):
+    inv_freq = 1.0 / (
+        base
+        ** (
+            torch.arange(0, rotary_dim, 2, dtype=torch.float32, device=device)
+            / rotary_dim
+        )
+    )
+    t = torch.arange(max_pos, dtype=torch.float32, device=device)
+    freqs = torch.einsum("i,j->ij", t, inv_freq)
+    return torch.cat((freqs.cos(), freqs.sin()), dim=-1)
+
+
+class Case:
+    """Everything a single (num_tokens, config) measurement needs."""
+
+    def __init__(self, num_tokens, nq, nkv, dtype, kv_cache_dtype, device):
+        self.num_tokens = num_tokens
+        self.nq, self.nkv = nq, nkv
+        self.dtype = dtype
+        self.kv_cache_dtype = kv_cache_dtype
+        self.eps = 1e-5
+        self.block_size = 128
+        self.rotary_dim = HEAD_DIM
+
+        self.qsz, self.kvsz = nq * HEAD_DIM, nkv * HEAD_DIM
+        self.row = self.qsz + 2 * self.kvsz + self.qsz  # [q | k | v | o_gate]
+        self.qkv = torch.randn(num_tokens, self.row, dtype=dtype, device=device)
+        self.pristine = self.qkv.clone()
+
+        max_pos = 32768
+        self.cos_sin = make_cos_sin_cache(max_pos, self.rotary_dim, 500_000.0, device)
+        self.positions = torch.randint(
+            0, max_pos, (num_tokens,), dtype=torch.int64, device=device
+        )
+
+        num_blocks = (num_tokens + self.block_size - 1) // self.block_size + 1
+        storage = torch.uint8 if kv_cache_dtype != "auto" else dtype
+        self.kv_cache = torch.zeros(
+            num_blocks,
+            nkv,
+            self.block_size,
+            2 * HEAD_DIM,
+            dtype=storage,
+            device=device,
+        )
+        self.k_cache, self.v_cache = self.kv_cache.transpose(1, 2).split(
+            HEAD_DIM, dim=-1
+        )
+        self.slot_mapping = torch.randperm(
+            num_blocks * self.block_size, dtype=torch.int64, device=device
+        )[:num_tokens]
+        self.k_scale = torch.tensor(1.0, dtype=torch.float32, device=device)
+        self.v_scale = torch.tensor(1.0, dtype=torch.float32, device=device)
+        self.ones = torch.ones(HEAD_DIM, dtype=dtype, device=device)
+        self.q_fp8 = torch.empty(
+            num_tokens, self.qsz, dtype=torch.float8_e4m3fn, device=device
+        )
+        self.q_scale = torch.tensor(0.37, dtype=torch.float32, device=device)
+
+    def views(self):
+        q, k, v, _ = self.qkv.split([self.qsz, self.kvsz, self.kvsz, self.qsz], dim=-1)
+        return q, k, v
+
+    def fused(self, use_rope=True, quant_query=False):
+        ops.fused_muse_spark_qknorm_rope_kv_insert(
+            self.qkv,
+            self.positions,
+            self.nq,
+            self.nkv,
+            HEAD_DIM,
+            self.eps,
+            cos_sin_cache=self.cos_sin if use_rope else None,
+            rotary_dim=self.rotary_dim if use_rope else 0,
+            slot_mapping=self.slot_mapping,
+            kv_cache=self.kv_cache,
+            block_size=self.block_size,
+            kv_cache_dtype=self.kv_cache_dtype,
+            k_scale=self.k_scale if self.kv_cache_dtype != "auto" else None,
+            v_scale=self.v_scale if self.kv_cache_dtype != "auto" else None,
+            q_fp8_out=self.q_fp8 if quant_query else None,
+            q_scale=self.q_scale if quant_query else None,
+        )
+
+    def separate_query_quant(self, use_rope=True):
+        self.fused(use_rope)
+        ops.scaled_fp8_quant(self.qkv[:, : self.qsz], scale=self.q_scale)
+
+    def _insert(self):
+        _, k, v = self.views()
+        ops.reshape_and_cache_flash(
+            k.view(-1, self.nkv, HEAD_DIM),
+            v.view(-1, self.nkv, HEAD_DIM),
+            self.k_cache,
+            self.v_cache,
+            self.slot_mapping,
+            self.kv_cache_dtype,
+            self.k_scale,
+            self.v_scale,
+        )
+
+    def unfused_rope(self, token_heads_per_warp: int = -1):
+        ops.fused_qk_norm_rope(
+            self.qkv,
+            self.nq,
+            self.nkv,
+            self.nkv + self.nq,  # v heads plus the o-gate
+            HEAD_DIM,
+            self.eps,
+            self.ones,
+            self.ones,
+            self.cos_sin,
+            False,
+            self.positions,
+            token_heads_per_warp,
+        )
+        self._insert()
+
+    def fused_bytes(self):
+        el = self.qkv.element_size()
+        cache_el = self.kv_cache.element_size()
+        read = (self.nq + 2 * self.nkv) * HEAD_DIM * el
+        write = (self.nq + self.nkv) * HEAD_DIM * el
+        cache = 2 * self.nkv * HEAD_DIM * cache_el
+        return self.num_tokens * (read + write + cache)
+
+
+def _l2_bytes(device) -> int:
+    return torch.cuda.get_device_properties(device).L2_cache_size
+
+
+def _peak_bandwidth(device) -> float:
+    """Measured large-copy bandwidth, in GB/s -- the practical ceiling."""
+    n = 1 << 28  # 256 Mi elements = 512 MiB each way at 2 bytes
+    a = torch.empty(n, dtype=torch.bfloat16, device=device)
+    b = torch.empty_like(a)
+    for _ in range(3):
+        b.copy_(a)
+    torch.accelerator.synchronize()
+    start, end = torch.cuda.Event(True), torch.cuda.Event(True)
+    start.record()
+    iters = 20
+    for _ in range(iters):
+        b.copy_(a)
+    end.record()
+    torch.accelerator.synchronize()
+    ms = start.elapsed_time(end) / iters
+    return (2 * a.numel() * a.element_size()) / (ms * 1e-3) / 1e9
+
+
+def time_eager(fn, warmup=20, iters=100) -> float:
+    for _ in range(warmup):
+        fn()
+    torch.accelerator.synchronize()
+    start, end = torch.cuda.Event(True), torch.cuda.Event(True)
+    start.record()
+    for _ in range(iters):
+        fn()
+    end.record()
+    torch.accelerator.synchronize()
+    return start.elapsed_time(end) / iters * 1e3  # us
+
+
+def time_graph(fn, warmup=10, iters=50, inner=32) -> float:
+    """GPU time per call; ``inner`` calls per graph amortize the ~2 us replay."""
+    side = torch.cuda.Stream()
+    side.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(side):
+        for _ in range(warmup):
+            fn()
+    torch.cuda.current_stream().wait_stream(side)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        for _ in range(inner):
+            fn()
+    for _ in range(3):
+        graph.replay()
+    torch.accelerator.synchronize()
+    start, end = torch.cuda.Event(True), torch.cuda.Event(True)
+    start.record()
+    for _ in range(iters):
+        graph.replay()
+    end.record()
+    torch.accelerator.synchronize()
+    return start.elapsed_time(end) / (iters * inner) * 1e3  # us
+
+
+def main():
+    p = argparse.ArgumentParser()
+    p.add_argument("--nq", type=int, default=32, help="query heads per rank")
+    p.add_argument("--nkv", type=int, default=4, help="kv heads per rank")
+    p.add_argument("--kv-cache-dtype", default="auto", choices=["auto", "fp8"])
+    p.add_argument("--dtype", default="bfloat16", choices=["bfloat16", "float16"])
+    p.add_argument("--tokens", type=int, nargs="*", default=DEFAULT_TOKENS)
+    p.add_argument("--csv", default=None)
+    p.add_argument("--no-graph", action="store_true")
+    p.add_argument(
+        "--query-fp8", action="store_true", help="compare fused Q quantization"
+    )
+    args = p.parse_args()
+
+    device = "cuda"
+    dtype = getattr(torch, args.dtype)
+    if args.query_fp8:
+        if args.kv_cache_dtype != "fp8":
+            p.error("--query-fp8 requires --kv-cache-dtype fp8")
+        timer = time_eager if args.no_graph else time_graph
+        print(f"device: {torch.cuda.get_device_name()}; nq={args.nq}, nkv={args.nkv}")
+        print("layer,tokens,separate_us,fused_us,speedup")
+        rows = []
+        for use_rope, nt in itertools.product([True, False], args.tokens):
+            case = Case(nt, args.nq, args.nkv, dtype, args.kv_cache_dtype, device)
+            separate_us = timer(lambda c=case, r=use_rope: c.separate_query_quant(r))
+            case.qkv.copy_(case.pristine)
+            fused_us = timer(lambda c=case, r=use_rope: c.fused(r, quant_query=True))
+            row = (
+                "rope" if use_rope else "nope",
+                nt,
+                separate_us,
+                fused_us,
+                separate_us / fused_us,
+            )
+            rows.append(row)
+            print(f"{row[0]},{nt},{separate_us:.3f},{fused_us:.3f},{row[4]:.3f}")
+        if args.csv:
+            import csv
+
+            with open(args.csv, "w", newline="") as f:
+                writer = csv.writer(f)
+                writer.writerow(
+                    ["layer", "tokens", "separate_us", "fused_us", "speedup"]
+                )
+                writer.writerows(rows)
+        return
+    peak = _peak_bandwidth(device)
+    print(
+        f"device: {torch.cuda.get_device_name()}   measured copy peak: {peak:.0f} GB/s"
+    )
+    print(
+        f"config: nq={args.nq} nkv={args.nkv} head_dim={HEAD_DIM} "
+        f"dtype={args.dtype} kv_cache={args.kv_cache_dtype}\n"
+    )
+
+    modes = ["rope"]
+    l2 = _l2_bytes(device)
+    print(
+        f"rows marked * have a working set inside the {l2 / 2**20:.0f} MiB L2, "
+        "so their bandwidth is not an HBM number\n"
+    )
+    header = (
+        f"{'layer':>5} {'tokens':>7} {'fused_us':>9} {'unfused_us':>11} "
+        f"{'tuned_us':>9} {'speedup':>8} {'GB/s':>8} {'SOL%':>6}"
+    )
+    rows = []
+    for mode, use_graph in itertools.product(modes, [False, True]):
+        if use_graph and args.no_graph:
+            continue
+        label = "graph" if use_graph else "eager"
+        print(f"── {mode} layer, {label} ──")
+        print(header)
+        timer = time_graph if use_graph else time_eager
+        for nt in args.tokens:
+            case = Case(nt, args.nq, args.nkv, dtype, args.kv_cache_dtype, device)
+            use_rope = mode == "rope"
+            fused_us = timer(lambda c=case, r=use_rope: c.fused(r))
+            baseline = case.unfused_rope
+            unfused_us = timer(baseline)
+            tuned_us = (
+                min(timer(lambda c=case, t=t: c.unfused_rope(t)) for t in (4, 8))
+                if use_rope
+                else unfused_us
+            )
+            nbytes = case.fused_bytes()
+            gbps = nbytes / (fused_us * 1e-6) / 1e9
+            sol = 100.0 * gbps / peak
+            resident = "*" if nbytes < l2 else " "
+            print(
+                f"{mode:>5} {nt:>7} {fused_us:>9.2f} {unfused_us:>11.2f} "
+                f"{tuned_us:>9.2f} {tuned_us / fused_us:>7.2f}x {gbps:>8.0f} "
+                f"{sol:>5.1f}%{resident}"
+            )
+            rows.append((mode, label, nt, fused_us, unfused_us, tuned_us, gbps, sol))
+        print()
+
+    if args.csv:
+        import csv
+
+        with open(args.csv, "w", newline="") as f:
+            w = csv.writer(f)
+            w.writerow(
+                [
+                    "layer",
+                    "timing",
+                    "tokens",
+                    "fused_us",
+                    "unfused_us",
+                    "tuned_us",
+                    "gbps",
+                    "sol",
+                ]
+            )
+            w.writerows(rows)
+        print(f"wrote {args.csv}")
+
+
+if __name__ == "__main__":
+    main()

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

@@ -27,14 +27,13 @@ if(DEEPGEMM_SRC_DIR)
   set(deepgemm_SOURCE_DIR "${DEEPGEMM_SRC_DIR}")
   message(STATUS "DeepGEMM using local DEEPGEMM_SRC_DIR: ${deepgemm_SOURCE_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
-  # port, the SM90 paged-MQA kv_block=32/next_n=4 port, configurable SwiGLU
-  # alpha/beta, and SiTU for FP8/FP4 Mega MoE. Also includes the CUDA 12.x
-  # layout header fix from vllm-project/DeepGEMM#12.
-  # TORCH_LIBRARY build with a CPython 3.10 stable-ABI floor.
-  set(_DEEPGEMM_UPSTREAM_TAG "1e1842a833699298f7afc02eefb2ed168fff6938")
+  # Pinned to the fork's dev branch: upstream's 26/09/30 release (2.8.1)
+  # plus the SM120 port, the SM90 paged-MQA kv_block=32/next_n=4 port,
+  # configurable SwiGLU alpha/beta, SiTU for FP8/FP4 Mega MoE, and NVFP4
+  # Mega MoE, plus swiglustep for BF16/FP8/NVFP4 Mega MoE.
+  # Uses the PyTorch 2.11 stable ABI and CPython 3.10 limited API.
+  set(_DEEPGEMM_UPSTREAM_TAG "a7a927731b109fdfec770df7c2a5be5d7e2216b5")
 
   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/flashmla.cmake

@@ -19,7 +19,7 @@ else()
   FetchContent_Declare(
         flashmla
         GIT_REPOSITORY https://github.com/vllm-project/FlashMLA
-        GIT_TAG 0eee43b12f034b657133cf2afca6a72ebb6efccf
+        GIT_TAG 5dceadbd9eabc9acfc547ef18ec8f7b14b980682
         GIT_PROGRESS TRUE
         CONFIGURE_COMMAND ""
         BUILD_COMMAND ""
@@ -81,72 +81,83 @@ if(FLASH_MLA_ARCHS)
         ${flashmla_SOURCE_DIR}/csrc/api/dense_decode.cpp
         ${flashmla_SOURCE_DIR}/csrc/api/fused_norm_rope_attn_rope_cast_fwd.cpp
 
-        # Misc kernels for decoding
-        ${flashmla_SOURCE_DIR}/csrc/kernels/smxx/decode/get_decoding_sched_meta/get_decoding_sched_meta.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/smxx/decode/combine/combine.cu
-
-        # sm90 dense decode
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/decode/dense/instantiations/fp16.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/decode/dense/instantiations/bf16.cu
-
-        # sm90 sparse decode
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/decode/sparse/instantiations/v4_persistent_h64.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/decode/sparse/instantiations/v4_persistent_h128.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/decode/sparse/instantiations/v32_persistent_h64.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/decode/sparse/instantiations/v32_persistent_h128.cu
-
-        # sm90 sparse prefill
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/prefill/sparse/instantiations/phase1_k512.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/prefill/sparse/instantiations/phase1_k512_topklen.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/prefill/sparse/instantiations/phase1_k576.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm90/prefill/sparse/instantiations/phase1_k576_topklen.cu
+        # Current kernels (DeepSeek V4.1, token-interleaved KV cache scales)
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/smxx/decode/get_decoding_sched_meta/get_decoding_sched_meta.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/smxx/decode/combine/combine.cu
 
         # sm100 dense prefill (inference-only: dense backward stays out)
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/dense/fmha_cutlass_fwd_sm100.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/dense/fmha_cutlass_fwd_sm100.cu
 
         # sm100 sparse prefill
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd/head64/instantiations/phase1_h64_k512.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd/head64/instantiations/phase1_h64_k576.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd/head128/instantiations/phase1_k512.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd/head128/instantiations/phase1_k576.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_k512.cu
-
-        # sm100 fused norm + rope + attn + rope + cast
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h64_prefill_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h64_prefill_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h128_prefill_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h128_prefill_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h64_decode_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h64_decode_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h128_decode_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h128_decode_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h64_decode_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h64_decode_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h128_decode_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h128_decode_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h64_decode_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h64_decode_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h128_decode_norm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h128_decode_nonorm.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/permute_q_b_proj/kernel.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/permute_wv_proj/kernel.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fwd/head64/instantiations/phase1_h64_k512.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fwd/head128/instantiations/phase1_k512.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_k512.cu
 
         # sm100 sparse decode
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v32_h64.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v32_h64_no_split.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v4_h64.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v4_h64_no_split.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v41_h64.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v41_h64_no_split.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v41fp4_h64.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/head64/instantiations/v41fp4_h64_no_split.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/decode/sparse/nvfp4_head64/instantiations/v32_nvfp4_fp8rope.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_splitkv.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41_splitkv.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41fp4.cu
-        ${flashmla_SOURCE_DIR}/csrc/kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41fp4_splitkv.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/decode/sparse/head64/instantiations/v41_h64.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/decode/sparse/head64/instantiations/v41_h64_no_split.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/decode/sparse/head64/instantiations/v41fp4_h64.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/decode/sparse/head64/instantiations/v41fp4_h64_no_split.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41_splitkv.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41fp4.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_v41fp4_splitkv.cu
+
+        # sm100 fused norm + rope + attn + rope + cast
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h64_prefill_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h64_prefill_nonorm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h64_decode_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h64_decode_nonorm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h64_decode_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h64_decode_nonorm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h128_prefill_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h128_prefill_nonorm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h128_decode_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41_h128_decode_nonorm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h128_decode_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v41fp4_h128_decode_nonorm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/permute_q_b_proj/kernel.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/permute_wv_proj/kernel.cu
+
+        # Legacy kernels (upstream keeps them under cuda_kernels/legacy/):
+        # V4, V3.2 and V3.2-NVFP4, and everything on Hopper
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/smxx/decode/get_decoding_sched_meta/get_decoding_sched_meta.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/smxx/decode/combine/combine.cu
+
+        # sm90 dense decode
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/decode/dense/instantiations/fp16.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/decode/dense/instantiations/bf16.cu
+
+        # sm90 sparse decode
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/decode/sparse/instantiations/v4_persistent_h64.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/decode/sparse/instantiations/v4_persistent_h128.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/decode/sparse/instantiations/v32_persistent_h64.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/decode/sparse/instantiations/v32_persistent_h128.cu
+
+        # sm90 sparse prefill
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/prefill/sparse/instantiations/phase1_k512.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/prefill/sparse/instantiations/phase1_k512_topklen.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/prefill/sparse/instantiations/phase1_k576.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm90/prefill/sparse/instantiations/phase1_k576_topklen.cu
+
+        # sm100 sparse prefill, d_qk = 576
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fwd/head64/instantiations/phase1_h64_k576.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fwd/head128/instantiations/phase1_k576.cu
+
+        # sm100 sparse decode
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/decode/sparse/head64/instantiations/v32_h64.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/decode/sparse/head64/instantiations/v32_h64_no_split.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/decode/sparse/head64/instantiations/v4_h64.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/decode/sparse/head64/instantiations/v4_h64_no_split.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/decode/sparse/nvfp4_head64/instantiations/v32_nvfp4_fp8rope.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fwd_for_small_topk/head128/instantiations/phase1_decode_k512_splitkv.cu
+
+        # sm100 fused norm + rope + attn + rope + cast, V4 decode
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h64_decode_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h64_decode_nonorm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h128_decode_norm.cu
+        ${flashmla_SOURCE_DIR}/csrc/cuda_kernels/legacy/sm100/prefill/sparse/fused_norm_rope_attn_rope_cast_fwd/core_attn/instantiations/v4_h128_decode_nonorm.cu
     )
 
     set(FlashMLA_Extension_SOURCES
@@ -158,17 +169,17 @@ if(FLASH_MLA_ARCHS)
 
     set(FlashMLA_INCLUDES
         ${flashmla_SOURCE_DIR}/csrc
-        ${flashmla_SOURCE_DIR}/csrc/kerutils/include
-        ${flashmla_SOURCE_DIR}/csrc/cutlass/include
-        ${flashmla_SOURCE_DIR}/csrc/cutlass/tools/util/include
+        ${flashmla_SOURCE_DIR}/csrc/3rdparty/kerutils/include
+        ${flashmla_SOURCE_DIR}/csrc/3rdparty/cutlass/include
+        ${flashmla_SOURCE_DIR}/csrc/3rdparty/cutlass/tools/util/include
     )
 
     set(FlashMLA_Extension_INCLUDES
         ${flashmla_SOURCE_DIR}/csrc
-        ${flashmla_SOURCE_DIR}/csrc/kerutils/include
+        ${flashmla_SOURCE_DIR}/csrc/3rdparty/kerutils/include
         ${flashmla_SOURCE_DIR}/csrc/extension/sm90/dense_fp8/
-        ${flashmla_SOURCE_DIR}/csrc/cutlass/include
-        ${flashmla_SOURCE_DIR}/csrc/cutlass/tools/util/include
+        ${flashmla_SOURCE_DIR}/csrc/3rdparty/cutlass/include
+        ${flashmla_SOURCE_DIR}/csrc/3rdparty/cutlass/tools/util/include
     )
 
     set_gencode_flags_for_srcs(
@@ -202,6 +213,12 @@ if(FLASH_MLA_ARCHS)
         target_compile_definitions(_flashmla_C PRIVATE USE_CUDA)
     endif()
 
+    # Select the CUDA branch of FlashMLA's params.h and kerutils; kerutils
+    # detects CUDA by itself in .cu files, so only host files need the flag.
+    target_compile_definitions(_flashmla_C PRIVATE
+        FLASH_MLA_IS_BUILD_ON_CUDA
+        $<$<COMPILE_LANGUAGE:CXX>:KERUTILS_IS_BUILD_ON_CUDA>)
+
     define_extension_target(
         _flashmla_extension_C
         DESTINATION vllm

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=["."],

cmake/utils.cmake

@@ -425,6 +425,7 @@ function(cuda_archs_loose_intersection OUT_CUDA_ARCHS SRC_CUDA_ARCHS TGT_CUDA_AR
       if("${_base}" IN_LIST _TGT_CUDA_ARCHS)
         list(REMOVE_ITEM _SRC_CUDA_ARCHS "${_arch}")
         list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_base}")
+        list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_base}a" "${_base}f")
         list(APPEND _CUDA_ARCHS "${_arch}")
       endif()
     endif()
@@ -464,8 +465,10 @@ function(cuda_archs_loose_intersection OUT_CUDA_ARCHS SRC_CUDA_ARCHS TGT_CUDA_AR
       string(REGEX REPLACE "[af]$" "" _base "${_arch}")
       if ("${_base}" IN_LIST _SRC_CUDA_ARCHS)
         list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_arch}")
-        list(REMOVE_ITEM _SRC_CUDA_ARCHS "${_base}")
         list(APPEND _CUDA_ARCHS "${_arch}")
+        if (NOT "${_base}" IN_LIST _TGT_CUDA_ARCHS)
+          list(REMOVE_ITEM _SRC_CUDA_ARCHS "${_base}")
+        endif()
       endif()
     endif()
   endforeach()

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"
@@ -1059,33 +1059,6 @@ void situ_and_mul(torch::stable::Tensor& out,    // [..., d]
       });
 }
 
-void masked_situ_and_mul(torch::stable::Tensor& out,    // [E, T, d]
-                         torch::stable::Tensor& input,  // [E, T, 2 * d]
-                         const torch::stable::Tensor& expert_num_tokens,
-                         double beta, double linear_beta) {
-  int num_experts = input.size(0);
-  int max_num_tokens = input.size(1);
-  int d = input.size(2) / 2;
-  if (num_experts == 0 || max_num_tokens == 0 || d == 0) {
-    return;
-  }
-  constexpr int block_size = 256;
-  dim3 grid((d + block_size - 1) / block_size, num_experts);
-  dim3 block(block_size);
-  const torch::stable::accelerator::DeviceGuard device_guard(
-      input.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream();
-  VLLM_STABLE_DISPATCH_FLOATING_TYPES(
-      input.scalar_type(), "masked_situ_and_mul_kernel", [&] {
-        vllm::masked_situ_and_mul_kernel<scalar_t, true>
-            <<<grid, block, 0, stream>>>(
-                out.mutable_data_ptr<scalar_t>(),
-                input.const_data_ptr<scalar_t>(),
-                expert_num_tokens.const_data_ptr<int>(), max_num_tokens, d,
-                (float)beta, (float)linear_beta);
-      });
-}
-
 #define LAUNCH_MASKED_ACT_AND_MUL(KERNEL, ACT_FIRST, HAS_CLAMP, LIMIT, ALPHA, \
                                   BETA)                                       \
   vllm::masked_act_and_mul_kernel<scalar_t, KERNEL<scalar_t>, ACT_FIRST,      \

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 "attention_utils.cuh"
-#include "../../quantization/w8a8/fp8/common.cuh"
+#include "core/vec_dtypes/vec_dtypes.h"
+#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 "../cuda_compat.h"
+#include "core/utils.hpp"
+#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>
@@ -273,7 +273,7 @@ __global__ void reshape_and_cache_kernel(
     const int64_t* __restrict__ slot_mapping,  // [num_tokens]
     const int key_stride, const int value_stride, const int num_heads,
     const int head_size, const int block_size, const int x,
-    const float* k_scale, const float* v_scale) {
+    const float* k_scale, const float* v_scale, const bool interleaved_v) {
   const int64_t token_idx = blockIdx.x;
   const int64_t slot_idx = slot_mapping[token_idx];
   if (slot_idx < 0) {
@@ -301,10 +301,20 @@ __global__ void reshape_and_cache_kernel(
       key_cache + block_idx * num_heads * h_block_count * block_size * x +
       head_idx * h_block_count * block_size * x + h_block * block_size * x +
       block_offset * x;
-  const int64_t tgt_value_start =
-      block_idx * num_heads * h_block_count * x * block_size +
-      head_idx * h_block_count * x * block_size + h_block * x * block_size +
-      block_offset;
+  // V cache target: standard [.., head_size, block_size], or the interleaved
+  // [.., block_size/x, head_size, x] the fused AITER writer/reader use.
+  const int64_t v_head_base = block_idx * num_heads * head_size * block_size +
+                              head_idx * head_size * block_size;
+  int64_t tgt_value_start;
+  int v_dim_stride;
+  if (interleaved_v) {
+    tgt_value_start = v_head_base + (block_offset / x) * (head_size * x) +
+                      (h_block * x) * x + (block_offset % x);
+    v_dim_stride = x;
+  } else {
+    tgt_value_start = v_head_base + h_block * x * block_size + block_offset;
+    v_dim_stride = block_size;
+  }
 
   constexpr int VEC_SIZE = (sizeof(scalar_t) == 2) ? 8 : 4;
   float k_scale_val = (kv_dt == Fp8KVCacheDataType::kAuto) ? 0.f : *k_scale;
@@ -318,7 +328,7 @@ __global__ void reshape_and_cache_kernel(
   cache_t* __restrict__ value_dst = value_cache + tgt_value_start;
 #pragma unroll
   for (int i = 0; i < x; i++) {
-    v_op(value_dst[i * block_size], value_src[i]);
+    v_op(value_dst[i * v_dim_stride], value_src[i]);
   }
 }
 
@@ -630,7 +640,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 +690,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 +729,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));
-       iter++) {
+  for (int iter = 0; iter < div_ceil(batch_size, int(blockDim.x)); iter++) {
     int tid = iter * blockDim.x + threadIdx.x;
     if (tid < batch_size) {
       const int seq_start = cu_seq_lens[tid];
@@ -769,7 +785,7 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
           slot_mapping.const_data_ptr<int64_t>(), key_stride, value_stride, \
           num_heads, head_size, block_size, x,                              \
           reinterpret_cast<const float*>(k_scale.data_ptr()),               \
-          reinterpret_cast<const float*>(v_scale.data_ptr()));
+          reinterpret_cast<const float*>(v_scale.data_ptr()), interleaved_v);
 
 void reshape_and_cache(
     torch::stable::Tensor& key,    // [num_tokens, num_heads, head_size]
@@ -777,10 +793,12 @@ void reshape_and_cache(
     torch::stable::Tensor&
         key_cache,  // [num_blocks, num_heads, head_size/x, block_size, x]
     torch::stable::Tensor&
-        value_cache,  // [num_blocks, num_heads, head_size, block_size]
+        value_cache,  // [num_blocks, num_heads, head_size, block_size], or the
+                      // interleaved [.., block_size/x, head_size, x] when
+                      // interleaved_v is set
     torch::stable::Tensor& slot_mapping,  // [num_tokens]
     const std::string& kv_cache_dtype, torch::stable::Tensor& k_scale,
-    torch::stable::Tensor& v_scale) {
+    torch::stable::Tensor& v_scale, bool interleaved_v) {
   int num_tokens = slot_mapping.size(0);
   int num_heads = key.size(1);
   int head_size = key.size(2);
@@ -1229,7 +1247,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 +1426,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) +
-      2 * num_reqs;
+      div_ceil(static_cast<int32_t>(num_tokens), block_size) + 2 * num_reqs;
   const dim3 grid(required_blocks);
   const dim3 block(page_threads);
 
@@ -1685,7 +1702,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 +1845,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) +
-                              2 * static_cast<int32_t>(batch_size);
+  const int required_blocks =
+      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 +1931,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 +1943,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)