Re: [patch,gomp4] Fix PR74600

[Cesar Philippidis <cesar_philippidis at mentor dot com>]

On 09/09/2016 03:34 PM, Cesar Philippidis wrote:
> By design, the libgomp OpenACC runtime prohibits data clauses with
> aliased addresses from being used in the same construct. E.g., the user
> is not allowed to specify
> 
>   #pragma acc parallel copy(var[0:10]) copy(pvar[0:10])
> 
> where pvar is a pointer to var or if those subarrays overlap. To a
> certain extent, this is what is happening in PR74600. The fortran FE
> implements internal subprograms as nested functions. In nested
> functions, the stack variables in the parent function are shared with
> the nested function. This stack sharing is handled by tree-nested.c
> immediately after gimplification but before omp lowering. Depending on
> if the function is nested or not, tree-nested.c may introduce a FRAME
> and CHAIN struct to represent the parent function's stack frame. Because
> FRAME overlays the parent function's stack frame, the runtime will
> trigger a similar error to the duplicate data clause error from the
> example above. This happens because tree-nested.c adds an implicit PCOPY
> clause for the FRAME and CHAIN variables.
> 
> This patch replaces those PCOPY clauses with PRIVATE clauses. The
> OpenACC spec does not mention how to handle data clauses in internal
> subprograms, but it's clear that those all of the variables specified in
> the data clauses shouldn't be handled altogether as a single monolithic
> unit as that would violate the CREATE/COPYIN/COPYOUT semantics that the
> user can specify. However, this may break a program when a subprogram
> uses a variable in the main program that has a) not passed to the
> subprogram or b) not present via some data construct. This solution does
> work for variables with explicit data clauses because those variables
> end up being remapped, and therefore bypass those FRAME and CHAIN structs.
> 
> Jakub, does OpenMP face a similar problem? If so, what do you think
> about this solution? It would have to be modified for OpenMP targets though.

Alexander, do you have any thoughts on the OpenMP side of this issue?

Thanks,
Cesar

2016-09-09  Cesar Philippidis  <cesar@codesourcery.com>

	PR fortran/74600

	gcc/
	* tree-nested.c (convert_nonlocal_reference_stmt): Create a private
	clause for the CHAIN stack frame.
	(convert_local_reference_stmt): Create a private clause for the
	FRAME stack frame.
	(convert_gimple_call): Look for OMP_CLAUSE_PRIVATE when scanning for
	FRAME and CHAIN data clauses.

	libgomp/
	* testsuite/libgomp.oacc-fortran/pr74600.f90: New test.


diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 984fa81..03706ed 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1433,10 +1433,15 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	{
 	  tree c, decl;
 	  decl = get_chain_decl (info);
-	  c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP);
+	  if (is_gimple_omp_oacc (stmt))
+	    c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_PRIVATE);
+	  else
+	    {
+	      c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+	      OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
+	    }
 	  OMP_CLAUSE_DECL (c) = decl;
-	  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
-	  OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
 	  OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt);
 	  gimple_omp_target_set_clauses (as_a <gomp_target *> (stmt), c);
 	}
@@ -2064,10 +2069,15 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	{
 	  tree c;
 	  (void) get_frame_type (info);
-	  c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP);
+	  if (is_gimple_omp_oacc (stmt))
+	    c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_PRIVATE);
+	  else
+	    {
+	      c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+	      OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (info->frame_decl);
+	    }
 	  OMP_CLAUSE_DECL (c) = info->frame_decl;
-	  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
-	  OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (info->frame_decl);
 	  OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt);
 	  gimple_omp_target_set_clauses (as_a <gomp_target *> (stmt), c);
 	}
@@ -2538,9 +2548,10 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	  for (c = gimple_omp_target_clauses (stmt);
 	       c;
 	       c = OMP_CLAUSE_CHAIN (c))
-	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-		&& OMP_CLAUSE_DECL (c) == decl)
-	      break;
+	    if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		 || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE)
+		  && OMP_CLAUSE_DECL (c) == decl)
+		break;
 	  if (c == NULL)
 	    {
 	      c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP);
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr74600.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr74600.f90
new file mode 100644
index 0000000..3b3a0fd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr74600.f90
@@ -0,0 +1,37 @@
+program test
+  implicit none
+
+  integer :: i
+  integer :: a, b
+
+  a = 1
+  b = 2
+
+  !$acc parallel copy (a)
+  call foo (a, b)
+  !$acc end parallel
+
+  if (a .ne. 2) call abort
+
+  b = 10
+  
+  call bar
+
+  if (a .ne. 10) call abort
+  
+contains
+  subroutine foo (aa, bb)
+    implicit none
+    integer :: aa, bb
+    !$acc routine
+
+    aa = bb
+  end subroutine foo
+
+  subroutine bar
+    implicit none
+    !$acc parallel copy (a)
+    call foo (a, b)
+    !$acc end parallel
+  end subroutine bar
+end program test