Prune removed funcs from offload table

2017-12-30  Tom de Vries  <tom@codesourcery.com>

	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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 18a92a980559af8d97f9d20445955507a152bacb..3ea3f5b5b8857539a750034d14832605282bb987 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,13 @@
+2017-12-30  Tom de Vries  <tom@codesourcery.com>
+
+       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  <jakub@redhat.com>
 
        * config/i386/sse.md (vgf2p8affineinvqb_<mode><mask_name>,

diff --git a/gcc/lto-streamer-out.c b/gcc/lto-streamer-out.c
index ba29bd088e6103d5d0f1fb79a3d2918ecf77ba97..ef170838fc0692fe701f8c4bda2732418ba9fc0a 100644 (file)
--- 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 02488339b401eb8145081d63ab819d5eadbf5afe..663711b3aa4913111180f15291c62b26beeaf4aa 100644 (file)
--- 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 77776ca8fca25c8b00c048b964202b6a0ccca4bb..0a53eaa63464ddef22faf7faeca537fec6f284c8 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2017-12-30  Tom de Vries  <tom@codesourcery.com>
+
+       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  <tom@codesourcery.com>
 
        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 (file)
index 0000000..90dcb70
--- /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 (file)
index 0000000..a2a085c
--- /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;
+}