I've backported this patch from trunk to gomp-4_0-branch. This patch
updates a previous patch to gomp4 in dealing with variables used
within a routine procedure.
Reference:
https://gcc.gnu.org/ml/gcc-patches/2016-01/msg01231.html
Thanks,
Jim
==== ChangeLog entries...
gcc/
* gimplify.c (omp_notice_variable): Add usage check.
gcc/testsuite/
* c-c++-common/goacc/routine-5.c: Add tests.
gcc/c/
* c-typeck.c (build_external_ref): Remove usage check.
gcc/cp/
* semantics.c (finish_id_expression): Remove usage check.
Index: gcc/ChangeLog.gomp
===================================================================
--- gcc/ChangeLog.gomp (revision 233104)
+++ gcc/ChangeLog.gomp (working copy)
@@ -1,3 +1,9 @@
+2016-02-03 James Norris <jnor...@codesourcery.com
+
+ Backport form trunk:
+ 2016-02-02 James Norris <jnor...@codesourcery.com>
+ * gimplify.c (omp_notice_variable): Add usage check.
+
2016-01-22 Nathan Sidwell <nat...@codesourcery.com>
* omp-low.c (struct oacc_loop): Add 'inner' field.
Index: gcc/c/ChangeLog.gomp
===================================================================
--- gcc/c/ChangeLog.gomp (revision 233104)
+++ gcc/c/ChangeLog.gomp (working copy)
@@ -1,3 +1,7 @@
+2016-02-03 James Norris <jnor...@codesourcery.com>
+
+ * c-typeck.c (build_external_ref): Remove usage check.
+
2016-01-14 James Norris <jnor...@codesourcery.com>
* c-parser.c (c_finish_oacc_routine): Remove attribute.
Index: gcc/c/c-typeck.c
===================================================================
--- gcc/c/c-typeck.c (revision 233104)
+++ gcc/c/c-typeck.c (working copy)
@@ -2677,26 +2677,6 @@
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)
- && get_oacc_fn_attrib (current_function_decl))
- {
- /* Validate data type for use with routine directive. */
- 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 233104)
+++ gcc/cp/ChangeLog.gomp (working copy)
@@ -1,3 +1,7 @@
+2016-02-03 James Norris <jnor...@codesourcery.com>
+
+ * semantics.c (finish_id_expression): Remove usage check.
+
2016-01-20 Cesar Philippidis <ce...@codesourcery.com>
* parser.c (cp_parser_oacc_all_clauses): Call finish_omp_clauses
Index: gcc/cp/semantics.c
===================================================================
--- gcc/cp/semantics.c (revision 233104)
+++ gcc/cp/semantics.c (working copy)
@@ -3712,25 +3712,6 @@
decl = convert_from_reference (decl);
}
-
- if (decl != error_mark_node
- && current_function_decl
- && TREE_CODE (decl) == VAR_DECL
- && is_global_var (decl)
- && get_oacc_fn_attrib (current_function_decl))
- {
- /* Validate data type for use with routine directive. */
- 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/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 233104)
+++ gcc/gimplify.c (working copy)
@@ -6095,9 +6095,9 @@
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);
@@ -6108,6 +6108,30 @@
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);
Index: gcc/testsuite/ChangeLog.gomp
===================================================================
--- gcc/testsuite/ChangeLog.gomp (revision 233104)
+++ gcc/testsuite/ChangeLog.gomp (working copy)
@@ -1,3 +1,9 @@
+2016-02-03 James Norris <jnor...@codesourcery.com>
+
+ Backport from trunk:
+ 2016-02-02 James Norris <jnor...@codesourcery.com>
+ * c-c++-common/goacc/routine-5.c: Add tests.
+
2016-01-22 Nathan Sidwell <nat...@codesourcery.com>
* c-c++-common/goacc/loop-auto-1.c: Adjust expected warnings.
Index: gcc/testsuite/c-c++-common/goacc/routine-5.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/routine-5.c (revision 233104)
+++ gcc/testsuite/c-c++-common/goacc/routine-5.c (working copy)
@@ -60,48 +60,96 @@
#pragma acc routine (Baz) // { dg-error "not been declared" }
-float vb1;
+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; /* { dg-error "invalid use in" } */
+ vb1 = a + 1;
+ vb2 = vb1 + 1;
+ vb3 = vb2 + 1;
- return vb1; /* { dg-error "invalid use in" } */
+ return vb3;
}
#pragma acc routine
int
func2 (int a)
{
- static int vb2;
+ extern int vb4; /* { dg-error "directive for use" } */
+ static int vb5; /* { dg-error "directive for use" } */
- vb2 = a + 1; /* { dg-error "invalid use in" } */
+ vb4 = a + 1;
+ vb5 = vb4 + 1;
- return vb2; /* { dg-error "invalid use in" } */
+ return vb5;
}
-float vb3;
-#pragma acc declare link (vb3)
+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)
{
- vb3 = a + 1; /* { dg-error "invalid use in" } */
+ vb6 = a + 1;
+ vb7 = vb6 + 1;
- return vb3; /* { dg-error "invalid use in" } */
+ return vb7;
}
-float vb4;
-#pragma acc declare create (vb4)
+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)
{
- vb4 = a + 1;
+ vb8 = a + 1;
+ vb9 = vb8 + 1;
+ vb10 = vb9 + 1;
- return vb4;
+ 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;
+}
