diff mbox series

[C/C++,OpenACC] Reject vars of different scope in acc declare (PR94120)

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

Commit Message

Tobias Burnus March 11, 2020, 1:28 p.m. UTC 
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

Comments

Frederik Harwath March 12, 2020, 2:43 p.m. UTC | #1 
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
Tobias Burnus March 24, 2020, 10:18 a.m. UTC | #2 
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
Thomas Schwinge April 1, 2020, 7:07 a.m. UTC | #3 
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
Tobias Burnus April 7, 2020, 7:21 a.m. UTC | #4 
*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
Thomas Schwinge April 7, 2020, 7:28 a.m. UTC | #5 
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
Tobias Burnus April 8, 2020, 7:55 a.m. UTC | #6 
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
Li, Pan2 via Gcc-patches April 8, 2020, 5:13 p.m. UTC | #7 
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
Tobias Burnus April 8, 2020, 6:17 p.m. UTC | #8 
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
Thomas Schwinge April 10, 2020, 1:11 p.m. UTC | #9 
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
Thomas Schwinge April 10, 2020, 1:20 p.m. UTC | #10 
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
Tobias Burnus April 20, 2020, 2:18 p.m. UTC | #11 
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
diff mbox series

Patch

[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;
+}