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.

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index ec22de0568d46c20b97c4f11929c6635743c208e..ce5ccabfcef7e17f3c5dcbc607cd718b4530795f 100644 (file)
--- 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 7cf774ed171d76856a358c7aca3b35d924621e3c..9ae494a45e7a79622fb3a50ebf8ba1119d0cd7b2 100644 (file)
--- 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 ae354e6af66ac0297133143b6a250961463adb00..5c8502bcceb270435e231082a489371bccfe98a1 100644 (file)
--- 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 a45657d09ade092ddd791a63a1881a32572ee539..5f5b31b24748224edacb5b76a54d563f458472ca 100644 (file)
--- 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 fc26a85f43aa00eda4f13799c9f0f4dae3f45e88..ede98929179d29d0a3a6aa4f4a942865e098dc5d 100644 (file)
--- 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 becfda4c5fa684ec986d3c58d5d3d2f9e63ec3bc..a3b673f7ed9cac274b0a6246bcce5c856c010476 100644 (file)
--- 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 8cf851eb6e30fdfd4d08612d1d7951200eaad3e9..e27ce41b7ce984457231a2e3bb60f620c64fcccf 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;
 
@@ -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 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 c1eb378e2a1c5795ac2d97c7b77321c429d2ff75..3e7012d649f7a6f17c2b3bfa4c95df387b7941d7 100644 (file)
--- 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 104c527f9c7290ea50a1c3db5886c991510a3f39..4bfce6920bfcda1b933f2ee91bc0f5d2cb671c5e 100644 (file)
--- 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 (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;
+}