commit d2c8c5bd2826851b727e93a8ea2141596e50a621
Author: Kwok Cheung Yeung <kcy@codesourcery.com>
Date:   Wed Oct 28 07:13:14 2020 -0700

    openmp: Implicitly add 'declare target' directives for dynamic static initializers in C++
    
    2020-10-28  Kwok Cheung Yeung  <kcy@codesourcery.com>
    
    	cp/
    	* decl2.c: Include omp-offload.h
    	(c_parse_final_cleanups): Call omp_mark_target_static_initializers.
    	* omp-offload.c (omp_discover_declare_target_var_r): Add all static
    	variables to worklist.
    	(omp_discover_implicit_declare_target): Check that worklist items
    	that are variable declarations have an initialization expression
    	before walking.
    	(omp_mark_target_static_initializers): New.
    	* omp-offload.h	(omp_mark_target_static_initializers): New prototype.
    
    	libgomp/
    	* testsuite/libgomp.c++/declare_target-3.C: New.

diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c
index 2f0d637..b207d58 100644
--- a/gcc/cp/decl2.c
+++ b/gcc/cp/decl2.c
@@ -48,6 +48,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "intl.h"
 #include "c-family/c-ada-spec.h"
 #include "asan.h"
+#include "omp-offload.h"
 
 /* Id for dumping the raw trees.  */
 int raw_dump_id;
@@ -4970,6 +4971,12 @@ c_parse_final_cleanups (void)
 	  /* Make sure the back end knows about all the variables.  */
 	  write_out_vars (vars);
 
+	  /* Mark functions and variables in static initializers as
+	     'omp declare target' if the initialized variable is marked
+	     as such.  */
+	  if (flag_openmp)
+	    omp_mark_target_static_initializers (vars);
+
 	  /* Set the line and file, so that it is obviously not from
 	     the source file.  */
 	  input_location = locus_at_end_of_parsing;
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 3e9c31d..8ecc181 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -296,7 +296,7 @@ omp_discover_declare_target_var_r (tree *tp, int *walk_subtrees, void *data)
 	  DECL_ATTRIBUTES (*tp)
 	    = remove_attribute ("omp declare target link", DECL_ATTRIBUTES (*tp));
 	}
-      if (TREE_STATIC (*tp) && DECL_INITIAL (*tp))
+      if (TREE_STATIC (*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);
@@ -348,7 +348,7 @@ omp_discover_implicit_declare_target (void)
   while (!worklist.is_empty ())
     {
       tree decl = worklist.pop ();
-      if (VAR_P (decl))
+      if (VAR_P (decl) && DECL_INITIAL (decl))
 	walk_tree_without_duplicates (&DECL_INITIAL (decl),
 				      omp_discover_declare_target_var_r,
 				      &worklist);
@@ -363,6 +363,33 @@ omp_discover_implicit_declare_target (void)
     }
 }
 
+void
+omp_mark_target_static_initializers (tree vars)
+{
+  tree node;
+  auto_vec<tree> worklist;
+
+  for (node = vars; node; node = TREE_CHAIN (node))
+    if (omp_declare_target_var_p (TREE_VALUE (node)))
+	worklist.safe_push (TREE_VALUE (node));
+
+  while (!worklist.is_empty ())
+    {
+      tree decl = worklist.pop ();
+
+      if (!VAR_P (decl) || !TREE_STATIC (decl))
+	continue;
+
+      for (node = vars; node; node = TREE_CHAIN (node))
+	if (TREE_VALUE (node) == decl)
+	  {
+	    walk_tree_without_duplicates (&TREE_PURPOSE (node),
+					  omp_discover_declare_target_var_r,
+					  &worklist);
+	    break;
+	  }
+    }
+}
 
 /* Create new symbols containing (address, size) pairs for global variables,
    marked with "omp declare target" attribute, as well as addresses for the
diff --git a/gcc/omp-offload.h b/gcc/omp-offload.h
index 0809189..a23a4d3 100644
--- a/gcc/omp-offload.h
+++ b/gcc/omp-offload.h
@@ -31,5 +31,6 @@ extern GTY(()) vec<tree, va_gc> *offload_vars;
 
 extern void omp_finish_file (void);
 extern void omp_discover_implicit_declare_target (void);
+extern void omp_mark_target_static_initializers (tree vars);
 
 #endif /* GCC_OMP_DEVICE_H */
diff --git a/libgomp/testsuite/libgomp.c++/declare_target-3.C b/libgomp/testsuite/libgomp.c++/declare_target-3.C
new file mode 100644
index 0000000..c545613
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare_target-3.C
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#include <stdlib.h>
+
+int f()
+{
+  return 54321;
+}
+
+int g()
+{
+  return 3333;
+}
+
+static int x = f() + 7777;
+static int y = g() + x + 12345;
+
+#pragma omp declare target(y)
+
+int main()
+{
+  int err = 0;
+  #pragma omp target map(from:err)
+  {
+    err |= x != 62098;
+    err |= y != 77776;
+    err |= f() != 54321;
+    err |= g() != 3333;
+  }
+  if (err)
+    abort ();
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)" 2 "gimple" } } */