* [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622] @ 2023-04-27 18:36 ` Julian Brown 2023-04-28 8:16 ` Tobias Burnus 2023-04-28 12:56 ` Thomas Schwinge 0 siblings, 2 replies; 8+ messages in thread From: Julian Brown @ 2023-04-27 18:36 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, tobias, jakub 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622] 2023-04-27 18:36 ` [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622] Julian Brown @ 2023-04-28 8:16 ` Tobias Burnus 2023-04-28 12:56 ` Thomas Schwinge 1 sibling, 0 replies; 8+ messages in thread From: Tobias Burnus @ 2023-04-28 8:16 UTC (permalink / raw) To: Julian Brown, gcc-patches; +Cc: fortran, tobias, jakub 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622] 2023-04-27 18:36 ` [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622] Julian Brown 2023-04-28 8:16 ` Tobias Burnus @ 2023-04-28 12:56 ` Thomas Schwinge 2023-04-29 10:57 ` [PATCH] OpenACC: Further " Julian Brown 1 sibling, 1 reply; 8+ messages in thread From: Thomas Schwinge @ 2023-04-28 12:56 UTC (permalink / raw) To: Julian Brown; +Cc: fortran, tobias, jakub, gcc-patches 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
* [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622] @ 2023-04-29 10:57 ` Julian Brown 2023-05-02 10:29 ` Tobias Burnus 2023-05-03 11:29 ` Thomas Schwinge 0 siblings, 2 replies; 8+ messages in thread From: Julian Brown @ 2023-04-29 10:57 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, tobias, jakub, thomas 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622] 2023-04-29 10:57 ` [PATCH] OpenACC: Further " Julian Brown @ 2023-05-02 10:29 ` Tobias Burnus 2023-05-03 12:59 ` Julian Brown 2023-05-03 11:29 ` Thomas Schwinge 1 sibling, 1 reply; 8+ messages in thread From: Tobias Burnus @ 2023-05-02 10:29 UTC (permalink / raw) To: Julian Brown, gcc-patches; +Cc: fortran, jakub, thomas 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622] 2023-05-02 10:29 ` Tobias Burnus @ 2023-05-03 12:59 ` Julian Brown 2023-05-03 13:50 ` Tobias Burnus 0 siblings, 1 reply; 8+ messages in thread From: Julian Brown @ 2023-05-03 12:59 UTC (permalink / raw) To: Tobias Burnus; +Cc: gcc-patches, fortran, jakub, thomas [-- 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622] 2023-05-03 12:59 ` Julian Brown @ 2023-05-03 13:50 ` Tobias Burnus 0 siblings, 0 replies; 8+ messages in thread From: Tobias Burnus @ 2023-05-03 13:50 UTC (permalink / raw) To: Julian Brown; +Cc: gcc-patches, fortran, jakub, thomas 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [PATCH] OpenACC: Further attach/detach clause fixes for Fortran [PR109622] 2023-04-29 10:57 ` [PATCH] OpenACC: Further " Julian Brown 2023-05-02 10:29 ` Tobias Burnus @ 2023-05-03 11:29 ` Thomas Schwinge 1 sibling, 0 replies; 8+ messages in thread From: Thomas Schwinge @ 2023-05-03 11:29 UTC (permalink / raw) To: Julian Brown; +Cc: fortran, tobias, jakub, gcc-patches 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 ^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2023-05-03 13:51 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <138a36f7-8b1e-5955-1d3a-5713a0fcf5b6@univ-grenoble-alpes.fr>
2023-04-27 18:36 ` [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622] Julian Brown
2023-04-28 8:16 ` Tobias Burnus
2023-04-28 12:56 ` Thomas Schwinge
2023-04-29 10:57 ` [PATCH] OpenACC: Further " Julian Brown
2023-05-02 10:29 ` Tobias Burnus
2023-05-03 12:59 ` Julian Brown
2023-05-03 13:50 ` Tobias Burnus
2023-05-03 11:29 ` Thomas Schwinge
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).