| Message ID | 874mek72lc.fsf@kepler.schwinge.homeip.net |
|---|---|
| State | New |
| Headers | show |
Hi! On 01/11/2016 05:55 AM, Thomas Schwinge wrote: > Hi! > > On Thu, 7 Jan 2016 12:57:32 -0600, James Norris <jnorris@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.) Yuck is right. In the interim I've found a function: get_oacc_fn_attrib, in omp-low.c, that seems to suit my needs. So I'll be revising to use that one instead. > 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)? This appears to be a situation where OpenMP and OpenACC diverge, so I need to discriminate between OpenACC routine and OpenMP declare target. > >> --- 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? No. C/C++ statics and extern's don't appear in 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 <thomas@codesourcery.com> > + > + * c-c++-common/goacc-gomp/nesting-fail-1.c: Add OpenACC declare > + directive for "i". > + > 2016-01-07 James Norris <jnorris@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 > Thank you, thank you, Jim
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 <thomas@codesourcery.com> + + * c-c++-common/goacc-gomp/nesting-fail-1.c: Add OpenACC declare + directive for "i". + 2016-01-07 James Norris <jnorris@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)