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 <frede...@codesourcery.com> wrote:
> Tobias Burnus <tob...@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
[C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)
2020-10-11 Tobias Burnus <tob...@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;
+}
