[Libomptarget] Run GPU offloading tests using the new driver

[jhuber6]

Summary

This patch adds a new target to the tests to run using the new driver as
the method for generating offloading code.

Depends on D116541

[jhuber6]

Adding XFAIL to test.

[JonChesterfield]

openmp/libomptarget/test/lit.cfg

114

You need to generalise / replace this function to avoid passing nvptx64-newDriver or similar as the triple

[jhuber6]

Addressing triple and adding test for static linking.

[jhuber6]

Add AMDGPU

[JonChesterfield]

Looks good. We can land this and revert with greater knowledge if the amdgpu buildbot fails on it.

openmp/libomptarget/CMakeLists.txt

40

Could drop the amdgcn-amd-amdhsa and nvptx64-nvidia-cuda lines in passing, I should have removed them previously. Or leave them and we'll remove the -newRTL ones instead of renaming them.

openmp/libomptarget/test/lit.cfg

121

LG. We should drop this newRTL thing at some point.

[jhuber6]

openmp/libomptarget/test/lit.cfg

121

Yeah, I'll soon make a patch that removes the old runtime.

openmp/libomptarget/CMakeLists.txt

 
 # LLVM source tree is required at build time for libomptarget
 if (NOT LIBOMPTARGET_LLVM_INCLUDE_DIRS)
   message(FATAL_ERROR "Missing definition for LIBOMPTARGET_LLVM_INCLUDE_DIRS")
 endif()
 
 # 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} amdgcn-amd-amdhsa")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} amdgcn-amd-amdhsa-newRTL")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} amdgcn-amd-amdhsa-newDriver")
 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} nvptx64-nvidia-cuda-newRTL")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda-newDriver")
 
 # 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/amdgpu/CMakeLists.txt

 get_property(AMDGPU_ARCH_COMMAND TARGET amdgpu-arch PROPERTY LOCATION)
 
 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-newRTL" PARENT_SCOPE)
+  set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} amdgcn-amd-amdhsa-newRTL amdgcn-amd-amdhsa-newDriver " PARENT_SCOPE)
 endif()
 

openmp/libomptarget/plugins/cuda/CMakeLists.txt

target_link_libraries(omptarget.rtl.cuda
   "-Wl,-z,defs")
 
 # 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-newRTL" PARENT_SCOPE)
+  set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS} nvptx64-nvidia-cuda-newRTL nvptx64-nvidia-cuda-newDriver" 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('-newRTL'):
         config.test_flags += " -fopenmp-target-new-runtime"
-    else:
+    elif not config.libomptarget_current_target.endswith('-newDriver'):
         config.test_flags += " -fno-openmp-target-new-runtime"
+    if config.libomptarget_current_target.endswith('-newDriver'):
+        config.test_flags += " -fopenmp-new-driver"
 
 def remove_newRTL_suffix_if_present(name):
     if name.endswith('-newRTL'):
         return name[:-7]
+    elif name.endswith('-newDriver'):
+        return name[:-10]
     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.
 
 # 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

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-newDriver
+// REQUIRES: amdgcn-amd-amdhsa-newDriver
+
+#ifdef LIBRARY
+int x = 42;
+#pragma omp declare target(x)
+
+int foo() {
+  int value;
+#pragma omp target map(from : value)
+  value = x;
+  return value;
+}
+#else
+#include <stdio.h>
+int foo();
+
+int main() {
+  int x = foo();
+
+  // CHECK: PASS
+  if (x == 42)
+    printf("PASS\n");
+}
+#endif

openmp/libomptarget/test/unified_shared_memory/api.c

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