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" diff --git a/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh new file mode 100755 index 000000000000..fc42cf7c3223 --- /dev/null +++ b/.github/scripts/run_rocjitsu_hipblaslt_race_check.sh @@ -0,0 +1,548 @@ +#!/usr/bin/env bash +set -euo pipefail + +# 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. 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_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}" +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) + # 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_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_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_CONFIG_DIR}/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() { + if [[ -f "${TIMING_FILE}" ]]; then + echo "" + 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}" + 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}" +} + + +# 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 +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 +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 + + +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_CONFIG_DIR}" >&2 + exit 1 +fi + +if [[ -z "${TENSILELITE_GPU_TARGET}" ]]; then + TENSILELITE_GPU_TARGET="${ROCJITSU_GPU_TARGET}" +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. 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 "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}" +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}" +echo "PYTHONPATH=${PYTHONPATH}" + + +show_mirage_version() { + echo "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" + # 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}" \ + "${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" \ + "${mirage_environment[@]}" \ + -- "${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 + echo "hipblaslt-bench race check command failed with status ${status}" >&2 + 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 + cat "${RACE_REPORT_DIR}/hipblaslt-bench/race.log" >&2 + return 1 + fi +} + + +write_tensilelite_check_yaml() { + local yaml="$1" + + 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. + # + # 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). 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 + 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 + +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. + # 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}" <&1 | tee "${RACE_REPORT_DIR}/tensilelite-client.log" + local status=$? + if [[ "${status}" -ne 0 ]]; then + echo "tensilelite-client race check command failed with status ${status}" >&2 + return "${status}" + fi + + 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, 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 + return 1 + fi +} + +run_timed "Mirage version" show_mirage_version + +check_status=0 + +# 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 check" run_hipblaslt_bench_check +hipblaslt_status=$? +run_timed "tensilelite-client race check" run_tensilelite_client_check +tensilelite_status=$? +set -e + +if [[ "${hipblaslt_status}" -ne 0 ]]; then + 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 check failed with status ${tensilelite_status}" >&2 + check_status=1 +fi + +if [[ "${check_status}" -ne 0 ]]; then + echo "one or more rocjitsu race checks failed" >&2 + exit "${check_status}" +fi + +echo "rocjitsu race checks completed without race reports" 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 0948b1e4ee59..4675b2f5176b 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 }} @@ -171,3 +171,16 @@ jobs: platform: "linux" test_type: ${{ inputs.test_type }} secrets: inherit + + rocjitsu-race-check-linux: + name: rocjitsu race check (${{ inputs.amdgpu_families }}) + needs: [therock-build-linux] + # 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 }} + build_runs_on: ${{ inputs.build_runs_on }} + amdgpu_families: ${{ inputs.amdgpu_families }} diff --git a/.github/workflows/therock-ci.yml b/.github/workflows/therock-ci.yml index c2379f99d73d..614db7983f99 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: 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..f6b4e846dcba --- /dev/null +++ b/.github/workflows/therock-rocjitsu-race-check-linux.yml @@ -0,0 +1,96 @@ +# 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: + 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 }}) + 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 }} + 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 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. + - 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 + sparse-checkout: .github/scripts + sparse-checkout-cone-mode: true + + # 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 + 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 --rocjitsu --mirage + + - 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