diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 71f71235756c72d096bc3637d1ea53a826c228dd..6ee22faa836a3cd8d5fcd3aa457d262404e5d892 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 8c2ee4a5cca449c8577ff8f4713dee46265637d2..734afbe6ca480d37840b005897eb42fbf03890b9 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 0000000000000000000000000000000000000000..8c5f373f39f7b5429d29e1eb4e74302d2a5c5d08
--- /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 0000000000000000000000000000000000000000..3ee1b43a74642fbd4cc352c7a02434fd62834672
--- /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 0000000000000000000000000000000000000000..5b8c4102f7687ddbbcf6a6716ad5767e5cdfff83
--- /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