Hi all,
This patch does two things:
(A) For OpenACC, only, it fixes the is-variable check. That check missed
to reject module names (as noted in the PR) but as my testing showed, it
also wrongly rejected function-result variables. (i.e. where the
return-value variable has the same name as the function). - For the
invalid input of the PR, gfortran gave an ICE in the gimplifier.
(B) Using such function-result variables did not work properly. OpenACC
used in both cases (see pr78260-2.f90) the function name – and at least
one variant failed with an ICE.
OpenMP used the result variable for "target data map" but not for
"target update". Additionally "task depend" had the same issue.
Bootstrapped and regtested on x86_64-gnu-linux w/o accelerator.
I intent to build/regtest it also applied to the OG9 (openacc-gnu-9)
branch and run the test case with actual nvptx+AMDGCN offloading, but I
have not done so, yet.
OK for the trunk?
Tobias
PS: Regtesting fails for continuation_6.f but that's PR fortran/91253
(fails to show a warning when a newer GLIBC is used). And for
gfortran.dg/vect/vect-8.f90 fails because 23 instead of 22 loops get
vectorized - probably someone (richi?) didn't update the expected value.
2019-09-20 Tobias Burnus <tob...@codesourcery.com>
PR fortran/78260
* 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 <tob...@codesourcery.com>
PR fortran/78260
* 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.
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 44fcb9db8c6..bda7f288989 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -6048,18 +6048,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 8eae7bc0a52..b4c77aebf4d 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2075,7 +2075,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)))
@@ -2136,7 +2136,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;
if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
@@ -2398,7 +2398,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/gfortran.dg/goacc/parameter.f95 b/gcc/testsuite/gfortran.dg/goacc/parameter.f95
index 84274611915..cbe67dba788 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 00000000000..e28564d6f70
--- /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 00000000000..21bde854919
--- /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 9c201b865b2..bae09de90ac 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 00000000000..c58ad93471c
--- /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 00000000000..4ca3e361a59
--- /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 00000000000..23acd4c1bf9
--- /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
