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  <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.

	* testsuite/libgomp.c/target-39.c: New test.

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index bd84f8f73f07e4f8e1b3049bb3e17b057638948c..ffa00559387be1289f765067c08e8216a4272c5b 100644 (file)
--- 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 0563932a7096492c367641ca646324331f042909..01b3f82a4b2a43d3c1b5153113613637cedf719a 100644 (file)
--- 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 c66f38b6f0cd3d16cba366a1ae5c22968fdce15c..c1eb378e2a1c5795ac2d97c7b77321c429d2ff75 100644 (file)
--- 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 6adc57663fbd7c0bc170e6debe0e4e8d413ca668..0809189db2556fb82c312ff419ae55addd3f3670 100644 (file)
--- 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 b6828adcbe3d1c991e6fde3128ad0a8f0ada473c..1265640a2c36fecd06f26ad4acf0526ba9e57278 100644 (file)
--- 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 (file)
index 0000000..4442f43
--- /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;
+}