[OpenMP] Enable more runtime tests and also run them with O3

[jdoerfert]

Summary

The test run fine on my AMD GPU machine, we should verify them on others
too and put them into our regular testing. Not testing O1/2/3 is really
bad and not testing all architecturs is similarly problematic.

[jplehr]

Works for me on gfx90a

[jhuber6]

In D148576#4277340, @jplehr wrote:

Works for me on gfx90a

Can you verify whether or not __builtin_trap or the s_trap 2 invocation hangs on those boxes? I can't tell if it's just my install or a wider problem.

[jplehr]

In D148576#4277344, @jhuber6 wrote:

In D148576#4277340, @jplehr wrote:

Works for me on gfx90a

Can you verify whether or not __builtin_trap or the s_trap 2 invocation hangs on those boxes? I can't tell if it's just my install or a wider problem.

I ran the tests several times and it seemed to work. Don't know if that counts exactly as verify.

[jdoerfert]

ping

[jhuber6]

Sure, a few tests already hang for me on AMDGPU for unknown reasons, so the extra one w/ __builtin_trap is the least of my worries.

openmp/libomptarget/test/api/omp_device_managed_memory_alloc.c

-// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
-// REQUIRES: nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-run-and-check-generic
+// RUN: %libomptarget-compileopt-run-and-check-generic
 
 #include <omp.h>
 #include <stdio.h>
 
 int main() {
   const int N = 64;
 
   // Allocates device managed memory that is shared between the host and device.

openmp/libomptarget/test/api/omp_dynamic_shared_memory.c

-// RUN: %libomptarget-compile-nvptx64-nvidia-cuda
+// RUN: %libomptarget-compile-generic
 // RUN: env LIBOMPTARGET_SHARED_MEMORY_SIZE=256 \
-// RUN:   %libomptarget-run-nvptx64-nvidia-cuda | %fcheck-nvptx64-nvidia-cuda
-// REQUIRES: nvptx64-nvidia-cuda
+// RUN:   %libomptarget-run-generic | %fcheck-generic
+
+// RUN: %libomptarget-compileopt-generic
+// RUN: env LIBOMPTARGET_SHARED_MEMORY_SIZE=256 \
+// RUN:   %libomptarget-run-generic | %fcheck-generic
+
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <omp.h>
 #include <stdio.h>
 
 int main() {
   int x;
 #pragma omp target parallel map(from : x)
   {

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
+// On AMDGPU we don't have malloc support yet. We need optimizations
+// to avoid a thread state which requires malloc.
+//
+// XUN: %libomptarget-compilexx-run-and-check-generic
+// RUN: %libomptarget-compileoptxx-run-and-check-generic
 
 #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/offloading/assert.cpp

-// RUN: %libomptarget-compilexx-nvptx64-nvidia-cuda && %libomptarget-run-fail-nvptx64-nvidia-cuda
-// REQUIRES: nvptx64-nvidia-cuda
+// RUN: %libomptarget-compilexx-generic && %libomptarget-run-fail-generic
+// RUN: %libomptarget-compileoptxx-generic && %libomptarget-run-fail-generic
 
 int main(int argc, char *argv[]) {
 #pragma omp target
   { __builtin_trap(); }
 
   return 0;
 }

openmp/libomptarget/test/offloading/barrier_fence.c

 // RUN: %libomptarget-compile-generic -fopenmp-offload-mandatory -O3
 // RUN: %libomptarget-run-generic
+// RUN: %libomptarget-compileopt-generic -fopenmp-offload-mandatory -O3
+// RUN: %libomptarget-run-generic
 
 // UNSUPPORTED: aarch64-unknown-linux-gnu
 // UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <omp.h>
 #include <stdio.h>

openmp/libomptarget/test/offloading/bug47654.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
+// RUN: %libomptarget-compileoptxx-run-and-check-generic
 
 #include <cassert>
 #include <iostream>
 
 int main(int argc, char *argv[]) {
   int i = 0, j = 0;
 
 #pragma omp target map(tofrom : i, j) nowait

openmp/libomptarget/test/offloading/bug49021.cpp

 // RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic
 // RUN: %libomptarget-compilexx-generic -O3 -ffast-math && %libomptarget-run-generic
+// RUN: %libomptarget-compileoptxx-generic -O3 && %libomptarget-run-generic
+// RUN: %libomptarget-compileoptxx-generic -O3 -ffast-math && %libomptarget-run-generic
 
 #include <iostream>
 
 template <typename T> int test_map() {
   std::cout << "map(T)" << std::endl;
   T a(0.2), a_check;
 #pragma omp target map(from : a_check)
   { a_check = a; }

openmp/libomptarget/test/offloading/bug49334.cpp

-// The test is disabled for now because of the flaky failure which can cause
-// Buildbot gives false negative results.
+// RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic
+// RUN: %libomptarget-compilexx-generic -O3 -ffast-math && \
+// RUN: %libomptarget-run-generic
+// RUN: %libomptarget-compileoptxx-generic -O3 && %libomptarget-run-generic
+// RUN: %libomptarget-compileoptxx-generic -O3 -ffast-math && \
+// RUN: %libomptarget-run-generic
 
-// RUN: nothing
-// REQUIRES: testToBeFixed
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// 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
+// RUN: %libomptarget-compileoptxx-generic && \
+// RUN:   env LIBOMPTARGET_STACK_SIZE=2048 %libomptarget-run-generic
 
+// We need malloc/global_alloc support
 // UNSUPPORTED: amdgcn-amd-amdhsa
 
 #include <cassert>
 #include <iostream>
 
 void work(int *C) {
 #pragma omp atomic
   ++(*C);

openmp/libomptarget/test/offloading/bug50022.cpp

 // RUN: %libomptarget-compilexx-and-run-generic
+// RUN: %libomptarget-compileoptxx-and-run-generic
 
 #include <cassert>
 #include <iostream>
 #include <stdexcept>
 
 int main(int argc, char *argv[]) {
   int a = 0;
   std::cout << "outside a = " << a << " addr " << &a << std::endl;

openmp/libomptarget/test/offloading/bug51982.c

 // RUN: %libomptarget-compile-generic -O1 && %libomptarget-run-generic
 // -O1 to run openmp-opt
+// RUN: %libomptarget-compileopt-generic -O1 && %libomptarget-run-generic
 
 int main(void) {
   long int aa = 0;
 
   int ng = 12;
   int nxyz = 5;
 
   const long exp = ng * nxyz;

openmp/libomptarget/test/offloading/bug53727.cpp

 // RUN: %libomptarget-compilexx-and-run-generic
+// RUN: %libomptarget-compileoptxx-and-run-generic
 
 #include <cassert>
 #include <iostream>
 
 constexpr const int N = 10;
 
 struct T {
   int a;

openmp/libomptarget/test/offloading/d2d_memcpy.c

-// RUN: %libomptarget-compile-generic && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | %fcheck-generic -allow-empty
+// RUN: %libomptarget-compile-generic && \
+// RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \
+// RUN: %fcheck-generic -allow-empty
+// RUN: %libomptarget-compileopt-generic && \
+// RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \
+// RUN: %fcheck-generic -allow-empty
 
 #include <assert.h>
 #include <omp.h>
 #include <stdio.h>
 #include <stdlib.h>
 
 const int magic_num = 7;
 

openmp/libomptarget/test/offloading/dynamic_module.c

-// RUN: %libomptarget-compile-generic -DSHARED -fPIC -shared -o %t.so && %libomptarget-compile-generic %t.so && %libomptarget-run-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-compile-generic -DSHARED -fPIC -shared -o %t.so && \
+// RUN: %libomptarget-compile-generic %t.so && %libomptarget-run-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-compileopt-generic -DSHARED -fPIC -shared -o %t.so && \
+// RUN: %libomptarget-compileopt-generic %t.so && %libomptarget-run-generic 2>&1 | %fcheck-generic
 
 #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/target_nowait_target.cpp

 // RUN: %libomptarget-compilexx-and-run-generic
-
-// UNSUPPORTED: amdgcn-amd-amdhsa
+// RUN: %libomptarget-compileoptxx-and-run-generic
 
 #include <cassert>
 
 int main(int argc, char *argv[]) {
   int data[1024];
   int sum = 0;
 
   for (int i = 0; i < 1024; ++i)