[Libomptarget] Add test config for compiling in LTO-mode

[jhuber6]

Summary

We are planning on making LTO the default compilation mode for
offloading. In order to make sure it works we should run these tests on
the test suite. AMDGPU already uses the LTO compilation path for its
linking, but in LTO mode it also links the static library late.

Performing LTO requires the static library to be built, if we make the
change this will be a hard requirement and the old bitcode library will
go away. This means users will need to use either a two-step build or a
runtimes build for libomptarget.

[JonChesterfield]

Looks good here. Will need to back out and change the xfail lines if the CI rejects it.

openmp/libomptarget/CMakeLists.txt

if (NOT LIBOMPTARGET_LLVM_INCLUDE_DIRS)
   message(FATAL_ERROR "Missing definition for LIBOMPTARGET_LLVM_INCLUDE_DIRS")
 endif()
 
 include_directories(${LIBOMPTARGET_LLVM_INCLUDE_DIRS})
 
 # 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} aarch64-unknown-linux-gnu-oldDriver")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} aarch64-unknown-linux-gnu-LTO")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} amdgcn-amd-amdhsa")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} amdgcn-amd-amdhsa-oldDriver")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} amdgcn-amd-amdhsa-LTO")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64le-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64le-ibm-linux-gnu-oldDriver")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64le-ibm-linux-gnu-LTO")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64-ibm-linux-gnu-oldDriver")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64-ibm-linux-gnu-LTO")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} x86_64-pc-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} x86_64-pc-linux-gnu-oldDriver")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} x86_64-pc-linux-gnu-LTO")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda-oldDriver")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda-LTO")
 
 # 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/plugins/CMakeLists.txt

if(LIBOMPTARGET_DEP_LIBFFI_FOUND)
         "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/../exports")
 
       list(APPEND LIBOMPTARGET_TESTED_PLUGINS
         "omptarget.rtl.${tmachine_libname}")
 
       # Report to the parent scope that we are building a plugin.
       set(LIBOMPTARGET_SYSTEM_TARGETS
         "${LIBOMPTARGET_SYSTEM_TARGETS} ${tmachine_triple} ${tmachine_triple}-oldDriver" PARENT_SCOPE)
+      set(LIBOMPTARGET_SYSTEM_TARGETS
+        "${LIBOMPTARGET_SYSTEM_TARGETS} ${tmachine_triple} ${tmachine_triple}-LTO" PARENT_SCOPE)
       set(LIBOMPTARGET_TESTED_PLUGINS
         "${LIBOMPTARGET_TESTED_PLUGINS}" PARENT_SCOPE)
 
     else(LIBOMPTARGET_DEP_LIBFFI_FOUND)
       libomptarget_say("Not building ${tmachine_name} offloading plugin: libffi dependency not found.")
     endif(LIBOMPTARGET_DEP_LIBFFI_FOUND)
   else(LIBOMPTARGET_DEP_LIBELF_FOUND)
     libomptarget_say("Not building ${tmachine_name} offloading plugin: libelf dependency not found.")

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

 
 execute_process(COMMAND ${AMDGPU_ARCH_COMMAND} RESULT_VARIABLE amdgpu_arch_result
                                                OUTPUT_VARIABLE amdgpu_arch_output)
 if (${amdgpu_arch_result})
   libomptarget_say("Not generating amdgcn test targets as amdgpu-arch exited with ${amdgpu_arch_result}")
 else()
   # Report to the parent scope that we are building a plugin for amdgpu
   set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} amdgcn-amd-amdhsa amdgcn-amd-amdhsa-oldDriver" PARENT_SCOPE)
+  set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} amdgcn-amd-amdhsa amdgcn-amd-amdhsa-LTO" PARENT_SCOPE)
   list(APPEND LIBOMPTARGET_TESTED_PLUGINS "omptarget.rtl.amdgpu")
   set(LIBOMPTARGET_TESTED_PLUGINS "${LIBOMPTARGET_TESTED_PLUGINS}" PARENT_SCOPE)
 endif()
 

openmp/libomptarget/plugins/cuda/CMakeLists.txt

 
 # Report to the parent scope that we are building a plugin for CUDA.
 # This controls whether tests are run for the nvptx offloading target
 # Run them if libcuda is available, or if the user explicitly asked for dlopen
 # Otherwise this plugin is being built speculatively and there may be no cuda available
 if (LIBOMPTARGET_CAN_LINK_LIBCUDA OR LIBOMPTARGET_FORCE_DLOPEN_LIBCUDA)
   libomptarget_say("Enable tests using CUDA plugin")
   set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} nvptx64-nvidia-cuda nvptx64-nvidia-cuda-oldDriver" PARENT_SCOPE)
+  set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} nvptx64-nvidia-cuda nvptx64-nvidia-cuda-LTO" PARENT_SCOPE)
   list(APPEND LIBOMPTARGET_TESTED_PLUGINS "omptarget.rtl.cuda")
   set(LIBOMPTARGET_TESTED_PLUGINS "${LIBOMPTARGET_TESTED_PLUGINS}" PARENT_SCOPE)
 else()
   libomptarget_say("Disabling tests using CUDA plugin as cuda may not be available")
 endif()

openmp/libomptarget/test/lit.cfg

else: # Unices
     if config.cuda_libdir:
         config.test_flags += " -Wl,-rpath," + config.cuda_libdir
     if config.libomptarget_current_target.startswith('amdgcn'):
         config.test_flags += " --libomptarget-amdgcn-bc-path=" + config.library_dir
     if config.libomptarget_current_target.startswith('nvptx'):
         config.test_flags += " --libomptarget-nvptx-bc-path=" + config.library_dir
     if config.libomptarget_current_target.endswith('-oldDriver'):
         config.test_flags += " -fno-openmp-new-driver"
+    if config.libomptarget_current_target.endswith('-LTO'):
+        config.test_flags += " -offload-lto"
 
 def remove_suffix_if_present(name):
     if name.endswith('-oldDriver'):
         return name[:-10]
+    if name.endswith('-LTO'):
+        return name[:-4]
     else:
         return name
 
 # 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.
 

openmp/libomptarget/test/mapping/data_member_ref.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-oldDriver
+// XFAIL: amdgcn-amd-amdhsa-LTO
 
 #include <stdio.h>
 
 struct View {
   int Data;
 };
 
 struct ViewPtr {

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-oldDriver
+// XFAIL: amdgcn-amd-amdhsa-LTO
 
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C1;

openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-oldDriver
+// XFAIL: amdgcn-amd-amdhsa-LTO
 
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C;

openmp/libomptarget/test/mapping/lambda_by_value.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-oldDriver
+// XFAIL: amdgcn-amd-amdhsa-LTO
 
 #include <stdio.h>
 #include <stdint.h>
 
 // CHECK: before: [[V1:111]] [[V2:222]] [[PX:0x[^ ]+]] [[PY:0x[^ ]+]]
 // CHECK: lambda: [[V1]] [[V2]] [[PX_TGT:0x[^ ]+]] 0x{{.*}}
 // CHECK: tgt   : [[V2]] [[PX_TGT]] 1
 // CHECK: out   : [[V2]] [[V2]] [[PX]] [[PY]]

openmp/libomptarget/test/mapping/lambda_mapping.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // Error on the gpu that crashes the host
 // UNSUPPORTED: amdgcn-amd-amdhsa
 // UNSUPPORTED: amdgcn-amd-amdhsa-oldDriver
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
 
 #include <iostream>
 
 template <typename LOOP_BODY>
 inline void forall(int Begin, int End, LOOP_BODY LoopBody) {
 #pragma omp target parallel for schedule(static)
   for (int I = Begin; I < End; ++I) {
     LoopBody(I);

openmp/libomptarget/test/mapping/map_back_race.cpp

 // RUN: %libomptarget-compilexx-and-run-generic
 
 // Taken from https://github.com/llvm/llvm-project/issues/54216
 
-
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <algorithm>
 #include <cstdlib>
 #include <iostream>
 
 bool almost_equal(float x, float gold, float rel_tol = 1e-09,
                   float abs_tol = 0.0) {
   return std::abs(x - gold) <=

openmp/libomptarget/test/mapping/ompx_hold/struct.c

 // RUN: %libomptarget-compile-generic -fopenmp-extensions
 // RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
 
 // Wrong results on amdgpu
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-oldDriver
+// XFAIL: amdgcn-amd-amdhsa-LTO
 
 #include <omp.h>
 #include <stdio.h>
 
 #define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
   printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \
          #Var1, #Var2, #Var3,                                                  \
          omp_target_is_present(&(Var1), omp_get_default_device()),             \

openmp/libomptarget/test/offloading/bug49021.cpp

 // RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic
 
 // Hangs
 // UNSUPPORTED: amdgcn-amd-amdhsa
 // UNSUPPORTED: amdgcn-amd-amdhsa-oldDriver
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
 
 #include <iostream>
 
 template <typename T> int test_map() {
   std::cout << "map(complex<>)" << std::endl;
   T a(0.2), a_check;
 #pragma omp target map(from : a_check)
   { a_check = a; }

openmp/libomptarget/test/offloading/bug49334.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // Currently hangs on amdgpu
 // UNSUPPORTED: amdgcn-amd-amdhsa
 // UNSUPPORTED: amdgcn-amd-amdhsa-oldDriver
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <cassert>
 #include <cmath>
 #include <iostream>
 #include <limits>
 #include <memory>
 #include <vector>
 

openmp/libomptarget/test/offloading/bug49779.cpp

 // RUN: %libomptarget-compilexx-generic && \
 // RUN:   env LIBOMPTARGET_STACK_SIZE=2048 %libomptarget-run-generic
 
 // UNSUPPORTED: amdgcn-amd-amdhsa
 // UNSUPPORTED: amdgcn-amd-amdhsa-oldDriver
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
 
 #include <cassert>
 #include <iostream>
 
 void work(int *C) {
 #pragma omp atomic
   ++(*C);
 }

openmp/libomptarget/test/offloading/bug51781.c

 // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom
 // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
 //
 // CUSTOM: Rewriting generic-mode kernel with a customized state machine.
 
 // Hangs
 // UNSUPPORTED: amdgcn-amd-amdhsa
 // UNSUPPORTED: amdgcn-amd-amdhsa-oldDriver
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
 
 #if ADD_REDUCTION
 # define REDUCTION(...) reduction(__VA_ARGS__)
 #else
 # define REDUCTION(...)
 #endif
 
 #include <stdio.h>

openmp/libomptarget/test/offloading/host_as_target.c

 // Check that specifying device as omp_get_initial_device():
 // - Doesn't cause the runtime to fail.
 // - Offloads code to the host.
 // - Doesn't transfer data.  In this case, just check that neither host data nor
 //   default device data are affected by the specified transfers.
 // - Works whether it's specified directly or as the default device.
 
 // RUN: %libomptarget-compile-run-and-check-generic
 
 // amdgpu does not have a working printf definition
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-oldDriver
+// XFAIL: amdgcn-amd-amdhsa-LTO
 
 #include <stdio.h>
 #include <omp.h>
 
 static void check(char *X, int Dev) {
   printf("  host X = %c\n", *X);
   #pragma omp target device(Dev)
   printf("device X = %c\n", *X);

openmp/libomptarget/test/offloading/memory_manager.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <omp.h>
 
 #include <cassert>
 #include <iostream>
 
 int main(int argc, char *argv[]) {
 #pragma omp parallel for

openmp/libomptarget/test/offloading/parallel_offloading_map.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <cassert>
 #include <iostream>
 
 int main(int argc, char *argv[]) {
   constexpr const int num_threads = 64, N = 128;
   int array[num_threads] = {0};
 

openmp/libomptarget/test/offloading/static_linking.c

 // RUN: %libomptarget-compile-generic -DLIBRARY -c -o %t.o
 // RUN: llvm-ar rcs %t.a %t.o
 // RUN: %libomptarget-compile-generic %t.a && %libomptarget-run-generic 2>&1 | %fcheck-generic
 
-// REQUIRES: nvptx64-nvidia-cuda-oldDriver
-// REQUIRES: amdgcn-amd-amdhsa-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: amdgcn-amd-amdhsa-oldDriver
 
 #ifdef LIBRARY
 int x = 42;
 #pragma omp declare target(x)
 
 int foo() {
   int value;
 #pragma omp target map(from : value)

openmp/libomptarget/test/offloading/taskloop_offload_nowait.cpp

 // RUN: %libomptarget-compilexx-and-run-generic
 
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <cmath>
 #include <cstdlib>
 #include <iostream>
 
 bool almost_equal(float x, float gold, float tol) {
   if (std::signbit(x) != std::signbit(gold))
     x = std::abs(gold) - std::abs(x);

openmp/libomptarget/test/unified_shared_memory/api.c

 // RUN: %libomptarget-compile-run-and-check-generic
 // XFAIL: nvptx64-nvidia-cuda
-// XFAIL: nvptx64-nvidia-cuda
 // XFAIL: nvptx64-nvidia-cuda-oldDriver
+// XFAIL: nvptx64-nvidia-cuda-LTO
 
 // Fails on amdgpu with error: GPU Memory Error
 // XFAIL: amdgcn-amd-amdhsa
 // XFAIL: amdgcn-amd-amdhsa-oldDriver
+// XFAIL: amdgcn-amd-amdhsa-LTO
 
 #include <stdio.h>
 #include <omp.h>
 
 // ---------------------------------------------------------------------------
 // Various definitions copied from OpenMP RTL
 
 extern void __tgt_register_requires(int64_t);