Enable OpenMP offloading to VE and enable tests for offloading to VE

[kosmynin]

Summary

OpenMP offloading to VE requires a cross-compiled OpenMP runtime and compiler-rt.

The compiler-rt files need to be placed in the build directory to run the tests, but the OpenMP runtime can be placed elsewhere with the CMake variable LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER pointing to it.
The patch uses this CMake variable to set the library search path and the runtime VE_LD_LIBRARY_PATH environment variable for VE tests.

[protze.joachim]

openmp/libomptarget/test/lit.cfg

67

Unrelated to this patch, but why are different separators used for Darwin?

78–83

I think, this should only be set for the specific device RUN-lines. I suggest to include the flags and export of the environment variable in the definition of %libomptarget-compile-ve-unknown-linux-unknown.

[manorom]

• Move VE specific vars into test run line

clang/lib/Frontend/CompilerInvocation.cpp

 //===- CompilerInvocation.cpp ---------------------------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
for (unsigned i = 0; i < A->getNumValues(); ++i) {
       if (TT.getArch() == llvm::Triple::UnknownArch ||
           !(TT.getArch() == llvm::Triple::aarch64 ||
             TT.getArch() == llvm::Triple::ppc ||
             TT.getArch() == llvm::Triple::ppc64 ||
             TT.getArch() == llvm::Triple::ppc64le ||
             TT.getArch() == llvm::Triple::nvptx ||
             TT.getArch() == llvm::Triple::nvptx64 ||
             TT.getArch() == llvm::Triple::amdgcn ||
+            TT.getArch() == llvm::Triple::ve ||
             TT.getArch() == llvm::Triple::x86 ||
             TT.getArch() == llvm::Triple::x86_64))
         Diags.Report(diag::err_drv_invalid_omp_target) << A->getValue(i);
       else
         Opts.OMPTargetTriples.push_back(TT);
     }
   }
 

openmp/libomptarget/CMakeLists.txt

 include(LibomptargetGetDependencies)
 
 # This is a list of all the targets that are supported/tested right now.
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} aarch64-unknown-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64le-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} x86_64-pc-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} ve-unknown-linux-unknown")
 
 # Once the plugins for the different targets are validated, they will be added to
 # the list of supported targets in the current system.
 set (LIBOMPTARGET_SYSTEM_TARGETS "")
 set (LIBOMPTARGET_TESTED_PLUGINS "")
 
 # Check whether using debug mode. In debug mode, allow dumping progress
 # messages at runtime by default. Otherwise, it can be enabled

openmp/libomptarget/test/api/omp_get_num_devices.c

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compile-ve-unknown-linux-unknown | env VE_LD_LIBRARY_PATH=%host-rtl-dir %libomptarget-run-ve-unknown-linux-unknown | %fcheck-ve-unknown-linux-unknown
 
 #include <stdio.h>
 #include <omp.h>
 
 int test_omp_get_num_devices()
 {
   /* checks that omp_get_num_devices() > 0 */
   int num_devices = omp_get_num_devices();

openmp/libomptarget/test/api/omp_get_num_devices_with_empty_target.c

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+// RUN: env VE_LD_LIBRARY_PATH=%host-rtl-dir %libomptarget-compile-run-and-check-ve-unknown-linux-unknown
 
 #include <omp.h>
 #include <stdio.h>
 
 static int test_omp_get_num_devices_with_empty_target() {
   /* checks that omp_get_num_devices() > 0 */
   return omp_get_num_devices() > 0;
 }

openmp/libomptarget/test/env/omp_target_debug.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=NDEBUG
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=NDEBUG
+// RUN: %libomptarget-compile-ve-unknown-linux-unknown && env VE_LD_LIBRARY_PATH=%host-rtl-dir LIBOMPTARGET_DEBUG=1 %libomptarget-run-ve-unknown-linux-unknown 2>&1 | %fcheck-ve-unknown-linux-unknown -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-ve-unknown-linux-unknown && env VE_LD_LIBRARY_PATH=%host-rtl-dir LIBOMPTARGET_DEBUG=0 %libomptarget-run-ve-unknown-linux-unknown 2>&1 | %fcheck-ve-unknown-linux-unknown -allow-empty -check-prefix=NDEBUG
 // REQUIRES: libomptarget-debug
 
 int main(void) {
 #pragma omp target
   {}
   return 0;
 }
 
 // DEBUG: Libomptarget
 // NDEBUG-NOT: Libomptarget
 // NDEBUG-NOT: Target
 

openmp/libomptarget/test/lit.cfg

 
 # Setup environment to find dynamic library at runtime
 if config.operating_system == 'Windows':
     append_dynamic_library_path('PATH', config.library_dir, ";")
     append_dynamic_library_path('PATH', config.omp_host_rtl_directory, ";")
 elif config.operating_system == 'Darwin':
     append_dynamic_library_path('DYLD_LIBRARY_PATH', config.library_dir, ":")
     append_dynamic_library_path('DYLD_LIBRARY_PATH', \
         config.omp_host_rtl_directory, ";")
     config.test_flags += " -Wl,-rpath," + config.library_dir
     config.test_flags += " -Wl,-rpath," + config.omp_host_rtl_directory
 else: # Unices
     append_dynamic_library_path('LD_LIBRARY_PATH', config.library_dir, ":")
     append_dynamic_library_path('LD_LIBRARY_PATH', \
         config.omp_host_rtl_directory, ":")
     append_dynamic_library_path('LIBRARY_PATH', config.library_dir, ":")
     append_dynamic_library_path('LIBRARY_PATH', \
         config.omp_host_rtl_directory, ":")
 
+# Target specific flags
+target_flags = {}
+target_flags['ve-unknown-linux-unknown'] = "-Xopenmp-target \"-L " + \
+    config.omp_host_rtl_directory + "\""
+
 # substitutions
 # - for targets that exist in the system create the actual command.
 # - for valid targets that do not exist in the system, return false, so that the
 #   same test can be used for different targets.
 
 # Scan all the valid targets.
 for libomptarget_target in config.libomptarget_all_targets:
     # Is this target in the current system? If so create a compile, run and test
     # command. Otherwise create command that return false.
if libomptarget_target in config.libomptarget_system_targets:
             libomptarget_target))
         config.substitutions.append(("%libomptarget-run-" + \
             libomptarget_target, \
             "%t-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-run-fail-" + \
             libomptarget_target, \
             "%not --crash %t-" + libomptarget_target))
         config.substitutions.append(("%clangxx-" + libomptarget_target, \
-            "%clangxx %openmp_flags %flags -fopenmp-targets=" + libomptarget_target))
+            "%clangxx %openmp_flags %flags-" + libomptarget_target + \
+            " -fopenmp-targets=" + libomptarget_target))
         config.substitutions.append(("%clang-" + libomptarget_target, \
-            "%clang %openmp_flags %flags -fopenmp-targets=" + libomptarget_target))
+            "%clang %openmp_flags %flags-" + libomptarget_target + \
+            " -fopenmp-targets=" + libomptarget_target))
+        config.substitutions.append(("%flags-" + libomptarget_target, \
+            "%flags " + target_flags.get(libomptarget_target, "")))
         config.substitutions.append(("%fcheck-" + libomptarget_target, \
             config.libomptarget_filecheck + " %s"))
     else:
         config.substitutions.append(("%libomptarget-compile-run-and-check-" + \
             libomptarget_target, \
             "echo ignored-command"))
         config.substitutions.append(("%libomptarget-compilexx-run-and-check-" + \
             libomptarget_target, \
else:
             "echo ignored-command"))
         config.substitutions.append(("%fcheck-" + libomptarget_target, \
             "echo ignored-command"))
 
 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))
 config.substitutions.append(("%flags", config.test_flags))
+config.substitutions.append(("%host-rtl-dir", config.omp_host_rtl_directory))
 config.substitutions.append(("%not", config.libomptarget_not))

openmp/libomptarget/test/mapping/alloc_fail.c

 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu
 // RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
 // RUN: | %fcheck-x86_64-pc-linux-gnu
 
 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda
 // RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 \
 // RUN: | %fcheck-nvptx64-nvidia-cuda
 
+// RUN: %libomptarget-compile-ve-unknown-linux-unknown
+// RUN: env VE_LD_LIBRARY_PATH=%host-rtl-dir %libomptarget-run-fail-ve-unknown-linux-unknown 2>&1 \
+// RUN: | %fcheck-ve-unknown-linux-unknown
+
 // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
 
 int main() {
   int arr[4] = {0, 1, 2, 3};
 #pragma omp target data map(alloc: arr[0:2])
 #pragma omp target data map(alloc: arr[1:2])
   ;
   return 0;
 }

openmp/libomptarget/test/mapping/target_implicit_partial_map.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
 // RUN: | %fcheck-aarch64-unknown-linux-gnu
 
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
 // RUN: | %fcheck-powerpc64-ibm-linux-gnu
 
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
 // RUN: | %fcheck-powerpc64le-ibm-linux-gnu
 
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu
 // RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
 // RUN: | %fcheck-x86_64-pc-linux-gnu
+
+// RUN: %libomptarget-compile-ve-unknown-linux-unknown
+// RUN: env VE_LD_LIBRARY_PATH=%host-rtl-dir %libomptarget-run-ve-unknown-linux-unknown 2>&1 \
+// RUN: | %fcheck-ve-unknown-linux-unknown
 //
 // END.
 
 #include <omp.h>
 #include <stdio.h>
 
 int main() {
   int arr[100];

openmp/libomptarget/test/offloading/offloading_success.c

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-ve-unknown-linux-unknown | env VE_LD_LIBRARY_PATH=%host-rtl-dir %libomptarget-run-ve-unknown-linux-unknown | %fcheck-ve-unknown-linux-unknown
 
 #include <stdio.h>
 #include <omp.h>
 
 int main(void) {
   int isHost = -1;
 
 #pragma omp target map(from: isHost)