Hi! On Thu, 7 Jan 2016 12:57:32 -0600, James Norris <jnor...@codesourcery.com> wrote: > 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.
Such a patch needs to go into trunk; see my report in <http://news.gmane.org/find-root.php?message_id=%3C87mvtapdp0.fsf%40kepler.schwinge.homeip.net%3E>. > --- 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)); > } Yuck, another attribute... (..., and it's not listed/documented in gcc/c-family/c-common.c:c_common_attribute_table.) You store clauses in the "oacc routine" here, but it's unused as far as I can tell. Given that we have the clauses available from the "omp declare target" attribute (subject to change to switching to a generic "omp clauses" attribute as suggested in <http://news.gmane.org/find-root.php?message_id=%3C87twns3ebs.fsf%40hertz.schwinge.homeip.net%3E>), could we maybe just look up some specific clause instead of detecting the presence of this extra attribute? (Jakub, any preference?) Anyway, have we verified that the desired behavior: > --- 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; > + } > + } ..., isn't applicable to OpenMP as well (thus, no "oacc routine" attribute conditional is needed here)? > --- 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)); > } You don't store clauses in the "oacc routine" here. > --- 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; > + } > + } Likewise. No equivalent change is needed for Fortran? > --- 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 I have not verified the details, but a very similar fix was required to get rid of a number of regressions: @@ -2637,18 +2637,18 @@ PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 350) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 356) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 358) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 360) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 371) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 378) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 386) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 395) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 402) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 409) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 416) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 42) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 423) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 430) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 432) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 434) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 47) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 52) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 57) @@ -2667,7 +2667,7 @@ PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 93) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 95) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 97) PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 99) [-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for excess errors) Same for C++. Committed to gomp-4_0-branch in r232219: commit 1ac87f2b59dd03cb305ec94a7c6b5657dbb54e66 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Mon Jan 11 11:51:57 2016 +0000 Resolve c-c++-common/goacc-gomp/nesting-fail-1.c regressions gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: Add OpenACC declare directive for "i". git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@232219 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/testsuite/ChangeLog.gomp | 5 +++++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c | 1 + 2 files changed, 6 insertions(+) diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 607ca8e..2db11df 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,8 @@ +2016-01-11 Thomas Schwinge <tho...@codesourcery.com> + + * c-c++-common/goacc-gomp/nesting-fail-1.c: Add OpenACC declare + directive for "i". + 2016-01-07 James Norris <jnor...@codesourcery.com> * c-c++-common/goacc/routine-5.c: Additional tests. diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 8e0f217..9011fcf 100644 --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -1,4 +1,5 @@ extern int i; +#pragma acc declare create(i) void f_omp (void) Grüße Thomas
signature.asc
Description: PGP signature
