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.

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index a9ecd96cb357eeec24e8d706432bbab2baf2edb4..9a3a8f63b5e28d5f27ff7133768bab524310dd84 100644 (file)
--- 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 c1267dcce2eaecfadb25c73a2441a325e88c34c8..df5b6cec58623e7fe9f1a45550ba837042ae250b 100644 (file)
--- 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 (file)
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 13f63647bda8d5198756947805254b0e21dc7257..bdd2890b277f6257649bb15c4ca08dc685d59888 100644 (file)
--- 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 (file)
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 (file)
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" } }