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

[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