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.

11 files changed, 128 insertions, 8 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index ec22de0..ce5ccab 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,17 @@
+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-05-14  Uroš Bizjak  <ubizjak@gmail.com>
 
 	PR target/95046
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 7cf774e..9ae494a 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,7 @@
+2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+	* c-parser.c (c_parser_omp_target): Set cfun->has_omp_target.
+
 2020-05-07  Richard Biener  <rguenther@suse.de>
 
 	PR middle-end/94703
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ae354e6..5c8502b 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -19866,6 +19866,7 @@ check_clauses:
 	  }
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
+  cfun->has_omp_target = true;
   return true;
 }
 
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index a45657d..5f5b31b 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,7 @@
+2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+	* cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target.
+
 2020-05-13  Patrick Palka  <ppalka@redhat.com>
 
 	PR c++/79706
diff --git a/gcc/cp/cp-gimplify.c b/gcc/cp/cp-gimplify.c
index fc26a85..ede9892 100644
--- a/gcc/cp/cp-gimplify.c
+++ b/gcc/cp/cp-gimplify.c
@@ -1558,6 +1558,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 b/gcc/fortran/ChangeLog
index becfda4..a3b673f 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,8 @@
+2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+	* trans-openmp.c: Include function.h.
+	(gfc_trans_omp_target): Set cfun->has_omp_target.
+
 2020-05-13  Steven G. Kargl  <kargl@gcc.gnu.org>
 
 	PR fortran/93497
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 8cf851e..e27ce41 100644
--- 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;
 
@@ -5392,6 +5393,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 1ee8ed3..d55cbdd 100644
--- 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 c1eb378..3e7012d 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -190,7 +190,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)
@@ -219,6 +219,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.  */
@@ -227,7 +245,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))
@@ -271,21 +289,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 b/libgomp/ChangeLog
index 104c527..4bfce69 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+	* testsuite/libgomp.c-c++-common/target-40.c: New test.
+
 2020-05-13  Tobias Burnus  <tobias@codesourcery.com>
 
 	PR fortran/94690
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
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;
+}