[OpenMP] Allow more tests to run on AMDGPU

[jdoerfert]

Summary

This basically works around the printf issue to increase test coverage.

[jhuber6]

LG, might make sense to move these off the device anyway. Fewer things to fail. I'm assuming these all pass as you've written them.

[jdoerfert]

@doru1004 There might be others that need the same treatment, I think. Might be worth to check.

[doru1004]

In D146838#4222757, @jdoerfert wrote:

@doru1004 There might be others that need the same treatment, I think. Might be worth to check.

I plan to look at the other tests and re-enable them and see if something breaks.

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
-
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C1;
 #pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b[0 : 2])
int main() {
 
   D *sp = &s;
   D **spp = &sp;
 
   printf("%d %d %d %4.5f %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.c.a,
          spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0);
   // CHECK: 111 222 777 20.00000 1
 
+  int spp00fa = -1, spp00fca = -1, spp00fb_r = -1;
   __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
-#pragma omp target map(tofrom : spp[0][0]) firstprivate(p)
+#pragma omp target map(tofrom: spp[0][0]) firstprivate(p)                           \
+                   map(from: spp00fa, spp00fca, spp00fb_r)
   {
-    printf("%d %d %d\n", spp[0][0].f.a, spp[0][0].f.c.a,
-           spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
-    // CHECK: 222 777 0
+    spp00fa = spp[0][0].f.a;
+    spp00fca = spp[0][0].f.c.a;
+    spp00fb_r = spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0;
+    printf("%d %d %d\n", spp00fa, spp00fca, spp00fb_r);
+    // XCHECK: 222 777 0
     spp[0][0].e = 333;
     spp[0][0].f.a = 444;
     spp[0][0].f.c.a = 555;
     spp[0][0].f.b[1] = 40;
   }
+  printf("%d %d %d\n", spp00fa, spp00fca, spp00fb_r);
+  // CHECK: 222 777 0
   printf("%d %d %d %4.5f %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.c.a,
          spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0);
   // CHECK: 333 222 777 40.00000 1
 }

openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// Wrong results on amdgpu
-// XFAIL: amdgcn-amd-amdhsa
-
 #include <cstdio>
 #include <cstdlib>
 
 typedef struct {
   int a;
   double *b;
 } C;
 #pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b[0 : 2])
int main() {
   D *sp = &s;
   D **spp = &sp;
 
   printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
          spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
          spp[0][0].g == &y[0] ? 1 : 0);
   // CHECK: 111 222 20.00000 1 30 1
 
+  int spp00fa = -1, spp00fb_r = -1, spp00fg1 = -1, spp00fg_r = -1;
   __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]),
              p1 = reinterpret_cast<__intptr_t>(&y[0]);
-#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1)
+#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1)                  \
+                   map(from: spp00fa, spp00fb_r, spp00fg1, spp00fg_r)
   {
-    printf("%d %d %d %d\n", spp[0][0].f.a,
-           spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0, spp[0][0].g[1],
-           spp[0][0].g == reinterpret_cast<void *>(p1) ? 1 : 0);
-    // CHECK: 222 0 30 0
+    spp00fa = spp[0][0].f.a;
+    spp00fb_r = spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0;
+    spp00fg1 = spp[0][0].g[1];
+    spp00fg_r = spp[0][0].g == reinterpret_cast<void *>(p1) ? 1 : 0;
+    printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
+    // XCHECK: 222 0 30 0
     spp[0][0].e = 333;
     spp[0][0].f.a = 444;
     spp[0][0].f.b[1] = 40;
     spp[0][0].g[1] = 50;
   }
+    printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
+  // CHECK: 222 0 30 0
   printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
          spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
          spp[0][0].g == &y[0] ? 1 : 0);
   // CHECK: 333 222 40.00000 1 50 1
 }

openmp/libomptarget/test/mapping/lambda_by_value.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// Wrong results on amdgpu
-// XFAIL: amdgcn-amd-amdhsa
-
 #include <stdint.h>
 #include <stdio.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]]
 
+#pragma omp begin declare target
+int a = -1, *c;
+long b = -1;
+const long *d;
+int e = -1, *f, g = -1;
+#pragma omp end declare target
+
 int main() {
   int x[10];
   long y[8];
   x[1] = 111;
   y[1] = 222;
 
   auto lambda = [&x, y]() {
+    a = x[1];
+    b = y[1];
+    c = &x[0];
+    d = &y[0];
     printf("lambda: %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]);
     x[1] = y[1];
   };
-
   printf("before: %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]);
 
   intptr_t xp = (intptr_t)&x[0];
 #pragma omp target firstprivate(xp)
   {
     lambda();
+    e = x[1];
+    f = &x[0];
+    g = (&x[0] != (int *)xp);
     printf("tgt   : %d %p %d\n", x[1], &x[0], (&x[0] != (int *)xp));
   }
+#pragma omp target update from(a, b, c, d, e, f, g)
+  printf("lambda: %d %ld %p %p\n", a, b, c, d);
+  printf("tgt   : %d %p %d\n", e, f, g);
   printf("out   : %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]);
 
   return 0;
 }

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
-
 #include <omp.h>
 #include <stdio.h>
 
+#pragma omp begin declare target
+char *N1, *N2;
+int V1, V2;
+#pragma omp declare target
+
 #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()),             \
          omp_target_is_present(&(Var2), omp_get_default_device()),             \
          omp_target_is_present(&(Var3), omp_get_default_device()))
 
+#define CHECK_VALUES_HELPER(N1, N2, Var1, Var2)                                \
+  printf("    values of %s, %s: %d, %d\n", N1, N2, (Var1), (Var2))
+
+#define CHECK_VALUES_DELAYED(Var1, Var2)                                       \
+  N1 = #Var1;                                                                  \
+  N2 = #Var2;                                                                  \
+  V1 = (Var1);                                                                 \
+  V2 = (Var2);
+
+#define CHECK_DELAYED_VALUS()                                                  \
+  _Pragma("omp target update from(N1, N2, V1, V2)")                            \
+      CHECK_VALUES_HELPER(N1, N2, V1, V2)
+
 #define CHECK_VALUES(Var1, Var2)                                               \
-  printf("    values of %s, %s: %d, %d\n", #Var1, #Var2, (Var1), (Var2))
+  CHECK_VALUES_HELPER(#Var1, #Var2, (Var1), (Var2))
 
 int main() {
   struct S {
     int i;
     int j;
   } s;
   // CHECK: presence of s, s.i, s.j: 0, 0, 0
   CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(from : s.i, s.j)
     } // No transfer here even though parent's DynRefCount=1.
     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
     // CHECK-NEXT: values of s.i, s.j: 20, 30
     CHECK_PRESENCE(s, s.i, s.j);
     CHECK_VALUES(s.i, s.j);
 #pragma omp target map(to : s.i, s.j)
     { // No transfer here even though parent's DynRefCount=1.
       // CHECK-NEXT: values of s.i, s.j: 21, 31
-      CHECK_VALUES(s.i, s.j);
+      CHECK_VALUES_DELAYED(s.i, s.j);
     }
+    CHECK_DELAYED_VALUS();
   }
   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
   // CHECK-NEXT: values of s.i, s.j: 21, 31
   CHECK_PRESENCE(s, s.i, s.j);
   CHECK_VALUES(s.i, s.j);
 
   // -----------------------------------------------------------------------
   // CHECK-LABEL: check:{{.*}}
#pragma omp target map(ompx_hold, from : s.i, s.j)
     } // No transfer here even though parent's HoldRefCount=1.
     // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
     // CHECK-NEXT: values of s.i, s.j: 20, 30
     CHECK_PRESENCE(s, s.i, s.j);
     CHECK_VALUES(s.i, s.j);
 #pragma omp target map(ompx_hold, to : s.i, s.j)
     { // No transfer here even though parent's HoldRefCount=1.
       // CHECK-NEXT: values of s.i, s.j: 21, 31
-      CHECK_VALUES(s.i, s.j);
+      CHECK_VALUES_DELAYED(s.i, s.j);
     }
+    CHECK_DELAYED_VALUS();
   }
   // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
   // CHECK-NEXT: values of s.i, s.j: 21, 31
   CHECK_PRESENCE(s, s.i, s.j);
   CHECK_VALUES(s.i, s.j);
 
   // -----------------------------------------------------------------------
   // CHECK-LABEL: check:{{.*}}

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
-
 #include <omp.h>
 #include <stdio.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);
+  char DV = -1;
+#pragma omp target device(Dev) map(from : DV)
+  DV = *X;
+  printf("device X = %c\n", DV);
 }
 
 #define CHECK_DATA() check(&X, DevDefault)
 
 int main(void) {
   int DevDefault = omp_get_default_device();
   int DevInit = omp_get_initial_device();