openmp: Also implicitly mark as declare target to functions mentioned in target regions

OpenMP 5.0 also specifies that functions referenced from target regions
(except for target regions with device(ancestor:)) are also implicitly declare target to.

This patch implements that.

2020-05-14  Jakub Jelinek  <jakub@redhat.com>

	* function.h (struct function): Add has_omp_target bit.
	* omp-offload.c (omp_discover_declare_target_fn_r): New function,
	old renamed to ...
	(omp_discover_declare_target_tgt_fn_r): ... this.
	(omp_discover_declare_target_var_r): Call
	omp_discover_declare_target_tgt_fn_r instead of
	omp_discover_declare_target_fn_r.
	(omp_discover_implicit_declare_target): Also queue functions with
	has_omp_target bit set, for those walk with
	omp_discover_declare_target_fn_r, for declare target to functions
	walk with omp_discover_declare_target_tgt_fn_r.
gcc/c/
	* c-parser.c (c_parser_omp_target): Set cfun->has_omp_target.
gcc/cp/
	* cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target.
gcc/fortran/
	* trans-openmp.c: Include function.h.
	(gfc_trans_omp_target): Set cfun->has_omp_target.
libgomp/
	* testsuite/libgomp.c-c++-common/target-40.c: New test.

(cherry picked from commit 49ddde69fc8e1e4c48d4b0027ea37ac862da0f1f)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 448a8c91ab46f12c3e2d35ca8abcb01d4582647a..63ccd8a3a797d90ec38c8ca04e99ff29d4513945 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,20 @@
+2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline
+       2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+       * function.h (struct function): Add has_omp_target bit.
+       * omp-offload.c (omp_discover_declare_target_fn_r): New function,
+       old renamed to ...
+       (omp_discover_declare_target_tgt_fn_r): ... this.
+       (omp_discover_declare_target_var_r): Call
+       omp_discover_declare_target_tgt_fn_r instead of
+       omp_discover_declare_target_fn_r.
+       (omp_discover_implicit_declare_target): Also queue functions with
+       has_omp_target bit set, for those walk with
+       omp_discover_declare_target_fn_r, for declare target to functions
+       walk with omp_discover_declare_target_tgt_fn_r.
+
 2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline

diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index 4a1630aed6343ab76b4e703a5a7c3493d68230ec..e520f95fd8993ee2c2ad6a3814ea5a9c683045dc 100644 (file)
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,10 @@
+2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline
+       2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+       * c-parser.c (c_parser_omp_target): Set cfun->has_omp_target.
+
 2020-09-16  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 988b6bd4091f5d54b7dfe89bd08856647586e56d..451e864c0109a8ab5ae66c3788003f9579566bf7 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -19878,6 +19878,7 @@ check_clauses:
          }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
+  cfun->has_omp_target = true;
   return true;
 }
 

diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index a9a94b3696a00bf1f4d68ec95405ee4dd4d87f00..1aaed1b8beae56724d0ad2559ef0af78ded368dd 100644 (file)
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,10 @@
+2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline
+       2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+       * cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target.
+
 2020-09-16  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline

diff --git a/gcc/cp/cp-gimplify.c b/gcc/cp/cp-gimplify.c
index 20648b910dd09b1c72b859303372fad94c97c647..e78dd3ca0c0afb8eba1ac5ac055e260b52f4163e 100644 (file)
--- a/gcc/cp/cp-gimplify.c
+++ b/gcc/cp/cp-gimplify.c
@@ -1278,6 +1278,10 @@ cp_genericize_r (tree *stmt_p, int *walk_subtrees, void *data)
       }
       break;
 
+    case OMP_TARGET:
+      cfun->has_omp_target = true;
+      break;
+
     case TRY_BLOCK:
       {
         *walk_subtrees = 0;

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index d35557a2503badddcb5f58cc65c520938c411f6d..e12c7ac4e9c71cbe9cc6bc2427da5413cb207850 100644 (file)
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,11 @@
+2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline
+       2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+       * trans-openmp.c: Include function.h.
+       (gfc_trans_omp_target): Set cfun->has_omp_target.
+
 2020-09-16  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 35cd2a83d4ab1fd7bb23d9a69e7545dcfc0b216e..77ecb957dfd805552a1948773e30295c3ef3bbc4 100644 (file)
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -44,6 +44,7 @@ along with GCC; see the file COPYING3.  If not see
 #undef GCC_DIAG_STYLE
 #define GCC_DIAG_STYLE __gcc_gfc__
 #include "attribs.h"
+#include "function.h"
 
 int ompws_flags;
 
@@ -5710,6 +5711,7 @@ gfc_trans_omp_target (gfc_code *code)
                         omp_clauses);
       if (code->op != EXEC_OMP_TARGET)
        OMP_TARGET_COMBINED (stmt) = 1;
+      cfun->has_omp_target = true;
     }
   gfc_add_expr_to_block (&block, stmt);
   return gfc_finish_block (&block);

diff --git a/gcc/function.h b/gcc/function.h
index 1ee8ed3de53f8d3c637ad5640d6360a0c66b3342..d55cbddd0b51f78cf272f9f660d5ee0fd006d5d6 100644 (file)
--- a/gcc/function.h
+++ b/gcc/function.h
@@ -421,6 +421,9 @@ struct GTY(()) function {
 
   /* Set if this is a coroutine-related function.  */
   unsigned int coroutine_component : 1;
+
+  /* Set if there are any OMP_TARGET regions in the function.  */
+  unsigned int has_omp_target : 1;
 };
 
 /* Add the decl D to the local_decls list of FUN.  */

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 2281b6747abae4fbcf237bc2502bde6843d99b40..9c0ca643a2ff25db859ab918ac18a7ac607e29dd 100644 (file)
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,7 +196,7 @@ omp_declare_target_var_p (tree decl)
    declare target to.  */
 
 static tree
-omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
+omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data)
 {
   if (TREE_CODE (*tp) == FUNCTION_DECL
       && !omp_declare_target_fn_p (*tp)
@@ -225,6 +225,24 @@ omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data)
   return NULL_TREE;
 }
 
+/* Similarly, but ignore references outside of OMP_TARGET regions.  */
+
+static tree
+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);
+      *walk_subtrees = 0;
+    }
+  else if (TYPE_P (*tp))
+    *walk_subtrees = 0;
+  return NULL_TREE;
+}
+
 /* Helper function for omp_discover_implicit_declare_target, called through
    walk_tree.  Mark referenced FUNCTION_DECLs implicitly as
    declare target to.  */
@@ -233,7 +251,7 @@ static tree
 omp_discover_declare_target_var_r (tree *tp, int *walk_subtrees, void *data)
 {
   if (TREE_CODE (*tp) == FUNCTION_DECL)
-    return omp_discover_declare_target_fn_r (tp, walk_subtrees, data);
+    return omp_discover_declare_target_tgt_fn_r (tp, walk_subtrees, data);
   else if (VAR_P (*tp)
           && is_global_var (*tp)
           && !omp_declare_target_var_p (*tp))
@@ -277,21 +295,31 @@ omp_discover_implicit_declare_target (void)
   auto_vec<tree> worklist;
 
   FOR_EACH_DEFINED_FUNCTION (node)
-    if (omp_declare_target_fn_p (node->decl) && DECL_SAVED_TREE (node->decl))
-      worklist.safe_push (node->decl);
+    if (DECL_SAVED_TREE (node->decl))
+      {
+        if (omp_declare_target_fn_p (node->decl))
+         worklist.safe_push (node->decl);
+       else if (DECL_STRUCT_FUNCTION (node->decl)
+                && DECL_STRUCT_FUNCTION (node->decl)->has_omp_target)
+         worklist.safe_push (node->decl);
+      }
   FOR_EACH_STATIC_INITIALIZER (vnode)
     if (omp_declare_target_var_p (vnode->decl))
       worklist.safe_push (vnode->decl);
   while (!worklist.is_empty ())
     {
       tree decl = worklist.pop ();
-      if (TREE_CODE (decl) == FUNCTION_DECL)
+      if (VAR_P (decl))
+       walk_tree_without_duplicates (&DECL_INITIAL (decl),
+                                     omp_discover_declare_target_var_r,
+                                     &worklist);
+      else if (omp_declare_target_fn_p (decl))
        walk_tree_without_duplicates (&DECL_SAVED_TREE (decl),
-                                     omp_discover_declare_target_fn_r,
+                                     omp_discover_declare_target_tgt_fn_r,
                                      &worklist);
       else
-       walk_tree_without_duplicates (&DECL_INITIAL (decl),
-                                     omp_discover_declare_target_var_r,
+       walk_tree_without_duplicates (&DECL_SAVED_TREE (decl),
+                                     omp_discover_declare_target_fn_r,
                                      &worklist);
     }
 }

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index c88312c5015a43d6d6d45bb3dc8eb6e7b22b069c..82e7cc2d941d572a116e3c92b120bf85e947d3e3 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,10 @@
+2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline
+       2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+       * testsuite/libgomp.c-c++-common/target-40.c: New test.
+
 2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline

diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-40.c b/libgomp/testsuite/libgomp.c-c++-common/target-40.c
new file mode 100644 (file)
index 0000000..22bbdd9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-40.c
@@ -0,0 +1,51 @@
+/* { dg-do run } */
+/* { dg-options "-O0" } */
+
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+volatile int v;
+#pragma omp declare target to (v)
+typedef void (*fnp1) (void);
+typedef fnp1 (*fnp2) (void);
+void f1 (void) { v++; }
+void f2 (void) { v += 4; }
+void f3 (void) { v += 16; f1 (); }
+fnp1 f4 (void) { v += 64; return f2; }
+int a = 1;
+int *b = &a;
+int **c = &b;
+fnp2 f5 (void) { f3 (); return f4; }
+#pragma omp declare target to (c)
+
+int
+main ()
+{
+  int err = 0;
+  #pragma omp target map(from:err)
+  {
+    volatile int xa;
+    int *volatile xb;
+    int **volatile xc;
+    fnp2 xd;
+    fnp1 xe;
+    err = 0;
+    xa = a;
+    err |= xa != 1;
+    xb = b;
+    err |= xb != &a;
+    xc = c;
+    err |= xc != &b;
+    xd = f5 ();
+    err |= v != 17;
+    xe = xd ();
+    err |= v != 81;
+    xe ();
+    err |= v != 85;
+  }
+  if (err)
+    abort ();
+  return 0;
+}