[gcc r11-6787] OpenMP/Fortran: Fixes for {use,is}_device_ptr
Tobias Burnus
burnus@gcc.gnu.org
Tue Jan 19 10:59:40 GMT 2021
https://gcc.gnu.org/g:049bfd186fae9fb764a3ec04acb20d3eaacda7a3
commit r11-6787-g049bfd186fae9fb764a3ec04acb20d3eaacda7a3
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Tue Jan 19 11:57:34 2021 +0100
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:
---
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" } }
More information about the Gcc-cvs
mailing list