[PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Jakub Jelinek]

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