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

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>,

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 ();
 

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);

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

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;
+}

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;
+}