On 13-12-18 14:44, Jakub Jelinek wrote: > On Thu, Dec 13, 2018 at 02:31:45PM +0100, Tom de Vries wrote: >>> 2015-07-24 Cesar Philippidis <ce...@codesourcery.com> > > Please use current date ;) > >>> >>> gcc/ >>> * lto-cgraph.c (input_overwrite_node): Error instead of assert >>> on missing cgraph partitions. >>> (input_varpool_node): Likewise. >>> >>> >>> diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c >>> index d70537d..7e2fc80 100644 >>> --- a/gcc/lto-cgraph.c >>> +++ b/gcc/lto-cgraph.c >>> @@ -1218,9 +1218,11 @@ input_overwrite_node (struct lto_file_decl_data >>> *file_data, >>> LDPR_NUM_KNOWN); >>> node->instrumentation_clone = bp_unpack_value (bp, 1); >>> node->split_part = bp_unpack_value (bp, 1); >>> - gcc_assert (flag_ltrans >>> - || (!node->in_other_partition >>> - && !node->used_from_other_partition)); >>> + >>> + int success = flag_ltrans || (!node->in_other_partition >>> + && !node->used_from_other_partition); > > I'd use internal_error (or internal_error_no_backtrace ?) here if it isn't > an offloading compiler (so #ifndef ACCEL_COMPILER), or just gcc_assert in > that case? Richard/Honza, your thoughts on that?
Left it as gcc_assert for now. > And error if ACCEL_COMPILER is defined. > Done. >>> + if (!success) >>> + error ("Missing %<%s%>", node->name ()); > > Diagnostics shouldn't start with capital letters. %<%s%> should be %qs. > Done. > The diagnostics for the non-accel case if any can be less verbose, but > perhaps should at least say whether it is a function (above case) or > variable. > I've left the non-accel case alone for now. > For the ACCEL_COMPILER case, I think we should be more verbose, say that > certain function or variable has been referenced in offloaded code but > corresponding function or variable definition hasn't been marked to be > included in the offloaded code. Done. > Would be nice to have a location where it has been referenced if > that can be dug from somewhere. It's not obvious to me how to do that, so I've used the location of the declaration of the symbol instead. > Or if e.g. from attribute we can figure out > whether it was OpenMP or OpenACC offloading and even suggest how to fix it > in the corresponding language extension. I left it language-agnostic for now. OK for trunk? Thanks, - Tom
[offloading] Error on missing symbols When compiling an OpenMP or OpenACC program containing a reference in the offloaded code to a symbol that has not been included in the offloaded code, the offloading compiler may ICE in lto1. Fix this by erroring out instead, mentioning the problematic symbol: ... error: variable 'results' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code lto1: fatal error: errors during merging of translation units compilation terminated. ... Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-13 Tom de Vries <tdevr...@suse.de> * lto-cgraph.c (verify_node_partition): New function. (input_overwrite_node, input_varpool_node): Use verify_node_partition. * testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test. * testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test. --- gcc/lto-cgraph.c | 42 ++++++++++++++++++---- .../libgomp.c-c++-common/function-not-offloaded.c | 21 +++++++++++ .../libgomp.c-c++-common/variable-not-offloaded.c | 21 +++++++++++ 3 files changed, 77 insertions(+), 7 deletions(-) diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c index 99998cc3c75..f10a96bd2a4 100644 --- a/gcc/lto-cgraph.c +++ b/gcc/lto-cgraph.c @@ -1091,6 +1091,37 @@ output_offload_tables (void) } } +/* Verify the partitioning of a varpool_node or cgraph_node with DECL and NAME, + as specified by IN_OTHER_PARTITION and USED_FROM_OTHER_PARTITION. */ + +static void +verify_node_partition (tree decl ATTRIBUTE_UNUSED, + const char *name ATTRIBUTE_UNUSED, + bool in_other_partition, + bool used_from_other_partition ATTRIBUTE_UNUSED) +{ + if (flag_ltrans) + return; + +#ifdef ACCEL_COMPILER + if (in_other_partition) + { + if (TREE_CODE (decl) == FUNCTION_DECL + || TREE_CODE (decl) == VAR_DECL) + error_at (DECL_SOURCE_LOCATION (decl), + "%s %qs has been referenced in offloaded code but" + " hasn't been marked to be included in the offloaded code", + TREE_CODE (decl) == FUNCTION_DECL ? "function" : "variable", + name); + else + gcc_unreachable (); + } +#else + gcc_assert (!in_other_partition + && !used_from_other_partition); +#endif +} + /* Overwrite the information in NODE based on FILE_DATA, TAG, FLAGS, STACK_SIZE, SELF_TIME and SELF_SIZE. This is called either to initialize NODE or to replace the values in it, for instance because the first @@ -1153,9 +1184,8 @@ input_overwrite_node (struct lto_file_decl_data *file_data, node->resolution = bp_unpack_enum (bp, ld_plugin_symbol_resolution, LDPR_NUM_KNOWN); node->split_part = bp_unpack_value (bp, 1); - gcc_assert (flag_ltrans - || (!node->in_other_partition - && !node->used_from_other_partition)); + verify_node_partition (node->decl, node->name (), node->in_other_partition, + node->used_from_other_partition); } /* Return string alias is alias of. */ @@ -1366,10 +1396,8 @@ input_varpool_node (struct lto_file_decl_data *file_data, node->set_section_for_node (section); node->resolution = streamer_read_enum (ib, ld_plugin_symbol_resolution, LDPR_NUM_KNOWN); - gcc_assert (flag_ltrans - || (!node->in_other_partition - && !node->used_from_other_partition)); - + verify_node_partition (node->decl, node->name (), node->in_other_partition, + node->used_from_other_partition); return node; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c new file mode 100644 index 00000000000..38b68cc28a6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c @@ -0,0 +1,21 @@ +/* { dg-do link } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" } */ + +#pragma omp declare target +int results[2000]; +#pragma omp end declare target + +void __attribute__((noinline, noclone)) +baz (int i) /* { dg-error "function 'baz' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" } */ +{ + results[i]++; +} + +int +main () +{ +#pragma omp target +#pragma omp for + for (int i = 0; i < 2000; i++) + baz (i); +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c new file mode 100644 index 00000000000..c2e1d57adea --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c @@ -0,0 +1,21 @@ +/* { dg-do link } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" } */ + +int results[2000]; /* { dg-error "variable 'results' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" } */ + +#pragma omp declare target +void __attribute__((noinline, noclone)) +baz (int i) +{ + results[i]++; +} +#pragma omp end declare target + +int +main () +{ +#pragma omp target +#pragma omp for + for (int i = 0; i < 2000; i++) + baz (i); +}