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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index e8338107f18cff315483afd72caf1556ce8a4a58..2cb845618556a68c05abade583988dd9a3fc9973 100644 (file)
--- 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 32bc1fdb2fb8dd11b25ddfca11548fdb90ce947e..b0ee27e9997f117f501045a79a04aba6170d8329 100644 (file)
--- 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 6da1f381b6c236177b517f2be90e58adb5365445..8277dff817e3ecfa50627092bbe4109bc0fe2403 100644 (file)
--- 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 ccda0976a04a306d7e0c4595005bf8752d9b780a..c34838f9d03a1b73c866870b0095f770606be586 100644 (file)
--- 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;
+}