OpenMP: Fix 'omp declare target' handling for vars [PR99509]

C/C++ use 'omp declare target implicit' when they call
varpool_node::get_create; make sure that a later
'omp declare target' (but not 'omp declare target link')
is honored.

gcc/ChangeLog:

	PR c++/99509
	* varpool.c (varpool_node::get_create): Re-check whether
	'omp declare target'.

libgomp/ChangeLog:

	PR c++/99509
	* testsuite/libgomp.c-c++-common/declare_target-1.c: New test.

 gcc/varpool.c                                      | 18 +++++++++++++++++-
 .../libgomp.c-c++-common/declare_target-1.c        | 22 ++++++++++++++++++++++
 2 files changed, 39 insertions(+), 1 deletion(-)

diff --git a/gcc/varpool.c b/gcc/varpool.c
index 4830df5..5ae6778 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -145,7 +145,23 @@ varpool_node::get_create (tree decl)
   varpool_node *node = varpool_node::get (decl);
   gcc_checking_assert (VAR_P (decl));
   if (node)
-    return node;
+    {
+      /* This happens with C/C++ when 'omp declare target implicit' is set
+	 which is only later replaced by 'omp declate target' or '... link'.  */
+      if (!node->offloadable
+	  && (flag_openacc || flag_openmp)
+	  && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
+	{
+	  node->offloadable = 1;
+	  if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl))
+	    {
+	      g->have_offload = true;
+	      if (!in_lto_p)
+		vec_safe_push (offload_vars, decl);
+	    }
+	}
+      return node;
+    }
 
   node = varpool_node::create_empty ();
   node->decl = decl;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c
new file mode 100644
index 0000000..c5670df
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c
@@ -0,0 +1,22 @@
+/* PR c++/99509  */
+
+#pragma omp declare target
+int data[] = {5};
+#pragma omp end declare target
+
+static inline int
+foo (int idx)
+{
+  return data[idx];
+}
+
+int
+main ()
+{
+  int i = -1;
+  #pragma omp target map(from:i)
+    i = foo(0);
+  if (i != 5)
+    __builtin_abort ();
+  return 0;
+}