From 341c3cf78c6284d88ad76944125478c050aa0c3f Mon Sep 17 00:00:00 2001 From: Ethan Luis McDonough Date: Wed, 28 Jun 2023 15:02:37 -0500 Subject: [PATCH] [flang][openmp] Fortran offloading test Flang currently supports offloading for AMD GPUs. This patch establishes a test structure for Fortran offloading tests in libomptarget. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D148778 --- openmp/README.rst | 6 +++ openmp/libomptarget/test/CMakeLists.txt | 6 ++- openmp/libomptarget/test/lit.cfg | 11 ++++++ openmp/libomptarget/test/lit.site.cfg.in | 5 +++ .../test/offloading/fortran/basic_array.c | 43 ++++++++++++++++++++++ .../test/offloading/fortran/basic_array.f90 | 6 +++ 6 files changed, 76 insertions(+), 1 deletion(-) create mode 100644 openmp/libomptarget/test/offloading/fortran/basic_array.c create mode 100644 openmp/libomptarget/test/offloading/fortran/basic_array.f90 diff --git a/openmp/README.rst b/openmp/README.rst index eb4bb04..f5d0cef 100644 --- a/openmp/README.rst +++ b/openmp/README.rst @@ -165,6 +165,12 @@ Options for ``libomp`` **LIBOMP_FORTRAN_MODULES** = ``OFF|ON`` Create the Fortran modules (requires Fortran compiler). + .. note:: + + If libomptarget is built in-tree with both flang and openmp in + `LLVM_ENABLE_PROJECTS`, flang will be used for Fortran offloading + tests. + macOS* Fat Libraries """""""""""""""""""" On macOS* machines, it is possible to build universal (or fat) libraries which diff --git a/openmp/libomptarget/test/CMakeLists.txt b/openmp/libomptarget/test/CMakeLists.txt index 5cc9d24..943dbd0 100644 --- a/openmp/libomptarget/test/CMakeLists.txt +++ b/openmp/libomptarget/test/CMakeLists.txt @@ -32,9 +32,13 @@ foreach(CURRENT_TARGET IN LISTS SYSTEM_TARGETS) configure_file(lit.site.cfg.in ${CURRENT_TARGET}/lit.site.cfg @ONLY) endforeach() +if ("flang" IN_LIST LLVM_ENABLE_PROJECTS) + SET(FORTRAN_TEST_DEPS flang-new) +endif() + add_openmp_testsuite(check-libomptarget "Running libomptarget tests" ${LIBOMPTARGET_LIT_TESTSUITES} EXCLUDE_FROM_CHECK_ALL - DEPENDS omptarget omp ${LIBOMPTARGET_TESTED_PLUGINS} + DEPENDS omptarget omp ${LIBOMPTARGET_TESTED_PLUGINS} ${FORTRAN_TEST_DEPS} ARGS ${LIBOMPTARGET_LIT_ARG_LIST}) diff --git a/openmp/libomptarget/test/lit.cfg b/openmp/libomptarget/test/lit.cfg index e3255b3..1fc01e2 100644 --- a/openmp/libomptarget/test/lit.cfg +++ b/openmp/libomptarget/test/lit.cfg @@ -4,6 +4,10 @@ import os import lit.formats +from lit.llvm import llvm_config +from lit.llvm.subst import ToolSubst +from lit.llvm.subst import FindTool + # Tell pylint that we know config and lit_config exist somewhere. if 'PYLINT_IMPORT' in os.environ: config = object() @@ -58,6 +62,8 @@ config.test_source_root = os.path.dirname(__file__) # test_exec_root: The root object directory where output is placed config.test_exec_root = config.libomptarget_obj_root +tools = [] + # test format config.test_format = lit.formats.ShTest() @@ -86,6 +92,10 @@ if config.has_libomptarget_ompt: config.available_features.add(config.libomptarget_current_target) +if 'flang' in config.llvm_enabled_projects: + config.available_features.add('flang') + tools.append(ToolSubst('%flang', command=FindTool('flang-new'), unresolved='fatal')) + # Determine whether the test system supports unified memory. # For CUDA, this is the case with compute capability 70 (Volta) or higher. # For all other targets, we currently assume it is. @@ -302,3 +312,4 @@ else: config.substitutions.append(("%cuda_flags", "")) config.substitutions.append(("%flags", config.test_flags)) config.substitutions.append(("%not", config.libomptarget_not)) +llvm_config.add_tool_substitutions(tools, config.bin_llvm_tools_dir) diff --git a/openmp/libomptarget/test/lit.site.cfg.in b/openmp/libomptarget/test/lit.site.cfg.in index e388a80..f4a63be 100644 --- a/openmp/libomptarget/test/lit.site.cfg.in +++ b/openmp/libomptarget/test/lit.site.cfg.in @@ -1,5 +1,6 @@ @AUTO_GEN_COMMENT@ +config.bin_llvm_tools_dir = "@CMAKE_BINARY_DIR@/bin" config.test_c_compiler = "@OPENMP_TEST_C_COMPILER@" config.test_cxx_compiler = "@OPENMP_TEST_CXX_COMPILER@" config.test_compiler_features = @OPENMP_TEST_COMPILER_FEATURES@ @@ -20,6 +21,10 @@ config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@" config.libomptarget_not = "@OPENMP_NOT_EXECUTABLE@" config.libomptarget_debug = @LIBOMPTARGET_DEBUG@ config.has_libomptarget_ompt = @LIBOMPTARGET_OMPT_SUPPORT@ +config.llvm_enabled_projects = "@LLVM_ENABLE_PROJECTS@".split(";") + +import lit.llvm +lit.llvm.initialize(lit_config, config) # Let the main config do the real work. lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg") diff --git a/openmp/libomptarget/test/offloading/fortran/basic_array.c b/openmp/libomptarget/test/offloading/fortran/basic_array.c new file mode 100644 index 0000000..20334e8 --- /dev/null +++ b/openmp/libomptarget/test/offloading/fortran/basic_array.c @@ -0,0 +1,43 @@ +// Basic offloading test for function compiled with flang +// REQUIRES: flang, amdgcn-amd-amdhsa + +// RUN: %flang -c -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ +// RUN: %S/basic_array.f90 -o basic_array.o +// RUN: %libomptarget-compile-generic basic_array.o +// RUN: %t | %fcheck-generic + +#include +#define TEST_ARR_LEN 10 + +#pragma omp declare target +void increment_at(int i, int *array); +#pragma omp end declare target + +void increment_array(int *b, int n) { +#pragma omp target map(tofrom : b [0:n]) + for (int i = 0; i < n; i++) { + increment_at(i, b); + } +} + +int main() { + int arr[TEST_ARR_LEN] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; + + increment_array(arr, TEST_ARR_LEN); + for (int i = 0; i < TEST_ARR_LEN; i++) { + printf("%d = %d\n", i, arr[i]); + } + + return 0; +} + +// CHECK: 0 = 1 +// CHECK-NEXT: 1 = 2 +// CHECK-NEXT: 2 = 3 +// CHECK-NEXT: 3 = 4 +// CHECK-NEXT: 4 = 5 +// CHECK-NEXT: 5 = 6 +// CHECK-NEXT: 6 = 7 +// CHECK-NEXT: 7 = 8 +// CHECK-NEXT: 8 = 9 +// CHECK-NEXT: 9 = 10 diff --git a/openmp/libomptarget/test/offloading/fortran/basic_array.f90 b/openmp/libomptarget/test/offloading/fortran/basic_array.f90 new file mode 100644 index 0000000..776316b --- /dev/null +++ b/openmp/libomptarget/test/offloading/fortran/basic_array.f90 @@ -0,0 +1,6 @@ +subroutine increment_at(c_index, arr) bind(C, name="increment_at") + use ISO_C_BINDING + integer (C_INT), dimension(*), intent(inout) :: arr + integer (C_INT), value :: c_index + arr(c_index+1) = arr(c_index+1) + 1 +end subroutine -- 2.7.4