From 9830ba76b47b9035039f8e967c29d4a8c58a4491 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Mon, 17 Aug 2026 09:56:26 -0500 Subject: [PATCH 01/19] [CI] Add HeCBench to CI --- .github/hecbench/bench.sh | 444 ++++++++++++++++++ .github/hecbench/ci_benchmarks.deps.txt | 93 ++++ .github/hecbench/ci_benchmarks.txt | 170 +++++++ .github/hecbench/lib/common.sh | 576 ++++++++++++++++++++++++ .github/hecbench/lib/presets.sh | 11 + .github/hecbench/modify_makefiles.sh | 438 ++++++++++++++++++ .github/hecbench/update_deps.sh | 100 ++++ .github/workflows/spirv-ci-linux.yml | 115 +++++ 8 files changed, 1947 insertions(+) create mode 100755 .github/hecbench/bench.sh create mode 100644 .github/hecbench/ci_benchmarks.deps.txt create mode 100644 .github/hecbench/ci_benchmarks.txt create mode 100755 .github/hecbench/lib/common.sh create mode 100644 .github/hecbench/lib/presets.sh create mode 100755 .github/hecbench/modify_makefiles.sh create mode 100755 .github/hecbench/update_deps.sh diff --git a/.github/hecbench/bench.sh b/.github/hecbench/bench.sh new file mode 100755 index 000000000..0a68e2178 --- /dev/null +++ b/.github/hecbench/bench.sh @@ -0,0 +1,444 @@ +#!/usr/bin/env bash +# Build and run HeCBench HIP benchmarks in one step (see show_usage below). +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +SCRIPT_NAME="$(basename "$0")" + +# shellcheck source=lib/common.sh +source "${SCRIPT_DIR}/lib/common.sh" +# shellcheck source=lib/presets.sh +source "${SCRIPT_DIR}/lib/presets.sh" + +# ============================================================================== +# DEFAULTS +# ============================================================================== + +TIMEOUT_SECONDS=${TIMEOUT_SECONDS:-300} +GPU_ID=${GPU_ID:-0,1} # Use both GPUs by default for multi-GPU benchmarks +BENCHMARK_FILTER=${BENCHMARK_FILTER:-} +DRY_RUN=false + +show_usage() { + cat </ + timings.csv benchmark,build_seconds,run_seconds + success.log benchmarks that built and ran cleanly + failed.log non-zero exit (build or run) + suspect.log exit 0 but output contains validation failures + (FAIL, MISMATCH, nan, missing input data, etc.) + timeout.log exceeded --timeout + skipped.log architecture-unsupported (e.g. saxpy-ompt-hip on amdgcnspirv) + .log per-benchmark combined build+run output +EOF +} + +# ============================================================================== +# ARG PARSING +# ============================================================================== + +ARCH="" +PRESET="" +while [[ $# -gt 0 ]]; do + case "$1" in + --help|-h) show_usage; exit 0 ;; + --preset) PRESET="$2"; shift 2 ;; + --dry-run) DRY_RUN=true; shift ;; + --filter) BENCHMARK_FILTER="$2"; shift 2 ;; + --timeout) TIMEOUT_SECONDS="$2"; shift 2 ;; + --gpu-id) GPU_ID="$2"; shift 2 ;; + -*) log_error "Unknown option: $1"; show_usage; exit 1 ;; + *) + if [[ -z "$ARCH" ]]; then + ARCH="$1" + else + log_error "Multiple HIP_ARCH values; this script accepts one" + exit 1 + fi + shift + ;; + esac +done + +# Resolve --preset into BENCHMARK_FILTER (--filter takes precedence). +if [[ -n "$PRESET" && -z "$BENCHMARK_FILTER" ]]; then + case "$PRESET" in + quick) BENCHMARK_FILTER="$PRESET_QUICK" ;; + standard) BENCHMARK_FILTER="$PRESET_STANDARD" ;; + extended) BENCHMARK_FILTER="$PRESET_EXTENDED" ;; + *) + log_error "Unknown preset: $PRESET (choose: quick, standard, extended)" + exit 1 + ;; + esac +fi + +if [[ -z "$ARCH" ]]; then + log_error "HIP_ARCH is required" + show_usage + exit 1 +fi + +# $ARCH is the user-facing token (log dir, prints); HIPCC_ARCH is what make gets +# as HIP_ARCH=. "amdgcnspirv-be" is our alias for the default SPIRV backend, so +# it maps to plain amdgcnspirv; every other arch passes through unchanged. +if [[ "$ARCH" == "amdgcnspirv-be" ]]; then + HIPCC_ARCH="amdgcnspirv" +else + HIPCC_ARCH="$ARCH" +fi + +# ============================================================================== +# VALIDATION & SETUP +# ============================================================================== + +require_rocm_path || exit 1 +HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR") || { + log_error "Cannot find HeCBench src directory" + exit 1 +} + +LOG_DIR="${SCRIPT_DIR}/bench_logs_$(date +%Y%m%d_%H%M%S)_${ARCH}" +mkdir -p "$LOG_DIR" + +export PATH="${ROCM_PATH}/bin:${PATH}" +export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${LD_LIBRARY_PATH:-}" +export HIP_PATH="${ROCM_PATH}" +# HIP_CLANG_PATH points hipcc at the clang to drive; CI sets it to the freshly +# built LLVM bin dir. Unset -> fall back to $ROCM_PATH/bin. +if [[ -n "${HIP_CLANG_PATH:-}" ]]; then + export HIP_CLANG_PATH + export PATH="${HIP_CLANG_PATH}:${PATH}" + export LD_LIBRARY_PATH="$(dirname "${HIP_CLANG_PATH}")/lib:${LD_LIBRARY_PATH}" +fi +# hipcc/amdclang++ honor this for AMDGCN bitcode lookup +export HIP_DEVICE_LIB_PATH="${ROCM_PATH}/lib/llvm/amdgcn/bitcode" +export ROCR_VISIBLE_DEVICES="$GPU_ID" +# prna-hip reads DATAPATH; overridable, harmless default for the rest. +export DATAPATH="${DATAPATH:-${HECBENCH_SRC}/prna-cuda/data_tables}" +ulimit -s unlimited 2>/dev/null || true + +echo "benchmark,build_seconds,run_seconds" > "$LOG_DIR/timings.csv" +touch "$LOG_DIR/success.log" \ + "$LOG_DIR/failed.log" \ + "$LOG_DIR/suspect.log" \ + "$LOG_DIR/timeout.log" \ + "$LOG_DIR/skipped.log" + +log_info "===========================================" +log_info "HeCBench bench (build+run merged)" +log_info "===========================================" +log_info "ROCm path: $ROCM_PATH" +log_info "HeCBench src: $HECBENCH_SRC" +log_info "Arch: $ARCH" +log_info "Timeout: ${TIMEOUT_SECONDS}s" +log_info "GPU ID: $GPU_ID" +log_info "Logs: $LOG_DIR" + +# Clear the COMGR JIT cache to prevent stale finalized kernels from being +# served after a compiler or runtime update. The cache keys are based on the +# SPIR-V input hash, so a runtime-only update (same compiler, same SPIR-V) +# would silently reuse the old (possibly buggy) native code. +if [[ -d "${HOME}/.cache/comgr" ]]; then + log_info "Clearing COMGR cache (${HOME}/.cache/comgr) ..." + rm -rf "${HOME}/.cache/comgr" +fi + +echo "" + +# ============================================================================== +# PER-BENCHMARK +# ============================================================================== + +# Architecture-incompatibility skip list. +should_skip() { + local name="$1" + # saxpy-ompt-hip needs a concrete GPU ISA; spirv64-amd-amdhsa not in toolchain. + [[ "$name" == "saxpy-ompt-hip" && "$HIPCC_ARCH" == "amdgcnspirv" ]] && return 0 + # TEMP: pingpong-hip hangs in MPI/NCCL (pre-existing issue, unrelated to + # current TheRock bring-up); burns the full timeout for nothing. + [[ "$name" == "pingpong-hip" ]] && return 0 + [[ "$name" == "assert-hip" ]] && return 0 + + return 1 +} + +# Benchmarks that need special environment overrides for `make run`. +apply_special_env() { + local name="$1" + case "$name" in + assert-hip) + # This benchmark intentionally triggers a device-side assertion to + # verify host-side error reporting. Suppress the ROCm runtime's GPU + # coredump on exception so the apport pipe in core_pattern is not + # invoked. Keep coredumps enabled for every other benchmark. + export HSA_DISABLE_COREDUMP_ON_EXCEPTION=1 + ;; + esac +} + +clear_special_env() { + local name="$1" + case "$name" in + assert-hip) + unset HSA_DISABLE_COREDUMP_ON_EXCEPTION + ;; + esac +} + +# Extra `make` command-line arguments specific to a benchmark. Used for +# benchmarks whose Makefile reads a variable other than HIP_ARCH for the +# offload arch (e.g. saxpy-ompt-hip uses ARCH for OpenMP `-march=$(ARCH)`). +# Returns the extra args via stdout, one per line. +extra_make_args() { + local name="$1" + case "$name" in + saxpy-ompt-hip) + echo "ARCH=$ARCH" + ;; + esac + # libhipcxx hard-errors on amdgcnspirv because no compile-time + # __gfx*__ macro is defined for SPIR-V; allow it through and silence the + # accompanying warning. Host-side std::chrono is unaffected. + # Applies to both SPIRV paths. Plain amdgcnspirv also opts out of the + # (now default) backend to exercise the legacy translator. + if [[ "$HIPCC_ARCH" == "amdgcnspirv" ]]; then + echo "EXTRA_CFLAGS=-D_LIBCUDACXX_ALLOW_UNSUPPORTED_ARCHITECTURE -DNDEBUG" + [[ "$ARCH" == "amdgcnspirv" ]] && echo "EXTRA_HIPCCFLAGS=-no-use-spirv-backend" + fi +} + +# Per-benchmark direct-run override. When non-empty, bench_one builds with +# `make` (default target) and then invokes each emitted command directly, +# bypassing the Makefile's `run:` recipe. Each line is "binary arg1 arg2 ..." +# parsed with `read -ra`; the binary is resolved relative to $dir. +# +# Use this when the Makefile's `run:` recipe includes a config that exceeds +# this system's resources (e.g. attention-paged-hip's 131072-block case OOMs +# on MI210). Keeps coverage of the configs that fit without an upstream patch. +direct_run_argsets() { + local name="$1" + local arch="${2:-$ARCH}" + case "$name" in + attention-paged-hip) + # 4th make-run config (131072 kv blocks) OOMs; substitute 65536. + echo "./main 8 32 128 4096 128 100" + echo "./main 8 32 128 4096 1024 100" + echo "./main 8 32 128 4096 8192 100" + echo "./main 8 32 128 4096 65536 100" + ;; + esac +} + +# check_output_validation() is defined in lib/common.sh. + +bench_one() { + local dir="$1" + local name; name=$(basename "$dir") + local log="$LOG_DIR/${name}.log" + + if should_skip "$name"; then + log_info "Skipped: $name (arch unsupported)" + echo "$name" >> "$LOG_DIR/skipped.log" + return + fi + + if [[ "$DRY_RUN" == "true" ]]; then + log_info "[DRY-RUN] Would build+run: $name" + return + fi + + # Reclaim leaked OpenMPI/RCCL backing files in /dev/shm so an MPI benchmark + # whose predecessor was SIGKILL'd doesn't fail with "not enough space". + cleanup_stale_shm + + cd "$dir" || { log_error "$name: cannot cd to $dir"; return; } + apply_special_env "$name" + + # Clean to ensure a deterministic build state. + make clean &>/dev/null || true + + local build_start build_elapsed run_start run_elapsed rc + local -a make_extra=() + while IFS= read -r arg; do + [[ -n "$arg" ]] && make_extra+=("$arg") + done < <(extra_make_args "$name") + + # ---- Phase 1: BUILD (default target only) ---- + build_start=$(date +%s.%N) + set +e + timeout "$TIMEOUT_SECONDS" make HIP_ARCH="$HIPCC_ARCH" "${make_extra[@]}" &>"$log" + rc=$? + set -e + build_elapsed=$(awk "BEGIN {printf \"%.3f\", $(date +%s.%N) - $build_start}") + + if [[ $rc -ne 0 ]]; then + if [[ $rc -eq 124 ]]; then + log_error "Timeout (build): $name (>${TIMEOUT_SECONDS}s)" + echo "$name" >> "$LOG_DIR/timeout.log" + else + log_error "Failed (build): $name (exit $rc)" + echo "$name" >> "$LOG_DIR/failed.log" + fi + clear_special_env "$name" + cd "$SCRIPT_DIR" + return + fi + + # ---- Phase 2: RUN (no make overhead) ---- + # Collect the run commands: either from direct_run_argsets overrides + # or by asking the Makefile what `make run` would execute. + local -a runcmds=() + while IFS= read -r line; do + [[ -n "$line" ]] && runcmds+=("$line") + done < <(direct_run_argsets "$name" "$HIPCC_ARCH") + + if [[ ${#runcmds[@]} -eq 0 ]]; then + # Extract commands from the Makefile's run recipe via dry-run. + # Binary is already built, so make -n run only prints run commands. + # Join backslash-continuation lines before splitting into commands. + local accum="" + while IFS= read -r line; do + if [[ "$line" == *'\' ]]; then + accum+="${line%\\} " + else + accum+="$line" + [[ -n "$accum" ]] && runcmds+=("$accum") + accum="" + fi + done < <(make -n run HIP_ARCH="$HIPCC_ARCH" "${make_extra[@]}" 2>/dev/null) + [[ -n "$accum" ]] && runcmds+=("$accum") + fi + + if [[ ${#runcmds[@]} -eq 0 ]]; then + log_error "Failed: $name (no run commands found)" + echo "$name" >> "$LOG_DIR/failed.log" + clear_special_env "$name" + cd "$SCRIPT_DIR" + return + fi + + run_start=$(date +%s.%N) + set +e + rc=0 + for cmd in "${runcmds[@]}"; do + echo "+ $cmd" >>"$log" + timeout "$TIMEOUT_SECONDS" bash -c "$cmd" &>>"$log" 2>&1 + rc=$? + [[ $rc -ne 0 ]] && break + done + set -e + run_elapsed=$(awk "BEGIN {printf \"%.3f\", $(date +%s.%N) - $run_start}") + + if [[ $rc -eq 0 ]]; then + local time_detail="build ${build_elapsed}s, run ${run_elapsed}s" + local validation_issues + if validation_issues=$(check_output_validation "$log" "$name"); then + log_success "$name (${time_detail})" + echo "$name" >> "$LOG_DIR/success.log" + else + log_warn "$name (${time_detail}) — exit 0 but output validation suspect" + echo "$name" >> "$LOG_DIR/suspect.log" + fi + echo "$name,$build_elapsed,$run_elapsed" >> "$LOG_DIR/timings.csv" + elif [[ $rc -eq 124 ]]; then + log_error "Timeout (run): $name (>${TIMEOUT_SECONDS}s)" + echo "$name" >> "$LOG_DIR/timeout.log" + else + log_error "Failed (run): $name (exit $rc)" + echo "$name" >> "$LOG_DIR/failed.log" + fi + + clear_special_env "$name" + cd "$SCRIPT_DIR" +} + +# ============================================================================== +# DISCOVER & ITERATE +# ============================================================================== + +if [[ -n "$BENCHMARK_FILTER" ]]; then + # Build dirs directly from the filter/preset list — no need to scan all dirs + declare -a dirs=() + IFS=',' read -ra wanted <<< "$BENCHMARK_FILTER" + for w in "${wanted[@]}"; do + w="${w## }"; w="${w%% }" + [[ -z "$w" ]] && continue + local_dir="$HECBENCH_SRC/$w" + [[ -d "$local_dir" ]] && dirs+=("$local_dir") + done + if [[ ${#dirs[@]} -eq 0 ]]; then + log_error "Filter matched no benchmarks: $BENCHMARK_FILTER" + exit 1 + fi +else + mapfile -t dirs < <(find "$HECBENCH_SRC" -maxdepth 1 -type d -name "*-hip" | sort) + if [[ ${#dirs[@]} -eq 0 ]]; then + log_error "No *-hip directories under $HECBENCH_SRC" + exit 1 + fi +fi + +log_info "Processing ${#dirs[@]} benchmarks..." +echo "" + +for d in "${dirs[@]}"; do + bench_one "$d" +done + +# ============================================================================== +# SUMMARY +# ============================================================================== + +count() { + local f="$LOG_DIR/$1" + [[ -f "$f" ]] && wc -l <"$f" | tr -d ' ' || echo 0 +} + +echo "" +log_info "===========================================" +log_info "Summary" +log_info "===========================================" +log_info " Success: $(count success.log)" +log_info " Suspect: $(count suspect.log) (exit 0 but output has failures)" +log_info " Failed: $(count failed.log)" +log_info " Timeout: $(count timeout.log)" +log_info " Skipped: $(count skipped.log)" +log_info "" +log_info "Logs: $LOG_DIR" +log_info "Timings: $LOG_DIR/timings.csv" diff --git a/.github/hecbench/ci_benchmarks.deps.txt b/.github/hecbench/ci_benchmarks.deps.txt new file mode 100644 index 000000000..4617c2274 --- /dev/null +++ b/.github/hecbench/ci_benchmarks.deps.txt @@ -0,0 +1,93 @@ +# GENERATED by update_deps.sh — do not edit by hand. +# In-tree HeCBench dirs referenced (transitively) via '../' by the benchmarks in +# ci_benchmarks.txt: sibling *-cuda data dirs, shared include/, etc. The +# test_hecbench CI job checks these out alongside the benchmarks so their builds +# resolve. Regenerate with .github/hecbench/update_deps.sh (see that script). +accuracy-cuda +ace-cuda +adam-cuda +adamw-cuda +addBiasQKV-cuda +addBiasResidualLayerNorm-cuda +affine-cuda +aidw-cuda +atomicReduction-cuda +attentionMultiHead-cuda +backprop-cuda +base64e-cuda +bfs-sycl +blockAccess-cuda +boxfilter-sycl +bsw-cuda +ccs-cuda +ccsd-trpdrv-cuda +cfd-cuda +channelShuffle-cuda +channelSum-cuda +che-cuda +cmp-cuda +cobahh-cuda +complex-cuda +damage-cuda +data +debayer-sycl +deredundancy-sycl +dpid-cuda +dxtc2-sycl +easyWave-omp +ecdh-cuda +egs-cuda +extend2-sycl +f16atomic-cuda +fft-cuda +fpdc-cuda +fwt-cuda +geodesic-cuda +geodesic-sycl +gibbs-cuda +groupnorm-cuda +haversine-cuda +hogbom-cuda +hotspot-cuda +include +inversek2j-sycl +is-cuda +laplace3d-cuda +log2-cuda +logan-cuda +lud-cuda +marchingCubes-cuda +medianfilter-cuda +medianfilter-sycl +minimap2-sycl +mis-cuda +mpc-cuda +mrc-cuda +mr-cuda +nlll-cuda +nms-cuda +opticalFlow-cuda +pad-cuda +particlefilter-cuda +particles-cuda +perplexity-cuda +phmm-cuda +projectile-cuda +pso-cuda +qkv-hip +resnet-kernels-cuda +rmsnorm-cuda +romberg-cuda +rsc-cuda +sc-cuda +slit-cuda +sobel-sycl +spaxpby-cuda +sptrsv-sycl +svd3x3-cuda +tqs-cuda +tsa-cuda +urng-cuda +urng-sycl +xlqc-cuda +zeropoint-cuda diff --git a/.github/hecbench/ci_benchmarks.txt b/.github/hecbench/ci_benchmarks.txt new file mode 100644 index 000000000..ff37fcb03 --- /dev/null +++ b/.github/hecbench/ci_benchmarks.txt @@ -0,0 +1,170 @@ +# HeCBench subset run by the test_hecbench CI job (amdgcnspirv-be backend). +# One benchmark directory name (under HeCBench/src) per line; '#' comments and +# blank lines are ignored by the workflow. +# +# Derivation: 2026-07-15 -be passing set (525), minus benchmarks needing +# external data (download_datasets.sh / DVC / extract_archives.sh), plus the 23 +# regression-sensitive ones, filled cheapest-first under a ~900s budget = 161 +# (miniWeather-hip dropped: needs OpenMPI, not worth the dependency for one bench). +# Timings are a gfx90a proxy; recalibrate on gfx942 before making this blocking. +accuracy-hip +ace-hip +adam-hip +adamw-hip +addBiasQKV-hip +addBiasResidualLayerNorm-hip +aes-hip +affine-hip +aidw-hip +aligned-types-hip +amgmk-hip +ans-hip +aobench-hip +asta-hip +atan2-hip +atomicReduction-hip +attentionMultiHead-hip +b+tree-hip +babelstream-hip +backprop-hip +base64e-hip +bfs-hip +binomial-hip +blockexchange-hip +bscan-hip +bsw-hip +ccs-hip +ccsd-trpdrv-hip +ced-hip +cfd-hip +channelShuffle-hip +channelSum-hip +che-hip +chemv-hip +clenergy-hip +clock-hip +cmp-hip +cobahh-hip +collision-hip +colorwheel-hip +complex-hip +concurrentKernels-hip +conversion-hip +cooling-hip +cross-hip +d2q9-bgk-hip +damage-hip +dct8x8-hip +debayer-hip +deredundancy-hip +determinant-hip +dispatch-hip +distort-hip +divergence-hip +dpid-hip +dropout-hip +dxtc2-hip +easyWave-hip +ecdh-hip +egs-hip +extend2-hip +f16atomic-hip +fft-hip +fhd-hip +floydwarshall-hip +fpc-hip +fpdc-hip +fwt-hip +gc-hip +gels-hip +gibbs-hip +graphExecution-hip +groupnorm-hip +haversine-hip +heartwall-hip +heat-hip +histogram-hip +hogbom-hip +hotspot-hip +hungarian-hip +interleave-hip +intrinsics-cast-hip +inversek2j-hip +is-hip +jenkins-hash-hip +langford-hip +laplace3d-hip +lavaMD-hip +layout-hip +lda-hip +libor-hip +log2-hip +logan-hip +lombscargle-hip +lud-hip +mallocFree-hip +mandelbrot-hip +marchingCubes-hip +md5hash-hip +medianfilter-hip +minimap2-hip +mis-hip +mixbench-hip +mpc-hip +mr-hip +mrc-hip +mrg32k3a-hip +murmurhash3-hip +myocyte-hip +ne-hip +nlll-hip +nms-hip +nn-hip +opticalFlow-hip +overlap-hip +pad-hip +particlefilter-hip +particles-hip +pathfinder-hip +permute-hip +perplexity-hip +phmm-hip +projectile-hip +pso-hip +qtclustering-hip +quant3MatMul-hip +quantAQLM-hip +radixsort-hip +resnet-kernels-hip +reverse-hip +reverse2D-hip +rmsnorm-hip +rodrigues-hip +romberg-hip +rsc-hip +sc-hip +scan2-hip +si-hip +slit-hip +sobel-hip +spaxpby-hip +split-hip +spm-hip +sptrsv-hip +srad-hip +stencil1d-hip +su3-hip +svd3x3-hip +threadfence-hip +tqs-hip +triad-hip +tsa-hip +unfold-hip +urng-hip +vmc-hip +warpexchange-hip +warpsort-hip +winograd-hip +xlqc-hip +zeropoint-hip +zmddft-hip diff --git a/.github/hecbench/lib/common.sh b/.github/hecbench/lib/common.sh new file mode 100755 index 000000000..32ddd0ad9 --- /dev/null +++ b/.github/hecbench/lib/common.sh @@ -0,0 +1,576 @@ +#!/usr/bin/env bash +# Common library for production-ready HeCBench scripts +# Provides shared logging, validation, and utility functions + +# ============================================================================== +# LOGGING FUNCTIONS +# ============================================================================== + +# Colors for output (disabled if NO_COLOR is set or not a terminal) +if [[ -t 1 ]] && [[ -z "${NO_COLOR:-}" ]]; then + readonly RED=$'\033[0;31m' + readonly GREEN=$'\033[0;32m' + readonly YELLOW=$'\033[1;33m' + readonly BLUE=$'\033[0;34m' + readonly NC=$'\033[0m' # No Color +else + readonly RED='' + readonly GREEN='' + readonly YELLOW='' + readonly BLUE='' + readonly NC='' +fi + +log_info() { + echo -e "[$(date '+%Y-%m-%d %H:%M:%S')] ${BLUE}[INFO]${NC} $*" +} + +log_success() { + echo -e "[$(date '+%Y-%m-%d %H:%M:%S')] ${GREEN}[SUCCESS]${NC} $*" +} + +log_warn() { + echo -e "[$(date '+%Y-%m-%d %H:%M:%S')] ${YELLOW}[WARN]${NC} $*" >&2 +} + +log_error() { + echo -e "[$(date '+%Y-%m-%d %H:%M:%S')] ${RED}[ERROR]${NC} $*" >&2 +} + +log_debug() { + if [[ "${VERBOSE:-false}" == "true" ]]; then + echo -e "[$(date '+%Y-%m-%d %H:%M:%S')] [DEBUG] $*" >&2 + fi +} + +# ============================================================================== +# ENVIRONMENT VARIABLE VALIDATION +# ============================================================================== + +# Require ROCM_PATH to be explicitly set +require_rocm_path() { + if [[ -z "${ROCM_PATH:-}" ]]; then + log_error "ROCM_PATH environment variable is not set" + echo "" + echo "Please set ROCM_PATH before running this script:" + echo " export ROCM_PATH=/path/to/TheRock-dist" + echo " # or" + echo " export ROCM_PATH=/opt/rocm" + echo "" + echo "Then run: $SCRIPT_NAME" + return 1 + fi + + # Strip trailing slashes to prevent double slashes in paths + ROCM_PATH="${ROCM_PATH%/}" + + # Validate the path + if ! validate_rocm_path "$ROCM_PATH"; then + log_error "ROCM_PATH is set but invalid: $ROCM_PATH" + log_error "ROCm installation must contain:" + log_error " - bin/hipcc (HIP compiler)" + log_error " - lib/ (runtime libraries)" + return 1 + fi + + log_success "Using ROCM_PATH: $ROCM_PATH" + return 0 +} + +# ============================================================================== +# VALIDATION FUNCTIONS +# ============================================================================== + +# Check if a command exists +check_command() { + local cmd="$1" + if ! command -v "$cmd" &>/dev/null; then + log_error "Required command not found: $cmd" + return 1 + fi + log_debug "Found command: $cmd" + return 0 +} + +# Validate directory exists and is readable +validate_dir() { + local dir="$1" + local description="${2:-Directory}" + + if [[ ! -d "$dir" ]]; then + log_error "$description not found: $dir" + return 1 + fi + + if [[ ! -r "$dir" ]]; then + log_error "$description not readable: $dir" + return 1 + fi + + log_debug "Validated directory: $dir" + return 0 +} + +# Validate file exists and is readable +validate_file() { + local file="$1" + local description="${2:-File}" + + if [[ ! -f "$file" ]]; then + log_error "$description not found: $file" + return 1 + fi + + if [[ ! -r "$file" ]]; then + log_error "$description not readable: $file" + return 1 + fi + + log_debug "Validated file: $file" + return 0 +} + +# ============================================================================== +# PATH DETECTION FUNCTIONS +# ============================================================================== + +# Find ROCm installation +find_rocm() { + # Try environment variable first + if [[ -n "${ROCM_PATH:-}" ]]; then + if validate_rocm_path "$ROCM_PATH"; then + echo "$ROCM_PATH" + return 0 + else + log_warn "ROCM_PATH is set but invalid: $ROCM_PATH" + fi + fi + + # Try common installation locations + local common_paths=( + "/opt/rocm" + "$HOME/rocm" + "/usr/local/rocm" + ) + + for path in "${common_paths[@]}"; do + if validate_rocm_path "$path"; then + log_debug "Found ROCm at: $path" + echo "$path" + return 0 + fi + done + + return 1 +} + +# Validate ROCm installation is complete +validate_rocm_path() { + local path="$1" + + [[ -d "$path" ]] || return 1 + [[ -x "$path/bin/hipcc" ]] || return 1 + [[ -d "$path/lib" ]] || return 1 + + return 0 +} + +# Find HeCBench source directory +find_hecbench_src() { + local script_dir="$1" + + # Try environment variable first + if [[ -n "${HECBENCH_SRC:-}" ]]; then + if validate_hecbench_src "$HECBENCH_SRC"; then + echo "$HECBENCH_SRC" + return 0 + else + log_warn "HECBENCH_SRC is set but invalid: $HECBENCH_SRC" + fi + fi + + # Try relative to script directory + local relative_paths=( + "$script_dir/../HeCBench/src" + "$script_dir/../../HeCBench/src" + "$script_dir/HeCBench/src" + ) + + for path in "${relative_paths[@]}"; do + local abs_path + abs_path=$(cd "$path" 2>/dev/null && pwd) + if [[ -n "$abs_path" ]] && validate_hecbench_src "$abs_path"; then + log_debug "Found HeCBench source at: $abs_path" + echo "$abs_path" + return 0 + fi + done + + return 1 +} + +# Validate HeCBench source directory +validate_hecbench_src() { + local path="$1" + + [[ -d "$path" ]] || return 1 + + # Check for characteristic HeCBench structure (HIP directories) + # Use a subshell to avoid pipefail issues with grep -q + local hip_dirs + hip_dirs=$(find "$path" -maxdepth 2 -name "*-hip" -type d 2>/dev/null | head -1) + [[ -n "$hip_dirs" ]] +} + +# Find MPI installation +find_mpi() { + # Try environment variable first + if [[ -n "${MPI_PATH:-}" ]]; then + if validate_mpi_path "$MPI_PATH"; then + echo "$MPI_PATH" + return 0 + else + log_warn "MPI_PATH is set but invalid: $MPI_PATH" + fi + fi + + # Try common installation locations + local common_paths=( + "/usr/lib/x86_64-linux-gnu/openmpi" + "/opt/openmpi" + "/usr/local/openmpi" + "/opt/mpi" + ) + + for path in "${common_paths[@]}"; do + if validate_mpi_path "$path"; then + log_debug "Found MPI at: $path" + echo "$path" + return 0 + fi + done + + return 1 +} + +# Validate MPI installation is complete +validate_mpi_path() { + local path="$1" + + [[ -d "$path" ]] || return 1 + [[ -d "$path/include" ]] || return 1 + [[ -d "$path/lib" ]] || return 1 + + # Check for mpi.h header + [[ -f "$path/include/mpi.h" ]] || return 1 + + return 0 +} + +# ============================================================================== +# UTILITY FUNCTIONS +# ============================================================================== + +# Execute command with dry-run support +execute_command() { + local description="$1" + shift + + if [[ "${DRY_RUN:-false}" == "true" ]]; then + log_info "[DRY-RUN] $description: $*" + return 0 + fi + + log_debug "Executing: $*" + "$@" +} + +# Retry a command with exponential backoff +retry_command() { + local max_attempts="$1" + local initial_delay="$2" + shift 2 + + local attempt=1 + local delay="$initial_delay" + + while [[ $attempt -le $max_attempts ]]; do + if "$@"; then + return 0 + fi + + if [[ $attempt -lt $max_attempts ]]; then + log_warn "Command failed (attempt $attempt/$max_attempts), retrying in ${delay}s..." + sleep "$delay" + delay=$((delay * 2)) + fi + + ((attempt++)) + done + + log_error "Command failed after $max_attempts attempts: $*" + return 1 +} + +# Show error and help message for missing path +show_path_error() { + local path_name="$1" + local env_var="$2" + local description="$3" + shift 3 + local required_contents=("$@") + + cat >&2 <&2 + done + + cat >&2 <&2 <&2 </dev/null; then + show_usage + exit 0 + else + log_error "show_usage function not defined" + exit 1 + fi + ;; + --dry-run) + DRY_RUN=true + log_info "Dry-run mode enabled" + shift + ;; + --verbose|-v) + VERBOSE=true + log_debug "Verbose mode enabled" + shift + ;; + *) + # Return remaining args for script-specific parsing + echo "$@" + return 0 + ;; + esac + done +} + +# ============================================================================== +# INTERACTIVE PROMPTS +# ============================================================================== + +# Prompt user for confirmation (yes/no) +# If INTERACTIVE=false, defaults to yes without prompting +# Usage: confirm_prompt "Do you want to continue?" && echo "Confirmed" +confirm_prompt() { + local prompt="$1" + local default="${2:-yes}" # yes or no + + # If not interactive, return based on default + if [[ "${INTERACTIVE:-true}" == "false" ]]; then + log_debug "Non-interactive mode: defaulting to '$default' for: $prompt" + [[ "$default" == "yes" ]] + return $? + fi + + # Interactive mode: ask user + local response + if [[ "$default" == "yes" ]]; then + echo -n "${BLUE}[?]${NC} $prompt [Y/n]: " + else + echo -n "${BLUE}[?]${NC} $prompt [y/N]: " + fi + read -r response < /dev/tty 2>/dev/null || response="" + + # Handle response + case "${response,,}" in + y|yes) + return 0 + ;; + n|no) + return 1 + ;; + "") + # Use default + [[ "$default" == "yes" ]] + return $? + ;; + *) + log_warn "Invalid response: '$response' (expected y/n)" + confirm_prompt "$prompt" "$default" # Retry + return $? + ;; + esac +} + +# ============================================================================== +# SHARED-MEMORY (/dev/shm) HOUSEKEEPING +# ============================================================================== +# OpenMPI's vader/sm BTL and RCCL leave per-rank backing files in /dev/shm +# (e.g. vader_segment...., nccl-XXXXXX). Normally these +# are unlinked when ranks exit cleanly, but a SIGKILL'd or timed-out run leaks +# them. On containers where /dev/shm is small (Docker default = 64 MiB), the +# leaks accumulate across runs until subsequent mpirun invocations fail with +# "not enough space for /dev/shm/...". +# +# Call this between benchmarks so every MPI/RCCL benchmark sees a clean slate. +# Scoped to the current user to avoid disturbing other tenants on shared +# systems. Silent on success; logs at debug level only. +cleanup_stale_shm() { + [[ -d /dev/shm ]] || return 0 + local user="${USER:-$(id -un)}" + # -user filter prevents touching other tenants' files on shared hosts. + # 2>/dev/null swallows the harmless ENOENT when the glob has no matches. + find /dev/shm -maxdepth 1 -user "$user" \ + \( -name 'vader_segment.*' -o -name 'nccl-*' \) \ + -delete 2>/dev/null || true +} + +# ============================================================================== +# BINARY VALIDATION +# ============================================================================== + +# Validate that a compiled binary actually contains the expected offload arch. +# Catches the catastrophic case where Makefiles ignore HIP_ARCH and hipcc falls +# back to the system default (e.g., building gfx90a when amdgcnspirv was requested). +# +# Usage: validate_arch +# Returns 0 if the binary contains the expected arch bundle, 1 otherwise. +validate_arch() { + local binary="$1" arch="$2" + if [[ ! -f "$binary" ]]; then + log_error "validate_arch: binary not found: $binary" + return 1 + fi + case "$arch" in + amdgcnspirv) + grep -qa 'hip-spirv64-amd-amdhsa--amdgcnspirv' "$binary" ;; + gfx*) + grep -qa "hip-amdgcn-amd-amdhsa--${arch}" "$binary" ;; + *) + log_warn "validate_arch: unknown arch '$arch', skipping validation" + return 0 ;; + esac +} + +# ============================================================================== +# BENCHMARK OUTPUT VALIDATION +# ============================================================================== + +# Scan a benchmark's log file for validation failures that the exit code missed. +# Many HeCBench benchmarks print PASS/FAIL but exit 0 regardless. +# Returns 0 if output looks clean, 1 if suspect lines were found. +# Writes the first few suspect lines to stdout for logging. +check_output_validation() { + local log="$1" + local name="${2:-}" + + [[ ! -s "$log" ]] && return 0 + + local hits + + hits=$( + grep -i -E \ + -e '\bFAIL(ED)?\b' \ + -e '\bMISMATCH\b' \ + -e '\bwrong result\b' \ + -e '\bdoes not match\b' \ + -e '\bTest Failed\b' \ + -e '\bincorrect\S*:\s*[1-9]' \ + -e 'hipblas\S+ failed' \ + -e 'rocblas\S+ failed' \ + -e ':\s*nan\b' \ + -e 'No such file or directory' \ + "$log" 2>/dev/null \ + | grep -v -i \ + -e '0 .* failed' \ + -e 'fail\(ures\?\|ed\):.\?[[:space:]]*0' \ + -e '\bno fail' \ + -e '\bwithout fail' \ + -e '\bif .*fail' \ + -e '\bon fail' \ + -e '#.*fail' \ + -e 'incorrect\S*:\s*0' \ + -e 'warning:' \ + -e 'note:' \ + -e '^\s*[0-9]\+\s*|' \ + -e '^/.*\.\(cpp\|c\|h\|hpp\|cc\)\b' \ + -e '^make\[' \ + -e '^hipcc' \ + -e '^amdclang' \ + -e '\.cpp:' \ + -e '\.cu:' \ + -e '\.h:' \ + -e 'hypre_assert' \ + -e 'hypre_error' \ + || true + ) + + if [[ -n "$hits" ]]; then + echo "$hits" | head -5 + return 1 + fi + + return 0 +} + +# ============================================================================== +# LIBRARY INITIALIZATION +# ============================================================================== + +# This library is now loaded +readonly COMMON_LIB_LOADED=true +log_debug "Common library loaded successfully" diff --git a/.github/hecbench/lib/presets.sh b/.github/hecbench/lib/presets.sh new file mode 100644 index 000000000..b764bc211 --- /dev/null +++ b/.github/hecbench/lib/presets.sh @@ -0,0 +1,11 @@ +# Preset benchmark lists based on gfx90a timings from 2026-07-15. +# Each is a cumulative set: quick ⊂ standard ⊂ extended ⊂ full. +# quick ~10min (179 benchmarks) +# standard ~30min (318 benchmarks) +# extended ~60min (417 benchmarks) + +PRESET_QUICK="accuracy-hip,ace-hip,adam-hip,addBiasQKV-hip,aes-hip,affine-hip,aligned-types-hip,amgmk-hip,ans-hip,aobench-hip,asta-hip,atan2-hip,atomicReduction-hip,babelstream-hip,backprop-hip,base64e-hip,bfs-hip,bicgstab-hip,binomial-hip,blockAccess-hip,blockexchange-hip,bscan-hip,bsw-hip,b+tree-hip,cc-hip,ccsd-trpdrv-hip,ced-hip,cfd-hip,chemv-hip,clenergy-hip,clock-hip,cmp-hip,cobahh-hip,collision-hip,colorwheel-hip,complex-hip,concurrentKernels-hip,conversion-hip,cooling-hip,cross-hip,d2q9-bgk-hip,damage-hip,daphne-hip,dct8x8-hip,debayer-hip,determinant-hip,dispatch-hip,distort-hip,divergence-hip,dropout-hip,dxtc2-hip,easyWave-hip,ecdh-hip,egs-hip,ert-hip,f16atomic-hip,face-hip,fft-hip,fhd-hip,floydwarshall-hip,fpc-hip,fpdc-hip,fwt-hip,gc-hip,gels-hip,geodesic-hip,gibbs-hip,graphExecution-hip,groupnorm-hip,haccmk-hip,haversine-hip,heartwall-hip,heat2d-hip,heat-hip,henry-hip,histogram-hip,hogbom-hip,hotspot-hip,hungarian-hip,interleave-hip,intrinsics-cast-hip,inversek2j-hip,is-hip,ising-hip,iso2dfd-hip,jenkins-hash-hip,lanczos-hip,langford-hip,laplace3d-hip,lavaMD-hip,layout-hip,lda-hip,libor-hip,local-ht-hip,log2-hip,logan-hip,logprob-hip,lombscargle-hip,lr-hip,lud-hip,mallocFree-hip,mandelbrot-hip,marchingCubes-hip,md5hash-hip,meanshift-hip,medianfilter-hip,minibude-hip,minimap2-hip,minimod-hip,miniWeather-hip,mis-hip,mixbench-hip,mpc-hip,mrc-hip,mrg32k3a-hip,mr-hip,murmurhash3-hip,myocyte-hip,ne-hip,nlll-hip,nms-hip,nn-hip,opticalFlow-hip,overlap-hip,pad-hip,particlefilter-hip,particles-hip,pathfinder-hip,permute-hip,perplexity-hip,phmm-hip,pitch-hip,projectile-hip,qtclustering-hip,quant3MatMul-hip,quantAQLM-hip,radixsort-hip,recursiveGaussian-hip,resnet-kernels-hip,reverse2D-hip,reverse-hip,ring-hip,rmsnorm-hip,rng-wallace-hip,rodrigues-hip,romberg-hip,rowwiseMoments-hip,rsc-hip,rtm8-hip,sad-hip,scan2-hip,scatterThrust-hip,sc-hip,seam-carving-hip,simplemoc-hip,sobel-hip,sph-hip,split-hip,spm-hip,sptrsv-hip,srad-hip,stencil1d-hip,streamcluster-hip,su3-hip,svd3x3-hip,tensorT-hip,threadfence-hip,tqs-hip,triad-hip,tsa-hip,unfold-hip,urng-hip,vmc-hip,warpexchange-hip,warpsort-hip,winograd-hip,xlqc-hip,zeropoint-hip,zmddft-hip" + +PRESET_STANDARD="accuracy-hip,ace-hip,adam-hip,adamw-hip,addBiasQKV-hip,aes-hip,affine-hip,aidw-hip,aligned-types-hip,allreduce-hip,amgmk-hip,ans-hip,aobench-hip,aop-hip,asmooth-hip,asta-hip,atan2-hip,atomicAggregate-hip,atomicPerf-hip,atomicReduction-hip,atomicSystemWide-hip,attentionMergeState-hip,attentionMultiHead-hip,axhelm-hip,babelstream-hip,backprop-hip,base64e-hip,bfs-hip,bgmv-hip,bh-hip,bicgstab-hip,bincount-hip,binomial-hip,bitpacking-hip,bitpermute-hip,black-scholes-hip,blas-fp8gemm-hip,blas-gemm-hip,blockAccess-hip,blockexchange-hip,bm3d-hip,bmf-hip,boxfilter-hip,bscan-hip,bsearch-hip,bspline-vgh-hip,bsw-hip,b+tree-hip,cbsfil-hip,cc-hip,ccsd-trpdrv-hip,ccs-hip,ced-hip,cfd-hip,chacha20-hip,che-hip,chemv-hip,clenergy-hip,clock-hip,cmembench-hip,cmp-hip,cobahh-hip,collision-hip,colorwheel-hip,complex-hip,compute-score-hip,concurrentKernels-hip,contract-hip,conversion-hip,convolutionSeparable-hip,cooling-hip,coordinates-hip,crc64-hip,cross-hip,crs-hip,d2q9-bgk-hip,damage-hip,daphne-hip,dct8x8-hip,ddbp-hip,debayer-hip,determinant-hip,dispatch-hip,distort-hip,divergence-hip,dpid-hip,dropout-hip,dslash-hip,dxtc2-hip,easyWave-hip,ecdh-hip,egs-hip,eigenvalue-hip,eikonal-hip,entropy-hip,ert-hip,expdist-hip,extend2-hip,extrema-hip,f16atomic-hip,f8cast-hip,face-hip,fdtd3d-hip,fft-hip,fhd-hip,flame-hip,floydwarshall2-hip,floydwarshall-hip,fma-hip,fpc-hip,fpdc-hip,fresnel-hip,fwt-hip,gabor-hip,ga-hip,gaussian-hip,gc-hip,gels-hip,gelu-hip,geodesic-hip,ge-spmm-hip,gerbil-hip,gibbs-hip,gmm-hip,gpp-hip,graphB+-hip,graphExecution-hip,groupnorm-hip,gru2-hip,haccmk-hip,halo-finder-hip,haversine-hip,heartwall-hip,heat2d-hip,heat-hip,hellinger-hip,henry-hip,hexciton-hip,histogram-hip,hmm-hip,hogbom-hip,hotspot-hip,hungarian-hip,hwt1d-hip,idivide-hip,interleave-hip,intrinsics-cast-hip,inversek2j-hip,is-hip,ising-hip,iso2dfd-hip,jacobi-hip,jenkins-hash-hip,kalman-hip,kiss-hip,kmc-hip,kurtosis-hip,lanczos-hip,langford-hip,laplace3d-hip,lavaMD-hip,layout-hip,lci-hip,lda-hip,ldpc-hip,leukocyte-hip,libor-hip,local-ht-hip,log2-hip,logan-hip,logprob-hip,lombscargle-hip,loopback-hip,lr-hip,ludb-hip,lud-hip,lulesh-hip,lzss-hip,mallocFree-hip,mandelbrot-hip,marchingCubes-hip,mask-hip,matrix-rotate-hip,maxpool3d-hip,mcpr-hip,md5hash-hip,mdh-hip,meanshift-hip,medianfilter-hip,memcpy-hip,mergeVS-hip,mf-sgd-hip,minibude-hip,miniFE-hip,minimap2-hip,minimod-hip,miniWeather-hip,minmax-hip,mis-hip,mixbench-hip,moe-hip,morphology-hip,mpc-hip,mrc-hip,mrg32k3a-hip,mr-hip,mriQ-hip,mt-hip,multinomial-hip,multimaterial-hip,murmurhash3-hip,myocyte-hip,nbnxm-hip,nbody-hip,ne-hip,nlll-hip,nms-hip,nn-hip,norm2-hip,ntt-hip,opticalFlow-hip,overlap-hip,overlay-hip,p2p-hip,p4-hip,pad-hip,particlefilter-hip,particles-hip,pathfinder-hip,permute-hip,perplexity-hip,phmm-hip,pitch-hip,pns-hip,popcount-hip,present-hip,projectile-hip,pso-hip,qkv-hip,qtclustering-hip,quant3MatMul-hip,quantAQLM-hip,quantBnB-hip,quantVLLM-hip,radixsort-hip,rainflow-hip,reaction-hip,recursiveGaussian-hip,resnet-kernels-hip,reverse2D-hip,reverse-hip,ring-hip,rle-hip,rmsnorm-hip,rng-wallace-hip,rodrigues-hip,romberg-hip,rowwiseMoments-hip,rsbench-hip,rsc-hip,rsmt-hip,rtm8-hip,rushlarsen-hip,s3d-hip,sad-hip,scan2-hip,scan3-hip,scatterThrust-hip,scel-hip,sc-hip,seam-carving-hip,secp256k1-hip,segment-reduce-hip,segsort-hip,sheath-hip,shmembench-hip,simplemoc-hip,simpleMultiDevice-hip,slit-hip,sobel-hip,sobol-hip,softmax-hip,sph-hip,split-hip,spm-hip,spmm-hip,spmv-hip,sptrsv-hip,srad-hip,ss-hip,sss-hip,stencil1d-hip,streamcluster-hip,su3-hip,svd3x3-hip,tensorAccessor-hip,tensorT-hip,threadfence-hip,tissue-hip,tonemapping-hip,tqs-hip,triad-hip,tsa-hip,tsp-hip,unfold-hip,urng-hip,vadd-hip,vanGenuchten-hip,vmc-hip,vote-hip,warpexchange-hip,warpsort-hip,winograd-hip,wlcpow-hip,wordcount-hip,wsm5-hip,wyllie-hip,xlqc-hip,zeropoint-hip,zmddft-hip,zoom-hip" + +PRESET_EXTENDED="accuracy-hip,ace-hip,adam-hip,adamw-hip,addBiasQKV-hip,adjacent-hip,adv-hip,aes-hip,affine-hip,aidw-hip,aligned-types-hip,all-pairs-distance-hip,allreduce-hip,amgmk-hip,ans-hip,aobench-hip,aop-hip,asmooth-hip,asta-hip,atan2-hip,atomicAggregate-hip,atomicCAS-hip,atomicCost-hip,atomicPerf-hip,atomicReduction-hip,atomicSystemWide-hip,attentionMergeState-hip,attentionMultiHead-hip,axhelm-hip,axpby-hip,babelstream-hip,background-subtract-hip,backprop-hip,base64e-hip,bezier-surface-hip,bfs-hip,bgmv-hip,bh-hip,bicgstab-hip,bincount-hip,binomial-hip,bitonic-sort-hip,bitpacking-hip,bitpermute-hip,black-scholes-hip,blas-dot-hip,blas-fp8gemm-hip,blas-gemm-hip,blas-mxfp8gemm-hip,blockAccess-hip,blockexchange-hip,bm3d-hip,bmf-hip,bonds-hip,boxfilter-hip,bscan-hip,bsearch-hip,bspline-vgh-hip,bsw-hip,b+tree-hip,burger-hip,bwt-hip,car-hip,cbsfil-hip,cc-hip,ccl-hip,ccsd-trpdrv-hip,ccs-hip,ced-hip,cfd-hip,chacha20-hip,che-hip,chemv-hip,clenergy-hip,clink-hip,clock-hip,cmembench-hip,cmp-hip,cobahh-hip,collision-hip,colorwheel-hip,complex-hip,compute-score-hip,concat-hip,concurrentKernels-hip,contract-hip,conversion-hip,convolutionSeparable-hip,cooling-hip,coordinates-hip,crc64-hip,crossEntropy-hip,cross-hip,crs-hip,d2q9-bgk-hip,d3q19-bgk-hip,damage-hip,daphne-hip,dct8x8-hip,ddbp-hip,debayer-hip,degrid-hip,depixel-hip,determinant-hip,dispatch-hip,distort-hip,divergence-hip,doh-hip,dpid-hip,dropout-hip,dslash-hip,dxtc2-hip,easyWave-hip,ecdh-hip,egs-hip,eigenvalue-hip,eikonal-hip,entropy-hip,ert-hip,expdist-hip,extend2-hip,extrema-hip,f16atomic-hip,f8cast-hip,face-hip,fdtd3d-hip,fft-hip,fhd-hip,filter-hip,flame-hip,flip-hip,floydwarshall2-hip,floydwarshall-hip,fluidSim-hip,fma-hip,fpc-hip,fpdc-hip,fresnel-hip,fwt-hip,gabor-hip,ga-hip,gamma-correction-hip,gaussian-hip,gc-hip,gd-hip,geam-hip,gels-hip,gelu-hip,gemv-hip,geodesic-hip,ge-spmm-hip,gerbil-hip,gibbs-hip,gmm-hip,goulash-hip,gpp-hip,graphB+-hip,graphExecution-hip,groupnorm-hip,gru2-hip,haccmk-hip,halo-finder-hip,hausdorff-hip,haversine-hip,hbc-hip,heartwall-hip,heat2d-hip,heat-hip,hellinger-hip,henry-hip,hexciton-hip,histogram-hip,hmm-hip,hogbom-hip,hotspot3D-hip,hotspot-hip,hungarian-hip,hwt1d-hip,hybridsort-hip,hypterm-hip,idivide-hip,interleave-hip,intrinsics-cast-hip,inversek2j-hip,is-hip,ising-hip,iso2dfd-hip,jacobi-hip,jenkins-hash-hip,kalman-hip,keccaktreehash-hip,keogh-hip,kernelLaunch-hip,kiss-hip,kmc-hip,knn-hip,kurtosis-hip,lanczos-hip,langevin-hip,langford-hip,laplace3d-hip,lavaMD-hip,layout-hip,lci-hip,lda-hip,ldpc-hip,lebesgue-hip,leukocyte-hip,libor-hip,lid-driven-cavity-hip,lif-hip,local-ht-hip,log2-hip,logan-hip,logic-resim-hip,logic-rewrite-hip,logprob-hip,lombscargle-hip,loopback-hip,lr-hip,ludb-hip,lud-hip,lulesh-hip,lzss-hip,mallocFree-hip,mandelbrot-hip,marchingCubes-hip,mask-hip,matrix-rotate-hip,matrixT-hip,maxpool3d-hip,mcpr-hip,md5hash-hip,mdh-hip,md-hip,meanshift-hip,medianfilter-hip,memcpy-hip,merge-hip,mergeVS-hip,merkle-hip,mf-sgd-hip,michalewicz-hip,minibude-hip,miniFE-hip,minimap2-hip,minimod-hip,miniWeather-hip,minkowski-hip,minmax-hip,mis-hip,mixbench-hip,mmcsf-hip,moe-hip,morphology-hip,mpc-hip,mrc-hip,mrg32k3a-hip,mr-hip,mriQ-hip,mtf-hip,multimaterial-hip,mt-hip,multinomial-hip,murmurhash3-hip,myocyte-hip,nbnxm-hip,nbody-hip,ne-hip,nlll-hip,nms-hip,nn-hip,norm2-hip,nosync-hip,ntt-hip,nw-hip,opticalFlow-hip,overlap-hip,overlay-hip,p2p-hip,p4-hip,pad-hip,page-rank-hip,particle-diffusion-hip,particlefilter-hip,particles-hip,pathfinder-hip,pcc-hip,perlin-hip,permute-hip,perplexity-hip,phmm-hip,pitch-hip,pns-hip,pointerchase-hip,pointwise-hip,pool-hip,popcount-hip,present-hip,projectile-hip,pso-hip,qem-hip,qkv-hip,qrg-hip,qtclustering-hip,quant3MatMul-hip,quantAQLM-hip,quantBnB-hip,quantVLLM-hip,quicksort-hip,radixsort-hip,rainflow-hip,randomAccess-hip,reaction-hip,recursiveGaussian-hip,relu-hip,reshapeKVCache-hip,resize-hip,resnet-kernels-hip,reverse2D-hip,reverse-hip,rfs-hip,ring-hip,rle-hip,rmsnorm-hip,rng-wallace-hip,rodrigues-hip,romberg-hip,rowwiseMoments-hip,rsbench-hip,rsc-hip,rsmt-hip,rtm8-hip,rushlarsen-hip,s3d-hip,sad-hip,sa-hip,sampling-hip,scan2-hip,scan3-hip,scatter-hip,scatterThrust-hip,scel-hip,sc-hip,score-hip,seam-carving-hip,secp256k1-hip,segment-reduce-hip,segsort-hip,sheath-hip,shmembench-hip,shuffle-hip,si-hip,simplemoc-hip,simpleMultiDevice-hip,simpleSpmv-hip,slit-hip,snake-hip,snicit-hip,sobel-hip,sobol-hip,softmax-hip,sort-hip,sparkler-hip,spgemm-hip,sph-hip,split-hip,spm-hip,spmm-hip,spmv-hip,sptrsv-hip,srad-hip,ss-hip,ssim-hip,sss-hip,sssp-hip,stddev-hip,stencil1d-hip,stencil3d-hip,streamcluster-hip,streamCreateCopyDestroy-hip,streamOrderedAllocation-hip,streamPriority-hip,streamUM-hip,su3-hip,surfel-hip,svd3x3-hip,tensorAccessor-hip,tensorT-hip,tgvnn-hip,thomas-hip,threadfence-hip,tissue-hip,tonemapping-hip,tpacf-hip,tqs-hip,triad-hip,tsa-hip,tsp-hip,twell-hip,unfold-hip,upsample-hip,urng-hip,vadd-hip,vanGenuchten-hip,vmc-hip,vol2col-hip,vote-hip,warpexchange-hip,warpsort-hip,wedford-hip,winograd-hip,wlcpow-hip,word2vec-hip,wordcount-hip,wsm5-hip,wyllie-hip,xlqc-hip,zeropoint-hip,zmddft-hip,zoom-hip" diff --git a/.github/hecbench/modify_makefiles.sh b/.github/hecbench/modify_makefiles.sh new file mode 100755 index 000000000..5179ec7cb --- /dev/null +++ b/.github/hecbench/modify_makefiles.sh @@ -0,0 +1,438 @@ +#!/usr/bin/env bash +# Script: modify_makefiles.sh +# Purpose: Modify HeCBench Makefiles to add HIP architecture support +# Usage: ./modify_makefiles.sh [OPTIONS] +# +# Replaces "hipcc" with "hipcc --offload-arch=$(HIP_ARCH) --rocm-device-lib-path=..." +# This enables multi-architecture builds for AMD GPUs. +# +# Uses HIP_ARCH (not ARCH) to avoid collisions with benchmarks that use ARCH for +# their own purposes (e.g. dp4a-hip uses ARCH = CDNA as a feature flag). + +set -euo pipefail + +# ============================================================================== +# CONSTANTS & DEFAULTS +# ============================================================================== + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +SCRIPT_NAME="$(basename "$0")" +DRY_RUN=false +VERBOSE=false + +# ============================================================================== +# SOURCE COMMON LIBRARY +# ============================================================================== + +# shellcheck source=lib/common.sh +source "${SCRIPT_DIR}/lib/common.sh" || { + echo "ERROR: Failed to load common library" >&2 + exit 1 +} + +# ============================================================================== +# CONFIGURATION +# ============================================================================== + +# REQUIRE ROCM_PATH to be explicitly set (no auto-detection) +if ! require_rocm_path; then + exit 1 +fi + +# Try to find HeCBench source (auto-detect is OK for this) +if ! HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR"); then + show_path_error "HeCBench" "HECBENCH_SRC" "HeCBench source directory" \ + "$SCRIPT_DIR/../HeCBench/src" "$SCRIPT_DIR/../../HeCBench/src" + exit 1 +fi + +DEVICE_LIB_PATH="${ROCM_PATH}/lib/llvm/amdgcn/bitcode" + +# ============================================================================== +# FUNCTIONS +# ============================================================================== + +show_usage() { + cat < "${makefile}.tmp" && mv "${makefile}.tmp" "$makefile" + elif ! grep -q "^HIPCC_BIN_DIR" "$makefile"; then + # Use printf to avoid shell expansion issues with sed + { + grep -B1000 "^HIP_ARCH.*=" "$makefile" | head -n -1 + grep "^HIP_ARCH.*=" "$makefile" + echo "" + echo "# HIP compiler location" + echo "HIPCC_BIN_DIR ?= $ROCM_PATH/bin" + grep -A1000 "^HIP_ARCH.*=" "$makefile" | tail -n +2 + } > "${makefile}.tmp" && mv "${makefile}.tmp" "$makefile" + fi + + # Escape device lib path for sed + local escaped_path="${DEVICE_LIB_PATH//\//\\/}" + + # Replace hipcc with $(HIPCC_BIN_DIR)/hipcc + flags + # Only replace standalone "hipcc" not already prefixed with path/variable + sed -i "s#\(^\|[^/})]\)\bhipcc\b#\1\$(HIPCC_BIN_DIR)/hipcc --offload-arch=\$(HIP_ARCH) --rocm-device-lib-path=${escaped_path} \$(EXTRA_HIPCCFLAGS)#g" "$makefile" + + # After the rewrite above, $(CC) / $(CXX) / $(HIPCC) typically expand to a + # multi-word string. Any nested-make recipe that passes `=$(CC)` + # unquoted would word-split, corrupting the inner build. Re-quote such + # assignments. Restricted to recipe lines (start with TAB) so top-level + # Makefile assignments are not altered. + local TAB + TAB=$'\t' + sed -i -E "s#^(${TAB}.*)([A-Z][A-Z_]+=)\\\$\\((CC|CXX|HIPCC)\\)#\\1\\2\"\\\$(\\3)\"#g" "$makefile" + + # Remove -save-temps flag + sed -i 's/-save-temps//g' "$makefile" + + # Fix HIP_PATH variable name conflict + local hip_path_renamed=false + if grep -q "^HIP_PATH\s*=" "$makefile"; then + sed -i -e 's/^\(HIP_PATH\s*=\)/HIP_SRC_PATH =/' \ + -e 's/\$(\(HIP_PATH\))/$(HIP_SRC_PATH)/g' \ + -e 's/\${\(HIP_PATH\)}/\${HIP_SRC_PATH}/g' "$makefile" + hip_path_renamed=true + fi + + # Log success with details + if [[ "$hip_path_renamed" == "true" ]]; then + log_success "Modified: $rel_path (renamed HIP_PATH → HIP_SRC_PATH)" + else + log_success "Modified: $rel_path" + fi + + return 0 +} + +process_makefiles() { + local -a makefiles + + # Get array of Makefiles + mapfile -t makefiles < <(find_makefiles) + if [[ ${#makefiles[@]} -eq 0 ]]; then + return 1 + fi + + log_info "Processing ${#makefiles[@]} Makefiles..." + echo "" + + local modified_count=0 + local skipped_already=0 + local skipped_no_hipcc=0 + local failed_count=0 + + for makefile in "${makefiles[@]}"; do + local rel_path="${makefile#$HECBENCH_SRC/}" + + modify_makefile "$makefile" + local result=$? + + case $result in + 0) + modified_count=$((modified_count + 1)) + ;; + 2) + log_info "Already modified: $rel_path" + skipped_already=$((skipped_already + 1)) + ;; + 3) + log_info "No hipcc usage: $rel_path" + skipped_no_hipcc=$((skipped_no_hipcc + 1)) + ;; + *) + log_error "Failed to modify: $rel_path" + failed_count=$((failed_count + 1)) + ;; + esac + done + + echo "" + log_info "========================================" + log_info "Modification Summary" + log_info "========================================" + log_info "Total Makefiles found: ${#makefiles[@]}" + log_success "Modified: $modified_count" + log_info "Already modified: $skipped_already" + log_info "No hipcc usage: $skipped_no_hipcc" + + if [[ $failed_count -gt 0 ]]; then + log_error "Failed: $failed_count" + fi + + # Show sample verification + if [[ $modified_count -gt 0 ]] && [[ "$DRY_RUN" == "false" ]]; then + echo "" + log_info "Sample verification:" + for makefile in "${makefiles[@]}"; do + if compgen -G "${makefile}.bak.*" >/dev/null; then + local backup + backup=$(ls -t "${makefile}.bak."* | head -1) + local rel_path="${makefile#$HECBENCH_SRC/}" + echo " File: $rel_path" + echo " Before: $(grep -m1 'hipcc' "$backup" || echo "(none)")" + echo " After: $(grep -m1 'hipcc' "$makefile" || echo "(none)")" + echo "" + break + fi + done + fi + + if [[ $failed_count -gt 0 ]]; then + return 2 + fi + + return 0 +} + +# ============================================================================== +# MAIN +# ============================================================================== + +main() { + # Parse arguments + local remaining_args + remaining_args=$(parse_common_flags "$SCRIPT_NAME" "$@") + + # No script-specific arguments expected + if [[ -n "$remaining_args" ]]; then + log_error "Unknown arguments: $remaining_args" + show_usage + exit 1 + fi + + log_info "========================================" + log_info "HeCBench Makefile Modifier" + log_info "========================================" + log_info "ROCm Path: $ROCM_PATH" + log_info "HeCBench Source: $HECBENCH_SRC" + log_info "Device Library: $DEVICE_LIB_PATH" + if [[ "$DRY_RUN" == "true" ]]; then + log_warn "DRY-RUN MODE: No files will be modified" + fi + echo "" + + # Validate environment + if ! validate_environment; then + exit 1 + fi + echo "" + + # Process Makefiles + if ! process_makefiles; then + exit 2 + fi + + echo "" + log_success "Makefile modification complete!" + echo "" +} + +main "$@" diff --git a/.github/hecbench/update_deps.sh b/.github/hecbench/update_deps.sh new file mode 100755 index 000000000..be65fd967 --- /dev/null +++ b/.github/hecbench/update_deps.sh @@ -0,0 +1,100 @@ +#!/usr/bin/env bash +# Regenerate ci_benchmarks.deps.txt from ci_benchmarks.txt. +# +# The benchmarks in ci_benchmarks.txt reference other in-tree HeCBench dirs via +# '../' (sibling *-cuda data dirs, shared include/, etc.). The test_hecbench +# CI job must check those out too, or the builds fail to resolve their inputs. +# This script resolves that set transitively at the pinned HECBENCH_REF and +# writes it to ci_benchmarks.deps.txt, so CI can do a single sparse checkout of +# ci_benchmarks.txt + ci_benchmarks.deps.txt instead of a discover-then-expand +# second pass. +# +# Re-run after editing ci_benchmarks.txt or bumping HECBENCH_REF, then commit the +# updated ci_benchmarks.deps.txt: +# .github/hecbench/update_deps.sh +# HECBENCH_REF defaults to the pin in ../workflows/spirv-ci-linux.yml; override +# via the env var. HECBENCH_REPO overrides the clone URL. +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +LIST="$SCRIPT_DIR/ci_benchmarks.txt" +OUT="$SCRIPT_DIR/ci_benchmarks.deps.txt" +REPO="${HECBENCH_REPO:-https://github.com/ORNL/HeCBench.git}" + +if [[ -z "${HECBENCH_REF:-}" ]]; then + HECBENCH_REF=$(grep -oP 'HECBENCH_REF:\s*\K\S+' \ + "$SCRIPT_DIR/../workflows/spirv-ci-linux.yml") +fi +[[ -n "$HECBENCH_REF" ]] || { echo "HECBENCH_REF not set and not found in workflow" >&2; exit 1; } + +mapfile -t benches < <(grep -vE '^[[:space:]]*(#|$)' "$LIST") +[[ ${#benches[@]} -gt 0 ]] || { echo "no benchmarks listed in $LIST" >&2; exit 1; } + +echo "Resolving deps for ${#benches[@]} benchmarks at $HECBENCH_REF ..." + +work=$(mktemp -d) +trap 'rm -rf "$work"' EXIT +git init -q "$work" +cd "$work" +git remote add origin "$REPO" +# Shallow, blob-less: fetch only the pinned commit's trees; blobs come lazily on +# checkout. Mirrors the CI checkout so this validates that path too. +git fetch -q --depth=1 --filter=blob:none origin "$HECBENCH_REF" +git sparse-checkout init --no-cone + +# want = set of src/ to materialize; seed with the benchmarks, then expand +# to their transitive '../' references until the set stops growing. +declare -A want=() +for b in "${benches[@]}"; do want["src/$b"]=1; done + +apply_sparse() { + local paths=() p + while IFS= read -r p; do paths+=("/$p/"); done \ + < <(printf '%s\n' "${!want[@]}" | sort) + git sparse-checkout set --no-cone "${paths[@]}" + git checkout -q FETCH_HEAD +} + +apply_sparse +changed=1 +while (( changed )); do + changed=0 + while IFS= read -r d; do + [[ -z "$d" ]] && continue + if [[ -z "${want[src/$d]:-}" ]]; then + want["src/$d"]=1 + changed=1 + fi + done < <( + for p in "${!want[@]}"; do + [[ -d "$p" ]] || continue + # '|| true': a dir with no '../' ref makes grep exit 1, which under + # pipefail+errexit would abort this subshell and truncate the scan. + grep -rhoE '\.\./[A-Za-z0-9_.-]+' "$p" 2>/dev/null | sed 's#\.\./##' || true + done | sort -u | grep -vE '^\.' || true ) + (( changed )) && apply_sparse +done + +declare -A isbench=() +for b in "${benches[@]}"; do isbench["$b"]=1; done + +{ + cat <<'EOF' +# GENERATED by update_deps.sh — do not edit by hand. +# In-tree HeCBench dirs referenced (transitively) via '../' by the benchmarks in +# ci_benchmarks.txt: sibling *-cuda data dirs, shared include/, etc. The +# test_hecbench CI job checks these out alongside the benchmarks so their builds +# resolve. Regenerate with .github/hecbench/update_deps.sh (see that script). +EOF + for p in $(printf '%s\n' "${!want[@]}" | sort); do + d="${p#src/}" + [[ -n "${isbench[$d]:-}" ]] && continue + # Only real dirs at this ref; drops false positives from '../' tokens in + # comments (e.g. NVIDIA SDK boilerplate paths in shrUtils.cu). + [[ -d "$p" ]] || continue + printf '%s\n' "$d" + done +} > "$OUT" + +ndeps=$(grep -cvE '^[[:space:]]*(#|$)' "$OUT" || true) +echo "Wrote $ndeps dependency dirs to $OUT" diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 6e0b61413..27090a895 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -14,6 +14,9 @@ env: # SHA (gated by .github/workflows/automerge-rocm-examples.yml). Keep this on # its own line: the bumper rewrites exactly this line via sed. ROCM_EXAMPLES_REF: 269e9068d6fd6e68b1cc5a76eb77744f88af42ec + # Pinned ORNL/HeCBench revision for the test_hecbench job (not a moving branch, + # so upstream churn can't fail unrelated PRs). Bump after re-validating the subset. + HECBENCH_REF: ab972efc2e3514bb7704c6685f2629f1638b8a50 LLVM_BUILD: build DEVLIBS_BUILD: build-device-libs COMGR_BUILD: build-comgr @@ -665,3 +668,115 @@ jobs: LD_LIBRARY_PATH: ${{ github.workspace }}/${{ env.STAGING }}/lib run: | ctest --test-dir examples-build-Libraries --output-on-failure -j $(nproc) + + # ===================================================================== + # Test - HeCBench (curated subset, amdgcnspirv-be backend) + # ===================================================================== + # Builds + runs a curated HeCBench subset through the in-tree SPIRV backend + # (the default codegen path) and gates on each benchmark's self-verification. + # Driver scripts and subset list are vendored under .github/hecbench. Budgeted + # at ~15 min to stay off the critical path. + test_hecbench: + name: Test HeCBench + needs: build + runs-on: linux-gfx942-1gpu-ccs-csp-ossci-rocm + timeout-minutes: 30 + container: + image: ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:702a5133851e6d1daf1207d2c9fbb01c2667914a5b6dc5a01faeb3ce66ea6421 + options: | + --device=/dev/kfd --device=/dev/dri --group-add video + + steps: + # Check out this repo at the PR commit for the vendored HeCBench scripts + # in .github/hecbench. Into a subdir so it doesn't collide with the build + # artifact, which untars build/ and staging/ at the workspace root. + - name: Checkout SPIRV-LLVM-Translator (PR head) + uses: actions/checkout@9c091bb21b7c1c1d1991bb908d89e4e9dddfe3e0 # v7.0.0 + with: + repository: ${{ github.event.pull_request.head.repo.full_name || github.repository }} + ref: ${{ github.event.pull_request.head.sha || github.sha }} + path: SPIRV-LLVM-Translator + fetch-depth: 1 + persist-credentials: false + + - name: Download build tree artifact + uses: actions/download-artifact@70fc10c6e5e1ce46ad2ea6f2b72d43f7d47b13c3 # v8.0.0 + with: + name: linux-build-tree + + - name: Untar build trees + run: tar -xmf linux-build-tree.tar + + - name: Install libnuma + run: dnf install -y numactl-libs + + # Single sparse, shallow blob:none pass over the subset + precomputed deps + # (ci_benchmarks.deps.txt from update_deps.sh): ~90 MB, not the full repo. + - name: Checkout HeCBench (pinned, sparse subset) + env: + SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench + run: | + git init -q HeCBench + cd HeCBench + git remote add origin https://github.com/ORNL/HeCBench.git + git fetch -q --depth=1 --filter=blob:none origin "$HECBENCH_REF" + git sparse-checkout init --no-cone + mapfile -t paths < <( + grep -hvE '^[[:space:]]*(#|$)' \ + "$SCRIPTS/ci_benchmarks.txt" "$SCRIPTS/ci_benchmarks.deps.txt" \ + | sed 's#^#/src/#; s#$#/#') + git sparse-checkout set --no-cone "${paths[@]}" + git checkout -q FETCH_HEAD + echo "Checked out $(find src -maxdepth 1 -type d -name '*-hip' | wc -l) -hip dirs" + + # Rewrite Makefiles for the staged ROCm, then build+run+self-verify. + # bench.sh always exits 0 (results land in its log dir); next step gates. + - name: Run HeCBench subset + env: + SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench + ROCM_PATH: ${{ github.workspace }}/${{ env.STAGING }} + HIP_CLANG_PATH: ${{ github.workspace }}/${{ env.LLVM_BUILD }}/bin + HECBENCH_SRC: ${{ github.workspace }}/HeCBench/src + run: | + # Surface the toolchain layout: does staging provide a usable hipcc? + echo "ROCM_PATH=$ROCM_PATH"; ls -la "$ROCM_PATH/bin" | grep -iE 'hip|clang' || true + "$SCRIPTS/modify_makefiles.sh" + FILTER=$(grep -vE '^[[:space:]]*(#|$)' "$SCRIPTS/ci_benchmarks.txt" | paste -sd,) + "$SCRIPTS/bench.sh" --gpu-id 0 --timeout 300 --filter "$FILTER" amdgcnspirv-be + + - name: Gate on HeCBench results + if: always() + env: + SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench + run: | + LOG_DIR=$(ls -dt "$SCRIPTS"/bench_logs_*_amdgcnspirv-be 2>/dev/null | head -1) + if [[ -z "$LOG_DIR" ]]; then + echo "::error::No HeCBench log dir produced (run step failed early)." + exit 1 + fi + echo "Log dir: $LOG_DIR" + for cat in success suspect failed timeout skipped; do + n=$( [[ -f "$LOG_DIR/$cat.log" ]] && wc -l < "$LOG_DIR/$cat.log" || echo 0 ) + echo " $cat: $n" + done + rc=0 + for cat in failed timeout suspect; do + if [[ -s "$LOG_DIR/$cat.log" ]]; then + echo "::group::$cat"; cat "$LOG_DIR/$cat.log"; echo "::endgroup::" + echo "::error::HeCBench $cat: $(wc -l < "$LOG_DIR/$cat.log") benchmark(s)." + rc=1 + fi + done + if [[ ! -s "$LOG_DIR/success.log" ]]; then + echo "::error::No HeCBench benchmark succeeded — toolchain/setup problem." + rc=1 + fi + exit $rc + + - name: Upload HeCBench logs + if: always() + uses: actions/upload-artifact@bbbca2ddaa5d8feaa63e36b76fdaad77386f024f # v7.0.0 + with: + name: hecbench-logs + path: SPIRV-LLVM-Translator/.github/hecbench/bench_logs_*_amdgcnspirv-be/ + if-no-files-found: ignore From ab5ba0c4935ca6e2bc733e7d1c0c3e2eb62556d9 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Mon, 17 Aug 2026 15:49:46 -0500 Subject: [PATCH 02/19] Fix error. --- .github/workflows/spirv-ci-linux.yml | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 27090a895..cb6848945 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -681,6 +681,11 @@ jobs: needs: build runs-on: linux-gfx942-1gpu-ccs-csp-ossci-rocm timeout-minutes: 30 + # This image defaults inline `run:` steps to `sh`; force bash so the steps + # below can use arrays, process substitution, and [[ ]]. + defaults: + run: + shell: bash container: image: ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:702a5133851e6d1daf1207d2c9fbb01c2667914a5b6dc5a01faeb3ce66ea6421 options: | From ddefa31fbbc9949fc1b1fb6b1a1f8b38510ccf64 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Tue, 18 Aug 2026 05:00:24 -0500 Subject: [PATCH 03/19] Build and stage hipcc, needed for HeCBench. --- .github/workflows/spirv-ci-linux.yml | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index cb6848945..4c9c56665 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -192,6 +192,15 @@ jobs: mkdir -p "$STAGING/bin" cp -a "$LLVM_BUILD/bin/amd-llvm-spirv" "$STAGING/bin/" + # Build hipcc from the amd-staging fork; test_hecbench compiles with it. + - name: Build and stage hipcc + run: | + cmake -G Ninja -S llvm-project/amd/hipcc -B build-hipcc \ + -DCMAKE_BUILD_TYPE=Release + ninja -C build-hipcc hipcc + mkdir -p "$STAGING/bin" + cp -a build-hipcc/hipcc "$STAGING/bin/hipcc" + # Create hipconfig script for rocPRIM/hipCUB detection. - name: Create hipconfig script run: | From c9d730d70ca3734b1046132ec5183132031cc319 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Tue, 18 Aug 2026 06:42:35 -0500 Subject: [PATCH 04/19] Set longer timeout --- .github/workflows/spirv-ci-linux.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 4c9c56665..4d222652e 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -689,7 +689,7 @@ jobs: name: Test HeCBench needs: build runs-on: linux-gfx942-1gpu-ccs-csp-ossci-rocm - timeout-minutes: 30 + timeout-minutes: 60 # This image defaults inline `run:` steps to `sh`; force bash so the steps # below can use arrays, process substitution, and [[ ]]. defaults: From ec2d2f5ba67fa0428f796b02a8732b73249d66fb Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Tue, 18 Aug 2026 09:14:15 -0500 Subject: [PATCH 05/19] Match rocm_examples' timeout. --- .github/workflows/spirv-ci-linux.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 4d222652e..5537618e8 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -689,7 +689,7 @@ jobs: name: Test HeCBench needs: build runs-on: linux-gfx942-1gpu-ccs-csp-ossci-rocm - timeout-minutes: 60 + timeout-minutes: 240 # This image defaults inline `run:` steps to `sh`; force bash so the steps # below can use arrays, process substitution, and [[ ]]. defaults: From 9f9c5f96eaae774a4f39d5734631b0a2e6574744 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Wed, 19 Aug 2026 01:05:59 -0500 Subject: [PATCH 06/19] Fix errors. --- .github/hecbench/bench.sh | 5 +- .github/hecbench/ci_benchmarks.deps.txt | 48 +++++--------- .github/hecbench/ci_benchmarks.txt | 88 +++++++++++++------------ .github/hecbench/modify_makefiles.sh | 2 +- 4 files changed, 67 insertions(+), 76 deletions(-) diff --git a/.github/hecbench/bench.sh b/.github/hecbench/bench.sh index 0a68e2178..13e3bcf5f 100755 --- a/.github/hecbench/bench.sh +++ b/.github/hecbench/bench.sh @@ -146,8 +146,9 @@ if [[ -n "${HIP_CLANG_PATH:-}" ]]; then export PATH="${HIP_CLANG_PATH}:${PATH}" export LD_LIBRARY_PATH="$(dirname "${HIP_CLANG_PATH}")/lib:${LD_LIBRARY_PATH}" fi -# hipcc/amdclang++ honor this for AMDGCN bitcode lookup -export HIP_DEVICE_LIB_PATH="${ROCM_PATH}/lib/llvm/amdgcn/bitcode" +# Runtime JIT (comgr) honors this for AMDGCN bitcode lookup; device libs are +# installed at the canonical ROCm layout ${ROCM_PATH}/amdgcn/bitcode. +export HIP_DEVICE_LIB_PATH="${ROCM_PATH}/amdgcn/bitcode" export ROCR_VISIBLE_DEVICES="$GPU_ID" # prna-hip reads DATAPATH; overridable, harmless default for the rest. export DATAPATH="${DATAPATH:-${HECBENCH_SRC}/prna-cuda/data_tables}" diff --git a/.github/hecbench/ci_benchmarks.deps.txt b/.github/hecbench/ci_benchmarks.deps.txt index 4617c2274..6613a4d69 100644 --- a/.github/hecbench/ci_benchmarks.deps.txt +++ b/.github/hecbench/ci_benchmarks.deps.txt @@ -8,65 +8,54 @@ ace-cuda adam-cuda adamw-cuda addBiasQKV-cuda -addBiasResidualLayerNorm-cuda affine-cuda -aidw-cuda +atomicPerf-cuda atomicReduction-cuda -attentionMultiHead-cuda backprop-cuda base64e-cuda -bfs-sycl blockAccess-cuda boxfilter-sycl -bsw-cuda -ccs-cuda ccsd-trpdrv-cuda -cfd-cuda -channelShuffle-cuda channelSum-cuda che-cuda -cmp-cuda cobahh-cuda complex-cuda damage-cuda data debayer-sycl -deredundancy-sycl dpid-cuda dxtc2-sycl easyWave-omp ecdh-cuda -egs-cuda extend2-sycl -f16atomic-cuda +extrema-cuda fft-cuda fpdc-cuda fwt-cuda -geodesic-cuda -geodesic-sycl -gibbs-cuda +gabor-cuda +ga-cuda groupnorm-cuda -haversine-cuda -hogbom-cuda -hotspot-cuda -include +hellinger-cuda +idivide-cuda inversek2j-sycl is-cuda +kiss-cuda laplace3d-cuda log2-cuda -logan-cuda lud-cuda marchingCubes-cuda medianfilter-cuda medianfilter-sycl -minimap2-sycl mis-cuda -mpc-cuda mrc-cuda mr-cuda +multinomial-cuda +nbnxm-cuda nlll-cuda nms-cuda opticalFlow-cuda +overlay-cuda +p4-cuda pad-cuda particlefilter-cuda particles-cuda @@ -74,20 +63,15 @@ perplexity-cuda phmm-cuda projectile-cuda pso-cuda -qkv-hip -resnet-kernels-cuda +reaction-cuda +recursiveGaussian-cuda rmsnorm-cuda romberg-cuda +rowwiseMoments-cuda rsc-cuda sc-cuda -slit-cuda -sobel-sycl -spaxpby-cuda -sptrsv-sycl -svd3x3-cuda tqs-cuda tsa-cuda -urng-cuda -urng-sycl -xlqc-cuda +wyllie-cuda zeropoint-cuda +zoom-cuda diff --git a/.github/hecbench/ci_benchmarks.txt b/.github/hecbench/ci_benchmarks.txt index ff37fcb03..49f6268f0 100644 --- a/.github/hecbench/ci_benchmarks.txt +++ b/.github/hecbench/ci_benchmarks.txt @@ -4,18 +4,20 @@ # # Derivation: 2026-07-15 -be passing set (525), minus benchmarks needing # external data (download_datasets.sh / DVC / extract_archives.sh), plus the 23 -# regression-sensitive ones, filled cheapest-first under a ~900s budget = 161 -# (miniWeather-hip dropped: needs OpenMPI, not worth the dependency for one bench). +# regression-sensitive ones, filled cheapest-first under a ~900s budget. # Timings are a gfx90a proxy; recalibrate on gfx942 before making this blocking. +# +# 2026-08-18: replaced 39 entries that could not run under the hermetic CI setup +# (DVC/tarball input data, or missing -lomp / hipblas / hipfft / hipsparse / +# hipsolver / gsl / boost / numpy, or the f16atomic backend crash) with 39 +# hermetic benchmarks validated to build + self-verify on the amdgcnspirv-be +# path in the pinned CI container (local gfx90a proxy). accuracy-hip ace-hip adam-hip adamw-hip addBiasQKV-hip -addBiasResidualLayerNorm-hip -aes-hip affine-hip -aidw-hip aligned-types-hip amgmk-hip ans-hip @@ -23,27 +25,18 @@ aobench-hip asta-hip atan2-hip atomicReduction-hip -attentionMultiHead-hip -b+tree-hip babelstream-hip backprop-hip base64e-hip -bfs-hip binomial-hip blockexchange-hip bscan-hip -bsw-hip -ccs-hip ccsd-trpdrv-hip -ced-hip -cfd-hip -channelShuffle-hip channelSum-hip che-hip chemv-hip clenergy-hip clock-hip -cmp-hip cobahh-hip collision-hip colorwheel-hip @@ -52,12 +45,9 @@ concurrentKernels-hip conversion-hip cooling-hip cross-hip -d2q9-bgk-hip damage-hip dct8x8-hip debayer-hip -deredundancy-hip -determinant-hip dispatch-hip distort-hip divergence-hip @@ -66,26 +56,17 @@ dropout-hip dxtc2-hip easyWave-hip ecdh-hip -egs-hip extend2-hip -f16atomic-hip fft-hip -fhd-hip floydwarshall-hip fpc-hip fpdc-hip fwt-hip gc-hip -gels-hip -gibbs-hip graphExecution-hip groupnorm-hip -haversine-hip -heartwall-hip heat-hip histogram-hip -hogbom-hip -hotspot-hip hungarian-hip interleave-hip intrinsics-cast-hip @@ -99,7 +80,6 @@ layout-hip lda-hip libor-hip log2-hip -logan-hip lombscargle-hip lud-hip mallocFree-hip @@ -107,10 +87,8 @@ mandelbrot-hip marchingCubes-hip md5hash-hip medianfilter-hip -minimap2-hip mis-hip mixbench-hip -mpc-hip mr-hip mrc-hip mrg32k3a-hip @@ -119,14 +97,12 @@ myocyte-hip ne-hip nlll-hip nms-hip -nn-hip opticalFlow-hip overlap-hip pad-hip particlefilter-hip particles-hip pathfinder-hip -permute-hip perplexity-hip phmm-hip projectile-hip @@ -135,7 +111,6 @@ qtclustering-hip quant3MatMul-hip quantAQLM-hip radixsort-hip -resnet-kernels-hip reverse-hip reverse2D-hip rmsnorm-hip @@ -144,27 +119,58 @@ romberg-hip rsc-hip sc-hip scan2-hip -si-hip -slit-hip -sobel-hip -spaxpby-hip split-hip spm-hip -sptrsv-hip srad-hip stencil1d-hip su3-hip -svd3x3-hip threadfence-hip tqs-hip triad-hip tsa-hip unfold-hip -urng-hip vmc-hip warpexchange-hip -warpsort-hip winograd-hip -xlqc-hip zeropoint-hip zmddft-hip +# --- 2026-08-18 hermetic replacements (cheapest-first) --- +asyncAllocation-hip +tensorT-hip +rowwiseMoments-hip +iso2dfd-hip +kiss-hip +zoom-hip +wyllie-hip +blockAccess-hip +haccmk-hip +pitch-hip +overlay-hip +scatterThrust-hip +ert-hip +popcount-hip +black-scholes-hip +present-hip +bsearch-hip +multinomial-hip +nbnxm-hip +p4-hip +ga-hip +maxpool3d-hip +idivide-hip +sheath-hip +crc64-hip +matrix-rotate-hip +f8cast-hip +atomicSystemWide-hip +reaction-hip +sph-hip +compute-score-hip +gabor-hip +ldpc-hip +extrema-hip +recursiveGaussian-hip +shmembench-hip +atomicPerf-hip +hellinger-hip +attentionMergeState-hip diff --git a/.github/hecbench/modify_makefiles.sh b/.github/hecbench/modify_makefiles.sh index 5179ec7cb..5b287ff07 100755 --- a/.github/hecbench/modify_makefiles.sh +++ b/.github/hecbench/modify_makefiles.sh @@ -46,7 +46,7 @@ if ! HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR"); then exit 1 fi -DEVICE_LIB_PATH="${ROCM_PATH}/lib/llvm/amdgcn/bitcode" +DEVICE_LIB_PATH="${ROCM_PATH}/amdgcn/bitcode" # ============================================================================== # FUNCTIONS From 749e2bd8ed662f796694c11bb4df8fb6d7adf41e Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Thu, 27 Aug 2026 09:45:54 -0500 Subject: [PATCH 07/19] Add temporary instrumentation to debug benchmark failures --- .github/workflows/spirv-ci-linux.yml | 32 ++++++++++++++++++++++++++++ 1 file changed, 32 insertions(+) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 5537618e8..7a6c40136 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -743,6 +743,25 @@ jobs: git checkout -q FETCH_HEAD echo "Checked out $(find src -maxdepth 1 -type d -name '*-hip' | wc -l) -hip dirs" + # TEMPORARY diagnostic instrumentation for two MI300-only failures that + # pass on gfx90a/gfx942 (shmembench-hip, ans-hip). Prints the actual + # checksum sum and the first decode mismatch so the artifact logs reveal + # how far off MI300 is. Remove once the divergence is root-caused. + - name: Instrument diagnostics (temporary) + env: + HECBENCH_SRC: ${{ github.workspace }}/HeCBench/src + run: | + shm="$HECBENCH_SRC/shmembench-hip/shmem_kernels.cu" + grep -q '\[DIAG shmembench\]' "$shm" || perl -pi -e \ + 's{^(\s*)if \(sum != 21256458760384741137729978368\.00\)}{$1printf("[DIAG shmembench] sum=%.6f diff=%.6e\\n", sum, sum - 21256458760384741137729978368.00);\n$1if (sum != 21256458760384741137729978368.00)}' \ + "$shm" + ans="$HECBENCH_SRC/ans-hip/src/main.cu" + grep -q '\[DIAG ans\]' "$ans" || perl -pi -e \ + 's{(else std::cout << "\*+ MISMATCH \*+" << std::endl;)}{$1\n { auto* _a=random_data->data(); auto* _b=output_buffer->get_decompressed_data().get(); size_t _n=0,_f=input_size; for(size_t _i=0;_i/dev/null && rocminfo | head -40 || echo "rocminfo: not available" + command -v rocm-smi >/dev/null && rocm-smi --showproductname --showfwinfo || echo "rocm-smi: not available" + cat /proc/driver/amdgpu/version 2>/dev/null || echo "amdgpu procfs: not available" + for n in /sys/class/kfd/kfd/topology/nodes/*/; do + [ -r "$n/name" ] || continue + echo "node $(basename "$n"): name=$(cat "$n/name" 2>/dev/null)" + grep -E 'gfx_target_version|simd_count|cu_|max_engine_clk' "$n/properties" 2>/dev/null || true + done + echo "::endgroup::" # Surface the toolchain layout: does staging provide a usable hipcc? echo "ROCM_PATH=$ROCM_PATH"; ls -la "$ROCM_PATH/bin" | grep -iE 'hip|clang' || true "$SCRIPTS/modify_makefiles.sh" From 971a86936c3cf383fbd8333fbdb96e18fd3a89e2 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Thu, 27 Aug 2026 13:46:23 -0500 Subject: [PATCH 08/19] Quarantine benchmark until fix lands upstream --- .github/hecbench/ci_benchmarks.deps.txt | 1 - .github/hecbench/ci_benchmarks.txt | 4 +++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/.github/hecbench/ci_benchmarks.deps.txt b/.github/hecbench/ci_benchmarks.deps.txt index 6613a4d69..e1414909b 100644 --- a/.github/hecbench/ci_benchmarks.deps.txt +++ b/.github/hecbench/ci_benchmarks.deps.txt @@ -50,7 +50,6 @@ mis-cuda mrc-cuda mr-cuda multinomial-cuda -nbnxm-cuda nlll-cuda nms-cuda opticalFlow-cuda diff --git a/.github/hecbench/ci_benchmarks.txt b/.github/hecbench/ci_benchmarks.txt index 49f6268f0..1f9e0d647 100644 --- a/.github/hecbench/ci_benchmarks.txt +++ b/.github/hecbench/ci_benchmarks.txt @@ -153,7 +153,9 @@ black-scholes-hip present-hip bsearch-hip multinomial-hip -nbnxm-hip +# QUARANTINED: times out on amdgcnspirv-be (hipMallocManaged). Un-quarantine +# after ORNL/HeCBench#326 lands and HECBENCH_REF is bumped past it. +#nbnxm-hip p4-hip ga-hip maxpool3d-hip From 404460aec16f17f3b3586110d5105519247cf5ee Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Thu, 27 Aug 2026 13:46:31 -0500 Subject: [PATCH 09/19] Revert "Add temporary instrumentation to debug benchmark failures" This reverts commit 749e2bd8ed662f796694c11bb4df8fb6d7adf41e. --- .github/workflows/spirv-ci-linux.yml | 32 ---------------------------- 1 file changed, 32 deletions(-) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 7e58ccd58..1ce149a2f 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -743,25 +743,6 @@ jobs: git checkout -q FETCH_HEAD echo "Checked out $(find src -maxdepth 1 -type d -name '*-hip' | wc -l) -hip dirs" - # TEMPORARY diagnostic instrumentation for two MI300-only failures that - # pass on gfx90a/gfx942 (shmembench-hip, ans-hip). Prints the actual - # checksum sum and the first decode mismatch so the artifact logs reveal - # how far off MI300 is. Remove once the divergence is root-caused. - - name: Instrument diagnostics (temporary) - env: - HECBENCH_SRC: ${{ github.workspace }}/HeCBench/src - run: | - shm="$HECBENCH_SRC/shmembench-hip/shmem_kernels.cu" - grep -q '\[DIAG shmembench\]' "$shm" || perl -pi -e \ - 's{^(\s*)if \(sum != 21256458760384741137729978368\.00\)}{$1printf("[DIAG shmembench] sum=%.6f diff=%.6e\\n", sum, sum - 21256458760384741137729978368.00);\n$1if (sum != 21256458760384741137729978368.00)}' \ - "$shm" - ans="$HECBENCH_SRC/ans-hip/src/main.cu" - grep -q '\[DIAG ans\]' "$ans" || perl -pi -e \ - 's{(else std::cout << "\*+ MISMATCH \*+" << std::endl;)}{$1\n { auto* _a=random_data->data(); auto* _b=output_buffer->get_decompressed_data().get(); size_t _n=0,_f=input_size; for(size_t _i=0;_i/dev/null && rocminfo | head -40 || echo "rocminfo: not available" - command -v rocm-smi >/dev/null && rocm-smi --showproductname --showfwinfo || echo "rocm-smi: not available" - cat /proc/driver/amdgpu/version 2>/dev/null || echo "amdgpu procfs: not available" - for n in /sys/class/kfd/kfd/topology/nodes/*/; do - [ -r "$n/name" ] || continue - echo "node $(basename "$n"): name=$(cat "$n/name" 2>/dev/null)" - grep -E 'gfx_target_version|simd_count|cu_|max_engine_clk' "$n/properties" 2>/dev/null || true - done - echo "::endgroup::" # Surface the toolchain layout: does staging provide a usable hipcc? echo "ROCM_PATH=$ROCM_PATH"; ls -la "$ROCM_PATH/bin" | grep -iE 'hip|clang' || true "$SCRIPTS/modify_makefiles.sh" From 95f4fbcc00c8ca29a5d68139698fbf12f36e20db Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Fri, 28 Aug 2026 05:22:44 -0500 Subject: [PATCH 10/19] Remove unused code. --- .github/hecbench/bench.sh | 299 +++++-------------- .github/hecbench/lib/common.sh | 421 ++------------------------- .github/hecbench/lib/presets.sh | 11 - .github/hecbench/modify_makefiles.sh | 343 +++------------------- .github/workflows/spirv-ci-linux.yml | 2 +- 5 files changed, 136 insertions(+), 940 deletions(-) delete mode 100644 .github/hecbench/lib/presets.sh diff --git a/.github/hecbench/bench.sh b/.github/hecbench/bench.sh index 13e3bcf5f..4d1e79be6 100755 --- a/.github/hecbench/bench.sh +++ b/.github/hecbench/bench.sh @@ -1,5 +1,7 @@ #!/usr/bin/env bash -# Build and run HeCBench HIP benchmarks in one step (see show_usage below). +# Build and run the HeCBench HIP subset for the amdgcnspirv SPIR-V backend, +# self-verifying each benchmark's output. Results land in a bench_logs_* dir; +# the CI "Gate on HeCBench results" step reads those logs. set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" @@ -7,64 +9,45 @@ SCRIPT_NAME="$(basename "$0")" # shellcheck source=lib/common.sh source "${SCRIPT_DIR}/lib/common.sh" -# shellcheck source=lib/presets.sh -source "${SCRIPT_DIR}/lib/presets.sh" # ============================================================================== # DEFAULTS # ============================================================================== +# "amdgcnspirv-be" is the user-facing token (log dir suffix, prints); the make +# HIP_ARCH is plain amdgcnspirv, built via the default SPIR-V backend. +ARCH="amdgcnspirv-be" +HIPCC_ARCH="amdgcnspirv" + TIMEOUT_SECONDS=${TIMEOUT_SECONDS:-300} -GPU_ID=${GPU_ID:-0,1} # Use both GPUs by default for multi-GPU benchmarks +GPU_ID=${GPU_ID:-0} BENCHMARK_FILTER=${BENCHMARK_FILTER:-} -DRY_RUN=false show_usage() { cat </ - timings.csv benchmark,build_seconds,run_seconds - success.log benchmarks that built and ran cleanly - failed.log non-zero exit (build or run) - suspect.log exit 0 but output contains validation failures - (FAIL, MISMATCH, nan, missing input data, etc.) - timeout.log exceeded --timeout - skipped.log architecture-unsupported (e.g. saxpy-ompt-hip on amdgcnspirv) - .log per-benchmark combined build+run output + bench_logs_YYYYMMDD_HHMMSS_amdgcnspirv-be/ + timings.csv benchmark,build_seconds,run_seconds + success.log built and ran cleanly + failed.log non-zero exit (build or run) + suspect.log exit 0 but output contains validation failures + timeout.log exceeded --timeout + .log per-benchmark combined build+run output EOF } @@ -72,57 +55,22 @@ EOF # ARG PARSING # ============================================================================== -ARCH="" -PRESET="" while [[ $# -gt 0 ]]; do case "$1" in - --help|-h) show_usage; exit 0 ;; - --preset) PRESET="$2"; shift 2 ;; - --dry-run) DRY_RUN=true; shift ;; - --filter) BENCHMARK_FILTER="$2"; shift 2 ;; - --timeout) TIMEOUT_SECONDS="$2"; shift 2 ;; - --gpu-id) GPU_ID="$2"; shift 2 ;; - -*) log_error "Unknown option: $1"; show_usage; exit 1 ;; - *) - if [[ -z "$ARCH" ]]; then - ARCH="$1" - else - log_error "Multiple HIP_ARCH values; this script accepts one" - exit 1 - fi - shift - ;; + --help|-h) show_usage; exit 0 ;; + --filter) BENCHMARK_FILTER="$2"; shift 2 ;; + --timeout) TIMEOUT_SECONDS="$2"; shift 2 ;; + --gpu-id) GPU_ID="$2"; shift 2 ;; + *) log_error "Unknown argument: $1"; show_usage; exit 1 ;; esac done -# Resolve --preset into BENCHMARK_FILTER (--filter takes precedence). -if [[ -n "$PRESET" && -z "$BENCHMARK_FILTER" ]]; then - case "$PRESET" in - quick) BENCHMARK_FILTER="$PRESET_QUICK" ;; - standard) BENCHMARK_FILTER="$PRESET_STANDARD" ;; - extended) BENCHMARK_FILTER="$PRESET_EXTENDED" ;; - *) - log_error "Unknown preset: $PRESET (choose: quick, standard, extended)" - exit 1 - ;; - esac -fi - -if [[ -z "$ARCH" ]]; then - log_error "HIP_ARCH is required" +if [[ -z "$BENCHMARK_FILTER" ]]; then + log_error "--filter is required" show_usage exit 1 fi -# $ARCH is the user-facing token (log dir, prints); HIPCC_ARCH is what make gets -# as HIP_ARCH=. "amdgcnspirv-be" is our alias for the default SPIRV backend, so -# it maps to plain amdgcnspirv; every other arch passes through unchanged. -if [[ "$ARCH" == "amdgcnspirv-be" ]]; then - HIPCC_ARCH="amdgcnspirv" -else - HIPCC_ARCH="$ARCH" -fi - # ============================================================================== # VALIDATION & SETUP # ============================================================================== @@ -150,16 +98,21 @@ fi # installed at the canonical ROCm layout ${ROCM_PATH}/amdgcn/bitcode. export HIP_DEVICE_LIB_PATH="${ROCM_PATH}/amdgcn/bitcode" export ROCR_VISIBLE_DEVICES="$GPU_ID" -# prna-hip reads DATAPATH; overridable, harmless default for the rest. -export DATAPATH="${DATAPATH:-${HECBENCH_SRC}/prna-cuda/data_tables}" ulimit -s unlimited 2>/dev/null || true +# libhipcxx hard-errors on amdgcnspirv because no compile-time __gfx*__ +# macro is defined for SPIR-V; allow it through and silence the warning. NDEBUG +# keeps host-side asserts out of timing paths. Passed to every benchmark build. +declare -a MAKE_ARGS=( + HIP_ARCH="$HIPCC_ARCH" + "EXTRA_CFLAGS=-D_LIBCUDACXX_ALLOW_UNSUPPORTED_ARCHITECTURE -DNDEBUG" +) + echo "benchmark,build_seconds,run_seconds" > "$LOG_DIR/timings.csv" touch "$LOG_DIR/success.log" \ "$LOG_DIR/failed.log" \ "$LOG_DIR/suspect.log" \ - "$LOG_DIR/timeout.log" \ - "$LOG_DIR/skipped.log" + "$LOG_DIR/timeout.log" log_info "===========================================" log_info "HeCBench bench (build+run merged)" @@ -186,86 +139,6 @@ echo "" # PER-BENCHMARK # ============================================================================== -# Architecture-incompatibility skip list. -should_skip() { - local name="$1" - # saxpy-ompt-hip needs a concrete GPU ISA; spirv64-amd-amdhsa not in toolchain. - [[ "$name" == "saxpy-ompt-hip" && "$HIPCC_ARCH" == "amdgcnspirv" ]] && return 0 - # TEMP: pingpong-hip hangs in MPI/NCCL (pre-existing issue, unrelated to - # current TheRock bring-up); burns the full timeout for nothing. - [[ "$name" == "pingpong-hip" ]] && return 0 - [[ "$name" == "assert-hip" ]] && return 0 - - return 1 -} - -# Benchmarks that need special environment overrides for `make run`. -apply_special_env() { - local name="$1" - case "$name" in - assert-hip) - # This benchmark intentionally triggers a device-side assertion to - # verify host-side error reporting. Suppress the ROCm runtime's GPU - # coredump on exception so the apport pipe in core_pattern is not - # invoked. Keep coredumps enabled for every other benchmark. - export HSA_DISABLE_COREDUMP_ON_EXCEPTION=1 - ;; - esac -} - -clear_special_env() { - local name="$1" - case "$name" in - assert-hip) - unset HSA_DISABLE_COREDUMP_ON_EXCEPTION - ;; - esac -} - -# Extra `make` command-line arguments specific to a benchmark. Used for -# benchmarks whose Makefile reads a variable other than HIP_ARCH for the -# offload arch (e.g. saxpy-ompt-hip uses ARCH for OpenMP `-march=$(ARCH)`). -# Returns the extra args via stdout, one per line. -extra_make_args() { - local name="$1" - case "$name" in - saxpy-ompt-hip) - echo "ARCH=$ARCH" - ;; - esac - # libhipcxx hard-errors on amdgcnspirv because no compile-time - # __gfx*__ macro is defined for SPIR-V; allow it through and silence the - # accompanying warning. Host-side std::chrono is unaffected. - # Applies to both SPIRV paths. Plain amdgcnspirv also opts out of the - # (now default) backend to exercise the legacy translator. - if [[ "$HIPCC_ARCH" == "amdgcnspirv" ]]; then - echo "EXTRA_CFLAGS=-D_LIBCUDACXX_ALLOW_UNSUPPORTED_ARCHITECTURE -DNDEBUG" - [[ "$ARCH" == "amdgcnspirv" ]] && echo "EXTRA_HIPCCFLAGS=-no-use-spirv-backend" - fi -} - -# Per-benchmark direct-run override. When non-empty, bench_one builds with -# `make` (default target) and then invokes each emitted command directly, -# bypassing the Makefile's `run:` recipe. Each line is "binary arg1 arg2 ..." -# parsed with `read -ra`; the binary is resolved relative to $dir. -# -# Use this when the Makefile's `run:` recipe includes a config that exceeds -# this system's resources (e.g. attention-paged-hip's 131072-block case OOMs -# on MI210). Keeps coverage of the configs that fit without an upstream patch. -direct_run_argsets() { - local name="$1" - local arch="${2:-$ARCH}" - case "$name" in - attention-paged-hip) - # 4th make-run config (131072 kv blocks) OOMs; substitute 65536. - echo "./main 8 32 128 4096 128 100" - echo "./main 8 32 128 4096 1024 100" - echo "./main 8 32 128 4096 8192 100" - echo "./main 8 32 128 4096 65536 100" - ;; - esac -} - # check_output_validation() is defined in lib/common.sh. bench_one() { @@ -273,37 +146,21 @@ bench_one() { local name; name=$(basename "$dir") local log="$LOG_DIR/${name}.log" - if should_skip "$name"; then - log_info "Skipped: $name (arch unsupported)" - echo "$name" >> "$LOG_DIR/skipped.log" - return - fi - - if [[ "$DRY_RUN" == "true" ]]; then - log_info "[DRY-RUN] Would build+run: $name" - return - fi - # Reclaim leaked OpenMPI/RCCL backing files in /dev/shm so an MPI benchmark # whose predecessor was SIGKILL'd doesn't fail with "not enough space". cleanup_stale_shm cd "$dir" || { log_error "$name: cannot cd to $dir"; return; } - apply_special_env "$name" # Clean to ensure a deterministic build state. make clean &>/dev/null || true local build_start build_elapsed run_start run_elapsed rc - local -a make_extra=() - while IFS= read -r arg; do - [[ -n "$arg" ]] && make_extra+=("$arg") - done < <(extra_make_args "$name") # ---- Phase 1: BUILD (default target only) ---- build_start=$(date +%s.%N) set +e - timeout "$TIMEOUT_SECONDS" make HIP_ARCH="$HIPCC_ARCH" "${make_extra[@]}" &>"$log" + timeout "$TIMEOUT_SECONDS" make "${MAKE_ARGS[@]}" &>"$log" rc=$? set -e build_elapsed=$(awk "BEGIN {printf \"%.3f\", $(date +%s.%N) - $build_start}") @@ -316,40 +173,30 @@ bench_one() { log_error "Failed (build): $name (exit $rc)" echo "$name" >> "$LOG_DIR/failed.log" fi - clear_special_env "$name" cd "$SCRIPT_DIR" return fi # ---- Phase 2: RUN (no make overhead) ---- - # Collect the run commands: either from direct_run_argsets overrides - # or by asking the Makefile what `make run` would execute. + # Ask the Makefile what `make run` would execute; the binary is already + # built, so `make -n run` only prints run commands. Join backslash- + # continuation lines before splitting into commands. local -a runcmds=() + local accum="" while IFS= read -r line; do - [[ -n "$line" ]] && runcmds+=("$line") - done < <(direct_run_argsets "$name" "$HIPCC_ARCH") - - if [[ ${#runcmds[@]} -eq 0 ]]; then - # Extract commands from the Makefile's run recipe via dry-run. - # Binary is already built, so make -n run only prints run commands. - # Join backslash-continuation lines before splitting into commands. - local accum="" - while IFS= read -r line; do - if [[ "$line" == *'\' ]]; then - accum+="${line%\\} " - else - accum+="$line" - [[ -n "$accum" ]] && runcmds+=("$accum") - accum="" - fi - done < <(make -n run HIP_ARCH="$HIPCC_ARCH" "${make_extra[@]}" 2>/dev/null) - [[ -n "$accum" ]] && runcmds+=("$accum") - fi + if [[ "$line" == *'\' ]]; then + accum+="${line%\\} " + else + accum+="$line" + [[ -n "$accum" ]] && runcmds+=("$accum") + accum="" + fi + done < <(make -n run "${MAKE_ARGS[@]}" 2>/dev/null) + [[ -n "$accum" ]] && runcmds+=("$accum") if [[ ${#runcmds[@]} -eq 0 ]]; then log_error "Failed: $name (no run commands found)" echo "$name" >> "$LOG_DIR/failed.log" - clear_special_env "$name" cd "$SCRIPT_DIR" return fi @@ -368,8 +215,7 @@ bench_one() { if [[ $rc -eq 0 ]]; then local time_detail="build ${build_elapsed}s, run ${run_elapsed}s" - local validation_issues - if validation_issues=$(check_output_validation "$log" "$name"); then + if check_output_validation "$log" >/dev/null; then log_success "$name (${time_detail})" echo "$name" >> "$LOG_DIR/success.log" else @@ -385,7 +231,6 @@ bench_one() { echo "$name" >> "$LOG_DIR/failed.log" fi - clear_special_env "$name" cd "$SCRIPT_DIR" } @@ -393,26 +238,17 @@ bench_one() { # DISCOVER & ITERATE # ============================================================================== -if [[ -n "$BENCHMARK_FILTER" ]]; then - # Build dirs directly from the filter/preset list — no need to scan all dirs - declare -a dirs=() - IFS=',' read -ra wanted <<< "$BENCHMARK_FILTER" - for w in "${wanted[@]}"; do - w="${w## }"; w="${w%% }" - [[ -z "$w" ]] && continue - local_dir="$HECBENCH_SRC/$w" - [[ -d "$local_dir" ]] && dirs+=("$local_dir") - done - if [[ ${#dirs[@]} -eq 0 ]]; then - log_error "Filter matched no benchmarks: $BENCHMARK_FILTER" - exit 1 - fi -else - mapfile -t dirs < <(find "$HECBENCH_SRC" -maxdepth 1 -type d -name "*-hip" | sort) - if [[ ${#dirs[@]} -eq 0 ]]; then - log_error "No *-hip directories under $HECBENCH_SRC" - exit 1 - fi +declare -a dirs=() +IFS=',' read -ra wanted <<< "$BENCHMARK_FILTER" +for w in "${wanted[@]}"; do + w="${w## }"; w="${w%% }" + [[ -z "$w" ]] && continue + local_dir="$HECBENCH_SRC/$w" + [[ -d "$local_dir" ]] && dirs+=("$local_dir") +done +if [[ ${#dirs[@]} -eq 0 ]]; then + log_error "Filter matched no benchmarks: $BENCHMARK_FILTER" + exit 1 fi log_info "Processing ${#dirs[@]} benchmarks..." @@ -439,7 +275,6 @@ log_info " Success: $(count success.log)" log_info " Suspect: $(count suspect.log) (exit 0 but output has failures)" log_info " Failed: $(count failed.log)" log_info " Timeout: $(count timeout.log)" -log_info " Skipped: $(count skipped.log)" log_info "" log_info "Logs: $LOG_DIR" log_info "Timings: $LOG_DIR/timings.csv" diff --git a/.github/hecbench/lib/common.sh b/.github/hecbench/lib/common.sh index 32ddd0ad9..87c8d6828 100755 --- a/.github/hecbench/lib/common.sh +++ b/.github/hecbench/lib/common.sh @@ -1,6 +1,5 @@ #!/usr/bin/env bash -# Common library for production-ready HeCBench scripts -# Provides shared logging, validation, and utility functions +# Shared logging, validation, and utility functions for the HeCBench CI scripts. # ============================================================================== # LOGGING FUNCTIONS @@ -44,32 +43,32 @@ log_debug() { } # ============================================================================== -# ENVIRONMENT VARIABLE VALIDATION +# PATH VALIDATION & DETECTION # ============================================================================== -# Require ROCM_PATH to be explicitly set +# Validate ROCm installation is complete. +validate_rocm_path() { + local path="$1" + [[ -d "$path" ]] || return 1 + [[ -x "$path/bin/hipcc" ]] || return 1 + [[ -d "$path/lib" ]] || return 1 + return 0 +} + +# Require ROCM_PATH to be explicitly set and valid. require_rocm_path() { if [[ -z "${ROCM_PATH:-}" ]]; then log_error "ROCM_PATH environment variable is not set" - echo "" - echo "Please set ROCM_PATH before running this script:" - echo " export ROCM_PATH=/path/to/TheRock-dist" - echo " # or" - echo " export ROCM_PATH=/opt/rocm" - echo "" - echo "Then run: $SCRIPT_NAME" + log_error " export ROCM_PATH=/path/to/ROCm" return 1 fi - # Strip trailing slashes to prevent double slashes in paths + # Strip trailing slashes to prevent double slashes in paths. ROCM_PATH="${ROCM_PATH%/}" - # Validate the path if ! validate_rocm_path "$ROCM_PATH"; then log_error "ROCM_PATH is set but invalid: $ROCM_PATH" - log_error "ROCm installation must contain:" - log_error " - bin/hipcc (HIP compiler)" - log_error " - lib/ (runtime libraries)" + log_error "ROCm installation must contain bin/hipcc and lib/" return 1 fi @@ -77,109 +76,19 @@ require_rocm_path() { return 0 } -# ============================================================================== -# VALIDATION FUNCTIONS -# ============================================================================== - -# Check if a command exists -check_command() { - local cmd="$1" - if ! command -v "$cmd" &>/dev/null; then - log_error "Required command not found: $cmd" - return 1 - fi - log_debug "Found command: $cmd" - return 0 -} - -# Validate directory exists and is readable -validate_dir() { - local dir="$1" - local description="${2:-Directory}" - - if [[ ! -d "$dir" ]]; then - log_error "$description not found: $dir" - return 1 - fi - - if [[ ! -r "$dir" ]]; then - log_error "$description not readable: $dir" - return 1 - fi - - log_debug "Validated directory: $dir" - return 0 -} - -# Validate file exists and is readable -validate_file() { - local file="$1" - local description="${2:-File}" - - if [[ ! -f "$file" ]]; then - log_error "$description not found: $file" - return 1 - fi - - if [[ ! -r "$file" ]]; then - log_error "$description not readable: $file" - return 1 - fi - - log_debug "Validated file: $file" - return 0 -} - -# ============================================================================== -# PATH DETECTION FUNCTIONS -# ============================================================================== - -# Find ROCm installation -find_rocm() { - # Try environment variable first - if [[ -n "${ROCM_PATH:-}" ]]; then - if validate_rocm_path "$ROCM_PATH"; then - echo "$ROCM_PATH" - return 0 - else - log_warn "ROCM_PATH is set but invalid: $ROCM_PATH" - fi - fi - - # Try common installation locations - local common_paths=( - "/opt/rocm" - "$HOME/rocm" - "/usr/local/rocm" - ) - - for path in "${common_paths[@]}"; do - if validate_rocm_path "$path"; then - log_debug "Found ROCm at: $path" - echo "$path" - return 0 - fi - done - - return 1 -} - -# Validate ROCm installation is complete -validate_rocm_path() { +# Validate HeCBench source directory (has *-hip benchmark dirs). +validate_hecbench_src() { local path="$1" - [[ -d "$path" ]] || return 1 - [[ -x "$path/bin/hipcc" ]] || return 1 - [[ -d "$path/lib" ]] || return 1 - - return 0 + local hip_dirs + hip_dirs=$(find "$path" -maxdepth 2 -name "*-hip" -type d 2>/dev/null | head -1) + [[ -n "$hip_dirs" ]] } -# Find HeCBench source directory +# Find HeCBench source directory: $HECBENCH_SRC, else relative to the script dir. find_hecbench_src() { local script_dir="$1" - # Try environment variable first if [[ -n "${HECBENCH_SRC:-}" ]]; then if validate_hecbench_src "$HECBENCH_SRC"; then echo "$HECBENCH_SRC" @@ -189,7 +98,6 @@ find_hecbench_src() { fi fi - # Try relative to script directory local relative_paths=( "$script_dir/../HeCBench/src" "$script_dir/../../HeCBench/src" @@ -209,253 +117,6 @@ find_hecbench_src() { return 1 } -# Validate HeCBench source directory -validate_hecbench_src() { - local path="$1" - - [[ -d "$path" ]] || return 1 - - # Check for characteristic HeCBench structure (HIP directories) - # Use a subshell to avoid pipefail issues with grep -q - local hip_dirs - hip_dirs=$(find "$path" -maxdepth 2 -name "*-hip" -type d 2>/dev/null | head -1) - [[ -n "$hip_dirs" ]] -} - -# Find MPI installation -find_mpi() { - # Try environment variable first - if [[ -n "${MPI_PATH:-}" ]]; then - if validate_mpi_path "$MPI_PATH"; then - echo "$MPI_PATH" - return 0 - else - log_warn "MPI_PATH is set but invalid: $MPI_PATH" - fi - fi - - # Try common installation locations - local common_paths=( - "/usr/lib/x86_64-linux-gnu/openmpi" - "/opt/openmpi" - "/usr/local/openmpi" - "/opt/mpi" - ) - - for path in "${common_paths[@]}"; do - if validate_mpi_path "$path"; then - log_debug "Found MPI at: $path" - echo "$path" - return 0 - fi - done - - return 1 -} - -# Validate MPI installation is complete -validate_mpi_path() { - local path="$1" - - [[ -d "$path" ]] || return 1 - [[ -d "$path/include" ]] || return 1 - [[ -d "$path/lib" ]] || return 1 - - # Check for mpi.h header - [[ -f "$path/include/mpi.h" ]] || return 1 - - return 0 -} - -# ============================================================================== -# UTILITY FUNCTIONS -# ============================================================================== - -# Execute command with dry-run support -execute_command() { - local description="$1" - shift - - if [[ "${DRY_RUN:-false}" == "true" ]]; then - log_info "[DRY-RUN] $description: $*" - return 0 - fi - - log_debug "Executing: $*" - "$@" -} - -# Retry a command with exponential backoff -retry_command() { - local max_attempts="$1" - local initial_delay="$2" - shift 2 - - local attempt=1 - local delay="$initial_delay" - - while [[ $attempt -le $max_attempts ]]; do - if "$@"; then - return 0 - fi - - if [[ $attempt -lt $max_attempts ]]; then - log_warn "Command failed (attempt $attempt/$max_attempts), retrying in ${delay}s..." - sleep "$delay" - delay=$((delay * 2)) - fi - - ((attempt++)) - done - - log_error "Command failed after $max_attempts attempts: $*" - return 1 -} - -# Show error and help message for missing path -show_path_error() { - local path_name="$1" - local env_var="$2" - local description="$3" - shift 3 - local required_contents=("$@") - - cat >&2 <&2 - done - - cat >&2 <&2 <&2 </dev/null; then - show_usage - exit 0 - else - log_error "show_usage function not defined" - exit 1 - fi - ;; - --dry-run) - DRY_RUN=true - log_info "Dry-run mode enabled" - shift - ;; - --verbose|-v) - VERBOSE=true - log_debug "Verbose mode enabled" - shift - ;; - *) - # Return remaining args for script-specific parsing - echo "$@" - return 0 - ;; - esac - done -} - -# ============================================================================== -# INTERACTIVE PROMPTS -# ============================================================================== - -# Prompt user for confirmation (yes/no) -# If INTERACTIVE=false, defaults to yes without prompting -# Usage: confirm_prompt "Do you want to continue?" && echo "Confirmed" -confirm_prompt() { - local prompt="$1" - local default="${2:-yes}" # yes or no - - # If not interactive, return based on default - if [[ "${INTERACTIVE:-true}" == "false" ]]; then - log_debug "Non-interactive mode: defaulting to '$default' for: $prompt" - [[ "$default" == "yes" ]] - return $? - fi - - # Interactive mode: ask user - local response - if [[ "$default" == "yes" ]]; then - echo -n "${BLUE}[?]${NC} $prompt [Y/n]: " - else - echo -n "${BLUE}[?]${NC} $prompt [y/N]: " - fi - read -r response < /dev/tty 2>/dev/null || response="" - - # Handle response - case "${response,,}" in - y|yes) - return 0 - ;; - n|no) - return 1 - ;; - "") - # Use default - [[ "$default" == "yes" ]] - return $? - ;; - *) - log_warn "Invalid response: '$response' (expected y/n)" - confirm_prompt "$prompt" "$default" # Retry - return $? - ;; - esac -} - # ============================================================================== # SHARED-MEMORY (/dev/shm) HOUSEKEEPING # ============================================================================== @@ -467,8 +128,7 @@ confirm_prompt() { # "not enough space for /dev/shm/...". # # Call this between benchmarks so every MPI/RCCL benchmark sees a clean slate. -# Scoped to the current user to avoid disturbing other tenants on shared -# systems. Silent on success; logs at debug level only. +# Scoped to the current user to avoid disturbing other tenants on shared systems. cleanup_stale_shm() { [[ -d /dev/shm ]] || return 0 local user="${USER:-$(id -un)}" @@ -479,33 +139,6 @@ cleanup_stale_shm() { -delete 2>/dev/null || true } -# ============================================================================== -# BINARY VALIDATION -# ============================================================================== - -# Validate that a compiled binary actually contains the expected offload arch. -# Catches the catastrophic case where Makefiles ignore HIP_ARCH and hipcc falls -# back to the system default (e.g., building gfx90a when amdgcnspirv was requested). -# -# Usage: validate_arch -# Returns 0 if the binary contains the expected arch bundle, 1 otherwise. -validate_arch() { - local binary="$1" arch="$2" - if [[ ! -f "$binary" ]]; then - log_error "validate_arch: binary not found: $binary" - return 1 - fi - case "$arch" in - amdgcnspirv) - grep -qa 'hip-spirv64-amd-amdhsa--amdgcnspirv' "$binary" ;; - gfx*) - grep -qa "hip-amdgcn-amd-amdhsa--${arch}" "$binary" ;; - *) - log_warn "validate_arch: unknown arch '$arch', skipping validation" - return 0 ;; - esac -} - # ============================================================================== # BENCHMARK OUTPUT VALIDATION # ============================================================================== @@ -516,12 +149,10 @@ validate_arch() { # Writes the first few suspect lines to stdout for logging. check_output_validation() { local log="$1" - local name="${2:-}" [[ ! -s "$log" ]] && return 0 local hits - hits=$( grep -i -E \ -e '\bFAIL(ED)?\b' \ @@ -566,11 +197,3 @@ check_output_validation() { return 0 } - -# ============================================================================== -# LIBRARY INITIALIZATION -# ============================================================================== - -# This library is now loaded -readonly COMMON_LIB_LOADED=true -log_debug "Common library loaded successfully" diff --git a/.github/hecbench/lib/presets.sh b/.github/hecbench/lib/presets.sh deleted file mode 100644 index b764bc211..000000000 --- a/.github/hecbench/lib/presets.sh +++ /dev/null @@ -1,11 +0,0 @@ -# Preset benchmark lists based on gfx90a timings from 2026-07-15. -# Each is a cumulative set: quick ⊂ standard ⊂ extended ⊂ full. -# quick ~10min (179 benchmarks) -# standard ~30min (318 benchmarks) -# extended ~60min (417 benchmarks) - -PRESET_QUICK="accuracy-hip,ace-hip,adam-hip,addBiasQKV-hip,aes-hip,affine-hip,aligned-types-hip,amgmk-hip,ans-hip,aobench-hip,asta-hip,atan2-hip,atomicReduction-hip,babelstream-hip,backprop-hip,base64e-hip,bfs-hip,bicgstab-hip,binomial-hip,blockAccess-hip,blockexchange-hip,bscan-hip,bsw-hip,b+tree-hip,cc-hip,ccsd-trpdrv-hip,ced-hip,cfd-hip,chemv-hip,clenergy-hip,clock-hip,cmp-hip,cobahh-hip,collision-hip,colorwheel-hip,complex-hip,concurrentKernels-hip,conversion-hip,cooling-hip,cross-hip,d2q9-bgk-hip,damage-hip,daphne-hip,dct8x8-hip,debayer-hip,determinant-hip,dispatch-hip,distort-hip,divergence-hip,dropout-hip,dxtc2-hip,easyWave-hip,ecdh-hip,egs-hip,ert-hip,f16atomic-hip,face-hip,fft-hip,fhd-hip,floydwarshall-hip,fpc-hip,fpdc-hip,fwt-hip,gc-hip,gels-hip,geodesic-hip,gibbs-hip,graphExecution-hip,groupnorm-hip,haccmk-hip,haversine-hip,heartwall-hip,heat2d-hip,heat-hip,henry-hip,histogram-hip,hogbom-hip,hotspot-hip,hungarian-hip,interleave-hip,intrinsics-cast-hip,inversek2j-hip,is-hip,ising-hip,iso2dfd-hip,jenkins-hash-hip,lanczos-hip,langford-hip,laplace3d-hip,lavaMD-hip,layout-hip,lda-hip,libor-hip,local-ht-hip,log2-hip,logan-hip,logprob-hip,lombscargle-hip,lr-hip,lud-hip,mallocFree-hip,mandelbrot-hip,marchingCubes-hip,md5hash-hip,meanshift-hip,medianfilter-hip,minibude-hip,minimap2-hip,minimod-hip,miniWeather-hip,mis-hip,mixbench-hip,mpc-hip,mrc-hip,mrg32k3a-hip,mr-hip,murmurhash3-hip,myocyte-hip,ne-hip,nlll-hip,nms-hip,nn-hip,opticalFlow-hip,overlap-hip,pad-hip,particlefilter-hip,particles-hip,pathfinder-hip,permute-hip,perplexity-hip,phmm-hip,pitch-hip,projectile-hip,qtclustering-hip,quant3MatMul-hip,quantAQLM-hip,radixsort-hip,recursiveGaussian-hip,resnet-kernels-hip,reverse2D-hip,reverse-hip,ring-hip,rmsnorm-hip,rng-wallace-hip,rodrigues-hip,romberg-hip,rowwiseMoments-hip,rsc-hip,rtm8-hip,sad-hip,scan2-hip,scatterThrust-hip,sc-hip,seam-carving-hip,simplemoc-hip,sobel-hip,sph-hip,split-hip,spm-hip,sptrsv-hip,srad-hip,stencil1d-hip,streamcluster-hip,su3-hip,svd3x3-hip,tensorT-hip,threadfence-hip,tqs-hip,triad-hip,tsa-hip,unfold-hip,urng-hip,vmc-hip,warpexchange-hip,warpsort-hip,winograd-hip,xlqc-hip,zeropoint-hip,zmddft-hip" - -PRESET_STANDARD="accuracy-hip,ace-hip,adam-hip,adamw-hip,addBiasQKV-hip,aes-hip,affine-hip,aidw-hip,aligned-types-hip,allreduce-hip,amgmk-hip,ans-hip,aobench-hip,aop-hip,asmooth-hip,asta-hip,atan2-hip,atomicAggregate-hip,atomicPerf-hip,atomicReduction-hip,atomicSystemWide-hip,attentionMergeState-hip,attentionMultiHead-hip,axhelm-hip,babelstream-hip,backprop-hip,base64e-hip,bfs-hip,bgmv-hip,bh-hip,bicgstab-hip,bincount-hip,binomial-hip,bitpacking-hip,bitpermute-hip,black-scholes-hip,blas-fp8gemm-hip,blas-gemm-hip,blockAccess-hip,blockexchange-hip,bm3d-hip,bmf-hip,boxfilter-hip,bscan-hip,bsearch-hip,bspline-vgh-hip,bsw-hip,b+tree-hip,cbsfil-hip,cc-hip,ccsd-trpdrv-hip,ccs-hip,ced-hip,cfd-hip,chacha20-hip,che-hip,chemv-hip,clenergy-hip,clock-hip,cmembench-hip,cmp-hip,cobahh-hip,collision-hip,colorwheel-hip,complex-hip,compute-score-hip,concurrentKernels-hip,contract-hip,conversion-hip,convolutionSeparable-hip,cooling-hip,coordinates-hip,crc64-hip,cross-hip,crs-hip,d2q9-bgk-hip,damage-hip,daphne-hip,dct8x8-hip,ddbp-hip,debayer-hip,determinant-hip,dispatch-hip,distort-hip,divergence-hip,dpid-hip,dropout-hip,dslash-hip,dxtc2-hip,easyWave-hip,ecdh-hip,egs-hip,eigenvalue-hip,eikonal-hip,entropy-hip,ert-hip,expdist-hip,extend2-hip,extrema-hip,f16atomic-hip,f8cast-hip,face-hip,fdtd3d-hip,fft-hip,fhd-hip,flame-hip,floydwarshall2-hip,floydwarshall-hip,fma-hip,fpc-hip,fpdc-hip,fresnel-hip,fwt-hip,gabor-hip,ga-hip,gaussian-hip,gc-hip,gels-hip,gelu-hip,geodesic-hip,ge-spmm-hip,gerbil-hip,gibbs-hip,gmm-hip,gpp-hip,graphB+-hip,graphExecution-hip,groupnorm-hip,gru2-hip,haccmk-hip,halo-finder-hip,haversine-hip,heartwall-hip,heat2d-hip,heat-hip,hellinger-hip,henry-hip,hexciton-hip,histogram-hip,hmm-hip,hogbom-hip,hotspot-hip,hungarian-hip,hwt1d-hip,idivide-hip,interleave-hip,intrinsics-cast-hip,inversek2j-hip,is-hip,ising-hip,iso2dfd-hip,jacobi-hip,jenkins-hash-hip,kalman-hip,kiss-hip,kmc-hip,kurtosis-hip,lanczos-hip,langford-hip,laplace3d-hip,lavaMD-hip,layout-hip,lci-hip,lda-hip,ldpc-hip,leukocyte-hip,libor-hip,local-ht-hip,log2-hip,logan-hip,logprob-hip,lombscargle-hip,loopback-hip,lr-hip,ludb-hip,lud-hip,lulesh-hip,lzss-hip,mallocFree-hip,mandelbrot-hip,marchingCubes-hip,mask-hip,matrix-rotate-hip,maxpool3d-hip,mcpr-hip,md5hash-hip,mdh-hip,meanshift-hip,medianfilter-hip,memcpy-hip,mergeVS-hip,mf-sgd-hip,minibude-hip,miniFE-hip,minimap2-hip,minimod-hip,miniWeather-hip,minmax-hip,mis-hip,mixbench-hip,moe-hip,morphology-hip,mpc-hip,mrc-hip,mrg32k3a-hip,mr-hip,mriQ-hip,mt-hip,multinomial-hip,multimaterial-hip,murmurhash3-hip,myocyte-hip,nbnxm-hip,nbody-hip,ne-hip,nlll-hip,nms-hip,nn-hip,norm2-hip,ntt-hip,opticalFlow-hip,overlap-hip,overlay-hip,p2p-hip,p4-hip,pad-hip,particlefilter-hip,particles-hip,pathfinder-hip,permute-hip,perplexity-hip,phmm-hip,pitch-hip,pns-hip,popcount-hip,present-hip,projectile-hip,pso-hip,qkv-hip,qtclustering-hip,quant3MatMul-hip,quantAQLM-hip,quantBnB-hip,quantVLLM-hip,radixsort-hip,rainflow-hip,reaction-hip,recursiveGaussian-hip,resnet-kernels-hip,reverse2D-hip,reverse-hip,ring-hip,rle-hip,rmsnorm-hip,rng-wallace-hip,rodrigues-hip,romberg-hip,rowwiseMoments-hip,rsbench-hip,rsc-hip,rsmt-hip,rtm8-hip,rushlarsen-hip,s3d-hip,sad-hip,scan2-hip,scan3-hip,scatterThrust-hip,scel-hip,sc-hip,seam-carving-hip,secp256k1-hip,segment-reduce-hip,segsort-hip,sheath-hip,shmembench-hip,simplemoc-hip,simpleMultiDevice-hip,slit-hip,sobel-hip,sobol-hip,softmax-hip,sph-hip,split-hip,spm-hip,spmm-hip,spmv-hip,sptrsv-hip,srad-hip,ss-hip,sss-hip,stencil1d-hip,streamcluster-hip,su3-hip,svd3x3-hip,tensorAccessor-hip,tensorT-hip,threadfence-hip,tissue-hip,tonemapping-hip,tqs-hip,triad-hip,tsa-hip,tsp-hip,unfold-hip,urng-hip,vadd-hip,vanGenuchten-hip,vmc-hip,vote-hip,warpexchange-hip,warpsort-hip,winograd-hip,wlcpow-hip,wordcount-hip,wsm5-hip,wyllie-hip,xlqc-hip,zeropoint-hip,zmddft-hip,zoom-hip" - -PRESET_EXTENDED="accuracy-hip,ace-hip,adam-hip,adamw-hip,addBiasQKV-hip,adjacent-hip,adv-hip,aes-hip,affine-hip,aidw-hip,aligned-types-hip,all-pairs-distance-hip,allreduce-hip,amgmk-hip,ans-hip,aobench-hip,aop-hip,asmooth-hip,asta-hip,atan2-hip,atomicAggregate-hip,atomicCAS-hip,atomicCost-hip,atomicPerf-hip,atomicReduction-hip,atomicSystemWide-hip,attentionMergeState-hip,attentionMultiHead-hip,axhelm-hip,axpby-hip,babelstream-hip,background-subtract-hip,backprop-hip,base64e-hip,bezier-surface-hip,bfs-hip,bgmv-hip,bh-hip,bicgstab-hip,bincount-hip,binomial-hip,bitonic-sort-hip,bitpacking-hip,bitpermute-hip,black-scholes-hip,blas-dot-hip,blas-fp8gemm-hip,blas-gemm-hip,blas-mxfp8gemm-hip,blockAccess-hip,blockexchange-hip,bm3d-hip,bmf-hip,bonds-hip,boxfilter-hip,bscan-hip,bsearch-hip,bspline-vgh-hip,bsw-hip,b+tree-hip,burger-hip,bwt-hip,car-hip,cbsfil-hip,cc-hip,ccl-hip,ccsd-trpdrv-hip,ccs-hip,ced-hip,cfd-hip,chacha20-hip,che-hip,chemv-hip,clenergy-hip,clink-hip,clock-hip,cmembench-hip,cmp-hip,cobahh-hip,collision-hip,colorwheel-hip,complex-hip,compute-score-hip,concat-hip,concurrentKernels-hip,contract-hip,conversion-hip,convolutionSeparable-hip,cooling-hip,coordinates-hip,crc64-hip,crossEntropy-hip,cross-hip,crs-hip,d2q9-bgk-hip,d3q19-bgk-hip,damage-hip,daphne-hip,dct8x8-hip,ddbp-hip,debayer-hip,degrid-hip,depixel-hip,determinant-hip,dispatch-hip,distort-hip,divergence-hip,doh-hip,dpid-hip,dropout-hip,dslash-hip,dxtc2-hip,easyWave-hip,ecdh-hip,egs-hip,eigenvalue-hip,eikonal-hip,entropy-hip,ert-hip,expdist-hip,extend2-hip,extrema-hip,f16atomic-hip,f8cast-hip,face-hip,fdtd3d-hip,fft-hip,fhd-hip,filter-hip,flame-hip,flip-hip,floydwarshall2-hip,floydwarshall-hip,fluidSim-hip,fma-hip,fpc-hip,fpdc-hip,fresnel-hip,fwt-hip,gabor-hip,ga-hip,gamma-correction-hip,gaussian-hip,gc-hip,gd-hip,geam-hip,gels-hip,gelu-hip,gemv-hip,geodesic-hip,ge-spmm-hip,gerbil-hip,gibbs-hip,gmm-hip,goulash-hip,gpp-hip,graphB+-hip,graphExecution-hip,groupnorm-hip,gru2-hip,haccmk-hip,halo-finder-hip,hausdorff-hip,haversine-hip,hbc-hip,heartwall-hip,heat2d-hip,heat-hip,hellinger-hip,henry-hip,hexciton-hip,histogram-hip,hmm-hip,hogbom-hip,hotspot3D-hip,hotspot-hip,hungarian-hip,hwt1d-hip,hybridsort-hip,hypterm-hip,idivide-hip,interleave-hip,intrinsics-cast-hip,inversek2j-hip,is-hip,ising-hip,iso2dfd-hip,jacobi-hip,jenkins-hash-hip,kalman-hip,keccaktreehash-hip,keogh-hip,kernelLaunch-hip,kiss-hip,kmc-hip,knn-hip,kurtosis-hip,lanczos-hip,langevin-hip,langford-hip,laplace3d-hip,lavaMD-hip,layout-hip,lci-hip,lda-hip,ldpc-hip,lebesgue-hip,leukocyte-hip,libor-hip,lid-driven-cavity-hip,lif-hip,local-ht-hip,log2-hip,logan-hip,logic-resim-hip,logic-rewrite-hip,logprob-hip,lombscargle-hip,loopback-hip,lr-hip,ludb-hip,lud-hip,lulesh-hip,lzss-hip,mallocFree-hip,mandelbrot-hip,marchingCubes-hip,mask-hip,matrix-rotate-hip,matrixT-hip,maxpool3d-hip,mcpr-hip,md5hash-hip,mdh-hip,md-hip,meanshift-hip,medianfilter-hip,memcpy-hip,merge-hip,mergeVS-hip,merkle-hip,mf-sgd-hip,michalewicz-hip,minibude-hip,miniFE-hip,minimap2-hip,minimod-hip,miniWeather-hip,minkowski-hip,minmax-hip,mis-hip,mixbench-hip,mmcsf-hip,moe-hip,morphology-hip,mpc-hip,mrc-hip,mrg32k3a-hip,mr-hip,mriQ-hip,mtf-hip,multimaterial-hip,mt-hip,multinomial-hip,murmurhash3-hip,myocyte-hip,nbnxm-hip,nbody-hip,ne-hip,nlll-hip,nms-hip,nn-hip,norm2-hip,nosync-hip,ntt-hip,nw-hip,opticalFlow-hip,overlap-hip,overlay-hip,p2p-hip,p4-hip,pad-hip,page-rank-hip,particle-diffusion-hip,particlefilter-hip,particles-hip,pathfinder-hip,pcc-hip,perlin-hip,permute-hip,perplexity-hip,phmm-hip,pitch-hip,pns-hip,pointerchase-hip,pointwise-hip,pool-hip,popcount-hip,present-hip,projectile-hip,pso-hip,qem-hip,qkv-hip,qrg-hip,qtclustering-hip,quant3MatMul-hip,quantAQLM-hip,quantBnB-hip,quantVLLM-hip,quicksort-hip,radixsort-hip,rainflow-hip,randomAccess-hip,reaction-hip,recursiveGaussian-hip,relu-hip,reshapeKVCache-hip,resize-hip,resnet-kernels-hip,reverse2D-hip,reverse-hip,rfs-hip,ring-hip,rle-hip,rmsnorm-hip,rng-wallace-hip,rodrigues-hip,romberg-hip,rowwiseMoments-hip,rsbench-hip,rsc-hip,rsmt-hip,rtm8-hip,rushlarsen-hip,s3d-hip,sad-hip,sa-hip,sampling-hip,scan2-hip,scan3-hip,scatter-hip,scatterThrust-hip,scel-hip,sc-hip,score-hip,seam-carving-hip,secp256k1-hip,segment-reduce-hip,segsort-hip,sheath-hip,shmembench-hip,shuffle-hip,si-hip,simplemoc-hip,simpleMultiDevice-hip,simpleSpmv-hip,slit-hip,snake-hip,snicit-hip,sobel-hip,sobol-hip,softmax-hip,sort-hip,sparkler-hip,spgemm-hip,sph-hip,split-hip,spm-hip,spmm-hip,spmv-hip,sptrsv-hip,srad-hip,ss-hip,ssim-hip,sss-hip,sssp-hip,stddev-hip,stencil1d-hip,stencil3d-hip,streamcluster-hip,streamCreateCopyDestroy-hip,streamOrderedAllocation-hip,streamPriority-hip,streamUM-hip,su3-hip,surfel-hip,svd3x3-hip,tensorAccessor-hip,tensorT-hip,tgvnn-hip,thomas-hip,threadfence-hip,tissue-hip,tonemapping-hip,tpacf-hip,tqs-hip,triad-hip,tsa-hip,tsp-hip,twell-hip,unfold-hip,upsample-hip,urng-hip,vadd-hip,vanGenuchten-hip,vmc-hip,vol2col-hip,vote-hip,warpexchange-hip,warpsort-hip,wedford-hip,winograd-hip,wlcpow-hip,word2vec-hip,wordcount-hip,wsm5-hip,wyllie-hip,xlqc-hip,zeropoint-hip,zmddft-hip,zoom-hip" diff --git a/.github/hecbench/modify_makefiles.sh b/.github/hecbench/modify_makefiles.sh index 5b287ff07..f7a023a99 100755 --- a/.github/hecbench/modify_makefiles.sh +++ b/.github/hecbench/modify_makefiles.sh @@ -1,28 +1,15 @@ #!/usr/bin/env bash -# Script: modify_makefiles.sh -# Purpose: Modify HeCBench Makefiles to add HIP architecture support -# Usage: ./modify_makefiles.sh [OPTIONS] -# -# Replaces "hipcc" with "hipcc --offload-arch=$(HIP_ARCH) --rocm-device-lib-path=..." -# This enables multi-architecture builds for AMD GPUs. +# Rewrite HeCBench HIP Makefiles so bench.sh can build them against the staged +# ROCm/LLVM toolchain. Replaces bare "hipcc" with +# $(HIPCC_BIN_DIR)/hipcc --offload-arch=$(HIP_ARCH) --rocm-device-lib-path=... +# and injects a $(EXTRA_CFLAGS) hook so per-benchmark flags can be passed via +# `make EXTRA_CFLAGS=...`. # # Uses HIP_ARCH (not ARCH) to avoid collisions with benchmarks that use ARCH for # their own purposes (e.g. dp4a-hip uses ARCH = CDNA as a feature flag). - set -euo pipefail -# ============================================================================== -# CONSTANTS & DEFAULTS -# ============================================================================== - SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -SCRIPT_NAME="$(basename "$0")" -DRY_RUN=false -VERBOSE=false - -# ============================================================================== -# SOURCE COMMON LIBRARY -# ============================================================================== # shellcheck source=lib/common.sh source "${SCRIPT_DIR}/lib/common.sh" || { @@ -30,21 +17,12 @@ source "${SCRIPT_DIR}/lib/common.sh" || { exit 1 } -# ============================================================================== -# CONFIGURATION -# ============================================================================== - -# REQUIRE ROCM_PATH to be explicitly set (no auto-detection) -if ! require_rocm_path; then - exit 1 -fi +require_rocm_path || exit 1 -# Try to find HeCBench source (auto-detect is OK for this) -if ! HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR"); then - show_path_error "HeCBench" "HECBENCH_SRC" "HeCBench source directory" \ - "$SCRIPT_DIR/../HeCBench/src" "$SCRIPT_DIR/../../HeCBench/src" +HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR") || { + log_error "Cannot find HeCBench src directory" exit 1 -fi +} DEVICE_LIB_PATH="${ROCM_PATH}/amdgcn/bitcode" @@ -52,81 +30,7 @@ DEVICE_LIB_PATH="${ROCM_PATH}/amdgcn/bitcode" # FUNCTIONS # ============================================================================== -show_usage() { - cat < "${makefile}.tmp" && mv "${makefile}.tmp" "$makefile" elif ! grep -q "^HIPCC_BIN_DIR" "$makefile"; then - # Use printf to avoid shell expansion issues with sed { grep -B1000 "^HIP_ARCH.*=" "$makefile" | head -n -1 grep "^HIP_ARCH.*=" "$makefile" @@ -274,14 +106,14 @@ modify_makefile() { } > "${makefile}.tmp" && mv "${makefile}.tmp" "$makefile" fi - # Escape device lib path for sed + # Escape device lib path for sed. local escaped_path="${DEVICE_LIB_PATH//\//\\/}" - # Replace hipcc with $(HIPCC_BIN_DIR)/hipcc + flags - # Only replace standalone "hipcc" not already prefixed with path/variable - sed -i "s#\(^\|[^/})]\)\bhipcc\b#\1\$(HIPCC_BIN_DIR)/hipcc --offload-arch=\$(HIP_ARCH) --rocm-device-lib-path=${escaped_path} \$(EXTRA_HIPCCFLAGS)#g" "$makefile" + # Replace standalone "hipcc" (not already prefixed with a path/variable) with + # $(HIPCC_BIN_DIR)/hipcc + offload-arch + device-lib-path. + sed -i "s#\(^\|[^/})]\)\bhipcc\b#\1\$(HIPCC_BIN_DIR)/hipcc --offload-arch=\$(HIP_ARCH) --rocm-device-lib-path=${escaped_path}#g" "$makefile" - # After the rewrite above, $(CC) / $(CXX) / $(HIPCC) typically expand to a + # After the rewrite, $(CC) / $(CXX) / $(HIPCC) typically expand to a # multi-word string. Any nested-make recipe that passes `=$(CC)` # unquoted would word-split, corrupting the inner build. Re-quote such # assignments. Restricted to recipe lines (start with TAB) so top-level @@ -290,20 +122,14 @@ modify_makefile() { TAB=$'\t' sed -i -E "s#^(${TAB}.*)([A-Z][A-Z_]+=)\\\$\\((CC|CXX|HIPCC)\\)#\\1\\2\"\\\$(\\3)\"#g" "$makefile" - # Remove -save-temps flag + # Remove -save-temps flag (disk clutter). sed -i 's/-save-temps//g' "$makefile" - # Fix HIP_PATH variable name conflict - local hip_path_renamed=false + # Rename HIP_PATH to avoid clobbering the environment variable of the same name. if grep -q "^HIP_PATH\s*=" "$makefile"; then sed -i -e 's/^\(HIP_PATH\s*=\)/HIP_SRC_PATH =/' \ -e 's/\$(\(HIP_PATH\))/$(HIP_SRC_PATH)/g' \ -e 's/\${\(HIP_PATH\)}/\${HIP_SRC_PATH}/g' "$makefile" - hip_path_renamed=true - fi - - # Log success with details - if [[ "$hip_path_renamed" == "true" ]]; then log_success "Modified: $rel_path (renamed HIP_PATH → HIP_SRC_PATH)" else log_success "Modified: $rel_path" @@ -314,8 +140,6 @@ modify_makefile() { process_makefiles() { local -a makefiles - - # Get array of Makefiles mapfile -t makefiles < <(find_makefiles) if [[ ${#makefiles[@]} -eq 0 ]]; then return 1 @@ -324,71 +148,23 @@ process_makefiles() { log_info "Processing ${#makefiles[@]} Makefiles..." echo "" - local modified_count=0 - local skipped_already=0 - local skipped_no_hipcc=0 - local failed_count=0 - + local modified=0 no_hipcc=0 failed=0 for makefile in "${makefiles[@]}"; do local rel_path="${makefile#$HECBENCH_SRC/}" - modify_makefile "$makefile" - local result=$? - - case $result in - 0) - modified_count=$((modified_count + 1)) - ;; - 2) - log_info "Already modified: $rel_path" - skipped_already=$((skipped_already + 1)) - ;; - 3) - log_info "No hipcc usage: $rel_path" - skipped_no_hipcc=$((skipped_no_hipcc + 1)) - ;; - *) - log_error "Failed to modify: $rel_path" - failed_count=$((failed_count + 1)) - ;; + case $? in + 0) modified=$((modified + 1)) ;; + 3) no_hipcc=$((no_hipcc + 1)) ;; + *) log_error "Failed to modify: $rel_path"; failed=$((failed + 1)) ;; esac done echo "" log_info "========================================" - log_info "Modification Summary" - log_info "========================================" - log_info "Total Makefiles found: ${#makefiles[@]}" - log_success "Modified: $modified_count" - log_info "Already modified: $skipped_already" - log_info "No hipcc usage: $skipped_no_hipcc" - - if [[ $failed_count -gt 0 ]]; then - log_error "Failed: $failed_count" - fi - - # Show sample verification - if [[ $modified_count -gt 0 ]] && [[ "$DRY_RUN" == "false" ]]; then - echo "" - log_info "Sample verification:" - for makefile in "${makefiles[@]}"; do - if compgen -G "${makefile}.bak.*" >/dev/null; then - local backup - backup=$(ls -t "${makefile}.bak."* | head -1) - local rel_path="${makefile#$HECBENCH_SRC/}" - echo " File: $rel_path" - echo " Before: $(grep -m1 'hipcc' "$backup" || echo "(none)")" - echo " After: $(grep -m1 'hipcc' "$makefile" || echo "(none)")" - echo "" - break - fi - done - fi - - if [[ $failed_count -gt 0 ]]; then - return 2 - fi - + log_info "Total Makefiles: ${#makefiles[@]}" + log_success "Modified: $modified" + log_info "No hipcc usage: $no_hipcc" + [[ $failed -gt 0 ]] && { log_error "Failed: $failed"; return 2; } return 0 } @@ -396,43 +172,16 @@ process_makefiles() { # MAIN # ============================================================================== -main() { - # Parse arguments - local remaining_args - remaining_args=$(parse_common_flags "$SCRIPT_NAME" "$@") +log_info "========================================" +log_info "HeCBench Makefile Modifier" +log_info "========================================" +log_info "ROCm Path: $ROCM_PATH" +log_info "HeCBench Source: $HECBENCH_SRC" +log_info "Device Library: $DEVICE_LIB_PATH" +[[ -d "$DEVICE_LIB_PATH" ]] || log_warn "Device library path not found: $DEVICE_LIB_PATH (builds may fail)" +echo "" - # No script-specific arguments expected - if [[ -n "$remaining_args" ]]; then - log_error "Unknown arguments: $remaining_args" - show_usage - exit 1 - fi - - log_info "========================================" - log_info "HeCBench Makefile Modifier" - log_info "========================================" - log_info "ROCm Path: $ROCM_PATH" - log_info "HeCBench Source: $HECBENCH_SRC" - log_info "Device Library: $DEVICE_LIB_PATH" - if [[ "$DRY_RUN" == "true" ]]; then - log_warn "DRY-RUN MODE: No files will be modified" - fi - echo "" - - # Validate environment - if ! validate_environment; then - exit 1 - fi - echo "" - - # Process Makefiles - if ! process_makefiles; then - exit 2 - fi - - echo "" - log_success "Makefile modification complete!" - echo "" -} +process_makefiles || exit $? -main "$@" +echo "" +log_success "Makefile modification complete!" diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 1ce149a2f..f5dd65866 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -756,7 +756,7 @@ jobs: echo "ROCM_PATH=$ROCM_PATH"; ls -la "$ROCM_PATH/bin" | grep -iE 'hip|clang' || true "$SCRIPTS/modify_makefiles.sh" FILTER=$(grep -vE '^[[:space:]]*(#|$)' "$SCRIPTS/ci_benchmarks.txt" | paste -sd,) - "$SCRIPTS/bench.sh" --gpu-id 0 --timeout 300 --filter "$FILTER" amdgcnspirv-be + "$SCRIPTS/bench.sh" --gpu-id 0 --timeout 300 --filter "$FILTER" - name: Gate on HeCBench results if: always() From c37e3e83a4c3b33030ee4987e5d0f50dd3a67edf Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Fri, 28 Aug 2026 06:01:57 -0500 Subject: [PATCH 11/19] Use single source of truth for log dir. --- .github/hecbench/bench.sh | 7 +++++-- .github/workflows/spirv-ci-linux.yml | 25 ++++++++++++++----------- 2 files changed, 19 insertions(+), 13 deletions(-) diff --git a/.github/hecbench/bench.sh b/.github/hecbench/bench.sh index 4d1e79be6..cd8550d4d 100755 --- a/.github/hecbench/bench.sh +++ b/.github/hecbench/bench.sh @@ -41,7 +41,7 @@ Environment Variables: HIP_CLANG_PATH clang bin dir for hipcc to drive (default: \$ROCM_PATH/bin) Output: - bench_logs_YYYYMMDD_HHMMSS_amdgcnspirv-be/ + $HECBENCH_LOG_DIR (CI), else bench_logs_YYYYMMDD_HHMMSS_amdgcnspirv-be/ timings.csv benchmark,build_seconds,run_seconds success.log built and ran cleanly failed.log non-zero exit (build or run) @@ -81,7 +81,10 @@ HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR") || { exit 1 } -LOG_DIR="${SCRIPT_DIR}/bench_logs_$(date +%Y%m%d_%H%M%S)_${ARCH}" +# In CI the workflow owns the log dir name (HECBENCH_LOG_DIR) so the gate and +# upload steps reference the same value — change it in one place. Outside CI it's +# unset, so fall back to a local timestamped dir next to this script. +LOG_DIR="${HECBENCH_LOG_DIR:-${SCRIPT_DIR}/bench_logs_$(date +%Y%m%d_%H%M%S)_${ARCH}}" mkdir -p "$LOG_DIR" export PATH="${ROCM_PATH}/bin:${PATH}" diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index f5dd65866..828a6cb87 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -700,6 +700,12 @@ jobs: options: | --device=/dev/kfd --device=/dev/dri --group-add video + # Single source of truth for the HeCBench log dir: bench.sh writes here, and + # the gate + upload steps read the same value. Tagged with the run id (+ + # attempt) so the uploaded logs map back to this exact workflow run. + env: + HECBENCH_LOG_DIR: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench/bench_logs_${{ github.run_id }}-${{ github.run_attempt }}_amdgcnspirv-be + steps: # Check out this repo at the PR commit for the vendored HeCBench scripts # in .github/hecbench. Into a subdir so it doesn't collide with the build @@ -760,28 +766,25 @@ jobs: - name: Gate on HeCBench results if: always() - env: - SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench run: | - LOG_DIR=$(ls -dt "$SCRIPTS"/bench_logs_*_amdgcnspirv-be 2>/dev/null | head -1) - if [[ -z "$LOG_DIR" ]]; then + if [[ ! -d "$HECBENCH_LOG_DIR" ]]; then echo "::error::No HeCBench log dir produced (run step failed early)." exit 1 fi - echo "Log dir: $LOG_DIR" + echo "Log dir: $HECBENCH_LOG_DIR" for cat in success suspect failed timeout skipped; do - n=$( [[ -f "$LOG_DIR/$cat.log" ]] && wc -l < "$LOG_DIR/$cat.log" || echo 0 ) + n=$( [[ -f "$HECBENCH_LOG_DIR/$cat.log" ]] && wc -l < "$HECBENCH_LOG_DIR/$cat.log" || echo 0 ) echo " $cat: $n" done rc=0 for cat in failed timeout suspect; do - if [[ -s "$LOG_DIR/$cat.log" ]]; then - echo "::group::$cat"; cat "$LOG_DIR/$cat.log"; echo "::endgroup::" - echo "::error::HeCBench $cat: $(wc -l < "$LOG_DIR/$cat.log") benchmark(s)." + if [[ -s "$HECBENCH_LOG_DIR/$cat.log" ]]; then + echo "::group::$cat"; cat "$HECBENCH_LOG_DIR/$cat.log"; echo "::endgroup::" + echo "::error::HeCBench $cat: $(wc -l < "$HECBENCH_LOG_DIR/$cat.log") benchmark(s)." rc=1 fi done - if [[ ! -s "$LOG_DIR/success.log" ]]; then + if [[ ! -s "$HECBENCH_LOG_DIR/success.log" ]]; then echo "::error::No HeCBench benchmark succeeded — toolchain/setup problem." rc=1 fi @@ -792,5 +795,5 @@ jobs: uses: actions/upload-artifact@bbbca2ddaa5d8feaa63e36b76fdaad77386f024f # v7.0.0 with: name: hecbench-logs - path: SPIRV-LLVM-Translator/.github/hecbench/bench_logs_*_amdgcnspirv-be/ + path: ${{ env.HECBENCH_LOG_DIR }}/ if-no-files-found: ignore From b89c9fb65494897b2a3abf90c58d836afca2fbf1 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Fri, 28 Aug 2026 06:40:29 -0500 Subject: [PATCH 12/19] Use single source of truth for SCRIPTS dir. --- .github/workflows/spirv-ci-linux.yml | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 828a6cb87..89dd7cec1 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -700,10 +700,13 @@ jobs: options: | --device=/dev/kfd --device=/dev/dri --group-add video - # Single source of truth for the HeCBench log dir: bench.sh writes here, and - # the gate + upload steps read the same value. Tagged with the run id (+ - # attempt) so the uploaded logs map back to this exact workflow run. + # Job-wide constants so each is defined once. SCRIPTS is the vendored driver + # dir; HECBENCH_LOG_DIR is the single source of truth for the log dir (bench.sh + # writes it, the gate + upload steps read it), tagged with the run id (+ + # attempt) so the uploaded logs map back to this exact workflow run. (An env + # entry can't reference another, so HECBENCH_LOG_DIR respells the SCRIPTS path.) env: + SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench HECBENCH_LOG_DIR: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench/bench_logs_${{ github.run_id }}-${{ github.run_attempt }}_amdgcnspirv-be steps: @@ -733,8 +736,6 @@ jobs: # Single sparse, shallow blob:none pass over the subset + precomputed deps # (ci_benchmarks.deps.txt from update_deps.sh): ~90 MB, not the full repo. - name: Checkout HeCBench (pinned, sparse subset) - env: - SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench run: | git init -q HeCBench cd HeCBench @@ -753,7 +754,6 @@ jobs: # bench.sh always exits 0 (results land in its log dir); next step gates. - name: Run HeCBench subset env: - SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench ROCM_PATH: ${{ github.workspace }}/${{ env.STAGING }} HIP_CLANG_PATH: ${{ github.workspace }}/${{ env.LLVM_BUILD }}/bin HECBENCH_SRC: ${{ github.workspace }}/HeCBench/src From f8b97267754d9f6ad46de87deb9fcd9cfecc88fb Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Fri, 28 Aug 2026 08:10:48 -0500 Subject: [PATCH 13/19] Trim comments. --- .github/hecbench/bench.sh | 20 +++++++------------- .github/hecbench/lib/common.sh | 13 ++++--------- .github/hecbench/modify_makefiles.sh | 20 ++++++-------------- .github/hecbench/update_deps.sh | 18 +++++------------- .github/workflows/spirv-ci-linux.yml | 16 ++++------------ 5 files changed, 26 insertions(+), 61 deletions(-) diff --git a/.github/hecbench/bench.sh b/.github/hecbench/bench.sh index cd8550d4d..4dc0b808a 100755 --- a/.github/hecbench/bench.sh +++ b/.github/hecbench/bench.sh @@ -81,9 +81,7 @@ HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR") || { exit 1 } -# In CI the workflow owns the log dir name (HECBENCH_LOG_DIR) so the gate and -# upload steps reference the same value — change it in one place. Outside CI it's -# unset, so fall back to a local timestamped dir next to this script. +# CI sets HECBENCH_LOG_DIR (single source of truth); local runs get a timestamp. LOG_DIR="${HECBENCH_LOG_DIR:-${SCRIPT_DIR}/bench_logs_$(date +%Y%m%d_%H%M%S)_${ARCH}}" mkdir -p "$LOG_DIR" @@ -103,9 +101,8 @@ export HIP_DEVICE_LIB_PATH="${ROCM_PATH}/amdgcn/bitcode" export ROCR_VISIBLE_DEVICES="$GPU_ID" ulimit -s unlimited 2>/dev/null || true -# libhipcxx hard-errors on amdgcnspirv because no compile-time __gfx*__ -# macro is defined for SPIR-V; allow it through and silence the warning. NDEBUG -# keeps host-side asserts out of timing paths. Passed to every benchmark build. +# libhipcxx hard-errors on amdgcnspirv (no compile-time __gfx*__ macro); +# the define lets it through. NDEBUG keeps host asserts out of timing paths. declare -a MAKE_ARGS=( HIP_ARCH="$HIPCC_ARCH" "EXTRA_CFLAGS=-D_LIBCUDACXX_ALLOW_UNSUPPORTED_ARCHITECTURE -DNDEBUG" @@ -127,10 +124,8 @@ log_info "Timeout: ${TIMEOUT_SECONDS}s" log_info "GPU ID: $GPU_ID" log_info "Logs: $LOG_DIR" -# Clear the COMGR JIT cache to prevent stale finalized kernels from being -# served after a compiler or runtime update. The cache keys are based on the -# SPIR-V input hash, so a runtime-only update (same compiler, same SPIR-V) -# would silently reuse the old (possibly buggy) native code. +# Clear the COMGR JIT cache: keys are the SPIR-V hash, so a runtime-only update +# (same SPIR-V) would otherwise serve stale native code. if [[ -d "${HOME}/.cache/comgr" ]]; then log_info "Clearing COMGR cache (${HOME}/.cache/comgr) ..." rm -rf "${HOME}/.cache/comgr" @@ -181,9 +176,8 @@ bench_one() { fi # ---- Phase 2: RUN (no make overhead) ---- - # Ask the Makefile what `make run` would execute; the binary is already - # built, so `make -n run` only prints run commands. Join backslash- - # continuation lines before splitting into commands. + # Ask the Makefile what `make run` runs (binary already built, so -n just + # prints); join backslash-continuation lines into whole commands. local -a runcmds=() local accum="" while IFS= read -r line; do diff --git a/.github/hecbench/lib/common.sh b/.github/hecbench/lib/common.sh index 87c8d6828..506b7b044 100755 --- a/.github/hecbench/lib/common.sh +++ b/.github/hecbench/lib/common.sh @@ -120,15 +120,10 @@ find_hecbench_src() { # ============================================================================== # SHARED-MEMORY (/dev/shm) HOUSEKEEPING # ============================================================================== -# OpenMPI's vader/sm BTL and RCCL leave per-rank backing files in /dev/shm -# (e.g. vader_segment...., nccl-XXXXXX). Normally these -# are unlinked when ranks exit cleanly, but a SIGKILL'd or timed-out run leaks -# them. On containers where /dev/shm is small (Docker default = 64 MiB), the -# leaks accumulate across runs until subsequent mpirun invocations fail with -# "not enough space for /dev/shm/...". -# -# Call this between benchmarks so every MPI/RCCL benchmark sees a clean slate. -# Scoped to the current user to avoid disturbing other tenants on shared systems. +# OpenMPI/RCCL leave per-rank backing files in /dev/shm; a SIGKILL'd or timed-out +# run leaks them, and on a small /dev/shm (Docker default 64 MiB) they accumulate +# until mpirun fails with "not enough space". Call between benchmarks. Scoped to +# the current user so other tenants on shared hosts are untouched. cleanup_stale_shm() { [[ -d /dev/shm ]] || return 0 local user="${USER:-$(id -un)}" diff --git a/.github/hecbench/modify_makefiles.sh b/.github/hecbench/modify_makefiles.sh index f7a023a99..8fa84d774 100755 --- a/.github/hecbench/modify_makefiles.sh +++ b/.github/hecbench/modify_makefiles.sh @@ -49,15 +49,9 @@ uses_hipcc() { grep -q "\bhipcc\b" "$1" } -# Inject $(EXTRA_CFLAGS) into the first matching assignment of each compile-flag -# variable that exists in the Makefile (CFLAGS, CXXFLAGS, HIPCC_FLAGS, NVCC_FLAGS). -# This lets bench.sh pass per-arch flags via `make EXTRA_CFLAGS=...` for the small -# subset of Makefiles whose CFLAGS line lacks the conventional `$(EXTRA_CFLAGS)` -# prefix that most HeCBench Makefiles already have. -# -# Idempotent: returns early if any $(EXTRA_CFLAGS)/${EXTRA_CFLAGS} reference is -# already present. Patches multiple flag vars per call so a Makefile that uses -# both CFLAGS and CXXFLAGS (e.g. halo-finder-hip) gets the hook in both rules. +# Prepend $(EXTRA_CFLAGS) to the first CFLAGS/CXXFLAGS/HIPCC_FLAGS/NVCC_FLAGS +# assignment so `make EXTRA_CFLAGS=...` reaches the few Makefiles lacking the +# conventional hook. Skips Makefiles that already have it. inject_extra_cflags() { local makefile="$1" if grep -qE '\$\(EXTRA_CFLAGS\)|\$\{EXTRA_CFLAGS\}' "$makefile"; then @@ -113,11 +107,9 @@ modify_makefile() { # $(HIPCC_BIN_DIR)/hipcc + offload-arch + device-lib-path. sed -i "s#\(^\|[^/})]\)\bhipcc\b#\1\$(HIPCC_BIN_DIR)/hipcc --offload-arch=\$(HIP_ARCH) --rocm-device-lib-path=${escaped_path}#g" "$makefile" - # After the rewrite, $(CC) / $(CXX) / $(HIPCC) typically expand to a - # multi-word string. Any nested-make recipe that passes `=$(CC)` - # unquoted would word-split, corrupting the inner build. Re-quote such - # assignments. Restricted to recipe lines (start with TAB) so top-level - # Makefile assignments are not altered. + # After the rewrite $(CC)/$(CXX)/$(HIPCC) expand to multi-word strings; a + # nested-make recipe passing `=$(CC)` unquoted would word-split. Re-quote, + # only on recipe lines (start with TAB) so top-level assignments are untouched. local TAB TAB=$'\t' sed -i -E "s#^(${TAB}.*)([A-Z][A-Z_]+=)\\\$\\((CC|CXX|HIPCC)\\)#\\1\\2\"\\\$(\\3)\"#g" "$makefile" diff --git a/.github/hecbench/update_deps.sh b/.github/hecbench/update_deps.sh index be65fd967..f38548b79 100755 --- a/.github/hecbench/update_deps.sh +++ b/.github/hecbench/update_deps.sh @@ -1,19 +1,11 @@ #!/usr/bin/env bash # Regenerate ci_benchmarks.deps.txt from ci_benchmarks.txt. # -# The benchmarks in ci_benchmarks.txt reference other in-tree HeCBench dirs via -# '../' (sibling *-cuda data dirs, shared include/, etc.). The test_hecbench -# CI job must check those out too, or the builds fail to resolve their inputs. -# This script resolves that set transitively at the pinned HECBENCH_REF and -# writes it to ci_benchmarks.deps.txt, so CI can do a single sparse checkout of -# ci_benchmarks.txt + ci_benchmarks.deps.txt instead of a discover-then-expand -# second pass. -# -# Re-run after editing ci_benchmarks.txt or bumping HECBENCH_REF, then commit the -# updated ci_benchmarks.deps.txt: -# .github/hecbench/update_deps.sh -# HECBENCH_REF defaults to the pin in ../workflows/spirv-ci-linux.yml; override -# via the env var. HECBENCH_REPO overrides the clone URL. +# Listed benchmarks reference sibling in-tree dirs via '../' (data dirs, shared +# include/, etc.); CI must check those out too or builds fail. This resolves that +# set transitively at HECBENCH_REF so CI does one sparse checkout, not two passes. +# Re-run after editing ci_benchmarks.txt or HECBENCH_REF, then commit the result. +# HECBENCH_REF defaults to the pin in the workflow; HECBENCH_REPO overrides the URL. set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 89dd7cec1..e24e15221 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -681,10 +681,8 @@ jobs: # ===================================================================== # Test - HeCBench (curated subset, amdgcnspirv-be backend) # ===================================================================== - # Builds + runs a curated HeCBench subset through the in-tree SPIRV backend - # (the default codegen path) and gates on each benchmark's self-verification. - # Driver scripts and subset list are vendored under .github/hecbench. Budgeted - # at ~15 min to stay off the critical path. + # Builds + runs a curated HeCBench subset through the in-tree SPIRV backend and + # gates on each benchmark's self-verification. Scripts vendored in .github/hecbench. test_hecbench: name: Test HeCBench needs: build @@ -700,19 +698,13 @@ jobs: options: | --device=/dev/kfd --device=/dev/dri --group-add video - # Job-wide constants so each is defined once. SCRIPTS is the vendored driver - # dir; HECBENCH_LOG_DIR is the single source of truth for the log dir (bench.sh - # writes it, the gate + upload steps read it), tagged with the run id (+ - # attempt) so the uploaded logs map back to this exact workflow run. (An env - # entry can't reference another, so HECBENCH_LOG_DIR respells the SCRIPTS path.) env: SCRIPTS: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench HECBENCH_LOG_DIR: ${{ github.workspace }}/SPIRV-LLVM-Translator/.github/hecbench/bench_logs_${{ github.run_id }}-${{ github.run_attempt }}_amdgcnspirv-be steps: - # Check out this repo at the PR commit for the vendored HeCBench scripts - # in .github/hecbench. Into a subdir so it doesn't collide with the build - # artifact, which untars build/ and staging/ at the workspace root. + # Check out this repo (PR head) for the vendored .github/hecbench scripts, + # into a subdir so it doesn't collide with the untarred build/ + staging/. - name: Checkout SPIRV-LLVM-Translator (PR head) uses: actions/checkout@9c091bb21b7c1c1d1991bb908d89e4e9dddfe3e0 # v7.0.0 with: From db1e7f6803abf32a1cd38932d072cd4bb8ea40cc Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Fri, 28 Aug 2026 12:42:37 -0500 Subject: [PATCH 14/19] Update HeCBench pin --- .github/hecbench/ci_benchmarks.deps.txt | 1 + .github/workflows/spirv-ci-linux.yml | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/.github/hecbench/ci_benchmarks.deps.txt b/.github/hecbench/ci_benchmarks.deps.txt index e1414909b..3388bb9a7 100644 --- a/.github/hecbench/ci_benchmarks.deps.txt +++ b/.github/hecbench/ci_benchmarks.deps.txt @@ -41,6 +41,7 @@ inversek2j-sycl is-cuda kiss-cuda laplace3d-cuda +ldpc-cuda log2-cuda lud-cuda marchingCubes-cuda diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index e24e15221..4d45c74fb 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -16,7 +16,7 @@ env: ROCM_EXAMPLES_REF: 33f07a01bd386a5213de5d8253df69b11dcf7bdd # Pinned ORNL/HeCBench revision for the test_hecbench job (not a moving branch, # so upstream churn can't fail unrelated PRs). Bump after re-validating the subset. - HECBENCH_REF: ab972efc2e3514bb7704c6685f2629f1638b8a50 + HECBENCH_REF: 10fb6ca670497bdaf3ba27780f4d8a0765989029 LLVM_BUILD: build DEVLIBS_BUILD: build-device-libs COMGR_BUILD: build-comgr From 9dfd81a09b9415f29d411c2f80d9de0d6a7dfcb3 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Fri, 28 Aug 2026 12:42:50 -0500 Subject: [PATCH 15/19] Further reduction --- .github/hecbench/modify_makefiles.sh | 78 +++------------------------- 1 file changed, 8 insertions(+), 70 deletions(-) diff --git a/.github/hecbench/modify_makefiles.sh b/.github/hecbench/modify_makefiles.sh index 8fa84d774..796ada847 100755 --- a/.github/hecbench/modify_makefiles.sh +++ b/.github/hecbench/modify_makefiles.sh @@ -1,12 +1,6 @@ #!/usr/bin/env bash -# Rewrite HeCBench HIP Makefiles so bench.sh can build them against the staged -# ROCm/LLVM toolchain. Replaces bare "hipcc" with -# $(HIPCC_BIN_DIR)/hipcc --offload-arch=$(HIP_ARCH) --rocm-device-lib-path=... -# and injects a $(EXTRA_CFLAGS) hook so per-benchmark flags can be passed via -# `make EXTRA_CFLAGS=...`. -# -# Uses HIP_ARCH (not ARCH) to avoid collisions with benchmarks that use ARCH for -# their own purposes (e.g. dp4a-hip uses ARCH = CDNA as a feature flag). +# Append --offload-arch=$(HIP_ARCH) to bare "hipcc" in HeCBench HIP Makefiles. +# HIP_ARCH (not ARCH) avoids clashing with benchmarks that use ARCH (e.g. dp4a-hip). set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" @@ -24,8 +18,6 @@ HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR") || { exit 1 } -DEVICE_LIB_PATH="${ROCM_PATH}/amdgcn/bitcode" - # ============================================================================== # FUNCTIONS # ============================================================================== @@ -49,84 +41,32 @@ uses_hipcc() { grep -q "\bhipcc\b" "$1" } -# Prepend $(EXTRA_CFLAGS) to the first CFLAGS/CXXFLAGS/HIPCC_FLAGS/NVCC_FLAGS -# assignment so `make EXTRA_CFLAGS=...` reaches the few Makefiles lacking the -# conventional hook. Skips Makefiles that already have it. -inject_extra_cflags() { - local makefile="$1" - if grep -qE '\$\(EXTRA_CFLAGS\)|\$\{EXTRA_CFLAGS\}' "$makefile"; then - return 0 - fi - local var - for var in CFLAGS CXXFLAGS HIPCC_FLAGS NVCC_FLAGS; do - if grep -qE "^[[:space:]]*${var}[[:space:]]*[:?+]?=" "$makefile"; then - sed -i -E "0,/^([[:space:]]*${var}[[:space:]]*[:?+]?=[[:space:]]*)/{s//\1\$(EXTRA_CFLAGS) /}" "$makefile" - fi - done -} - # Returns: 0 modified, 3 no hipcc usage. modify_makefile() { local makefile="$1" local rel_path="${makefile#$HECBENCH_SRC/}" - # Skip wrapper / pytorch Makefiles that don't compile with hipcc. + # Skip Makefiles that don't use hipcc (wrappers, pytorch). if ! uses_hipcc "$makefile"; then return 3 fi - # Inject the $(EXTRA_CFLAGS) hook so bench.sh can pass per-arch flags. - inject_extra_cflags "$makefile" - - # Add HIP_ARCH and HIPCC_BIN_DIR variables if not present. + # Provide a HIP_ARCH default unless the Makefile sets its own (e.g. dp-hip). if ! grep -q "^HIP_ARCH.*=" "$makefile"; then { echo "# User-configurable architecture (default: gfx906)" echo "HIP_ARCH ?= gfx906" echo "" - echo "# HIP compiler location" - echo "HIPCC_BIN_DIR ?= $ROCM_PATH/bin" - echo "" cat "$makefile" } > "${makefile}.tmp" && mv "${makefile}.tmp" "$makefile" - elif ! grep -q "^HIPCC_BIN_DIR" "$makefile"; then - { - grep -B1000 "^HIP_ARCH.*=" "$makefile" | head -n -1 - grep "^HIP_ARCH.*=" "$makefile" - echo "" - echo "# HIP compiler location" - echo "HIPCC_BIN_DIR ?= $ROCM_PATH/bin" - grep -A1000 "^HIP_ARCH.*=" "$makefile" | tail -n +2 - } > "${makefile}.tmp" && mv "${makefile}.tmp" "$makefile" fi - # Escape device lib path for sed. - local escaped_path="${DEVICE_LIB_PATH//\//\\/}" - - # Replace standalone "hipcc" (not already prefixed with a path/variable) with - # $(HIPCC_BIN_DIR)/hipcc + offload-arch + device-lib-path. - sed -i "s#\(^\|[^/})]\)\bhipcc\b#\1\$(HIPCC_BIN_DIR)/hipcc --offload-arch=\$(HIP_ARCH) --rocm-device-lib-path=${escaped_path}#g" "$makefile" - - # After the rewrite $(CC)/$(CXX)/$(HIPCC) expand to multi-word strings; a - # nested-make recipe passing `=$(CC)` unquoted would word-split. Re-quote, - # only on recipe lines (start with TAB) so top-level assignments are untouched. - local TAB - TAB=$'\t' - sed -i -E "s#^(${TAB}.*)([A-Z][A-Z_]+=)\\\$\\((CC|CXX|HIPCC)\\)#\\1\\2\"\\\$(\\3)\"#g" "$makefile" - - # Remove -save-temps flag (disk clutter). - sed -i 's/-save-temps//g' "$makefile" - - # Rename HIP_PATH to avoid clobbering the environment variable of the same name. - if grep -q "^HIP_PATH\s*=" "$makefile"; then - sed -i -e 's/^\(HIP_PATH\s*=\)/HIP_SRC_PATH =/' \ - -e 's/\$(\(HIP_PATH\))/$(HIP_SRC_PATH)/g' \ - -e 's/\${\(HIP_PATH\)}/\${HIP_SRC_PATH}/g' "$makefile" - log_success "Modified: $rel_path (renamed HIP_PATH → HIP_SRC_PATH)" - else - log_success "Modified: $rel_path" + # Append the flag to bare hipcc; the grep guard makes re-runs idempotent. + if ! grep -qF 'hipcc --offload-arch=$(HIP_ARCH)' "$makefile"; then + sed -i "s#\(^\|[^/})]\)\bhipcc\b#\1hipcc --offload-arch=\$(HIP_ARCH)#g" "$makefile" fi + log_success "Modified: $rel_path" return 0 } @@ -169,8 +109,6 @@ log_info "HeCBench Makefile Modifier" log_info "========================================" log_info "ROCm Path: $ROCM_PATH" log_info "HeCBench Source: $HECBENCH_SRC" -log_info "Device Library: $DEVICE_LIB_PATH" -[[ -d "$DEVICE_LIB_PATH" ]] || log_warn "Device library path not found: $DEVICE_LIB_PATH (builds may fail)" echo "" process_makefiles || exit $? From feed0c234c3cdd44bb383458a8467748903c2b0d Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Tue, 1 Sep 2026 09:40:01 -0500 Subject: [PATCH 16/19] Debug slow download step --- .github/workflows/spirv-ci-linux.yml | 29 ++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 32585cc2f..cd5d01455 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -714,6 +714,35 @@ jobs: fetch-depth: 1 persist-credentials: false + # TEMP DIAGNOSTIC: the "Download build tree artifact" step takes ~95 min on + # this gfx942 pool but <1 min on the aws/azure runners for the identical + # artifact. Localize the bottleneck (network path vs disk) before the real + # download, then hand off to the owning team if neither test explains it. + # Only two benign, non-disclosing probes: a GitHub-hosted single-stream + # download and a local disk write. Never fails the job (continue-on-error). + - name: Diagnose slow artifact download + continue-on-error: true + run: | + set +e + echo "===== Single-stream network throughput (GitHub-hosted file) =====" + # ~180 MB release asset on GitHub's own CDN -- same trust boundary the + # job already uses; isolates raw single-stream speed from the artifact + # action's parallelism. Written to /dev/null so disk is not involved. + URL="https://github.com/actions/runner/releases/download/v2.319.1/actions-runner-linux-x64-2.319.1.tar.gz" + if command -v curl >/dev/null 2>&1; then + curl -sSL -o /dev/null -w 'dl_bytes=%{size_download} time=%{time_total}s speed=%{speed_download} B/s\n' "$URL" + elif command -v wget >/dev/null 2>&1; then + wget -O /dev/null "$URL" + else + echo "no curl/wget available" + fi + + echo "===== Disk write throughput (2 GB into the artifact workdir) =====" + dd if=/dev/zero of=./_disktest.bin bs=1M count=2000 oflag=direct 2>&1 || \ + dd if=/dev/zero of=./_disktest.bin bs=1M count=2000 conv=fdatasync 2>&1 + rm -f ./_disktest.bin + echo "===== END DIAGNOSTIC =====" + - name: Download build tree artifact uses: actions/download-artifact@70fc10c6e5e1ce46ad2ea6f2b72d43f7d47b13c3 # v8.0.0 with: From cc38b8ec89623087f0292663d512810100bf7283 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Mon, 14 Sep 2026 09:39:24 -0500 Subject: [PATCH 17/19] Revert "Debug slow download step" This reverts commit feed0c234c3cdd44bb383458a8467748903c2b0d. --- .github/workflows/spirv-ci-linux.yml | 29 ---------------------------- 1 file changed, 29 deletions(-) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index cd5d01455..32585cc2f 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -714,35 +714,6 @@ jobs: fetch-depth: 1 persist-credentials: false - # TEMP DIAGNOSTIC: the "Download build tree artifact" step takes ~95 min on - # this gfx942 pool but <1 min on the aws/azure runners for the identical - # artifact. Localize the bottleneck (network path vs disk) before the real - # download, then hand off to the owning team if neither test explains it. - # Only two benign, non-disclosing probes: a GitHub-hosted single-stream - # download and a local disk write. Never fails the job (continue-on-error). - - name: Diagnose slow artifact download - continue-on-error: true - run: | - set +e - echo "===== Single-stream network throughput (GitHub-hosted file) =====" - # ~180 MB release asset on GitHub's own CDN -- same trust boundary the - # job already uses; isolates raw single-stream speed from the artifact - # action's parallelism. Written to /dev/null so disk is not involved. - URL="https://github.com/actions/runner/releases/download/v2.319.1/actions-runner-linux-x64-2.319.1.tar.gz" - if command -v curl >/dev/null 2>&1; then - curl -sSL -o /dev/null -w 'dl_bytes=%{size_download} time=%{time_total}s speed=%{speed_download} B/s\n' "$URL" - elif command -v wget >/dev/null 2>&1; then - wget -O /dev/null "$URL" - else - echo "no curl/wget available" - fi - - echo "===== Disk write throughput (2 GB into the artifact workdir) =====" - dd if=/dev/zero of=./_disktest.bin bs=1M count=2000 oflag=direct 2>&1 || \ - dd if=/dev/zero of=./_disktest.bin bs=1M count=2000 conv=fdatasync 2>&1 - rm -f ./_disktest.bin - echo "===== END DIAGNOSTIC =====" - - name: Download build tree artifact uses: actions/download-artifact@70fc10c6e5e1ce46ad2ea6f2b72d43f7d47b13c3 # v8.0.0 with: From 5b443e029a257d6da695531625bf86c9ab5ef606 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Wed, 16 Sep 2026 13:14:56 -0500 Subject: [PATCH 18/19] Bump HeCBench pin --- .github/workflows/spirv-ci-linux.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/spirv-ci-linux.yml b/.github/workflows/spirv-ci-linux.yml index 63eaaa0c8..471fe23fd 100644 --- a/.github/workflows/spirv-ci-linux.yml +++ b/.github/workflows/spirv-ci-linux.yml @@ -16,7 +16,7 @@ env: ROCM_EXAMPLES_REF: fa0d558d6f92107f54e242dd473f675bcf613e2a # Pinned ORNL/HeCBench revision for the test_hecbench job (not a moving branch, # so upstream churn can't fail unrelated PRs). Bump after re-validating the subset. - HECBENCH_REF: 10fb6ca670497bdaf3ba27780f4d8a0765989029 + HECBENCH_REF: f2bb6769225d89e6982e60d8f628cc1ae961d9c3 LLVM_BUILD: build DEVLIBS_BUILD: build-device-libs COMGR_BUILD: build-comgr From bf87888314c813dd2a733cb5d501b146ac652790 Mon Sep 17 00:00:00 2001 From: Marcos Maronas Date: Thu, 17 Sep 2026 02:50:29 -0500 Subject: [PATCH 19/19] Update benchmark deps. --- .github/hecbench/ci_benchmarks.deps.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/hecbench/ci_benchmarks.deps.txt b/.github/hecbench/ci_benchmarks.deps.txt index 3388bb9a7..5878de40b 100644 --- a/.github/hecbench/ci_benchmarks.deps.txt +++ b/.github/hecbench/ci_benchmarks.deps.txt @@ -70,6 +70,7 @@ romberg-cuda rowwiseMoments-cuda rsc-cuda sc-cuda +shmembench-cuda tqs-cuda tsa-cuda wyllie-cuda