On Tue, Feb 10, 2015 at 11:16:22AM +0100, Martin Jambor wrote:
> On Mon, Jan 12, 2015 at 12:22:44AM +0300, Ilya Verbin wrote:
> > Currently if a target* pragma appears within a target region, GCC
> > successfully
> > compiles such code (with a warning). But the binary fails at run-time,
> > since it
> > tries to call GOMP_target* functions on target.
> >
> > The spec says: "If a target, target update, or target data construct appears
> > within a target region then the behavior is unspecified."
> >
> > I see 2 options to make the behavior more user-friendly:
> > 1. To return an error at compile-time.
> > 2. To check at run-time in libgomp whether GOMP_target* is called on
> > target, and
> > perform target-fallback if so.
> >
>
> What actually happens when an accelerator calls a libgomp function?
> Is a target libgomp port invoked? If so, it should easily know it
> runs on a target even without a run-time check, I suppose. Or do you
> somehow bring that call back to the host?
The spec says that it is undefined behavior to invoke
#pragma omp target {,data,update} from within #pragma omp target region.
For intelmic, the offloading shared libraries are normally linked against
-lgomp and thus can call any functions from there.
For nvptx, libgomp still needs to be ported to that target.
So, what we can do is e.g. ignore the nested #pragma omp target* regions
inside of #pragma omp target, or turn them into __builtin_trap ().
Jakub
