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
