merge aomp-dev into amd-staging

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/<bench>-{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

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"

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

test/smoke-fort-fails/flang-use-dev-addr-performance/bar.hip

@@ -0,0 +1,27 @@
+#include <hip/hip_runtime.h>
+#include <stdlib.h>
+
+__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" */

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
+

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
+