diff --git a/.github/hecbench/bench.sh b/.github/hecbench/bench.sh new file mode 100755 index 000000000..4dc0b808a --- /dev/null +++ b/.github/hecbench/bench.sh @@ -0,0 +1,277 @@ +#!/usr/bin/env bash +# 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)" +SCRIPT_NAME="$(basename "$0")" + +# shellcheck source=lib/common.sh +source "${SCRIPT_DIR}/lib/common.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} +BENCHMARK_FILTER=${BENCHMARK_FILTER:-} + +show_usage() { + cat <.log per-benchmark combined build+run output +EOF +} + +# ============================================================================== +# ARG PARSING +# ============================================================================== + +while [[ $# -gt 0 ]]; do + case "$1" in + --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 + +if [[ -z "$BENCHMARK_FILTER" ]]; then + log_error "--filter is required" + show_usage + exit 1 +fi + +# ============================================================================== +# VALIDATION & SETUP +# ============================================================================== + +require_rocm_path || exit 1 +HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR") || { + log_error "Cannot find HeCBench src directory" + exit 1 +} + +# 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" + +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 +# 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" +ulimit -s unlimited 2>/dev/null || true + +# 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" +) + +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_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: 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" +fi + +echo "" + +# ============================================================================== +# PER-BENCHMARK +# ============================================================================== + +# 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" + + # 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; } + + # Clean to ensure a deterministic build state. + make clean &>/dev/null || true + + local build_start build_elapsed run_start run_elapsed rc + + # ---- Phase 1: BUILD (default target only) ---- + build_start=$(date +%s.%N) + set +e + timeout "$TIMEOUT_SECONDS" make "${MAKE_ARGS[@]}" &>"$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 + cd "$SCRIPT_DIR" + return + fi + + # ---- Phase 2: RUN (no make overhead) ---- + # 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 + 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" + 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" + if check_output_validation "$log" >/dev/null; 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 + + cd "$SCRIPT_DIR" +} + +# ============================================================================== +# DISCOVER & ITERATE +# ============================================================================== + +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..." +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 "" +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..5878de40b --- /dev/null +++ b/.github/hecbench/ci_benchmarks.deps.txt @@ -0,0 +1,78 @@ +# 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 +affine-cuda +atomicPerf-cuda +atomicReduction-cuda +backprop-cuda +base64e-cuda +blockAccess-cuda +boxfilter-sycl +ccsd-trpdrv-cuda +channelSum-cuda +che-cuda +cobahh-cuda +complex-cuda +damage-cuda +data +debayer-sycl +dpid-cuda +dxtc2-sycl +easyWave-omp +ecdh-cuda +extend2-sycl +extrema-cuda +fft-cuda +fpdc-cuda +fwt-cuda +gabor-cuda +ga-cuda +groupnorm-cuda +hellinger-cuda +idivide-cuda +inversek2j-sycl +is-cuda +kiss-cuda +laplace3d-cuda +ldpc-cuda +log2-cuda +lud-cuda +marchingCubes-cuda +medianfilter-cuda +medianfilter-sycl +mis-cuda +mrc-cuda +mr-cuda +multinomial-cuda +nlll-cuda +nms-cuda +opticalFlow-cuda +overlay-cuda +p4-cuda +pad-cuda +particlefilter-cuda +particles-cuda +perplexity-cuda +phmm-cuda +projectile-cuda +pso-cuda +reaction-cuda +recursiveGaussian-cuda +rmsnorm-cuda +romberg-cuda +rowwiseMoments-cuda +rsc-cuda +sc-cuda +shmembench-cuda +tqs-cuda +tsa-cuda +wyllie-cuda +zeropoint-cuda +zoom-cuda diff --git a/.github/hecbench/ci_benchmarks.txt b/.github/hecbench/ci_benchmarks.txt new file mode 100644 index 000000000..1f9e0d647 --- /dev/null +++ b/.github/hecbench/ci_benchmarks.txt @@ -0,0 +1,178 @@ +# 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. +# 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 +affine-hip +aligned-types-hip +amgmk-hip +ans-hip +aobench-hip +asta-hip +atan2-hip +atomicReduction-hip +babelstream-hip +backprop-hip +base64e-hip +binomial-hip +blockexchange-hip +bscan-hip +ccsd-trpdrv-hip +channelSum-hip +che-hip +chemv-hip +clenergy-hip +clock-hip +cobahh-hip +collision-hip +colorwheel-hip +complex-hip +concurrentKernels-hip +conversion-hip +cooling-hip +cross-hip +damage-hip +dct8x8-hip +debayer-hip +dispatch-hip +distort-hip +divergence-hip +dpid-hip +dropout-hip +dxtc2-hip +easyWave-hip +ecdh-hip +extend2-hip +fft-hip +floydwarshall-hip +fpc-hip +fpdc-hip +fwt-hip +gc-hip +graphExecution-hip +groupnorm-hip +heat-hip +histogram-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 +lombscargle-hip +lud-hip +mallocFree-hip +mandelbrot-hip +marchingCubes-hip +md5hash-hip +medianfilter-hip +mis-hip +mixbench-hip +mr-hip +mrc-hip +mrg32k3a-hip +murmurhash3-hip +myocyte-hip +ne-hip +nlll-hip +nms-hip +opticalFlow-hip +overlap-hip +pad-hip +particlefilter-hip +particles-hip +pathfinder-hip +perplexity-hip +phmm-hip +projectile-hip +pso-hip +qtclustering-hip +quant3MatMul-hip +quantAQLM-hip +radixsort-hip +reverse-hip +reverse2D-hip +rmsnorm-hip +rodrigues-hip +romberg-hip +rsc-hip +sc-hip +scan2-hip +split-hip +spm-hip +srad-hip +stencil1d-hip +su3-hip +threadfence-hip +tqs-hip +triad-hip +tsa-hip +unfold-hip +vmc-hip +warpexchange-hip +winograd-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 +# 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 +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/lib/common.sh b/.github/hecbench/lib/common.sh new file mode 100755 index 000000000..506b7b044 --- /dev/null +++ b/.github/hecbench/lib/common.sh @@ -0,0 +1,194 @@ +#!/usr/bin/env bash +# Shared logging, validation, and utility functions for the HeCBench CI scripts. + +# ============================================================================== +# 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 +} + +# ============================================================================== +# PATH VALIDATION & DETECTION +# ============================================================================== + +# 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" + log_error " export ROCM_PATH=/path/to/ROCm" + return 1 + fi + + # Strip trailing slashes to prevent double slashes in paths. + ROCM_PATH="${ROCM_PATH%/}" + + if ! validate_rocm_path "$ROCM_PATH"; then + log_error "ROCM_PATH is set but invalid: $ROCM_PATH" + log_error "ROCm installation must contain bin/hipcc and lib/" + return 1 + fi + + log_success "Using ROCM_PATH: $ROCM_PATH" + return 0 +} + +# Validate HeCBench source directory (has *-hip benchmark dirs). +validate_hecbench_src() { + local path="$1" + [[ -d "$path" ]] || return 1 + 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: $HECBENCH_SRC, else relative to the script dir. +find_hecbench_src() { + local script_dir="$1" + + 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 + + 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 +} + +# ============================================================================== +# SHARED-MEMORY (/dev/shm) HOUSEKEEPING +# ============================================================================== +# 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)}" + # -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 +} + +# ============================================================================== +# 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" + + [[ ! -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 +} diff --git a/.github/hecbench/modify_makefiles.sh b/.github/hecbench/modify_makefiles.sh new file mode 100755 index 000000000..796ada847 --- /dev/null +++ b/.github/hecbench/modify_makefiles.sh @@ -0,0 +1,117 @@ +#!/usr/bin/env bash +# 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)" + +# shellcheck source=lib/common.sh +source "${SCRIPT_DIR}/lib/common.sh" || { + echo "ERROR: Failed to load common library" >&2 + exit 1 +} + +require_rocm_path || exit 1 + +HECBENCH_SRC=$(find_hecbench_src "$SCRIPT_DIR") || { + log_error "Cannot find HeCBench src directory" + exit 1 +} + +# ============================================================================== +# FUNCTIONS +# ============================================================================== + +find_makefiles() { + local makefiles=() + while IFS= read -r -d '' makefile; do + makefiles+=("$makefile") + done < <(find "$HECBENCH_SRC" -path "*-hip/*" -name "Makefile" -type f ! -name "*.bak*" -print0 | sort -z) + + if [[ ${#makefiles[@]} -eq 0 ]]; then + log_error "No Makefiles found in $HECBENCH_SRC" + return 1 + fi + + printf '%s\n' "${makefiles[@]}" + return 0 +} + +uses_hipcc() { + grep -q "\bhipcc\b" "$1" +} + +# Returns: 0 modified, 3 no hipcc usage. +modify_makefile() { + local makefile="$1" + local rel_path="${makefile#$HECBENCH_SRC/}" + + # Skip Makefiles that don't use hipcc (wrappers, pytorch). + if ! uses_hipcc "$makefile"; then + return 3 + fi + + # 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 "" + cat "$makefile" + } > "${makefile}.tmp" && mv "${makefile}.tmp" "$makefile" + fi + + # 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 +} + +process_makefiles() { + local -a makefiles + mapfile -t makefiles < <(find_makefiles) + if [[ ${#makefiles[@]} -eq 0 ]]; then + return 1 + fi + + log_info "Processing ${#makefiles[@]} Makefiles..." + echo "" + + local modified=0 no_hipcc=0 failed=0 + for makefile in "${makefiles[@]}"; do + local rel_path="${makefile#$HECBENCH_SRC/}" + modify_makefile "$makefile" + 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 "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 +} + +# ============================================================================== +# MAIN +# ============================================================================== + +log_info "========================================" +log_info "HeCBench Makefile Modifier" +log_info "========================================" +log_info "ROCm Path: $ROCM_PATH" +log_info "HeCBench Source: $HECBENCH_SRC" +echo "" + +process_makefiles || exit $? + +echo "" +log_success "Makefile modification complete!" diff --git a/.github/hecbench/update_deps.sh b/.github/hecbench/update_deps.sh new file mode 100755 index 000000000..f38548b79 --- /dev/null +++ b/.github/hecbench/update_deps.sh @@ -0,0 +1,92 @@ +#!/usr/bin/env bash +# Regenerate ci_benchmarks.deps.txt from ci_benchmarks.txt. +# +# 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)" +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 839c8e381..471fe23fd 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: 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: f2bb6769225d89e6982e60d8f628cc1ae961d9c3 LLVM_BUILD: build DEVLIBS_BUILD: build-device-libs COMGR_BUILD: build-comgr @@ -189,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: | @@ -665,3 +677,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 and + # gates on each benchmark's self-verification. Scripts vendored in .github/hecbench. + test_hecbench: + name: Test HeCBench + needs: build + runs-on: linux-gfx942-1gpu-ccs-csp-ossci-rocm + timeout-minutes: 240 + # 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: | + --device=/dev/kfd --device=/dev/dri --group-add video + + 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 (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: + 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) + 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: + 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" + + - name: Gate on HeCBench results + if: always() + run: | + if [[ ! -d "$HECBENCH_LOG_DIR" ]]; then + echo "::error::No HeCBench log dir produced (run step failed early)." + exit 1 + fi + echo "Log dir: $HECBENCH_LOG_DIR" + for cat in success suspect failed timeout skipped; do + 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 "$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 "$HECBENCH_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: ${{ env.HECBENCH_LOG_DIR }}/ + if-no-files-found: ignore