455 war room findings

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

@@ -5,21 +5,20 @@
 # Generate Buildkite annotation for ROCm wheel release
 set -ex
 
-# Extract build configuration from Dockerfile.rocm_base (single source of truth)
+# Get build configuration from meta-data
 # 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=//')
+PYTHON_VERSION=$(buildkite-agent meta-data get rocm-python-version 2>/dev/null || echo "3.12")
+PYTORCH_ROCM_ARCH=$(buildkite-agent meta-data get rocm-pytorch-rocm-arch 2>/dev/null || echo "gfx90a;gfx942;gfx950;gfx1250;gfx1100;gfx1101;gfx1200;gfx1201;gfx1150;gfx1151")
 
+# TODO: Enable the nightly build for ROCm
 # 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 "")
 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}"
@@ -69,7 +68,7 @@ aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/triton
 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}/aiter-*.whl .
 aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/flash-attn-*.whl .
 \`\`\`
 
@@ -81,7 +80,7 @@ aws s3 cp s3://${S3_BUCKET}/rocm/${BUILDKITE_COMMIT}/${ROCM_VERSION_PATH}/flash-
 - **torchvision**: TorchVision for ROCm PyTorch
 - **torchaudio**: Torchaudio for ROCm PyTorch
 - **amdsmi**: AMD SMI Python bindings
-- **amd_aiter**: Aiter for ROCm
+- **aiter**: Aiter for ROCm
 - **flash-attn**: Flash Attention for ROCm
 
 ### :warning: Notes
@@ -97,7 +96,7 @@ 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 public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-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

CMakeLists.txt

@@ -41,7 +41,7 @@ install(CODE "set(CMAKE_INSTALL_LOCAL_ONLY TRUE)" ALL_COMPONENTS)
 set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13" "3.14")
 
 # Supported AMD GPU architectures.
-set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
+set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1250;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
 
 # ROCm installation prefix. Default to /opt/rocm but allow override via
 # -DROCM_PATH=/your/rocm/path when invoking cmake.
@@ -1267,13 +1267,24 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
     "csrc/rocm/torch_bindings.cpp"
     "csrc/rocm/skinny_gemms.cu"
     "csrc/rocm/attention.cu")
+  set(VLLM_ROCM_EXT_FLAGS ${VLLM_GPU_FLAGS})
+
+  # skinny_gemms.cu is built on gfx9/gfx11 ISA (MFMA, dot2/dot4, legacy
+  # s_waitcnt asm) that gfx1250 (gfx12) does not provide. Exclude it from the
+  # gfx1250 build and disable its op registrations (VLLM_SKIP_SKINNY_GEMMS);
+  # vLLM falls back to default/Triton GEMM for those ops on gfx1250.
+  if(VLLM_GPU_ARCHES MATCHES "gfx1250")
+    message(STATUS "gfx1250 detected: excluding csrc/rocm/skinny_gemms.cu from _rocm_C")
+    list(REMOVE_ITEM VLLM_ROCM_EXT_SRC "csrc/rocm/skinny_gemms.cu")
+    list(APPEND VLLM_ROCM_EXT_FLAGS "-DVLLM_SKIP_SKINNY_GEMMS")
+  endif()
 
   define_extension_target(
     _rocm_C
     DESTINATION vllm
     LANGUAGE ${VLLM_GPU_LANG}
     SOURCES ${VLLM_ROCM_EXT_SRC}
-    COMPILE_FLAGS ${VLLM_GPU_FLAGS}
+    COMPILE_FLAGS ${VLLM_ROCM_EXT_FLAGS}
     ARCHITECTURES ${VLLM_GPU_ARCHES}
     USE_SABI 3
     WITH_SOABI)

benchmarks/curl_gpt.sh

@@ -0,0 +1,7 @@
+MODEL="${MODEL:-/data/models/gpt-oss-120b-w-mxfp4-a-fp8}" \
+curl -s http://localhost:8000/v1/completions -H "Content-Type: application/json" -d '{
+  "model": ${MODEL},
+  "prompt": "What is the capital of France, and what is it known for?",
+  "temperature": 0.0,
+  "max_tokens": 100
+}'

benchmarks/gpt_fp4_serve.sh

@@ -0,0 +1,7 @@
+#ROCR_VISIBLE_DEVICE=0 \
+MODEL="${MODEL:-/data/models/gpt-oss-120b-w-mxfp4-a-fp8}" \
+HSA_ENABLE_SDMA=0 USE_SVM=0 HSA_XNACK=0 \
+VLLM_ROCM_AITER_FUSED_MOE_TRITON_GEMM_A4W4=1 \
+VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_UNIFIED_ATTENTION=1 \
+VLLM_ROCM_USE_SKINNY_GEMM=0 VLLM_ROCM_USE_AITER_RMSNORM=0 \
+vllm serve --model ${MODEL} --host localhost --port 8000 --tensor-parallel-size 1 --gpu_memory_utilization 0.7 #--compilation-config '{"mode":"None","cudagraph_mode": "FULL", "cudagraph_capture_sizes": [1]}'

benchmarks/vllm_smoketest.sh

@@ -0,0 +1,103 @@
+#!/bin/bash
+set -euo pipefail
+
+# add model args here
+declare -A MODEL_PATHS=(
+  [gpt-oss-120b-mxfp4]="/data/amd/gpt-oss-120b-w-mxfp4-a-fp8"
+  [DeepSeek-R1-0528-MXFP4]="/data/amd/DeepSeek-R1-0528-MXFP4"
+)
+declare -A MODEL_SERVE_ARGS=(
+  [gpt-oss-120b-mxfp4]="--tensor-parallel-size 1 --gpu_memory_utilization 0.9 --attention-backend ROCM_AITER_UNIFIED_ATTN --compilation-config {\"mode\":\"None\",\"cudagraph_mode\":\"FULL\",\"cudagraph_capture_sizes\":[1]}"
+  [DeepSeek-R1-0528-MXFP4]="--tensor-parallel-size 1 --gpu_memory_utilization 0.9 --dtype auto --no-enable-prefix-caching --disable-uvicorn-access-log --trust-remote-code"
+)
+declare -A MODEL_ENV=(
+  [gpt-oss-120b-mxfp4]="HSA_ENABLE_SDMA=0 USE_SVM=0 HSA_XNACK=0 VLLM_ROCM_AITER_FUSED_MOE_TRITON_GEMM_A4W4=1 VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_UNIFIED_ATTENTION=1 VLLM_ROCM_USE_SKINNY_GEMM=0 VLLM_ROCM_USE_AITER_RMSNORM=0"
+  [DeepSeek-R1-0528-MXFP4]="HSA_ENABLE_SDMA=0 USE_SVM=0 HSA_XNACK=0 VLLM_ROCM_AITER_FUSED_MOE_TRITON_GEMM_A4W4=1 VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_UNIFIED_ATTENTION=1 VLLM_ROCM_USE_SKINNY_GEMM=0 VLLM_ROCM_USE_AITER_RMSNORM=0"
+)
+
+MODEL_NAME="gpt-oss-120b-mxfp4" # default
+MODEL_PATH_OVERRIDE=""         # overrided path
+MODEL_SET=0                    # whether --model was explicitly passed
+RUN_LM_EVAL=0
+PORT=8000
+
+usage() {
+  cat <<'EOF'
+Usage: ./vllm_smoketest.sh [--model NAME] [--model-path PATH] [--lm-eval] [--port PORT] [--list]
+  --model NAME       Registry key to run (default: gpt-oss-120b-mxfp4)
+  --model-path PATH  Override model path for this run
+  --lm-eval          Run lm_eval after curl check
+  --port PORT        Server port (default: 8000)
+  --list             List available models
+EOF
+}
+
+while [[ $# -gt 0 ]]; do
+  case "$1" in
+    --model)      MODEL_NAME="$2"; MODEL_SET=1; shift 2 ;;
+    --model-path) MODEL_PATH_OVERRIDE="$2"; shift 2 ;;
+    --lm-eval)    RUN_LM_EVAL=1; shift ;;
+    --port)       PORT="$2"; shift 2 ;;
+    --list)       for n in "${!MODEL_PATHS[@]}"; do echo "$n -> ${MODEL_PATHS[$n]}"; done; exit 0 ;;
+    -h|--help)    usage; exit 0 ;;
+    *) echo "Unknown argument: $1" >&2; usage; exit 1 ;;
+  esac
+done
+
+if [[ -n "$MODEL_PATH_OVERRIDE" && $MODEL_SET -eq 0 ]]; then
+  echo "ERROR: --model-path requires --model (which model's serve args/env to use)." >&2
+  usage; exit 1
+fi
+
+MODEL="${MODEL_PATHS[$MODEL_NAME]:-$MODEL_NAME}"
+[[ -n "$MODEL_PATH_OVERRIDE" ]] && MODEL="$MODEL_PATH_OVERRIDE"
+SERVE_ARGS="${MODEL_SERVE_ARGS[$MODEL_NAME]:-}"
+MODEL_ENV_ARGS="${MODEL_ENV[$MODEL_NAME]:-}"
+echo "Model: $MODEL | Port: $PORT | lm_eval: $RUN_LM_EVAL"
+
+# temp
+pip3 install "fastapi<0.137"
+
+LOG_DIR="$(pwd)/vllm_smoke_test_logs"
+mkdir -p "$LOG_DIR"
+echo "Logs: $LOG_DIR"
+
+server_pid=""
+cleanup() { [[ -n "$server_pid" ]] && kill "$server_pid" 2>/dev/null && wait "$server_pid" 2>/dev/null; true; }
+trap cleanup EXIT
+trap 'exit 130' INT TERM
+
+# --- start server ---
+env $MODEL_ENV_ARGS \
+vllm serve --model "$MODEL" --host localhost --port "$PORT" $SERVE_ARGS \
+  > "$LOG_DIR/server.log" 2>&1 &
+server_pid=$!
+echo "Waiting for server (pid $server_pid)..."
+
+until curl -s "http://localhost:$PORT/health" &>/dev/null; do
+  ps -p "$server_pid" >/dev/null || { echo "ERROR: server died"; tail -n 50 "$LOG_DIR/server.log"; exit 1; }
+  sleep 5
+done
+echo "Server ready."
+
+# --- check if text was generated ---
+echo "=== Curl completion check ==="
+http_code=$(curl -s -o "$LOG_DIR/curl_completion.log" -w '%{http_code}' \
+  "http://localhost:$PORT/v1/completions" \
+  -H "Content-Type: application/json" \
+  -d "{\"model\":\"$MODEL\",\"prompt\":\"San Francisco is a\",\"max_tokens\":32,\"temperature\":0}")
+cat "$LOG_DIR/curl_completion.log"; echo
+[[ "$http_code" == "200" ]] || { echo "FAIL: HTTP $http_code" >&2; exit 1; }
+grep -q '"text"' "$LOG_DIR/curl_completion.log" || { echo "FAIL: no text in response" >&2; exit 1; }
+echo "PASS"
+
+# --- lm_eval ---
+if [[ $RUN_LM_EVAL -eq 1 ]]; then
+  echo "=== lm_eval ==="
+  pip install "lm-eval[api]"
+  lm_eval --model local-completions \
+    --model_args "model=$MODEL,base_url=http://localhost:$PORT/v1/completions,num_concurrent=64,max_retries=3,tokenized_requests=False" \
+    --tasks gsm8k --num_fewshot 3 --limit 100 2>&1 | tee "$LOG_DIR/lm_eval.log"
+fi
+
+echo "Done."

csrc/quickreduce/base.h

@@ -81,11 +81,11 @@ union BufferResource {
 
 __quickreduce_device_inline__ static int32x4_t buffer_load_dwordx4(
     int32x4_t srsrc, int32_t voffset, int32_t soffset,
-    int32_t aux) __asm("llvm.amdgcn.raw.buffer.load.v4i32");
+    int32_t aux) {}
 
 __quickreduce_device_inline__ static void buffer_store_dwordx4(
     int32x4_t data, int32x4_t srsrc, int32_t voffset, int32_t soffset,
-    int32_t aux) __asm("llvm.amdgcn.raw.buffer.store.v4i32");
+    int32_t aux) {}
 
 __quickreduce_device_inline__ static void set_fp16_ovfl(bool const value) {
 #if defined(__gfx942__)

csrc/rocm/attention.cu

@@ -2405,13 +2405,23 @@ template <typename T, int absz, int cbid, int blgp>
 __device__ __forceinline__ floatx8 gcn_wmma16x16x16_instr(const bit16x8& inpA,
                                                           const bit16x8& inpB,
                                                           const floatx8& inpC) {
+#if defined(__gfx1250__)
+  // gfx1250 (gfx12 family) does not provide the gfx12 WMMA variant used by
+  // gfx1200/1201 (needs wmma-128b-insts). This custom-attention WMMA path is
+  // unsupported on gfx1250; trap if ever launched (fail loud, not silent-wrong).
+  (void)inpA;
+  (void)inpB;
+  __builtin_trap();
+  return inpC;
+#else
   if constexpr (std::is_same<T, _Float16>::value) {
     return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(inpA, inpB, inpC);
   } else if constexpr (std::is_same<T, __hip_bfloat16>::value) {
     return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(inpA, inpB, inpC);
   } else {
     static_assert(false, "unsupported 16b dtype");
   }
+#endif
 }
 
 template <typename T>

csrc/rocm/torch_bindings.cpp

@@ -14,6 +14,10 @@
 TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, rocm_ops) {
   // vLLM custom ops for rocm
 
+// skinny_gemms.cu (LLMM1/wvSplitK/wvSplitKrc/wvSplitKQ) is excluded on gfx1250
+// (gfx9/gfx11 ISA, unsupported there); skip these registrations to avoid
+// undefined symbols. vLLM uses default/Triton GEMM for these ops on gfx1250.
+#ifndef VLLM_SKIP_SKINNY_GEMMS
   // Custom gemm op for matrix-vector multiplication
   rocm_ops.def(
       "LLMM1(Tensor in_a, Tensor in_b, int rows_per_block) -> "
@@ -38,6 +42,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, rocm_ops) {
       "Tensor scale_a, "
       "          Tensor scale_b, int CuCount) -> ()");
   rocm_ops.impl("wvSplitKQ", torch::kCUDA, &wvSplitKQ);
+#endif  // VLLM_SKIP_SKINNY_GEMMS
 
   // Custom attention op
   // Compute the attention between an input query and the cached