From 60bf575ccb787310f5f3063b036498de7e9c2c6b Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Sat, 30 Dec 2017 17:02:00 +0000 Subject: [PATCH] Prune removed funcs from offload table 2017-12-30 Tom de Vries PR libgomp/83046 * omp-expand.c (expand_omp_target): If in_lto_p, mark offload_funcs with DECL_PRESERVE_P. * lto-streamer-out.c (prune_offload_funcs): New function. Remove offload_funcs entries that no longer have a corresponding cgraph_node. Mark the remaining ones as DECL_PRESERVE_P. (output_lto): Call prune_offload_funcs. * testsuite/libgomp.oacc-c-c++-common/pr83046.c: New test. * testsuite/libgomp.c-c++-common/pr83046.c: New test. From-SVN: r256045 --- gcc/ChangeLog | 10 ++++++ gcc/lto-streamer-out.c | 32 +++++++++++++++++++ gcc/omp-expand.c | 6 +++- libgomp/ChangeLog | 6 ++++ .../testsuite/libgomp.c-c++-common/pr83046.c | 25 +++++++++++++++ .../libgomp.oacc-c-c++-common/pr83046.c | 25 +++++++++++++++ 6 files changed, 103 insertions(+), 1 deletion(-) create mode 100644 libgomp/testsuite/libgomp.c-c++-common/pr83046.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr83046.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 18a92a98055..3ea3f5b5b88 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,13 @@ +2017-12-30 Tom de Vries + + PR libgomp/83046 + * omp-expand.c (expand_omp_target): If in_lto_p, mark offload_funcs with + DECL_PRESERVE_P. + * lto-streamer-out.c (prune_offload_funcs): New function. Remove + offload_funcs entries that no longer have a corresponding cgraph_node. + Mark the remaining ones as DECL_PRESERVE_P. + (output_lto): Call prune_offload_funcs. + 2017-12-30 Jakub Jelinek * config/i386/sse.md (vgf2p8affineinvqb_, diff --git a/gcc/lto-streamer-out.c b/gcc/lto-streamer-out.c index ba29bd088e6..ef170838fc0 100644 --- a/gcc/lto-streamer-out.c +++ b/gcc/lto-streamer-out.c @@ -41,6 +41,7 @@ along with GCC; see the file COPYING3. If not see #include "builtins.h" #include "gomp-constants.h" #include "debug.h" +#include "omp-offload.h" static void lto_write_tree (struct output_block*, tree, bool); @@ -2345,6 +2346,35 @@ wrap_refs (tree *tp, int *ws, void *) return NULL_TREE; } +/* Remove functions that are no longer used from offload_funcs, and mark the + remaining ones with DECL_PRESERVE_P. */ + +static void +prune_offload_funcs (void) +{ + if (!offload_funcs) + return; + + unsigned int write_index = 0; + for (unsigned read_index = 0; read_index < vec_safe_length (offload_funcs); + read_index++) + { + tree fn_decl = (*offload_funcs)[read_index]; + bool remove_p = cgraph_node::get (fn_decl) == NULL; + if (remove_p) + continue; + + DECL_PRESERVE_P (fn_decl) = 1; + + if (write_index != read_index) + (*offload_funcs)[write_index] = (*offload_funcs)[read_index]; + + write_index++; + } + + offload_funcs->truncate (write_index); +} + /* Main entry point from the pass manager. */ void @@ -2355,6 +2385,8 @@ lto_output (void) int i, n_nodes; lto_symtab_encoder_t encoder = lto_get_out_decl_state ()->symtab_node_encoder; + prune_offload_funcs (); + if (flag_checking) output = lto_bitmap_alloc (); diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 02488339b40..663711b3aa4 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -7058,7 +7058,11 @@ expand_omp_target (struct omp_region *region) /* Add the new function to the offload table. */ if (ENABLE_OFFLOADING) - vec_safe_push (offload_funcs, child_fn); + { + if (in_lto_p) + DECL_PRESERVE_P (child_fn) = 1; + vec_safe_push (offload_funcs, child_fn); + } bool need_asm = DECL_ASSEMBLER_NAME_SET_P (current_function_decl) && !DECL_ASSEMBLER_NAME_SET_P (child_fn); diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 77776ca8fca..0a53eaa6346 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,9 @@ +2017-12-30 Tom de Vries + + PR libgomp/83046 + * testsuite/libgomp.oacc-c-c++-common/pr83046.c: New test. + * testsuite/libgomp.c-c++-common/pr83046.c: New test. + 2017-12-27 Tom de Vries PR c++/83046 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; +} -- 2.30.2