From: Jakub Jelinek Date: Thu, 14 May 2020 07:48:32 +0000 (+0200) Subject: openmp: Also implicitly mark as declare target to functions mentioned in target regions X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=49ddde69fc8e1e4c48d4b0027ea37ac862da0f1f;p=gcc.git 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 * 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 ec22de0568d..ce5ccabfcef 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,17 @@ +2020-05-14 Jakub Jelinek + + * 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 PR target/95046 diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 7cf774ed171..9ae494a45e7 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,7 @@ +2020-05-14 Jakub Jelinek + + * c-parser.c (c_parser_omp_target): Set cfun->has_omp_target. + 2020-05-07 Richard Biener PR middle-end/94703 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index ae354e6af66..5c8502bcceb 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 a45657d09ad..5f5b31b2474 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,7 @@ +2020-05-14 Jakub Jelinek + + * cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target. + 2020-05-13 Patrick Palka PR c++/79706 diff --git a/gcc/cp/cp-gimplify.c b/gcc/cp/cp-gimplify.c index fc26a85f43a..ede98929179 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 becfda4c5fa..a3b673f7ed9 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,8 @@ +2020-05-14 Jakub Jelinek + + * trans-openmp.c: Include function.h. + (gfc_trans_omp_target): Set cfun->has_omp_target. + 2020-05-13 Steven G. Kargl PR fortran/93497 diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 8cf851eb6e3..e27ce41b7ce 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 1ee8ed3de53..d55cbddd0b5 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 c1eb378e2a1..3e7012d649f 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 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 104c527f9c7..4bfce6920bf 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2020-05-14 Jakub Jelinek + + * testsuite/libgomp.c-c++-common/target-40.c: New test. + 2020-05-13 Tobias Burnus 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 00000000000..22bbdd97b7e --- /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; +}