@@ -1,3 +1,7 @@
+2016-01-XX James Norris <jnorris@codesourcery.com>
+
+ * gimplify.c (omp_notice_variable): Add usage check.
+
2016-01-29 Jakub Jelinek <jakub@redhat.com>
PR target/69551
@@ -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);
@@ -8223,7 +8247,7 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
if (oacc_declare_returns == NULL)
oacc_declare_returns = new hash_map<tree, tree>;
- oacc_declare_returns->put (decl, c);
+ oacc_declare_returns->put (decl, c);
}
}
@@ -1,3 +1,7 @@
+2016-01-XX James Norris <jnorris@codesourcery.com>
+
+ * c-c++-common/goacc/routine-5.c: Add tests.
+
2016-01-29 Jakub Jelinek <jakub@redhat.com>
PR target/69551
@@ -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;
+}