diff mbox series

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

Message ID a1e08b8d-a3fa-a6a5-38da-57577ea9be2a@codesourcery.com
State New
Headers show
Series OpenMP/Fortran: Fixes for {use,is}_device_ptr [PR98476] | expand

Commit Message

Tobias Burnus Jan. 18, 2021, 4:56 p.m. UTC 
use_device_ptr and is_device_ptr are underspecified for Fortran in OpenMP <= 5.0
and, hence, also not properly implemented in GCC.

OpenMP 5.1 cleaned this up by mapping (for Fortran) use_device_ptr to
the existing and well-defined use_device_addr (except for type(c_ptr)) and
is_device_ptr (with the same exception) to the new has_device_addr.

The attached testcase gave a ME ICE for 'use_device_ptr(cc,dd)' –
which was be fixed by applying the new OpenMP 5.1 mapping to
'use_device_addr'. → gcc/fortran/openmp.c.

is_device_ptr(aa) is from real-world code (see PR) and could be fixed by
adding two 'pointer_type_p(type)' checks. → omp-low.c

While testing, it turned out that 'is_device_ptr(aa,bb,cc,dd) was accepted
as only the first list item was checked – giving later an ICE in the ME.
(Otherwise deferred to OpenMP 5.1's has_device_addr.)

OK?

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Comments

Jakub Jelinek Jan. 19, 2021, 10:18 a.m. UTC | #1 
On Mon, Jan 18, 2021 at 05:56:21PM +0100, Tobias Burnus wrote:
> 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-1.f90: New test.
> 	* gfortran.dg/gomp/use_device_ptr-2.f90: New test.

Ok, thanks.

	Jakub
Jakub Jelinek Jan. 19, 2021, 5:49 p.m. UTC | #2 
On Mon, Jan 18, 2021 at 05:56:21PM +0100, Tobias Burnus wrote:
> gcc/testsuite/ChangeLog:
> 
> 	PR fortran/98476
> 	* gfortran.dg/gomp/map-3.f90: Update expected scan-dump-tree.
> 	* gfortran.dg/gomp/is_device_ptr-1.f90: New test.
> 	* gfortran.dg/gomp/use_device_ptr-2.f90: New test.

I'm getting
/usr/src/gcc/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90:11:36: Error: Non-dummy object 'dd' in IS_DEVICE_PTR clause at (1)
compiler exited with status 1
FAIL: gfortran.dg/gomp/is_device_ptr-2.f90   -O   (test for errors, line 11)
FAIL: gfortran.dg/gomp/is_device_ptr-2.f90   -O  (test for excess errors)
Excess errors:
/usr/src/gcc/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90:11:36: Error: Non-dummy object 'dd' in IS_DEVICE_PTR clause at (1)
failure everywhere, the test expects cc instead of dd to be printed.
Do we want it to diagnose both, or should the regexp accept any of them?

	Jakub
Tobias Burnus Jan. 20, 2021, 7:38 a.m. UTC | #3 
On 18.01.21 17:56, Tobias Burnus wrote:

> While testing, it turned out that 'is_device_ptr(aa,bb,cc,dd) was
> accepted
> as only the first list item was checked – giving later an ICE in the ME.
... and as PR fortran/98757 showed, I forgot a 'git add' before committing
the associated testcase, given that I already had locally the following
patch ...

Committed as Rev. r11-6809-gc05cdfb3f6335d55226cef7917a783498aa41244

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
diff mbox series

Patch

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-1.f90: New test.
	* gfortran.dg/gomp/use_device_ptr-2.f90: New test.

 gcc/fortran/openmp.c                               | 67 ++++++++++++++++------
 gcc/omp-low.c                                      |  6 +-
 gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90 | 21 +++++++
 gcc/testsuite/gfortran.dg/gomp/map-3.f90           | 10 ++--
 .../gfortran.dg/gomp/use_device_ptr-1.f90          | 25 ++++++++
 .../testsuite/libgomp.fortran/is_device_ptr-1.f90  | 54 +++++++++++++++++
 6 files changed, 160 insertions(+), 23 deletions(-)

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index a9ecd96cb35..9a3a8f63b5e 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 c1267dcce2e..df5b6cec586 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 00000000000..bf498208aa8
--- /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 13f63647bda..bdd2890b277 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 00000000000..6f47fddf7cb
--- /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 00000000000..30a927a19ba
--- /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" } }