* [patch,gomp4] Fix PR74600
@ 2016-09-09 22:40 Cesar Philippidis
2016-09-16 3:29 ` Cesar Philippidis
0 siblings, 1 reply; 4+ messages in thread
From: Cesar Philippidis @ 2016-09-09 22:40 UTC (permalink / raw)
To: gcc-patches, Fortran List, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 1886 bytes --]
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
[-- Attachment #2: gomp4-pr74600.diff --]
[-- Type: text/x-patch, Size: 3516 bytes --]
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
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [patch,gomp4] Fix PR74600
2016-09-09 22:40 [patch,gomp4] Fix PR74600 Cesar Philippidis
@ 2016-09-16 3:29 ` Cesar Philippidis
2016-09-16 8:38 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Cesar Philippidis @ 2016-09-16 3:29 UTC (permalink / raw)
To: gcc-patches, Fortran List, Jakub Jelinek, Alexander Monakov
[-- Attachment #1: Type: text/plain, Size: 2077 bytes --]
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
[-- Attachment #2: gomp4-pr74600.diff --]
[-- Type: text/x-patch, Size: 3516 bytes --]
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
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [patch,gomp4] Fix PR74600
2016-09-16 3:29 ` Cesar Philippidis
@ 2016-09-16 8:38 ` Jakub Jelinek
2016-09-16 8:39 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Jakub Jelinek @ 2016-09-16 8:38 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches, Fortran List, Alexander Monakov
On Thu, Sep 15, 2016 at 05:37:52PM -0700, Cesar Philippidis wrote:
> 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.
I'll need to discuss it on omp-lang. I think we just need to declare
accesses to the vars living in parent frame when the parent is partially or
fully not offloaded and the nested function is as unspecified behavior.
Of course that means we shouldn't ICE, but runtime errors are just fine.
CHAIN is just a pointer, making it private doesn't really work well,
it needs to be pointing to something reasonable.
And FRAME is a structure holding all the vars that could be accessed by the
nested functions, mapping it as whole without the user knowing is just
wrong, first of all it can result in failures, because one could e.g.
target data map just a couple of vars from it, and then implicitly map it
all (which is impossible). And in many cases, e.g. if pointers are used
etc., the compiler really doesn't know what exactly out of the FRAME vars is
mapped.
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [patch,gomp4] Fix PR74600
2016-09-16 8:38 ` Jakub Jelinek
@ 2016-09-16 8:39 ` Jakub Jelinek
0 siblings, 0 replies; 4+ messages in thread
From: Jakub Jelinek @ 2016-09-16 8:39 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches, Fortran List, Alexander Monakov
On Fri, Sep 16, 2016 at 10:28:49AM +0200, Jakub Jelinek wrote:
> I'll need to discuss it on omp-lang. I think we just need to declare
> accesses to the vars living in parent frame when the parent is partially or
> fully not offloaded and the nested function is as unspecified behavior.
> Of course that means we shouldn't ICE, but runtime errors are just fine.
>
> CHAIN is just a pointer, making it private doesn't really work well,
> it needs to be pointing to something reasonable.
> And FRAME is a structure holding all the vars that could be accessed by the
> nested functions, mapping it as whole without the user knowing is just
> wrong, first of all it can result in failures, because one could e.g.
> target data map just a couple of vars from it, and then implicitly map it
> all (which is impossible). And in many cases, e.g. if pointers are used
> etc., the compiler really doesn't know what exactly out of the FRAME vars is
> mapped.
C++11 lambdas are another issue.
Anyway, looks like I've already raised this issue on omp-lang 2 years ago,
but there were no responses to that and I haven't pinged. Guess I'll create
a ticket then.
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2016-09-16 8:35 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2016-09-09 22:40 [patch,gomp4] Fix PR74600 Cesar Philippidis
2016-09-16 3:29 ` Cesar Philippidis
2016-09-16 8:38 ` Jakub Jelinek
2016-09-16 8:39 ` Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).