[ROCM] Add rocjitsu e2e validation (#24913)
I want to enable e2e testing across all recent amdgpu / rcom targets via
rocjitsu, an amdgpu simulator, so that anyone with an x86 machine can
run those tests regardless of the local GPU or lack thereof.
Add a CPU-hosted ROCm e2e test helper for gfx942, gfx950, gfx1100,
gfx1201, and gfx1250. Reconfigure a regular IREE build and rebuild
iree-test-deps for each target, then run HIP tests with the simulator
from a pinned TheRock nightly.
Run each test in a fresh Mirage session with a watchdog and per-test
JUnit output. Propagate configure/build failures so stale artifacts
cannot produce a successful run.
Make the default native-test timeout configurable and label tests that
are too expensive to simulate. Keep those tests enabled for hardware
runs. Exclude TopK v2 from ROCm tests because its missing thread
distribution causes overlapping output writes on all five targets.
Use the target's preferred subgroup size when an executable does not
specify one, so gfx1250 defaults to wave32. Add regression coverage for
the default and explicit subgroup sizes.
Assisted-by: codex (gpt-6-astra)
diff --git a/.github/workflows/ci_linux_x64_clang.yml b/.github/workflows/ci_linux_x64_clang.yml
index 96dd2e7..42092df 100644
--- a/.github/workflows/ci_linux_x64_clang.yml
+++ b/.github/workflows/ci_linux_x64_clang.yml
@@ -37,9 +37,42 @@
run: |
source ./build_tools/cmake/setup_sccache.sh
./build_tools/cmake/build_all.sh "${BUILD_DIR}"
+ env:
+ # Build the first rocjitsu target's test deps along with everything else.
+ IREE_ROCM_TEST_TARGET_CHIP: gfx942
- name: Run CTest
run: ./build_tools/cmake/ctest_all.sh "${BUILD_DIR}"
env:
CTEST_PARALLEL_LEVEL: 32
+ - name: Install TheRock SDK with rocjitsu
+ run: |
+ sdk_venv="${BUILD_DIR}/rocjitsu/therock-venv"
+ python3 -m venv "${sdk_venv}"
+ "${sdk_venv}/bin/python" -m pip install --pre \
+ --index-url https://nightly.repo.amd.com/rocm/whl-next/ \
+ "rocm[libraries,devel]==10.1.0a20260909"
+ "${sdk_venv}/bin/rocm-sdk" init
+ "${sdk_venv}/bin/rocm-sdk" version
+ rocm_root="$("${sdk_venv}/bin/rocm-sdk" path --root)"
+ echo "ROCM_ROOT=${rocm_root}" >> "${GITHUB_ENV}"
+ echo "ROCJITSU_BIN=${rocm_root}/bin/mirage" >> "${GITHUB_ENV}"
+
+ - name: Run ROCm e2e tests under rocjitsu
+ timeout-minutes: 120
+ run: |
+ ./build_tools/cmake/test_rocm_targets_with_rocjitsu.sh \
+ "${BUILD_DIR}" gfx942 gfx950 gfx1100 gfx1201 gfx1250
+ env:
+ IREE_ROCJITSU_RUN_ID: ci
+ CTEST_PARALLEL_LEVEL: 8
+
+ - name: Upload rocjitsu test results
+ if: always()
+ uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7.0.1
+ with:
+ name: rocjitsu-linux-x64-clang
+ path: ${{ env.BUILD_DIR }}/rocjitsu/logs/ci/
+ if-no-files-found: ignore
+
# Alerting on failure is the responsibility of the calling job.
diff --git a/CMakeLists.txt b/CMakeLists.txt
index d940c82..f59326f 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -79,6 +79,8 @@
option(IREE_BUILD_COMPILER "Builds the IREE compiler." ON)
option(IREE_BUILD_TESTS "Builds IREE unit tests." ON)
+set(IREE_NATIVE_TEST_TIMEOUT_DEFAULT 60 CACHE STRING
+ "Default timeout in seconds for native tests")
option(IREE_BUILD_DOCS "Builds IREE documentation files." OFF)
option(IREE_BUILD_SAMPLES "Builds IREE sample projects." ON)
option(IREE_BUILD_PYTHON_BINDINGS "Builds the IREE python bindings" OFF)
diff --git a/build_tools/cmake/build_all.sh b/build_tools/cmake/build_all.sh
index c427be7..44a93a3 100755
--- a/build_tools/cmake/build_all.sh
+++ b/build_tools/cmake/build_all.sh
@@ -76,6 +76,12 @@
"-DIREE_TARGET_BACKEND_WEBGPU_SPIRV=${IREE_TARGET_BACKEND_WEBGPU_SPIRV}"
)
+# Include the first simulator target's test modules in the initial build.
+# Leave existing cache settings unchanged when no target is requested.
+if [[ -n "${IREE_ROCM_TEST_TARGET_CHIP:-}" ]]; then
+ CMAKE_ARGS+=("-DIREE_ROCM_TEST_TARGET_CHIP=${IREE_ROCM_TEST_TARGET_CHIP}")
+fi
+
# Force /Z7 (embedded per-.obj debug info) instead of /Zi (shared per-target
# .pdb serialized through mspdbsrv.exe via /FS) on Windows.
#
diff --git a/build_tools/cmake/iree_native_test.cmake b/build_tools/cmake/iree_native_test.cmake
index 082ba05..60ee858 100644
--- a/build_tools/cmake/iree_native_test.cmake
+++ b/build_tools/cmake/iree_native_test.cmake
@@ -176,7 +176,7 @@
endif()
if (NOT DEFINED _RULE_TIMEOUT OR "${_RULE_TIMEOUT}" STREQUAL "")
- set(_RULE_TIMEOUT 60)
+ set(_RULE_TIMEOUT ${IREE_NATIVE_TEST_TIMEOUT_DEFAULT})
endif()
list(APPEND _RULE_LABELS "${_PACKAGE_PATH}")
diff --git a/build_tools/cmake/test_rocm_targets_with_rocjitsu.sh b/build_tools/cmake/test_rocm_targets_with_rocjitsu.sh
new file mode 100755
index 0000000..4c92f15
--- /dev/null
+++ b/build_tools/cmake/test_rocm_targets_with_rocjitsu.sh
@@ -0,0 +1,296 @@
+#!/usr/bin/env bash
+
+# Copyright 2026 The IREE Authors
+#
+# Licensed under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+# Reconfigures a regular IREE build for one or more ROCm test targets and runs
+# its HIP e2e tests under rocjitsu. Expects to be run from the root of the IREE
+# repository after build_tools/cmake/build_all.sh has populated BUILD_DIR.
+#
+# Requires an initialized ROCm SDK containing rocjitsu. Install the
+# SDK separately and set ROCM_ROOT; this script does not install dependencies.
+# Set IREE_ROCM_TEST_TARGET_CHIP to the first target when running build_all.sh
+# to include its test modules in the initial build.
+
+set -euo pipefail
+
+SOURCE_DIR="$(pwd)"
+BUILD_DIR="${1:-${IREE_BUILD_DIR:-build}}"
+if (($#)); then
+ shift
+fi
+
+if [[ ! -f "${SOURCE_DIR}/CMakeLists.txt" ||
+ ! -x "${SOURCE_DIR}/build_tools/cmake/build_all.sh" ]]; then
+ echo "error: run this script from the root of the IREE repository" >&2
+ exit 2
+fi
+
+readonly -a ALL_TARGETS=(gfx942 gfx950 gfx1100 gfx1201 gfx1250)
+declare -Ar ROCJITSU_CONFIGS=(
+ [gfx942]="gfx942_cdna3.json"
+ [gfx950]="gfx950_mi355x.json"
+ [gfx1100]="gfx1100_w7900.json"
+ [gfx1201]="gfx1201_r9700.json"
+ [gfx1250]="gfx1250_mi455x.json"
+)
+declare -Ar INCOMPATIBLE_GPU_LABELS=(
+ [gfx942]='^requires-gpu-(cdna4|rdna3|rdna4|gfx1250)$'
+ [gfx950]='^requires-gpu-(cdna3|rdna3|rdna4|gfx1250)$'
+ [gfx1100]='^requires-gpu-(cdna3|cdna4|rdna4|gfx1250)$'
+ [gfx1201]='^requires-gpu-(cdna3|cdna4|rdna3|gfx1250)$'
+ [gfx1250]='^requires-gpu-(cdna3|cdna4|rdna3|rdna4|wave64)$'
+)
+
+usage() {
+ cat <<'EOF'
+Usage: build_tools/cmake/test_rocm_targets_with_rocjitsu.sh [BUILD_DIR] [TARGET ...]
+
+Reconfigures an existing regular IREE build and runs its ROCm e2e tests under
+rocjitsu. Supported targets:
+ gfx942 gfx950 gfx1100 gfx1201 gfx1250
+
+BUILD_DIR defaults to IREE_BUILD_DIR or "build". With no target arguments, all
+supported targets are built and tested sequentially in the same build tree.
+
+Requires an initialized ROCm SDK containing rocjitsu. Install the
+SDK separately, run rocm-sdk init, and set ROCM_ROOT to its root directory.
+Set IREE_ROCM_TEST_TARGET_CHIP to the first target when running build_all.sh
+to include its test modules in the initial build.
+
+Runs one parallel CTest invocation per target. CTest prints test progress and
+output on failure; the target's log directory contains the log and JUnit report.
+Each test process uses an independent simulator so concurrent tests do not
+contend for a shared emulated GPU.
+
+Configuration:
+ IREE_ROCJITSU_WORK_DIR Logs and simulator runtime state.
+ IREE_ROCJITSU_RUN_ID Override the log and simulator runtime run ID.
+ ROCM_ROOT Required: root of an initialized ROCm SDK.
+ ROCJITSU_BIN Override the rocjitsu CLI executable.
+ ROCJITSU_CONFIG_DIR Override the rocjitsu config directory.
+ CTEST_PARALLEL_LEVEL Test concurrency (default: 8).
+ IREE_ROCJITSU_TEST_TIMEOUT Default timeout in seconds for tests without an
+ explicit CMake TIMEOUT value (default: 180).
+ IREE_ROCJITSU_SESSION_TIMEOUT
+ Whole-session watchdog in seconds (default: 1800).
+ IREE_ROCJITSU_TESTS_REGEX Include only tests matching this CTest regex.
+ IREE_ROCJITSU_EXCLUDE_TESTS_REGEX
+ Exclude tests matching this CTest regex.
+ IREE_ROCJITSU_EXCLUDE_LABEL_REGEX
+ Exclude tests with labels matching this regex
+ (default: ^very-expensive$; use ^$ to disable).
+
+Examples:
+ export ROCM_ROOT="$(rocm-sdk path --root)"
+ build_tools/cmake/test_rocm_targets_with_rocjitsu.sh build gfx1250
+ build_tools/cmake/test_rocm_targets_with_rocjitsu.sh build
+ ROCM_ROOT=/path/to/rocm-sdk \
+ build_tools/cmake/test_rocm_targets_with_rocjitsu.sh build gfx942 gfx950
+EOF
+}
+
+if [[ "${BUILD_DIR}" == "--help" || "${BUILD_DIR}" == "-h" ]]; then
+ usage
+ exit 0
+fi
+
+if [[ -z "${ROCM_ROOT:-}" ]]; then
+ echo "error: set ROCM_ROOT to an initialized ROCm SDK containing rocjitsu" >&2
+ exit 2
+fi
+
+if [[ ! -f "${BUILD_DIR}/CMakeCache.txt" ]]; then
+ echo "error: ${BUILD_DIR} is not a configured IREE build directory" >&2
+ echo "Run build_tools/cmake/build_all.sh ${BUILD_DIR} first." >&2
+ exit 2
+fi
+BUILD_DIR="$(realpath "${BUILD_DIR}")"
+
+WORK_DIR="${IREE_ROCJITSU_WORK_DIR:-${BUILD_DIR}/rocjitsu}"
+CTEST_PARALLEL_LEVEL="${CTEST_PARALLEL_LEVEL:-8}"
+TEST_TIMEOUT="${IREE_ROCJITSU_TEST_TIMEOUT:-180}"
+SESSION_TIMEOUT="${IREE_ROCJITSU_SESSION_TIMEOUT:-1800}"
+EXCLUDE_LABEL_REGEX="${IREE_ROCJITSU_EXCLUDE_LABEL_REGEX:-^very-expensive$}"
+RUN_ID="${IREE_ROCJITSU_RUN_ID:-$(date +%Y%m%d-%H%M%S)-$$}"
+
+if (($#)); then
+ TARGETS=("$@")
+else
+ TARGETS=("${ALL_TARGETS[@]}")
+fi
+
+for target in "${TARGETS[@]}"; do
+ if [[ -z "${ROCJITSU_CONFIGS[${target}]+x}" ]]; then
+ echo "error: unsupported target '${target}'" >&2
+ usage >&2
+ exit 2
+ fi
+done
+
+declare -a CTEST_BASE_FILTER_ARGS=(
+ -L '^driver=hip$'
+ -L '^iree/tests/e2e/'
+)
+if [[ -n "${IREE_ROCJITSU_TESTS_REGEX:-}" ]]; then
+ CTEST_BASE_FILTER_ARGS+=(-R "${IREE_ROCJITSU_TESTS_REGEX}")
+fi
+if [[ -n "${IREE_ROCJITSU_EXCLUDE_TESTS_REGEX:-}" ]]; then
+ CTEST_BASE_FILTER_ARGS+=(-E "${IREE_ROCJITSU_EXCLUDE_TESTS_REGEX}")
+fi
+
+require_cache_setting() {
+ local setting="$1"
+ if ! grep -q "^${setting}:BOOL=ON$" "${BUILD_DIR}/CMakeCache.txt"; then
+ echo "error: regular build must have ${setting}=ON" >&2
+ exit 2
+ fi
+}
+
+require_cache_setting IREE_BUILD_TESTS
+require_cache_setting IREE_BUILD_COMPILER
+require_cache_setting IREE_HAL_DRIVER_HIP
+require_cache_setting IREE_TARGET_BACKEND_ROCM
+if ! grep -q '^CMAKE_GENERATOR:INTERNAL=Ninja$' "${BUILD_DIR}/CMakeCache.txt"; then
+ echo "error: this script expects the regular CI Ninja build" >&2
+ exit 2
+fi
+
+mkdir -p "${WORK_DIR}"
+
+ROCM_ROOT="$(realpath "${ROCM_ROOT}")"
+export ROCM_ROOT
+
+if [[ -z "${ROCJITSU_BIN:-}" ]]; then
+ if [[ -x "${ROCM_ROOT}/bin/rocjitsu" ]]; then
+ ROCJITSU_BIN="${ROCM_ROOT}/bin/rocjitsu"
+ else
+ ROCJITSU_BIN="${ROCM_ROOT}/bin/mirage"
+ fi
+fi
+ROCJITSU_CONFIG_DIR="${ROCJITSU_CONFIG_DIR:-${ROCM_ROOT}/share/rocjitsu/configs}"
+if [[ ! -x "${ROCJITSU_BIN}" ]]; then
+ echo "error: rocjitsu CLI not found at ${ROCJITSU_BIN}" >&2
+ echo "Install a nightly that includes a simulator CLI or set ROCJITSU_BIN." >&2
+ exit 2
+fi
+case "$(basename "${ROCJITSU_BIN}")" in
+ mirage) ROCJITSU_CLI_KIND="session" ;;
+ *) ROCJITSU_CLI_KIND="legacy" ;;
+esac
+if [[ ! -f "${ROCM_ROOT}/lib/libamdhip64.so" ]]; then
+ echo "error: HIP runtime not found under ${ROCM_ROOT}/lib" >&2
+ exit 2
+fi
+
+echo "IREE build: ${BUILD_DIR}"
+echo "ROCm SDK root: ${ROCM_ROOT}"
+echo "Simulator CLI: ${ROCJITSU_BIN} ($("${ROCJITSU_BIN}" --version))"
+echo "Targets: ${TARGETS[*]}"
+
+declare -a passed_targets=()
+declare -a build_failed_targets=()
+declare -a test_failed_targets=()
+
+for target in "${TARGETS[@]}"; do
+ config_path="${ROCJITSU_CONFIG_DIR}/${ROCJITSU_CONFIGS[${target}]}"
+ runtime_dir="${WORK_DIR}/runtime/${RUN_ID}/${target}"
+ log_dir="${WORK_DIR}/logs/${RUN_ID}/${target}"
+ mkdir -p "${runtime_dir}" "${log_dir}"
+
+ label_exclude_regex="${INCOMPATIBLE_GPU_LABELS[${target}]}"
+ if [[ -n "${EXCLUDE_LABEL_REGEX}" ]]; then
+ label_exclude_regex="(${label_exclude_regex})|(${EXCLUDE_LABEL_REGEX})"
+ fi
+ target_ctest_filter_args=(
+ "${CTEST_BASE_FILTER_ARGS[@]}"
+ -LE "${label_exclude_regex}"
+ )
+
+ if [[ ! -f "${config_path}" ]]; then
+ echo "error: rocjitsu config not found: ${config_path}" >&2
+ build_failed_targets+=("${target} (missing config)")
+ continue
+ fi
+
+ echo
+ echo "=== Retargeting regular IREE build for ${target} ==="
+ # Bash disables errexit inside conditional subshells, so propagate failures
+ # explicitly instead of letting the final CTest listing mask a failed build.
+ if (
+ cmake \
+ -S "${SOURCE_DIR}" \
+ -B "${BUILD_DIR}" \
+ "-DIREE_NATIVE_TEST_TIMEOUT_DEFAULT=${TEST_TIMEOUT}" \
+ "-DIREE_ROCM_TEST_TARGET_CHIP=${target}" || exit 1
+ cmake --build "${BUILD_DIR}" --target iree-test-deps -- -k 0 || exit 1
+ ctest \
+ --test-dir "${BUILD_DIR}" \
+ --show-only=human \
+ "${target_ctest_filter_args[@]}"
+ ) 2>&1 | tee "${log_dir}/build.log"; then
+ echo "Built ${target} ROCm test dependencies"
+ else
+ build_failed_targets+=("${target}")
+ continue
+ fi
+
+ echo
+ if (
+ export IREE_HIP_DYLIB_PATH="${ROCM_ROOT}/lib"
+ export LD_LIBRARY_PATH="${ROCM_ROOT}/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}"
+ ctest_bin="$(command -v ctest)" || exit 1
+ ctest_command=(
+ "${ctest_bin}"
+ --test-dir "${BUILD_DIR}"
+ --parallel "${CTEST_PARALLEL_LEVEL}"
+ --timeout "${TEST_TIMEOUT}"
+ --output-on-failure
+ --no-tests=error
+ --output-junit "${log_dir}/ctest.xml"
+ "${target_ctest_filter_args[@]}"
+ )
+ if [[ "${ROCJITSU_CLI_KIND}" == "session" ]]; then
+ session_runtime_dir="$(mktemp -d "${TMPDIR:-/tmp}/iree-rocjitsu-${target}.XXXXXX")" || exit 1
+ cleanup_session_runtime() {
+ rm -rf -- "${session_runtime_dir}" || exit 1
+ }
+ trap cleanup_session_runtime EXIT
+ export MIRAGE_RUNTIME="${session_runtime_dir}"
+ # Independent tests need no shared GPU state. Give each CTest child its
+ # own simulator instead of contending for a single daemon's workers.
+ timeout \
+ --signal=TERM \
+ --kill-after=30 \
+ "${SESSION_TIMEOUT}" \
+ "${ROCJITSU_BIN}" run \
+ --in-process \
+ --config "${config_path}" \
+ --workdir "${BUILD_DIR}" \
+ --env "IREE_HIP_DYLIB_PATH=${IREE_HIP_DYLIB_PATH}" \
+ --env "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}" \
+ -- "${ctest_command[@]}"
+ else
+ export ROCJITSU_RUNTIME_DIR="${runtime_dir}"
+ "${ROCJITSU_BIN}" --config "${config_path}" -- "${ctest_command[@]}"
+ fi
+ ) 2>&1 | tee "${log_dir}/ctest.log"; then
+ passed_targets+=("${target}")
+ else
+ test_failed_targets+=("${target}")
+ fi
+done
+
+echo
+echo "=== Rocjitsu ROCm e2e summary ==="
+echo "Passed: ${passed_targets[*]:-none}"
+echo "Build failed: ${build_failed_targets[*]:-none}"
+echo "Test failed: ${test_failed_targets[*]:-none}"
+echo "Logs: ${WORK_DIR}/logs/${RUN_ID}"
+
+if ((${#build_failed_targets[@]} || ${#test_failed_targets[@]})); then
+ exit 1
+fi
diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp
index f6cdb12..2775a1d 100644
--- a/compiler/plugins/target/ROCM/ROCMTarget.cpp
+++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp
@@ -701,9 +701,11 @@
}
StringRef targetArch = targetOptions.target;
StringRef targetFeatures = targetOptions.targetFeatures;
+ uint32_t preferredSubgroupSize = 64;
if (auto attr = getGPUTargetAttr(variantOp.getContext(), targetAttr)) {
targetArch = attr.getArch();
targetFeatures = attr.getFeatures();
+ preferredSubgroupSize = attr.getPreferredSubgroupSize();
}
// We name our files after the executable name so that they are easy to
@@ -828,7 +830,7 @@
: llvm::GlobalISelAbortMode::Disable;
SmallVector<std::string> features;
if (chipset.majorVersion >= 10 && chipset.majorVersion <= 12) {
- switch (subgroupSize.value_or(64)) {
+ switch (subgroupSize.value_or(preferredSubgroupSize)) {
case 32:
isWave64 = false;
features.emplace_back("+wavefrontsize32");
diff --git a/compiler/plugins/target/ROCM/test/CMakeLists.txt b/compiler/plugins/target/ROCM/test/CMakeLists.txt
index 7487c87..ba3403f 100644
--- a/compiler/plugins/target/ROCM/test/CMakeLists.txt
+++ b/compiler/plugins/target/ROCM/test/CMakeLists.txt
@@ -45,6 +45,7 @@
NAME
lit
SRCS
+ "default_subgroup_size.mlir"
"emit_debug_info.mlir"
"external_function_validation.mlir"
"module_target_triple.mlir"
diff --git a/compiler/plugins/target/ROCM/test/default_subgroup_size.mlir b/compiler/plugins/target/ROCM/test/default_subgroup_size.mlir
new file mode 100644
index 0000000..7c8c553
--- /dev/null
+++ b/compiler/plugins/target/ROCM/test/default_subgroup_size.mlir
@@ -0,0 +1,25 @@
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=gfx1250 --iree-hal-dump-executable-intermediates-to=%t %s -o /dev/null
+// RUN: cat %t/*.rocmasm | FileCheck %s
+
+// Verify that serialization uses the target's preferred subgroup size when an
+// export does not specify one. gfx1250 only supports wave32; asking LLVM for
+// wave64 silently produces an executable with no kernels.
+
+// CHECK: .globl empty
+// CHECK: amdhsa.kernels:
+// CHECK: .name: empty
+
+#pipeline_layout = #hal.pipeline.layout<bindings = []>
+hal.executable.source public @exe {
+ hal.executable.export public @empty ordinal(0) layout(#pipeline_layout) count(%arg0: !hal.device) -> (index, index, index) {
+ %c1 = arith.constant 1 : index
+ hal.return %c1, %c1, %c1 : index, index, index
+ } attributes {workgroup_size = [1 : index, 1 : index, 1 : index]}
+ builtin.module {
+ llvm.func @empty() attributes {rocdl.kernel} {
+ llvm.return
+ }
+ }
+}
diff --git a/tests/e2e/encoding/CMakeLists.txt b/tests/e2e/encoding/CMakeLists.txt
index a7ea684..4b737ba 100644
--- a/tests/e2e/encoding/CMakeLists.txt
+++ b/tests/e2e/encoding/CMakeLists.txt
@@ -24,3 +24,13 @@
)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
+
+if(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx1250")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/encoding/check_rocm_hip_encoding.mlir"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+endif()
diff --git a/tests/e2e/linalg/CMakeLists.txt b/tests/e2e/linalg/CMakeLists.txt
index 3bc34d7..1a7b049 100644
--- a/tests/e2e/linalg/CMakeLists.txt
+++ b/tests/e2e/linalg/CMakeLists.txt
@@ -217,3 +217,15 @@
)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
+
+if(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx1250")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/linalg/check_rocm_hip_pack_i8.mlir"
+ "iree/tests/e2e/linalg/check_rocm_hip_softmax.mlir"
+ "iree/tests/e2e/linalg/check_rocm_hip_unpack.mlir"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+endif()
diff --git a/tests/e2e/linalg_ext_ops/BUILD.bazel b/tests/e2e/linalg_ext_ops/BUILD.bazel
index 104d751..57238d5 100644
--- a/tests/e2e/linalg_ext_ops/BUILD.bazel
+++ b/tests/e2e/linalg_ext_ops/BUILD.bazel
@@ -129,10 +129,8 @@
"map_load.mlir",
"map_store.mlir",
"scan.mlir",
- "scan_configured.mlir",
"scatter.mlir",
"sort.mlir",
- "topk_v2.mlir",
"winograd_input.mlir",
"winograd_output.mlir",
],
@@ -141,6 +139,10 @@
"top-k.mlir",
"attention_i1_mask.mlir",
"dynamic_attention.mlir",
+ "scan_configured.mlir",
+ # The Distribute pipeline skips TopK v2 thread distribution, causing
+ # overlapping output writes on all ROCm targets even if the test passes.
+ "topk_v2.mlir",
],
)
@@ -152,6 +154,14 @@
)
iree_check_single_backend_test_suite(
+ name = "check_rocm_hip_wave64",
+ srcs = ["scan_configured.mlir"],
+ driver = "hip",
+ tags = ["requires-gpu-wave64"],
+ target_backend = "rocm",
+)
+
+iree_check_single_backend_test_suite(
name = "check_metal-spirv_vulkan",
srcs = enforce_glob(
# keep sorted
diff --git a/tests/e2e/linalg_ext_ops/CMakeLists.txt b/tests/e2e/linalg_ext_ops/CMakeLists.txt
index 819c95c..aba7ceb 100644
--- a/tests/e2e/linalg_ext_ops/CMakeLists.txt
+++ b/tests/e2e/linalg_ext_ops/CMakeLists.txt
@@ -97,10 +97,8 @@
"map_load.mlir"
"map_store.mlir"
"scan.mlir"
- "scan_configured.mlir"
"scatter.mlir"
"sort.mlir"
- "topk_v2.mlir"
"winograd_input.mlir"
"winograd_output.mlir"
TARGET_BACKEND
@@ -111,6 +109,19 @@
iree_check_single_backend_test_suite(
NAME
+ check_rocm_hip_wave64
+ SRCS
+ "scan_configured.mlir"
+ TARGET_BACKEND
+ "rocm"
+ DRIVER
+ "hip"
+ LABELS
+ "requires-gpu-wave64"
+)
+
+iree_check_single_backend_test_suite(
+ NAME
check_metal-spirv_vulkan
SRCS
"arg_compare.mlir"
diff --git a/tests/e2e/math/CMakeLists.txt b/tests/e2e/math/CMakeLists.txt
index 3533fdc..148bd22 100644
--- a/tests/e2e/math/CMakeLists.txt
+++ b/tests/e2e/math/CMakeLists.txt
@@ -65,3 +65,13 @@
)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
+
+if(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx1250")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/math/math_ops_rocm_math_ops_rocm.mlir"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+endif()
diff --git a/tests/e2e/matmul/CMakeLists.txt b/tests/e2e/matmul/CMakeLists.txt
index 5e93079..b521082 100644
--- a/tests/e2e/matmul/CMakeLists.txt
+++ b/tests/e2e/matmul/CMakeLists.txt
@@ -3648,6 +3648,8 @@
"notsan"
"noubsan"
"requires-gpu-rdna4"
+ TIMEOUT
+ 120
)
iree_generated_e2e_runner_test(
@@ -3806,6 +3808,8 @@
"notsan"
"noubsan"
"requires-gpu-rdna4"
+ TIMEOUT
+ 120
)
iree_generated_e2e_runner_test(
@@ -4367,3 +4371,83 @@
)
endif()
+
+# These complete quickly on hardware but take at least 20 seconds per test (or
+# time out) when every GPU instruction is simulated. The RocJITsu CPU-only
+# suite excludes this label; hardware suites continue to run them.
+if(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx94")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_coalesced_dma_f32_2k_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_coalesced_dma_f32_4k_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_coalesced_dma_f32_8k_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_coalesced_dma_f32_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna3_vecdistmfma_block_f16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna3_tileandfusemfma_block_f16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna3_vecdistmfma_block_bf16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna3_tileandfusemfma_block_bf16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna3_vecdistmfma_block_i8_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna3_tileandfusemfma_block_i8_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_dt_f8E4M3FNUZ_tensor_ukernel_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_tensor_ukernel_f8E4M3FNUZ_large_expanded_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_dt_f32_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_dt_f64_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_pad_f32_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_pad_i8_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_tensor_ukernel_f16f16f32_large_expanded_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_dt_tensor_ukernel_f16f16f32_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_tensor_ukernel_bf16bf16f32_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna3_tensor_ukernel_bf16bf16f32_large_expanded_rocm_hip"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+elseif(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx95")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_coalesced_dma_f32_2k_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_coalesced_dma_f32_4k_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_coalesced_dma_f32_8k_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_coalesced_dma_f32_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna4_vecdistmfma_block_f16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna4_vecdistmfma_block_bf16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna4_vecdistmfma_block_i8_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_batch_matmul_cdna4_tileandfusemfma_block_i8_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_dt_f8E4M3FN_tensor_ukernel_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_tensor_ukernel_f8E4M3FN_large_expanded_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_dt_f64_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_pad_f32_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_pad_i8_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_tensor_ukernel_f16f16f32_large_expanded_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_dt_tensor_ukernel_f16f16f32_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_tensor_ukernel_bf16bf16f32_large_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_tensor_ukernel_bf16bf16f32_large_expanded_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_mxfp4_llama_0_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_mxfp4_llama_1_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_mxfp4_llama_2_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_cdna4_mxfp4_llama_3_rocm_hip"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+elseif(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx1250")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_wmma_f16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_tileandfusewmma_f16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_wmma_f16_tb_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_tileandfusewmma_f16_tb_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_wmma_f8E4M3FN_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_tileandfusewmma_f8E4M3FN_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_wmma_f8E4M3FN_tb_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_tileandfusewmma_f8E4M3FN_tb_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_wmma_i8_tb_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_tileandfusewmma_i8_tb_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_dt_f16_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_dt_f8E4M3FN_rocm_hip"
+ "iree/tests/e2e/matmul/e2e_matmul_gfx1250_dt_i8_rocm_hip"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+endif()
diff --git a/tests/e2e/regression/CMakeLists.txt b/tests/e2e/regression/CMakeLists.txt
index b8a5fc2..7ca3e61 100644
--- a/tests/e2e/regression/CMakeLists.txt
+++ b/tests/e2e/regression/CMakeLists.txt
@@ -265,3 +265,17 @@
)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
+
+if(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx1250")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/regression/check_regression_hip_dynamic_batch_matmul_gfx1100.mlir"
+ "iree/tests/e2e/regression/check_regression_hip_dynamic_gather_attention.mlir"
+ "iree/tests/e2e/regression/check_regression_hip_linalg_ops_dynamic.mlir"
+ "iree/tests/e2e/regression/check_regression_hip_split_reduction_using_tiling.mlir"
+ "iree/tests/e2e/regression/check_regression_vector_distribute_64bit_hip_vector_distribute_64bit_amdgpu.mlir"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+endif()
diff --git a/tests/e2e/stablehlo_ops/CMakeLists.txt b/tests/e2e/stablehlo_ops/CMakeLists.txt
index 9ff2976..f7ee488 100644
--- a/tests/e2e/stablehlo_ops/CMakeLists.txt
+++ b/tests/e2e/stablehlo_ops/CMakeLists.txt
@@ -798,3 +798,14 @@
COMPILER_FLAGS
"--iree-input-type=stablehlo"
)
+
+if(IREE_HAL_DRIVER_HIP AND IREE_ROCM_TEST_TARGET_CHIP MATCHES "^gfx1250")
+ foreach(_TEST IN ITEMS
+ "iree/tests/e2e/stablehlo_ops/check_rocm_hip_stream_scatter.mlir"
+ "iree/tests/e2e/stablehlo_ops/check_rocm-rocm_scatter.mlir"
+ )
+ if(TEST "${_TEST}")
+ set_property(TEST "${_TEST}" APPEND PROPERTY LABELS "very-expensive")
+ endif()
+ endforeach()
+endif()
