[OpenMP] Enable more tests

[jdoerfert]

Summary

This enables tests to run on AMDGPU and also enables optimization for the compilation.

[jdoerfert]

They all pass for me on an X86 + gfx90a, can people try them on different setups and report results here, plz.

[ronlieb]

@doru1004 is interested in these tests being enabled.
thx for adding them.

[jdoerfert]

In D146927#4222736, @ronlieb wrote:

@doru1004 is interested in these tests being enabled.
thx for adding them.

@ronlieb Can you test this commit on the builedbot too?

[ronlieb]

i applied the patch to a gfx908 and a gfx906 system.
tests all passed , and the gfx906 should be the same as the amdgpu buildbot.

one odd thing i ran into was that the raw patch did not apply to
openmp/libomptarget/test/offloading/barrier_fence.c

perhaps it just needs rebase ?

any way, seems like its worth landing after you get some feedback from other platforms.
or if none seems forthcoming ,land it. so i gave it the old LGTM

[jdoerfert]

one odd thing i ran into was that the raw patch did not apply to
openmp/libomptarget/test/offloading/barrier_fence.c

That one is included in the fence patch. I will ensure it works once I upstream.

[jdoerfert]

Subsumed by D148576

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: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
 #include <omp.h>
 #include <stdio.h>
 
 struct IdentTy;

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)