diff --git a/bin/run_HeCBench.sh b/bin/run_HeCBench.sh new file mode 100755 index 000000000..1c434bae3 --- /dev/null +++ b/bin/run_HeCBench.sh @@ -0,0 +1,169 @@ +#!/usr/bin/env bash + +# +#Copyright © Advanced Micro Devices, Inc., or its affiliates. +# +#SPDX-License-Identifier: MIT +# + +# run_HeCBench.sh - runs HeCBench benchmarks in the $AOMP_REPOS_TEST dir. +# +# Environment variables (set before running; none are required unless noted): +# +# Compiler / ROCm layout: +# AOMP LLVM compiler tree (clang++, libomp) +# if unset: /opt/rocm/lib/llvm +# if set: use as-is +# ROCM_PATH HIP/ROCm install root (hipcc, libamdhip64) +# if unset: realpath(AOMP/../..) (e.g. /opt/rocm) +# AOMP_GPU GPU arch for OpenMP builds (ARCH= in Makefile.aomp); +# auto-detected via rocm_agent_enumerator if unset +# +# Test tree (from aomp_common_vars): +# AOMP_REPOS_TEST parent of cloned HeCBench; default: $HOME/git/aomp-test +# expects: $AOMP_REPOS_TEST/HeCBench/src/-{omp,hip} +# +# Run control: +# PROGRAMMING_MODELS space-separated list of build variants to run; +# default: "openmp hip" (both) +# openmp - src/*-omp dirs, build with Makefile.aomp +# (clang++, ARCH=$AOMP_GPU) +# hip - src/*-hip dirs, build with Makefile (hipcc) +# examples: +# PROGRAMMING_MODELS=openmp +# PROGRAMMING_MODELS="openmp hip" +# PROGRAMMING_MODELS=hip +# HECBENCH_LIST space-separated benchmark dirs to run (default: all) +# HECBENCH_TIMEOUT per-benchmark timeout in seconds (default: 180) +# LAUNCHER passed to Makefile run target (e.g. "gpurun time -p") +# +# Compiler flags: +# EXTRA_CFLAGS extra compiler flags (Makefile.aomp / Makefile); not set +# by this script — export before running, e.g.: +# export EXTRA_CFLAGS='-fopenmp-target-fast' + +# --- Start standard header to set AOMP environment variables ---- +realpath=$(realpath "$0") +thisdir=$(dirname "$realpath") + +# shellcheck disable=SC1091 +. "$thisdir/aomp_common_vars" + +# If AOMP and ROCM_PATH are already set, use them. If not, use defaults. +# The default for AOMP is /opt/rocm/llvm. The default for ROCM_PATH is $AOMP/../.. +export AOMP="${AOMP:-/opt/rocm/lib/llvm}" +export ROCM_PATH="${ROCM_PATH:-$(realpath -m "${AOMP}/../..")}" +export PATH=$AOMP/bin:$ROCM_PATH/bin:$PATH +export LD_LIBRARY_PATH=$AOMP/lib:$ROCM_PATH/lib:$LD_LIBRARY_PATH + +PROGRAMMING_MODELS=${PROGRAMMING_MODELS:-"openmp hip"} +HECBENCH_TIMEOUT=${HECBENCH_TIMEOUT:-180} +HECBENCH_LIST=${HECBENCH_LIST:-""} +LAUNCHER=${LAUNCHER:-} + +hecbench_root=$AOMP_REPOS_TEST/HeCBench +hecbench_src=$hecbench_root/src + +check_hipcc_clang_mismatch() { + local hipcc_bin clang_bin hipcc_clang_line clang_ver + hipcc_bin=$(PATH="$AOMP/bin:$ROCM_PATH/bin:$PATH" command -v hipcc 2>/dev/null) + clang_bin=$(PATH="$AOMP/bin:$PATH" command -v clang 2>/dev/null) + if [ -z "$hipcc_bin" ] || [ -z "$clang_bin" ]; then + return 0 + fi + hipcc_clang_line=$("$hipcc_bin" --version 2>&1 | grep -i 'clang version') + clang_ver=$("$clang_bin" --version 2>&1 | grep -i 'clang version') + if [ -n "$hipcc_clang_line" ] && [ -n "$clang_ver" ]; then + if [ "$hipcc_clang_line" == "$clang_ver" ]; then + echo "INFO: hipcc and clang compiler versions match." >&2 + else + echo "WARNING: hipcc and clang report different compiler versions:" >&2 + echo " hipcc ($hipcc_bin):" >&2 + printf ' %s\n' "$hipcc_clang_line" >&2 + echo " clang ($clang_bin): $clang_ver" >&2 + fi + else + echo "WARNING: hipcc and clang compiler versions unverified." >&2 + fi +} + +# Use function to set and test AOMP_GPU +setaompgpu + +if [ ! -d "$hecbench_root" ]; then + echo "ERROR: HeCBench not found in $AOMP_REPOS_TEST." + exit 1 +elif [ ! -d "$hecbench_src" ]; then + echo "ERROR: HeCBench src not found: $hecbench_src" + exit 1 +fi + +cd "$hecbench_src" || exit 1 + +results=$hecbench_root/results.txt +rm -f "$results" + +# Check for a mismatch. +check_hipcc_clang_mismatch + +echo PROGRAMMING_MODELS: "$PROGRAMMING_MODELS" +for model in $PROGRAMMING_MODELS; do + if [ "$model" == "openmp" ]; then + suffix="-omp" + makefile="Makefile.aomp" + elif [ "$model" == "hip" ]; then + suffix="-hip" + makefile="Makefile" + else + echo "ERROR: Option not recognized: $model." + exit 1 + fi + + if [ -n "$HECBENCH_LIST" ]; then + dirs="$HECBENCH_LIST" + else + dirs=$(find . -maxdepth 1 -type d -name "*$suffix" | sort | sed 's|^\./||') + fi + + if [ -z "$dirs" ]; then + echo "WARNING: No benchmark dirs found for model=$model suffix=$suffix in $(pwd)" + continue + fi + + NumTestsRun=0 + NumTestsSkipped=0 + for d in $dirs; do + if [ ! -d "$d" ]; then + NumTestsSkipped=$((NumTestsSkipped + 1)) + continue + fi + if [ ! -f "$d/$makefile" ]; then + NumTestsSkipped=$((NumTestsSkipped + 1)) + continue + fi + NumTestsRun=$((NumTestsRun + 1)) + echo "=== [$model] $d ===" | tee -a "$results" + ( + cd "$d" || exit 1 + if [ "$model" == "openmp" ]; then + make_clean=(make -f "$makefile" "ARCH=$AOMP_GPU" clean) + make_run=(make -f "$makefile" "ARCH=$AOMP_GPU" "LAUNCHER=$LAUNCHER" run) + else + make_clean=(make -f "$makefile" clean) + make_run=(make -f "$makefile" "LAUNCHER=$LAUNCHER" run) + fi + "${make_clean[@]}" >/dev/null 2>&1 + if timeout "$HECBENCH_TIMEOUT" "${make_run[@]}" >>"$results" 2>&1; then + echo "STATUS $d: PASS" | tee -a "$results" + "${make_clean[@]}" >/dev/null 2>&1 + else + echo "STATUS $d: FAIL(rc=$?)" | tee -a "$results" + fi + ) + done + echo "[$model] NumTestsRun=$NumTestsRun NumTestsSkipped=$NumTestsSkipped" + echo >> "$results" + echo "=== SUMMARY [$model] ===" | tee -a "$results" + echo "NumTestsRun=$NumTestsRun" | tee -a "$results" + echo "NumTestsSkipped=$NumTestsSkipped" | tee -a "$results" +done diff --git a/srock-bin/srock_common_vars b/srock-bin/srock_common_vars index c55b400c5..82848353f 100644 --- a/srock-bin/srock_common_vars +++ b/srock-bin/srock_common_vars @@ -226,6 +226,7 @@ _cmake_args=(-B build -GNinja -DTHEROCK_AMDGPU_DIST_BUNDLE_NAME=srock -DTHEROCK_BACKGROUND_BUILD_JOBS=1 -DTHEROCK_BUILD_LLVM_TESTS=1 + -DTHEROCK_FLAG_LLVM_ENABLE_ASSERTIONS=ON "${_cmake_enable[@]}" $SROCK_CMAKE_EXTRA "$SROCK_THEROCK_DIR" diff --git a/test/smoke-fort-fails/flang-use-dev-addr-performance/Makefile b/test/smoke-fort-fails/flang-use-dev-addr-performance/Makefile new file mode 100644 index 000000000..c2150aaf2 --- /dev/null +++ b/test/smoke-fort-fails/flang-use-dev-addr-performance/Makefile @@ -0,0 +1,26 @@ +include ../../Makefile.defs + +TESTNAME = use-dev-addr-performance +TESTSRC_MAIN = main.f90 +TESTSRC_AUX = bar.o +TESTSRC_ALL = $(TESTSRC_MAIN) $(TESTSRC_AUX) +AOMPHIP ?= $(AOMP) +HIPCC ?= $(AOMPHIP)/bin/hipcc +HIP_CLANG_PATH ?= $(AOMP)/bin + +CFLAGS = -O3 +FLANG ?= flang +OMP_BIN = $(AOMP)/bin/$(FLANG) +CC = $(OMP_BIN) $(VERBOSE) +EXTRA_CFLAGS = -L$(AOMPHIP)/lib -lamdhip64 -Wl,-rpath,$(AOMPHIP)/lib -fPIC +#-ccc-print-phases +#"-\#\#\#" + +include ../Makefile.rules +all: $(TESTNAME) + +bar.o : bar.hip + HIP_CLANG_PATH=$(HIP_CLANG_PATH) $(HIPCC) -c --offload-arch=$(AOMP_GPU) -fPIC $^ -o $@ + +run: $(TESTNAME) + LIBOMPTARGET_INFO=8 ./$(TESTNAME) 2>&1 | $(AOMP)/bin/FileCheck check.txt diff --git a/test/smoke-fort-fails/flang-use-dev-addr-performance/bar.hip b/test/smoke-fort-fails/flang-use-dev-addr-performance/bar.hip new file mode 100644 index 000000000..5ca383c51 --- /dev/null +++ b/test/smoke-fort-fails/flang-use-dev-addr-performance/bar.hip @@ -0,0 +1,27 @@ +#include +#include + +__global__ void bar_kernel(int *x, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) { + x[i]++; + } +} + +extern "C" { + +void bar_GPU(int *x, int n) +{ + int num_threads = 256; + int num_blocks = n / num_threads + 1; + hipLaunchKernelGGL(bar_kernel, dim3(num_blocks), dim3(num_threads), 0, 0, x, n); + (void)hipDeviceSynchronize(); +} + +void print_ptr(int * x) { + static int cnt; + printf("pointer x %lx number of calls: %d\n",(ulong)x, ++cnt); +} +} /* extern "C" */ diff --git a/test/smoke-fort-fails/flang-use-dev-addr-performance/check.txt b/test/smoke-fort-fails/flang-use-dev-addr-performance/check.txt new file mode 100644 index 000000000..c2c05551f --- /dev/null +++ b/test/smoke-fort-fails/flang-use-dev-addr-performance/check.txt @@ -0,0 +1,13 @@ +; CHECK: Creating new map entry +; CHECK: Creating new map entry +; CHECK-NOT: Creating new map entry +; CHECK: Removing map entry +; CHECK: Removing map entry +; CHECK-NOT: Removing map entry +; CHECK: Success +; CHECK: pointer x [[ADDR:[0-9a-f]+]] number of calls: 1 +; CHECK: pointer x [[ADDR]] number of calls: 2 +; CHECK: pointer x [[GPU_ADDR:[0-9a-f]+]] number of calls: 3 +; CHECK: pointer x [[GPU_ADDR]] number of calls: 4 +; CHECK: pointer x [[ADDR]] number of calls: 5 + diff --git a/test/smoke-fort-fails/flang-use-dev-addr-performance/main.f90 b/test/smoke-fort-fails/flang-use-dev-addr-performance/main.f90 new file mode 100644 index 000000000..2ff3e3d00 --- /dev/null +++ b/test/smoke-fort-fails/flang-use-dev-addr-performance/main.f90 @@ -0,0 +1,65 @@ +MODULE foo + USE iso_c_binding + USE omp_lib + IMPLICIT NONE + PRIVATE + PUBLIC :: bar_device_addr, print_ptr, bar + + INTERFACE + SUBROUTINE bar(x, n) BIND(C, name="bar_GPU") + USE iso_c_binding + TYPE(C_PTR), VALUE, INTENT(IN) :: x + INTEGER(C_INT), VALUE, INTENT(IN) :: n + END SUBROUTINE + SUBROUTINE print_ptr(x) BIND(C, name="print_ptr") + USE iso_c_binding + TYPE(C_PTR), VALUE, INTENT(IN) :: x + END SUBROUTINE + + END INTERFACE + +CONTAINS + + SUBROUTINE bar_device_addr(x, n) + INTEGER, TARGET, INTENT(IN) :: x(:) + INTEGER(C_INT), INTENT(IN) :: n + !$omp target data use_device_addr (x) + CALL print_ptr(c_loc(x)) + CALL bar(c_loc(x), n) + !$omp end target data + END SUBROUTINE + +END MODULE foo + +PROGRAM test_ptr + USE iso_c_binding + USE, intrinsic :: iso_fortran_env, only: error_unit + USE omp_lib + USE foo + IMPLICIT NONE + + INTEGER, ALLOCATABLE, TARGET :: x(:) + INTEGER(C_INT) :: i, n + n = 1000 + ALLOCATE(x(n)) + x = 1 + CALL print_ptr(c_loc(x)) + !$omp target enter data map(to: x) + CALL print_ptr(c_loc(x)) + !$omp target data use_device_addr (x) + CALL print_ptr(c_loc(x)) + CALL bar(c_loc(x), n) + !$omp end target data + CALL bar_device_addr(x,n) + !$omp target exit data map(from: x) + CALL print_ptr(c_loc(x)) + DO i = 1,n + IF (x(i) .ne. 3) then + PRINT *, "Bad result for use_device_addr!" + STOP 1 + ENDIF + END DO + DEALLOCATE(x) + write(error_unit, *) 'Success' +END PROGRAM test_ptr +