OpenMP: Handle ancestor:1 with discover_declare_target

gcc/
	* omp-offload.cc (omp_discover_declare_target_tgt_fn_r,
	omp_discover_declare_target_fn_r): Don't walk reverse-offload
	target regions.

gcc/testsuite/
	* c-c++-common/gomp/reverse-offload-1.c: New.

diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index ad4e772..fcbe6cf 100644 (file)
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -268,12 +268,12 @@ omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data)
     }
   else if (TYPE_P (*tp))
     *walk_subtrees = 0;
-  /* else if (TREE_CODE (*tp) == OMP_TARGET)
-       {
-        if (tree dev = omp_find_clause (OMP_TARGET_CLAUSES (*tp)))
-          if (OMP_DEVICE_ANCESTOR (dev))
-            *walk_subtrees = 0;
-       } */
+  else if (TREE_CODE (*tp) == OMP_TARGET)
+    {
+      tree c = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE);
+      if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c))
+       *walk_subtrees = 0;
+    }
   return NULL_TREE;
 }
 
@@ -284,10 +284,11 @@ omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
 {
   if (TREE_CODE (*tp) == OMP_TARGET)
     {
-      /* And not OMP_DEVICE_ANCESTOR.  */
-      walk_tree_without_duplicates (&OMP_TARGET_BODY (*tp),
-                                   omp_discover_declare_target_tgt_fn_r,
-                                   data);
+      tree c = omp_find_clause (OMP_CLAUSES (*tp), OMP_CLAUSE_DEVICE);
+      if (!c || !OMP_CLAUSE_DEVICE_ANCESTOR (c))
+       walk_tree_without_duplicates (&OMP_TARGET_BODY (*tp),
+                                     omp_discover_declare_target_tgt_fn_r,
+                                     data);
       *walk_subtrees = 0;
     }
   else if (TYPE_P (*tp))

diff --git a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
new file mode 100644 (file)
index 0000000..9a3fa52
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
@@ -0,0 +1,93 @@
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+/* { dg-final { scan-tree-dump-times "omp declare target\[^ \]" 3 "omplower" } }  */
+
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)\[\n\r\]*int called_in_target1" 1 "omplower" } }  */
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)\[\n\r\]*int called_in_target2" 1 "omplower" } }  */
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target, omp declare target block\\)\\)\[\n\r\]*void tg_fn" 1 "omplower" } }  */
+
+/* { dg-prune-output "'reverse_offload' clause on 'requires' directive not supported yet" } */
+
+#pragma omp requires reverse_offload
+
+extern int add_3 (int);
+
+static int global_var = 5;
+
+void
+check_offload (int *x, int *y)
+{
+  *x = add_3 (*x);
+  *y = add_3 (*y);
+}
+
+int
+called_in_target1 ()
+{
+  return 42;
+}
+
+int
+called_in_target2 ()
+{
+  return -6;
+}
+
+#pragma omp declare target
+void
+tg_fn (int *x, int *y)
+{
+  int x2 = *x, y2 = *y;
+  if (x2 != 2 || y2 != 3)
+    __builtin_abort ();
+  x2 = x2 + 2 + called_in_target1 ();
+  y2 = y2 + 7;
+
+  #pragma omp target device(ancestor : 1) map(tofrom: x2)
+    check_offload(&x2, &y2);
+
+  if (x2 != 2+2+3+42 || y2 != 3 + 7)
+    __builtin_abort ();
+  *x = x2, *y = y2;
+}
+#pragma omp end declare target
+
+void
+my_func (int *x, int *y)
+{
+  if (global_var != 5)
+    __builtin_abort ();
+  global_var = 242;
+  *x = 2*add_3(*x);
+  *y = 3*add_3(*y);
+}
+
+int
+main ()
+{
+  #pragma omp target
+  {
+     int x = 2, y = 3;
+     tg_fn (&x, &y);
+  }
+
+  #pragma omp target
+  {
+     int x = -2, y = -1;
+     x += called_in_target2 ();
+     #pragma omp target device ( ancestor:1 ) firstprivate(y) map(tofrom:x)
+     {
+       if (x != -2-6 || y != -1)
+         __builtin_abort ();
+       my_func (&x, &y);
+       if (x != 2*(3-2) || y != 3*(3-1))
+         __builtin_abort ();
+     }
+     if (x != 2*(3-2) || y != -1)
+       __builtin_abort ();
+  }
+
+  if (global_var != 242)
+    __builtin_abort ();
+  return 0;
+}