aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorTobias Burnus <tobias@codesourcery.com>2019-09-20 18:23:38 +0200
committerThomas Schwinge <thomas@codesourcery.com>2020-03-03 12:51:25 +0100
commitbb911c6f61333cf8e1018f303a66c35add30011c (patch)
tree078aac60b3bb09c4600c0aee2688197174ca4dbf
parentff39b8b9248bcd24bc30d64fae8085322564d8ff (diff)
downloadgcc-bb911c6f61333cf8e1018f303a66c35add30011c.zip
gcc-bb911c6f61333cf8e1018f303a66c35add30011c.tar.gz
gcc-bb911c6f61333cf8e1018f303a66c35add30011c.tar.bz2
Backport from mainline
2019-09-20 Tobias Burnus <tobias@codesourcery.com> * openmp.c (gfc_resolve_oacc_declare): Reject all non variables but accept function result variables. * trans-openmp.c (gfc_trans_omp_clauses): Handle function-result variables for remaing cases. 2019-09-20 Tobias Burnus <tobias@codesourcery.com> * gfortran.dg/goacc/parameter.f95: Change dg-error as it is now detected earlier. * gfortran.dg/goacc/pr85701.f90: Modify to use a separate result variable. * gfortran.dg/goacc/pr78260.f90: New. * gfortran.dg/goacc/pr78260-2.f90: New. * gfortran.dg/gomp/pr78260.f90: New. * gfortran.dg/gomp/pr78260-2.f90: New. * gfortran.dg/gomp/pr78260-3.f90: New. (cherry picked from openacc-gcc-9-branch commit 7d0d394c1c9bfb1b489c4ad3bd606d7a4765ae1d)
-rw-r--r--gcc/fortran/ChangeLog.omp10
-rw-r--r--gcc/fortran/openmp.c10
-rw-r--r--gcc/fortran/trans-openmp.c6
-rw-r--r--gcc/testsuite/ChangeLog.omp15
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/parameter.f952
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/pr78260-2.f9020
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/pr78260.f9036
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/pr85701.f904
-rw-r--r--gcc/testsuite/gfortran.dg/gomp/pr78260-2.f9059
-rw-r--r--gcc/testsuite/gfortran.dg/gomp/pr78260-3.f9074
-rw-r--r--gcc/testsuite/gfortran.dg/gomp/pr78260.f9033
11 files changed, 256 insertions, 13 deletions
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index a54fb4e..87766c9 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,13 @@
+2019-09-20 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline
+ 2019-09-20 Tobias Burnus <tobias@codesourcery.com>
+
+ * openmp.c (gfc_resolve_oacc_declare): Reject all
+ non variables but accept function result variables.
+ * trans-openmp.c (gfc_trans_omp_clauses): Handle
+ function-result variables for remaing cases.
+
2019-09-17 Tobias Burnus <tobias@codesourcery.com>
* trans-expr.c (gfc_auto_dereference_var): Use passed loc argument.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index fa0cd6f..f8b4cb0 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -6178,18 +6178,14 @@ gfc_resolve_oacc_declare (gfc_namespace *ns)
for (n = oc->clauses->lists[list]; n; n = n->next)
{
n->sym->mark = 0;
- if (n->sym->attr.function || n->sym->attr.subroutine)
+ if (n->sym->attr.flavor != FL_VARIABLE
+ && (n->sym->attr.flavor != FL_PROCEDURE
+ || n->sym->result != n->sym))
{
gfc_error ("Object %qs is not a variable at %L",
n->sym->name, &oc->loc);
continue;
}
- if (n->sym->attr.flavor == FL_PARAMETER)
- {
- gfc_error ("PARAMETER object %qs is not allowed at %L",
- n->sym->name, &oc->loc);
- continue;
- }
if (n->expr && n->expr->ref->type == REF_ARRAY)
{
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d1799d7..5d0d6d2c 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2280,7 +2280,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
tree node = build_omp_clause (input_location, OMP_CLAUSE_DEPEND);
if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
{
- tree decl = gfc_get_symbol_decl (n->sym);
+ tree decl = gfc_trans_omp_variable (n->sym, false);
if (gfc_omp_privatize_by_reference (decl))
decl = build_fold_indirect_ref (decl);
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
@@ -2341,7 +2341,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
tree node2 = NULL_TREE;
tree node3 = NULL_TREE;
tree node4 = NULL_TREE;
- tree decl = gfc_get_symbol_decl (n->sym);
+ tree decl = gfc_trans_omp_variable (n->sym, false);
if (DECL_P (decl))
TREE_ADDRESSABLE (decl) = 1;
@@ -2749,7 +2749,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
tree node = build_omp_clause (input_location, clause_code);
if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
{
- tree decl = gfc_get_symbol_decl (n->sym);
+ tree decl = gfc_trans_omp_variable (n->sym, false);
if (gfc_omp_privatize_by_reference (decl))
decl = build_fold_indirect_ref (decl);
else if (DECL_P (decl))
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 5f452e6..8282d96 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,18 @@
+2019-09-20 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline
+ 2019-09-20 Tobias Burnus <tobias@codesourcery.com>
+
+ * gfortran.dg/goacc/parameter.f95: Change
+ dg-error as it is now detected earlier.
+ * gfortran.dg/goacc/pr85701.f90: Modify to
+ use a separate result variable.
+ * gfortran.dg/goacc/pr78260.f90: New.
+ * gfortran.dg/goacc/pr78260-2.f90: New.
+ * gfortran.dg/gomp/pr78260.f90: New.
+ * gfortran.dg/gomp/pr78260-2.f90: New.
+ * gfortran.dg/gomp/pr78260-3.f90: New.
+
2019-09-19 Tobias Burnus <tobias@codesourcery.com>
* gfortran.dg/goacc/classify-kernels-unparallelized.f95: Add
diff --git a/gcc/testsuite/gfortran.dg/goacc/parameter.f95 b/gcc/testsuite/gfortran.dg/goacc/parameter.f95
index 8427461..cbe67db 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parameter.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parameter.f95
@@ -6,7 +6,7 @@ contains
implicit none
integer :: i
integer, parameter :: a = 1
- !$acc declare device_resident (a) ! { dg-error "PARAMETER" }
+ !$acc declare device_resident (a) ! { dg-error "is not a variable" }
!$acc data copy (a) ! { dg-error "not a variable" }
!$acc end data
!$acc data deviceptr (a) ! { dg-error "not a variable" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90
new file mode 100644
index 0000000..e28564d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90
@@ -0,0 +1,20 @@
+! { dg-do compile }
+! { dg-options "-fopenacc -fdump-tree-original" }
+! { dg-require-effective-target fopenacc }
+
+! PR fortran/78260
+
+module m
+ implicit none
+ integer :: n = 0
+contains
+ integer function f1()
+ !$acc declare present(f1)
+ !$acc kernels copyin(f1)
+ f1 = 5
+ !$acc end kernels
+ end function f1
+end module m
+! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } }
+
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78260.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78260.f90
new file mode 100644
index 0000000..21bde85
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr78260.f90
@@ -0,0 +1,36 @@
+! { dg-do compile }
+! { dg-options "-fopenacc" }
+! { dg-require-effective-target fopenacc }
+
+! PR fortran/78260
+! Contributed by Gerhard Steinmetz
+
+module m
+ implicit none
+ integer :: n = 0
+contains
+ subroutine s
+ !$acc declare present(m) ! { dg-error "Object .m. is not a variable" }
+ !$acc kernels copyin(m) ! { dg-error "Object .m. is not a variable" }
+ n = n + 1
+ !$acc end kernels
+ end subroutine s
+ subroutine s2
+ !$acc declare present(s2) ! { dg-error "Object .s2. is not a variable" }
+ !$acc kernels copyin(s2) ! { dg-error "Object .s2. is not a variable" }
+ n = n + 1
+ !$acc end kernels
+ end subroutine s2
+ integer function f1()
+ !$acc declare present(f1) ! OK, f1 is also the result variable
+ !$acc kernels copyin(f1) ! OK, f1 is also the result variable
+ f1 = 5
+ !$acc end kernels
+ end function f1
+ integer function f2() result(res)
+ !$acc declare present(f2) ! { dg-error "Object .f2. is not a variable" }
+ !$acc kernels copyin(f2) ! { dg-error "Object .f2. is not a variable" }
+ res = 5
+ !$acc end kernels
+ end function f2
+end module m
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr85701.f90 b/gcc/testsuite/gfortran.dg/goacc/pr85701.f90
index 9c201b8..bae09de 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pr85701.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/pr85701.f90
@@ -9,11 +9,11 @@ subroutine s2
!$acc declare present(s2) ! { dg-error "is not a variable" }
end
-function f1 ()
+function f1 () result(res)
!$acc declare copy(f1) ! { dg-error "is not a variable" }
end
-function f2 ()
+function f2 () result(res)
!$acc declare present(f2) ! { dg-error "is not a variable" }
end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
new file mode 100644
index 0000000..c58ad93
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90
@@ -0,0 +1,59 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-original" }
+
+! PR fortran/78260
+
+module m
+ implicit none
+ integer :: n = 0
+contains
+ integer function f1()
+ !$omp target data map(f1)
+ !$omp target update to(f1)
+ f1 = 5
+ !$omp end target data
+ end function f1
+
+ integer function f2()
+ dimension :: f2(1)
+ !$omp target data map(f2)
+ !$omp target update to(f2)
+ f2(1) = 5
+ !$omp end target data
+ end function f2
+
+ integer function f3() result(res)
+ dimension :: res(1)
+ !$omp target data map(res)
+ !$omp target update to(res)
+ res(1) = 5
+ !$omp end target data
+ end function f3
+
+ integer function f4() result(res)
+ allocatable :: res
+ dimension :: res(:)
+ !$omp target data map(res)
+ !$omp target update to(res)
+ res = [5]
+ !$omp end target data
+ end function f4
+
+ subroutine sub()
+ integer, allocatable :: arr(:)
+ !$omp target data map(arr)
+ !$omp target update to(arr)
+ arr = [5]
+ !$omp end target data
+ end subroutine sub
+end module m
+
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(c_char \\*\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(c_char \\*\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*__result.0\\) map\\(alloc:__result.0 \\\[pointer assign, bias: 0\\\]\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*__result.0\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:__result_f1\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(__result_f1\\)" 1 "original" } }
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90
new file mode 100644
index 0000000..4ca3e36
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90
@@ -0,0 +1,74 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-original" }
+
+! PR fortran/78260
+
+integer function f1()
+ implicit none
+
+ f1 = 0
+
+ !$omp task depend(inout:f1)
+ !$omp end task
+
+ !$omp task depend(inout:f1)
+ !$omp end task
+end function f1
+
+integer function f2()
+ implicit none
+ dimension :: f2(1)
+
+ f2(1) = 0
+
+ !$omp task depend(inout:f2)
+ !$omp end task
+
+ !$omp task depend(inout:f2)
+ !$omp end task
+end function f2
+
+integer function f3() result(res)
+ implicit none
+ dimension :: res(1)
+
+ res(1) = 0
+
+ !$omp task depend(inout:res)
+ !$omp end task
+
+ !$omp task depend(inout:res)
+ !$omp end task
+end function f3
+
+integer function f4() result(res)
+ implicit none
+ allocatable :: res
+ dimension :: res(:)
+
+ res = [0]
+
+ !$omp task depend(inout:res)
+ !$omp end task
+
+ !$omp task depend(inout:res)
+ !$omp end task
+end function f4
+
+subroutine sub()
+ implicit none
+ integer, allocatable :: arr(:)
+
+ arr = [3]
+
+ !$omp task depend(inout:arr)
+ !$omp end task
+
+ !$omp task depend(inout:arr)
+ !$omp end task
+end subroutine sub
+
+! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:__result_f1\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*__result.0\\)" 4 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(c_char \\*\\) __result->data\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(c_char \\*\\) arr.data\\)" 2 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260.f90
new file mode 100644
index 0000000..23acd4c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/pr78260.f90
@@ -0,0 +1,33 @@
+! { dg-do compile }
+
+! PR fortran/78260
+
+module m
+ implicit none
+ integer :: n = 0
+contains
+ subroutine s
+ !$omp target data map(m) ! { dg-error "Object .m. is not a variable" }
+ !$omp target update to(m) ! { dg-error "Object .m. is not a variable" }
+ n = n + 1
+ !$omp end target data
+ end subroutine s
+ subroutine s2
+ !$omp target data map(s2) ! { dg-error "Object .s2. is not a variable" }
+ !$omp target update to(s2) ! { dg-error "Object .s2. is not a variable" }
+ n = n + 1
+ !$omp end target data
+ end subroutine s2
+ integer function f1()
+ !$omp target data map(f1) ! OK, f1 is also the result variable
+ !$omp target update to(f1) ! OK, f1 is also the result variable
+ f1 = 5
+ !$omp end target data
+ end function f1
+ integer function f2() result(res)
+ !$omp target data map(f2) ! { dg-error "Object .f2. is not a variable" }
+ !$omp target update to(f2) ! { dg-error "Object .f2. is not a variable" }
+ res = 5
+ !$omp end target data
+ end function f2
+end module m