gimplify.c (omp_notice_variable): Add usage check.

	gcc/
	* gimplify.c (omp_notice_variable): Add usage check.

	gcc/testsuite/
	* c-c++-common/goacc/routine-5.c: Add tests.

From-SVN: r233089

4 files changed, 127 insertions, 1 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index e833810..2cb8456 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,7 @@
+2016-02-02  James Norris  <jnorris@codesourcery.com
+
+	* gimplify.c (omp_notice_variable): Add usage check.
+
 2016-02-02  Alexander Monakov  <amonakov@ispras.ru>
 
 	* config/nvptx/nvptx.c (nvptx_print_operand): Treat LEU, GEU, LTU, GTU
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 32bc1fd..b0ee27e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6087,9 +6087,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
   if (ctx->region_type == ORT_NONE)
     return lang_hooks.decls.omp_disregard_value_expr (decl, false);
 
-  /* Threadprivate variables are predetermined.  */
   if (is_global_var (decl))
     {
+      /* Threadprivate variables are predetermined.  */
       if (DECL_THREAD_LOCAL_P (decl))
 	return omp_notice_threadprivate_variable (ctx, decl, NULL_TREE);
 
@@ -6100,6 +6100,30 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	  if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
 	    return omp_notice_threadprivate_variable (ctx, decl, value);
 	}
+
+      if (gimplify_omp_ctxp->outer_context == NULL
+	  && VAR_P (decl)
+	  && get_oacc_fn_attrib (current_function_decl))
+	{
+	  location_t loc = DECL_SOURCE_LOCATION (decl);
+
+	  if (lookup_attribute ("omp declare target link",
+				DECL_ATTRIBUTES (decl)))
+	    {
+	      error_at (loc,
+			"%qE with %<link%> clause used in %<routine%> function",
+			DECL_NAME (decl));
+	      return false;
+	    }
+	  else if (!lookup_attribute ("omp declare target",
+				      DECL_ATTRIBUTES (decl)))
+	    {
+	      error_at (loc,
+			"%qE requires a %<declare%> directive for use "
+			"in a %<routine%> function", DECL_NAME (decl));
+	      return false;
+	    }
+	}
     }
 
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 6da1f38..8277dff 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,7 @@
+2016-02-02  James Norris  <jnorris@codesourcery.com>
+
+	* c-c++-common/goacc/routine-5.c: Add tests.
+
 2016-02-02  Alexander Monakov  <amonakov@ispras.ru>
 
 	* gcc.target/nvptx/unsigned-cmp.c: New test.
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index ccda097..c34838f 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -45,3 +45,97 @@ using namespace g;
 #pragma acc routine (a) /* { dg-error "does not refer to" } */
   
 #pragma acc routine (c) /* { dg-error "does not refer to" } */
+
+int vb1;		/* { dg-error "directive for use" } */
+extern int vb2;		/* { dg-error "directive for use" } */
+static int vb3;		/* { dg-error "directive for use" } */
+
+#pragma acc routine
+int
+func1 (int a)
+{
+  vb1 = a + 1;
+  vb2 = vb1 + 1;
+  vb3 = vb2 + 1;
+
+  return vb3;
+}
+
+#pragma acc routine
+int
+func2 (int a)
+{
+  extern int vb4;	/* { dg-error "directive for use" } */
+  static int vb5;	/* { dg-error "directive for use" } */
+
+  vb4 = a + 1;
+  vb5 = vb4 + 1;
+
+  return vb5;
+}
+
+extern int vb6;			/* { dg-error "clause used in" } */
+#pragma acc declare link (vb6)
+static int vb7;			/* { dg-error "clause used in" } */
+#pragma acc declare link (vb7)
+
+#pragma acc routine
+int
+func3 (int a)
+{
+  vb6 = a + 1;
+  vb7 = vb6 + 1;
+
+  return vb7;
+}
+
+int vb8;
+#pragma acc declare create (vb8)
+extern int vb9;
+#pragma acc declare create (vb9)
+static int vb10;
+#pragma acc declare create (vb10)
+
+#pragma acc routine
+int
+func4 (int a)
+{
+  vb8 = a + 1;
+  vb9 = vb8 + 1;
+  vb10 = vb9 + 1;
+
+  return vb10;
+}
+
+int vb11;
+#pragma acc declare device_resident (vb11)
+extern int vb12;
+#pragma acc declare device_resident (vb12)
+extern int vb13;
+#pragma acc declare device_resident (vb13)
+
+#pragma acc routine
+int
+func5 (int a)
+{
+  vb11 = a + 1;
+  vb12 = vb11 + 1;
+  vb13 = vb12 + 1;
+
+  return vb13;
+}
+
+#pragma acc routine
+int
+func6 (int a)
+{
+  extern int vb14;
+#pragma acc declare create (vb14)
+  static int vb15;
+#pragma acc declare create (vb15)
+
+  vb14 = a + 1;
+  vb15 = vb14 + 1;
+
+  return vb15;
+}