Fwd: Re: [Fortran, OpenACC] Reject vars of different scope in $acc declare (PR94120)

Fwd: Re: [Fortran, OpenACC] Reject vars of different scope in $acc declare (PR94120)

[Tobias Burnus]

(I assume that should have gone to gcc-patches@ and fortran@ as well.)


-------- Forwarded Message --------
Subject:        Re: [Fortran, OpenACC] Reject vars of different scope in $acc
declare (PR94120)
Date:   Tue, 10 Mar 2020 18:02:41 +0000
From:   Paul Richard Thomas <paul.richard.thomas@gmail.com>
To:     Tobias Burnus <tobias@codesourcery.com>



Hi Tobias,

This looks to be straightforward enough to me :-) OK for trunk.

Thanks

Paul

On Tue, 10 Mar 2020 at 17:18, Tobias Burnus <tobias@codesourcery.com> wrote:
> [This fixes a bunch of issues found when actually only
> wanting to add a check for the following restriction.]
>
> OpenACC's "Declare Directive" has the following restriction:
>
> "A declare directive must be in the same scope
>    as the declaration of any var that appears in
>    the data clauses of the directive."
>
> The gfortran now rejects a "var" declared is in a different
> namespace than the "$acc declare". (Use-associated variables
> are already rejected.) Testing showed that a straight-forward
> check fails if the result variable is the function name – as
> then the function symbol is in the parent scope. — Extending
> the failing test to use a result variable showed that the
> current is-a-module-variable check didn't work and when fixing,
> one was running into a likewise issue.
>
> The reason that I exclude 's' being a module is that at
> resolution time, an is-variable check is done.
>
>
> The other changes are because the following restriction was
> mishandled:
>
> "In a Fortran module declaration section, only
>    create, copyin, device_resident, and link clauses are allowed."
>
> But all examples where for variables using those in a module
> procedure …
>
>
> OK for the trunk?
>
> Cheers,
>
> Tobias
>
> PS: For those who wounder what happens in a BLOCK DATA construct:
> gfortran outrightly rejects 'acc declare'. (It probably should
> work for COMMON blocks with 'declare device_resident' – but
> currently the spec only permits it in program/subroutine/function
> + declaration part of a module.)
>
> PPS: The PR shows (for C) that one can construct a test case,
> which violates the OpenACC restriction and actually fails with
> an ICE. I have a draft patch for C (see PR) but not yet one for
> C++, hence, I start with the Fortran side. – I currently still
> struggle to write a same-scope check in C++.
> [The C test case in turn was a fallout of debugging an
> ICE-on-valid-code issue …]
>
> -----------------
> Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
> Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter



--
"If you can't explain it simply, you don't understand it well enough"
- Albert Einstein

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Tobias Burnus]

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

Fortran patch: https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html

Like for Fortran, it also fixes some other issues – here for C++
related to namespaces. (For class, see PR c++/94140.)

Test case of the PR yields an ICE in the middle end and the
namespace tests an ICE in cc1plus. Additionally, invalid code
is not diagnosed.

The OpenACC spec has under "Declare Directive" has the following restriction:

"A declare directive must be in the same scope
  as the declaration of any var that appears in
  the data clauses of the directive."

("A declare directive is used […] following a variable
   declaration in C or C++".)

NOTE for C++: This patch assumes that variables in a namespace
are handled in the same way as those which are at
global (namespace) scope; however, the OpenACC specification's
wording currently is "In C or C++ global scope, only …".
Hence, one can argue about this part of the patch; but as
it fixes an ICE and is a very sensible extension – the other
option is to reject it – I believe it is fine.
(On the OpenACC side, this is now Issue 288.)

OK for the trunk?

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[-- Attachment #2: same-scope-c-cxx.diff --]
[-- Type: text/x-patch, Size: 7214 bytes --]

[C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

2020-10-11  Tobias Burnus  <tobias@codesourcery.com>

	gcc/c/
	PR middle-end/94120
	* c-decl.c (c_check_oacc_same_scope): New function.
	* c-tree.h (c_check_oacc_same_scope): Declare it.
	* c-parser.c (c_parser_oacc_declare): Add check that variables
	are declared in the same scope as the directive. Fix handling
	of namespace vars.

	gcc/c/
	PR middle-end/94120
	* paser.c (cp_parser_oacc_declare): Add check that variables
        are declared in the same scope as the directive.

	gcc/testsuite/
	PR middle-end/94120
	* c-c++-common/goacc/declare-pr94120.c: New.
	* g++.dg/declare-pr94120.C: New.

	libgomp/testsuite/
	PR middle-end/94120
	* libgomp.oacc-c++/declare-pr94120.C: New.
	

 gcc/c/c-decl.c                                     |  8 +++
 gcc/c/c-parser.c                                   |  9 ++++
 gcc/c/c-tree.h                                     |  1 +
 gcc/cp/parser.c                                    | 21 +++++++-
 gcc/testsuite/c-c++-common/goacc/declare-pr94120.c | 23 +++++++++
 gcc/testsuite/g++.dg/declare-pr94120.C             | 30 ++++++++++++
 .../testsuite/libgomp.oacc-c++/declare-pr94120.C   | 57 ++++++++++++++++++++++
 7 files changed, 147 insertions(+), 2 deletions(-)

diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c
index c819fd0d0d5..eda95d664de 100644
--- a/gcc/c/c-decl.c
+++ b/gcc/c/c-decl.c
@@ -12016,4 +12016,12 @@ c_check_omp_declare_reduction_r (tree *tp, int *, void *data)
   return NULL_TREE;
 }
 
+
+bool
+c_check_oacc_same_scope (tree decl)
+{
+  struct c_binding *b = I_SYMBOL_BINDING (DECL_NAME (decl));
+  return b != NULL && B_IN_CURRENT_SCOPE (b);
+}
+
 #include "gt-c-c-decl.h"
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 1e8f2f7108d..63e8ab0ad17 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -16573,6 +16573,15 @@ c_parser_oacc_declare (c_parser *parser)
 	  break;
 	}
 
+      if (!c_check_oacc_same_scope (decl))
+	{
+	  error_at (loc,
+		    "%qD must be a variable declared in the same scope as "
+		    "%<#pragma acc declare%>", decl);
+	  error = true;
+	  continue;
+	}
+
       if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
 	  || lookup_attribute ("omp declare target link",
 			       DECL_ATTRIBUTES (decl)))
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 71229927cb6..6d578705d77 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -789,6 +789,7 @@ extern tree c_omp_reduction_id (enum tree_code, tree);
 extern tree c_omp_reduction_decl (tree);
 extern tree c_omp_reduction_lookup (tree, tree);
 extern tree c_check_omp_declare_reduction_r (tree *, int *, void *);
+extern bool c_check_oacc_same_scope (tree);
 extern void c_pushtag (location_t, tree, tree);
 extern void c_bind (location_t, tree, bool);
 extern bool tag_exists_p (enum tree_code, tree);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 24f71671469..8f09eb0d375 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40722,6 +40722,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 {
   tree clauses, stmt;
   bool error = false;
+  bool found_in_scope = global_bindings_p ();
 
   clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
 					"#pragma acc declare", pragma_tok, true);
@@ -40794,6 +40795,22 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	  break;
 	}
 
+      if (!found_in_scope)
+	for (tree d = current_binding_level->names; d; d = TREE_CHAIN (d))
+	  if (d == decl)
+	    {
+	      found_in_scope = true;
+	      break;
+	    }
+      if (!found_in_scope)
+	{
+	  error_at (loc,
+		    "%qD must be a variable declared in the same scope as "
+		    "%<#pragma acc declare%>", decl);
+	  error = true;
+	  continue;
+	}
+
       if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
 	  || lookup_attribute ("omp declare target link",
 			       DECL_ATTRIBUTES (decl)))
@@ -40815,7 +40832,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 
 	  DECL_ATTRIBUTES (decl)
 	    = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
-	  if (global_bindings_p ())
+	  if (current_binding_level->kind == sk_namespace)
 	    {
 	      symtab_node *node = symtab_node::get (decl);
 	      if (node != NULL)
@@ -40832,7 +40849,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	}
     }
 
-  if (error || global_bindings_p ())
+  if (error || current_binding_level->kind == sk_namespace)
     return NULL_TREE;
 
   stmt = make_node (OACC_DECLARE);
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c b/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c
new file mode 100644
index 00000000000..21b2cc14fc7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c
@@ -0,0 +1,23 @@
+/* { dg-do compile }  */
+
+/* PR middle-end/94120  */
+
+void foo()
+{
+  int foo;
+  {
+    #pragma acc declare copy(foo)  /* { dg-error "'foo' must be a variable declared in the same scope as '#pragma acc declare'" }  */
+  }
+}
+
+void
+f_data (void)
+{
+  int B[10];
+#pragma acc data
+  {
+# pragma acc declare copy(B)  /* { dg-error "'B' must be a variable declared in the same scope as '#pragma acc declare'" }  */
+    for (int i = 0; i < 10; i++)
+      B[i] = -i;
+  }
+}
diff --git a/gcc/testsuite/g++.dg/declare-pr94120.C b/gcc/testsuite/g++.dg/declare-pr94120.C
new file mode 100644
index 00000000000..8515c4ff875
--- /dev/null
+++ b/gcc/testsuite/g++.dg/declare-pr94120.C
@@ -0,0 +1,30 @@
+/* { dg-do compile }  */
+
+/* PR middle-end/94120  */
+
+int b[8];
+#pragma acc declare create (b)
+ 
+namespace my {
+ int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+ #pragma acc declare copyin (d)
+};
+
+namespace outer {
+  namespace inner {
+    int e[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+    #pragma acc declare copyin (e)
+  };
+};
+
+int f[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+namespace my {
+ #pragma acc declare copyin (f)
+};
+
+namespace outer {
+  int g[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+  namespace inner {
+    #pragma acc declare copyin (g)
+  };
+};
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
new file mode 100644
index 00000000000..1e1254187ea
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
@@ -0,0 +1,57 @@
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 8
+
+namespace one {
+  int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+  #pragma acc declare copyin (A)
+};
+
+namespace outer {
+  namespace inner {
+    int B[N];
+    #pragma acc declare create (B)
+  };
+};
+
+static void
+f (void)
+{
+  int i;
+  int C[N];
+  #pragma acc declare copyout (C)
+
+  if (!acc_is_present (&one::A, sizeof (one::A)))
+    abort ();
+
+  if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B)))
+    abort ();
+
+#pragma acc parallel
+  for (i = 0; i < N; i++)
+    {
+      outer::inner::B[i] = one::A[i];
+      C[i] = outer::inner::B[i];
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      if (C[i] != i + 1)
+	abort ();
+    }
+
+#pragma acc parallel
+  for (i = 0; i < N; i++)
+    if (outer::inner::B[i] != i + 1)
+      abort ();
+}
+
+
+int
+main (int argc, char **argv)
+{
+  f ();
+
+  return 0;
+}

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Frederik Harwath]

Tobias Burnus <tobias@codesourcery.com> writes:

Hi Tobias,

> Fortran patch: https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html
>
> "A declare directive must be in the same scope
>   as the declaration of any var that appears in
>   the data clauses of the directive."
>
> ("A declare directive is used […] following a variable
>    declaration in C or C++".)
>
> NOTE for C++: This patch assumes that variables in a namespace
> are handled in the same way as those which are at
> global (namespace) scope; however, the OpenACC specification's
> wording currently is "In C or C++ global scope, only …".
> Hence, one can argue about this part of the patch; but as
> it fixes an ICE and is a very sensible extension – the other
> option is to reject it – I believe it is fine.
> (On the OpenACC side, this is now Issue 288.)

Sounds reasonable to me.

> +bool
> +c_check_oacc_same_scope (tree decl)
> +{
> +  struct c_binding *b = I_SYMBOL_BINDING (DECL_NAME (decl));
> +  return b != NULL && B_IN_CURRENT_SCOPE (b);
> +}

Is the function really specific to OpenACC? If not, then "_oacc"
could be dropped from its name. How about "c_check_current_scope"?

> diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
> index 24f71671469..8f09eb0d375 100644
> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c
> [...]
> -       if (global_bindings_p ())
> +       if (current_binding_level->kind == sk_namespace)
> [...]
> -  if (error || global_bindings_p ())
> +  if (error || current_binding_level->kind == sk_namespace)
>      return NULL_TREE;

So - just to be sure - the new namespace condition subsumes the old
"global_bindings_p" condition because the global scope is also a namespace,
right? Yes, now I see that you have a test case that demonstrates that
the declare directive still works for global variables with those changes.

> diff --git a/gcc/testsuite/g++.dg/declare-pr94120.C b/gcc/testsuite/g++.dg/declare-pr94120.C
> new file mode 100644
> index 00000000000..8515c4ff875
> --- /dev/null
> +++ b/gcc/testsuite/g++.dg/declare-pr94120.C
> @@ -0,0 +1,30 @@
> +/* { dg-do compile }  */
> +
> +/* PR middle-end/94120  */
> +
> +int b[8];
> +#pragma acc declare create (b)

Looks good to me.

Frederik
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Thomas Schwinge]

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

Hi!

Oh, the OpenACC 'declare' implementation in GCC -- be careful to not look
too carefully, or you'll discover one problem after the other...  ;-/
..., and also, a number of issues are open in the OpenACC tracker
regarding the 'declare' clause.


On 2020-03-12T15:43:03+0100, Frederik Harwath <frederik@codesourcery.com> wrote:
> Tobias Burnus <tobias@codesourcery.com> writes:
>> Fortran patch: https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html
>>
>> "A declare directive must be in the same scope
>>   as the declaration of any var that appears in
>>   the data clauses of the directive."
>>
>> ("A declare directive is used […] following a variable
>>    declaration in C or C++".)
>>
>> NOTE for C++: This patch assumes that variables in a namespace
>> are handled in the same way as those which are at
>> global (namespace) scope; however, the OpenACC specification's
>> wording currently is "In C or C++ global scope, only …".
>> Hence, one can argue about this part of the patch; but as
>> it fixes an ICE and is a very sensible extension – the other
>> option is to reject it – I believe it is fine.
>> (On the OpenACC side, this is now Issue 288.)

Please in GCC PRs use the "See Also" field "to refer to bugs in other
installations".  That makes it easy to cross-reference GCC with OpenACC
tickets.

> Sounds reasonable to me.

Even if the ICE is then fixed, should we still keep
<https://gcc.gnu.org/PR94120> open (with a note) until
<https://github.com/OpenACC/openacc-spec/issues/288> is
resolved/published?

Regarding the C/C++ patch you posted: I'm not at all familiar with the
front ends' scoping implementation.  But, given that your patch really
only touches the OpenACC 'declare' code paths, it can't cause any harm
otherwise, so we shall accept it.  Maybe Jakub has any comments, though?
(Patch attached again, for easy reference.)

>> --- a/gcc/c/c-decl.c
>> +++ b/gcc/c/c-decl.c

>> +bool
>> +c_check_oacc_same_scope (tree decl)
>> +{
>> +  struct c_binding *b = I_SYMBOL_BINDING (DECL_NAME (decl));
>> +  return b != NULL && B_IN_CURRENT_SCOPE (b);
>> +}
>
> Is the function really specific to OpenACC? If not, then "_oacc"
> could be dropped from its name. How about "c_check_current_scope"?

Yeah, that may be a good idea.  Similar constructs seem to be used in a
few other places, though without the 'DECL_NAME' indirection.

>> --- a/gcc/cp/parser.c
>> +++ b/gcc/cp/parser.c
>> [...]
>> -       if (global_bindings_p ())
>> +       if (current_binding_level->kind == sk_namespace)
>> [...]
>> -  if (error || global_bindings_p ())
>> +  if (error || current_binding_level->kind == sk_namespace)
>>      return NULL_TREE;
>
> So - just to be sure - the new namespace condition subsumes the old
> "global_bindings_p" condition because the global scope is also a namespace,
> right? Yes, now I see that you have a test case that demonstrates that
> the declare directive still works for global variables with those changes.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: same-scope-c-cxx.diff --]
[-- Type: text/x-diff, Size: 7214 bytes --]

[C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

2020-10-11  Tobias Burnus  <tobias@codesourcery.com>

	gcc/c/
	PR middle-end/94120
	* c-decl.c (c_check_oacc_same_scope): New function.
	* c-tree.h (c_check_oacc_same_scope): Declare it.
	* c-parser.c (c_parser_oacc_declare): Add check that variables
	are declared in the same scope as the directive. Fix handling
	of namespace vars.

	gcc/c/
	PR middle-end/94120
	* paser.c (cp_parser_oacc_declare): Add check that variables
        are declared in the same scope as the directive.

	gcc/testsuite/
	PR middle-end/94120
	* c-c++-common/goacc/declare-pr94120.c: New.
	* g++.dg/declare-pr94120.C: New.

	libgomp/testsuite/
	PR middle-end/94120
	* libgomp.oacc-c++/declare-pr94120.C: New.
	

 gcc/c/c-decl.c                                     |  8 +++
 gcc/c/c-parser.c                                   |  9 ++++
 gcc/c/c-tree.h                                     |  1 +
 gcc/cp/parser.c                                    | 21 +++++++-
 gcc/testsuite/c-c++-common/goacc/declare-pr94120.c | 23 +++++++++
 gcc/testsuite/g++.dg/declare-pr94120.C             | 30 ++++++++++++
 .../testsuite/libgomp.oacc-c++/declare-pr94120.C   | 57 ++++++++++++++++++++++
 7 files changed, 147 insertions(+), 2 deletions(-)

diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c
index c819fd0d0d5..eda95d664de 100644
--- a/gcc/c/c-decl.c
+++ b/gcc/c/c-decl.c
@@ -12016,4 +12016,12 @@ c_check_omp_declare_reduction_r (tree *tp, int *, void *data)
   return NULL_TREE;
 }
 
+
+bool
+c_check_oacc_same_scope (tree decl)
+{
+  struct c_binding *b = I_SYMBOL_BINDING (DECL_NAME (decl));
+  return b != NULL && B_IN_CURRENT_SCOPE (b);
+}
+
 #include "gt-c-c-decl.h"
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 1e8f2f7108d..63e8ab0ad17 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -16573,6 +16573,15 @@ c_parser_oacc_declare (c_parser *parser)
 	  break;
 	}
 
+      if (!c_check_oacc_same_scope (decl))
+	{
+	  error_at (loc,
+		    "%qD must be a variable declared in the same scope as "
+		    "%<#pragma acc declare%>", decl);
+	  error = true;
+	  continue;
+	}
+
       if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
 	  || lookup_attribute ("omp declare target link",
 			       DECL_ATTRIBUTES (decl)))
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 71229927cb6..6d578705d77 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -789,6 +789,7 @@ extern tree c_omp_reduction_id (enum tree_code, tree);
 extern tree c_omp_reduction_decl (tree);
 extern tree c_omp_reduction_lookup (tree, tree);
 extern tree c_check_omp_declare_reduction_r (tree *, int *, void *);
+extern bool c_check_oacc_same_scope (tree);
 extern void c_pushtag (location_t, tree, tree);
 extern void c_bind (location_t, tree, bool);
 extern bool tag_exists_p (enum tree_code, tree);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 24f71671469..8f09eb0d375 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40722,6 +40722,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 {
   tree clauses, stmt;
   bool error = false;
+  bool found_in_scope = global_bindings_p ();
 
   clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
 					"#pragma acc declare", pragma_tok, true);
@@ -40794,6 +40795,22 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	  break;
 	}
 
+      if (!found_in_scope)
+	for (tree d = current_binding_level->names; d; d = TREE_CHAIN (d))
+	  if (d == decl)
+	    {
+	      found_in_scope = true;
+	      break;
+	    }
+      if (!found_in_scope)
+	{
+	  error_at (loc,
+		    "%qD must be a variable declared in the same scope as "
+		    "%<#pragma acc declare%>", decl);
+	  error = true;
+	  continue;
+	}
+
       if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
 	  || lookup_attribute ("omp declare target link",
 			       DECL_ATTRIBUTES (decl)))
@@ -40815,7 +40832,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 
 	  DECL_ATTRIBUTES (decl)
 	    = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
-	  if (global_bindings_p ())
+	  if (current_binding_level->kind == sk_namespace)
 	    {
 	      symtab_node *node = symtab_node::get (decl);
 	      if (node != NULL)
@@ -40832,7 +40849,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	}
     }
 
-  if (error || global_bindings_p ())
+  if (error || current_binding_level->kind == sk_namespace)
     return NULL_TREE;
 
   stmt = make_node (OACC_DECLARE);
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c b/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c
new file mode 100644
index 00000000000..21b2cc14fc7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c
@@ -0,0 +1,23 @@
+/* { dg-do compile }  */
+
+/* PR middle-end/94120  */
+
+void foo()
+{
+  int foo;
+  {
+    #pragma acc declare copy(foo)  /* { dg-error "'foo' must be a variable declared in the same scope as '#pragma acc declare'" }  */
+  }
+}
+
+void
+f_data (void)
+{
+  int B[10];
+#pragma acc data
+  {
+# pragma acc declare copy(B)  /* { dg-error "'B' must be a variable declared in the same scope as '#pragma acc declare'" }  */
+    for (int i = 0; i < 10; i++)
+      B[i] = -i;
+  }
+}
diff --git a/gcc/testsuite/g++.dg/declare-pr94120.C b/gcc/testsuite/g++.dg/declare-pr94120.C
new file mode 100644
index 00000000000..8515c4ff875
--- /dev/null
+++ b/gcc/testsuite/g++.dg/declare-pr94120.C
@@ -0,0 +1,30 @@
+/* { dg-do compile }  */
+
+/* PR middle-end/94120  */
+
+int b[8];
+#pragma acc declare create (b)
+ 
+namespace my {
+ int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+ #pragma acc declare copyin (d)
+};
+
+namespace outer {
+  namespace inner {
+    int e[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+    #pragma acc declare copyin (e)
+  };
+};
+
+int f[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+namespace my {
+ #pragma acc declare copyin (f)
+};
+
+namespace outer {
+  int g[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+  namespace inner {
+    #pragma acc declare copyin (g)
+  };
+};
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
new file mode 100644
index 00000000000..1e1254187ea
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
@@ -0,0 +1,57 @@
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 8
+
+namespace one {
+  int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+  #pragma acc declare copyin (A)
+};
+
+namespace outer {
+  namespace inner {
+    int B[N];
+    #pragma acc declare create (B)
+  };
+};
+
+static void
+f (void)
+{
+  int i;
+  int C[N];
+  #pragma acc declare copyout (C)
+
+  if (!acc_is_present (&one::A, sizeof (one::A)))
+    abort ();
+
+  if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B)))
+    abort ();
+
+#pragma acc parallel
+  for (i = 0; i < N; i++)
+    {
+      outer::inner::B[i] = one::A[i];
+      C[i] = outer::inner::B[i];
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      if (C[i] != i + 1)
+	abort ();
+    }
+
+#pragma acc parallel
+  for (i = 0; i < N; i++)
+    if (outer::inner::B[i] != i + 1)
+      abort ();
+}
+
+
+int
+main (int argc, char **argv)
+{
+  f ();
+
+  return 0;
+}

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Tobias Burnus]

I have now committed this patch
as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373,
reading "so we shall accept it" as approval …

On 4/1/20 9:07 AM, Thomas Schwinge wrote:

> Even if the ICE is then fixed, should we still keep
> <https://gcc.gnu.org/PR94120> open (with a note) until
> <https://github.com/OpenACC/openacc-spec/issues/288> is
> resolved/published?

I decided to close it and mention the namelist issue and
the now-fixed PR in PR84140 (which is about (non)static
class member variables, which is also covered in Issue 288).

> Regarding the C/C++ patch you posted: I'm not at all familiar with the
> front ends' scoping implementation.  But, given that your patch really
> only touches the OpenACC 'declare' code paths, it can't cause any harm
> otherwise, so we shall accept it.  Maybe Jakub has any comments, though?

>>> +c_check_oacc_same_scope (tree decl)
>> Is the function really specific to OpenACC? If not, then "_oacc"
>> could be dropped from its name. How about "c_check_current_scope"?
> Yeah, that may be a good idea.  Similar constructs seem to be used in a
> few other places, though without the 'DECL_NAME' indirection.

I now use c_check_in_current_scope.

Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[H.J. Lu]

On Wed, Apr 8, 2020 at 12:55 AM Tobias Burnus <tobias@codesourcery.com> wrote:
>
> I have now committed this patch
> as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373,

On Linux/x86, I got

FAIL: g++.dg/declare-pr94120.C  -std=c++14 (test for excess errors)
FAIL: g++.dg/declare-pr94120.C  -std=c++17 (test for excess errors)
FAIL: g++.dg/declare-pr94120.C  -std=c++2a (test for excess errors)
FAIL: g++.dg/declare-pr94120.C  -std=c++98 (test for excess errors)

> reading "so we shall accept it" as approval …
>
> On 4/1/20 9:07 AM, Thomas Schwinge wrote:
>
> > Even if the ICE is then fixed, should we still keep
> > <https://gcc.gnu.org/PR94120> open (with a note) until
> > <https://github.com/OpenACC/openacc-spec/issues/288> is
> > resolved/published?
>
> I decided to close it and mention the namelist issue and
> the now-fixed PR in PR84140 (which is about (non)static
> class member variables, which is also covered in Issue 288).
>
> > Regarding the C/C++ patch you posted: I'm not at all familiar with the
> > front ends' scoping implementation.  But, given that your patch really
> > only touches the OpenACC 'declare' code paths, it can't cause any harm
> > otherwise, so we shall accept it.  Maybe Jakub has any comments, though?
>
> >>> +c_check_oacc_same_scope (tree decl)
> >> Is the function really specific to OpenACC? If not, then "_oacc"
> >> could be dropped from its name. How about "c_check_current_scope"?
> > Yeah, that may be a good idea.  Similar constructs seem to be used in a
> > few other places, though without the 'DECL_NAME' indirection.
>
> I now use c_check_in_current_scope.
>
> Tobias
> -----------------
> Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
> Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter



-- 
H.J.

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Tobias Burnus]

There was a glitch in the test case (location + dg-error),
already fixed by Jakub (thanks!)
in commit r10-7637-g08d1e7a5aabcf7eeac48bfd99deb80451b8f9974

Sorry,

Tobias

On 4/8/20 7:13 PM, H.J. Lu wrote:

> On Wed, Apr 8, 2020 at 12:55 AM Tobias Burnus <tobias@codesourcery.com> wrote:
>> I have now committed this patch
>> as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373,
> On Linux/x86, I got
>
> FAIL: g++.dg/declare-pr94120.C  -std=c++14 (test for excess errors)
> FAIL: g++.dg/declare-pr94120.C  -std=c++17 (test for excess errors)
> FAIL: g++.dg/declare-pr94120.C  -std=c++2a (test for excess errors)
> FAIL: g++.dg/declare-pr94120.C  -std=c++98 (test for excess errors)
>
>> reading "so we shall accept it" as approval …
>>
>> On 4/1/20 9:07 AM, Thomas Schwinge wrote:
>>
>>> Even if the ICE is then fixed, should we still keep
>>> <https://gcc.gnu.org/PR94120> open (with a note) until
>>> <https://github.com/OpenACC/openacc-spec/issues/288> is
>>> resolved/published?
>> I decided to close it and mention the namelist issue and
>> the now-fixed PR in PR84140 (which is about (non)static
>> class member variables, which is also covered in Issue 288).
>>
>>> Regarding the C/C++ patch you posted: I'm not at all familiar with the
>>> front ends' scoping implementation.  But, given that your patch really
>>> only touches the OpenACC 'declare' code paths, it can't cause any harm
>>> otherwise, so we shall accept it.  Maybe Jakub has any comments, though?
>>>>> +c_check_oacc_same_scope (tree decl)
>>>> Is the function really specific to OpenACC? If not, then "_oacc"
>>>> could be dropped from its name. How about "c_check_current_scope"?
>>> Yeah, that may be a good idea.  Similar constructs seem to be used in a
>>> few other places, though without the 'DECL_NAME' indirection.
>> I now use c_check_in_current_scope.
>>
>> Tobias
>> -----------------
>> Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
>> Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
>
>
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Thomas Schwinge]

Hi Tobias!

On 2020-04-08T09:55:24+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
> I have now committed this patch
> as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373,
> reading "so we shall accept it" as approval …

That's OK.  As I said: "I'm not at all familiar with the front ends'
scoping implementation", and don't currently have time to learn about
that.  So, either you're confident that you're doing the right things
there (which I shall assume, given that you didn't explicitly ask whether
you're doing the right things), or you need to wait for somebody else to
review the patch.

Reviewers don't know everything -- certainly I don't ;-) -- and don't
have bandwidth to learn everything.  And even if a proper review is done,
often enough the reviewer fails to foresee the one detail that should've
been caught.  Hence, I'm happy to incrementally improve things, as long
as a patch isn't regressing any (or, too many) other things, and isn't
conceptually questionable (both of which doesn't apply here).  That's why
I said "we shall accept [the patch]" -- meaning "approved for commit,
without my proper review".


Grüße
 Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

*PING* Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Tobias Burnus]

On 3/11/20 2:28 PM, Tobias Burnus wrote:
> Fortran patch:
> https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html
>
> Like for Fortran, it also fixes some other issues – here for C++
> related to namespaces. (For class, see PR c++/94140.)
>
> Test case of the PR yields an ICE in the middle end and the
> namespace tests an ICE in cc1plus. Additionally, invalid code
> is not diagnosed.
>
> The OpenACC spec has under "Declare Directive" has the following
> restriction:
>
> "A declare directive must be in the same scope
>  as the declaration of any var that appears in
>  the data clauses of the directive."
>
> ("A declare directive is used […] following a variable
>   declaration in C or C++".)
>
> NOTE for C++: This patch assumes that variables in a namespace
> are handled in the same way as those which are at
> global (namespace) scope; however, the OpenACC specification's
> wording currently is "In C or C++ global scope, only …".
> Hence, one can argue about this part of the patch; but as
> it fixes an ICE and is a very sensible extension – the other
> option is to reject it – I believe it is fine.
> (On the OpenACC side, this is now Issue 288.)
>
> OK for the trunk?
>
> Tobias
>
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

*PING**2 Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Tobias Burnus]

*PING**2

On 3/24/20 11:18 AM, Tobias Burnus wrote:
> On 3/11/20 2:28 PM, Tobias Burnus wrote:
>> Fortran patch:
>> https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html
>>
>> Like for Fortran, it also fixes some other issues – here for C++
>> related to namespaces. (For class, see PR c++/94140.)
>>
>> Test case of the PR yields an ICE in the middle end and the
>> namespace tests an ICE in cc1plus. Additionally, invalid code
>> is not diagnosed.
>>
>> The OpenACC spec has under "Declare Directive" has the following
>> restriction:
>>
>> "A declare directive must be in the same scope
>>  as the declaration of any var that appears in
>>  the data clauses of the directive."
>>
>> ("A declare directive is used […] following a variable
>>   declaration in C or C++".)
>>
>> NOTE for C++: This patch assumes that variables in a namespace
>> are handled in the same way as those which are at
>> global (namespace) scope; however, the OpenACC specification's
>> wording currently is "In C or C++ global scope, only …".
>> Hence, one can argue about this part of the patch; but as
>> it fixes an ICE and is a very sensible extension – the other
>> option is to reject it – I believe it is fine.
>> (On the OpenACC side, this is now Issue 288.)
>>
>> OK for the trunk?
>>
>> Tobias
>>
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Re: *PING**2 Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Thomas Schwinge]

Hi Tobias!

On 2020-04-07T09:21:27+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
> *PING**2

<http://mid.mail-archive.com/87h7y3ofad.fsf@euler.schwinge.homeip.net>?


Grüße
 Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[committed] openacc: Fix up declare-pr94120.C testcase [PR94533]

[Jakub Jelinek]

Hi!

On Wed, Mar 11, 2020 at 02:28:44PM +0100, Tobias Burnus wrote:
> 	gcc/testsuite/
> 	PR middle-end/94120
> 	* c-c++-common/goacc/declare-pr94120.c: New.
> 	* g++.dg/declare-pr94120.C: New.

This test has been put in a wrong directory, where OpenACC tests aren't
tested with -fopenacc, and also contained trailing semicolons.
I've moved it where it belongs, added dg-error directives and removed
the extra semicolons.

Tested on x86_64-linux, committed to trunk as obvious.

2020-04-08  Jakub Jelinek  <jakub@redhat.com>

	PR middle-end/94120
	PR testsuite/94533
	* g++.dg/declare-pr94120.C: Move test to ...
	* g++.dg/goacc/declare-pr94120.C: ... here.  Add dg-error directives.
	Remove semicolons after } at the end of namespaces.

diff --git a/gcc/testsuite/g++.dg/declare-pr94120.C b/gcc/testsuite/g++.dg/goacc/declare-pr94120.C
similarity index 60%
rename from gcc/testsuite/g++.dg/declare-pr94120.C
rename to gcc/testsuite/g++.dg/goacc/declare-pr94120.C
index 8515c4ff875..7aa56121e38 100644
--- a/gcc/testsuite/g++.dg/declare-pr94120.C
+++ b/gcc/testsuite/g++.dg/goacc/declare-pr94120.C
@@ -8,23 +8,23 @@ int b[8];
 namespace my {
  int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
  #pragma acc declare copyin (d)
-};
+}
 
 namespace outer {
   namespace inner {
     int e[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
     #pragma acc declare copyin (e)
-  };
-};
+  }
+}
 
 int f[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
 namespace my {
- #pragma acc declare copyin (f)
-};
+ #pragma acc declare copyin (f)		/* { dg-error "'f' must be a variable declared in the same scope as '#pragma acc declare'" }  */
+}
 
 namespace outer {
   int g[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
   namespace inner {
-    #pragma acc declare copyin (g)
-  };
-};
+    #pragma acc declare copyin (g)	/* { dg-error "'outer::g' must be a variable declared in the same scope as '#pragma acc declare'" }  */
+  }
+}


	Jakub

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)

[Thomas Schwinge]

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

Hi Tobias!

I do see the new 'libgomp.oacc-c++/declare-pr94120.C' (see below, for
reference) FAIL for '-foffload=nvptx-none' execution testing.  The reason
is that the 'C' array doesn't have the content it's checked to have.

Now, my question, shouldn't it be changed like the attached patch, or
similar, because we actually first need to return from function 'f' in
order for the 'copyout(C)' to be executed?

Or, otherwise, what's the use of the 'copyout' clause with OpenACC
'declare'?  I suppose it's only meant to be used with function arguments,
because for every locally-defined variable (like 'C' in
'libgomp.oacc-c++/declare-pr94120.C'), that variable will be gone just
after the 'copyout' is executed, so the 'copyout' conceptually
equals/acts as 'create' in that case.  However, OpenACC explicitly does
allow 'copyout' in certain scenarios.  Do you think my interpretation is
correct, or what am I missing?  (In case the answer isn't obvious to you,
too, please file an issue with OpenACC, linking to this email for
reference.)

However, GCC doesn't like my changes either; actually two errors are
diagnosed:

    [...]/libgomp.oacc-c++/declare-pr94120.C: In function ‘void f(int*)’:
    [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: array section in ‘#pragma acc declare’
       21 | #pragma acc declare copyout (C[0:N])
          |                              ^
    [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: ‘C’ must be a variable declared in the same scope as ‘#pragma acc declare’

Please have a look, and as necessary file GCC PRs, and XFAIL the
'libgomp.oacc-c++/declare-pr94120.C' execution testing for
'-DACC_MEM_SHARED=0'.


Grüße
 Thomas


On 2020-03-11T13:28:44+0100, Tobias Burnus <tobias@codesourcery.com> wrote:
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
> @@ -0,0 +1,57 @@
> +#include <openacc.h>
> +#include <stdlib.h>
> +
> +#define N 8
> +
> +namespace one {
> +  int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 };
> +  #pragma acc declare copyin (A)
> +};
> +
> +namespace outer {
> +  namespace inner {
> +    int B[N];
> +    #pragma acc declare create (B)
> +  };
> +};
> +
> +static void
> +f (void)
> +{
> +  int i;
> +  int C[N];
> +  #pragma acc declare copyout (C)
> +
> +  if (!acc_is_present (&one::A, sizeof (one::A)))
> +    abort ();
> +
> +  if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B)))
> +    abort ();
> +
> +#pragma acc parallel
> +  for (i = 0; i < N; i++)
> +    {
> +      outer::inner::B[i] = one::A[i];
> +      C[i] = outer::inner::B[i];
> +    }
> +
> +  for (i = 0; i < N; i++)
> +    {
> +      if (C[i] != i + 1)
> +     abort ();
> +    }
> +
> +#pragma acc parallel
> +  for (i = 0; i < N; i++)
> +    if (outer::inner::B[i] != i + 1)
> +      abort ();
> +}
> +
> +
> +int
> +main (int argc, char **argv)
> +{
> +  f ();
> +
> +  return 0;
> +}


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-WIP-Fix-libgomp.oacc-c-declare-pr94120.C-PR94120.patch --]
[-- Type: text/x-diff, Size: 1568 bytes --]

From cdb045dbfaec6dfca1cc49aa4d9b1fc79fc7837e Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 10 Apr 2020 14:50:28 +0200
Subject: [PATCH] [WIP] Fix 'libgomp.oacc-c++/declare-pr94120.C' [PR94120]

---
 .../libgomp.oacc-c++/declare-pr94120.C        | 25 +++++++++++--------
 1 file changed, 15 insertions(+), 10 deletions(-)

diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
index 1e1254187ea..bad74865135 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
@@ -16,11 +16,11 @@ namespace outer {
 };
 
 static void
-f (void)
+f (int *C)
 {
+#pragma acc declare copyout (C[0:N])
+
   int i;
-  int C[N];
-  #pragma acc declare copyout (C)
 
   if (!acc_is_present (&one::A, sizeof (one::A)))
     abort ();
@@ -28,6 +28,9 @@ f (void)
   if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B)))
     abort ();
 
+  if (!acc_is_present (C, N * sizeof *C))
+    abort ();
+
 #pragma acc parallel
   for (i = 0; i < N; i++)
     {
@@ -35,12 +38,6 @@ f (void)
       C[i] = outer::inner::B[i];
     }
 
-  for (i = 0; i < N; i++)
-    {
-      if (C[i] != i + 1)
-	abort ();
-    }
-
 #pragma acc parallel
   for (i = 0; i < N; i++)
     if (outer::inner::B[i] != i + 1)
@@ -51,7 +48,15 @@ f (void)
 int
 main (int argc, char **argv)
 {
-  f ();
+  int C[N];
+
+  f (C);
+
+  for (int i = 0; i < N; i++)
+    {
+      if (C[i] != i + 1)
+	abort ();
+    }
 
   return 0;
 }
-- 
2.17.1

Re: [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120) | Fix declare copyout in libgomp.oacc-c++/declare-pr94120.C

[Tobias Burnus]

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

Hi Thomas,

I have now fixed it temporarily differently – without actually
testing that feature. – See attachment.

For the proper fix, it would help to get the semantic right in
OpenACC itself (→ ongoing discussion). I think the patch
for PR94120 is probably not completely right – and should be
fixed for GCC 10, but I think one should wait for the next two
OpenACC meetings to get converged. One presumably needs to handle
'device_resident' and 'link' differently from 'copy* etc.

I added a note to PR 94140.

Regards,

Tobias

On 4/10/20 3:20 PM, Thomas Schwinge wrote:

> Hi Tobias!
>
> I do see the new 'libgomp.oacc-c++/declare-pr94120.C' (see below, for
> reference) FAIL for '-foffload=nvptx-none' execution testing.  The reason
> is that the 'C' array doesn't have the content it's checked to have.
>
> Now, my question, shouldn't it be changed like the attached patch, or
> similar, because we actually first need to return from function 'f' in
> order for the 'copyout(C)' to be executed?
>
> Or, otherwise, what's the use of the 'copyout' clause with OpenACC
> 'declare'?  I suppose it's only meant to be used with function arguments,
> because for every locally-defined variable (like 'C' in
> 'libgomp.oacc-c++/declare-pr94120.C'), that variable will be gone just
> after the 'copyout' is executed, so the 'copyout' conceptually
> equals/acts as 'create' in that case.  However, OpenACC explicitly does
> allow 'copyout' in certain scenarios.  Do you think my interpretation is
> correct, or what am I missing?  (In case the answer isn't obvious to you,
> too, please file an issue with OpenACC, linking to this email for
> reference.)
>
> However, GCC doesn't like my changes either; actually two errors are
> diagnosed:
>
>      [...]/libgomp.oacc-c++/declare-pr94120.C: In function ‘void f(int*)’:
>      [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: array section in ‘#pragma acc declare’
>         21 | #pragma acc declare copyout (C[0:N])
>            |                              ^
>      [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: ‘C’ must be a variable declared in the same scope as ‘#pragma acc declare’
>
> Please have a look, and as necessary file GCC PRs, and XFAIL the
> 'libgomp.oacc-c++/declare-pr94120.C' execution testing for
> '-DACC_MEM_SHARED=0'.
>
>
> Grüße
>   Thomas
>
>
> On 2020-03-11T13:28:44+0100, Tobias Burnus<tobias@codesourcery.com>  wrote:
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
>> @@ -0,0 +1,57 @@
>> +#include <openacc.h>
>> +#include <stdlib.h>
>> +
>> +#define N 8
>> +
>> +namespace one {
>> +  int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 };
>> +  #pragma acc declare copyin (A)
>> +};
>> +
>> +namespace outer {
>> +  namespace inner {
>> +    int B[N];
>> +    #pragma acc declare create (B)
>> +  };
>> +};
>> +
>> +static void
>> +f (void)
>> +{
>> +  int i;
>> +  int C[N];
>> +  #pragma acc declare copyout (C)
>> +
>> +  if (!acc_is_present (&one::A, sizeof (one::A)))
>> +    abort ();
>> +
>> +  if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B)))
>> +    abort ();
>> +
>> +#pragma acc parallel
>> +  for (i = 0; i < N; i++)
>> +    {
>> +      outer::inner::B[i] = one::A[i];
>> +      C[i] = outer::inner::B[i];
>> +    }
>> +
>> +  for (i = 0; i < N; i++)
>> +    {
>> +      if (C[i] != i + 1)
>> +    abort ();
>> +    }
>> +
>> +#pragma acc parallel
>> +  for (i = 0; i < N; i++)
>> +    if (outer::inner::B[i] != i + 1)
>> +      abort ();
>> +}
>> +
>> +
>> +int
>> +main (int argc, char **argv)
>> +{
>> +  f ();
>> +
>> +  return 0;
>> +}
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[-- Attachment #2: committed.diff --]
[-- Type: text/x-patch, Size: 1557 bytes --]

r10-7810-g85d8c05a02bf7d1b256f806582a11e3fd8970a32
commit 85d8c05a02bf7d1b256f806582a11e3fd8970a32
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Mon Apr 20 12:38:50 2020 +0200

    Fix declare copyout in libgomp.oacc-c++/declare-pr94120.C
    
    Testing on the host does not make sense for 'declare copyout' for
    a same-scope stack-allocated variable. Once the copyout is done,
    the variable is gone. Hence, test the variable on the device. This
    can be revisit after the OpenACC semantic has been fixed; but with
    that fix, the test PASSes again with devices.
    
            PR middle-end/94120
            * testsuite/libgomp.oacc-c++/declare-pr94120.C: Fix 'declare copy(out)'
            test case.

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index ce71ac6e783..b1cf297a0d7 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2020-04-20  Tobias Burnus  <tobias@codesourcery.com>
+
+	PR middle-end/94120
+	* testsuite/libgomp.oacc-c++/declare-pr94120.C: Fix 'declare copy(out)'
+	test case.
+
 2020-04-17  Tobias Burnus  <tobias@codesourcery.com>
 
 	PR middle-end/94635
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
index 1e1254187ea..ed69359b533 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
@@ -35,6 +35,7 @@ f (void)
       C[i] = outer::inner::B[i];
     }
 
+#pragma acc parallel
   for (i = 0; i < N; i++)
     {
       if (C[i] != i + 1)

Re: Re: [Fortran, OpenACC] Reject vars of different scope in $acc declare (PR94120)

[Paul Richard Thomas]

I would swear that I selected "reply all" but... Never mind, you put it right!

Cheers

Paul

On Thu, 12 Mar 2020 at 08:44, Tobias Burnus <tobias@codesourcery.com> wrote:
>
> (I assume that should have gone to gcc-patches@ and fortran@ as well.)
>
>
> -------- Forwarded Message --------
> Subject:        Re: [Fortran, OpenACC] Reject vars of different scope in $acc
> declare (PR94120)
> Date:   Tue, 10 Mar 2020 18:02:41 +0000
> From:   Paul Richard Thomas <paul.richard.thomas@gmail.com>
> To:     Tobias Burnus <tobias@codesourcery.com>
>
>
>
> Hi Tobias,
>
> This looks to be straightforward enough to me :-) OK for trunk.
>
> Thanks
>
> Paul
>
> On Tue, 10 Mar 2020 at 17:18, Tobias Burnus <tobias@codesourcery.com> wrote:
> > [This fixes a bunch of issues found when actually only
> > wanting to add a check for the following restriction.]
> >
> > OpenACC's "Declare Directive" has the following restriction:
> >
> > "A declare directive must be in the same scope
> >    as the declaration of any var that appears in
> >    the data clauses of the directive."
> >
> > The gfortran now rejects a "var" declared is in a different
> > namespace than the "$acc declare". (Use-associated variables
> > are already rejected.) Testing showed that a straight-forward
> > check fails if the result variable is the function name – as
> > then the function symbol is in the parent scope. — Extending
> > the failing test to use a result variable showed that the
> > current is-a-module-variable check didn't work and when fixing,
> > one was running into a likewise issue.
> >
> > The reason that I exclude 's' being a module is that at
> > resolution time, an is-variable check is done.
> >
> >
> > The other changes are because the following restriction was
> > mishandled:
> >
> > "In a Fortran module declaration section, only
> >    create, copyin, device_resident, and link clauses are allowed."
> >
> > But all examples where for variables using those in a module
> > procedure …
> >
> >
> > OK for the trunk?
> >
> > Cheers,
> >
> > Tobias
> >
> > PS: For those who wounder what happens in a BLOCK DATA construct:
> > gfortran outrightly rejects 'acc declare'. (It probably should
> > work for COMMON blocks with 'declare device_resident' – but
> > currently the spec only permits it in program/subroutine/function
> > + declaration part of a module.)
> >
> > PPS: The PR shows (for C) that one can construct a test case,
> > which violates the OpenACC restriction and actually fails with
> > an ICE. I have a draft patch for C (see PR) but not yet one for
> > C++, hence, I start with the Fortran side. – I currently still
> > struggle to write a same-scope check in C++.
> > [The C test case in turn was a fallout of debugging an
> > ICE-on-valid-code issue …]
> >
> > -----------------
> > Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
> > Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
>
>
>
> --
> "If you can't explain it simply, you don't understand it well enough"
> - Albert Einstein
>
> -----------------
> Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
> Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter



-- 
"If you can't explain it simply, you don't understand it well enough"
- Albert Einstein

[Fortran, OpenACC] Reject vars of different scope in $acc declare (PR94120)

[Tobias Burnus]

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

[This fixes a bunch of issues found when actually only
wanting to add a check for the following restriction.]

OpenACC's "Declare Directive" has the following restriction:

"A declare directive must be in the same scope
  as the declaration of any var that appears in
  the data clauses of the directive."

The gfortran now rejects a "var" declared is in a different
namespace than the "$acc declare". (Use-associated variables
are already rejected.) Testing showed that a straight-forward
check fails if the result variable is the function name – as
then the function symbol is in the parent scope. — Extending
the failing test to use a result variable showed that the
current is-a-module-variable check didn't work and when fixing,
one was running into a likewise issue.

The reason that I exclude 's' being a module is that at
resolution time, an is-variable check is done.


The other changes are because the following restriction was
mishandled:

"In a Fortran module declaration section, only
  create, copyin, device_resident, and link clauses are allowed."

But all examples where for variables using those in a module
procedure …


OK for the trunk?

Cheers,

Tobias

PS: For those who wounder what happens in a BLOCK DATA construct:
gfortran outrightly rejects 'acc declare'. (It probably should
work for COMMON blocks with 'declare device_resident' – but
currently the spec only permits it in program/subroutine/function
+ declaration part of a module.)

PPS: The PR shows (for C) that one can construct a test case,
which violates the OpenACC restriction and actually fails with
an ICE. I have a draft patch for C (see PR) but not yet one for
C++, hence, I start with the Fortran side. – I currently still
struggle to write a same-scope check in C++.
[The C test case in turn was a fallout of debugging an
ICE-on-valid-code issue …]

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[-- Attachment #2: same-scope.diff --]
[-- Type: text/x-patch, Size: 7302 bytes --]

[Fortran, OpenACC] Reject vars of different scope in $acc declare (PR94120)

2020-10-03  Tobias Burnus  <tobias@codesourcery.com>

	PR middle-end/94120
	* openmp.c (gfc_match_oacc_declare): Accept function-result
	variables; reject variables declared in a different scoping unit.

2020-10-03  Tobias Burnus  <tobias@codesourcery.com>

	PR middle-end/94120
	* gfortran.dg/goacc/pr78260-2.f90: Correct scan-tree-dump-times.
	Extend test case to result variables.
	* gfortran.dg/goacc/declare-2.f95: Actually check module-declaration
	restriction of OpenACC.
	* gfortran.dg/goacc/declare-3.f95: Remove case where this
	restriction is violated.
	* gfortran.dg/goacc/pr94120-1.f90: New.
	* gfortran.dg/goacc/pr94120-2.f90: New.
	* gfortran.dg/goacc/pr94120-3.f90: New.

 gcc/fortran/openmp.c                          | 12 +++++++++++-
 gcc/testsuite/gfortran.dg/goacc/declare-2.f95 | 21 ++++++++++++++++-----
 gcc/testsuite/gfortran.dg/goacc/declare-3.f95 | 10 +---------
 gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 | 13 +++++++++++--
 gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90 | 11 +++++++++++
 gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90 | 12 ++++++++++++
 gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90 | 13 +++++++++++++
 7 files changed, 75 insertions(+), 17 deletions(-)

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 35f6b2f4938..930bca541b9 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2155,7 +2155,8 @@ gfc_match_oacc_declare (void)
     {
       gfc_symbol *s = n->sym;
 
-      if (s->ns->proc_name && s->ns->proc_name->attr.proc == PROC_MODULE)
+      if (gfc_current_ns->proc_name
+	  && gfc_current_ns->proc_name->attr.flavor == FL_MODULE)
 	{
 	  if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
 	    {
@@ -2174,6 +2175,15 @@ gfc_match_oacc_declare (void)
 	  return MATCH_ERROR;
 	}
 
+      if ((s->result == s && s->ns->contained != gfc_current_ns)
+	  || ((s->attr.flavor == FL_UNKNOWN || s->attr.flavor == FL_VARIABLE)
+	      && s->ns != gfc_current_ns))
+	{
+	  gfc_error ("Variable %qs shall be declared in the same scoping unit "
+		     "as !$ACC DECLARE at %L", s->name, &where);
+	  return MATCH_ERROR;
+	}
+
       if ((s->attr.dimension || s->attr.codimension)
 	  && s->attr.dummy && s->as->type != AS_EXPLICIT)
 	{
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
index 7aa3dab4707..bad5de9d757 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
@@ -1,9 +1,5 @@
 
 module amod
-
-contains
-
-subroutine asubr (b)
   implicit none
   integer :: b(8)
 
@@ -16,9 +12,24 @@ subroutine asubr (b)
   !$acc declare present_or_create (b) ! { dg-error "present on multiple" }
   !$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" }
   !$acc declare create (b) copyin (b) ! { dg-error "present on multiple" }
+end module
 
+module amod2
+contains
+subroutine asubr (a, b, c, d, e, f, g, h, i, j, k)
+  implicit none
+  integer, dimension(8) :: a, b, c, d, e, f, g, h, i, j, k
+
+  !$acc declare copy (a)
+  !$acc declare copyout (b)
+  !$acc declare present (c)
+  !$acc declare present_or_copy (d)
+  !$acc declare present_or_copyin (e)
+  !$acc declare present_or_copyout (f)
+  !$acc declare present_or_create (g)
+  !$acc declare deviceptr (h)
+  !$acc declare create (j) copyin (k)
 end subroutine
-
 end module
 
 module bmod
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
index 80d9903a9dc..9127cba6600 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
@@ -14,12 +14,6 @@ module mod_b
   !$acc declare copyin (b)
 end module
 
-module mod_c
-  implicit none
-  integer :: c
-  !$acc declare deviceptr (c)
-end module
-
 module mod_d
   implicit none
   integer :: d
@@ -35,7 +29,6 @@ end module
 subroutine sub1
   use mod_a
   use mod_b
-  use mod_c
   use mod_d
   use mod_e
 end subroutine sub1
@@ -43,11 +36,10 @@ end subroutine sub1
 program test
   use mod_a
   use mod_b
-  use mod_c
   use mod_d
   use mod_e
 
-  ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_deviceptr:c\) map\(force_to:b\) map\(force_alloc:a\)$} original } }
+  ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_to:b\) map\(force_alloc:a\)$} original } }
 end program test
 
 ! { dg-final { scan-tree-dump-times {#pragma acc data} 1 original } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90
index e28564d6f70..f8a3dc864cc 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90
@@ -4,6 +4,8 @@
 
 ! PR fortran/78260
 
+! Loosely related to PR fortran/94120
+
 module m
   implicit none
   integer :: n = 0
@@ -14,7 +16,14 @@ contains
     f1 = 5 
     !$acc end kernels
   end function f1
+  integer function g1() result(g1res)
+    !$acc declare present(g1res)
+    !$acc kernels copyin(g1res)
+    g1res = 5 
+    !$acc end kernels
+  end function g1
 end module m
 ! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } }
-
+! { dg-final { scan-tree-dump-times "#pragma acc kernels map\\(to:__result_f1\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:g1res\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma acc kernels map\\(to:g1res\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90 b/gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90
new file mode 100644
index 00000000000..8d1fc16a5ff
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90
@@ -0,0 +1,11 @@
+! { dg-do compile }
+!
+! PR fortran/94120
+!
+implicit none
+integer :: i
+contains
+  subroutine f()
+    !$acc declare copy(i)  ! { dg-error "Variable 'i' shall be declared in the same scoping unit as !.ACC DECLARE" }
+  end
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90 b/gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90
new file mode 100644
index 00000000000..216c04b844d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90
@@ -0,0 +1,12 @@
+! { dg-do compile }
+!
+! PR fortran/94120
+!
+! BLOCK is not supported in OpenACC <= 3.0
+!
+subroutine f()
+  block
+    integer :: k
+    !$acc declare copy(j)  ! { dg-error "Sorry, !.ACC DECLARE at .1. is not allowed in BLOCK construct" }
+  end block
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90 b/gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90
new file mode 100644
index 00000000000..1eec90ae7ed
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90
@@ -0,0 +1,13 @@
+! { dg-do compile }
+!
+! PR fortran/94120
+!
+! Note: BLOCK is not supported in OpenACC <= 3.0 – but the following check comes earlier:
+! It is also invalid because the variable is in a different scoping unit
+!
+subroutine g()
+  integer :: k
+  block
+    !$acc declare copy(k)  ! { dg-error "Variable 'k' shall be declared in the same scoping unit as !.ACC DECLARE" }
+  end block
+end