(GCC) 13.0.1: internal compiler error

(GCC) 13.0.1: internal compiler error

[Patrick Begou]

[-- Attachment #1: Type: text/plain, Size: 3389 bytes --]

Hi,

I have built this morning the latest gfortran from a git clone:

GNU Fortran (GCC) 13.0.1 20230421 (prerelease)

I'm trying this compiler on a large and complexe Fortran90 code with 
offloading using OpenACC.

At this time:

- code compiles with nvfortran and runs on A100 GPU.

- code compiles with Cray Fortran (with some difficulties) but do not 
run on MI250 GPU (we are tacking the problem, a segfault if openacc is 
set on)

- code compile with GNU GCC 13 without -fopenacc option and runs on cpu 
(Epyc2 7302)

- a basic test-code using OpenACC compiles and run on the GPU.

- compiling my large code with gcc 13.0.1 using -fopenacc for A100 GPU 
produce an internal error in the compiler :


transforms_defs_m.f90:354:53:

   354 |             !$acc enter data attach(atransform2%next)
       |                                                     ^
internal compiler error: in omp_group_base, at gimplify.cc:9412
0xa830c6 omp_group_base
     ../../gcc/gcc/gimplify.cc:9412
0xa830c6 omp_index_mapping_groups_1
     ../../gcc/gcc/gimplify.cc:9441
0xa833c7 omp_index_mapping_groups
     ../../gcc/gcc/gimplify.cc:9502
0xa96a9a gimplify_scan_omp_clauses
     ../../gcc/gcc/gimplify.cc:10802
0xa8660d gimplify_omp_target_update
     ../../gcc/gcc/gimplify.cc:15563
0xa8660d gimplify_expr(tree_node**, gimple**, gimple**, bool 
(*)(tree_node*), int)
     ../../gcc/gcc/gimplify.cc:16928
0xa89826 gimplify_stmt(tree_node**, gimple**)
     ../../gcc/gcc/gimplify.cc:7219
0xa875a3 gimplify_statement_list
     ../../gcc/gcc/gimplify.cc:2019
0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool 
(*)(tree_node*), int)
     ../../gcc/gcc/gimplify.cc:16821
0xa89826 gimplify_stmt(tree_node**, gimple**)
     ../../gcc/gcc/gimplify.cc:7219
0xa86e8a gimplify_and_add(tree_node*, gimple**)
     ../../gcc/gcc/gimplify.cc:492
0xa86e8a gimplify_loop_expr
     ../../gcc/gcc/gimplify.cc:1993
0xa86e8a gimplify_expr(tree_node**, gimple**, gimple**, bool 
(*)(tree_node*), int)
     ../../gcc/gcc/gimplify.cc:16581
0xa89826 gimplify_stmt(tree_node**, gimple**)
     ../../gcc/gcc/gimplify.cc:7219
0xa875a3 gimplify_statement_list
     ../../gcc/gcc/gimplify.cc:2019
0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool 
(*)(tree_node*), int)
     ../../gcc/gcc/gimplify.cc:16821
0xa89826 gimplify_stmt(tree_node**, gimple**)
     ../../gcc/gcc/gimplify.cc:7219
0xa89d2b gimplify_bind_expr
     ../../gcc/gcc/gimplify.cc:1430
0xa86d8e gimplify_expr(tree_node**, gimple**, gimple**, bool 
(*)(tree_node*), int)
     ../../gcc/gcc/gimplify.cc:16577
0xa89826 gimplify_stmt(tree_node**, gimple**)
     ../../gcc/gcc/gimplify.cc:7219
Please submit a full bug report, with preprocessed source (by using 
-freport-bug).


Options used (I've just added  -fopenacc for moving from cpu version to 
OpenACC):

-fopenacc -freport-bug -g -fpic -x f95-cpp-input -std=gnu -ffree-form 
-fall-intrinsics -fallow-argument-mismatch -Wall -Wextra -W 
-Wno-unused-function -Wno-compare-reals -fno-omit-frame-pointer  -O3  
-ftree-vectorize -ffast-math -funroll-loops -pipe

No additionnal files a produced with -freport-bug.

In attachment the script used to build the compiler.

Let me know how I can help with informations to improve Gnu fortran 
compilers.

Patrick

[-- Attachment #2: build_all.sh --]
[-- Type: application/x-shellscript, Size: 1854 bytes --]

Re: (GCC) 13.0.1: internal compiler error

[Harald Anlauf]

Hi Patrick,

I did not see any similar report in bugzilla, so could you please
open a PR and attach a self-contained reproducer?  Ideally the
reproducer would be reduced to simplify the analysis for those
familiar with the status of the OpenACC implementation.

Thanks,
Harald

Am 21.04.23 um 17:13 schrieb Patrick Begou:
> Hi,
>
> I have built this morning the latest gfortran from a git clone:
>
> GNU Fortran (GCC) 13.0.1 20230421 (prerelease)
>
> I'm trying this compiler on a large and complexe Fortran90 code with
> offloading using OpenACC.
>
> At this time:
>
> - code compiles with nvfortran and runs on A100 GPU.
>
> - code compiles with Cray Fortran (with some difficulties) but do not
> run on MI250 GPU (we are tacking the problem, a segfault if openacc is
> set on)
>
> - code compile with GNU GCC 13 without -fopenacc option and runs on cpu
> (Epyc2 7302)
>
> - a basic test-code using OpenACC compiles and run on the GPU.
>
> - compiling my large code with gcc 13.0.1 using -fopenacc for A100 GPU
> produce an internal error in the compiler :
>
>
> transforms_defs_m.f90:354:53:
>
>    354 |             !$acc enter data attach(atransform2%next)
>        |                                                     ^
> internal compiler error: in omp_group_base, at gimplify.cc:9412
> 0xa830c6 omp_group_base
>      ../../gcc/gcc/gimplify.cc:9412
> 0xa830c6 omp_index_mapping_groups_1
>      ../../gcc/gcc/gimplify.cc:9441
> 0xa833c7 omp_index_mapping_groups
>      ../../gcc/gcc/gimplify.cc:9502
> 0xa96a9a gimplify_scan_omp_clauses
>      ../../gcc/gcc/gimplify.cc:10802
> 0xa8660d gimplify_omp_target_update
>      ../../gcc/gcc/gimplify.cc:15563
> 0xa8660d gimplify_expr(tree_node**, gimple**, gimple**, bool
> (*)(tree_node*), int)
>      ../../gcc/gcc/gimplify.cc:16928
> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>      ../../gcc/gcc/gimplify.cc:7219
> 0xa875a3 gimplify_statement_list
>      ../../gcc/gcc/gimplify.cc:2019
> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
> (*)(tree_node*), int)
>      ../../gcc/gcc/gimplify.cc:16821
> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>      ../../gcc/gcc/gimplify.cc:7219
> 0xa86e8a gimplify_and_add(tree_node*, gimple**)
>      ../../gcc/gcc/gimplify.cc:492
> 0xa86e8a gimplify_loop_expr
>      ../../gcc/gcc/gimplify.cc:1993
> 0xa86e8a gimplify_expr(tree_node**, gimple**, gimple**, bool
> (*)(tree_node*), int)
>      ../../gcc/gcc/gimplify.cc:16581
> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>      ../../gcc/gcc/gimplify.cc:7219
> 0xa875a3 gimplify_statement_list
>      ../../gcc/gcc/gimplify.cc:2019
> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
> (*)(tree_node*), int)
>      ../../gcc/gcc/gimplify.cc:16821
> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>      ../../gcc/gcc/gimplify.cc:7219
> 0xa89d2b gimplify_bind_expr
>      ../../gcc/gcc/gimplify.cc:1430
> 0xa86d8e gimplify_expr(tree_node**, gimple**, gimple**, bool
> (*)(tree_node*), int)
>      ../../gcc/gcc/gimplify.cc:16577
> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>      ../../gcc/gcc/gimplify.cc:7219
> Please submit a full bug report, with preprocessed source (by using
> -freport-bug).
>
>
> Options used (I've just added  -fopenacc for moving from cpu version to
> OpenACC):
>
> -fopenacc -freport-bug -g -fpic -x f95-cpp-input -std=gnu -ffree-form
> -fall-intrinsics -fallow-argument-mismatch -Wall -Wextra -W
> -Wno-unused-function -Wno-compare-reals -fno-omit-frame-pointer  -O3
> -ftree-vectorize -ffast-math -funroll-loops -pipe
>
> No additionnal files a produced with -freport-bug.
>
> In attachment the script used to build the compiler.
>
> Let me know how I can help with informations to improve Gnu fortran
> compilers.
>
> Patrick

Re: (GCC) 13.0.1: internal compiler error

[Patrick Begou]

Hi Harald

as I said, it is a large, massively parallel fortran code: more than 700 
files, some with several thousands of lines. It could be difficult to 
create a small reproducer but I will try if the problem is not known as 
an other git branch of this code also create an internal error on 
another file.

Best regards

Patrick

Le 24/04/2023 à 19:27, Harald Anlauf a écrit :
> Hi Patrick,
>
> I did not see any similar report in bugzilla, so could you please
> open a PR and attach a self-contained reproducer?  Ideally the
> reproducer would be reduced to simplify the analysis for those
> familiar with the status of the OpenACC implementation.
>
> Thanks,
> Harald
>
> Am 21.04.23 um 17:13 schrieb Patrick Begou:
>> Hi,
>>
>> I have built this morning the latest gfortran from a git clone:
>>
>> GNU Fortran (GCC) 13.0.1 20230421 (prerelease)
>>
>> I'm trying this compiler on a large and complexe Fortran90 code with
>> offloading using OpenACC.
>>
>> At this time:
>>
>> - code compiles with nvfortran and runs on A100 GPU.
>>
>> - code compiles with Cray Fortran (with some difficulties) but do not
>> run on MI250 GPU (we are tacking the problem, a segfault if openacc is
>> set on)
>>
>> - code compile with GNU GCC 13 without -fopenacc option and runs on cpu
>> (Epyc2 7302)
>>
>> - a basic test-code using OpenACC compiles and run on the GPU.
>>
>> - compiling my large code with gcc 13.0.1 using -fopenacc for A100 GPU
>> produce an internal error in the compiler :
>>
>>
>> transforms_defs_m.f90:354:53:
>>
>>    354 |             !$acc enter data attach(atransform2%next)
>>        |                                                     ^
>> internal compiler error: in omp_group_base, at gimplify.cc:9412
>> 0xa830c6 omp_group_base
>>      ../../gcc/gcc/gimplify.cc:9412
>> 0xa830c6 omp_index_mapping_groups_1
>>      ../../gcc/gcc/gimplify.cc:9441
>> 0xa833c7 omp_index_mapping_groups
>>      ../../gcc/gcc/gimplify.cc:9502
>> 0xa96a9a gimplify_scan_omp_clauses
>>      ../../gcc/gcc/gimplify.cc:10802
>> 0xa8660d gimplify_omp_target_update
>>      ../../gcc/gcc/gimplify.cc:15563
>> 0xa8660d gimplify_expr(tree_node**, gimple**, gimple**, bool
>> (*)(tree_node*), int)
>>      ../../gcc/gcc/gimplify.cc:16928
>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>      ../../gcc/gcc/gimplify.cc:7219
>> 0xa875a3 gimplify_statement_list
>>      ../../gcc/gcc/gimplify.cc:2019
>> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
>> (*)(tree_node*), int)
>>      ../../gcc/gcc/gimplify.cc:16821
>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>      ../../gcc/gcc/gimplify.cc:7219
>> 0xa86e8a gimplify_and_add(tree_node*, gimple**)
>>      ../../gcc/gcc/gimplify.cc:492
>> 0xa86e8a gimplify_loop_expr
>>      ../../gcc/gcc/gimplify.cc:1993
>> 0xa86e8a gimplify_expr(tree_node**, gimple**, gimple**, bool
>> (*)(tree_node*), int)
>>      ../../gcc/gcc/gimplify.cc:16581
>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>      ../../gcc/gcc/gimplify.cc:7219
>> 0xa875a3 gimplify_statement_list
>>      ../../gcc/gcc/gimplify.cc:2019
>> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
>> (*)(tree_node*), int)
>>      ../../gcc/gcc/gimplify.cc:16821
>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>      ../../gcc/gcc/gimplify.cc:7219
>> 0xa89d2b gimplify_bind_expr
>>      ../../gcc/gcc/gimplify.cc:1430
>> 0xa86d8e gimplify_expr(tree_node**, gimple**, gimple**, bool
>> (*)(tree_node*), int)
>>      ../../gcc/gcc/gimplify.cc:16577
>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>      ../../gcc/gcc/gimplify.cc:7219
>> Please submit a full bug report, with preprocessed source (by using
>> -freport-bug).
>>
>>
>> Options used (I've just added  -fopenacc for moving from cpu version to
>> OpenACC):
>>
>> -fopenacc -freport-bug -g -fpic -x f95-cpp-input -std=gnu -ffree-form
>> -fall-intrinsics -fallow-argument-mismatch -Wall -Wextra -W
>> -Wno-unused-function -Wno-compare-reals -fno-omit-frame-pointer -O3
>> -ftree-vectorize -ffast-math -funroll-loops -pipe
>>
>> No additionnal files a produced with -freport-bug.
>>
>> In attachment the script used to build the compiler.
>>
>> Let me know how I can help with informations to improve Gnu fortran
>> compilers.
>>
>> Patrick
>

Re: (GCC) 13.0.1: internal compiler error

[Bernhard Reutner-Fischer]

Cc:ing Thomas, who knows openacc better.

There is a devel/omp/gcc-12 branch you might want to try. I don't know
how different that branch is wrt openacc.

HTH,

On Mon, 24 Apr 2023 19:39:15
+0200 Patrick Begou via Fortran <fortran@gcc.gnu.org> wrote:

> Hi Harald
> 
> as I said, it is a large, massively parallel fortran code: more than 700 
> files, some with several thousands of lines. It could be difficult to 
> create a small reproducer but I will try if the problem is not known as 
> an other git branch of this code also create an internal error on 
> another file.
> 
> Best regards
> 
> Patrick
> 
> Le 24/04/2023 à 19:27, Harald Anlauf a écrit :
> > Hi Patrick,
> >
> > I did not see any similar report in bugzilla, so could you please
> > open a PR and attach a self-contained reproducer?  Ideally the
> > reproducer would be reduced to simplify the analysis for those
> > familiar with the status of the OpenACC implementation.
> >
> > Thanks,
> > Harald
> >
> > Am 21.04.23 um 17:13 schrieb Patrick Begou:  
> >> Hi,
> >>
> >> I have built this morning the latest gfortran from a git clone:
> >>
> >> GNU Fortran (GCC) 13.0.1 20230421 (prerelease)
> >>
> >> I'm trying this compiler on a large and complexe Fortran90 code with
> >> offloading using OpenACC.
> >>
> >> At this time:
> >>
> >> - code compiles with nvfortran and runs on A100 GPU.
> >>
> >> - code compiles with Cray Fortran (with some difficulties) but do not
> >> run on MI250 GPU (we are tacking the problem, a segfault if openacc is
> >> set on)
> >>
> >> - code compile with GNU GCC 13 without -fopenacc option and runs on cpu
> >> (Epyc2 7302)
> >>
> >> - a basic test-code using OpenACC compiles and run on the GPU.
> >>
> >> - compiling my large code with gcc 13.0.1 using -fopenacc for A100 GPU
> >> produce an internal error in the compiler :
> >>
> >>
> >> transforms_defs_m.f90:354:53:
> >>
> >>    354 |             !$acc enter data attach(atransform2%next)
> >>        |                                                     ^
> >> internal compiler error: in omp_group_base, at gimplify.cc:9412
> >> 0xa830c6 omp_group_base
> >>      ../../gcc/gcc/gimplify.cc:9412
> >> 0xa830c6 omp_index_mapping_groups_1
> >>      ../../gcc/gcc/gimplify.cc:9441
> >> 0xa833c7 omp_index_mapping_groups
> >>      ../../gcc/gcc/gimplify.cc:9502
> >> 0xa96a9a gimplify_scan_omp_clauses
> >>      ../../gcc/gcc/gimplify.cc:10802
> >> 0xa8660d gimplify_omp_target_update
> >>      ../../gcc/gcc/gimplify.cc:15563
> >> 0xa8660d gimplify_expr(tree_node**, gimple**, gimple**, bool
> >> (*)(tree_node*), int)
> >>      ../../gcc/gcc/gimplify.cc:16928
> >> 0xa89826 gimplify_stmt(tree_node**, gimple**)
> >>      ../../gcc/gcc/gimplify.cc:7219
> >> 0xa875a3 gimplify_statement_list
> >>      ../../gcc/gcc/gimplify.cc:2019
> >> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
> >> (*)(tree_node*), int)
> >>      ../../gcc/gcc/gimplify.cc:16821
> >> 0xa89826 gimplify_stmt(tree_node**, gimple**)
> >>      ../../gcc/gcc/gimplify.cc:7219
> >> 0xa86e8a gimplify_and_add(tree_node*, gimple**)
> >>      ../../gcc/gcc/gimplify.cc:492
> >> 0xa86e8a gimplify_loop_expr
> >>      ../../gcc/gcc/gimplify.cc:1993
> >> 0xa86e8a gimplify_expr(tree_node**, gimple**, gimple**, bool
> >> (*)(tree_node*), int)
> >>      ../../gcc/gcc/gimplify.cc:16581
> >> 0xa89826 gimplify_stmt(tree_node**, gimple**)
> >>      ../../gcc/gcc/gimplify.cc:7219
> >> 0xa875a3 gimplify_statement_list
> >>      ../../gcc/gcc/gimplify.cc:2019
> >> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
> >> (*)(tree_node*), int)
> >>      ../../gcc/gcc/gimplify.cc:16821
> >> 0xa89826 gimplify_stmt(tree_node**, gimple**)
> >>      ../../gcc/gcc/gimplify.cc:7219
> >> 0xa89d2b gimplify_bind_expr
> >>      ../../gcc/gcc/gimplify.cc:1430
> >> 0xa86d8e gimplify_expr(tree_node**, gimple**, gimple**, bool
> >> (*)(tree_node*), int)
> >>      ../../gcc/gcc/gimplify.cc:16577
> >> 0xa89826 gimplify_stmt(tree_node**, gimple**)
> >>      ../../gcc/gcc/gimplify.cc:7219
> >> Please submit a full bug report, with preprocessed source (by using
> >> -freport-bug).
> >>
> >>
> >> Options used (I've just added  -fopenacc for moving from cpu version to
> >> OpenACC):
> >>
> >> -fopenacc -freport-bug -g -fpic -x f95-cpp-input -std=gnu -ffree-form
> >> -fall-intrinsics -fallow-argument-mismatch -Wall -Wextra -W
> >> -Wno-unused-function -Wno-compare-reals -fno-omit-frame-pointer -O3
> >> -ftree-vectorize -ffast-math -funroll-loops -pipe
> >>
> >> No additionnal files a produced with -freport-bug.
> >>
> >> In attachment the script used to build the compiler.
> >>
> >> Let me know how I can help with informations to improve Gnu fortran
> >> compilers.
> >>
> >> Patrick  
> >  
> 

Re: (GCC) 13.0.1: internal compiler error

[Patrick Begou]

Hi,

I've juste created a bug report with a small fortran code to reproduce 
the internal error in gfortran.
Building a new compiler version from devel/omp/gcc-12 branch is in progress.

Patrick

Le 24/04/2023 à 20:29, Bernhard Reutner-Fischer a écrit :
> Cc:ing Thomas, who knows openacc better.
>
> There is a devel/omp/gcc-12 branch you might want to try. I don't know
> how different that branch is wrt openacc.
>
> HTH,
>
> On Mon, 24 Apr 2023 19:39:15
> +0200 Patrick Begou via Fortran <fortran@gcc.gnu.org> wrote:
>
>> Hi Harald
>>
>> as I said, it is a large, massively parallel fortran code: more than 700
>> files, some with several thousands of lines. It could be difficult to
>> create a small reproducer but I will try if the problem is not known as
>> an other git branch of this code also create an internal error on
>> another file.
>>
>> Best regards
>>
>> Patrick
>>
>> Le 24/04/2023 à 19:27, Harald Anlauf a écrit :
>>> Hi Patrick,
>>>
>>> I did not see any similar report in bugzilla, so could you please
>>> open a PR and attach a self-contained reproducer?  Ideally the
>>> reproducer would be reduced to simplify the analysis for those
>>> familiar with the status of the OpenACC implementation.
>>>
>>> Thanks,
>>> Harald
>>>
>>> Am 21.04.23 um 17:13 schrieb Patrick Begou:
>>>> Hi,
>>>>
>>>> I have built this morning the latest gfortran from a git clone:
>>>>
>>>> GNU Fortran (GCC) 13.0.1 20230421 (prerelease)
>>>>
>>>> I'm trying this compiler on a large and complexe Fortran90 code with
>>>> offloading using OpenACC.
>>>>
>>>> At this time:
>>>>
>>>> - code compiles with nvfortran and runs on A100 GPU.
>>>>
>>>> - code compiles with Cray Fortran (with some difficulties) but do not
>>>> run on MI250 GPU (we are tacking the problem, a segfault if openacc is
>>>> set on)
>>>>
>>>> - code compile with GNU GCC 13 without -fopenacc option and runs on cpu
>>>> (Epyc2 7302)
>>>>
>>>> - a basic test-code using OpenACC compiles and run on the GPU.
>>>>
>>>> - compiling my large code with gcc 13.0.1 using -fopenacc for A100 GPU
>>>> produce an internal error in the compiler :
>>>>
>>>>
>>>> transforms_defs_m.f90:354:53:
>>>>
>>>>     354 |             !$acc enter data attach(atransform2%next)
>>>>         |                                                     ^
>>>> internal compiler error: in omp_group_base, at gimplify.cc:9412
>>>> 0xa830c6 omp_group_base
>>>>       ../../gcc/gcc/gimplify.cc:9412
>>>> 0xa830c6 omp_index_mapping_groups_1
>>>>       ../../gcc/gcc/gimplify.cc:9441
>>>> 0xa833c7 omp_index_mapping_groups
>>>>       ../../gcc/gcc/gimplify.cc:9502
>>>> 0xa96a9a gimplify_scan_omp_clauses
>>>>       ../../gcc/gcc/gimplify.cc:10802
>>>> 0xa8660d gimplify_omp_target_update
>>>>       ../../gcc/gcc/gimplify.cc:15563
>>>> 0xa8660d gimplify_expr(tree_node**, gimple**, gimple**, bool
>>>> (*)(tree_node*), int)
>>>>       ../../gcc/gcc/gimplify.cc:16928
>>>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>>>       ../../gcc/gcc/gimplify.cc:7219
>>>> 0xa875a3 gimplify_statement_list
>>>>       ../../gcc/gcc/gimplify.cc:2019
>>>> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
>>>> (*)(tree_node*), int)
>>>>       ../../gcc/gcc/gimplify.cc:16821
>>>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>>>       ../../gcc/gcc/gimplify.cc:7219
>>>> 0xa86e8a gimplify_and_add(tree_node*, gimple**)
>>>>       ../../gcc/gcc/gimplify.cc:492
>>>> 0xa86e8a gimplify_loop_expr
>>>>       ../../gcc/gcc/gimplify.cc:1993
>>>> 0xa86e8a gimplify_expr(tree_node**, gimple**, gimple**, bool
>>>> (*)(tree_node*), int)
>>>>       ../../gcc/gcc/gimplify.cc:16581
>>>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>>>       ../../gcc/gcc/gimplify.cc:7219
>>>> 0xa875a3 gimplify_statement_list
>>>>       ../../gcc/gcc/gimplify.cc:2019
>>>> 0xa875a3 gimplify_expr(tree_node**, gimple**, gimple**, bool
>>>> (*)(tree_node*), int)
>>>>       ../../gcc/gcc/gimplify.cc:16821
>>>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>>>       ../../gcc/gcc/gimplify.cc:7219
>>>> 0xa89d2b gimplify_bind_expr
>>>>       ../../gcc/gcc/gimplify.cc:1430
>>>> 0xa86d8e gimplify_expr(tree_node**, gimple**, gimple**, bool
>>>> (*)(tree_node*), int)
>>>>       ../../gcc/gcc/gimplify.cc:16577
>>>> 0xa89826 gimplify_stmt(tree_node**, gimple**)
>>>>       ../../gcc/gcc/gimplify.cc:7219
>>>> Please submit a full bug report, with preprocessed source (by using
>>>> -freport-bug).
>>>>
>>>>
>>>> Options used (I've just added  -fopenacc for moving from cpu version to
>>>> OpenACC):
>>>>
>>>> -fopenacc -freport-bug -g -fpic -x f95-cpp-input -std=gnu -ffree-form
>>>> -fall-intrinsics -fallow-argument-mismatch -Wall -Wextra -W
>>>> -Wno-unused-function -Wno-compare-reals -fno-omit-frame-pointer -O3
>>>> -ftree-vectorize -ffast-math -funroll-loops -pipe
>>>>
>>>> No additionnal files a produced with -freport-bug.
>>>>
>>>> In attachment the script used to build the compiler.
>>>>
>>>> Let me know how I can help with informations to improve Gnu fortran
>>>> compilers.
>>>>
>>>> Patrick
>>>   

[PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622]

[Julian Brown]

This patch fixes several cases where multiple attach or detach mapping
nodes were being created for stand-alone attach or detach clauses
in Fortran.  After the introduction of stricter checking later during
compilation, these extra nodes could cause ICEs, as seen in the PR.

The patch also fixes cases that "happened to work" previously where
the user attaches/detaches a pointer to array using a descriptor, and
(I think!) the "_data" field has offset zero, hence the same address as
the descriptor as a whole.

Tested with offloading to nvptx. OK?

Thanks,

Julian

2023-04-27  Julian Brown  <julian@codesourcery.com>

	PR fortran/109622

gcc/fortran/
	* trans-openmp.cc (gfc_trans_omp_clauses): Attach/detach clause fixes.

gcc/testsuite/
	* gfortran.dg/goacc/attach-descriptor.f90: Adjust expected output.

libgomp/
	* testsuite/libgomp.fortran/pr109622.f90: New test.
	* testsuite/libgomp.fortran/pr109622-2.f90: New test.
	* testsuite/libgomp.fortran/pr109622-3.f90: New test.
---
 gcc/fortran/trans-openmp.cc                   | 36 +++++++++++++++++--
 .../gfortran.dg/goacc/attach-descriptor.f90   | 12 +++----
 .../testsuite/libgomp.fortran/pr109622-2.f90  | 32 +++++++++++++++++
 .../testsuite/libgomp.fortran/pr109622-3.f90  | 32 +++++++++++++++++
 .../testsuite/libgomp.fortran/pr109622.f90    | 32 +++++++++++++++++
 5 files changed, 135 insertions(+), 9 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.fortran/pr109622-2.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/pr109622-3.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/pr109622.f90

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 4ff9c59df5cb..dbb4a335ab57 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3388,6 +3388,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  gfc_add_block_to_block (block, &se.post);
 		  if (pointer || allocatable)
 		    {
+		      /* If it's a bare attach/detach clause, we just want
+			 to perform a single attach/detach operation, of the
+			 pointer itself, not of the pointed-to object.  */
+		      if (openacc
+			  && (n->u.map_op == OMP_MAP_ATTACH
+			      || n->u.map_op == OMP_MAP_DETACH))
+			{
+			  OMP_CLAUSE_SIZE (node) = size_zero_node;
+			  goto finalize_map_clause;
+			}
+
 		      node2 = build_omp_clause (input_location,
 						OMP_CLAUSE_MAP);
 		      gomp_map_kind kind
@@ -3458,6 +3469,19 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		    {
 		      if (pointer || (openacc && allocatable))
 			{
+			  /* If it's a bare attach/detach clause, we just want
+			     to perform a single attach/detach operation, of the
+			     pointer itself, not of the pointed-to object.  */
+			  if (openacc
+			      && (n->u.map_op == OMP_MAP_ATTACH
+				  || n->u.map_op == OMP_MAP_DETACH))
+			    {
+			      OMP_CLAUSE_DECL (node)
+				= build_fold_addr_expr (inner);
+			      OMP_CLAUSE_SIZE (node) = size_zero_node;
+			      goto finalize_map_clause;
+			    }
+
 			  tree data, size;
 
 			  if (lastref->u.c.component->ts.type == BT_CLASS)
@@ -3494,12 +3518,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  else if (lastref->type == REF_ARRAY
 			   && lastref->u.ar.type == AR_FULL)
 		    {
-		      /* Just pass the (auto-dereferenced) decl through for
-			 bare attach and detach clauses.  */
+		      /* Bare attach and detach clauses don't want any
+			 additional nodes.  */
 		      if (n->u.map_op == OMP_MAP_ATTACH
 			  || n->u.map_op == OMP_MAP_DETACH)
 			{
-			  OMP_CLAUSE_DECL (node) = inner;
+			  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
+			    {
+			      tree ptr = gfc_conv_descriptor_data_get (inner);
+			      OMP_CLAUSE_DECL (node) = ptr;
+			    }
+			  else
+			    OMP_CLAUSE_DECL (node) = inner;
 			  OMP_CLAUSE_SIZE (node) = size_zero_node;
 			  goto finalize_map_clause;
 			}
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 8c2ee4a5cca4..734afbe6ca48 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -11,19 +11,19 @@ program att
   integer, pointer :: myptr(:)
 
   !$acc enter data attach(myvar%arr2, myptr)
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_data map\\(attach:myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
 
   !$acc exit data detach(myvar%arr2, myptr)
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(detach:myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
 
   ! Test valid usage and processing of the finalize clause.
   !$acc exit data detach(myvar%arr2, myptr) finalize
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
   ! For array-descriptor detaches, we no longer generate a "release" mapping
   ! for the pointed-to data for gimplify.c to turn into "delete".  Make sure
   ! the mapping still isn't there.
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_detach:myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
 
 end program att
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-2.f90 b/libgomp/testsuite/libgomp.fortran/pr109622-2.f90
new file mode 100644
index 000000000000..8c5f373f39f7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr109622-2.f90
@@ -0,0 +1,32 @@
+! { dg-do run }
+
+type t
+integer :: foo
+integer, pointer :: bar
+end type t
+
+type(t) :: var
+integer, target :: tgt
+
+var%bar => tgt
+
+var%foo = 99
+tgt = 199
+
+!$acc enter data copyin(var)
+
+!$acc enter data attach(var%bar)
+
+!$acc serial
+var%foo = 5
+var%bar = 7
+!$acc end serial
+
+!$acc exit data detach(var%bar)
+
+!$acc exit data copyout(var)
+
+if (var%foo.ne.5) stop 1
+if (tgt.ne.7) stop 2
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-3.f90 b/libgomp/testsuite/libgomp.fortran/pr109622-3.f90
new file mode 100644
index 000000000000..3ee1b43a7464
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr109622-3.f90
@@ -0,0 +1,32 @@
+! { dg-do run }
+
+type t
+integer :: foo
+integer, pointer :: bar(:)
+end type t
+
+type(t) :: var
+integer, target :: tgt(20)
+
+var%bar => tgt
+
+var%foo = 99
+tgt = 199
+
+!$acc enter data copyin(var, tgt)
+
+!$acc enter data attach(var%bar)
+
+!$acc serial
+var%foo = 5
+var%bar = 7
+!$acc end serial
+
+!$acc exit data detach(var%bar)
+
+!$acc exit data copyout(var, tgt)
+
+if (var%foo.ne.5) stop 1
+if (any(tgt.ne.7)) stop 2
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622.f90 b/libgomp/testsuite/libgomp.fortran/pr109622.f90
new file mode 100644
index 000000000000..5b8c4102f768
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr109622.f90
@@ -0,0 +1,32 @@
+! { dg-do run }
+
+type t
+integer :: value
+type(t), pointer :: chain
+end type t
+
+type(t), target :: var, var2
+
+var%value = 99
+var2%value = 199
+
+var%chain => var2
+nullify(var2%chain)
+
+!$acc enter data copyin(var, var2)
+
+!$acc enter data attach(var%chain)
+
+!$acc serial
+var%value = 5
+var%chain%value = 7
+!$acc end serial
+
+!$acc exit data detach(var%chain)
+
+!$acc exit data copyout(var, var2)
+
+if (var%value.ne.5) stop 1
+if (var2%value.ne.7) stop 2
+
+end
-- 
2.29.2

Re: [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622]

[Tobias Burnus]

On 27.04.23 20:36, Julian Brown wrote:
> This patch fixes several cases where multiple attach or detach mapping
> nodes were being created for stand-alone attach or detach clauses
> in Fortran.  After the introduction of stricter checking later during
> compilation, these extra nodes could cause ICEs, as seen in the PR.
>
> The patch also fixes cases that "happened to work" previously where
> the user attaches/detaches a pointer to array using a descriptor, and
> (I think!) the "_data" field has offset zero, hence the same address as
> the descriptor as a whole.
>
> Tested with offloading to nvptx. OK?

LGTM for mainline and, after some grace period, for GCC 13.

Thanks,

Tobias

> 2023-04-27  Julian Brown  <julian@codesourcery.com>
>
>       PR fortran/109622
>
> gcc/fortran/
>       * trans-openmp.cc (gfc_trans_omp_clauses): Attach/detach clause fixes.
>
> gcc/testsuite/
>       * gfortran.dg/goacc/attach-descriptor.f90: Adjust expected output.
>
> libgomp/
>       * testsuite/libgomp.fortran/pr109622.f90: New test.
>       * testsuite/libgomp.fortran/pr109622-2.f90: New test.
>       * testsuite/libgomp.fortran/pr109622-3.f90: New test.
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

Re: [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622]

[Thomas Schwinge]

Hi Julian!

On 2023-04-27T11:36:47-0700, Julian Brown <julian@codesourcery.com> wrote:
> This patch fixes several cases where multiple attach or detach mapping
> nodes were being created for stand-alone attach or detach clauses
> in Fortran.  After the introduction of stricter checking later during
> compilation, these extra nodes could cause ICEs, as seen in the PR.
>
> The patch also fixes cases that "happened to work" previously where
> the user attaches/detaches a pointer to array using a descriptor, and
> (I think!) the "_data" field has offset zero, hence the same address as
> the descriptor as a whole.

Thanks for looking into this.

I haven't reviewed the patch itself, but noticed one thing:

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.fortran/pr109622-2.f90

> +!$acc enter data copyin(var)

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.fortran/pr109622-3.f90

> +!$acc enter data copyin(var, tgt)

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.fortran/pr109622.f90

> +!$acc enter data copyin(var, var2)

You'll want to move these into 'libgomp/testsuite/libgomp.oacc-fortran/'
to actually test them with '-fopenacc' instead of '-fopenmp'.  ;-)


Chalk up one for the idea that I once had, to have '-fopenacc',
'-fopenmp', '-fopenmp-simd' enable '-Wunknown-pragmas' by default.


Grüße
 Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622]

[Julian Brown]

This patch moves several tests introduced by the following patch:

  https://gcc.gnu.org/pipermail/gcc-patches/2023-April/616939.html

into the proper location for OpenACC testing (thanks to Thomas for
spotting my mistake!), and also fixes a few additional problems --
missing diagnostics for non-pointer attaches, and a case where a pointer
was incorrectly dereferenced. Tests are also adjusted for vector-length
warnings on nvidia accelerators.

Tested with offloading to nvptx. OK?

2023-04-29  Julian Brown  <julian@codesourcery.com>

	PR fortran/109622

gcc/fortran/
	* trans-openmp.cc (gfc_trans_omp_clauses): Add diagnostic for
	non-pointer/non-allocatable attach/detach.  Remove dereference for
	pointer-to-scalar derived type component attach/detach.

gcc/testsuite/
	* gfortran.dg/goacc/pr109622-5.f90: New test.

libgomp/
	* testsuite/libgomp.fortran/pr109622.f90: Move test...
	* testsuite/libgomp.oacc-fortran/pr109622.f90: ...to here. Ignore
	vector length warning.
	* testsuite/libgomp.fortran/pr109622-2.f90: Move test...
	* testsuite/libgomp.oacc-fortran/pr109622-2.f90: ...to here.  Add
	missing copyin/copyout variable. Ignore vector length warnings.
	* testsuite/libgomp.fortran/pr109622-3.f90: Move test...
	* testsuite/libgomp.oacc-fortran/pr109622-3.f90: ...to here.  Ignore
	vector length warnings.
	* testsuite/libgomp.oacc-fortran/pr109622-4.f90: New test.
---
 gcc/fortran/trans-openmp.cc                   | 38 ++++++++++++---
 .../gfortran.dg/goacc/pr109622-5.f90          | 45 ++++++++++++++++++
 .../pr109622-2.f90                            |  7 ++-
 .../pr109622-3.f90                            |  3 ++
 .../libgomp.oacc-fortran/pr109622-4.f90       | 47 +++++++++++++++++++
 .../pr109622.f90                              |  3 ++
 6 files changed, 135 insertions(+), 8 deletions(-)
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90
 rename libgomp/testsuite/{libgomp.fortran => libgomp.oacc-fortran}/pr109622-2.f90 (63%)
 rename libgomp/testsuite/{libgomp.fortran => libgomp.oacc-fortran}/pr109622-3.f90 (76%)
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90
 rename libgomp/testsuite/{libgomp.fortran => libgomp.oacc-fortran}/pr109622.f90 (78%)

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 6ee22faa836a..b9a4ae3e53a8 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3395,6 +3395,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  && (n->u.map_op == OMP_MAP_ATTACH
 			      || n->u.map_op == OMP_MAP_DETACH))
 			{
+			  OMP_CLAUSE_DECL (node)
+			    = build_fold_addr_expr (OMP_CLAUSE_DECL (node));
 			  OMP_CLAUSE_SIZE (node) = size_zero_node;
 			  goto finalize_map_clause;
 			}
@@ -3430,6 +3432,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			    = TYPE_SIZE_UNIT (gfc_charlen_type_node);
 			}
 		    }
+		  else if (openacc
+			   && (n->u.map_op == OMP_MAP_ATTACH
+			       || n->u.map_op == OMP_MAP_DETACH))
+		    gfc_error ("%qs clause argument not pointer or "
+			       "allocatable at %L",
+			       (n->u.map_op == OMP_MAP_ATTACH)
+			       ? "attach" : "detach", &where);
 		}
 	      else if (n->expr
 		       && n->expr->expr_type == EXPR_VARIABLE
@@ -3510,6 +3519,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			}
 		      else
 			{
+			  if (openacc
+				&& (n->u.map_op == OMP_MAP_ATTACH
+				    || n->u.map_op == OMP_MAP_DETACH))
+			    gfc_error ("%qs clause argument not pointer or "
+				       "allocatable at %L",
+				       (n->u.map_op == OMP_MAP_ATTACH)
+				       ? "attach" : "detach", &where);
 			  OMP_CLAUSE_DECL (node) = inner;
 			  OMP_CLAUSE_SIZE (node)
 			    = TYPE_SIZE_UNIT (TREE_TYPE (inner));
@@ -3523,15 +3539,25 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      if (n->u.map_op == OMP_MAP_ATTACH
 			  || n->u.map_op == OMP_MAP_DETACH)
 			{
-			  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
+			  if (POINTER_TYPE_P (TREE_TYPE (inner))
+			      || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
 			    {
-			      tree ptr = gfc_conv_descriptor_data_get (inner);
-			      OMP_CLAUSE_DECL (node) = ptr;
+			      if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
+				{
+				  tree ptr
+				    = gfc_conv_descriptor_data_get (inner);
+				  OMP_CLAUSE_DECL (node) = ptr;
+				}
+			      else
+				OMP_CLAUSE_DECL (node) = inner;
+			      OMP_CLAUSE_SIZE (node) = size_zero_node;
+			      goto finalize_map_clause;
 			    }
 			  else
-			    OMP_CLAUSE_DECL (node) = inner;
-			  OMP_CLAUSE_SIZE (node) = size_zero_node;
-			  goto finalize_map_clause;
+			    gfc_error ("%qs clause argument not pointer or "
+				       "allocatable at %L",
+				       (n->u.map_op == OMP_MAP_ATTACH)
+				       ? "attach" : "detach", &where);
 			}
 
 		      if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90 b/gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90
new file mode 100644
index 000000000000..e2748964a1c2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90
@@ -0,0 +1,45 @@
+! { dg-do compile }
+
+use openacc
+implicit none
+
+type t
+integer :: foo
+character(len=8) :: bar
+integer :: qux(5)
+end type t
+
+type(t) :: var
+
+var%foo = 3
+var%bar = "HELLOOMP"
+var%qux = (/ 1, 2, 3, 4, 5 /) 
+
+!$acc enter data copyin(var)
+
+!$acc enter data attach(var%foo)
+! { dg-error "'attach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
+!$acc enter data attach(var%bar)
+! { dg-error "'attach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
+!$acc enter data attach(var%qux)
+! { dg-error "'attach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
+
+!$acc serial
+var%foo = 5
+var%bar = "GOODBYE!"
+var%qux = (/ 6, 7, 8, 9, 10 /)
+!$acc end serial
+
+!$acc exit data detach(var%qux)
+! { dg-error "'detach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
+!$acc exit data detach(var%bar)
+! { dg-error "'detach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
+!$acc exit data detach(var%foo)
+! { dg-error "'detach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
+
+!$acc exit data copyout(var)
+
+if (var%foo.ne.5) stop 1
+if (var%bar.ne."GOODBYE!") stop 2
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90
similarity index 63%
rename from libgomp/testsuite/libgomp.fortran/pr109622-2.f90
rename to libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90
index 8c5f373f39f7..d3cbebea6892 100644
--- a/libgomp/testsuite/libgomp.fortran/pr109622-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90
@@ -1,5 +1,7 @@
 ! { dg-do run }
 
+implicit none
+
 type t
 integer :: foo
 integer, pointer :: bar
@@ -13,18 +15,19 @@ var%bar => tgt
 var%foo = 99
 tgt = 199
 
-!$acc enter data copyin(var)
+!$acc enter data copyin(var, tgt)
 
 !$acc enter data attach(var%bar)
 
 !$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
 var%foo = 5
 var%bar = 7
 !$acc end serial
 
 !$acc exit data detach(var%bar)
 
-!$acc exit data copyout(var)
+!$acc exit data copyout(var, tgt)
 
 if (var%foo.ne.5) stop 1
 if (tgt.ne.7) stop 2
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90
similarity index 76%
rename from libgomp/testsuite/libgomp.fortran/pr109622-3.f90
rename to libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90
index 3ee1b43a7464..a25b1a814143 100644
--- a/libgomp/testsuite/libgomp.fortran/pr109622-3.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90
@@ -1,5 +1,7 @@
 ! { dg-do run }
 
+implicit none
+
 type t
 integer :: foo
 integer, pointer :: bar(:)
@@ -18,6 +20,7 @@ tgt = 199
 !$acc enter data attach(var%bar)
 
 !$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
 var%foo = 5
 var%bar = 7
 !$acc end serial
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90
new file mode 100644
index 000000000000..3198a0bbf79f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90
@@ -0,0 +1,47 @@
+! { dg-do run }
+
+use openacc
+implicit none
+
+type t
+integer :: foo
+character(len=8), pointer :: bar
+character(len=4), allocatable :: qux
+end type t
+
+type(t) :: var
+character(len=8), target :: tgt
+
+allocate(var%qux)
+
+var%bar => tgt
+
+var%foo = 99
+tgt = "Octopus!"
+var%qux = "Fish"
+
+!$acc enter data copyin(var, tgt)
+
+! Avoid automatic attach (i.e. with "enter data")
+call acc_copyin (var%qux)
+
+!$acc enter data attach(var%bar, var%qux)
+
+!$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
+var%foo = 5
+var%bar = "Plankton"
+var%qux = "Pond"
+!$acc end serial
+
+!$acc exit data detach(var%bar, var%qux)
+
+call acc_copyout (var%qux)
+
+!$acc exit data copyout(var, tgt)
+
+if (var%foo.ne.5) stop 1
+if (tgt.ne."Plankton") stop 2
+if (var%qux.ne."Pond") stop 3
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90
similarity index 78%
rename from libgomp/testsuite/libgomp.fortran/pr109622.f90
rename to libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90
index 5b8c4102f768..a17c4f627147 100644
--- a/libgomp/testsuite/libgomp.fortran/pr109622.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90
@@ -1,5 +1,7 @@
 ! { dg-do run }
 
+implicit none
+
 type t
 integer :: value
 type(t), pointer :: chain
@@ -18,6 +20,7 @@ nullify(var2%chain)
 !$acc enter data attach(var%chain)
 
 !$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
 var%value = 5
 var%chain%value = 7
 !$acc end serial
-- 
2.29.2

Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622]

[Tobias Burnus]

On 29.04.23 12:57, Julian Brown wrote:
> This patch moves several tests introduced by the following patch:
>
>    https://gcc.gnu.org/pipermail/gcc-patches/2023-April/616939.html

I believe you intent this as git log entry. Can you add
   ... commit r14-325-gcacf65d74463600815773255e8b82b4043432bd7
as this makes looking at the git history easier.

> into the proper location for OpenACC testing (thanks to Thomas for
> spotting my mistake!), and also fixes a few additional problems --
> missing diagnostics for non-pointer attaches, and a case where a pointer
> was incorrectly dereferenced. Tests are also adjusted for vector-length
> warnings on nvidia accelerators.
>
> Tested with offloading to nvptx. OK?
>
> 2023-04-29  Julian Brown  <julian@codesourcery.com>
>
>       PR fortran/109622
>
> gcc/fortran/
>       * trans-openmp.cc (gfc_trans_omp_clauses): Add diagnostic for
>       non-pointer/non-allocatable attach/detach.  Remove dereference for
>       pointer-to-scalar derived type component attach/detach.

In general, we prefer resolution-time diagnostic to tree-translation diagnostic,
unless there is a good reason to do the latter.
At a glance, it should be even sufficient to have a single diagnostic
instead of two when placed into openmp.cc.

Search for lastref in resolve_omp_clauses; I think it should do,
but otherwise something like:
  symbol_attr attr = gfc_expr_attr(e);
  if (attr.pointer || attr.allocatable)
should work.

You currently have:

> @@ -3430,6 +3432,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>                           = TYPE_SIZE_UNIT (gfc_charlen_type_node);
>                       }
>                   }
> +               else if (openacc
> +                        && (n->u.map_op == OMP_MAP_ATTACH
> +                            || n->u.map_op == OMP_MAP_DETACH))
> +                 gfc_error ("%qs clause argument not pointer or "
> +                            "allocatable at %L",
> +                            (n->u.map_op == OMP_MAP_ATTACH)
> +                            ? "attach" : "detach", &where);

Additionally, I think we we usually have wording like: 'must be ALLOCATABLE or a POINTER'.
(Which avoids also the question whether 'neither' instead of 'not should be used
and/besides to 'nor' instead of 'or'.)

Additionally, I think there should be a also an error for:

integer :: a
!$acc enter data attach(a)
end

(Well, in that case, just looking at n->expr won't work, but also n->sym needs to be
handled for list == OMP_LIST_MAP && n->u.map_op == OMP_MAP_(DE)ATTACH.

The other changes look fine. Thanks!

Tobias

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622]

[Thomas Schwinge]

Hi Julian!

On 2023-04-29T03:57:41-0700, Julian Brown <julian@codesourcery.com> wrote:
> This patch moves several tests introduced by the following patch:
>
>   https://gcc.gnu.org/pipermail/gcc-patches/2023-April/616939.html
>
> into the proper location for OpenACC testing (thanks to Thomas for
> spotting my mistake!), and also fixes a few additional problems --
> missing diagnostics for non-pointer attaches, and a case where a pointer
> was incorrectly dereferenced. Tests are also adjusted for vector-length
> warnings on nvidia accelerators.
>
> Tested with offloading to nvptx. OK?

Thanks for looking into this.

I haven't reviewed the patch itself, but noticed one thing:

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90
> @@ -0,0 +1,45 @@
> +! { dg-do compile }
> +
> +use openacc

    [...]/gfortran.dg/goacc/pr109622-5.f90:3:5: Fatal Error: Cannot open module file 'openacc.mod' for reading at (1): No such file or directory

... for GCC build-tree testing.  Just remove the 'use openacc'; it's not
necessary here.


Grüße
 Thomas


> +implicit none
> +
> +type t
> +integer :: foo
> +character(len=8) :: bar
> +integer :: qux(5)
> +end type t
> +
> +type(t) :: var
> +
> +var%foo = 3
> +var%bar = "HELLOOMP"
> +var%qux = (/ 1, 2, 3, 4, 5 /)
> +
> +!$acc enter data copyin(var)
> +
> +!$acc enter data attach(var%foo)
> +! { dg-error "'attach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
> +!$acc enter data attach(var%bar)
> +! { dg-error "'attach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
> +!$acc enter data attach(var%qux)
> +! { dg-error "'attach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
> +
> +!$acc serial
> +var%foo = 5
> +var%bar = "GOODBYE!"
> +var%qux = (/ 6, 7, 8, 9, 10 /)
> +!$acc end serial
> +
> +!$acc exit data detach(var%qux)
> +! { dg-error "'detach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
> +!$acc exit data detach(var%bar)
> +! { dg-error "'detach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
> +!$acc exit data detach(var%foo)
> +! { dg-error "'detach' clause argument not pointer or allocatable" "" { target *-*-* } .-1 }
> +
> +!$acc exit data copyout(var)
> +
> +if (var%foo.ne.5) stop 1
> +if (var%bar.ne."GOODBYE!") stop 2
> +
> +end
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622]

[Julian Brown]

[-- Attachment #1: Type: text/plain, Size: 1502 bytes --]

On Tue, 2 May 2023 12:29:22 +0200
Tobias Burnus <tobias@codesourcery.com> wrote:

> On 29.04.23 12:57, Julian Brown wrote:
> > This patch moves several tests introduced by the following patch:
> >
> >    https://gcc.gnu.org/pipermail/gcc-patches/2023-April/616939.html
> >  
> 
> I believe you intent this as git log entry. Can you add
>    ... commit r14-325-gcacf65d74463600815773255e8b82b4043432bd7
> as this makes looking at the git history easier.

Added.

> > into the proper location for OpenACC testing (thanks to Thomas for
> > spotting my mistake!), and also fixes a few additional problems --
> > missing diagnostics for non-pointer attaches, and a case where a
> > pointer was incorrectly dereferenced. Tests are also adjusted for
> > vector-length warnings on nvidia accelerators.
> >
> > Tested with offloading to nvptx. OK?
> >
> > 2023-04-29  Julian Brown  <julian@codesourcery.com>
> >
> >       PR fortran/109622
> >
> > gcc/fortran/
> >       * trans-openmp.cc (gfc_trans_omp_clauses): Add diagnostic for
> >       non-pointer/non-allocatable attach/detach.  Remove
> > dereference for pointer-to-scalar derived type component
> > attach/detach.  
> 
> In general, we prefer resolution-time diagnostic to tree-translation
> diagnostic, unless there is a good reason to do the latter.
> At a glance, it should be even sufficient to have a single diagnostic
> instead of two when placed into openmp.cc.

How does this version look?

Retested with offloading to nvptx.

Thanks,

Julian

[-- Attachment #2: fortran-attach-detach-fixes-2.diff --]
[-- Type: text/x-patch, Size: 11901 bytes --]

commit 43be8cd7a3e86af421e611c72f714b6c40f35bba
Author: Julian Brown <julian@codesourcery.com>
Date:   Fri Apr 28 22:27:54 2023 +0000

    OpenACC: Further attach/detach clause fixes for Fortran [PR109622]
    
    This patch moves several tests introduced by the following patch:
    
      https://gcc.gnu.org/pipermail/gcc-patches/2023-April/616939.html
      commit r14-325-gcacf65d74463600815773255e8b82b4043432bd7
    
    into the proper location for OpenACC testing (thanks to Thomas for
    spotting my mistake!), and also fixes a few additional problems --
    missing diagnostics for non-pointer attaches, and a case where a pointer
    was incorrectly dereferenced. Tests are also adjusted for vector-length
    warnings on nvidia accelerators.
    
    2023-04-29  Julian Brown  <julian@codesourcery.com>
    
    	PR fortran/109622
    
    gcc/fortran/
    	* openmp.cc (resolve_omp_clauses): Add diagnostic for
    	non-pointer/non-allocatable attach/detach.
    	* trans-openmp.cc (gfc_trans_omp_clauses): Remove dereference for
    	pointer-to-scalar derived type component attach/detach.  Fix
    	attach/detach handling for descriptors.
    
    gcc/testsuite/
    	* gfortran.dg/goacc/pr109622-5.f90: New test.
    	* gfortran.dg/goacc/pr109622-6.f90: New test.
    
    libgomp/
    	* testsuite/libgomp.fortran/pr109622.f90: Move test...
    	* testsuite/libgomp.oacc-fortran/pr109622.f90: ...to here. Ignore
    	vector length warning.
    	* testsuite/libgomp.fortran/pr109622-2.f90: Move test...
    	* testsuite/libgomp.oacc-fortran/pr109622-2.f90: ...to here.  Add
    	missing copyin/copyout variable. Ignore vector length warnings.
    	* testsuite/libgomp.fortran/pr109622-3.f90: Move test...
    	* testsuite/libgomp.oacc-fortran/pr109622-3.f90: ...to here.  Ignore
    	vector length warnings.
    	* testsuite/libgomp.oacc-fortran/pr109622-4.f90: New test.

diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 86e4515..322856a 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -7711,6 +7711,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 				     &n->where);
 		      }
 		  }
+		if (openacc
+		    && list == OMP_LIST_MAP
+		    && (n->u.map_op == OMP_MAP_ATTACH
+			|| n->u.map_op == OMP_MAP_DETACH))
+		  {
+		    symbol_attribute attr;
+		    gfc_clear_attr (&attr);
+		    if (n->expr)
+		      attr = gfc_expr_attr (n->expr);
+		    else if (n->sym)
+		      attr = n->sym->attr;
+		    if (!attr.pointer && !attr.allocatable)
+		      gfc_error ("%qs clause argument must be ALLOCATABLE or "
+				 "a POINTER at %L",
+				 (n->u.map_op == OMP_MAP_ATTACH) ? "attach"
+				 : "detach", &n->where);
+		  }
 		if (lastref
 		    || (n->expr
 			&& (!resolved || n->expr->expr_type != EXPR_VARIABLE)))
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 6ee22fa..e5de85b 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3395,6 +3395,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			  && (n->u.map_op == OMP_MAP_ATTACH
 			      || n->u.map_op == OMP_MAP_DETACH))
 			{
+			  OMP_CLAUSE_DECL (node)
+			    = build_fold_addr_expr (OMP_CLAUSE_DECL (node));
 			  OMP_CLAUSE_SIZE (node) = size_zero_node;
 			  goto finalize_map_clause;
 			}
@@ -3523,15 +3525,20 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      if (n->u.map_op == OMP_MAP_ATTACH
 			  || n->u.map_op == OMP_MAP_DETACH)
 			{
-			  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
+			  if (POINTER_TYPE_P (TREE_TYPE (inner))
+			      || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
 			    {
-			      tree ptr = gfc_conv_descriptor_data_get (inner);
-			      OMP_CLAUSE_DECL (node) = ptr;
+			      if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
+				{
+				  tree ptr
+				    = gfc_conv_descriptor_data_get (inner);
+				  OMP_CLAUSE_DECL (node) = ptr;
+				}
+			      else
+				OMP_CLAUSE_DECL (node) = inner;
+			      OMP_CLAUSE_SIZE (node) = size_zero_node;
+			      goto finalize_map_clause;
 			    }
-			  else
-			    OMP_CLAUSE_DECL (node) = inner;
-			  OMP_CLAUSE_SIZE (node) = size_zero_node;
-			  goto finalize_map_clause;
 			}
 
 		      if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90 b/gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90
new file mode 100644
index 0000000..5c483eb
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr109622-5.f90
@@ -0,0 +1,45 @@
+! { dg-do compile }
+
+use openacc
+implicit none
+
+type t
+integer :: foo
+character(len=8) :: bar
+integer :: qux(5)
+end type t
+
+type(t) :: var
+
+var%foo = 3
+var%bar = "HELLOOMP"
+var%qux = (/ 1, 2, 3, 4, 5 /) 
+
+!$acc enter data copyin(var)
+
+!$acc enter data attach(var%foo)
+! { dg-error "'attach' clause argument must be ALLOCATABLE or a POINTER" "" { target *-*-* } .-1 }
+!$acc enter data attach(var%bar)
+! { dg-error "'attach' clause argument must be ALLOCATABLE or a POINTER" "" { target *-*-* } .-1 }
+!$acc enter data attach(var%qux)
+! { dg-error "'attach' clause argument must be ALLOCATABLE or a POINTER" "" { target *-*-* } .-1 }
+
+!$acc serial
+var%foo = 5
+var%bar = "GOODBYE!"
+var%qux = (/ 6, 7, 8, 9, 10 /)
+!$acc end serial
+
+!$acc exit data detach(var%qux)
+! { dg-error "'detach' clause argument must be ALLOCATABLE or a POINTER" "" { target *-*-* } .-1 }
+!$acc exit data detach(var%bar)
+! { dg-error "'detach' clause argument must be ALLOCATABLE or a POINTER" "" { target *-*-* } .-1 }
+!$acc exit data detach(var%foo)
+! { dg-error "'detach' clause argument must be ALLOCATABLE or a POINTER" "" { target *-*-* } .-1 }
+
+!$acc exit data copyout(var)
+
+if (var%foo.ne.5) stop 1
+if (var%bar.ne."GOODBYE!") stop 2
+
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr109622-6.f90 b/gcc/testsuite/gfortran.dg/goacc/pr109622-6.f90
new file mode 100644
index 0000000..256ab90
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr109622-6.f90
@@ -0,0 +1,8 @@
+! { dg-do compile }
+
+implicit none
+integer :: x
+!$acc enter data attach(x)
+! { dg-error "'attach' clause argument must be ALLOCATABLE or a POINTER" "" { target *-*-* } .-1 }
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-2.f90 b/libgomp/testsuite/libgomp.fortran/pr109622-2.f90
deleted file mode 100644
index 8c5f373..0000000
--- a/libgomp/testsuite/libgomp.fortran/pr109622-2.f90
+++ /dev/null
@@ -1,32 +0,0 @@
-! { dg-do run }
-
-type t
-integer :: foo
-integer, pointer :: bar
-end type t
-
-type(t) :: var
-integer, target :: tgt
-
-var%bar => tgt
-
-var%foo = 99
-tgt = 199
-
-!$acc enter data copyin(var)
-
-!$acc enter data attach(var%bar)
-
-!$acc serial
-var%foo = 5
-var%bar = 7
-!$acc end serial
-
-!$acc exit data detach(var%bar)
-
-!$acc exit data copyout(var)
-
-if (var%foo.ne.5) stop 1
-if (tgt.ne.7) stop 2
-
-end
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-3.f90 b/libgomp/testsuite/libgomp.fortran/pr109622-3.f90
deleted file mode 100644
index 3ee1b43..0000000
--- a/libgomp/testsuite/libgomp.fortran/pr109622-3.f90
+++ /dev/null
@@ -1,32 +0,0 @@
-! { dg-do run }
-
-type t
-integer :: foo
-integer, pointer :: bar(:)
-end type t
-
-type(t) :: var
-integer, target :: tgt(20)
-
-var%bar => tgt
-
-var%foo = 99
-tgt = 199
-
-!$acc enter data copyin(var, tgt)
-
-!$acc enter data attach(var%bar)
-
-!$acc serial
-var%foo = 5
-var%bar = 7
-!$acc end serial
-
-!$acc exit data detach(var%bar)
-
-!$acc exit data copyout(var, tgt)
-
-if (var%foo.ne.5) stop 1
-if (any(tgt.ne.7)) stop 2
-
-end
diff --git a/libgomp/testsuite/libgomp.fortran/pr109622.f90 b/libgomp/testsuite/libgomp.fortran/pr109622.f90
deleted file mode 100644
index 5b8c410..0000000
--- a/libgomp/testsuite/libgomp.fortran/pr109622.f90
+++ /dev/null
@@ -1,32 +0,0 @@
-! { dg-do run }
-
-type t
-integer :: value
-type(t), pointer :: chain
-end type t
-
-type(t), target :: var, var2
-
-var%value = 99
-var2%value = 199
-
-var%chain => var2
-nullify(var2%chain)
-
-!$acc enter data copyin(var, var2)
-
-!$acc enter data attach(var%chain)
-
-!$acc serial
-var%value = 5
-var%chain%value = 7
-!$acc end serial
-
-!$acc exit data detach(var%chain)
-
-!$acc exit data copyout(var, var2)
-
-if (var%value.ne.5) stop 1
-if (var2%value.ne.7) stop 2
-
-end
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90
new file mode 100644
index 0000000..d3cbebe
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90
@@ -0,0 +1,35 @@
+! { dg-do run }
+
+implicit none
+
+type t
+integer :: foo
+integer, pointer :: bar
+end type t
+
+type(t) :: var
+integer, target :: tgt
+
+var%bar => tgt
+
+var%foo = 99
+tgt = 199
+
+!$acc enter data copyin(var, tgt)
+
+!$acc enter data attach(var%bar)
+
+!$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
+var%foo = 5
+var%bar = 7
+!$acc end serial
+
+!$acc exit data detach(var%bar)
+
+!$acc exit data copyout(var, tgt)
+
+if (var%foo.ne.5) stop 1
+if (tgt.ne.7) stop 2
+
+end
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90
new file mode 100644
index 0000000..a25b1a8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90
@@ -0,0 +1,35 @@
+! { dg-do run }
+
+implicit none
+
+type t
+integer :: foo
+integer, pointer :: bar(:)
+end type t
+
+type(t) :: var
+integer, target :: tgt(20)
+
+var%bar => tgt
+
+var%foo = 99
+tgt = 199
+
+!$acc enter data copyin(var, tgt)
+
+!$acc enter data attach(var%bar)
+
+!$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
+var%foo = 5
+var%bar = 7
+!$acc end serial
+
+!$acc exit data detach(var%bar)
+
+!$acc exit data copyout(var, tgt)
+
+if (var%foo.ne.5) stop 1
+if (any(tgt.ne.7)) stop 2
+
+end
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90
new file mode 100644
index 0000000..3198a0b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90
@@ -0,0 +1,47 @@
+! { dg-do run }
+
+use openacc
+implicit none
+
+type t
+integer :: foo
+character(len=8), pointer :: bar
+character(len=4), allocatable :: qux
+end type t
+
+type(t) :: var
+character(len=8), target :: tgt
+
+allocate(var%qux)
+
+var%bar => tgt
+
+var%foo = 99
+tgt = "Octopus!"
+var%qux = "Fish"
+
+!$acc enter data copyin(var, tgt)
+
+! Avoid automatic attach (i.e. with "enter data")
+call acc_copyin (var%qux)
+
+!$acc enter data attach(var%bar, var%qux)
+
+!$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
+var%foo = 5
+var%bar = "Plankton"
+var%qux = "Pond"
+!$acc end serial
+
+!$acc exit data detach(var%bar, var%qux)
+
+call acc_copyout (var%qux)
+
+!$acc exit data copyout(var, tgt)
+
+if (var%foo.ne.5) stop 1
+if (tgt.ne."Plankton") stop 2
+if (var%qux.ne."Pond") stop 3
+
+end
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90
new file mode 100644
index 0000000..a17c4f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90
@@ -0,0 +1,35 @@
+! { dg-do run }
+
+implicit none
+
+type t
+integer :: value
+type(t), pointer :: chain
+end type t
+
+type(t), target :: var, var2
+
+var%value = 99
+var2%value = 199
+
+var%chain => var2
+nullify(var2%chain)
+
+!$acc enter data copyin(var, var2)
+
+!$acc enter data attach(var%chain)
+
+!$acc serial
+! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 }
+var%value = 5
+var%chain%value = 7
+!$acc end serial
+
+!$acc exit data detach(var%chain)
+
+!$acc exit data copyout(var, var2)
+
+if (var%value.ne.5) stop 1
+if (var2%value.ne.7) stop 2
+
+end

Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622]

[Tobias Burnus]

On 03.05.23 14:59, Julian Brown wrote:
> How does this version look?
> Retested with offloading to nvptx.
LGTM (for mainline + GCC 13 backporting) but I have two tiny comments:
> diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
> index 86e4515..322856a 100644
> --- a/gcc/fortran/openmp.cc
> +++ b/gcc/fortran/openmp.cc
> @@ -7711,6 +7711,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>                                    &n->where);
>                     }
>                 }
> +             if (openacc
> +                 && list == OMP_LIST_MAP
> +                 && (n->u.map_op == OMP_MAP_ATTACH
> +                     || n->u.map_op == OMP_MAP_DETACH))
> +               {
> +                 symbol_attribute attr;
> +                 gfc_clear_attr (&attr);
> +                 if (n->expr)
> +                   attr = gfc_expr_attr (n->expr);
> +                 else if (n->sym)
> +                   attr = n->sym->attr;

Note that n->sym == NULL is only the case if the argument was
omp_all_memory (→ gfc_match_omp_variable_list); that can only be the
case for OMP_CLAUSE_DEPEND.

As OpenMP's 'depend' clause is handled before and you additionally deal
with OpenACC, only, you could just use 'else' instead of 'else if
(n->sym)' – and also get rid of the 'gfc_clear_attr' as the values get
overwritten by the assignment and by the function call.

Later code (e.g. line 7785 in the current code) also assumes (for
OpenACC + MAP) that n->sym != NULL by bluntly dereferencing it.

@@ -3523,15 +3525,20 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
>                     if (n->u.map_op == OMP_MAP_ATTACH
>                         || n->u.map_op == OMP_MAP_DETACH)
>                       {
> -                       if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
> +                       if (POINTER_TYPE_P (TREE_TYPE (inner))
> +                           || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
>                           {
...
>                           }
> -                       else
> -                         OMP_CLAUSE_DECL (node) = inner;
> -                       OMP_CLAUSE_SIZE (node) = size_zero_node;
> -                       goto finalize_map_clause;
>                       }

You can now combine the two if conditions, which avoids some indenting
and should permit to put 'tree ptr' / ' = ...' again on the same line.

Thanks for the patch!

Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955