[flang][openmp] Fortran offloading test

[elmcdonough]

Summary

Flang currently supports offloading for AMD GPUs. This patch establishes a test structure for Fortran offloading tests in libomptarget.

[elmcdonough]

Fix build issues with past flang existence check.

[jdoerfert]

@tianshilei1992 @jhuber6, works for me, anything jumps out on your end?

[vzakhari]

openmp/libomptarget/test/lit.cfg

97

I think you have to make sure that check-libomptarget CMake target has a dependency on whatever target produces flang-new tool. Alternatively, you may follow the same logic that is used for clang with config.test_*_compiler properties, but then it does not really matter if Flang is in LLVM_ENABLE_PROJECTS.

[elmcdonough]

openmp/libomptarget/test/lit.cfg

97

I wanted to make it so that the flang dependency is optional. The offloading tests should be run if flang is an enabled project, otherwise, they should be ignored. I had a config.test_fortran_compiler in my original build that inherited from CMAKE_Fortran_COMPILER in standalone builds and looked for flang-new in the binary dir with CMake's find_program for in-tree builds. Unfortunately, that version didn't work in certain circumstances.

The current revision seems to work better and is cleaner IMO, but I'm open to the idea of adding an extra build flag for fortran tests if others think it would work better. @jdoerfert Any thoughts on this?

[vzakhari]

openmp/libomptarget/test/lit.cfg

97

... and looked for flang-new in the binary dir with CMake's find_program for in-tree builds. Unfortunately, that version didn't work in certain circumstances.

Right, without the CMake dependency between check-libomptarget and flang-new this will not work when flang-new is built later than all prerequisites for check-libomptarget are satisfied.
I believe the current version has the same issue.

One of the solutions is to add the target dependency in libomptarget/test/CMakeLists.txt when flang is present in LLVM_ENABLE_PROJECTS.

[elmcdonough]

Ah okay, thanks for clarifying. I haven't had any issues when testing it locally, but I'll make sure to look into this.

[elmcdonough]

Add flang-new as dependency when flang is enabled

[jdoerfert]

Looks reasonable, @jhuber6 , @tianshilei1992, wdyt?

openmp/README.rst

**LIBOMP_ENABLE_SHARED** = ``ON|OFF``
 
   .. note::
 
     Static libraries are not supported on Windows*.
 
 **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
 include both i386 and x86_64 architecture objects in a single archive.
 
 .. code-block:: console
 
   $ cmake -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -DCMAKE_OSX_ARCHITECTURES='i386;x86_64' ..

openmp/libomptarget/test/CMakeLists.txt

foreach(CURRENT_TARGET IN LISTS SYSTEM_TARGETS)
   list(APPEND LIBOMPTARGET_LIT_TESTSUITES ${CMAKE_CURRENT_BINARY_DIR}/${CURRENT_TARGET})
 
   # Configure the lit.site.cfg.in file
   set(AUTO_GEN_COMMENT "## Autogenerated by libomptarget configuration.\n# Do not edit!")
   get_filename_component(CUDA_LIBDIR "${CUDA_cudart_static_LIBRARY}" DIRECTORY)
   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})

openmp/libomptarget/test/lit.cfg

 # -*- Python -*- vim: set ft=python ts=4 sw=4 expandtab tw=79:
 # Configuration file for the 'lit' test runner.
 
 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()
     lit_config = object()
 
 # Use the CUDA device as suggested by the env
 if 'CUDA_VISIBLE_DEVICES' in os.environ:
     config.environment['CUDA_VISIBLE_DEVICES'] = os.environ['CUDA_VISIBLE_DEVICES']

 config.suffixes = ['.c', '.cpp', '.cc']
 
 # test_source_root: The root path where tests are located.
 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()
 
 # compiler flags
 config.test_flags = " -I " + config.test_source_root + \
     " -I " + config.omp_header_directory + \
     " -L " + config.library_dir + \
     " -L " + config.llvm_lib_directory

 if config.libomptarget_debug:
   config.available_features.add('libomptarget-debug')
 
 if config.has_libomptarget_ompt:
   config.available_features.add('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.
 supports_unified_shared_memory = True
 if config.libomptarget_current_target.startswith('nvptx'):
   try:
     cuda_arch = int(config.cuda_test_arch[:3])
     if cuda_arch < 70:

 config.substitutions.append(("%clang", config.test_c_compiler))
 config.substitutions.append(("%openmp_flags", config.test_openmp_flags))
 if config.libomptarget_current_target.startswith('nvptx') and config.cuda_path:
     config.substitutions.append(("%cuda_flags", "--cuda-path=" + config.cuda_path))
 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)

openmp/libomptarget/test/lit.site.cfg.in

 @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@
 config.test_openmp_flags = "@OPENMP_TEST_OPENMP_FLAGS@"
 config.test_extra_flags = "@OPENMP_TEST_FLAGS@"
 config.cuda_path = "@CUDA_TOOLKIT_ROOT_DIR@"
 config.cuda_libdir = "@CUDA_LIBDIR@"
 config.cuda_test_arch = "@LIBOMPTARGET_DEP_CUDA_ARCH@"
 config.libomptarget_obj_root = "@CMAKE_CURRENT_BINARY_DIR@/@CURRENT_TARGET@"
 config.library_dir = "@LIBOMPTARGET_LIBRARY_DIR@"
 config.omp_header_directory = "@LIBOMPTARGET_OPENMP_HEADER_FOLDER@"
 config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
 config.llvm_lib_directory = "@LIBOMPTARGET_LLVM_LIBRARY_DIR@"
 config.operating_system = "@CMAKE_SYSTEM_NAME@"
 config.libomptarget_all_targets = "@LIBOMPTARGET_ALL_TARGETS@".split()
 config.libomptarget_current_target = "@CURRENT_TARGET@"
 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")

openmp/libomptarget/test/offloading/fortran/basic_array.c

+// 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 <stdio.h>
+#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

openmp/libomptarget/test/offloading/fortran/basic_array.f90

+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