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

[Ilya Verbin]

On 02 Feb 13:15, Jakub Jelinek 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.
> > 
> > If we will select option #1, the patch is ready.
> 
> Option #1 is just wrong.  There is nothing wrong with such constructs
> appearing in #pragma omp declare target functions etc., the problem is
> if you hit them at runtime.  You can very well have say #pragma omp declare
> target function, that optionally invokes #pragma omp target region e.g.
> based on its parameters, state of global variables, what other functions
> return etc. - and the program can be written so that that condition just
> never happens if the function is already offloaded.

I thought that "If a target, target update, or target data construct appears
within a target region then the behavior is unspecified." applies to '#pragma
omp declare target' functions as well, but evidently this applies only to
'#pragma omp target' regions.

But there is another issue, I forgot to mention it in the first mail.
Here is a testcase:

int main ()
{
  #pragma omp target
    {
      int x;
      #pragma omp target map(to: x)
	x;
    }
}

This causes an ICE in the offload compiler, since .omp_data_sizes.3 and
.omp_data_kinds.4 are used in main._omp_fn.0, which should be compiled for
target, but these variables are static without 'declare target' attribute.

main ()
{
  struct .omp_data_t.1 .omp_data_arr.2;
  static long unsigned int .omp_data_sizes.3[1] = {4};
  static unsigned char .omp_data_kinds.4[1] = {17};
  GOMP_target (-1, main._omp_fn.0, 0B, 0, 0B, 0B, 0B);
}

main._omp_fn.0 (void * .omp_data_i)
{
  struct .omp_data_t.1 .omp_data_arr.2;
  int x;
  .omp_data_arr.2.x = &x;
  GOMP_target (-1, main._omp_fn.1, 0B, 1, &.omp_data_arr.2,
	       &.omp_data_sizes.3, &.omp_data_kinds.4);
}

main._omp_fn.1 (struct .omp_data_t.1 & restrict .omp_data_i)
{
  int x [value-expr: *.omp_data_i->x];
}

Therefore I wanted just to forbid nested target regions.  Or should we make
omp_data_sizes and omp_data_kinds non-static?

  -- Ilya