From 049bfd186fae9fb764a3ec04acb20d3eaacda7a3 Mon Sep 17 00:00:00 2001 From: Tobias Burnus Date: Tue, 19 Jan 2021 11:57:34 +0100 Subject: [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-2.f90: New test. * gfortran.dg/gomp/use_device_ptr-1.f90: New test. --- gcc/fortran/openmp.c | 67 ++++++++++++++----- gcc/omp-low.c | 6 +- .../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 +++++++ .../libgomp.fortran/is_device_ptr-1.f90 | 54 +++++++++++++++ 6 files changed, 160 insertions(+), 23 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/use_device_ptr-1.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/is_device_ptr-1.f90 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" } } -- 2.30.2