[OpenMP][Tests] Enable nvptx64 testing for most libomptarget tests

[protze.joachim]

Summary

This patch adds a run line to most libomptarget tests to also test the nvptx64 target, if available.

Adding $BUILD/lib to the LIBRARY_PATH fixes https://bugs.llvm.org/show_bug.cgi?id=46836.

I moved the existing runline for nvptx64 to the end. Lit will stop for the first failing runline and mark the test as failed. If the GPU is configured for exclusive mode, this line tends to fail quite often.

None of the newly enabled tests fail on my system.

[JonChesterfield]

Not Johannes, but LGTM. I didn't realise our tests could pick up a different runtime to the one just built, that was a nasty failure mode. Thanks for the patch

[protze.joachim]

Thanks for reviewing.

At the time of the commit 9 tests are failing during offloading to V100:

Broken because of D83963, @jdenny proposed a fix, which works for me:

libomptarget :: mapping/alloc_fail.c
libomptarget :: mapping/present/target.c
libomptarget :: mapping/present/target_data.c
libomptarget :: mapping/present/target_enter_data.c
libomptarget :: mapping/present/target_exit_data.c
libomptarget :: mapping/present/zero_length_array_section.c
libomptarget :: mapping/present/zero_length_array_section_exit.c

Broken since D68100, solution seems to be unclear:

libomptarget :: mapping/declare_mapper_target.cpp

Fix is discussed in D84470, patch works for me:

libomptarget :: offloading/target_depend_nowait.cpp

[jdoerfert]

In D84557#2178019, @protze.joachim wrote:

Thanks for reviewing.

At the time of the commit 9 tests are failing during offloading to V100:

Broken because of D83963, @jdenny proposed a fix, which works for me:

libomptarget :: mapping/alloc_fail.c
libomptarget :: mapping/present/target.c
libomptarget :: mapping/present/target_data.c
libomptarget :: mapping/present/target_enter_data.c
libomptarget :: mapping/present/target_exit_data.c
libomptarget :: mapping/present/zero_length_array_section.c
libomptarget :: mapping/present/zero_length_array_section_exit.c

Broken since D68100, solution seems to be unclear:

libomptarget :: mapping/declare_mapper_target.cpp

Fix is discussed in D84470, patch works for me:

libomptarget :: offloading/target_depend_nowait.cpp

Thanks. Could you ask @hans to include this in the release (via your bug) if you haven't already?

[jdenny]

In D84557#2178019, @protze.joachim wrote:

Broken because of D83963, @jdenny proposed a fix, which works for me:

libomptarget :: mapping/alloc_fail.c
libomptarget :: mapping/present/target.c
libomptarget :: mapping/present/target_data.c
libomptarget :: mapping/present/target_enter_data.c
libomptarget :: mapping/present/target_exit_data.c
libomptarget :: mapping/present/zero_length_array_section.c
libomptarget :: mapping/present/zero_length_array_section_exit.c

I've pushed it to master as 9b4826d18b5f.

[protze.joachim]

In D84557#2178384, @jdoerfert wrote:

Thanks. Could you ask @hans to include this in the release (via your bug) if you haven't already?

The release branch has no libomptarget tests with run lines for nvptx64. Therefore the release branch is not affected.

[JonChesterfield]

I'm having second thoughts about this.

It works because clang searches LIBRARY_PATH first when looking for the deviceRTL. That means a system that has LIBRARY_PATH accidentally/incidentally pointing to a different toolchain to the one that contains the running clang will pick up the wrong library. That's relatively likely to fail in opaque fashion.

Would you accept changing the tests to pass the exact location of the deviceRTL to use instead? The command line flag is already there, and wins out over LIBRARY_PATH. Closely related, I'd like to delete the behaviour of checking the LIBRARY_PATH to find the deviceRTL to use with a particular clang, so that only the people who consciously choose to point it at a different device runtime experience the result. LIBRARY_PATH is too easy to set accidentally and/or forget about.

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
 // REQUIRES: libomptarget-debug
 
 int main(void) {
 #pragma omp target
   {}
   return 0;
 }
 
 // DEBUG: Libomptarget
 // NDEBUG-NOT: Libomptarget
 // NDEBUG-NOT: Target
 

openmp/libomptarget/test/lit.cfg

elif config.operating_system == 'Darwin':
     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, ":")
 
 # 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:

openmp/libomptarget/test/mapping/alloc_fail.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
 // RUN: | %fcheck-aarch64-unknown-linux-gnu
 
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
 // RUN: | %fcheck-powerpc64-ibm-linux-gnu
 
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
 // RUN: | %fcheck-powerpc64le-ibm-linux-gnu
 
 // 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
+
 // 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/declare_mapper_api.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
 
 #include <cstdio>
 #include <cstdlib>
 #include <vector>
 #include <cinttypes>
 
 // Data structure definitions copied from OpenMP RTL.
 struct MapComponentInfoTy {

openmp/libomptarget/test/mapping/declare_mapper_target.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
-// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
 // 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
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/declare_mapper_target_data.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
-// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
 // 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
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/declare_mapper_target_data_enter_exit.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
-// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
 // 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
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/declare_mapper_target_update.cpp

 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
-// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
 // 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
 
 #include <cstdio>
 #include <cstdlib>
 
 #define NUM 1024
 
 class C {
 public:

openmp/libomptarget/test/mapping/delete_inf_refcount.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
 
 #include <stdio.h>
 #include <omp.h>
 
 #pragma omp declare target
 int isHost;
 #pragma omp end declare target
 

openmp/libomptarget/test/mapping/pr38704.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
 
 // Clang 6.0 doesn't use the new map interface, undefined behavior when
 // the compiler emits "old" interface code for structures.
 // UNSUPPORTED: clang-6
 
 #include <stdio.h>
 #include <stdlib.h>
 

openmp/libomptarget/test/offloading/d2d_memcpy.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-aarch64-unknown-linux-gnu | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64-ibm-linux-gnu | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64le-ibm-linux-gnu | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-x86_64-pc-linux-gnu | %fcheck-x86_64-pc-linux-gnu -allow-empty
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-nvptx64-nvidia-cuda | %fcheck-nvptx64-nvidia-cuda -allow-empty
 
 #include <assert.h>
 #include <omp.h>
 #include <stdio.h>
 #include <stdlib.h>
 
 const int magic_num = 7;
 
 int main(int argc, char *argv[]) {
   const int N = 128;
   const int num_devices = omp_get_num_devices();
 
   // No target device, just return
   if (num_devices == 0) {
     printf("PASS\n");
     return 0;
   }
 
   const int src_device = 0;
-  int dst_device = 1;
-  if (dst_device >= num_devices)
-    dst_device = num_devices - 1;
+  int dst_device = num_devices - 1;
 
   int length = N * sizeof(int);
   int *src_ptr = omp_target_alloc(length, src_device);
   int *dst_ptr = omp_target_alloc(length, dst_device);
 
   assert(src_ptr && "src_ptr is NULL");
   assert(dst_ptr && "dst_ptr is NULL");
 

openmp/libomptarget/test/offloading/dynamic_module.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-aarch64-unknown-linux-gnu %t.so && %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-powerpc64-ibm-linux-gnu %t.so && %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-powerpc64le-ibm-linux-gnu %t.so && %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-x86_64-pc-linux-gnu %t.so && %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-nvptx64-nvidia-cuda %t.so && %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda
 
 #ifdef SHARED
 void foo() {}
 #else
 #include <stdio.h>
 int main() {
 #pragma omp target
   ;
   // CHECK: DONE.
   printf("%s\n", "DONE.");
   return 0;
 }
 #endif

openmp/libomptarget/test/offloading/dynamic_module_load.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-aarch64-unknown-linux-gnu -ldl && %libomptarget-run-aarch64-unknown-linux-gnu %t.so 2>&1 | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-powerpc64-ibm-linux-gnu -ldl && %libomptarget-run-powerpc64-ibm-linux-gnu %t.so 2>&1 | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-powerpc64le-ibm-linux-gnu -ldl && %libomptarget-run-powerpc64le-ibm-linux-gnu %t.so 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-x86_64-pc-linux-gnu -ldl && %libomptarget-run-x86_64-pc-linux-gnu %t.so 2>&1 | %fcheck-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t-nvptx64-nvidia-cuda -ldl && %libomptarget-run-nvptx64-nvidia-cuda %t.so 2>&1 | %fcheck-nvptx64-nvidia-cuda
 
 #ifdef SHARED
 #include <stdio.h>
 int foo() {
 #pragma omp target
   ;
   printf("%s\n", "DONE.");
   return 0;

openmp/libomptarget/test/offloading/looptripcnt.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-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-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-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-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
 // REQUIRES: libomptarget-debug
 
 /*
   Test for looptripcount being popped from runtime stack.
 */
 #include <stdio.h>
 #include <omp.h>
 int main()

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
 
 #include <stdio.h>
 #include <omp.h>
 
 int main(void) {
   int isHost = -1;
 
 #pragma omp target map(from: isHost)

openmp/libomptarget/test/offloading/offloading_success.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
 
 #include <stdio.h>
 #include <omp.h>
 
 int main(void) {
   int isHost = 0;
 
 #pragma omp target map(from: isHost)

openmp/libomptarget/test/offloading/parallel_offloading_map.c

 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-aarch64-unknown-linux-gnu | %fcheck-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64-ibm-linux-gnu | %fcheck-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64le-ibm-linux-gnu | %fcheck-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-x86_64-pc-linux-gnu | %fcheck-x86_64-pc-linux-gnu -allow-empty
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-nvptx64-nvidia-cuda | %fcheck-nvptx64-nvidia-cuda -allow-empty
 #include <assert.h>
 #include <stdio.h>
 
 int main(int argc, char *argv[]) {
   const int num_threads = 64, N = 128;
   int array[num_threads] = {0};
 
 #pragma omp parallel for

openmp/libomptarget/test/offloading/requires.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-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-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-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-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=DEBUG
 // REQUIRES: libomptarget-debug
 
 /*
   Test for the 'requires' clause check.
   When a target region is used, the requires flags are set in the
   runtime for the entire compilation unit. If the flags are set again,
   (for whatever reason) the set must be consistent with previously
   set values.

openmp/libomptarget/test/offloading/target_depend_nowait.cpp

 // RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
 // RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
 
 #include <omp.h>
 #include <stdio.h>
 
 #define N 1024
 
 int A[N];
 int B[N];