[Patch] OpenMP: Handle ancestor:1 with discover_declare_target

[Patch] OpenMP: Handle ancestor:1 with discover_declare_target

[Tobias Burnus]

[-- Attachment #1: Type: text/plain, Size: 744 bytes --]

Another minor step to getting reverse offloading to work ...

OK for mainline?

Tobias

PS: As attached, this patch is a stand-alone patch, which fails
due to the requires sorry (see dg-prune-output).
With the requires patch, it should fail with the next sorry:
   'ancestor' not yet supported"
(For Fortran, it currently fails already in the FE as the
sorry prevents generating tree code. Thus, there cannot be
a Fortran check until the requires patch is in.)
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Attachment #2: omp-reverse-tag.diff --]
[-- Type: text/x-patch, Size: 3941 bytes --]

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 ad4e772015e..fcbe6cf83d8 100644
--- 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
index 00000000000..9a3fa5230f8
--- /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;
+}

Re: [Patch] OpenMP: Handle ancestor:1 with discover_declare_target

[Jakub Jelinek]

On Thu, Jun 09, 2022 at 12:47:17PM +0200, Tobias Burnus wrote:
> Another minor step to getting reverse offloading to work ...
> 
> OK for mainline?
> 
> Tobias
> 
> PS: As attached, this patch is a stand-alone patch, which fails
> due to the requires sorry (see dg-prune-output).
> With the requires patch, it should fail with the next sorry:
>   'ancestor' not yet supported"
> (For Fortran, it currently fails already in the FE as the
> sorry prevents generating tree code. Thus, there cannot be
> a Fortran check until the requires patch is in.)
> -----------------
> Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

> 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.

LGTM, thanks.

	Jakub