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 <ce...@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
