[gcc r11-7665] OpenMP: Fix 'omp declare target' handling for vars [PR99509]

[Tobias Burnus]

https://gcc.gnu.org/g:f20fe2cb213dffba47ec1b62c625590b5bbe50d7

commit r11-7665-gf20fe2cb213dffba47ec1b62c625590b5bbe50d7
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Mon Mar 15 10:12:58 2021 +0100

    OpenMP: Fix 'omp declare target' handling for vars [PR99509]
    
    For variables with 'declare target' attribute,
    varpool_node::get_create marks variables as offload; however,
    if the node already exists, it is not updated. C/C++ may tag
    decl with 'declare target implicit', which may only be after
    varpool creation turned into 'declare target' or 'declare target link';
    in this case, the tagging has to happen in the FE.
    
    gcc/c/ChangeLog:
    
            PR c++/99509
            * c-decl.c (finish_decl): For 'omp declare target implicit' vars,
            ensure that the varpool node is marked as offloadable.
    
    gcc/cp/ChangeLog:
    
            PR c++/99509
            * decl.c (cp_finish_decl): For 'omp declare target implicit' vars,
            ensure that the varpool node is marked as offloadable.
    
    libgomp/ChangeLog:
    
            PR c++/99509
            * testsuite/libgomp.c-c++-common/declare_target-1.c: New test.

Diff:
---
 gcc/c/c-decl.c                                     | 22 +++++++++++++++++++---
 gcc/cp/decl.c                                      | 21 ++++++++++++++++++---
 .../libgomp.c-c++-common/declare_target-1.c        | 22 ++++++++++++++++++++++
 3 files changed, 59 insertions(+), 6 deletions(-)

diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c
index b559ed5d76a..3b2241bfd97 100644
--- a/gcc/c/c-decl.c
+++ b/gcc/c/c-decl.c
@@ -58,6 +58,9 @@ along with GCC; see the file COPYING3.  If not see
 #include "c-family/name-hint.h"
 #include "c-family/known-headers.h"
 #include "c-family/c-spellcheck.h"
+#include "context.h"  /* For 'g'.  */
+#include "omp-general.h"
+#include "omp-offload.h"  /* For offload_vars.  */
 
 #include "tree-pretty-print.h"
 
@@ -5658,9 +5661,22 @@ finish_decl (tree decl, location_t init_loc, tree init,
 				  DECL_ATTRIBUTES (decl))
 	       && !lookup_attribute ("omp declare target link",
 				     DECL_ATTRIBUTES (decl)))
-	DECL_ATTRIBUTES (decl)
-	  = tree_cons (get_identifier ("omp declare target"),
-		       NULL_TREE, DECL_ATTRIBUTES (decl));
+	{
+	  DECL_ATTRIBUTES (decl)
+	    = tree_cons (get_identifier ("omp declare target"),
+			 NULL_TREE, DECL_ATTRIBUTES (decl));
+	    symtab_node *node = symtab_node::get (decl);
+	    if (node != NULL)
+	      {
+		node->offloadable = 1;
+		if (ENABLE_OFFLOADING)
+		  {
+		    g->have_offload = true;
+		    if (is_a <varpool_node *> (node))
+		      vec_safe_push (offload_vars, decl);
+		  }
+	      }
+	}
     }
 
   /* This is the last point we can lower alignment so give the target the
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index 9c7f6e59bbb..1b671cec9a3 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -53,7 +53,9 @@ along with GCC; see the file COPYING3.  If not see
 #include "asan.h"
 #include "gcc-rich-location.h"
 #include "langhooks.h"
+#include "context.h"  /* For 'g'.  */
 #include "omp-general.h"
+#include "omp-offload.h"  /* For offload_vars.  */
 
 /* Possible cases of bad specifiers type used by bad_specifiers. */
 enum bad_spec_place {
@@ -8176,9 +8178,22 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p,
 				  DECL_ATTRIBUTES (decl))
 	       && !lookup_attribute ("omp declare target link",
 				     DECL_ATTRIBUTES (decl)))
-	DECL_ATTRIBUTES (decl)
-	  = tree_cons (get_identifier ("omp declare target"),
-		       NULL_TREE, DECL_ATTRIBUTES (decl));
+	{
+	  DECL_ATTRIBUTES (decl)
+	    = tree_cons (get_identifier ("omp declare target"),
+			 NULL_TREE, DECL_ATTRIBUTES (decl));
+	  symtab_node *node = symtab_node::get (decl);
+	  if (node != NULL)
+	    {
+	      node->offloadable = 1;
+	      if (ENABLE_OFFLOADING)
+		{
+		  g->have_offload = true;
+		  if (is_a <varpool_node *> (node))
+		    vec_safe_push (offload_vars, decl);
+		}
+	    }
+	}
     }
 
   /* This is the last point we can lower alignment so give the target the
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 00000000000..c5670dfb7db
--- /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;
+}