diff options
| -rw-r--r-- | gcc/ChangeLog | 10 | ||||
| -rw-r--r-- | gcc/cgraphunit.c | 4 | ||||
| -rw-r--r-- | gcc/omp-offload.c | 133 | ||||
| -rw-r--r-- | gcc/omp-offload.h | 1 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 4 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/target-39.c | 47 |
6 files changed, 199 insertions, 0 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index bd84f8f73f0..ffa00559387 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,13 @@ +2020-05-12 Jakub Jelinek <jakub@redhat.com> + + * omp-offload.h (omp_discover_implicit_declare_target): Declare. + * omp-offload.c: Include context.h. + (omp_declare_target_fn_p, omp_declare_target_var_p, + omp_discover_declare_target_fn_r, omp_discover_declare_target_var_r, + omp_discover_implicit_declare_target): New functions. + * cgraphunit.c (analyze_functions): Call + omp_discover_implicit_declare_target. + 2020-05-12 Richard Biener <rguenther@suse.de> * gimple-fold.c (maybe_canonicalize_mem_ref_addr): Canonicalize diff --git a/gcc/cgraphunit.c b/gcc/cgraphunit.c index 0563932a709..01b3f82a4b2 100644 --- a/gcc/cgraphunit.c +++ b/gcc/cgraphunit.c @@ -206,6 +206,7 @@ along with GCC; see the file COPYING3. If not see #include "stringpool.h" #include "attribs.h" #include "ipa-inline.h" +#include "omp-offload.h" /* Queue of cgraph nodes scheduled to be added into cgraph. This is a secondary queue used during optimization to accommodate passes that @@ -1160,6 +1161,9 @@ analyze_functions (bool first_time) node->fixup_same_cpp_alias_visibility (node->get_alias_target ()); build_type_inheritance_graph (); + if (flag_openmp && first_time) + omp_discover_implicit_declare_target (); + /* Analysis adds static variables that in turn adds references to new functions. So we need to iterate the process until it stabilize. */ while (changed) diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index c66f38b6f0c..c1eb378e2a1 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -52,6 +52,7 @@ along with GCC; see the file COPYING3. If not see #include "stringpool.h" #include "attribs.h" #include "cfgloop.h" +#include "context.h" /* Describe the OpenACC looping structure of a function. The entire function is held in a 'NULL' loop. */ @@ -158,6 +159,138 @@ add_decls_addresses_to_decl_constructor (vec<tree, va_gc> *v_decls, } } +/* Return true if DECL is a function for which its references should be + analyzed. */ + +static bool +omp_declare_target_fn_p (tree decl) +{ + return (TREE_CODE (decl) == FUNCTION_DECL + && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) + && !lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (decl)) + && (!flag_openacc + || oacc_get_fn_attrib (decl) == NULL_TREE)); +} + +/* Return true if DECL Is a variable for which its initializer references + should be analyzed. */ + +static bool +omp_declare_target_var_p (tree decl) +{ + return (VAR_P (decl) + && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) + && !lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (decl))); +} + +/* Helper function for omp_discover_implicit_declare_target, called through + walk_tree. Mark referenced FUNCTION_DECLs implicitly as + declare target to. */ + +static tree +omp_discover_declare_target_fn_r (tree *tp, int *walk_subtrees, void *data) +{ + if (TREE_CODE (*tp) == FUNCTION_DECL + && !omp_declare_target_fn_p (*tp) + && !lookup_attribute ("omp declare target host", DECL_ATTRIBUTES (*tp))) + { + tree id = get_identifier ("omp declare target"); + if (!DECL_EXTERNAL (*tp) && DECL_SAVED_TREE (*tp)) + ((vec<tree> *) data)->safe_push (*tp); + DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp)); + symtab_node *node = symtab_node::get (*tp); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + g->have_offload = true; + } + } + else if (TYPE_P (*tp)) + *walk_subtrees = 0; + /* else if (TREE_CODE (*tp) == OMP_TARGET) + { + if (tree dev = omp_find_clause (OMP_TARGET_CLAUSES (*tp))) + if (OMP_DEVICE_ANCESTOR (dev)) + *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. */ + +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); + else if (VAR_P (*tp) + && is_global_var (*tp) + && !omp_declare_target_var_p (*tp)) + { + tree id = get_identifier ("omp declare target"); + if (lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (*tp))) + { + error_at (DECL_SOURCE_LOCATION (*tp), + "%qD specified both in declare target %<link%> and " + "implicitly in %<to%> clauses", *tp); + DECL_ATTRIBUTES (*tp) + = remove_attribute ("omp declare target link", DECL_ATTRIBUTES (*tp)); + } + if (TREE_STATIC (*tp) && DECL_INITIAL (*tp)) + ((vec<tree> *) data)->safe_push (*tp); + DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp)); + symtab_node *node = symtab_node::get (*tp); + if (node != NULL && !node->offloadable) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a <varpool_node *> (node)) + vec_safe_push (offload_vars, node->decl); + } + } + } + else if (TYPE_P (*tp)) + *walk_subtrees = 0; + return NULL_TREE; +} + +/* Perform the OpenMP implicit declare target to discovery. */ + +void +omp_discover_implicit_declare_target (void) +{ + cgraph_node *node; + varpool_node *vnode; + 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); + 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) + walk_tree_without_duplicates (&DECL_SAVED_TREE (decl), + omp_discover_declare_target_fn_r, + &worklist); + else + walk_tree_without_duplicates (&DECL_INITIAL (decl), + omp_discover_declare_target_var_r, + &worklist); + } +} + + /* Create new symbols containing (address, size) pairs for global variables, marked with "omp declare target" attribute, as well as addresses for the functions, which are outlined offloading regions. */ diff --git a/gcc/omp-offload.h b/gcc/omp-offload.h index 6adc57663fb..0809189db25 100644 --- a/gcc/omp-offload.h +++ b/gcc/omp-offload.h @@ -30,5 +30,6 @@ extern GTY(()) vec<tree, va_gc> *offload_funcs; extern GTY(()) vec<tree, va_gc> *offload_vars; extern void omp_finish_file (void); +extern void omp_discover_implicit_declare_target (void); #endif /* GCC_OMP_DEVICE_H */ diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index b6828adcbe3..1265640a2c3 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2020-05-12 Jakub Jelinek <jakub@redhat.com> + + * testsuite/libgomp.c/target-39.c: New test. + 2020-04-29 Thomas Schwinge <thomas@codesourcery.com> * config/accel/openacc.f90 (acc_device_current): Set to '-1'. diff --git a/libgomp/testsuite/libgomp.c/target-39.c b/libgomp/testsuite/libgomp.c/target-39.c new file mode 100644 index 00000000000..4442f43c8ef --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-39.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-options "-O0" } */ + +extern 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, f5) + +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; +} |