Hi,
Consider this openmp example:
...
/* { dg-do link } */
#define N 100
int
main ()
{
int a[N];
int i, x;
int c;
c = 1;
#pragma omp target
for (i = 0; i < 100; i++)
a[i] = 0;
if (c)
__builtin_unreachable ();
#pragma omp target
for (i = 0; i < 100; i++)
a[i] = 1;
return 0;
}
...
At ompexp, there are two offloaded functions, main._omp_fn.0 and
main._omp_fn.1:
...
<bb 2> :
c = 1;
...
__builtin_GOMP_target_ext (-1, main._omp_fn.0, 2, &.omp_data_arr.2,
&.omp_data_sizes.3, &.omp_data_kinds.4,
0, 0B, &.omp_target_args.9);
...
if (c != 0)
goto <bb 3>; [INV]
else
goto <bb 4>; [INV]
<bb 3> :
__builtin_unreachable ();
<bb 4> :
...
__builtin_GOMP_target_ext (-1, main._omp_fn.1, 2, &.omp_data_arr.5,
&.omp_data_sizes.6, &.omp_data_kinds.7, 0,
0B, &.omp_target_args.8);
...
But after cpp1, the reference to main._omp_fn.1 in main is removed:
...
__builtin_GOMP_target_ext (-1, main._omp_fn.0, 2, &.omp_data_arr.2,
&.omp_data_sizes.3, &.omp_data_kinds.4,
0, 0B, &.omp_target_args.9);
__builtin_unreachable ();
...
Consequently, during free-fnsummary, the cgraph_node for main._omp_fn.1
is removed.
However, the main._omp_fn.1 function is still present in the offload
table offload_funcs. This causes an ICE in lto1 when we're trying
access the cgraph_node* for main._omp_fn.1, which is NULL:
....
lto1: internal compiler error: Segmentation fault
0xab73cf crash_signal
gcc/toplev.c:325
0x94f694 cgraph_node::mark_force_output()
gcc/cgraph.h:3140
0x94dfda input_offload_tables(bool)
gcc/lto-cgraph.c:1940
0x5aa19f read_cgraph_and_symbols
gcc/lto/lto.c:2872
0x5aa19f lto_main()
gcc/lto/lto.c:3323
...
The ICE can be triggered for both openmp and openacc.
This patch fixes the ICE by removing entries from offload_funcs that no
longer have corresponding cgraph_nodes.
Bootstrapped and reg-tested on x86_64.
Build and reg-tested on x86_64 with nvptx accelerator.
OK for trunk?
Thanks,
- Tom
Prune removed funcs from offload table
2017-12-27 Tom de Vries <t...@codesourcery.com>
PR libgomp/83046
* lto-cgraph.c (output_offload_tables): Remove offload_funcs entries
that no longer have a corresponding cgraph_node.
* testsuite/libgomp.oacc-c-c++-common/pr83046.c: New test.
* testsuite/libgomp.c-c++-common/pr83046.c: New test.
---
gcc/lto-cgraph.c | 10 +++++++++
libgomp/testsuite/libgomp.c-c++-common/pr83046.c | 25 ++++++++++++++++++++++
.../testsuite/libgomp.oacc-c-c++-common/pr83046.c | 25 ++++++++++++++++++++++
3 files changed, 60 insertions(+)
diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
index ed3df15b143..6bef2d974a6 100644
--- a/gcc/lto-cgraph.c
+++ b/gcc/lto-cgraph.c
@@ -1111,6 +1111,16 @@ output_offload_tables (void)
struct lto_simple_output_block *ob
= lto_create_simple_output_block (LTO_section_offload_table);
+ for (unsigned i = 0; i < vec_safe_length (offload_funcs);)
+ {
+ if (!cgraph_node::get ((*offload_funcs)[i]))
+ {
+ offload_funcs->ordered_remove (i);
+ continue;
+ }
+ i++;
+ }
+
for (unsigned i = 0; i < vec_safe_length (offload_funcs); i++)
{
streamer_write_enum (ob->main_stream, LTO_symtab_tags,
diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr83046.c b/libgomp/testsuite/libgomp.c-c++-common/pr83046.c
new file mode 100644
index 00000000000..90dcb704fb3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr83046.c
@@ -0,0 +1,25 @@
+/* { dg-do link } */
+
+#define N 100
+
+int
+main ()
+{
+ int a[N];
+ int i, x;
+ int c;
+
+ c = 1;
+#pragma omp target
+ for (i = 0; i < 100; i++)
+ a[i] = 0;
+
+ if (c)
+ __builtin_unreachable ();
+
+#pragma omp target
+ for (i = 0; i < 100; i++)
+ a[i] = 1;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr83046.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr83046.c
new file mode 100644
index 00000000000..a2a085c5fb2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr83046.c
@@ -0,0 +1,25 @@
+/* { dg-do link } */
+
+#define N 100
+
+int
+main ()
+{
+ int a[N];
+ int i, x;
+ int c;
+
+ c = 1;
+#pragma acc parallel loop
+ for (i = 0; i < 100; i++)
+ a[i] = 0;
+
+ if (c)
+ __builtin_unreachable ();
+
+#pragma acc parallel loop
+ for (i = 0; i < 100; i++)
+ a[i] = 1;
+
+ return 0;
+}
