Hi!
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.
Bootstrapped/regtested on x86_64-linux and i686-linux and also tested with
x86_64-linux -> nvptx-none offloading, committed to trunk.
2020-05-12 Jakub Jelinek <ja...@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.
* testsuite/libgomp.c/target-39.c: New test.
--- gcc/omp-offload.h.jj 2020-01-15 11:05:19.315140331 +0100
+++ gcc/omp-offload.h 2020-05-11 19:45:04.752660397 +0200
@@ -30,5 +30,6 @@ extern GTY(()) vec<tree, va_gc> *offload
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 */
--- gcc/omp-offload.c.jj 2020-05-11 18:33:23.032680781 +0200
+++ gcc/omp-offload.c 2020-05-11 20:04:19.473126701 +0200
@@ -52,6 +52,7 @@ along with GCC; see the file COPYING3.
#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
}
}
+/* 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. */
--- gcc/cgraphunit.c.jj 2020-05-11 18:33:22.699685732 +0200
+++ gcc/cgraphunit.c 2020-05-11 19:45:04.754660367 +0200
@@ -206,6 +206,7 @@ along with GCC; see the file COPYING3.
#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)
--- libgomp/testsuite/libgomp.c/target-39.c.jj 2020-05-11 18:46:40.067800364
+0200
+++ libgomp/testsuite/libgomp.c/target-39.c 2020-05-11 18:46:40.067800364
+0200
@@ -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;
+}
Jakub
