Message ID | 0e98c5c9-7536-e11f-c42e-4e8060b147a5@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | [C/C++,OpenACC] Reject vars of different scope in acc declare (PR94120) | expand |
Tobias Burnus <tobias@codesourcery.com> writes: Hi Tobias, > Fortran patch: https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html > > "A declare directive must be in the same scope > as the declaration of any var that appears in > the data clauses of the directive." > > ("A declare directive is used […] following a variable > declaration in C or C++".) > > NOTE for C++: This patch assumes that variables in a namespace > are handled in the same way as those which are at > global (namespace) scope; however, the OpenACC specification's > wording currently is "In C or C++ global scope, only …". > Hence, one can argue about this part of the patch; but as > it fixes an ICE and is a very sensible extension – the other > option is to reject it – I believe it is fine. > (On the OpenACC side, this is now Issue 288.) Sounds reasonable to me. > +bool > +c_check_oacc_same_scope (tree decl) > +{ > + struct c_binding *b = I_SYMBOL_BINDING (DECL_NAME (decl)); > + return b != NULL && B_IN_CURRENT_SCOPE (b); > +} Is the function really specific to OpenACC? If not, then "_oacc" could be dropped from its name. How about "c_check_current_scope"? > diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c > index 24f71671469..8f09eb0d375 100644 > --- a/gcc/cp/parser.c > +++ b/gcc/cp/parser.c > [...] > - if (global_bindings_p ()) > + if (current_binding_level->kind == sk_namespace) > [...] > - if (error || global_bindings_p ()) > + if (error || current_binding_level->kind == sk_namespace) > return NULL_TREE; So - just to be sure - the new namespace condition subsumes the old "global_bindings_p" condition because the global scope is also a namespace, right? Yes, now I see that you have a test case that demonstrates that the declare directive still works for global variables with those changes. > diff --git a/gcc/testsuite/g++.dg/declare-pr94120.C b/gcc/testsuite/g++.dg/declare-pr94120.C > new file mode 100644 > index 00000000000..8515c4ff875 > --- /dev/null > +++ b/gcc/testsuite/g++.dg/declare-pr94120.C > @@ -0,0 +1,30 @@ > +/* { dg-do compile } */ > + > +/* PR middle-end/94120 */ > + > +int b[8]; > +#pragma acc declare create (b) Looks good to me. Frederik ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On 3/11/20 2:28 PM, Tobias Burnus wrote: > Fortran patch: > https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html > > Like for Fortran, it also fixes some other issues – here for C++ > related to namespaces. (For class, see PR c++/94140.) > > Test case of the PR yields an ICE in the middle end and the > namespace tests an ICE in cc1plus. Additionally, invalid code > is not diagnosed. > > The OpenACC spec has under "Declare Directive" has the following > restriction: > > "A declare directive must be in the same scope > as the declaration of any var that appears in > the data clauses of the directive." > > ("A declare directive is used […] following a variable > declaration in C or C++".) > > NOTE for C++: This patch assumes that variables in a namespace > are handled in the same way as those which are at > global (namespace) scope; however, the OpenACC specification's > wording currently is "In C or C++ global scope, only …". > Hence, one can argue about this part of the patch; but as > it fixes an ICE and is a very sensible extension – the other > option is to reject it – I believe it is fine. > (On the OpenACC side, this is now Issue 288.) > > OK for the trunk? > > Tobias > ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi! Oh, the OpenACC 'declare' implementation in GCC -- be careful to not look too carefully, or you'll discover one problem after the other... ;-/ ..., and also, a number of issues are open in the OpenACC tracker regarding the 'declare' clause. On 2020-03-12T15:43:03+0100, Frederik Harwath <frederik@codesourcery.com> wrote: > Tobias Burnus <tobias@codesourcery.com> writes: >> Fortran patch: https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html >> >> "A declare directive must be in the same scope >> as the declaration of any var that appears in >> the data clauses of the directive." >> >> ("A declare directive is used […] following a variable >> declaration in C or C++".) >> >> NOTE for C++: This patch assumes that variables in a namespace >> are handled in the same way as those which are at >> global (namespace) scope; however, the OpenACC specification's >> wording currently is "In C or C++ global scope, only …". >> Hence, one can argue about this part of the patch; but as >> it fixes an ICE and is a very sensible extension – the other >> option is to reject it – I believe it is fine. >> (On the OpenACC side, this is now Issue 288.) Please in GCC PRs use the "See Also" field "to refer to bugs in other installations". That makes it easy to cross-reference GCC with OpenACC tickets. > Sounds reasonable to me. Even if the ICE is then fixed, should we still keep <https://gcc.gnu.org/PR94120> open (with a note) until <https://github.com/OpenACC/openacc-spec/issues/288> is resolved/published? Regarding the C/C++ patch you posted: I'm not at all familiar with the front ends' scoping implementation. But, given that your patch really only touches the OpenACC 'declare' code paths, it can't cause any harm otherwise, so we shall accept it. Maybe Jakub has any comments, though? (Patch attached again, for easy reference.) >> --- a/gcc/c/c-decl.c >> +++ b/gcc/c/c-decl.c >> +bool >> +c_check_oacc_same_scope (tree decl) >> +{ >> + struct c_binding *b = I_SYMBOL_BINDING (DECL_NAME (decl)); >> + return b != NULL && B_IN_CURRENT_SCOPE (b); >> +} > > Is the function really specific to OpenACC? If not, then "_oacc" > could be dropped from its name. How about "c_check_current_scope"? Yeah, that may be a good idea. Similar constructs seem to be used in a few other places, though without the 'DECL_NAME' indirection. >> --- a/gcc/cp/parser.c >> +++ b/gcc/cp/parser.c >> [...] >> - if (global_bindings_p ()) >> + if (current_binding_level->kind == sk_namespace) >> [...] >> - if (error || global_bindings_p ()) >> + if (error || current_binding_level->kind == sk_namespace) >> return NULL_TREE; > > So - just to be sure - the new namespace condition subsumes the old > "global_bindings_p" condition because the global scope is also a namespace, > right? Yes, now I see that you have a test case that demonstrates that > the declare directive still works for global variables with those changes. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
*PING**2 On 3/24/20 11:18 AM, Tobias Burnus wrote: > On 3/11/20 2:28 PM, Tobias Burnus wrote: >> Fortran patch: >> https://gcc.gnu.org/pipermail/gcc-patches/current/541774.html >> >> Like for Fortran, it also fixes some other issues – here for C++ >> related to namespaces. (For class, see PR c++/94140.) >> >> Test case of the PR yields an ICE in the middle end and the >> namespace tests an ICE in cc1plus. Additionally, invalid code >> is not diagnosed. >> >> The OpenACC spec has under "Declare Directive" has the following >> restriction: >> >> "A declare directive must be in the same scope >> as the declaration of any var that appears in >> the data clauses of the directive." >> >> ("A declare directive is used […] following a variable >> declaration in C or C++".) >> >> NOTE for C++: This patch assumes that variables in a namespace >> are handled in the same way as those which are at >> global (namespace) scope; however, the OpenACC specification's >> wording currently is "In C or C++ global scope, only …". >> Hence, one can argue about this part of the patch; but as >> it fixes an ICE and is a very sensible extension – the other >> option is to reject it – I believe it is fine. >> (On the OpenACC side, this is now Issue 288.) >> >> OK for the trunk? >> >> Tobias >> ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi Tobias!
On 2020-04-07T09:21:27+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
> *PING**2
<http://mid.mail-archive.com/87h7y3ofad.fsf@euler.schwinge.homeip.net>?
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
I have now committed this patch as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373, reading "so we shall accept it" as approval … On 4/1/20 9:07 AM, Thomas Schwinge wrote: > Even if the ICE is then fixed, should we still keep > <https://gcc.gnu.org/PR94120> open (with a note) until > <https://github.com/OpenACC/openacc-spec/issues/288> is > resolved/published? I decided to close it and mention the namelist issue and the now-fixed PR in PR84140 (which is about (non)static class member variables, which is also covered in Issue 288). > Regarding the C/C++ patch you posted: I'm not at all familiar with the > front ends' scoping implementation. But, given that your patch really > only touches the OpenACC 'declare' code paths, it can't cause any harm > otherwise, so we shall accept it. Maybe Jakub has any comments, though? >>> +c_check_oacc_same_scope (tree decl) >> Is the function really specific to OpenACC? If not, then "_oacc" >> could be dropped from its name. How about "c_check_current_scope"? > Yeah, that may be a good idea. Similar constructs seem to be used in a > few other places, though without the 'DECL_NAME' indirection. I now use c_check_in_current_scope. Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On Wed, Apr 8, 2020 at 12:55 AM Tobias Burnus <tobias@codesourcery.com> wrote: > > I have now committed this patch > as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373, On Linux/x86, I got FAIL: g++.dg/declare-pr94120.C -std=c++14 (test for excess errors) FAIL: g++.dg/declare-pr94120.C -std=c++17 (test for excess errors) FAIL: g++.dg/declare-pr94120.C -std=c++2a (test for excess errors) FAIL: g++.dg/declare-pr94120.C -std=c++98 (test for excess errors) > reading "so we shall accept it" as approval … > > On 4/1/20 9:07 AM, Thomas Schwinge wrote: > > > Even if the ICE is then fixed, should we still keep > > <https://gcc.gnu.org/PR94120> open (with a note) until > > <https://github.com/OpenACC/openacc-spec/issues/288> is > > resolved/published? > > I decided to close it and mention the namelist issue and > the now-fixed PR in PR84140 (which is about (non)static > class member variables, which is also covered in Issue 288). > > > Regarding the C/C++ patch you posted: I'm not at all familiar with the > > front ends' scoping implementation. But, given that your patch really > > only touches the OpenACC 'declare' code paths, it can't cause any harm > > otherwise, so we shall accept it. Maybe Jakub has any comments, though? > > >>> +c_check_oacc_same_scope (tree decl) > >> Is the function really specific to OpenACC? If not, then "_oacc" > >> could be dropped from its name. How about "c_check_current_scope"? > > Yeah, that may be a good idea. Similar constructs seem to be used in a > > few other places, though without the 'DECL_NAME' indirection. > > I now use c_check_in_current_scope. > > Tobias > ----------------- > Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany > Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
There was a glitch in the test case (location + dg-error), already fixed by Jakub (thanks!) in commit r10-7637-g08d1e7a5aabcf7eeac48bfd99deb80451b8f9974 Sorry, Tobias On 4/8/20 7:13 PM, H.J. Lu wrote: > On Wed, Apr 8, 2020 at 12:55 AM Tobias Burnus <tobias@codesourcery.com> wrote: >> I have now committed this patch >> as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373, > On Linux/x86, I got > > FAIL: g++.dg/declare-pr94120.C -std=c++14 (test for excess errors) > FAIL: g++.dg/declare-pr94120.C -std=c++17 (test for excess errors) > FAIL: g++.dg/declare-pr94120.C -std=c++2a (test for excess errors) > FAIL: g++.dg/declare-pr94120.C -std=c++98 (test for excess errors) > >> reading "so we shall accept it" as approval … >> >> On 4/1/20 9:07 AM, Thomas Schwinge wrote: >> >>> Even if the ICE is then fixed, should we still keep >>> <https://gcc.gnu.org/PR94120> open (with a note) until >>> <https://github.com/OpenACC/openacc-spec/issues/288> is >>> resolved/published? >> I decided to close it and mention the namelist issue and >> the now-fixed PR in PR84140 (which is about (non)static >> class member variables, which is also covered in Issue 288). >> >>> Regarding the C/C++ patch you posted: I'm not at all familiar with the >>> front ends' scoping implementation. But, given that your patch really >>> only touches the OpenACC 'declare' code paths, it can't cause any harm >>> otherwise, so we shall accept it. Maybe Jakub has any comments, though? >>>>> +c_check_oacc_same_scope (tree decl) >>>> Is the function really specific to OpenACC? If not, then "_oacc" >>>> could be dropped from its name. How about "c_check_current_scope"? >>> Yeah, that may be a good idea. Similar constructs seem to be used in a >>> few other places, though without the 'DECL_NAME' indirection. >> I now use c_check_in_current_scope. >> >> Tobias >> ----------------- >> Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany >> Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter > > ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi Tobias! On 2020-04-08T09:55:24+0200, Tobias Burnus <tobias@codesourcery.com> wrote: > I have now committed this patch > as r10-7614-g13e41d8b9d3d7598c72c38acc86a3d97046c8373, > reading "so we shall accept it" as approval … That's OK. As I said: "I'm not at all familiar with the front ends' scoping implementation", and don't currently have time to learn about that. So, either you're confident that you're doing the right things there (which I shall assume, given that you didn't explicitly ask whether you're doing the right things), or you need to wait for somebody else to review the patch. Reviewers don't know everything -- certainly I don't ;-) -- and don't have bandwidth to learn everything. And even if a proper review is done, often enough the reviewer fails to foresee the one detail that should've been caught. Hence, I'm happy to incrementally improve things, as long as a patch isn't regressing any (or, too many) other things, and isn't conceptually questionable (both of which doesn't apply here). That's why I said "we shall accept [the patch]" -- meaning "approved for commit, without my proper review". Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi Tobias! I do see the new 'libgomp.oacc-c++/declare-pr94120.C' (see below, for reference) FAIL for '-foffload=nvptx-none' execution testing. The reason is that the 'C' array doesn't have the content it's checked to have. Now, my question, shouldn't it be changed like the attached patch, or similar, because we actually first need to return from function 'f' in order for the 'copyout(C)' to be executed? Or, otherwise, what's the use of the 'copyout' clause with OpenACC 'declare'? I suppose it's only meant to be used with function arguments, because for every locally-defined variable (like 'C' in 'libgomp.oacc-c++/declare-pr94120.C'), that variable will be gone just after the 'copyout' is executed, so the 'copyout' conceptually equals/acts as 'create' in that case. However, OpenACC explicitly does allow 'copyout' in certain scenarios. Do you think my interpretation is correct, or what am I missing? (In case the answer isn't obvious to you, too, please file an issue with OpenACC, linking to this email for reference.) However, GCC doesn't like my changes either; actually two errors are diagnosed: [...]/libgomp.oacc-c++/declare-pr94120.C: In function ‘void f(int*)’: [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: array section in ‘#pragma acc declare’ 21 | #pragma acc declare copyout (C[0:N]) | ^ [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: ‘C’ must be a variable declared in the same scope as ‘#pragma acc declare’ Please have a look, and as necessary file GCC PRs, and XFAIL the 'libgomp.oacc-c++/declare-pr94120.C' execution testing for '-DACC_MEM_SHARED=0'. Grüße Thomas On 2020-03-11T13:28:44+0100, Tobias Burnus <tobias@codesourcery.com> wrote: > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C > @@ -0,0 +1,57 @@ > +#include <openacc.h> > +#include <stdlib.h> > + > +#define N 8 > + > +namespace one { > + int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 }; > + #pragma acc declare copyin (A) > +}; > + > +namespace outer { > + namespace inner { > + int B[N]; > + #pragma acc declare create (B) > + }; > +}; > + > +static void > +f (void) > +{ > + int i; > + int C[N]; > + #pragma acc declare copyout (C) > + > + if (!acc_is_present (&one::A, sizeof (one::A))) > + abort (); > + > + if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B))) > + abort (); > + > +#pragma acc parallel > + for (i = 0; i < N; i++) > + { > + outer::inner::B[i] = one::A[i]; > + C[i] = outer::inner::B[i]; > + } > + > + for (i = 0; i < N; i++) > + { > + if (C[i] != i + 1) > + abort (); > + } > + > +#pragma acc parallel > + for (i = 0; i < N; i++) > + if (outer::inner::B[i] != i + 1) > + abort (); > +} > + > + > +int > +main (int argc, char **argv) > +{ > + f (); > + > + return 0; > +} ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
Hi Thomas, I have now fixed it temporarily differently – without actually testing that feature. – See attachment. For the proper fix, it would help to get the semantic right in OpenACC itself (→ ongoing discussion). I think the patch for PR94120 is probably not completely right – and should be fixed for GCC 10, but I think one should wait for the next two OpenACC meetings to get converged. One presumably needs to handle 'device_resident' and 'link' differently from 'copy* etc. I added a note to PR 94140. Regards, Tobias On 4/10/20 3:20 PM, Thomas Schwinge wrote: > Hi Tobias! > > I do see the new 'libgomp.oacc-c++/declare-pr94120.C' (see below, for > reference) FAIL for '-foffload=nvptx-none' execution testing. The reason > is that the 'C' array doesn't have the content it's checked to have. > > Now, my question, shouldn't it be changed like the attached patch, or > similar, because we actually first need to return from function 'f' in > order for the 'copyout(C)' to be executed? > > Or, otherwise, what's the use of the 'copyout' clause with OpenACC > 'declare'? I suppose it's only meant to be used with function arguments, > because for every locally-defined variable (like 'C' in > 'libgomp.oacc-c++/declare-pr94120.C'), that variable will be gone just > after the 'copyout' is executed, so the 'copyout' conceptually > equals/acts as 'create' in that case. However, OpenACC explicitly does > allow 'copyout' in certain scenarios. Do you think my interpretation is > correct, or what am I missing? (In case the answer isn't obvious to you, > too, please file an issue with OpenACC, linking to this email for > reference.) > > However, GCC doesn't like my changes either; actually two errors are > diagnosed: > > [...]/libgomp.oacc-c++/declare-pr94120.C: In function ‘void f(int*)’: > [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: array section in ‘#pragma acc declare’ > 21 | #pragma acc declare copyout (C[0:N]) > | ^ > [...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: ‘C’ must be a variable declared in the same scope as ‘#pragma acc declare’ > > Please have a look, and as necessary file GCC PRs, and XFAIL the > 'libgomp.oacc-c++/declare-pr94120.C' execution testing for > '-DACC_MEM_SHARED=0'. > > > Grüße > Thomas > > > On 2020-03-11T13:28:44+0100, Tobias Burnus<tobias@codesourcery.com> wrote: >> --- /dev/null >> +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C >> @@ -0,0 +1,57 @@ >> +#include <openacc.h> >> +#include <stdlib.h> >> + >> +#define N 8 >> + >> +namespace one { >> + int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 }; >> + #pragma acc declare copyin (A) >> +}; >> + >> +namespace outer { >> + namespace inner { >> + int B[N]; >> + #pragma acc declare create (B) >> + }; >> +}; >> + >> +static void >> +f (void) >> +{ >> + int i; >> + int C[N]; >> + #pragma acc declare copyout (C) >> + >> + if (!acc_is_present (&one::A, sizeof (one::A))) >> + abort (); >> + >> + if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B))) >> + abort (); >> + >> +#pragma acc parallel >> + for (i = 0; i < N; i++) >> + { >> + outer::inner::B[i] = one::A[i]; >> + C[i] = outer::inner::B[i]; >> + } >> + >> + for (i = 0; i < N; i++) >> + { >> + if (C[i] != i + 1) >> + abort (); >> + } >> + >> +#pragma acc parallel >> + for (i = 0; i < N; i++) >> + if (outer::inner::B[i] != i + 1) >> + abort (); >> +} >> + >> + >> +int >> +main (int argc, char **argv) >> +{ >> + f (); >> + >> + return 0; >> +} ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120) 2020-10-11 Tobias Burnus <tobias@codesourcery.com> gcc/c/ PR middle-end/94120 * c-decl.c (c_check_oacc_same_scope): New function. * c-tree.h (c_check_oacc_same_scope): Declare it. * c-parser.c (c_parser_oacc_declare): Add check that variables are declared in the same scope as the directive. Fix handling of namespace vars. gcc/c/ PR middle-end/94120 * paser.c (cp_parser_oacc_declare): Add check that variables are declared in the same scope as the directive. gcc/testsuite/ PR middle-end/94120 * c-c++-common/goacc/declare-pr94120.c: New. * g++.dg/declare-pr94120.C: New. libgomp/testsuite/ PR middle-end/94120 * libgomp.oacc-c++/declare-pr94120.C: New. gcc/c/c-decl.c | 8 +++ gcc/c/c-parser.c | 9 ++++ gcc/c/c-tree.h | 1 + gcc/cp/parser.c | 21 +++++++- gcc/testsuite/c-c++-common/goacc/declare-pr94120.c | 23 +++++++++ gcc/testsuite/g++.dg/declare-pr94120.C | 30 ++++++++++++ .../testsuite/libgomp.oacc-c++/declare-pr94120.C | 57 ++++++++++++++++++++++ 7 files changed, 147 insertions(+), 2 deletions(-) diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c index c819fd0d0d5..eda95d664de 100644 --- a/gcc/c/c-decl.c +++ b/gcc/c/c-decl.c @@ -12016,4 +12016,12 @@ c_check_omp_declare_reduction_r (tree *tp, int *, void *data) return NULL_TREE; } + +bool +c_check_oacc_same_scope (tree decl) +{ + struct c_binding *b = I_SYMBOL_BINDING (DECL_NAME (decl)); + return b != NULL && B_IN_CURRENT_SCOPE (b); +} + #include "gt-c-c-decl.h" diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 1e8f2f7108d..63e8ab0ad17 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -16573,6 +16573,15 @@ c_parser_oacc_declare (c_parser *parser) break; } + if (!c_check_oacc_same_scope (decl)) + { + error_at (loc, + "%qD must be a variable declared in the same scope as " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) || lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index 71229927cb6..6d578705d77 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -789,6 +789,7 @@ extern tree c_omp_reduction_id (enum tree_code, tree); extern tree c_omp_reduction_decl (tree); extern tree c_omp_reduction_lookup (tree, tree); extern tree c_check_omp_declare_reduction_r (tree *, int *, void *); +extern bool c_check_oacc_same_scope (tree); extern void c_pushtag (location_t, tree, tree); extern void c_bind (location_t, tree, bool); extern bool tag_exists_p (enum tree_code, tree); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 24f71671469..8f09eb0d375 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40722,6 +40722,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) { tree clauses, stmt; bool error = false; + bool found_in_scope = global_bindings_p (); clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK, "#pragma acc declare", pragma_tok, true); @@ -40794,6 +40795,22 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) break; } + if (!found_in_scope) + for (tree d = current_binding_level->names; d; d = TREE_CHAIN (d)) + if (d == decl) + { + found_in_scope = true; + break; + } + if (!found_in_scope) + { + error_at (loc, + "%qD must be a variable declared in the same scope as " + "%<#pragma acc declare%>", decl); + error = true; + continue; + } + if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)) || lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) @@ -40815,7 +40832,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) DECL_ATTRIBUTES (decl) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl)); - if (global_bindings_p ()) + if (current_binding_level->kind == sk_namespace) { symtab_node *node = symtab_node::get (decl); if (node != NULL) @@ -40832,7 +40849,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) } } - if (error || global_bindings_p ()) + if (error || current_binding_level->kind == sk_namespace) return NULL_TREE; stmt = make_node (OACC_DECLARE); diff --git a/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c b/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c new file mode 100644 index 00000000000..21b2cc14fc7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/declare-pr94120.c @@ -0,0 +1,23 @@ +/* { dg-do compile } */ + +/* PR middle-end/94120 */ + +void foo() +{ + int foo; + { + #pragma acc declare copy(foo) /* { dg-error "'foo' must be a variable declared in the same scope as '#pragma acc declare'" } */ + } +} + +void +f_data (void) +{ + int B[10]; +#pragma acc data + { +# pragma acc declare copy(B) /* { dg-error "'B' must be a variable declared in the same scope as '#pragma acc declare'" } */ + for (int i = 0; i < 10; i++) + B[i] = -i; + } +} diff --git a/gcc/testsuite/g++.dg/declare-pr94120.C b/gcc/testsuite/g++.dg/declare-pr94120.C new file mode 100644 index 00000000000..8515c4ff875 --- /dev/null +++ b/gcc/testsuite/g++.dg/declare-pr94120.C @@ -0,0 +1,30 @@ +/* { dg-do compile } */ + +/* PR middle-end/94120 */ + +int b[8]; +#pragma acc declare create (b) + +namespace my { + int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; + #pragma acc declare copyin (d) +}; + +namespace outer { + namespace inner { + int e[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; + #pragma acc declare copyin (e) + }; +}; + +int f[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; +namespace my { + #pragma acc declare copyin (f) +}; + +namespace outer { + int g[8] = { 1, 2, 3, 4, 5, 6, 7, 8 }; + namespace inner { + #pragma acc declare copyin (g) + }; +}; diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C new file mode 100644 index 00000000000..1e1254187ea --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C @@ -0,0 +1,57 @@ +#include <openacc.h> +#include <stdlib.h> + +#define N 8 + +namespace one { + int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 }; + #pragma acc declare copyin (A) +}; + +namespace outer { + namespace inner { + int B[N]; + #pragma acc declare create (B) + }; +}; + +static void +f (void) +{ + int i; + int C[N]; + #pragma acc declare copyout (C) + + if (!acc_is_present (&one::A, sizeof (one::A))) + abort (); + + if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B))) + abort (); + +#pragma acc parallel + for (i = 0; i < N; i++) + { + outer::inner::B[i] = one::A[i]; + C[i] = outer::inner::B[i]; + } + + for (i = 0; i < N; i++) + { + if (C[i] != i + 1) + abort (); + } + +#pragma acc parallel + for (i = 0; i < N; i++) + if (outer::inner::B[i] != i + 1) + abort (); +} + + +int +main (int argc, char **argv) +{ + f (); + + return 0; +}