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
