Add A4W16 GPT OSS MoE routing for gfx1250

.agents/skills/triton-kernel-writing/SKILL.md

@@ -22,10 +22,14 @@ description: Write or review Triton kernels for vLLM, with practical guidance fo
   a small, legible heuristic when workloads need different choices. Use
   `triton.autotune` only when tuning is critical to performance, such as for a
   matrix multiplication. Otherwise prioritize simple code and fast startup.
-- Be careful to avoid unintended runtime JIT compilation. For example, put
-  unimportant runtime integer scalars in `do_not_specialize`, especially those
-  that may alternate between values such as 0 and 1, which can produce
-  different specialization keys.
+- Avoid unintended runtime JIT compilation. Put unimportant runtime integer
+  scalars in `do_not_specialize`, especially values that alternate between 0
+  and 1. Pointer alignment also affects specialization: an offset tensor view
+  can trigger a new variant after warmup used an aligned allocation. For
+  scalar-only pointers whose alignment provides no useful optimization, use
+  `@triton.jit(do_not_specialize_on_alignment=["end_ptr"])`. Retain alignment
+  specialization when vectorized accesses benefit from it. See the
+  [Triton JIT options](https://triton-lang.org/main/python-api/generated/triton.jit.html).
 - The Triton compiler does not guarantee safe ordering when a kernel writes to
   a pointer and subsequently reads from the same pointer. This pattern must
   have a `tl.debug_barrier()` between the write and read. The barrier
@@ -49,7 +53,9 @@ description: Write or review Triton kernels for vLLM, with practical guidance fo
 ## Validation
 
 - Check correctness at boundary shapes and at sizes that exercise masks and
-  large offsets.
+  large offsets. Include offset tensor views, not only fresh allocations.
+  Warm aligned buffers, then exercise offset views with vLLM's
+  `--jit-monitor-mode error --jit-monitor-verbose` to catch missing variants.
 - Choose accumulation and intermediate dtypes explicitly. Test numerically
   difficult inputs, not only random, well-scaled tensors.
 - Use `$kernel-microbenchmark` for benchmark construction, measurement, and

.buildkite/hardware_tests/cpu.yaml

@@ -52,6 +52,52 @@ steps:
       tests/v1/sample/test_topk_topp_sampler.py::TestTritonTopkTopp"
   parallelism: 2
 
+# Speculative decoding under a CPU platform. The unit tests already pass on
+# CPU but no CPU job runs them; the e2e tests gate the engine path itself,
+# which unit tests cannot reach.
+- label: CPU-Spec Decode Tests
+  key: cpu-spec-decode-tests
+  depends_on: []
+  soft_fail: true
+  device: intel_cpu
+  no_plugin: true
+  source_file_dependencies:
+  - vllm/v1/spec_decode/
+  - vllm/v1/worker/cpu/
+  - vllm/v1/worker/cpu_model_runner.py
+  - vllm/v1/worker/gpu/spec_decode/
+  - vllm/v1/attention/backends/cpu_attn.py
+  - vllm/config/speculative.py
+  - vllm/model_executor/models/qwen3_5.py
+  - vllm/model_executor/models/qwen3_5_mtp.py
+  - vllm/model_executor/models/llama_eagle3.py
+  - vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py
+  - vllm/model_executor/layers/mamba/mamba_mixer2.py
+  - vllm/model_executor/layers/mamba/ops/cpu/
+  - tests/v1/e2e/test_cpu_spec_decode.py
+  - tests/v1/spec_decode/
+  - tests/v1/worker/
+  commands:
+    - |
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
+      pytest -x -v -s \
+      tests/v1/worker/test_gpu_autoregressive_speculator.py \
+      tests/v1/worker/test_gpu_extract_hidden_states_speculator.py \
+      tests/v1/worker/test_gpu_batch_ordering.py \
+      tests/v1/worker/test_gpu_rejection_sampler_chunking.py \
+      tests/v1/worker/test_eagle3_aux_hidden_states_pp.py \
+      tests/v1/worker/test_spec_decode_embed_sharing_pp.py \
+      tests/v1/spec_decode/test_adaptive_verification.py \
+      tests/v1/spec_decode/test_dflash2.py \
+      tests/v1/spec_decode/test_dflash_causality.py \
+      tests/v1/spec_decode/test_eagle_draft_attn_metadata.py \
+      tests/v1/spec_decode/test_draft_attention_backend_override.py \
+      tests/v1/spec_decode/test_draft_moe_backend_override.py \
+      tests/v1/spec_decode/test_ngram.py \
+      tests/v1/spec_decode/test_llm_base_proposer_sampling.py
+      unset VLLM_CPU_CI_ENV
+      pytest -x -v -s tests/v1/e2e/test_cpu_spec_decode.py"
+
 # Note: SDE can't be downloaded from CI host because of AWS WAF
 # - label: CPU-Compatibility Tests
 #   depends_on: []

.buildkite/image_build/image_build.sh

@@ -79,6 +79,54 @@ setup_buildx_builder() {
     docker buildx ls | grep -E '^\*|^NAME' || docker buildx ls
 }
 
+download_ci_hcl() {
+    # Returns non-zero instead of exiting: the optional export path must be
+    # able to give up without failing a build whose image already exists.
+    echo "--- :arrow_down: Downloading ci.hcl"
+    curl -sSfL -o "${CI_HCL_PATH}" "${CI_HCL_URL}" || return $?
+    echo "Downloaded to ${CI_HCL_PATH}"
+
+    if [[ ! -f "${CI_HCL_PATH}" ]]; then
+        echo "Error: ci.hcl not found at ${CI_HCL_PATH}"
+        return 1
+    fi
+}
+
+export_kernel_symbol_map_from_cache() {
+    # The existing-image path: everything here is best effort. Called as
+    # `... || echo`, so errexit is suspended inside and any failure just
+    # returns; the build keeps the image it already has.
+    local rc=0
+    download_ci_hcl \
+        && setup_buildx_builder \
+        && resolve_parent_commit \
+        && export PARENT_COMMIT \
+        && BUILD_TMP_DIR="$(mktemp -d)" \
+        && export_kernel_symbol_map || rc=$?
+    rm -rf -- "${BUILD_TMP_DIR:-}"
+    return "${rc}"
+}
+
+export_kernel_symbol_map() {
+    # Only when the build asked for it (nightly/post-merge). The map was
+    # produced inside the csrc-build stage during the main bake; this second
+    # bake of the scratch stage is a cache hit that just writes the file out.
+    if [[ "${VLLM_KERNEL_SYMBOL_MAP:-0}" != "1" ]]; then
+        return 0
+    fi
+    echo "--- :world_map: Exporting kernel symbol map"
+    local out_dir="${BUILD_TMP_DIR}/kernel-symbol-map"
+    mkdir -p "${out_dir}"
+    if docker buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --progress plain \
+        --set "kernel-symbol-map.output=type=local,dest=${out_dir}" kernel-symbol-map \
+        && [[ -s "${out_dir}/kernel_symbol_map.json.gz" ]]; then
+        ls -la "${out_dir}"
+        (cd "${out_dir}" && buildkite-agent artifact upload "kernel_symbol_map.json.gz")
+    else
+        echo "kernel symbol map export failed; continuing without it" >&2
+    fi
+}
+
 annotate_image_tags() {
     .buildkite/scripts/annotate-image-build.sh \
         "${IMAGE_TAG:-}" "${IMAGE_TAG_LATEST:-}"
@@ -90,6 +138,15 @@ check_and_skip_if_image_exists() {
         if docker manifest inspect "${IMAGE_TAG}" >/dev/null 2>&1; then
             echo "Image already exists: ${IMAGE_TAG}"
             echo "Skipping build"
+            if [[ "${VLLM_KERNEL_SYMBOL_MAP:-0}" == "1" ]]; then
+                # The image is reused, but the symbol map is an artifact of
+                # this build. Bake just the export target: its layers come
+                # from the registry cache, and the map step itself reruns
+                # if the cached image was built without the arg. Never
+                # fails the build: the image is already there.
+                export_kernel_symbol_map_from_cache \
+                    || echo "kernel symbol map: export from the cached image failed; continuing" >&2
+            fi
             annotate_image_tags
             exit 0
         fi
@@ -154,7 +211,7 @@ print_bake_config() {
     local bake_tmp
     bake_tmp="$(mktemp -d)"
     BAKE_CONFIG_FILE="${bake_tmp}/bake-config-build-${BUILDKITE_BUILD_NUMBER:-local}.json"
-    docker buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --print "${TARGET}" | tee "${BAKE_CONFIG_FILE}" || true
+    docker buildx bake "${BAKE_FILES[@]}" --print "${TARGET}" | tee "${BAKE_CONFIG_FILE}" || true
     echo "Saved bake config to ${BAKE_CONFIG_FILE}"
     echo "--- :arrow_down: Uploading bake config to Buildkite"
     (cd "$(dirname "${BAKE_CONFIG_FILE}")" && buildkite-agent artifact upload "$(basename "${BAKE_CONFIG_FILE}")")
@@ -230,8 +287,10 @@ TARGET="test-ci"
 VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"
 BUILDER_NAME="${BUILDER_NAME:-vllm-builder}"
 CI_HCL_URL="${CI_HCL_URL:-https://raw.githubusercontent.com/vllm-project/ci-infra/main/docker/ci.hcl}"
-CI_HCL_PATH="/tmp/ci.hcl"
+CI_HCL_PATH="${CI_HCL_PATH:-/tmp/ci.hcl}"
+ZSTD_HCL_PATH="${ZSTD_HCL_PATH:-.buildkite/image_build/zstd.hcl}"
 BUILDKIT_SOCKET="/run/buildkit/buildkitd.sock"
+BAKE_FILES=(-f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" -f "${ZSTD_HCL_PATH}")
 
 prepare_cache_tags
 ecr_login
@@ -288,14 +347,7 @@ if [[ ! -f "${VLLM_BAKE_FILE_PATH}" ]]; then
     exit 1
 fi
 
-echo "--- :arrow_down: Downloading ci.hcl"
-curl -sSfL -o "${CI_HCL_PATH}" "${CI_HCL_URL}"
-echo "Downloaded to ${CI_HCL_PATH}"
-
-if [[ ! -f "${CI_HCL_PATH}" ]]; then
-    echo "Error: ci.hcl not found at ${CI_HCL_PATH}"
-    exit 1
-fi
+download_ci_hcl || exit 1
 
 setup_buildx_builder
 
@@ -309,7 +361,7 @@ BUILD_TMP_DIR="$(mktemp -d)"
 trap 'rm -rf -- "${BUILD_TMP_DIR}"' EXIT
 BUILD_METADATA_FILE="${BUILD_TMP_DIR}/build-metadata.json"
 BUILD_STATUS=0
-docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" \
+docker --debug buildx bake "${BAKE_FILES[@]}" \
     --progress plain --metadata-file "${BUILD_METADATA_FILE}" "${TARGET}" || BUILD_STATUS=$?
 
 record_buildkit_trace "${BUILD_METADATA_FILE}" || true
@@ -319,5 +371,6 @@ if [[ "${BUILD_STATUS}" -ne 0 ]]; then
 fi
 
 echo "--- :white_check_mark: Build complete"
+export_kernel_symbol_map || true
 
 annotate_image_tags

.buildkite/image_build/image_build_torch_nightly.sh

@@ -49,7 +49,16 @@ echo "Image not found, proceeding with build..."
 # Nightly CI uses CUDA 13.0 while regular CI stays on CUDA 12.9
 NIGHTLY_CUDA_VERSION="13.0.2"
 NIGHTLY_BUILD_BASE_IMAGE="pytorch/manylinux2_28-builder:cuda13.0-78e737ad29420ffc4800e677c51e2a852caf8359"
-NIGHTLY_FINAL_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-base-ubuntu22.04"
+# Must track ARG UBUNTU_VERSION in docker/Dockerfile. The nightly lane exists to
+# run the normal pipeline against nightly torch, so it has to differ from the
+# regular image in torch only -- not in the OS underneath it. This was pinned to
+# 22.04 while everything else moved to 24.04, and the runtime dependency list in
+# tools/install_snapshot_runtime.sh is 24.04-only (libgnutls30t64 is the 64-bit
+# time_t rename and does not exist on jammy), so the nightly build failed with
+# "E: Unable to locate package libgnutls30t64" while regular CI stayed green on
+# a cached layer.
+NIGHTLY_UBUNTU_VERSION="24.04"
+NIGHTLY_FINAL_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-base-ubuntu${NIGHTLY_UBUNTU_VERSION}"
 
 echo "--- :docker: Building torch nightly image (CUDA ${NIGHTLY_CUDA_VERSION})"
 docker buildx build --file docker/Dockerfile \

.buildkite/image_build/zstd.hcl

@@ -0,0 +1,3 @@
+target "test-ci" {
+  output = ["type=registry,compression=zstd,compression-level=3,force-compression=true,oci-mediatypes=true"]
+}

.buildkite/intel_jobs/basic_correctness_intel.yaml

@@ -3,6 +3,7 @@ depends_on:
   - image-build-xpu
 steps:
 - label: XPU Sleep Mode
+  key: xpu-sleep-mode
   timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:

.buildkite/intel_jobs/engine_intel.yaml

@@ -41,6 +41,7 @@ steps:
       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)
+  key: engine-1-gpu
   timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:
@@ -62,9 +63,10 @@ steps:
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
       pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py &&
-      VLLM_XPU_ENABLE_XPU_GRAPH=1 pytest -v -s test_config.py'
+      pytest -v -s test_config.py'
 
 - label: V1 e2e (2 GPUs)
+  key: v1-e2e-2-gpus
   timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:

.buildkite/intel_jobs/entrypoints_intel.yaml

@@ -134,7 +134,8 @@ steps:
     'cd tests &&
     pytest -v -s tool_use
     --deselect="tests/tool_use/mistral/test_mistral_tool_calls.py::test_tool_call_none_with_tools[ministral-3b]"
-    --deselect="tests/tool_use/test_chat_completions.py::test_chat_completion_with_tools[granite-3.0-8b]" &&
+    --deselect="tests/tool_use/test_chat_completions.py::test_chat_completion_with_tools[granite-3.0-8b]"
+    --deselect="tests/tool_use/test_chat_completions.py::test_chat_completion_with_tools[llama3.2]" &&
     pytest -v -s entrypoints/tool_parsers &&
     pytest -v -s entrypoints/generate &&
     pytest -v -s entrypoints/anthropic &&

.buildkite/intel_jobs/kernels_intel.yaml

@@ -3,6 +3,7 @@ depends_on:
   - image-build-xpu
 steps:
 - label: vLLM IR Tests
+  key: vllm-ir-tests
   timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:

.buildkite/intel_jobs/lm_eval_intel.yaml

@@ -23,6 +23,30 @@ steps:
   commands:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install "gpt-oss[eval]==0.0.5" &&
+      'uv pip install "gpt-oss[eval]==0.0.5" &&
       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)
+  key: lm-eval-large-models-2-gpus
+  timeout_in_minutes: 45
+  device: intel_gpu
+  agent_tags:
+    label: production
+    gpu: 2+
+    mem: 24+
+  no_plugin: true
+  optional: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+    - vllm/
+    - .buildkite/lm-eval-harness/
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'cd .buildkite/lm-eval-harness &&
+       export VLLM_USE_DEEP_GEMM=0 &&
+      pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-xpu-fp8.txt --tp-size=2'

.buildkite/intel_jobs/lora_intel.yaml

@@ -3,6 +3,7 @@ depends_on:
   - image-build-xpu
 steps:
 - label: LoRA Runtime + Utils
+  key: lora-runtime-utils
   timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
@@ -36,6 +37,7 @@ steps:
       pytest -v -s lora/test_worker.py'
 
 - label: LoRA Fused/MoE Kernels
+  key: lora-fused-moe-kernels
   timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
@@ -60,6 +62,7 @@ steps:
       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
+  key: lora-punica-kernels
   timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
@@ -86,6 +89,7 @@ steps:
       pytest -v -s lora/test_punica_ops.py::test_add_lora_fused_moe_early_exit'
 
 - label: LoRA Punica FP8/XPU Ops
+  key: lora-punica-fp8-xpu-ops
   timeout_in_minutes: 60
   device: intel_gpu
   agent_tags:
@@ -110,6 +114,7 @@ steps:
       pytest -v -s lora/test_punica_xpu_ops.py'
 
 - label: LoRA Models
+  key: lora-models
   timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
@@ -137,6 +142,7 @@ steps:
       pytest -s -v lora/test_minicpmv_tp.py'
 
 - label: LoRA Multimodal
+  key: lora-multimodal
   timeout_in_minutes: 45
   device: intel_gpu
   agent_tags: