[OpenMP] Allow more tests to run on AMDGPU

jdoerfert · jdoerfert · commit 747af2415519 · 2023-03-27T21:30:22.000-07:00
This basically works around the printf issue to increase test coverage. Differential Revision: https://reviews.llvm.org/D146838
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
@@ -1,8 +1,5 @@
 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// Wrong results on amdgpu
-// XFAIL: amdgcn-amd-amdhsa
-
 #include <cstdio>
 #include <cstdlib>
 
@@ -45,17 +42,23 @@ int main() {
          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
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
@@ -1,8 +1,5 @@
 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// Wrong results on amdgpu
-// XFAIL: amdgcn-amd-amdhsa
-
 #include <cstdio>
 #include <cstdlib>
 
@@ -42,19 +39,25 @@ int main() {
          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);
diff --git a/openmp/libomptarget/test/mapping/lambda_by_value.cpp b/openmp/libomptarget/test/mapping/lambda_by_value.cpp
@@ -1,8 +1,5 @@
 // RUN: %libomptarget-compilexx-run-and-check-generic
 
-// Wrong results on amdgpu
-// XFAIL: amdgcn-amd-amdhsa
-
 #include <stdint.h>
 #include <stdio.h>
 
@@ -11,25 +8,41 @@
 // 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;
diff --git a/openmp/libomptarget/test/mapping/ompx_hold/struct.c b/openmp/libomptarget/test/mapping/ompx_hold/struct.c
@@ -1,20 +1,35 @@
 // 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 {
@@ -132,8 +147,9 @@ int main() {
 #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
@@ -162,8 +178,9 @@ int main() {
 #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
diff --git a/openmp/libomptarget/test/offloading/host_as_target.c b/openmp/libomptarget/test/offloading/host_as_target.c
@@ -7,16 +7,15 @@
 
 // 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)
