Hi!
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.
Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk.
2020-05-14 Jakub Jelinek <ja...@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.
--- gcc/function.h.jj 2020-01-18 13:47:09.506357871 +0100
+++ gcc/function.h 2020-05-13 12:46:03.424600627 +0200
@@ -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. */
--- gcc/omp-offload.c.jj 2020-05-13 15:13:53.156139627 +0200
+++ gcc/omp-offload.c 2020-05-13 15:07:01.591378851 +0200
@@ -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 *
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 (vo
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);
}
}
--- gcc/c/c-parser.c.jj 2020-04-19 12:13:12.914268947 +0200
+++ gcc/c/c-parser.c 2020-05-13 13:00:59.723008745 +0200
@@ -19866,6 +19866,7 @@ check_clauses:
}
pc = &OMP_CLAUSE_CHAIN (*pc);
}
+ cfun->has_omp_target = true;
return true;
}
--- gcc/cp/cp-gimplify.c.jj 2020-05-05 08:57:55.450641796 +0200
+++ gcc/cp/cp-gimplify.c 2020-05-13 14:04:23.255359955 +0200
@@ -1558,6 +1558,10 @@ cp_genericize_r (tree *stmt_p, int *walk
}
break;
+ case OMP_TARGET:
+ cfun->has_omp_target = true;
+ break;
+
case TRY_BLOCK:
{
*walk_subtrees = 0;
--- gcc/fortran/trans-openmp.c.jj 2020-05-13 11:17:40.688618996 +0200
+++ gcc/fortran/trans-openmp.c 2020-05-13 14:25:10.255453261 +0200
@@ -44,6 +44,7 @@ along with GCC; see the file COPYING3.
#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);
--- libgomp/testsuite/libgomp.c-c++-common/target-40.c.jj 2020-05-13
15:12:02.753813301 +0200
+++ libgomp/testsuite/libgomp.c-c++-common/target-40.c 2020-05-13
15:11:57.095899077 +0200
@@ -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;
+}
Jakub
