From dc703151d4f4560e647649506d5b4ceb0ee11e90 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Tue, 12 May 2020 09:17:09 +0200 Subject: [PATCH] openmp: Implement discovery of implicit declare target to clauses This attempts to implement what the OpenMP 5.0 spec in declare target section says as ammended by the 5.1 changes so far (related to device_type(host)), except that it doesn't have the device(ancestor: ...) handling yet because we do not support it yet, and I've left so far out the except lambda note, because I need that clarified. 2020-05-12 Jakub Jelinek * 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. * testsuite/libgomp.c/target-39.c: New test. --- gcc/ChangeLog | 10 ++ gcc/cgraphunit.c | 4 + gcc/omp-offload.c | 133 ++++++++++++++++++++++++ gcc/omp-offload.h | 1 + libgomp/ChangeLog | 4 + libgomp/testsuite/libgomp.c/target-39.c | 47 +++++++++ 6 files changed, 199 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/target-39.c 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 + + * 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 * 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 *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 *) 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 % and " + "implicitly in % clauses", *tp); + DECL_ATTRIBUTES (*tp) + = remove_attribute ("omp declare target link", DECL_ATTRIBUTES (*tp)); + } + if (TREE_STATIC (*tp) && DECL_INITIAL (*tp)) + ((vec *) 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 (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 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 *offload_funcs; extern GTY(()) vec *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 + + * testsuite/libgomp.c/target-39.c: New test. + 2020-04-29 Thomas Schwinge * 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; +} -- 2.30.2