This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [PATCH] error on missing LTO symbols
- From: Tom de Vries <tdevries at suse dot de>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: Richard Biener <rguenther at suse dot de>, Jan Hubicka <jh at suse dot cz>, Thomas Schwinge <thomas at codesourcery dot com>, "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, Martin Jambor <mjambor at suse dot cz>
- Date: Fri, 14 Dec 2018 14:17:49 +0100
- Subject: Re: [PATCH] error on missing LTO symbols
- References: <e1cce527-a509-de3d-5cc7-3ad3a9278064@suse.de> <3b6a5c74-8a3e-d347-f07b-5407380e7429@suse.de> <20181213134425.GX12380@tucnak> <324441e7-136e-d60d-4431-e730f342f6e9@suse.de> <20181214095416.GK12380@tucnak> <a5e53d12-ab5f-da19-5bcf-a859f894d935@suse.de> <20181214130841.GP12380@tucnak>
On 14-12-18 14:08, Jakub Jelinek wrote:
> On Fri, Dec 14, 2018 at 02:07:18PM +0100, Tom de Vries wrote:
>> Done, using offload_device_nonshared_as for
>> libgomp.c-c++-common/variable-not-offloaded.c and
>> openacc_nvidia_accel_configured for
>> libgomp.oacc-c-c++-common/function-not-offloaded.c.
>>
>>> Otherwise LGTM.
>>
>> Updated patch OK?
>
> ENOPATCH
Sorry, here it is.
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 '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-13 Tom de Vries <tdevries@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-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.
---
gcc/lto-cgraph.c | 40 ++++++++++++++++++----
.../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 +++++++++
6 files changed, 115 insertions(+), 7 deletions(-)
diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
index 99998cc3c75..546abaeff48 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 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 +1184,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 +1395,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/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 00000000000..b8aa3da48a1
--- /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 00000000000..9e59ef8864e
--- /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 00000000000..bc4b916e9a4
--- /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 00000000000..c94f268462f
--- /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 00000000000..8e10271b97f
--- /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 ();
+}