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.

Author: jakubjelinek

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

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)

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

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 */

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

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;
+}