From b0aba46ca6668ef16fa83ff146654a9bf946ff16 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Fri, 14 Dec 2018 13:48:56 +0000 Subject: [PATCH] [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 'var' 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 x86_64 with nvptx accelerator and reg-tested libgomp. Build x86_64 and reg-tested libgomp. 2018-12-14 Tom de Vries * 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-aux.c: New test. * testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test. * testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test. * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test. * testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test. From-SVN: r267134 --- gcc/ChangeLog | 5 +++ gcc/lto-cgraph.c | 39 ++++++++++++++++++---- libgomp/ChangeLog | 8 +++++ .../function-not-offloaded-aux.c | 12 +++++++ .../libgomp.c-c++-common/function-not-offloaded.c | 16 +++++++++ .../libgomp.c-c++-common/variable-not-offloaded.c | 19 +++++++++++ .../function-not-offloaded.c | 18 ++++++++++ .../variable-not-offloaded.c | 17 ++++++++++ 8 files changed, 127 insertions(+), 7 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index e5c2264..afbdf75 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2018-12-14 Tom de Vries + + * lto-cgraph.c (verify_node_partition): New function. + (input_overwrite_node, input_varpool_node): Use verify_node_partition. + 2018-12-14 H.J. Lu PR target/88483 diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c index 99998cc..536e73c 100644 --- a/gcc/lto-cgraph.c +++ b/gcc/lto-cgraph.c @@ -1091,6 +1091,36 @@ output_offload_tables (void) } } +/* Verify the partitioning of NODE. */ + +static inline void +verify_node_partition (symtab_node *node) +{ + if (flag_ltrans) + return; + +#ifdef ACCEL_COMPILER + if (node->in_other_partition) + { + if (TREE_CODE (node->decl) == FUNCTION_DECL) + error_at (DECL_SOURCE_LOCATION (node->decl), + "function %qs has been referenced in offloaded code but" + " hasn%'t been marked to be included in the offloaded code", + node->name ()); + else if (VAR_P (node->decl)) + error_at (DECL_SOURCE_LOCATION (node->decl), + "variable %qs has been referenced in offloaded code but" + " hasn%'t been marked to be included in the offloaded code", + node->name ()); + else + gcc_unreachable (); + } +#else + gcc_assert (!node->in_other_partition + && !node->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 +1183,7 @@ 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); } /* Return string alias is alias of. */ @@ -1366,10 +1394,7 @@ 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); return node; } diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 7b42e87..4c66021 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,11 @@ +2018-12-14 Tom de Vries + + * testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c: New test. + * testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test. + * testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test. + * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test. + * testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test. + 2018-12-13 Tom de Vries * affinity-fmt.c (gomp_print_string): New function, factored out of ... diff --git a/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c new file mode 100644 index 0000000..b8aa3da --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c @@ -0,0 +1,12 @@ +/* { dg-skip-if "" { *-*-* } } */ + +#pragma omp declare target +extern int var; +#pragma omp end declare target + +void __attribute__((noinline, noclone)) +foo (void) +{ + var++; +} + 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 0000000..9e59ef8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c @@ -0,0 +1,16 @@ +/* { dg-do link } */ +/* { dg-excess-errors "unresolved symbol foo, lto1, mkoffload and lto-wrapper fatal errors" { target offload_device_nonshared_as } } */ +/* { dg-additional-sources "function-not-offloaded-aux.c" } */ + +#pragma omp declare target +int var; +#pragma omp end declare target + +extern void foo (); + +int +main () +{ +#pragma omp target + foo (); +} 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 0000000..bc4b916 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c @@ -0,0 +1,19 @@ +/* { dg-do link } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target offload_device_nonshared_as } } */ + +int var; /* { dg-error "variable 'var' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target offload_device_nonshared_as } } */ + +#pragma omp declare target +void __attribute__((noinline, noclone)) +foo (void) +{ + var++; +} +#pragma omp end declare target + +int +main () +{ +#pragma omp target + foo (); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c new file mode 100644 index 0000000..c94f268 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c @@ -0,0 +1,18 @@ +/* { dg-do link } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target openacc_nvidia_accel_configured } } */ + +int var; +#pragma acc declare create (var) + +void __attribute__((noinline, noclone)) +foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target openacc_nvidia_accel_configured } } */ +{ + var++; +} + +int +main () +{ +#pragma acc parallel + foo (); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c new file mode 100644 index 0000000..8e10271 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c @@ -0,0 +1,17 @@ +/* { dg-do link } */ + +int var; /* { dg-error "'var' requires a 'declare' directive for use in a 'routine' function" } */ + +#pragma acc routine +void __attribute__((noinline, noclone)) +foo (void) +{ + var++; +} + +int +main () +{ +#pragma acc parallel + foo (); +} -- 2.7.4