Hi! If one (by mistake) calls a non-target function from the target region, the offload compiler crashes in input_overwrite_node. This is because compute_ltrans_boundary during streaming-out includes to SET the non-offloadable nodes, called from offloadable nodes. Probably it's possible to ignore such incorrect nodes (and edges) in streaming-out, but such a situation can not appear in a correct OpenMP 4.0 program, therefore I've added a check to scan_omp_1_stmt.
Bootstrapped/regtested on x86_64-linux and i686-linux, ok for trunk? gcc/ * omp-low.c (scan_omp_1_stmt): Forbid usage of non-target functions in target regions. gcc/testsuite/ * gcc.dg/gomp/target-3.c: New test. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 8f88d5e..021f86f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2818,6 +2818,19 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, default: break; } + else if (!DECL_EXTERNAL (fndecl) + && !cgraph_node::get_create (fndecl)->offloadable) + { + omp_context *octx; + if (cgraph_node::get (current_function_decl)->offloadable) + remove = true; + for (octx = ctx; octx && !remove; octx = octx->outer) + if (is_targetreg_ctx (octx)) + remove = true; + if (remove) + error_at (gimple_location (stmt), "function called from " + "target region, but not marked as 'declare target'"); + } } } if (remove) diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c new file mode 100644 index 0000000..7473d08 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/target-3.c @@ -0,0 +1,24 @@ +/* { dg-do compile } */ + +int +bar () +{ + return 1; +} + +int +foo () +{ + int x = 0; + #pragma omp target + x = bar (); /* { dg-error "function called from target region, but not marked as 'declare target'" } */ + return x; +} + +#pragma omp declare target +int +baz () +{ + return bar (); /* { dg-error "function called from target region, but not marked as 'declare target'" } */ +} +#pragma omp end declare target -- Ilya