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

[Tobias Burnus <tobias@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 668 bytes --]

The C/C++ FE sets for an 'omp declare target' ... 'omp end declare target'
the attribute 'omp declare target implicit'.

That's later processed (for C++) in decl.c - which remove that attribute
and either keeps and explicit 'omp declare target' or 'omp declare target link'
attribute.

Unfortunately, adding 'omp declare target' comes too late as the varpool
has been generated.

I am not sure what's the best way, but the following seems to work.

OK for mainline?

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf

[-- Attachment #2: omp-decl-tgt.diff --]
[-- Type: text/x-patch, Size: 2107 bytes --]

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