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

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

[Tobias Burnus]

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

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

[Jakub Jelinek]

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:
          symtab_node *node = symtab_node::get (t);
          if (node != NULL)
            {
              node->offloadable = 1;
              if (ENABLE_OFFLOADING)
                {
                  g->have_offload = true;
                  if (is_a <varpool_node *> (node))
                    vec_safe_push (offload_vars, t);
                }
            }
(e.g. from cp/parser.c).
So, I think it would be better to do the same thing when we turn
an "omp declare target implicit" into "omp declare target", i.e. in
cp/decl.c (cp_finish_decl) and in c/c-decl.c (finish_decl) - right
after adding the "omp declare target" attribute in there.

	Jakub

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

[Tobias Burnus]

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

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

[Jakub Jelinek]

On Mon, Mar 15, 2021 at 09:17:47AM +0100, Tobias Burnus wrote:
> 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.

Ok, thanks.

	Jakub