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

[Tobias Burnus <tobias@codesourcery.com>]

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

On 12.03.21 13:20, Jakub Jelinek wrote:
> On Wed, Mar 10, 2021 at 03:20:42PM +0100, Tobias Burnus wrote:
>> 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.
>>
>> Looking attributes on every get rather than just once is IMHO too expensive.
>> For explicit declare target, we use:
>> ...
>> (e.g. from cp/parser.c).
>> So, I think it would be better to do the same thing

I concur. I was already wondering whether it should be in the FE or
not, thinking about some separation between FE and ME. But given
the precedent and the performance issue, it surely makes sense to
solve it in c*/*decl.c.

Done now. 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-v2.diff --]
[-- Type: text/x-patch, Size: 4284 bytes --]

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.

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