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