[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

90

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

90

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

90

... 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/CMakeLists.txt

if (OPENMP_STANDALONE_BUILD)
   # Do not use OPENMP_LIBDIR_SUFFIX directly, use OPENMP_INSTALL_LIBDIR.
   set(OPENMP_INSTALL_LIBDIR "lib${OPENMP_LIBDIR_SUFFIX}")
 
   # Group test settings.
   set(OPENMP_TEST_C_COMPILER ${CMAKE_C_COMPILER} CACHE STRING
     "C compiler to use for testing OpenMP runtime libraries.")
   set(OPENMP_TEST_CXX_COMPILER ${CMAKE_CXX_COMPILER} CACHE STRING
     "C++ compiler to use for testing OpenMP runtime libraries.")
+  set(OPENMP_TEST_Fortran_COMPILER ${CMAKE_Fortran_COMPILER} CACHE STRING
+    "Fortran compiler to use for testing OpenMP runtime libraries.")
   set(OPENMP_LLVM_TOOLS_DIR "" CACHE PATH "Path to LLVM tools for testing.")
 else()
   set(OPENMP_ENABLE_WERROR ${LLVM_ENABLE_WERROR})
   # If building in tree, we honor the same install suffix LLVM uses.
   set(OPENMP_INSTALL_LIBDIR "lib${LLVM_LIBDIR_SUFFIX}")
 
   if (NOT MSVC)
     set(OPENMP_TEST_C_COMPILER ${LLVM_RUNTIME_OUTPUT_INTDIR}/clang)
     set(OPENMP_TEST_CXX_COMPILER ${LLVM_RUNTIME_OUTPUT_INTDIR}/clang++)
   else()
     set(OPENMP_TEST_C_COMPILER ${LLVM_RUNTIME_OUTPUT_INTDIR}/clang.exe)
     set(OPENMP_TEST_CXX_COMPILER ${LLVM_RUNTIME_OUTPUT_INTDIR}/clang++.exe)
   endif()
+
+  # Check if flang is in build bin folder
+  find_program(OPENMP_TEST_Fortran_COMPILER NAMES flang-new flang
+    PATHS ${LLVM_RUNTIME_OUTPUT_INTDIR} NO_DEFAULT_PATH)
 endif()
 
 # Check and set up common compiler flags.
 include(config-ix)
 include(HandleOpenMPOptions)
 
 # Set up testing infrastructure.
 include(OpenMPTesting)

openmp/README.rst

 **LIBOMPTARGET_OPENMP_HEADER_FOLDER** = ``""``
   Path of the folder that contains ``omp.h``.  This is required for testing
   out-of-tree builds.
 
 **LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER** = ``""``
   Path of the folder that contains ``libomp.so``, and ``libLLVMSupport.so``
   when profiling is enabled.  This is required for testing.
 
+**OPENMP_TEST_Fortran_COMPILER** = ``${CMAKE_Fortran_COMPILER}``
+  Compiler to use for testing. Defaults to the compiler that was also used for
+  building when building libomptarget standalone and flang for in-tree builds.
+
 Options for ``NVPTX device RTL``
 --------------------------------
 
 **LIBOMPTARGET_NVPTX_ENABLE_BCLIB** = ``ON|OFF``
   Enable CUDA LLVM bitcode offloading device RTL. This is used for link time
   optimization of the OMP runtime and application code. This option is enabled
   by default if the build system determines that `CMAKE_C_COMPILER` is able to
   compile and link the library.

openmp/libomptarget/test/lit.cfg

 
 if config.has_libomptarget_ompt:
   config.available_features.add('ompt')
 
 config.available_features.add(config.libomptarget_current_target)
 
 # 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:
       supports_unified_shared_memory = False
   except ValueError:
     # If the architecture is invalid, assume it is supported.
else:
             "echo ignored-command"))
         config.substitutions.append(("%clang-" + libomptarget_target, \
             "echo ignored-command"))
         config.substitutions.append(("%clangxx-" + libomptarget_target, \
             "echo ignored-command"))
         config.substitutions.append(("%fcheck-" + libomptarget_target, \
             "echo ignored-command"))
 
+if config.test_fortran_compiler:
+    config.available_features.add("flang")
+    config.substitutions.append(("%flang", config.test_fortran_compiler))
+
 config.substitutions.append(("%clangxx", config.test_cxx_compiler))
 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(("%llvm_lib", config.llvm_lib_directory))
 config.substitutions.append(("%flags", config.test_flags))
 config.substitutions.append(("%not", config.libomptarget_not))

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

 @AUTO_GEN_COMMENT@
 
 config.test_c_compiler = "@OPENMP_TEST_C_COMPILER@"
 config.test_cxx_compiler = "@OPENMP_TEST_CXX_COMPILER@"
+config.test_fortran_compiler = "@OPENMP_TEST_Fortran_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@"

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: %clang-amdgcn-amd-amdhsa -L%llvm_lib basic_array.o %s -o %t
+// 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