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

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.

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.

(cherry picked from commit cacf65d74463600815773255e8b82b4043432bd7)

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 84c0184f48efd46ae51c8498d255c3bbcd34a59b..be2261302cc494c0df709a22a9958360e6d661fc 100644 (file)
--- 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 (file)
--- 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 (file)
index 0000000..8c5f373
--- /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 (file)
index 0000000..3ee1b43
--- /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 (file)
index 0000000..5b8c410
--- /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