OpenMP/Fortran: Fixes for {use,is}_device_ptr

gcc/fortran/ChangeLog:

	PR fortran/98476
	* openmp.c (resolve_omp_clauses): Change use_device_ptr
	to use_device_addr for unless type(c_ptr); check all
	list item for is_device_ptr.

gcc/ChangeLog:

	PR fortran/98476
	* omp-low.c (lower_omp_target): Handle nonpointer is_device_ptr.

libgomp/ChangeLog:

	PR fortran/98476
	* testsuite/libgomp.fortran/is_device_ptr-1.f90: New test.

gcc/testsuite/ChangeLog:

	PR fortran/98476
	* gfortran.dg/gomp/map-3.f90: Update expected scan-dump-tree.
	* gfortran.dg/gomp/is_device_ptr-2.f90: New test.
	* gfortran.dg/gomp/use_device_ptr-1.f90: New test.

6 files changed, 160 insertions, 23 deletions

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index a9ecd96..9a3a8f6 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -5345,22 +5345,25 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		}
 	    break;
 	  case OMP_LIST_IS_DEVICE_PTR:
-	    if (!n->sym->attr.dummy)
-	      gfc_error ("Non-dummy object %qs in %s clause at %L",
-			 n->sym->name, name, &n->where);
-	    if (n->sym->attr.allocatable
-		|| (n->sym->ts.type == BT_CLASS
-		    && CLASS_DATA (n->sym)->attr.allocatable))
-	      gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
-			 n->sym->name, name, &n->where);
-	    if (n->sym->attr.pointer
-		|| (n->sym->ts.type == BT_CLASS
-		    && CLASS_DATA (n->sym)->attr.pointer))
-	      gfc_error ("POINTER object %qs in %s clause at %L",
-			 n->sym->name, name, &n->where);
-	    if (n->sym->attr.value)
-	      gfc_error ("VALUE object %qs in %s clause at %L",
-			 n->sym->name, name, &n->where);
+	    for (n = omp_clauses->lists[list]; n != NULL; n = n->next)
+	      {
+		if (!n->sym->attr.dummy)
+		  gfc_error ("Non-dummy object %qs in %s clause at %L",
+			     n->sym->name, name, &n->where);
+		if (n->sym->attr.allocatable
+		    || (n->sym->ts.type == BT_CLASS
+			&& CLASS_DATA (n->sym)->attr.allocatable))
+		  gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
+			     n->sym->name, name, &n->where);
+		if (n->sym->attr.pointer
+		    || (n->sym->ts.type == BT_CLASS
+			&& CLASS_DATA (n->sym)->attr.pointer))
+		  gfc_error ("POINTER object %qs in %s clause at %L",
+			     n->sym->name, name, &n->where);
+		if (n->sym->attr.value)
+		  gfc_error ("VALUE object %qs in %s clause at %L",
+			     n->sym->name, name, &n->where);
+	      }
 	    break;
 	  case OMP_LIST_USE_DEVICE_PTR:
 	  case OMP_LIST_USE_DEVICE_ADDR:
@@ -5657,6 +5660,38 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	    break;
 	  }
       }
+  /* OpenMP 5.1: use_device_ptr acts like use_device_addr, except for
+     type(c_ptr).  */
+  if (omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR])
+    {
+      gfc_omp_namelist *n_prev, *n_next, *n_addr;
+      n_addr = omp_clauses->lists[OMP_LIST_USE_DEVICE_ADDR];
+      for (; n_addr && n_addr->next; n_addr = n_addr->next)
+	;
+      n_prev = NULL;
+      n = omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR];
+      while (n)
+	{
+	  n_next = n->next;
+	  if (n->sym->ts.type != BT_DERIVED
+	      || n->sym->ts.u.derived->ts.f90_type != BT_VOID)
+	    {
+	      n->next = NULL;
+	      if (n_addr)
+		n_addr->next = n;
+	      else
+		omp_clauses->lists[OMP_LIST_USE_DEVICE_ADDR] = n;
+	      n_addr = n;
+	      if (n_prev)
+		n_prev->next = n_next;
+	      else
+		omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR] = n_next;
+	    }
+	  else
+	    n_prev = n;
+	  n = n_next;
+	}
+    }
   if (omp_clauses->safelen_expr)
     resolve_positive_int_expr (omp_clauses->safelen_expr, "SAFELEN");
   if (omp_clauses->simdlen_expr)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index c1267dc..df5b6ce 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12520,7 +12520,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    || omp_is_allocatable_or_ptr (ovar))
 		  {
 		    type = TREE_TYPE (type);
-		    if (TREE_CODE (type) != ARRAY_TYPE
+		    if (POINTER_TYPE_P (type)
+			&& TREE_CODE (type) != ARRAY_TYPE
 			&& ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
 			    && !omp_is_allocatable_or_ptr (ovar))
 			   || (omp_is_reference (ovar)
@@ -12784,7 +12785,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		if (omp_is_reference (var))
 		  {
 		    type = TREE_TYPE (type);
-		    if (TREE_CODE (type) != ARRAY_TYPE
+		    if (POINTER_TYPE_P (type)
+			&& TREE_CODE (type) != ARRAY_TYPE
 			&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
 			    || (omp_is_reference (var)
 				&& omp_is_allocatable_or_ptr (var))))
diff --git a/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90 b/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90
new file mode 100644
index 0000000..bf49820
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90
@@ -0,0 +1,21 @@
+! PR fortran/98476
+
+subroutine abc(cc)
+    integer, target :: cc, dd
+    cc = 131
+    dd = 484
+
+    !$omp target enter data map(to: cc, dd)
+
+    !$omp target data use_device_addr(cc) use_device_ptr(dd)
+      !$omp target is_device_ptr(cc, dd)  ! { dg-error "Non-dummy object 'cc' in IS_DEVICE_PTR clause at" }
+        if (cc /= 131 .or. dd /= 484) stop 1
+        cc = 44
+        dd = 45
+      !$omp end target
+    !$omp end target data
+
+    !$omp target exit data map(from:cc, dd)
+
+    if (cc /= 44 .or. dd /= 45) stop 5
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-3.f90 b/gcc/testsuite/gfortran.dg/gomp/map-3.f90
index 13f6364..bdd2890 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-3.f90
@@ -1,10 +1,10 @@
 ! { dg-additional-options "-fdump-tree-original" }
 
 subroutine bar
-integer, target :: x
+integer, target :: x, x2
 integer, allocatable, target :: y(:,:), z(:,:)
 x = 7
-!$omp target enter data map(to:x)
+!$omp target enter data map(to:x, x2)
 
 x = 8
 !$omp target data map(always, to: x)
@@ -15,7 +15,7 @@ call foo(x)
 call foo2(x)
 !$omp end target data
 
-!$omp target data use_device_addr(x)
+!$omp target data use_device_addr(x2)
 call foo2(x)
 !$omp end target data
 !$omp target exit data map(release:x)
@@ -31,8 +31,8 @@ end
 
 ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:x\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,to:x\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_ptr\\(x\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x2\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(release:x\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) y.data \\\[len: .*\\) map\\(to:y \\\[pointer set, len: .*\\) map\\(alloc:.*y.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(y\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) z.data \\\[len: .*\\) map\\(to:z \\\[pointer set, len: .*\\) map\\(alloc:.*z.data \\\[pointer assign, bias: 0\\\]\\) use_device_ptr\\(z\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) z.data \\\[len: .*\\) map\\(to:z \\\[pointer set, len: .*\\) map\\(alloc:.*z.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(z\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/use_device_ptr-1.f90 b/gcc/testsuite/gfortran.dg/gomp/use_device_ptr-1.f90
new file mode 100644
index 0000000..6f47fdd
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/use_device_ptr-1.f90
@@ -0,0 +1,25 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+! PR fortran/98476
+
+use iso_c_binding, only: c_ptr
+implicit none (external, type)
+
+interface
+  subroutine bar(x)
+    import
+    type(c_ptr), value :: x
+  end
+end interface
+
+type(c_ptr) :: x
+
+!$omp target data map(alloc: x)
+!$omp target data use_device_ptr(x)
+  call bar(x)
+!$omp end target data
+!$omp end target data
+end
+
+! { dg-final { scan-tree-dump-times "pragma omp target data use_device_ptr\\(x\\)" 1 "original" } }
diff --git a/libgomp/testsuite/libgomp.fortran/is_device_ptr-1.f90 b/libgomp/testsuite/libgomp.fortran/is_device_ptr-1.f90
new file mode 100644
index 0000000..30a927a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/is_device_ptr-1.f90
@@ -0,0 +1,54 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+! PR fortran/98476
+
+program abc
+  implicit none
+  integer a, b
+
+  a = 83
+  b = 73
+  call test(a, b)
+
+contains
+  subroutine test(aa, bb)
+    use iso_c_binding, only: c_ptr, c_loc, c_f_pointer
+    integer :: aa, bb
+    integer, target :: cc, dd
+    type(c_ptr) :: pcc, pdd
+    cc = 131
+    dd = 484
+
+    !$omp target enter data map(to: aa, bb, cc, dd)
+
+    !$omp target data use_device_ptr(aa, cc) use_device_addr(bb, dd)
+      pcc = c_loc(cc)
+      pdd = c_loc(dd)
+
+      ! TODO: has_device_addr(cc, dd)
+      !$omp target is_device_ptr(aa, bb)
+        if (aa /= 83 .or. bb /= 73) stop 1
+        aa = 42
+        bb = 43
+        block
+          integer, pointer :: c2, d2
+          call c_f_pointer(pcc, c2)
+          call c_f_pointer(pdd, d2)
+          if (c2 /= 131 .or. d2 /= 484) stop 2
+          c2 = 44
+          d2 = 45
+        end block
+      !$omp end target
+    !$omp end target data
+
+    !$omp target exit data map(from:aa, bb, cc, dd)
+
+    if (aa /= 42 .or. bb /= 43) stop 3
+    if (cc /= 44 .or. dd /= 45) stop 5
+  endsubroutine
+end program
+
+! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(aa\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(bb\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(cc\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(dd\\)" 1 "original" } }