[gomp4] Fix use of declare'd vars by routine procedures.

[James Norris <jnorris at codesourcery dot com>]

Hi!

The checking of variables declared by OpenACC declare directives
used within an OpenACC routine procedure was not being done correctly.
This fix adds the checking required and also adds to the testing.

This fix resolves the issue pointed out by Cesar with declare-4.c
(https://gcc.gnu.org/ml/gcc-patches/2016-01/msg00339.html).

This patch has been applied to gomp-4_0-branch.

Thanks,
Jim

Index: gcc/c/ChangeLog.gomp
===================================================================
--- gcc/c/ChangeLog.gomp	(revision 232138)
+++ gcc/c/ChangeLog.gomp	(working copy)
@@ -1,3 +1,8 @@
+2016-01-07  James Norris  <jnorris@codesourcery.com>
+
+	* c-parser.c (c_finish_oacc_routine): Add new attribute.
+	* c-typeck.c (build_external_ref): Add usage check.
+
 2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* c-parser.c (c_parser_oacc_clause_bind, c_parser_oacc_routine)
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c	(revision 232138)
+++ gcc/c/c-parser.c	(working copy)
@@ -14115,6 +14115,10 @@
   /* Also add an "omp declare target" attribute, with clauses.  */
   DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"),
 					clauses, DECL_ATTRIBUTES (fndecl));
+
+  DECL_ATTRIBUTES (fndecl)
+    = tree_cons (get_identifier ("oacc routine"),
+		 clauses, DECL_ATTRIBUTES (fndecl));
 }
 
 /* OpenACC 2.0:
Index: gcc/c/c-typeck.c
===================================================================
--- gcc/c/c-typeck.c	(revision 232138)
+++ gcc/c/c-typeck.c	(working copy)
@@ -2664,6 +2664,26 @@
   tree ref;
   tree decl = lookup_name (id);
 
+  if (decl
+      && decl != error_mark_node
+      && current_function_decl
+      && TREE_CODE (decl) == VAR_DECL
+      && is_global_var (decl)
+      && lookup_attribute ("oacc routine",
+			   DECL_ATTRIBUTES (current_function_decl)))
+    {
+      if (lookup_attribute ("omp declare target link",
+			    DECL_ATTRIBUTES (decl))
+	  || ((!lookup_attribute ("omp declare target",
+				  DECL_ATTRIBUTES (decl))
+	       && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
+		   || (!TREE_STATIC (decl) && DECL_EXTERNAL (decl))))))
+	{
+	  error_at (loc, "invalid use in %<routine%> function");
+	  return error_mark_node;
+	}
+    }
+
   /* In Objective-C, an instance variable (ivar) may be preferred to
      whatever lookup_name() found.  */
   decl = objc_lookup_ivar (decl, id);
Index: gcc/cp/ChangeLog.gomp
===================================================================
--- gcc/cp/ChangeLog.gomp	(revision 232138)
+++ gcc/cp/ChangeLog.gomp	(working copy)
@@ -1,3 +1,8 @@
+2016-01-07  James Norris  <jnorris@codesourcery.com>
+
+	* parser.c (cp_finalize_oacc_routine): Add new attribute.
+	* semantics.c (finish_id_expression): Add usage check.
+
 2016-01-07  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* cp-tree.h (bind_decls_match): Declare.
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c	(revision 232138)
+++ gcc/cp/parser.c	(working copy)
@@ -36732,6 +36732,10 @@
       DECL_ATTRIBUTES (fndecl)
 	= tree_cons (get_identifier ("omp declare target"),
 		     clauses, DECL_ATTRIBUTES (fndecl));
+
+      DECL_ATTRIBUTES (fndecl)
+	= tree_cons (get_identifier ("oacc routine"),
+		     NULL_TREE, DECL_ATTRIBUTES (fndecl));
     }
 }
 
Index: gcc/cp/semantics.c
===================================================================
--- gcc/cp/semantics.c	(revision 232138)
+++ gcc/cp/semantics.c	(working copy)
@@ -3700,6 +3700,25 @@
 
 	  decl = convert_from_reference (decl);
 	}
+
+      if (decl != error_mark_node
+	  && current_function_decl
+	  && TREE_CODE (decl) == VAR_DECL
+	  && is_global_var (decl)
+	  && lookup_attribute ("oacc routine",
+			       DECL_ATTRIBUTES (current_function_decl)))
+	{
+	  if (lookup_attribute ("omp declare target link",
+				DECL_ATTRIBUTES (decl))
+	      || ((!lookup_attribute ("omp declare target",
+				  DECL_ATTRIBUTES (decl))
+		   && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
+			|| (!TREE_STATIC (decl) && DECL_EXTERNAL (decl))))))
+	    {
+	      *error_msg = "invalid use in %<routine%> function";
+	      return error_mark_node;
+	    }
+	}
     }
 
   return cp_expr (decl, location);
Index: gcc/testsuite/ChangeLog.gomp
===================================================================
--- gcc/testsuite/ChangeLog.gomp	(revision 232138)
+++ gcc/testsuite/ChangeLog.gomp	(working copy)
@@ -1,3 +1,7 @@
+2016-01-07  James Norris  <jnorris@codesourcery.com>
+
+	* c-c++-common/goacc/routine-5.c: Additional tests.
+
 2016-01-07  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* g++.dg/goacc/routine-2.C: Add more coverage.
Index: gcc/testsuite/c-c++-common/goacc/routine-5.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/routine-5.c	(revision 232138)
+++ gcc/testsuite/c-c++-common/goacc/routine-5.c	(working copy)
@@ -59,3 +59,49 @@
 #pragma acc routine (Foo) gang // { dg-error "must be applied before definition" }
 
 #pragma acc routine (Baz) // { dg-error "not been declared" }
+
+float vb1;
+
+#pragma acc routine
+int
+func1 (int a)
+{
+  vb1 = a + 1; /* { dg-error "invalid use in" } */
+
+  return vb1; /* { dg-error "invalid use in" } */
+}
+
+#pragma acc routine
+int
+func2 (int a)
+{
+  static int vb2;
+
+  vb2 = a + 1; /* { dg-error "invalid use in" } */
+
+  return vb2; /* { dg-error "invalid use in" } */
+}
+
+float vb3;
+#pragma acc declare link (vb3)
+
+#pragma acc routine
+int
+func3 (int a)
+{
+  vb3 = a + 1; /* { dg-error "invalid use in" } */
+
+  return vb3; /* { dg-error "invalid use in" } */
+}
+
+float vb4;
+#pragma acc declare create (vb4)
+
+#pragma acc routine
+int
+func4 (int a)
+{
+  vb4 = a + 1;
+
+  return vb4;
+}
Index: libgomp/ChangeLog.gomp
===================================================================
--- libgomp/ChangeLog.gomp	(revision 232138)
+++ libgomp/ChangeLog.gomp	(working copy)
@@ -1,3 +1,7 @@
+2016-01-07  James Norris  <jnorris@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/declare-4.c: Fix test.
+
 2016-01-06  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* testsuite/libgomp.oacc-fortran/pr68813.f90: New test.
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c	(revision 232138)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c	(working copy)
@@ -4,7 +4,7 @@
 #include <openacc.h>
 
 float b;
-#pragma acc declare link (b)
+#pragma acc declare create (b)
 
 #pragma acc routine
 int