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