From 872f51ea4db8ac0f595039d5d8430a2343d67488 Mon Sep 17 00:00:00 2001 From: James Newling Date: Mon, 29 Jun 2026 14:52:34 -0700 Subject: [PATCH 01/19] ci(hipblaslt): add rocjitsu race smoke for gfx950 --- .../run_rocjitsu_hipblaslt_race_smoke.sh | 375 ++++++++++++++++++ .github/workflows/therock-ci-linux.yml | 98 +++++ 2 files changed, 473 insertions(+) create mode 100755 .github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh new file mode 100755 index 000000000000..eac3636f48ec --- /dev/null +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh @@ -0,0 +1,375 @@ +#!/usr/bin/env bash +set -euo pipefail + +# Temporary bridge for rocjitsu race detection in rocm-libraries CI. +# +# The calling workflow fetches hipBLASLt/TensileLite artifacts from the current +# TheRock build, then checks out rocm-systems. Until rocjitsu is packaged as a +# complete runnable TheRock artifact, build the rocjitsu CLI locally and run +# small gfx950 hipBLASLt and TensileLite smokes under the race detector. + +ROCM_PATH="${ROCM_PATH:-${PWD}/build}" +ROCJITSU_SOURCE_DIR="${ROCJITSU_SOURCE_DIR:-${PWD}/rocm-systems/emulation/rocjitsu}" +ROCJITSU_BUILD_DIR="${ROCJITSU_BUILD_DIR:-${PWD}/rocjitsu-build}" +ROCJITSU_CONFIG="${ROCJITSU_CONFIG:-}" +RACE_REPORT_DIR="${RACE_REPORT_DIR:-${PWD}/race-reports}" +HIPBLASLT_BENCH="${HIPBLASLT_BENCH:-${ROCM_PATH}/bin/hipblaslt-bench}" +TENSILELITE_ROOT="${TENSILELITE_ROOT:-${ROCM_PATH}/share/hipblaslt/tensilelite}" +TENSILE_DRIVER="${TENSILE_DRIVER:-${TENSILELITE_ROOT}/Tensile/bin/Tensile}" +TENSILELITE_CLIENT="${TENSILELITE_CLIENT:-${ROCM_PATH}/libexec/hipblaslt/tensilelite/tensilelite-client}" +RACE_TIMEOUT_SECONDS="${RACE_TIMEOUT_SECONDS:-180}" +TENSILELITE_TIMEOUT_SECONDS="${TENSILELITE_TIMEOUT_SECONDS:-420}" +TIMING_FILE="${RACE_REPORT_DIR}/timing.tsv" + +print_timing_summary() { + if [[ -f "${TIMING_FILE}" ]]; then + echo "" + echo "=== rocjitsu race smoke timing summary ===" + printf "%-36s %10s %6s\n" "stage" "seconds" "status" + while IFS=$'\t' read -r label seconds status; do + printf "%-36s %10s %6s\n" "${label}" "${seconds}" "${status}" + done <"${TIMING_FILE}" + fi +} +trap print_timing_summary EXIT + +run_timed() { + local label="$1" + shift + + echo "::group::${label}" + local start + start="$(date +%s)" + local had_errexit=0 + case "$-" in + *e*) had_errexit=1 ;; + esac + + set +e + "$@" + local status=$? + + local end + end="$(date +%s)" + local elapsed=$((end - start)) + echo "::endgroup::" + printf "%s\t%s\t%s\n" "${label}" "${elapsed}" "${status}" | tee -a "${TIMING_FILE}" + + if [[ "${had_errexit}" -ne 0 ]]; then + set -e + else + set +e + fi + return "${status}" +} + +if [[ ! -d "${ROCM_PATH}" ]]; then + echo "ROCM_PATH does not exist: ${ROCM_PATH}" >&2 + exit 1 +fi + +if [[ ! -x "${HIPBLASLT_BENCH}" ]]; then + echo "hipblaslt-bench not found or not executable: ${HIPBLASLT_BENCH}" >&2 + exit 1 +fi + +if [[ ! -d "${TENSILELITE_ROOT}" ]]; then + echo "TensileLite artifacts not found: ${TENSILELITE_ROOT}" >&2 + echo "Expected TheRock BLAS test artifacts to contain share/hipblaslt/tensilelite" >&2 + exit 1 +fi + +if [[ ! -f "${TENSILE_DRIVER}" ]]; then + echo "Tensile driver not found: ${TENSILE_DRIVER}" >&2 + exit 1 +fi + +if [[ ! -x "${TENSILELITE_CLIENT}" ]]; then + echo "tensilelite-client not found or not executable: ${TENSILELITE_CLIENT}" >&2 + exit 1 +fi + +if [[ -z "${ROCJITSU_CONFIG}" ]]; then + for candidate in \ + "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" \ + "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna4_kmd.json"; do + if [[ -f "${candidate}" ]]; then + ROCJITSU_CONFIG="${candidate}" + break + fi + done +fi + +if [[ -z "${ROCJITSU_CONFIG}" || ! -f "${ROCJITSU_CONFIG}" ]]; then + echo "rocjitsu gfx950 config not found under ${ROCJITSU_SOURCE_DIR}/configs" >&2 + exit 1 +fi + +mkdir -p "${RACE_REPORT_DIR}" +: >"${TIMING_FILE}" + +export ROCM_PATH +export PATH="${ROCM_PATH}/bin:${ROCM_PATH}/lib/llvm/bin:${PATH}" +export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/llvm/lib:${LD_LIBRARY_PATH:-}" +export PYTHONPATH="${TENSILELITE_ROOT}${PYTHONPATH:+:${PYTHONPATH}}" + +echo "ROCM_PATH=${ROCM_PATH}" +echo "ROCJITSU_SOURCE_DIR=${ROCJITSU_SOURCE_DIR}" +echo "ROCJITSU_BUILD_DIR=${ROCJITSU_BUILD_DIR}" +echo "ROCJITSU_CONFIG=${ROCJITSU_CONFIG}" +echo "HIPBLASLT_BENCH=${HIPBLASLT_BENCH}" +echo "TENSILELITE_ROOT=${TENSILELITE_ROOT}" +echo "TENSILE_DRIVER=${TENSILE_DRIVER}" +echo "TENSILELITE_CLIENT=${TENSILELITE_CLIENT}" +echo "RACE_REPORT_DIR=${RACE_REPORT_DIR}" +echo "PATH=${PATH}" +echo "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}" +echo "PYTHONPATH=${PYTHONPATH}" + +cmake_args=( + -S "${ROCJITSU_SOURCE_DIR}" + -B "${ROCJITSU_BUILD_DIR}" + -G Ninja + -DCMAKE_BUILD_TYPE=Release + -DBUILD_TESTING=OFF + -DROCM_PATH="${ROCM_PATH}" + -DCMAKE_PREFIX_PATH="${ROCM_PATH}" + -DCMAKE_CXX_FLAGS="-Wno-error=unknown-warning-option -Wno-error=nested-anon-types" +) + +if command -v amdclang >/dev/null 2>&1 && command -v amdclang++ >/dev/null 2>&1; then + cmake_args+=( + -DCMAKE_C_COMPILER="$(command -v amdclang)" + -DCMAKE_CXX_COMPILER="$(command -v amdclang++)" + ) +fi + +run_timed "configure rocjitsu" cmake "${cmake_args[@]}" + +cmake_target_exists() { + local target="$1" + cmake --build "${ROCJITSU_BUILD_DIR}" --target help \ + | grep -Eq "(^|[[:space:]])${target}([:[:space:]]|$)" +} + +build_targets=(rocjitsu_bin) +if cmake_target_exists rocjitsu_shared; then + build_targets+=(rocjitsu_shared) +fi +if cmake_target_exists rocjitsu_kmd_shim; then + build_targets+=(rocjitsu_kmd_shim) +fi +run_timed "build rocjitsu" cmake --build "${ROCJITSU_BUILD_DIR}" --target "${build_targets[@]}" + +ROCJITSU_BIN="${ROCJITSU_BUILD_DIR}/tools/rocjitsu/rocjitsu" +if [[ ! -x "${ROCJITSU_BIN}" ]]; then + echo "rocjitsu binary not found after build: ${ROCJITSU_BIN}" >&2 + exit 1 +fi + +rm -f "${RACE_REPORT_DIR}/race.log" + +show_rocjitsu_version() { + echo "rocjitsu version:" + "${ROCJITSU_BIN}" --version +} + +run_sanity_check() { + echo "rocjitsu /bin/true sanity check:" + mkdir -p "${RACE_REPORT_DIR}/sanity" + timeout 30 \ + env \ + RJ_LOG=1 \ + RJ_SINKS=stderr,file \ + RJ_SINK_DIR="${RACE_REPORT_DIR}/sanity" \ + "${ROCJITSU_BIN}" \ + --config "${ROCJITSU_CONFIG}" \ + -- /bin/true \ + 2>&1 | tee "${RACE_REPORT_DIR}/sanity.log" + local status=$? + if [[ "${status}" -ne 0 ]]; then + echo "rocjitsu /bin/true sanity check failed with status ${status}" >&2 + return "${status}" + fi +} + +run_hipblaslt_bench_smoke() { + mkdir -p "${RACE_REPORT_DIR}/hipblaslt-bench" + echo "running hipblaslt-bench under rocjitsu race detection" + timeout "${RACE_TIMEOUT_SECONDS}" \ + env \ + HSA_ENABLE_SDMA=1 \ + RJ_RACE=1 \ + RJ_LOG=1 \ + RJ_SINKS=stderr,file \ + RJ_SINK_DIR="${RACE_REPORT_DIR}/hipblaslt-bench" \ + "${ROCJITSU_BIN}" \ + --config "${ROCJITSU_CONFIG}" \ + -- "${HIPBLASLT_BENCH}" \ + --precision f32_r \ + -m 128 \ + -n 128 \ + -k 128 \ + --iters 1 \ + --cold_iters 0 \ + 2>&1 | tee "${RACE_REPORT_DIR}/hipblaslt-bench.log" + local status=$? + if [[ "${status}" -ne 0 ]]; then + echo "hipblaslt-bench smoke command failed with status ${status}" >&2 + return "${status}" + fi + + if [[ -f "${RACE_REPORT_DIR}/hipblaslt-bench/race.log" ]] && + grep -q '^RACE ' "${RACE_REPORT_DIR}/hipblaslt-bench/race.log"; then + echo "rocjitsu race detector reported a race in hipblaslt-bench:" >&2 + cat "${RACE_REPORT_DIR}/hipblaslt-bench/race.log" >&2 + return 1 + fi +} + +write_tensilelite_smoke_yaml() { + local yaml="$1" + cat >"${yaml}" <<'YAML' +GlobalParameters: + NumElementsToValidate: -1 + DataInitTypeBeta: 0 + DataInitTypeAlpha: 1 + Device: 0 + CpuThreads: 1 + +BenchmarkProblems: + - + - OperationType: GEMM + DataType: s + TransposeA: False + TransposeB: True + UseBeta: True + Batched: True + - InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - MatrixInstruction: + - [16, 16, 4, 1, 1, 2, 3, 2, 2] + - PrefetchGlobalRead: [2] + - PrefetchLocalRead: [1] + - DepthU: [32] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - TransposeLDS: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - StaggerU: [16] + - StaggerUStride: [-1] + - WorkGroupMapping: [1] + - 1LDSBuffer: [-1] + - WorkGroupMappingXCC: [8] + - WorkGroupMappingXCCGroup: [-1] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: ["MultipleBuffer"] + - GlobalReadPerMfma: [1.0] + - LocalWritePerMfma: [-1] + - StoreRemapVectorWidth: [0] + - StoreVectorWidth: [-1] + - SourceSwap: [True] + - NumElementsPerBatchStore: [16] + - ClusterLocalRead: [1] + - DirectToVgprA: [True] + - DirectToVgprB: [False] + - WorkGroup: [[32, 4, 4]] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [32, 32, 1, 32] +YAML +} + +run_tensilelite_client_smoke() { + local test_dir="${RACE_REPORT_DIR}/tensilelite" + local output_dir="${test_dir}/output" + local sink_dir="${test_dir}/sinks" + local yaml="${test_dir}/f32_nt_32x32.yaml" + mkdir -p "${output_dir}" "${sink_dir}" + write_tensilelite_smoke_yaml "${yaml}" + + local tensile_args=( + "${yaml}" + "${output_dir}" + --prebuilt-client "${TENSILELITE_CLIENT}" + --gpu-targets gfx950 + --library-format msgpack + ) + if [[ -x "${ROCM_PATH}/bin/amdclang++" ]]; then + tensile_args+=(--cxx-compiler "${ROCM_PATH}/bin/amdclang++") + fi + + echo "running tensilelite-client under rocjitsu race detection" + timeout "${TENSILELITE_TIMEOUT_SECONDS}" \ + env \ + HSA_ENABLE_SDMA=1 \ + RJ_RACE=1 \ + RJ_LOG=1 \ + RJ_SINKS=stderr,file \ + RJ_SINK_DIR="${sink_dir}" \ + "${ROCJITSU_BIN}" \ + --config "${ROCJITSU_CONFIG}" \ + -- python3 "${TENSILE_DRIVER}" "${tensile_args[@]}" \ + 2>&1 | tee "${RACE_REPORT_DIR}/tensilelite-client.log" + local status=$? + if [[ "${status}" -ne 0 ]]; then + echo "tensilelite-client smoke command failed with status ${status}" >&2 + return "${status}" + fi + + if ! grep -q "PASSED" "${RACE_REPORT_DIR}/tensilelite-client.log"; then + echo "tensilelite-client smoke did not report validation success" >&2 + return 1 + fi + + if [[ -f "${sink_dir}/race.log" ]] && grep -q '^RACE ' "${sink_dir}/race.log"; then + echo "rocjitsu race detector reported a race in tensilelite-client:" >&2 + cat "${sink_dir}/race.log" >&2 + return 1 + fi +} + +run_timed "rocjitsu version" show_rocjitsu_version + +set +e +run_timed "rocjitsu sanity check" run_sanity_check +sanity_status=$? +set -e +if [[ "${sanity_status}" -ne 0 ]]; then + echo "WARNING: rocjitsu /bin/true sanity check failed; continuing to GPU workload smokes" >&2 +fi + +smoke_status=0 + +set +e +run_timed "hipblaslt-bench race smoke" run_hipblaslt_bench_smoke +hipblaslt_status=$? +run_timed "tensilelite-client race smoke" run_tensilelite_client_smoke +tensilelite_status=$? +set -e + +if [[ "${hipblaslt_status}" -ne 0 ]]; then + echo "hipblaslt-bench race smoke failed with status ${hipblaslt_status}" >&2 + smoke_status=1 +fi + +if [[ "${tensilelite_status}" -ne 0 ]]; then + echo "tensilelite-client race smoke failed with status ${tensilelite_status}" >&2 + smoke_status=1 +fi + +if [[ "${smoke_status}" -ne 0 ]]; then + echo "one or more rocjitsu race smokes failed" >&2 + exit "${smoke_status}" +fi + +echo "rocjitsu race smokes completed without race reports" diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index 0948b1e4ee59..2821a33ad67e 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -171,3 +171,101 @@ jobs: platform: "linux" test_type: ${{ inputs.test_type }} secrets: inherit + + rocjitsu-race-smoke-linux: + name: rocjitsu race smoke (${{ inputs.amdgpu_families }}) + needs: [therock-build-linux] + if: ${{ inputs.test_runs_on != '' && startsWith(inputs.amdgpu_families, 'gfx950') && contains(format(',{0},', inputs.projects_to_test), ',hipblaslt,') }} + runs-on: ${{ inputs.test_runs_on }} + timeout-minutes: 90 + container: + image: ghcr.io/rocm/no_rocm_image_ubuntu24_04@sha256:4150afe4759d14822f0e3f8930e1124f26e11f68b5c7b91ec9a02b20b1ebbb98 + options: --ipc host + --user 0:0 + --ulimit memlock=-1:-1 + --security-opt seccomp=unconfined + defaults: + run: + shell: bash + env: + VENV_DIR: ${{ github.workspace }}/.venv + ARTIFACT_RUN_ID: ${{ github.run_id }} + OUTPUT_ARTIFACTS_DIR: ./build + ROCM_PATH: ${{ github.workspace }}/build + ROCJITSU_SOURCE_DIR: ${{ github.workspace }}/rocm-systems/emulation/rocjitsu + ROCJITSU_BUILD_DIR: ${{ github.workspace }}/rocjitsu-build + RACE_REPORT_DIR: ${{ github.workspace }}/race-reports + RACE_TIMEOUT_SECONDS: "180" + TENSILELITE_TIMEOUT_SECONDS: "420" + steps: + - name: Checkout TheRock repository + uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 + with: + repository: "ROCm/TheRock" + ref: ${{ inputs.therock_ref }} + + - name: Checkout rocm-libraries repository for scripts + uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 + with: + path: rocm-libraries + + - name: Checkout rocm-systems for rocjitsu + uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 + with: + repository: newling/rocm-systems + ref: rocjitsu-early-mmap-fix + path: rocm-systems + sparse-checkout: | + emulation/rocjitsu + shared/machine-readable-isa/isa + sparse-checkout-cone-mode: true + + - name: Install rocjitsu build dependencies + run: | + apt-get update + apt-get install -y --no-install-recommends \ + build-essential \ + ca-certificates \ + cmake \ + git \ + libdrm-dev \ + ninja-build + + - name: Run setup test environment workflow + timeout-minutes: 15 + uses: ./.github/actions/setup_test_environment + with: + ARTIFACT_RUN_ID: ${{ env.ARTIFACT_RUN_ID }} + ARTIFACT_GROUP: ${{ inputs.amdgpu_families }} + OUTPUT_ARTIFACTS_DIR: ${{ env.OUTPUT_ARTIFACTS_DIR }} + VENV_DIR: ${{ env.VENV_DIR }} + FETCH_ARTIFACT_ARGS: --blas --tests + + - name: Verify TensileLite Python environment + run: | + export PYTHONPATH="${ROCM_PATH}/share/hipblaslt/tensilelite${PYTHONPATH:+:${PYTHONPATH}}" + export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/llvm/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}" + "${VENV_DIR}/bin/python" - <<'PY' + import filelock + import joblib + import msgpack + import numpy + import yaml + import Tensile + import rocisa + import rocisa.instruction + print("Tensile:", Tensile.ROOT_PATH) + print("rocisa:", rocisa.__file__) + PY + + - name: Run rocjitsu race smoke + timeout-minutes: 75 + run: bash rocm-libraries/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh + + - name: Upload race reports + if: ${{ always() }} + uses: actions/upload-artifact@de65e23aa2b7e23d713bb51fbfcb6d502f8667d8 # v4.6.2 + with: + name: rocjitsu-race-reports-${{ inputs.amdgpu_families }} + path: race-reports/ + if-no-files-found: ignore From b5c7d453af2d7e304db9d75f3c981c6ac1543606 Mon Sep 17 00:00:00 2001 From: James Newling Date: Thu, 9 Jul 2026 12:16:46 -0700 Subject: [PATCH 02/19] ci(hipblaslt): add sysdeps lib path to race smoke --- .github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh | 2 +- .github/workflows/therock-ci-linux.yml | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh index eac3636f48ec..2db8a7356739 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh @@ -110,7 +110,7 @@ mkdir -p "${RACE_REPORT_DIR}" export ROCM_PATH export PATH="${ROCM_PATH}/bin:${ROCM_PATH}/lib/llvm/bin:${PATH}" -export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/llvm/lib:${LD_LIBRARY_PATH:-}" +export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROCM_PATH}/lib/llvm/lib:${LD_LIBRARY_PATH:-}" export PYTHONPATH="${TENSILELITE_ROOT}${PYTHONPATH:+:${PYTHONPATH}}" echo "ROCM_PATH=${ROCM_PATH}" diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index 2821a33ad67e..e7820e26daf6 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -244,7 +244,7 @@ jobs: - name: Verify TensileLite Python environment run: | export PYTHONPATH="${ROCM_PATH}/share/hipblaslt/tensilelite${PYTHONPATH:+:${PYTHONPATH}}" - export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/llvm/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}" + export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROCM_PATH}/lib/llvm/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}" "${VENV_DIR}/bin/python" - <<'PY' import filelock import joblib From 097fa78929e4a7b744b63e6fe25d3be56956d32c Mon Sep 17 00:00:00 2001 From: James Newling Date: Thu, 9 Jul 2026 15:24:07 -0700 Subject: [PATCH 03/19] ci(hipblaslt): use zero init for race smoke --- .github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh index 2db8a7356739..f064707bd611 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh @@ -207,6 +207,7 @@ run_hipblaslt_bench_smoke() { --config "${ROCJITSU_CONFIG}" \ -- "${HIPBLASLT_BENCH}" \ --precision f32_r \ + --initialization zero \ -m 128 \ -n 128 \ -k 128 \ From 973e061c1103160d57b845a54ac74a92a90bd391 Mon Sep 17 00:00:00 2001 From: James Newling Date: Fri, 10 Jul 2026 09:59:29 -0700 Subject: [PATCH 04/19] ci(hipblaslt): document rocjitsu race smoke flow --- .../run_rocjitsu_hipblaslt_race_smoke.sh | 50 +++++++++++++++++++ .github/workflows/therock-ci-linux.yml | 25 ++++++++++ 2 files changed, 75 insertions(+) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh index f064707bd611..a1d54d34f832 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh @@ -7,6 +7,12 @@ set -euo pipefail # TheRock build, then checks out rocm-systems. Until rocjitsu is packaged as a # complete runnable TheRock artifact, build the rocjitsu CLI locally and run # small gfx950 hipBLASLt and TensileLite smokes under the race detector. +# +# Keep this script deliberately narrow. It is not a replacement for the normal +# hipBLASLt or TensileLite test suites; it is a post-build instrumentation pass +# that proves the TheRock BLAS artifacts can be launched through rocjitsu and do +# not emit race reports for one small, deterministic gfx950 GEMM path. The +# normal package tests still own broad functional coverage. ROCM_PATH="${ROCM_PATH:-${PWD}/build}" ROCJITSU_SOURCE_DIR="${ROCJITSU_SOURCE_DIR:-${PWD}/rocm-systems/emulation/rocjitsu}" @@ -21,6 +27,10 @@ RACE_TIMEOUT_SECONDS="${RACE_TIMEOUT_SECONDS:-180}" TENSILELITE_TIMEOUT_SECONDS="${TENSILELITE_TIMEOUT_SECONDS:-420}" TIMING_FILE="${RACE_REPORT_DIR}/timing.tsv" +# The rocjitsu smoke currently has several expensive pieces: artifact setup in +# the caller, local rocjitsu configure/build here, and the two instrumented GPU +# workloads. Keep a durable timing file as well as grouped console output so +# slow CI runs are diagnosable without re-running the workflow. print_timing_summary() { if [[ -f "${TIMING_FILE}" ]]; then echo "" @@ -63,6 +73,9 @@ run_timed() { return "${status}" } +# Fail before doing any rocjitsu work if the expected TheRock BLAS test +# artifacts are missing. Most failures here mean the artifact fetch/layout +# changed, not that the race detector found a race. if [[ ! -d "${ROCM_PATH}" ]]; then echo "ROCM_PATH does not exist: ${ROCM_PATH}" >&2 exit 1 @@ -89,6 +102,9 @@ if [[ ! -x "${TENSILELITE_CLIENT}" ]]; then exit 1 fi +# Prefer the architecture-specific config name when present, but accept the +# older generic CDNA4 config while rocm-systems and TheRock packaging are still +# converging. if [[ -z "${ROCJITSU_CONFIG}" ]]; then for candidate in \ "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" \ @@ -108,6 +124,10 @@ fi mkdir -p "${RACE_REPORT_DIR}" : >"${TIMING_FILE}" +# Everything below must resolve against the fetched ROCm payload, not whatever +# happens to be installed in the CI image. In particular, lib/rocm_sysdeps/lib +# is needed by some packaged dependencies, and the TensileLite Python tree comes +# from the BLAS test artifacts. export ROCM_PATH export PATH="${ROCM_PATH}/bin:${ROCM_PATH}/lib/llvm/bin:${PATH}" export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROCM_PATH}/lib/llvm/lib:${LD_LIBRARY_PATH:-}" @@ -126,6 +146,9 @@ echo "PATH=${PATH}" echo "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}" echo "PYTHONPATH=${PYTHONPATH}" +# rocjitsu is still consumed from a source checkout in this workflow. Build only +# the CLI and optional runtime/shim targets needed to launch the smoke workloads; +# this keeps the job independent of full rocm-systems packaging. cmake_args=( -S "${ROCJITSU_SOURCE_DIR}" -B "${ROCJITSU_BUILD_DIR}" @@ -146,6 +169,9 @@ fi run_timed "configure rocjitsu" cmake "${cmake_args[@]}" +# rocm-systems has been renaming/splitting rocjitsu build targets while this +# bridge job is being developed. Probe for optional targets instead of assuming +# one exact target set, so the CI job survives small target-layout changes. cmake_target_exists() { local target="$1" cmake --build "${ROCJITSU_BUILD_DIR}" --target help \ @@ -174,6 +200,9 @@ show_rocjitsu_version() { "${ROCJITSU_BIN}" --version } +# This is a cheap launch-path check for rocjitsu itself. Treat it as diagnostic +# rather than authoritative: the real signal is whether HIP/HSA workloads from +# the fetched ROCm payload run and whether their race logs contain reports. run_sanity_check() { echo "rocjitsu /bin/true sanity check:" mkdir -p "${RACE_REPORT_DIR}/sanity" @@ -196,6 +225,10 @@ run_sanity_check() { run_hipblaslt_bench_smoke() { mkdir -p "${RACE_REPORT_DIR}/hipblaslt-bench" echo "running hipblaslt-bench under rocjitsu race detection" + # Use zero initialization so this smoke measures the hipBLASLt GEMM dispatch + # path rather than spending most of its signal on device-side fill kernels. + # The default HPL initialization can generate separate fill kernels that are + # useful for compiler/detector investigation, but they obscure this CI gate. timeout "${RACE_TIMEOUT_SECONDS}" \ env \ HSA_ENABLE_SDMA=1 \ @@ -220,6 +253,9 @@ run_hipblaslt_bench_smoke() { return "${status}" fi + # rocjitsu writes race findings to its sink directory. The application may + # still exit normally, so explicitly inspect race.log and make any report a CI + # failure with the report echoed into the job log. if [[ -f "${RACE_REPORT_DIR}/hipblaslt-bench/race.log" ]] && grep -q '^RACE ' "${RACE_REPORT_DIR}/hipblaslt-bench/race.log"; then echo "rocjitsu race detector reported a race in hipblaslt-bench:" >&2 @@ -230,6 +266,10 @@ run_hipblaslt_bench_smoke() { write_tensilelite_smoke_yaml() { local yaml="$1" + # This reduced TensileLite config is intentionally embedded instead of checked + # in as a test-data file. The CI smoke only needs one small f32 assembly GEMM + # to exercise the Python driver, generated code object, prebuilt + # tensilelite-client, and rocjitsu race detector together. cat >"${yaml}" <<'YAML' GlobalParameters: NumElementsToValidate: -1 @@ -310,6 +350,10 @@ run_tensilelite_client_smoke() { fi echo "running tensilelite-client under rocjitsu race detection" + # Drive the normal Tensile front end with the client from the TheRock artifact. + # This covers a different surface from hipblaslt-bench: Python-side Tensile + # setup, rocisa imports, generated client config, and the standalone + # tensilelite-client runtime path. timeout "${TENSILELITE_TIMEOUT_SECONDS}" \ env \ HSA_ENABLE_SDMA=1 \ @@ -332,6 +376,9 @@ run_tensilelite_client_smoke() { return 1 fi + # As above, validation success and race-detector success are distinct signals. + # Require both: the client must pass numerics, and rocjitsu must leave the race + # sink free of RACE records. if [[ -f "${sink_dir}/race.log" ]] && grep -q '^RACE ' "${sink_dir}/race.log"; then echo "rocjitsu race detector reported a race in tensilelite-client:" >&2 cat "${sink_dir}/race.log" >&2 @@ -351,6 +398,9 @@ fi smoke_status=0 +# Run both workload smokes even if the first one fails. That gives the uploaded +# race-reports artifact a complete picture for the failing CI attempt instead of +# forcing a second long run just to learn whether the other workload also broke. set +e run_timed "hipblaslt-bench race smoke" run_hipblaslt_bench_smoke hipblaslt_status=$? diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index e7820e26daf6..cc458d73e5e4 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -175,6 +175,11 @@ jobs: rocjitsu-race-smoke-linux: name: rocjitsu race smoke (${{ inputs.amdgpu_families }}) needs: [therock-build-linux] + # This is a narrow post-build instrumentation job for hipBLASLt on gfx950. + # Normal package tests cover broad functionality; this job fetches the BLAS + # artifacts from the just-finished TheRock build and runs two small workloads + # under rocjitsu's race detector. Keep it gated to hipblaslt + gfx950 while + # rocjitsu packaging and runner capacity are still evolving. if: ${{ inputs.test_runs_on != '' && startsWith(inputs.amdgpu_families, 'gfx950') && contains(format(',{0},', inputs.projects_to_test), ',hipblaslt,') }} runs-on: ${{ inputs.test_runs_on }} timeout-minutes: 90 @@ -198,6 +203,11 @@ jobs: RACE_TIMEOUT_SECONDS: "180" TENSILELITE_TIMEOUT_SECONDS: "420" steps: + # The job needs three inputs in one workspace: + # 1. TheRock actions/scripts to fetch the exact artifacts from this run. + # 2. This rocm-libraries checkout for the race-smoke driver script. + # 3. rocm-systems source for rocjitsu until rocjitsu is packaged as a + # complete runnable artifact. - name: Checkout TheRock repository uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 with: @@ -220,6 +230,9 @@ jobs: shared/machine-readable-isa/isa sparse-checkout-cone-mode: true + # The base image intentionally does not contain a full ROCm install. Keep + # the apt dependency set to ordinary host build tools; the ROCm runtime, + # clients, and TensileLite Python tree come from TheRock artifacts below. - name: Install rocjitsu build dependencies run: | apt-get update @@ -231,6 +244,10 @@ jobs: libdrm-dev \ ninja-build + # Pull the BLAS runtime and test payload from the current TheRock build. + # This is what keeps the smoke aligned with the exact hipBLASLt and + # TensileLite artifacts produced by the PR, rather than an image-local or + # previously installed ROCm tree. - name: Run setup test environment workflow timeout-minutes: 15 uses: ./.github/actions/setup_test_environment @@ -241,6 +258,10 @@ jobs: VENV_DIR: ${{ env.VENV_DIR }} FETCH_ARTIFACT_ARGS: --blas --tests + # Fail early with import paths printed if the test artifact is missing + # Python dependencies or if PYTHONPATH/LD_LIBRARY_PATH no longer match the + # artifact layout. That makes packaging/layout regressions distinct from + # real race-detector findings. - name: Verify TensileLite Python environment run: | export PYTHONPATH="${ROCM_PATH}/share/hipblaslt/tensilelite${PYTHONPATH:+:${PYTHONPATH}}" @@ -262,6 +283,10 @@ jobs: timeout-minutes: 75 run: bash rocm-libraries/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh + # Upload even on failure. The report directory contains the per-workload + # stdout logs, rocjitsu sink files, race.log files, generated reduced + # TensileLite input/output, and timing.tsv; without this artifact a failed + # detector run is usually not actionable from the console log alone. - name: Upload race reports if: ${{ always() }} uses: actions/upload-artifact@de65e23aa2b7e23d713bb51fbfcb6d502f8667d8 # v4.6.2 From f4927f84851e02d6b6e6e788d9ab2ae7cc3128e9 Mon Sep 17 00:00:00 2001 From: James Newling Date: Fri, 10 Jul 2026 10:54:38 -0700 Subject: [PATCH 05/19] ci(hipblaslt): clarify rocjitsu race check flow --- ...h => run_rocjitsu_hipblaslt_race_check.sh} | 123 +++++++++--------- .github/workflows/therock-ci-linux.yml | 14 +- 2 files changed, 71 insertions(+), 66 deletions(-) rename .github/scripts/{run_rocjitsu_hipblaslt_race_smoke.sh => run_rocjitsu_hipblaslt_race_check.sh} (76%) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh similarity index 76% rename from .github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh rename to .github/scripts/run_rocjitsu_hipblaslt_race_check.sh index a1d54d34f832..2a40344d21f6 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -1,18 +1,13 @@ #!/usr/bin/env bash set -euo pipefail -# Temporary bridge for rocjitsu race detection in rocm-libraries CI. +# This script's calling workflow fetches hipBLASLt/TensileLite artifacts from +# the current TheRock build, then checks out rocm-systems. # -# The calling workflow fetches hipBLASLt/TensileLite artifacts from the current -# TheRock build, then checks out rocm-systems. Until rocjitsu is packaged as a -# complete runnable TheRock artifact, build the rocjitsu CLI locally and run -# small gfx950 hipBLASLt and TensileLite smokes under the race detector. +# TODO(newling) Until rocjitsu is packaged as a complete runnable TheRock +# artifact, we build the rocjitsu CLI locally. # -# Keep this script deliberately narrow. It is not a replacement for the normal -# hipBLASLt or TensileLite test suites; it is a post-build instrumentation pass -# that proves the TheRock BLAS artifacts can be launched through rocjitsu and do -# not emit race reports for one small, deterministic gfx950 GEMM path. The -# normal package tests still own broad functional coverage. +# The script runs small gfx950 hipBLASLt and TensileLite GEMMs under the race detector. ROCM_PATH="${ROCM_PATH:-${PWD}/build}" ROCJITSU_SOURCE_DIR="${ROCJITSU_SOURCE_DIR:-${PWD}/rocm-systems/emulation/rocjitsu}" @@ -27,14 +22,12 @@ RACE_TIMEOUT_SECONDS="${RACE_TIMEOUT_SECONDS:-180}" TENSILELITE_TIMEOUT_SECONDS="${TENSILELITE_TIMEOUT_SECONDS:-420}" TIMING_FILE="${RACE_REPORT_DIR}/timing.tsv" -# The rocjitsu smoke currently has several expensive pieces: artifact setup in -# the caller, local rocjitsu configure/build here, and the two instrumented GPU -# workloads. Keep a durable timing file as well as grouped console output so -# slow CI runs are diagnosable without re-running the workflow. +# Record the duration and exit status of each stage in a file that is uploaded +# with the other race reports. The same information is also printed at exit. print_timing_summary() { if [[ -f "${TIMING_FILE}" ]]; then echo "" - echo "=== rocjitsu race smoke timing summary ===" + echo "=== rocjitsu race check timing summary ===" printf "%-36s %10s %6s\n" "stage" "seconds" "status" while IFS=$'\t' read -r label seconds status; do printf "%-36s %10s %6s\n" "${label}" "${seconds}" "${status}" @@ -73,9 +66,10 @@ run_timed() { return "${status}" } -# Fail before doing any rocjitsu work if the expected TheRock BLAS test -# artifacts are missing. Most failures here mean the artifact fetch/layout -# changed, not that the race detector found a race. + +# These checks validate the artifact layout before any workload is launched +# under rocjitsu. A failure here means the expected TheRock BLAS test payload is +# not available to this job. if [[ ! -d "${ROCM_PATH}" ]]; then echo "ROCM_PATH does not exist: ${ROCM_PATH}" >&2 exit 1 @@ -102,9 +96,10 @@ if [[ ! -x "${TENSILELITE_CLIENT}" ]]; then exit 1 fi -# Prefer the architecture-specific config name when present, but accept the -# older generic CDNA4 config while rocm-systems and TheRock packaging are still -# converging. +# Use the gfx950/CDNA4 KMD config from the checked-out rocm-systems tree. The +# workflow pins the rocm-systems ref; if that ref renames configs, update this +# list or pass ROCJITSU_CONFIG from the workflow. Keep the selection explicit so +# future architectures can add their own config candidates here. if [[ -z "${ROCJITSU_CONFIG}" ]]; then for candidate in \ "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" \ @@ -125,9 +120,8 @@ mkdir -p "${RACE_REPORT_DIR}" : >"${TIMING_FILE}" # Everything below must resolve against the fetched ROCm payload, not whatever -# happens to be installed in the CI image. In particular, lib/rocm_sysdeps/lib -# is needed by some packaged dependencies, and the TensileLite Python tree comes -# from the BLAS test artifacts. +# happens to be installed in the CI image. ROCM_PATH is exported for subprocesses +# that use it to find the ROCm install root. export ROCM_PATH export PATH="${ROCM_PATH}/bin:${ROCM_PATH}/lib/llvm/bin:${PATH}" export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROCM_PATH}/lib/llvm/lib:${LD_LIBRARY_PATH:-}" @@ -146,9 +140,18 @@ echo "PATH=${PATH}" echo "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}" echo "PYTHONPATH=${PYTHONPATH}" + +# TODO(newling): Track migration to a packaged rocjitsu once TheRock provides a +# complete runnable artifact. Until then, build rocjitsu from the pinned +# rocm-systems checkout so this job controls the tool version. # rocjitsu is still consumed from a source checkout in this workflow. Build only -# the CLI and optional runtime/shim targets needed to launch the smoke workloads; +# the CLI and optional runtime/shim targets needed to launch the test workloads; # this keeps the job independent of full rocm-systems packaging. +# +# The warning suppressions keep the local rocjitsu build from failing on +# compiler/header warning mismatches. `nested-anon-types` is a Clang warning for +# anonymous structs/unions nested inside another type; use -Wno-error so those +# warnings remain visible but do not fail this bridge build. cmake_args=( -S "${ROCJITSU_SOURCE_DIR}" -B "${ROCJITSU_BUILD_DIR}" @@ -160,12 +163,14 @@ cmake_args=( -DCMAKE_CXX_FLAGS="-Wno-error=unknown-warning-option -Wno-error=nested-anon-types" ) -if command -v amdclang >/dev/null 2>&1 && command -v amdclang++ >/dev/null 2>&1; then - cmake_args+=( - -DCMAKE_C_COMPILER="$(command -v amdclang)" - -DCMAKE_CXX_COMPILER="$(command -v amdclang++)" - ) +if ! command -v amdclang >/dev/null 2>&1 || ! command -v amdclang++ >/dev/null 2>&1; then + echo "amdclang and amdclang++ must be available from the fetched ROCm payload" >&2 + exit 1 fi +cmake_args+=( + -DCMAKE_C_COMPILER="$(command -v amdclang)" + -DCMAKE_CXX_COMPILER="$(command -v amdclang++)" +) run_timed "configure rocjitsu" cmake "${cmake_args[@]}" @@ -193,8 +198,6 @@ if [[ ! -x "${ROCJITSU_BIN}" ]]; then exit 1 fi -rm -f "${RACE_REPORT_DIR}/race.log" - show_rocjitsu_version() { echo "rocjitsu version:" "${ROCJITSU_BIN}" --version @@ -222,13 +225,12 @@ run_sanity_check() { fi } -run_hipblaslt_bench_smoke() { +run_hipblaslt_bench_check() { mkdir -p "${RACE_REPORT_DIR}/hipblaslt-bench" echo "running hipblaslt-bench under rocjitsu race detection" - # Use zero initialization so this smoke measures the hipBLASLt GEMM dispatch - # path rather than spending most of its signal on device-side fill kernels. - # The default HPL initialization can generate separate fill kernels that are - # useful for compiler/detector investigation, but they obscure this CI gate. + # Use zero initialization because the default HPL initialization currently + # triggers a separate device-fill kernel race report that needs independent + # investigation. This check is scoped to the hipBLASLt GEMM dispatch path. timeout "${RACE_TIMEOUT_SECONDS}" \ env \ HSA_ENABLE_SDMA=1 \ @@ -249,7 +251,7 @@ run_hipblaslt_bench_smoke() { 2>&1 | tee "${RACE_REPORT_DIR}/hipblaslt-bench.log" local status=$? if [[ "${status}" -ne 0 ]]; then - echo "hipblaslt-bench smoke command failed with status ${status}" >&2 + echo "hipblaslt-bench race check command failed with status ${status}" >&2 return "${status}" fi @@ -264,12 +266,13 @@ run_hipblaslt_bench_smoke() { fi } -write_tensilelite_smoke_yaml() { + +write_tensilelite_check_yaml() { local yaml="$1" - # This reduced TensileLite config is intentionally embedded instead of checked - # in as a test-data file. The CI smoke only needs one small f32 assembly GEMM - # to exercise the Python driver, generated code object, prebuilt - # tensilelite-client, and rocjitsu race detector together. + # One reduced config is embedded here to keep this bridge job self-contained. + # If the race check grows to multiple configs, move them into a small checked-in + # data directory or install them with the TensileLite test artifacts, then have + # this script iterate over that list. cat >"${yaml}" <<'YAML' GlobalParameters: NumElementsToValidate: -1 @@ -330,13 +333,13 @@ BenchmarkProblems: YAML } -run_tensilelite_client_smoke() { +run_tensilelite_client_check() { local test_dir="${RACE_REPORT_DIR}/tensilelite" local output_dir="${test_dir}/output" local sink_dir="${test_dir}/sinks" local yaml="${test_dir}/f32_nt_32x32.yaml" mkdir -p "${output_dir}" "${sink_dir}" - write_tensilelite_smoke_yaml "${yaml}" + write_tensilelite_check_yaml "${yaml}" local tensile_args=( "${yaml}" @@ -367,12 +370,12 @@ run_tensilelite_client_smoke() { 2>&1 | tee "${RACE_REPORT_DIR}/tensilelite-client.log" local status=$? if [[ "${status}" -ne 0 ]]; then - echo "tensilelite-client smoke command failed with status ${status}" >&2 + echo "tensilelite-client race check command failed with status ${status}" >&2 return "${status}" fi if ! grep -q "PASSED" "${RACE_REPORT_DIR}/tensilelite-client.log"; then - echo "tensilelite-client smoke did not report validation success" >&2 + echo "tensilelite-client race check did not report validation success" >&2 return 1 fi @@ -393,34 +396,34 @@ run_timed "rocjitsu sanity check" run_sanity_check sanity_status=$? set -e if [[ "${sanity_status}" -ne 0 ]]; then - echo "WARNING: rocjitsu /bin/true sanity check failed; continuing to GPU workload smokes" >&2 + echo "WARNING: rocjitsu /bin/true sanity check failed; continuing to GPU workload checks" >&2 fi -smoke_status=0 +check_status=0 -# Run both workload smokes even if the first one fails. That gives the uploaded +# Run both workload checks even if the first one fails. That gives the uploaded # race-reports artifact a complete picture for the failing CI attempt instead of # forcing a second long run just to learn whether the other workload also broke. set +e -run_timed "hipblaslt-bench race smoke" run_hipblaslt_bench_smoke +run_timed "hipblaslt-bench race check" run_hipblaslt_bench_check hipblaslt_status=$? -run_timed "tensilelite-client race smoke" run_tensilelite_client_smoke +run_timed "tensilelite-client race check" run_tensilelite_client_check tensilelite_status=$? set -e if [[ "${hipblaslt_status}" -ne 0 ]]; then - echo "hipblaslt-bench race smoke failed with status ${hipblaslt_status}" >&2 - smoke_status=1 + echo "hipblaslt-bench race check failed with status ${hipblaslt_status}" >&2 + check_status=1 fi if [[ "${tensilelite_status}" -ne 0 ]]; then - echo "tensilelite-client race smoke failed with status ${tensilelite_status}" >&2 - smoke_status=1 + echo "tensilelite-client race check failed with status ${tensilelite_status}" >&2 + check_status=1 fi -if [[ "${smoke_status}" -ne 0 ]]; then - echo "one or more rocjitsu race smokes failed" >&2 - exit "${smoke_status}" +if [[ "${check_status}" -ne 0 ]]; then + echo "one or more rocjitsu race checks failed" >&2 + exit "${check_status}" fi -echo "rocjitsu race smokes completed without race reports" +echo "rocjitsu race checks completed without race reports" diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index cc458d73e5e4..19b769875389 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -172,8 +172,8 @@ jobs: test_type: ${{ inputs.test_type }} secrets: inherit - rocjitsu-race-smoke-linux: - name: rocjitsu race smoke (${{ inputs.amdgpu_families }}) + rocjitsu-race-check-linux: + name: rocjitsu race check (${{ inputs.amdgpu_families }}) needs: [therock-build-linux] # This is a narrow post-build instrumentation job for hipBLASLt on gfx950. # Normal package tests cover broad functionality; this job fetches the BLAS @@ -200,12 +200,14 @@ jobs: ROCJITSU_SOURCE_DIR: ${{ github.workspace }}/rocm-systems/emulation/rocjitsu ROCJITSU_BUILD_DIR: ${{ github.workspace }}/rocjitsu-build RACE_REPORT_DIR: ${{ github.workspace }}/race-reports + # Per-workload timeouts are below the job timeout so the script can report + # which workload failed and still upload race-reports on exit. RACE_TIMEOUT_SECONDS: "180" TENSILELITE_TIMEOUT_SECONDS: "420" steps: # The job needs three inputs in one workspace: # 1. TheRock actions/scripts to fetch the exact artifacts from this run. - # 2. This rocm-libraries checkout for the race-smoke driver script. + # 2. This rocm-libraries checkout for the race-check driver script. # 3. rocm-systems source for rocjitsu until rocjitsu is packaged as a # complete runnable artifact. - name: Checkout TheRock repository @@ -245,7 +247,7 @@ jobs: ninja-build # Pull the BLAS runtime and test payload from the current TheRock build. - # This is what keeps the smoke aligned with the exact hipBLASLt and + # This is what keeps the check aligned with the exact hipBLASLt and # TensileLite artifacts produced by the PR, rather than an image-local or # previously installed ROCm tree. - name: Run setup test environment workflow @@ -279,9 +281,9 @@ jobs: print("rocisa:", rocisa.__file__) PY - - name: Run rocjitsu race smoke + - name: Run rocjitsu race check timeout-minutes: 75 - run: bash rocm-libraries/.github/scripts/run_rocjitsu_hipblaslt_race_smoke.sh + run: bash rocm-libraries/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh # Upload even on failure. The report directory contains the per-workload # stdout logs, rocjitsu sink files, race.log files, generated reduced From 0bca60698b167050087fcf0edae21a9a9625ffc7 Mon Sep 17 00:00:00 2001 From: James Newling Date: Fri, 10 Jul 2026 11:03:19 -0700 Subject: [PATCH 06/19] comment adjustments --- .github/scripts/run_rocjitsu_hipblaslt_race_check.sh | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index 2a40344d21f6..d64f4155a4fa 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -1,13 +1,14 @@ #!/usr/bin/env bash set -euo pipefail -# This script's calling workflow fetches hipBLASLt/TensileLite artifacts from +# This workflow that calls this script fetches hipBLASLt/TensileLite artifacts from # the current TheRock build, then checks out rocm-systems. # # TODO(newling) Until rocjitsu is packaged as a complete runnable TheRock -# artifact, we build the rocjitsu CLI locally. +# artifact, we build the rocjitsu CLI locally. Monitor progress on packaging rocjitsu. # # The script runs small gfx950 hipBLASLt and TensileLite GEMMs under the race detector. +# TODO(newling) extend to different architectures and expand GEMM-space tested. ROCM_PATH="${ROCM_PATH:-${PWD}/build}" ROCJITSU_SOURCE_DIR="${ROCJITSU_SOURCE_DIR:-${PWD}/rocm-systems/emulation/rocjitsu}" @@ -96,10 +97,7 @@ if [[ ! -x "${TENSILELITE_CLIENT}" ]]; then exit 1 fi -# Use the gfx950/CDNA4 KMD config from the checked-out rocm-systems tree. The -# workflow pins the rocm-systems ref; if that ref renames configs, update this -# list or pass ROCJITSU_CONFIG from the workflow. Keep the selection explicit so -# future architectures can add their own config candidates here. + if [[ -z "${ROCJITSU_CONFIG}" ]]; then for candidate in \ "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" \ From 5603f2573ab8d7664b5caa6d9f348f05cf8cc2c7 Mon Sep 17 00:00:00 2001 From: James Newling Date: Fri, 10 Jul 2026 15:16:37 -0700 Subject: [PATCH 07/19] ci(hipblaslt): use rocm-systems develop for rocjitsu --- .github/scripts/run_rocjitsu_hipblaslt_race_check.sh | 4 ++-- .github/workflows/therock-ci-linux.yml | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index d64f4155a4fa..bba50833f8a4 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -140,8 +140,8 @@ echo "PYTHONPATH=${PYTHONPATH}" # TODO(newling): Track migration to a packaged rocjitsu once TheRock provides a -# complete runnable artifact. Until then, build rocjitsu from the pinned -# rocm-systems checkout so this job controls the tool version. +# complete runnable artifact. Until then, build rocjitsu from the rocm-systems +# checkout selected by the workflow so this job controls the tool source. # rocjitsu is still consumed from a source checkout in this workflow. Build only # the CLI and optional runtime/shim targets needed to launch the test workloads; # this keeps the job independent of full rocm-systems packaging. diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index 19b769875389..aba41ceaf3b6 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -224,8 +224,8 @@ jobs: - name: Checkout rocm-systems for rocjitsu uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 with: - repository: newling/rocm-systems - ref: rocjitsu-early-mmap-fix + repository: ROCm/rocm-systems + ref: develop path: rocm-systems sparse-checkout: | emulation/rocjitsu From af599c1514c75708b12ef64109c68b40d032d4b8 Mon Sep 17 00:00:00 2001 From: James Newling Date: Mon, 13 Jul 2026 13:14:21 -0700 Subject: [PATCH 08/19] ci(hipblaslt): explore CPU rocjitsu race-check targets --- .../run_rocjitsu_hipblaslt_race_check.sh | 96 ++++++++++++++++--- .github/workflows/therock-ci-linux.yml | 13 ++- .github/workflows/therock-ci.yml | 7 +- 3 files changed, 95 insertions(+), 21 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index bba50833f8a4..427655246908 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -7,10 +7,12 @@ set -euo pipefail # TODO(newling) Until rocjitsu is packaged as a complete runnable TheRock # artifact, we build the rocjitsu CLI locally. Monitor progress on packaging rocjitsu. # -# The script runs small gfx950 hipBLASLt and TensileLite GEMMs under the race detector. -# TODO(newling) extend to different architectures and expand GEMM-space tested. +# The script runs small hipBLASLt and TensileLite GEMMs under the race detector. +# TODO(newling) expand the GEMM-space tested. ROCM_PATH="${ROCM_PATH:-${PWD}/build}" +AMDGPU_FAMILIES="${AMDGPU_FAMILIES:-}" +ROCJITSU_GPU_TARGET="${ROCJITSU_GPU_TARGET:-}" ROCJITSU_SOURCE_DIR="${ROCJITSU_SOURCE_DIR:-${PWD}/rocm-systems/emulation/rocjitsu}" ROCJITSU_BUILD_DIR="${ROCJITSU_BUILD_DIR:-${PWD}/rocjitsu-build}" ROCJITSU_CONFIG="${ROCJITSU_CONFIG:-}" @@ -19,10 +21,77 @@ HIPBLASLT_BENCH="${HIPBLASLT_BENCH:-${ROCM_PATH}/bin/hipblaslt-bench}" TENSILELITE_ROOT="${TENSILELITE_ROOT:-${ROCM_PATH}/share/hipblaslt/tensilelite}" TENSILE_DRIVER="${TENSILE_DRIVER:-${TENSILELITE_ROOT}/Tensile/bin/Tensile}" TENSILELITE_CLIENT="${TENSILELITE_CLIENT:-${ROCM_PATH}/libexec/hipblaslt/tensilelite/tensilelite-client}" +TENSILELITE_GPU_TARGET="${TENSILELITE_GPU_TARGET:-}" RACE_TIMEOUT_SECONDS="${RACE_TIMEOUT_SECONDS:-180}" TENSILELITE_TIMEOUT_SECONDS="${TENSILELITE_TIMEOUT_SECONDS:-420}" TIMING_FILE="${RACE_REPORT_DIR}/timing.tsv" +derive_rocjitsu_gpu_target() { + if [[ -n "${ROCJITSU_GPU_TARGET}" ]]; then + return + fi + + case "${AMDGPU_FAMILIES}" in + gfx94*) + ROCJITSU_GPU_TARGET="gfx942" + ;; + gfx950*) + ROCJITSU_GPU_TARGET="gfx950" + ;; + gfx1151) + ROCJITSU_GPU_TARGET="gfx1151" + ;; + *) + echo "Unsupported rocjitsu race-check artifact group: ${AMDGPU_FAMILIES}" >&2 + echo "Supported groups are gfx94*, gfx950*, and gfx1151." >&2 + exit 1 + ;; + esac +} + +select_rocjitsu_config() { + if [[ -n "${ROCJITSU_CONFIG}" ]]; then + return + fi + + local candidates=() + case "${ROCJITSU_GPU_TARGET}" in + gfx942) + candidates=( + "${ROCJITSU_SOURCE_DIR}/configs/gfx942_cdna3_kmd.json" + "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna3_kmd.json" + "${ROCJITSU_SOURCE_DIR}/configs/gfx942_cdna3.json" + "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna3.json" + ) + ;; + gfx950) + candidates=( + "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" + "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna4_kmd.json" + "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4.json" + "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna4.json" + ) + ;; + gfx1151) + candidates=( + "${ROCJITSU_SOURCE_DIR}/configs/gfx1151.json" + ) + ;; + *) + echo "Unsupported rocjitsu target: ${ROCJITSU_GPU_TARGET}" >&2 + exit 1 + ;; + esac + + local candidate + for candidate in "${candidates[@]}"; do + if [[ -f "${candidate}" ]]; then + ROCJITSU_CONFIG="${candidate}" + return + fi + done +} + # Record the duration and exit status of each stage in a file that is uploaded # with the other race reports. The same information is also printed at exit. print_timing_summary() { @@ -98,22 +167,18 @@ if [[ ! -x "${TENSILELITE_CLIENT}" ]]; then fi -if [[ -z "${ROCJITSU_CONFIG}" ]]; then - for candidate in \ - "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" \ - "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna4_kmd.json"; do - if [[ -f "${candidate}" ]]; then - ROCJITSU_CONFIG="${candidate}" - break - fi - done -fi +derive_rocjitsu_gpu_target +select_rocjitsu_config if [[ -z "${ROCJITSU_CONFIG}" || ! -f "${ROCJITSU_CONFIG}" ]]; then - echo "rocjitsu gfx950 config not found under ${ROCJITSU_SOURCE_DIR}/configs" >&2 + echo "rocjitsu ${ROCJITSU_GPU_TARGET} config not found under ${ROCJITSU_SOURCE_DIR}/configs" >&2 exit 1 fi +if [[ -z "${TENSILELITE_GPU_TARGET}" ]]; then + TENSILELITE_GPU_TARGET="${ROCJITSU_GPU_TARGET}" +fi + mkdir -p "${RACE_REPORT_DIR}" : >"${TIMING_FILE}" @@ -126,6 +191,8 @@ export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROC export PYTHONPATH="${TENSILELITE_ROOT}${PYTHONPATH:+:${PYTHONPATH}}" echo "ROCM_PATH=${ROCM_PATH}" +echo "AMDGPU_FAMILIES=${AMDGPU_FAMILIES}" +echo "ROCJITSU_GPU_TARGET=${ROCJITSU_GPU_TARGET}" echo "ROCJITSU_SOURCE_DIR=${ROCJITSU_SOURCE_DIR}" echo "ROCJITSU_BUILD_DIR=${ROCJITSU_BUILD_DIR}" echo "ROCJITSU_CONFIG=${ROCJITSU_CONFIG}" @@ -133,6 +200,7 @@ echo "HIPBLASLT_BENCH=${HIPBLASLT_BENCH}" echo "TENSILELITE_ROOT=${TENSILELITE_ROOT}" echo "TENSILE_DRIVER=${TENSILE_DRIVER}" echo "TENSILELITE_CLIENT=${TENSILELITE_CLIENT}" +echo "TENSILELITE_GPU_TARGET=${TENSILELITE_GPU_TARGET}" echo "RACE_REPORT_DIR=${RACE_REPORT_DIR}" echo "PATH=${PATH}" echo "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}" @@ -343,7 +411,7 @@ run_tensilelite_client_check() { "${yaml}" "${output_dir}" --prebuilt-client "${TENSILELITE_CLIENT}" - --gpu-targets gfx950 + --gpu-targets "${TENSILELITE_GPU_TARGET}" --library-format msgpack ) if [[ -x "${ROCM_PATH}/bin/amdclang++" ]]; then diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index aba41ceaf3b6..fa8b708025be 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -161,7 +161,7 @@ jobs: name: Test (${{ inputs.amdgpu_families }}) needs: [therock-build-linux] # TODO(TheRock#3288): Re-enable gfx950-dcgpu runners once we get more capacity - if: ${{ inputs.amdgpu_families != 'gfx950-dcgpu' }} + if: ${{ inputs.test_runs_on != '' && inputs.amdgpu_families != 'gfx950-dcgpu' }} uses: ./.github/workflows/therock-test-packages.yml with: therock_ref: ${{ inputs.therock_ref }} @@ -175,13 +175,15 @@ jobs: rocjitsu-race-check-linux: name: rocjitsu race check (${{ inputs.amdgpu_families }}) needs: [therock-build-linux] - # This is a narrow post-build instrumentation job for hipBLASLt on gfx950. + # This is a narrow post-build instrumentation job for hipBLASLt. # Normal package tests cover broad functionality; this job fetches the BLAS # artifacts from the just-finished TheRock build and runs two small workloads - # under rocjitsu's race detector. Keep it gated to hipblaslt + gfx950 while + # under rocjitsu's race detector. The workloads run entirely under CPU-side + # emulation; use the CPU build runner instead of consuming GPU test capacity. + # Keep it gated to hipblaslt + known rocjitsu race-check targets while # rocjitsu packaging and runner capacity are still evolving. - if: ${{ inputs.test_runs_on != '' && startsWith(inputs.amdgpu_families, 'gfx950') && contains(format(',{0},', inputs.projects_to_test), ',hipblaslt,') }} - runs-on: ${{ inputs.test_runs_on }} + if: ${{ inputs.build_runs_on != '' && (startsWith(inputs.amdgpu_families, 'gfx94') || startsWith(inputs.amdgpu_families, 'gfx950') || inputs.amdgpu_families == 'gfx1151') && contains(format(',{0},', inputs.projects_to_test), ',hipblaslt,') }} + runs-on: ${{ inputs.build_runs_on }} timeout-minutes: 90 container: image: ghcr.io/rocm/no_rocm_image_ubuntu24_04@sha256:4150afe4759d14822f0e3f8930e1124f26e11f68b5c7b91ec9a02b20b1ebbb98 @@ -197,6 +199,7 @@ jobs: ARTIFACT_RUN_ID: ${{ github.run_id }} OUTPUT_ARTIFACTS_DIR: ./build ROCM_PATH: ${{ github.workspace }}/build + AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }} ROCJITSU_SOURCE_DIR: ${{ github.workspace }}/rocm-systems/emulation/rocjitsu ROCJITSU_BUILD_DIR: ${{ github.workspace }}/rocjitsu-build RACE_REPORT_DIR: ${{ github.workspace }}/race-reports diff --git a/.github/workflows/therock-ci.yml b/.github/workflows/therock-ci.yml index c2379f99d73d..3f374b16e925 100644 --- a/.github/workflows/therock-ci.yml +++ b/.github/workflows/therock-ci.yml @@ -126,7 +126,7 @@ jobs: - name: Fetch Linux targets for build and test env: THEROCK_PACKAGE_PLATFORM: "linux" - AMDGPU_FAMILIES: ${{ inputs.linux_amdgpu_families || 'gfx94X, gfx950' }} + AMDGPU_FAMILIES: ${{ inputs.linux_amdgpu_families || 'gfx94X, gfx950, gfx1151' }} CI_CONFIG_PATH: ci-config id: configure_linux run: python ./TheRock/build_tools/github_actions/fetch_package_targets.py @@ -161,7 +161,10 @@ jobs: projects_to_test: ${{ matrix.projects.projects_to_test }} test_type: ${{ needs.setup.outputs.test_type }} amdgpu_families: ${{ matrix.target_bundle.amdgpu_family }} - test_runs_on: ${{ matrix.target_bundle.test_machine }} + # The rocjitsu race check runs under CPU-only emulation. For gfx1151, + # avoid scheduling normal package tests on a gfx1151 GPU runner while we + # are using that target to validate the race-check sidecar. + test_runs_on: ${{ matrix.target_bundle.amdgpu_family == 'gfx1151' && '' || matrix.target_bundle.test_machine }} build_runs_on: ${{ needs.setup.outputs.linux_build_runs_on }} therock-ci-windows: From cf0be640f072602f248b6f69d21b71c2d961fa5d Mon Sep 17 00:00:00 2001 From: James Newling Date: Mon, 13 Jul 2026 15:05:55 -0700 Subject: [PATCH 09/19] ci(therock): suppress normal gfx1151 Linux tests --- .github/workflows/therock-ci.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/therock-ci.yml b/.github/workflows/therock-ci.yml index 3f374b16e925..614db7983f99 100644 --- a/.github/workflows/therock-ci.yml +++ b/.github/workflows/therock-ci.yml @@ -164,7 +164,7 @@ jobs: # The rocjitsu race check runs under CPU-only emulation. For gfx1151, # avoid scheduling normal package tests on a gfx1151 GPU runner while we # are using that target to validate the race-check sidecar. - test_runs_on: ${{ matrix.target_bundle.amdgpu_family == 'gfx1151' && '' || matrix.target_bundle.test_machine }} + test_runs_on: ${{ matrix.target_bundle.amdgpu_family != 'gfx1151' && matrix.target_bundle.test_machine || '' }} build_runs_on: ${{ needs.setup.outputs.linux_build_runs_on }} therock-ci-windows: From cdcd49a2b5c0d7d50016bc2cba5fb422b2467ea6 Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 14 Jul 2026 09:15:27 -0700 Subject: [PATCH 10/19] ci(rocjitsu): use sudo for race-check deps --- .github/workflows/therock-ci-linux.yml | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index fa8b708025be..e49fcb2c324a 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -180,8 +180,7 @@ jobs: # artifacts from the just-finished TheRock build and runs two small workloads # under rocjitsu's race detector. The workloads run entirely under CPU-side # emulation; use the CPU build runner instead of consuming GPU test capacity. - # Keep it gated to hipblaslt + known rocjitsu race-check targets while - # rocjitsu packaging and runner capacity are still evolving. + # Gate on hipBLASLt and artifact families with checked-in rocjitsu configs. if: ${{ inputs.build_runs_on != '' && (startsWith(inputs.amdgpu_families, 'gfx94') || startsWith(inputs.amdgpu_families, 'gfx950') || inputs.amdgpu_families == 'gfx1151') && contains(format(',{0},', inputs.projects_to_test), ',hipblaslt,') }} runs-on: ${{ inputs.build_runs_on }} timeout-minutes: 90 @@ -235,13 +234,13 @@ jobs: shared/machine-readable-isa/isa sparse-checkout-cone-mode: true - # The base image intentionally does not contain a full ROCm install. Keep - # the apt dependency set to ordinary host build tools; the ROCm runtime, - # clients, and TensileLite Python tree come from TheRock artifacts below. + # The base image contains the shell and source-control tools used before + # artifact fetch; install only the host libraries needed for the local + # rocjitsu source build. ROCm runtime pieces come from TheRock artifacts. - name: Install rocjitsu build dependencies run: | - apt-get update - apt-get install -y --no-install-recommends \ + sudo apt-get update + sudo apt-get install -y --no-install-recommends \ build-essential \ ca-certificates \ cmake \ From e898481809fd1283509d43776b722c7beefeffd3 Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 14 Jul 2026 11:30:44 -0700 Subject: [PATCH 11/19] ci(rocjitsu): adjust gfx1151 TensileLite smoke --- .github/scripts/run_rocjitsu_hipblaslt_race_check.sh | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index 427655246908..520e6fed10a9 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -335,11 +335,19 @@ run_hipblaslt_bench_check() { write_tensilelite_check_yaml() { local yaml="$1" + local direct_to_vgpr_a=True + if [[ "${TENSILELITE_GPU_TARGET}" == "gfx1151" ]]; then + # gfx1151 rejects this TLUA + f32 + MIInputPerThread > 1 combination when + # DirectToVgprA is enabled. Keep the rest of the reduced problem identical + # so the target still exercises the standalone TensileLite client path. + direct_to_vgpr_a=False + fi + # One reduced config is embedded here to keep this bridge job self-contained. # If the race check grows to multiple configs, move them into a small checked-in # data directory or install them with the TensileLite test artifacts, then have # this script iterate over that list. - cat >"${yaml}" <<'YAML' + cat >"${yaml}" < Date: Tue, 14 Jul 2026 11:39:38 -0700 Subject: [PATCH 12/19] ci(rocjitsu): print TensileLite rejection reasons --- .github/scripts/run_rocjitsu_hipblaslt_race_check.sh | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index 520e6fed10a9..70ef2846e9b2 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -354,6 +354,7 @@ GlobalParameters: DataInitTypeAlpha: 1 Device: 0 CpuThreads: 1 + PrintSolutionRejectionReason: True BenchmarkProblems: - From 0ad39293a02ebd6e62421528e40482830c69b740 Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 14 Jul 2026 13:28:31 -0700 Subject: [PATCH 13/19] ci(rocjitsu): use gfx1151 VALU Tensile smoke --- .../run_rocjitsu_hipblaslt_race_check.sh | 61 +++++++++++++++++-- 1 file changed, 55 insertions(+), 6 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index 70ef2846e9b2..37bcd6c449b5 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -335,12 +335,61 @@ run_hipblaslt_bench_check() { write_tensilelite_check_yaml() { local yaml="$1" - local direct_to_vgpr_a=True + if [[ "${TENSILELITE_GPU_TARGET}" == "gfx1151" ]]; then - # gfx1151 rejects this TLUA + f32 + MIInputPerThread > 1 combination when - # DirectToVgprA is enabled. Keep the rest of the reduced problem identical - # so the target still exercises the standalone TensileLite client path. - direct_to_vgpr_a=False + # gfx1151 f32 does not support the WMMA instruction shape used by the + # gfx94x/gfx950 smoke below. Use a small VALU assembly config that still + # exercises the Tensile front end, kernel writer, and standalone client. + cat >"${yaml}" <<'YAML' +GlobalParameters: + NumElementsToValidate: -1 + DataInitTypeBeta: 0 + DataInitTypeAlpha: 1 + Device: 0 + CpuThreads: 1 + PrintSolutionRejectionReason: True + +BenchmarkProblems: + - + - OperationType: GEMM + DataType: s + TransposeA: False + TransposeB: True + UseBeta: True + Batched: True + - InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - ThreadTile: + - [1, 1] + - WorkGroup: + - [8, 8, 1] + - DepthU: [8] + - VectorWidthA: [1] + - VectorWidthB: [1] + - GlobalReadVectorWidthA: [1] + - GlobalReadVectorWidthB: [1] + - LocalReadVectorWidth: [1] + - PrefetchGlobalRead: [1] + - PrefetchLocalRead: [0] + - ScheduleIterAlg: [1] + - StaggerU: [32] + - StaggerUStride: [256] + - WorkGroupMapping: [1] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: ["MultipleBuffer"] + - StoreRemapVectorWidth: [0] + - StoreVectorWidth: [1] + - SourceSwap: [False] + - NumElementsPerBatchStore: [0] + - ClusterLocalRead: [0] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [32, 32, 1, 32] +YAML + return fi # One reduced config is embedded here to keep this bridge job self-contained. @@ -398,7 +447,7 @@ BenchmarkProblems: - SourceSwap: [True] - NumElementsPerBatchStore: [16] - ClusterLocalRead: [1] - - DirectToVgprA: [${direct_to_vgpr_a}] + - DirectToVgprA: [True] - DirectToVgprB: [False] - WorkGroup: [[32, 4, 4]] BenchmarkJoinParameters: From 43d1b7f134432da2eda830b56c49601b4e981653 Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 14 Jul 2026 16:35:44 -0700 Subject: [PATCH 14/19] ci(rocjitsu): make gfx1151 TensileLite a weak smoke --- .../run_rocjitsu_hipblaslt_race_check.sh | 48 +++++++++++++++---- 1 file changed, 38 insertions(+), 10 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index 37bcd6c449b5..9530c88ee62b 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -338,13 +338,34 @@ write_tensilelite_check_yaml() { if [[ "${TENSILELITE_GPU_TARGET}" == "gfx1151" ]]; then # gfx1151 f32 does not support the WMMA instruction shape used by the - # gfx94x/gfx950 smoke below. Use a small VALU assembly config that still - # exercises the Tensile front end, kernel writer, and standalone client. + # gfx94x/gfx950 smoke below. Use a small VALU assembly config. + # + # Keep this as a weak smoke for now: running the benchmark path currently + # trips a suspected rocjitsu false positive in the TensileLite client helper + # kernels (`global_load_d16_u8` destination-merge reads racing with the same + # instruction's outstanding VGPR write). `NumBenchmarks: 0` still exercises + # the Tensile front end, solution filtering, KernelWriter, code-object build, + # rocjitsu launch, and standalone client startup. The real executed GEMM + # coverage for gfx1151 remains the hipblaslt-bench check above until the + # rocjitsu D16 VGPR self-race is fixed. cat >"${yaml}" <<'YAML' GlobalParameters: - NumElementsToValidate: -1 - DataInitTypeBeta: 0 + NumBenchmarks: 0 + NumElementsToValidate: 0 + DataInitTypeAB: 0 + DataInitTypeA: 0 + DataInitTypeB: 0 + DataInitTypeC: 0 + DataInitTypeD: 0 + DataInitTypeE: 0 + DataInitTypeBias: 0 + DataInitTypeScaleA: 0 + DataInitTypeScaleB: 0 + DataInitTypeScaleC: 0 + DataInitTypeScaleD: 0 + DataInitTypeScaleAlphaVec: 0 DataInitTypeAlpha: 1 + DataInitTypeBeta: 0 Device: 0 CpuThreads: 1 PrintSolutionRejectionReason: True @@ -498,14 +519,21 @@ run_tensilelite_client_check() { return "${status}" fi - if ! grep -q "PASSED" "${RACE_REPORT_DIR}/tensilelite-client.log"; then - echo "tensilelite-client race check did not report validation success" >&2 - return 1 + if [[ "${TENSILELITE_GPU_TARGET}" == "gfx1151" ]]; then + if ! grep -q "clientExit=0 (PASS)" "${RACE_REPORT_DIR}/tensilelite-client.log"; then + echo "tensilelite-client smoke did not report a clean client exit" >&2 + return 1 + fi + else + if ! grep -q "PASSED" "${RACE_REPORT_DIR}/tensilelite-client.log"; then + echo "tensilelite-client race check did not report validation success" >&2 + return 1 + fi fi - # As above, validation success and race-detector success are distinct signals. - # Require both: the client must pass numerics, and rocjitsu must leave the race - # sink free of RACE records. + # As above, client success and race-detector success are distinct signals. + # Require both: the client must finish cleanly, and rocjitsu must leave the + # race sink free of RACE records. if [[ -f "${sink_dir}/race.log" ]] && grep -q '^RACE ' "${sink_dir}/race.log"; then echo "rocjitsu race detector reported a race in tensilelite-client:" >&2 cat "${sink_dir}/race.log" >&2 From c9c480bf31fb19fd085767df8099a566edcd9de5 Mon Sep 17 00:00:00 2001 From: James Newling Date: Wed, 15 Jul 2026 13:31:44 -0700 Subject: [PATCH 15/19] ci(rocjitsu): split race-check sidecar workflow --- .../run_rocjitsu_hipblaslt_race_check.sh | 75 +++++----- .github/workflows/therock-ci-linux.yml | 130 ++---------------- .../therock-rocjitsu-race-check-linux.yml | 124 +++++++++++++++++ 3 files changed, 166 insertions(+), 163 deletions(-) create mode 100644 .github/workflows/therock-rocjitsu-race-check-linux.yml diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index 9530c88ee62b..ab11b3f7db18 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -1,15 +1,28 @@ #!/usr/bin/env bash set -euo pipefail -# This workflow that calls this script fetches hipBLASLt/TensileLite artifacts from -# the current TheRock build, then checks out rocm-systems. +# The workflow that calls this script fetches hipBLASLt/TensileLite artifacts +# from the current TheRock build, then checks out rocm-systems. # # TODO(newling) Until rocjitsu is packaged as a complete runnable TheRock # artifact, we build the rocjitsu CLI locally. Monitor progress on packaging rocjitsu. # # The script runs small hipBLASLt and TensileLite GEMMs under the race detector. # TODO(newling) expand the GEMM-space tested. - +# +# Basic flow: +# 1. Use the TheRock artifact tree unpacked at ROCM_PATH. +# 2. Build rocjitsu from the rocm-systems checkout in ROCJITSU_SOURCE_DIR. +# 3. Select a rocjitsu config for gfx942, gfx950, or gfx1151. +# 4. Run hipblaslt-bench and a reduced TensileLite smoke under RJ_RACE=1. + +# These defaults match the GitHub Actions workspace layout: ROCM_PATH is the +# unpacked TheRock artifact tree, ROCJITSU_SOURCE_DIR is the checked-out +# rocm-systems source tree, ROCJITSU_BUILD_DIR is a local build directory, and +# RACE_REPORT_DIR is uploaded at the end of the job. They can be overridden for +# local reproduction. The brittle part is not the path names themselves, but the +# artifact layout underneath ROCM_PATH; the checks below fail early if +# hipBLASLt/TensileLite files move in TheRock artifacts. ROCM_PATH="${ROCM_PATH:-${PWD}/build}" AMDGPU_FAMILIES="${AMDGPU_FAMILIES:-}" ROCJITSU_GPU_TARGET="${ROCJITSU_GPU_TARGET:-}" @@ -57,6 +70,9 @@ select_rocjitsu_config() { local candidates=() case "${ROCJITSU_GPU_TARGET}" in gfx942) + # Four names cover two rocm-systems config naming choices that are still + # in flight: target-specific vs family-generic, and KMD vs non-KMD. + # Prefer the target-specific KMD config when it exists. candidates=( "${ROCJITSU_SOURCE_DIR}/configs/gfx942_cdna3_kmd.json" "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna3_kmd.json" @@ -65,6 +81,8 @@ select_rocjitsu_config() { ) ;; gfx950) + # Same four-name compatibility pattern as gfx942: target-specific before + # family-generic, and KMD before non-KMD. candidates=( "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna4_kmd.json" @@ -240,9 +258,9 @@ cmake_args+=( run_timed "configure rocjitsu" cmake "${cmake_args[@]}" -# rocm-systems has been renaming/splitting rocjitsu build targets while this -# bridge job is being developed. Probe for optional targets instead of assuming -# one exact target set, so the CI job survives small target-layout changes. +# The CLI is required. The shared library and KMD shim target names have existed +# in some rocm-systems revisions and not others; build them when available, but +# do not make this bridge job depend on optional target names. cmake_target_exists() { local target="$1" cmake --build "${ROCJITSU_BUILD_DIR}" --target help \ @@ -269,28 +287,6 @@ show_rocjitsu_version() { "${ROCJITSU_BIN}" --version } -# This is a cheap launch-path check for rocjitsu itself. Treat it as diagnostic -# rather than authoritative: the real signal is whether HIP/HSA workloads from -# the fetched ROCm payload run and whether their race logs contain reports. -run_sanity_check() { - echo "rocjitsu /bin/true sanity check:" - mkdir -p "${RACE_REPORT_DIR}/sanity" - timeout 30 \ - env \ - RJ_LOG=1 \ - RJ_SINKS=stderr,file \ - RJ_SINK_DIR="${RACE_REPORT_DIR}/sanity" \ - "${ROCJITSU_BIN}" \ - --config "${ROCJITSU_CONFIG}" \ - -- /bin/true \ - 2>&1 | tee "${RACE_REPORT_DIR}/sanity.log" - local status=$? - if [[ "${status}" -ne 0 ]]; then - echo "rocjitsu /bin/true sanity check failed with status ${status}" >&2 - return "${status}" - fi -} - run_hipblaslt_bench_check() { mkdir -p "${RACE_REPORT_DIR}/hipblaslt-bench" echo "running hipblaslt-bench under rocjitsu race detection" @@ -343,11 +339,16 @@ write_tensilelite_check_yaml() { # Keep this as a weak smoke for now: running the benchmark path currently # trips a suspected rocjitsu false positive in the TensileLite client helper # kernels (`global_load_d16_u8` destination-merge reads racing with the same - # instruction's outstanding VGPR write). `NumBenchmarks: 0` still exercises - # the Tensile front end, solution filtering, KernelWriter, code-object build, - # rocjitsu launch, and standalone client startup. The real executed GEMM - # coverage for gfx1151 remains the hipblaslt-bench check above until the - # rocjitsu D16 VGPR self-race is fixed. + # instruction's outstanding VGPR write). ROCm/rocm-systems#8628 fixes that + # in rocjitsu. Until that fix is available in this CI path, + # `NumBenchmarks: 0` still exercises the Tensile front end, solution + # filtering, KernelWriter, code-object build, rocjitsu launch, and + # standalone client startup. The real executed GEMM coverage for gfx1151 + # remains the hipblaslt-bench check above. + # + # TODO(newling): Move these inline YAMLs into checked-in TensileLite test + # data, or reuse existing TensileLite YAMLs once the component path can + # consume product-owned workload definitions. cat >"${yaml}" <<'YAML' GlobalParameters: NumBenchmarks: 0 @@ -543,14 +544,6 @@ run_tensilelite_client_check() { run_timed "rocjitsu version" show_rocjitsu_version -set +e -run_timed "rocjitsu sanity check" run_sanity_check -sanity_status=$? -set -e -if [[ "${sanity_status}" -ne 0 ]]; then - echo "WARNING: rocjitsu /bin/true sanity check failed; continuing to GPU workload checks" >&2 -fi - check_status=0 # Run both workload checks even if the first one fails. That gives the uploaded diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index e49fcb2c324a..b870a2d47e84 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -175,126 +175,12 @@ jobs: rocjitsu-race-check-linux: name: rocjitsu race check (${{ inputs.amdgpu_families }}) needs: [therock-build-linux] - # This is a narrow post-build instrumentation job for hipBLASLt. - # Normal package tests cover broad functionality; this job fetches the BLAS - # artifacts from the just-finished TheRock build and runs two small workloads - # under rocjitsu's race detector. The workloads run entirely under CPU-side - # emulation; use the CPU build runner instead of consuming GPU test capacity. - # Gate on hipBLASLt and artifact families with checked-in rocjitsu configs. + # Run the narrow rocjitsu sidecar only for hipBLASLt rows and GPU families + # that rocjitsu can emulate here. if: ${{ inputs.build_runs_on != '' && (startsWith(inputs.amdgpu_families, 'gfx94') || startsWith(inputs.amdgpu_families, 'gfx950') || inputs.amdgpu_families == 'gfx1151') && contains(format(',{0},', inputs.projects_to_test), ',hipblaslt,') }} - runs-on: ${{ inputs.build_runs_on }} - timeout-minutes: 90 - container: - image: ghcr.io/rocm/no_rocm_image_ubuntu24_04@sha256:4150afe4759d14822f0e3f8930e1124f26e11f68b5c7b91ec9a02b20b1ebbb98 - options: --ipc host - --user 0:0 - --ulimit memlock=-1:-1 - --security-opt seccomp=unconfined - defaults: - run: - shell: bash - env: - VENV_DIR: ${{ github.workspace }}/.venv - ARTIFACT_RUN_ID: ${{ github.run_id }} - OUTPUT_ARTIFACTS_DIR: ./build - ROCM_PATH: ${{ github.workspace }}/build - AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }} - ROCJITSU_SOURCE_DIR: ${{ github.workspace }}/rocm-systems/emulation/rocjitsu - ROCJITSU_BUILD_DIR: ${{ github.workspace }}/rocjitsu-build - RACE_REPORT_DIR: ${{ github.workspace }}/race-reports - # Per-workload timeouts are below the job timeout so the script can report - # which workload failed and still upload race-reports on exit. - RACE_TIMEOUT_SECONDS: "180" - TENSILELITE_TIMEOUT_SECONDS: "420" - steps: - # The job needs three inputs in one workspace: - # 1. TheRock actions/scripts to fetch the exact artifacts from this run. - # 2. This rocm-libraries checkout for the race-check driver script. - # 3. rocm-systems source for rocjitsu until rocjitsu is packaged as a - # complete runnable artifact. - - name: Checkout TheRock repository - uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 - with: - repository: "ROCm/TheRock" - ref: ${{ inputs.therock_ref }} - - - name: Checkout rocm-libraries repository for scripts - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 - with: - path: rocm-libraries - - - name: Checkout rocm-systems for rocjitsu - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 - with: - repository: ROCm/rocm-systems - ref: develop - path: rocm-systems - sparse-checkout: | - emulation/rocjitsu - shared/machine-readable-isa/isa - sparse-checkout-cone-mode: true - - # The base image contains the shell and source-control tools used before - # artifact fetch; install only the host libraries needed for the local - # rocjitsu source build. ROCm runtime pieces come from TheRock artifacts. - - name: Install rocjitsu build dependencies - run: | - sudo apt-get update - sudo apt-get install -y --no-install-recommends \ - build-essential \ - ca-certificates \ - cmake \ - git \ - libdrm-dev \ - ninja-build - - # Pull the BLAS runtime and test payload from the current TheRock build. - # This is what keeps the check aligned with the exact hipBLASLt and - # TensileLite artifacts produced by the PR, rather than an image-local or - # previously installed ROCm tree. - - name: Run setup test environment workflow - timeout-minutes: 15 - uses: ./.github/actions/setup_test_environment - with: - ARTIFACT_RUN_ID: ${{ env.ARTIFACT_RUN_ID }} - ARTIFACT_GROUP: ${{ inputs.amdgpu_families }} - OUTPUT_ARTIFACTS_DIR: ${{ env.OUTPUT_ARTIFACTS_DIR }} - VENV_DIR: ${{ env.VENV_DIR }} - FETCH_ARTIFACT_ARGS: --blas --tests - - # Fail early with import paths printed if the test artifact is missing - # Python dependencies or if PYTHONPATH/LD_LIBRARY_PATH no longer match the - # artifact layout. That makes packaging/layout regressions distinct from - # real race-detector findings. - - name: Verify TensileLite Python environment - run: | - export PYTHONPATH="${ROCM_PATH}/share/hipblaslt/tensilelite${PYTHONPATH:+:${PYTHONPATH}}" - export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROCM_PATH}/lib/llvm/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}" - "${VENV_DIR}/bin/python" - <<'PY' - import filelock - import joblib - import msgpack - import numpy - import yaml - import Tensile - import rocisa - import rocisa.instruction - print("Tensile:", Tensile.ROOT_PATH) - print("rocisa:", rocisa.__file__) - PY - - - name: Run rocjitsu race check - timeout-minutes: 75 - run: bash rocm-libraries/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh - - # Upload even on failure. The report directory contains the per-workload - # stdout logs, rocjitsu sink files, race.log files, generated reduced - # TensileLite input/output, and timing.tsv; without this artifact a failed - # detector run is usually not actionable from the console log alone. - - name: Upload race reports - if: ${{ always() }} - uses: actions/upload-artifact@de65e23aa2b7e23d713bb51fbfcb6d502f8667d8 # v4.6.2 - with: - name: rocjitsu-race-reports-${{ inputs.amdgpu_families }} - path: race-reports/ - if-no-files-found: ignore + uses: ./.github/workflows/therock-rocjitsu-race-check-linux.yml + with: + therock_ref: ${{ inputs.therock_ref }} + build_runs_on: ${{ inputs.build_runs_on }} + amdgpu_families: ${{ inputs.amdgpu_families }} + secrets: inherit diff --git a/.github/workflows/therock-rocjitsu-race-check-linux.yml b/.github/workflows/therock-rocjitsu-race-check-linux.yml new file mode 100644 index 000000000000..b547fd582d4e --- /dev/null +++ b/.github/workflows/therock-rocjitsu-race-check-linux.yml @@ -0,0 +1,124 @@ +name: TheRock rocjitsu race check Linux + +on: + workflow_call: + inputs: + therock_ref: + type: string + required: true + build_runs_on: + type: string + required: true + amdgpu_families: + type: string + required: true + +permissions: + contents: read + +jobs: + rocjitsu-race-check-linux: + name: rocjitsu race check (${{ inputs.amdgpu_families }}) + # This is a narrow post-build instrumentation job for hipBLASLt. + # Normal package tests cover broad functionality; this job fetches the BLAS + # artifacts from the just-finished TheRock build and runs two small workloads + # under rocjitsu's race detector. The workloads run entirely under CPU-side + # emulation; use the CPU build runner instead of consuming GPU test capacity. + runs-on: ${{ inputs.build_runs_on }} + timeout-minutes: 90 + container: + # Use a plain Ubuntu image without ROCm preinstalled. The ROCm runtime, + # compiler, hipBLASLt, and TensileLite all come from TheRock artifacts + # fetched below, so this job tests the PR-produced payload. + image: ghcr.io/rocm/no_rocm_image_ubuntu24_04@sha256:4150afe4759d14822f0e3f8930e1124f26e11f68b5c7b91ec9a02b20b1ebbb98 + options: --ipc host + --user 0:0 + --ulimit memlock=-1:-1 + --security-opt seccomp=unconfined + defaults: + run: + shell: bash + env: + VENV_DIR: ${{ github.workspace }}/.venv + ARTIFACT_RUN_ID: ${{ github.run_id }} + OUTPUT_ARTIFACTS_DIR: ./build + ROCM_PATH: ${{ github.workspace }}/build + AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }} + ROCJITSU_SOURCE_DIR: ${{ github.workspace }}/rocm-systems/emulation/rocjitsu + ROCJITSU_BUILD_DIR: ${{ github.workspace }}/rocjitsu-build + RACE_REPORT_DIR: ${{ github.workspace }}/race-reports + # Per-workload timeouts are below the job timeout so the script can report + # which workload failed and still upload race-reports on exit. + RACE_TIMEOUT_SECONDS: "180" + TENSILELITE_TIMEOUT_SECONDS: "420" + steps: + # The job needs three inputs in one workspace: + # 1. TheRock actions/scripts to fetch the exact artifacts from this run. + # 2. This rocm-libraries checkout for the race-check driver script. + # 3. rocm-systems source for rocjitsu until rocjitsu is packaged as a + # complete runnable artifact. + - name: Checkout TheRock repository + uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 + with: + repository: "ROCm/TheRock" + ref: ${{ inputs.therock_ref }} + + - name: Checkout rocm-libraries repository for scripts + uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 + with: + path: rocm-libraries + + - name: Checkout rocm-systems for rocjitsu + uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 + with: + repository: ROCm/rocm-systems + ref: develop + path: rocm-systems + sparse-checkout: | + emulation/rocjitsu + shared/machine-readable-isa/isa + sparse-checkout-cone-mode: true + + # The base image contains the shell and source-control tools used before + # artifact fetch; install only the host libraries needed for the local + # rocjitsu source build. ROCm runtime pieces come from TheRock artifacts. + - name: Install rocjitsu build dependencies + run: | + sudo apt-get update + sudo apt-get install -y --no-install-recommends \ + build-essential \ + ca-certificates \ + cmake \ + git \ + libdrm-dev \ + ninja-build + + # Pull the BLAS runtime and test payload from the current TheRock build. + # This is what keeps the check aligned with the exact hipBLASLt and + # TensileLite artifacts produced by the PR, rather than an image-local or + # previously installed ROCm tree. + - name: Run setup test environment workflow + timeout-minutes: 15 + uses: ./.github/actions/setup_test_environment + with: + ARTIFACT_RUN_ID: ${{ env.ARTIFACT_RUN_ID }} + ARTIFACT_GROUP: ${{ inputs.amdgpu_families }} + OUTPUT_ARTIFACTS_DIR: ${{ env.OUTPUT_ARTIFACTS_DIR }} + VENV_DIR: ${{ env.VENV_DIR }} + FETCH_ARTIFACT_ARGS: --blas --tests + + - name: Run rocjitsu race check + timeout-minutes: 75 + run: bash rocm-libraries/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh + + # Upload even on failure. The report directory contains the per-workload + # stdout logs, rocjitsu sink files, race.log files, generated reduced + # TensileLite input/output, and timing.tsv; without this artifact a failed + # detector run is usually not actionable from the console log alone. + - name: Upload race reports + if: ${{ always() }} + uses: actions/upload-artifact@de65e23aa2b7e23d713bb51fbfcb6d502f8667d8 # v4.6.2 + with: + name: rocjitsu-race-reports-${{ inputs.amdgpu_families }} + path: race-reports/ + if-no-files-found: ignore From 0ec20968e925633e4e2d00d5d1b509df0d877b25 Mon Sep 17 00:00:00 2001 From: James Newling Date: Mon, 20 Jul 2026 10:13:18 -0700 Subject: [PATCH 16/19] address Geo's review comments --- .github/workflows/therock-ci-linux.yml | 3 +-- .../therock-rocjitsu-race-check-linux.yml | 26 +++++++++---------- 2 files changed, 14 insertions(+), 15 deletions(-) diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index b870a2d47e84..c9b8e67c293d 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -177,10 +177,9 @@ jobs: needs: [therock-build-linux] # Run the narrow rocjitsu sidecar only for hipBLASLt rows and GPU families # that rocjitsu can emulate here. - if: ${{ inputs.build_runs_on != '' && (startsWith(inputs.amdgpu_families, 'gfx94') || startsWith(inputs.amdgpu_families, 'gfx950') || inputs.amdgpu_families == 'gfx1151') && contains(format(',{0},', inputs.projects_to_test), ',hipblaslt,') }} + if: ${{ inputs.build_runs_on != '' && (startsWith(inputs.amdgpu_families, 'gfx94') || startsWith(inputs.amdgpu_families, 'gfx950') || inputs.amdgpu_families == 'gfx1151') && contains(inputs.projects_to_test, 'hipblaslt') }} uses: ./.github/workflows/therock-rocjitsu-race-check-linux.yml with: therock_ref: ${{ inputs.therock_ref }} build_runs_on: ${{ inputs.build_runs_on }} amdgpu_families: ${{ inputs.amdgpu_families }} - secrets: inherit diff --git a/.github/workflows/therock-rocjitsu-race-check-linux.yml b/.github/workflows/therock-rocjitsu-race-check-linux.yml index b547fd582d4e..d81509f9b541 100644 --- a/.github/workflows/therock-rocjitsu-race-check-linux.yml +++ b/.github/workflows/therock-rocjitsu-race-check-linux.yml @@ -1,3 +1,9 @@ +# This reusable workflow is a narrow post-build instrumentation job for +# hipBLASLt. It fetches BLAS artifacts from the just-finished TheRock build and +# runs two small workloads under rocjitsu's race detector. The workloads run +# entirely under CPU-side emulation, so the workflow uses the CPU build runner +# instead of consuming GPU test capacity. + name: TheRock rocjitsu race check Linux on: @@ -19,11 +25,6 @@ permissions: jobs: rocjitsu-race-check-linux: name: rocjitsu race check (${{ inputs.amdgpu_families }}) - # This is a narrow post-build instrumentation job for hipBLASLt. - # Normal package tests cover broad functionality; this job fetches the BLAS - # artifacts from the just-finished TheRock build and runs two small workloads - # under rocjitsu's race detector. The workloads run entirely under CPU-side - # emulation; use the CPU build runner instead of consuming GPU test capacity. runs-on: ${{ inputs.build_runs_on }} timeout-minutes: 90 container: @@ -67,6 +68,8 @@ jobs: uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 with: path: rocm-libraries + sparse-checkout: .github/scripts + sparse-checkout-cone-mode: true - name: Checkout rocm-systems for rocjitsu uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 @@ -79,19 +82,16 @@ jobs: shared/machine-readable-isa/isa sparse-checkout-cone-mode: true - # The base image contains the shell and source-control tools used before - # artifact fetch; install only the host libraries needed for the local - # rocjitsu source build. ROCm runtime pieces come from TheRock artifacts. + # The pinned base image already provides CA certificates, CMake, Git, and + # Ninja. Install only the remaining host development files needed for the + # local rocjitsu source build. ROCm runtime pieces come from TheRock + # artifacts. - name: Install rocjitsu build dependencies run: | sudo apt-get update sudo apt-get install -y --no-install-recommends \ build-essential \ - ca-certificates \ - cmake \ - git \ - libdrm-dev \ - ninja-build + libdrm-dev # Pull the BLAS runtime and test payload from the current TheRock build. # This is what keeps the check aligned with the exact hipBLASLt and From b136bc0c8783e30a7b5ac48a9a2e9cd66365a80f Mon Sep 17 00:00:00 2001 From: James Newling Date: Mon, 20 Jul 2026 17:02:13 -0700 Subject: [PATCH 17/19] ci(hipblaslt): consume packaged Mirage and rocjitsu --- .../run_rocjitsu_hipblaslt_race_check.sh | 189 +++++++----------- .github/scripts/tests/therock_matrix_test.py | 8 + .github/scripts/therock_matrix.py | 8 +- .github/workflows/therock-ci-linux.yml | 7 +- .../therock-rocjitsu-race-check-linux.yml | 36 +--- 5 files changed, 96 insertions(+), 152 deletions(-) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index ab11b3f7db18..386fc698a903 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -1,35 +1,29 @@ #!/usr/bin/env bash set -euo pipefail -# The workflow that calls this script fetches hipBLASLt/TensileLite artifacts -# from the current TheRock build, then checks out rocm-systems. -# -# TODO(newling) Until rocjitsu is packaged as a complete runnable TheRock -# artifact, we build the rocjitsu CLI locally. Monitor progress on packaging rocjitsu. +# The workflow that calls this script fetches hipBLASLt/TensileLite, Mirage, +# and rocjitsu artifacts from the current TheRock build. # # The script runs small hipBLASLt and TensileLite GEMMs under the race detector. # TODO(newling) expand the GEMM-space tested. # # Basic flow: # 1. Use the TheRock artifact tree unpacked at ROCM_PATH. -# 2. Build rocjitsu from the rocm-systems checkout in ROCJITSU_SOURCE_DIR. -# 3. Select a rocjitsu config for gfx942, gfx950, or gfx1151. -# 4. Run hipblaslt-bench and a reduced TensileLite smoke under RJ_RACE=1. - -# These defaults match the GitHub Actions workspace layout: ROCM_PATH is the -# unpacked TheRock artifact tree, ROCJITSU_SOURCE_DIR is the checked-out -# rocm-systems source tree, ROCJITSU_BUILD_DIR is a local build directory, and -# RACE_REPORT_DIR is uploaded at the end of the job. They can be overridden for -# local reproduction. The brittle part is not the path names themselves, but the -# artifact layout underneath ROCM_PATH; the checks below fail early if -# hipBLASLt/TensileLite files move in TheRock artifacts. +# 2. Select a packaged rocjitsu config for gfx942, gfx950, or gfx1151. +# 3. Run hipblaslt-bench and a reduced TensileLite smoke through Mirage with +# rocjitsu race detection enabled. + +# These defaults match the GitHub Actions workspace layout. They can be +# overridden for local reproduction. The checks below fail early if an expected +# TheRock artifact moves or is missing. ROCM_PATH="${ROCM_PATH:-${PWD}/build}" AMDGPU_FAMILIES="${AMDGPU_FAMILIES:-}" ROCJITSU_GPU_TARGET="${ROCJITSU_GPU_TARGET:-}" -ROCJITSU_SOURCE_DIR="${ROCJITSU_SOURCE_DIR:-${PWD}/rocm-systems/emulation/rocjitsu}" -ROCJITSU_BUILD_DIR="${ROCJITSU_BUILD_DIR:-${PWD}/rocjitsu-build}" +ROCJITSU_CONFIG_DIR="${ROCJITSU_CONFIG_DIR:-${ROCM_PATH}/share/rocjitsu/configs}" ROCJITSU_CONFIG="${ROCJITSU_CONFIG:-}" RACE_REPORT_DIR="${RACE_REPORT_DIR:-${PWD}/race-reports}" +MIRAGE_BIN="${MIRAGE_BIN:-${ROCM_PATH}/bin/mirage}" +ROCJITSU_LIB="${ROCJITSU_LIB:-${ROCM_PATH}/lib/librocjitsu.so}" HIPBLASLT_BENCH="${HIPBLASLT_BENCH:-${ROCM_PATH}/bin/hipblaslt-bench}" TENSILELITE_ROOT="${TENSILELITE_ROOT:-${ROCM_PATH}/share/hipblaslt/tensilelite}" TENSILE_DRIVER="${TENSILE_DRIVER:-${TENSILELITE_ROOT}/Tensile/bin/Tensile}" @@ -74,25 +68,25 @@ select_rocjitsu_config() { # in flight: target-specific vs family-generic, and KMD vs non-KMD. # Prefer the target-specific KMD config when it exists. candidates=( - "${ROCJITSU_SOURCE_DIR}/configs/gfx942_cdna3_kmd.json" - "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna3_kmd.json" - "${ROCJITSU_SOURCE_DIR}/configs/gfx942_cdna3.json" - "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna3.json" + "${ROCJITSU_CONFIG_DIR}/gfx942_cdna3_kmd.json" + "${ROCJITSU_CONFIG_DIR}/amdgpu_cdna3_kmd.json" + "${ROCJITSU_CONFIG_DIR}/gfx942_cdna3.json" + "${ROCJITSU_CONFIG_DIR}/amdgpu_cdna3.json" ) ;; gfx950) # Same four-name compatibility pattern as gfx942: target-specific before # family-generic, and KMD before non-KMD. candidates=( - "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4_kmd.json" - "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna4_kmd.json" - "${ROCJITSU_SOURCE_DIR}/configs/gfx950_cdna4.json" - "${ROCJITSU_SOURCE_DIR}/configs/amdgpu_cdna4.json" + "${ROCJITSU_CONFIG_DIR}/gfx950_cdna4_kmd.json" + "${ROCJITSU_CONFIG_DIR}/amdgpu_cdna4_kmd.json" + "${ROCJITSU_CONFIG_DIR}/gfx950_cdna4.json" + "${ROCJITSU_CONFIG_DIR}/amdgpu_cdna4.json" ) ;; gfx1151) candidates=( - "${ROCJITSU_SOURCE_DIR}/configs/gfx1151.json" + "${ROCJITSU_CONFIG_DIR}/gfx1151.json" ) ;; *) @@ -163,6 +157,21 @@ if [[ ! -d "${ROCM_PATH}" ]]; then exit 1 fi +if [[ ! -x "${MIRAGE_BIN}" ]]; then + echo "Mirage executable not found or not executable: ${MIRAGE_BIN}" >&2 + exit 1 +fi + +if [[ ! -f "${ROCJITSU_LIB}" ]]; then + echo "rocjitsu runtime library not found: ${ROCJITSU_LIB}" >&2 + exit 1 +fi + +if [[ ! -d "${ROCJITSU_CONFIG_DIR}" ]]; then + echo "rocjitsu config directory not found: ${ROCJITSU_CONFIG_DIR}" >&2 + exit 1 +fi + if [[ ! -x "${HIPBLASLT_BENCH}" ]]; then echo "hipblaslt-bench not found or not executable: ${HIPBLASLT_BENCH}" >&2 exit 1 @@ -189,7 +198,7 @@ derive_rocjitsu_gpu_target select_rocjitsu_config if [[ -z "${ROCJITSU_CONFIG}" || ! -f "${ROCJITSU_CONFIG}" ]]; then - echo "rocjitsu ${ROCJITSU_GPU_TARGET} config not found under ${ROCJITSU_SOURCE_DIR}/configs" >&2 + echo "rocjitsu ${ROCJITSU_GPU_TARGET} config not found under ${ROCJITSU_CONFIG_DIR}" >&2 exit 1 fi @@ -204,15 +213,18 @@ mkdir -p "${RACE_REPORT_DIR}" # happens to be installed in the CI image. ROCM_PATH is exported for subprocesses # that use it to find the ROCm install root. export ROCM_PATH +export ROCM_HOME="${ROCM_PATH}" export PATH="${ROCM_PATH}/bin:${ROCM_PATH}/lib/llvm/bin:${PATH}" export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROCM_PATH}/lib/llvm/lib:${LD_LIBRARY_PATH:-}" export PYTHONPATH="${TENSILELITE_ROOT}${PYTHONPATH:+:${PYTHONPATH}}" echo "ROCM_PATH=${ROCM_PATH}" +echo "ROCM_HOME=${ROCM_HOME}" echo "AMDGPU_FAMILIES=${AMDGPU_FAMILIES}" echo "ROCJITSU_GPU_TARGET=${ROCJITSU_GPU_TARGET}" -echo "ROCJITSU_SOURCE_DIR=${ROCJITSU_SOURCE_DIR}" -echo "ROCJITSU_BUILD_DIR=${ROCJITSU_BUILD_DIR}" +echo "MIRAGE_BIN=${MIRAGE_BIN}" +echo "ROCJITSU_LIB=${ROCJITSU_LIB}" +echo "ROCJITSU_CONFIG_DIR=${ROCJITSU_CONFIG_DIR}" echo "ROCJITSU_CONFIG=${ROCJITSU_CONFIG}" echo "HIPBLASLT_BENCH=${HIPBLASLT_BENCH}" echo "TENSILELITE_ROOT=${TENSILELITE_ROOT}" @@ -225,66 +237,9 @@ echo "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}" echo "PYTHONPATH=${PYTHONPATH}" -# TODO(newling): Track migration to a packaged rocjitsu once TheRock provides a -# complete runnable artifact. Until then, build rocjitsu from the rocm-systems -# checkout selected by the workflow so this job controls the tool source. -# rocjitsu is still consumed from a source checkout in this workflow. Build only -# the CLI and optional runtime/shim targets needed to launch the test workloads; -# this keeps the job independent of full rocm-systems packaging. -# -# The warning suppressions keep the local rocjitsu build from failing on -# compiler/header warning mismatches. `nested-anon-types` is a Clang warning for -# anonymous structs/unions nested inside another type; use -Wno-error so those -# warnings remain visible but do not fail this bridge build. -cmake_args=( - -S "${ROCJITSU_SOURCE_DIR}" - -B "${ROCJITSU_BUILD_DIR}" - -G Ninja - -DCMAKE_BUILD_TYPE=Release - -DBUILD_TESTING=OFF - -DROCM_PATH="${ROCM_PATH}" - -DCMAKE_PREFIX_PATH="${ROCM_PATH}" - -DCMAKE_CXX_FLAGS="-Wno-error=unknown-warning-option -Wno-error=nested-anon-types" -) - -if ! command -v amdclang >/dev/null 2>&1 || ! command -v amdclang++ >/dev/null 2>&1; then - echo "amdclang and amdclang++ must be available from the fetched ROCm payload" >&2 - exit 1 -fi -cmake_args+=( - -DCMAKE_C_COMPILER="$(command -v amdclang)" - -DCMAKE_CXX_COMPILER="$(command -v amdclang++)" -) - -run_timed "configure rocjitsu" cmake "${cmake_args[@]}" - -# The CLI is required. The shared library and KMD shim target names have existed -# in some rocm-systems revisions and not others; build them when available, but -# do not make this bridge job depend on optional target names. -cmake_target_exists() { - local target="$1" - cmake --build "${ROCJITSU_BUILD_DIR}" --target help \ - | grep -Eq "(^|[[:space:]])${target}([:[:space:]]|$)" -} - -build_targets=(rocjitsu_bin) -if cmake_target_exists rocjitsu_shared; then - build_targets+=(rocjitsu_shared) -fi -if cmake_target_exists rocjitsu_kmd_shim; then - build_targets+=(rocjitsu_kmd_shim) -fi -run_timed "build rocjitsu" cmake --build "${ROCJITSU_BUILD_DIR}" --target "${build_targets[@]}" - -ROCJITSU_BIN="${ROCJITSU_BUILD_DIR}/tools/rocjitsu/rocjitsu" -if [[ ! -x "${ROCJITSU_BIN}" ]]; then - echo "rocjitsu binary not found after build: ${ROCJITSU_BIN}" >&2 - exit 1 -fi - -show_rocjitsu_version() { - echo "rocjitsu version:" - "${ROCJITSU_BIN}" --version +show_mirage_version() { + echo "Mirage version:" + "${MIRAGE_BIN}" --version } run_hipblaslt_bench_check() { @@ -294,22 +249,23 @@ run_hipblaslt_bench_check() { # triggers a separate device-fill kernel race report that needs independent # investigation. This check is scoped to the hipBLASLt GEMM dispatch path. timeout "${RACE_TIMEOUT_SECONDS}" \ - env \ - HSA_ENABLE_SDMA=1 \ - RJ_RACE=1 \ - RJ_LOG=1 \ - RJ_SINKS=stderr,file \ - RJ_SINK_DIR="${RACE_REPORT_DIR}/hipblaslt-bench" \ - "${ROCJITSU_BIN}" \ - --config "${ROCJITSU_CONFIG}" \ - -- "${HIPBLASLT_BENCH}" \ - --precision f32_r \ - --initialization zero \ - -m 128 \ - -n 128 \ - -k 128 \ - --iters 1 \ - --cold_iters 0 \ + "${MIRAGE_BIN}" run \ + --emulator rocjitsu \ + --config "${ROCJITSU_CONFIG}" \ + --in-process \ + --env HSA_ENABLE_SDMA=1 \ + --env RJ_RACE=1 \ + --env RJ_LOG=1 \ + --env RJ_SINKS=stderr,file \ + --env "RJ_SINK_DIR=${RACE_REPORT_DIR}/hipblaslt-bench" \ + -- "${HIPBLASLT_BENCH}" \ + --precision f32_r \ + --initialization zero \ + -m 128 \ + -n 128 \ + -k 128 \ + --iters 1 \ + --cold_iters 0 \ 2>&1 | tee "${RACE_REPORT_DIR}/hipblaslt-bench.log" local status=$? if [[ "${status}" -ne 0 ]]; then @@ -504,15 +460,16 @@ run_tensilelite_client_check() { # setup, rocisa imports, generated client config, and the standalone # tensilelite-client runtime path. timeout "${TENSILELITE_TIMEOUT_SECONDS}" \ - env \ - HSA_ENABLE_SDMA=1 \ - RJ_RACE=1 \ - RJ_LOG=1 \ - RJ_SINKS=stderr,file \ - RJ_SINK_DIR="${sink_dir}" \ - "${ROCJITSU_BIN}" \ - --config "${ROCJITSU_CONFIG}" \ - -- python3 "${TENSILE_DRIVER}" "${tensile_args[@]}" \ + "${MIRAGE_BIN}" run \ + --emulator rocjitsu \ + --config "${ROCJITSU_CONFIG}" \ + --in-process \ + --env HSA_ENABLE_SDMA=1 \ + --env RJ_RACE=1 \ + --env RJ_LOG=1 \ + --env RJ_SINKS=stderr,file \ + --env "RJ_SINK_DIR=${sink_dir}" \ + -- python3 "${TENSILE_DRIVER}" "${tensile_args[@]}" \ 2>&1 | tee "${RACE_REPORT_DIR}/tensilelite-client.log" local status=$? if [[ "${status}" -ne 0 ]]; then @@ -542,7 +499,7 @@ run_tensilelite_client_check() { fi } -run_timed "rocjitsu version" show_rocjitsu_version +run_timed "Mirage version" show_mirage_version check_status=0 diff --git a/.github/scripts/tests/therock_matrix_test.py b/.github/scripts/tests/therock_matrix_test.py index 8bea16e1ff81..92623745e7da 100644 --- a/.github/scripts/tests/therock_matrix_test.py +++ b/.github/scripts/tests/therock_matrix_test.py @@ -19,6 +19,14 @@ def test_collect_projects_to_run_without_additional_option(self): "hipsparselt", blas_entry["projects_to_test"].split(","), ) + self.assertIn( + "-DTHEROCK_ENABLE_ROCJITSU=ON", + blas_entry["cmake_options"].split(), + ) + self.assertIn( + "-DTHEROCK_ENABLE_MIRAGE=ON", + blas_entry["cmake_options"].split(), + ) def test_collect_projects_to_run(self): subtrees = ["projects/rocsparse", "projects/hipblaslt"] diff --git a/.github/scripts/therock_matrix.py b/.github/scripts/therock_matrix.py index c9d0f61f614d..58ec56944808 100644 --- a/.github/scripts/therock_matrix.py +++ b/.github/scripts/therock_matrix.py @@ -48,7 +48,13 @@ "projects_to_test": ["rocrand", "hiprand"], }, "blas": { - "cmake_options": ["-DTHEROCK_ENABLE_BLAS=ON"], + # The hipBLASLt race-check sidecar consumes these target-neutral + # artifacts from the same TheRock run instead of rebuilding rocjitsu. + "cmake_options": [ + "-DTHEROCK_ENABLE_BLAS=ON", + "-DTHEROCK_ENABLE_ROCJITSU=ON", + "-DTHEROCK_ENABLE_MIRAGE=ON", + ], "projects_to_test": ["hipblaslt", "rocblas", "hipblas", "tensilelite"], }, "miopen": { diff --git a/.github/workflows/therock-ci-linux.yml b/.github/workflows/therock-ci-linux.yml index c9b8e67c293d..4675b2f5176b 100644 --- a/.github/workflows/therock-ci-linux.yml +++ b/.github/workflows/therock-ci-linux.yml @@ -175,9 +175,10 @@ jobs: rocjitsu-race-check-linux: name: rocjitsu race check (${{ inputs.amdgpu_families }}) needs: [therock-build-linux] - # Run the narrow rocjitsu sidecar only for hipBLASLt rows and GPU families - # that rocjitsu can emulate here. - if: ${{ inputs.build_runs_on != '' && (startsWith(inputs.amdgpu_families, 'gfx94') || startsWith(inputs.amdgpu_families, 'gfx950') || inputs.amdgpu_families == 'gfx1151') && contains(inputs.projects_to_test, 'hipblaslt') }} + # Start the packaged Mirage/rocjitsu experiment on gfx950. The current + # TheRock rocm-systems pin predates the gfx1151 D16 false-positive fix; + # expand the matrix after the packaged path is proven and the pin advances. + if: ${{ inputs.build_runs_on != '' && inputs.amdgpu_families == 'gfx950-dcgpu' && contains(inputs.projects_to_test, 'hipblaslt') }} uses: ./.github/workflows/therock-rocjitsu-race-check-linux.yml with: therock_ref: ${{ inputs.therock_ref }} diff --git a/.github/workflows/therock-rocjitsu-race-check-linux.yml b/.github/workflows/therock-rocjitsu-race-check-linux.yml index d81509f9b541..f6b4e846dcba 100644 --- a/.github/workflows/therock-rocjitsu-race-check-linux.yml +++ b/.github/workflows/therock-rocjitsu-race-check-linux.yml @@ -45,19 +45,15 @@ jobs: OUTPUT_ARTIFACTS_DIR: ./build ROCM_PATH: ${{ github.workspace }}/build AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }} - ROCJITSU_SOURCE_DIR: ${{ github.workspace }}/rocm-systems/emulation/rocjitsu - ROCJITSU_BUILD_DIR: ${{ github.workspace }}/rocjitsu-build RACE_REPORT_DIR: ${{ github.workspace }}/race-reports # Per-workload timeouts are below the job timeout so the script can report # which workload failed and still upload race-reports on exit. RACE_TIMEOUT_SECONDS: "180" TENSILELITE_TIMEOUT_SECONDS: "420" steps: - # The job needs three inputs in one workspace: + # The job needs two inputs in one workspace: # 1. TheRock actions/scripts to fetch the exact artifacts from this run. # 2. This rocm-libraries checkout for the race-check driver script. - # 3. rocm-systems source for rocjitsu until rocjitsu is packaged as a - # complete runnable artifact. - name: Checkout TheRock repository uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2 with: @@ -71,32 +67,8 @@ jobs: sparse-checkout: .github/scripts sparse-checkout-cone-mode: true - - name: Checkout rocm-systems for rocjitsu - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0 - with: - repository: ROCm/rocm-systems - ref: develop - path: rocm-systems - sparse-checkout: | - emulation/rocjitsu - shared/machine-readable-isa/isa - sparse-checkout-cone-mode: true - - # The pinned base image already provides CA certificates, CMake, Git, and - # Ninja. Install only the remaining host development files needed for the - # local rocjitsu source build. ROCm runtime pieces come from TheRock - # artifacts. - - name: Install rocjitsu build dependencies - run: | - sudo apt-get update - sudo apt-get install -y --no-install-recommends \ - build-essential \ - libdrm-dev - - # Pull the BLAS runtime and test payload from the current TheRock build. - # This is what keeps the check aligned with the exact hipBLASLt and - # TensileLite artifacts produced by the PR, rather than an image-local or - # previously installed ROCm tree. + # Pull the BLAS runtime/test payload and the target-neutral Mirage + + # rocjitsu artifacts from the current TheRock build. - name: Run setup test environment workflow timeout-minutes: 15 uses: ./.github/actions/setup_test_environment @@ -105,7 +77,7 @@ jobs: ARTIFACT_GROUP: ${{ inputs.amdgpu_families }} OUTPUT_ARTIFACTS_DIR: ${{ env.OUTPUT_ARTIFACTS_DIR }} VENV_DIR: ${{ env.VENV_DIR }} - FETCH_ARTIFACT_ARGS: --blas --tests + FETCH_ARTIFACT_ARGS: --blas --tests --rocjitsu --mirage - name: Run rocjitsu race check timeout-minutes: 75 From 9720aa09c9525df52d14ae4c387c6b612fad1f9d Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 21 Jul 2026 09:40:03 -0700 Subject: [PATCH 18/19] ci(therock): use Rust-enabled build image --- .github/actions/ci-env/action.yml | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/.github/actions/ci-env/action.yml b/.github/actions/ci-env/action.yml index 1a9f001f872a..b401283ad035 100644 --- a/.github/actions/ci-env/action.yml +++ b/.github/actions/ci-env/action.yml @@ -7,7 +7,8 @@ outputs: value: "2d0da43cdb34bc59ea309834329546d96ea730f2" # 2026-07-10 commit docker-image: description: "Docker container image for Linux builds" - value: "ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:48492540591673fdc8d51beb89bde41e7ba13cbb528f643c0a481ba42c4058f2" + # Current TheRock build image, including Rust/Cargo for Mirage. + value: "ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:a382085df3ba2419b58aa9051350883a0d0b732a4bc0a4ef60458f8161bb08c6" linux-runner: description: "GitHub Actions runner label for Linux jobs" value: "aws-linux-scale-rocm-prod" From b3b0515154a01f0f68cfb7086c559366598e29de Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 21 Jul 2026 10:51:42 -0700 Subject: [PATCH 19/19] ci(rocjitsu): forward artifact environment through Mirage --- .../run_rocjitsu_hipblaslt_race_check.sh | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh index 386fc698a903..fc42cf7c3223 100755 --- a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -242,6 +242,21 @@ show_mirage_version() { "${MIRAGE_BIN}" --version } +# Mirage intentionally launches non-containerized workloads with a minimal +# inherited environment. Pass the fetched artifact and Python environment +# explicitly so the workload resolves ROCm libraries and the setup-python +# interpreter's shared library. +mirage_environment=( + --env "PATH=${PATH}" + --env "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}" + --env "PYTHONPATH=${PYTHONPATH}" + --env "ROCM_PATH=${ROCM_PATH}" + --env "ROCM_HOME=${ROCM_HOME}" +) +if [[ -n "${VIRTUAL_ENV:-}" ]]; then + mirage_environment+=(--env "VIRTUAL_ENV=${VIRTUAL_ENV}") +fi + run_hipblaslt_bench_check() { mkdir -p "${RACE_REPORT_DIR}/hipblaslt-bench" echo "running hipblaslt-bench under rocjitsu race detection" @@ -258,6 +273,7 @@ run_hipblaslt_bench_check() { --env RJ_LOG=1 \ --env RJ_SINKS=stderr,file \ --env "RJ_SINK_DIR=${RACE_REPORT_DIR}/hipblaslt-bench" \ + "${mirage_environment[@]}" \ -- "${HIPBLASLT_BENCH}" \ --precision f32_r \ --initialization zero \ @@ -469,6 +485,7 @@ run_tensilelite_client_check() { --env RJ_LOG=1 \ --env RJ_SINKS=stderr,file \ --env "RJ_SINK_DIR=${sink_dir}" \ + "${mirage_environment[@]}" \ -- python3 "${TENSILE_DRIVER}" "${tensile_args[@]}" \ 2>&1 | tee "${RACE_REPORT_DIR}/tensilelite-client.log" local status=$?