PR fortran/78260 - OpenACC + OpenMP target fixes - esp. with function-result variables
diff mbox series

Message ID 64ef5bde-759c-9a41-985f-6f1ccb02529d@codesourcery.com
State New
Headers show
Series
  • PR fortran/78260 - OpenACC + OpenMP target fixes - esp. with function-result variables
Related show

Commit Message

Tobias Burnus Sept. 20, 2019, 3:47 p.m. UTC
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.

Comments

Jakub Jelinek Sept. 20, 2019, 4 p.m. UTC | #1
On Fri, Sep 20, 2019 at 05:47:59PM +0200, Tobias Burnus wrote:
> 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?

LGTM, thanks.
I'm not entirely sure about the detailed #pragma matches in the dumps, if
they turn out to be too hard to maintain, we can always remove those or
adjust so that they match with fewer details.

> 2019-09-20  Tobias Burnus  <tobias@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  <tobias@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.

	Jakub
Steve Kargl Sept. 20, 2019, 4:01 p.m. UTC | #2
On Fri, Sep 20, 2019 at 05:47:59PM +0200, Tobias Burnus wrote:
> (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?

Patch looks ok to me, but I'll defer to Jakub (opemp) and/or
Thomas Schwinge (openacc).

PS: Hopefully, Welcome back!
Tobias Burnus Sept. 20, 2019, 4:27 p.m. UTC | #3
I have now committed this patch also to the OG9 (openacc-gcc-9) branch.

Thanks for the reviews of the trunk patch!

Tobias

On 9/20/19 5:47 PM, Tobias Burnus wrote:
> 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.
>

Patch
diff mbox series

2019-09-20  Tobias Burnus  <tobias@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  <tobias@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