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 --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c
index b559ed5..3b2241b 100644 (file)
--- 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 9c7f6e5..1b671ce 100644 (file)
--- 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 (file)
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;
+}