Fix ICE for static vars in offloaded functions

2018-03-06  Tom de Vries  <tom@codesourcery.com>

	PR lto/84592
	* varpool.c (varpool_node::get_create): Mark static variables in
	offloaded functions as offloadable.

	* testsuite/libgomp.c/pr84592-2.c: New test.
	* testsuite/libgomp.c/pr84592.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test.

---
 gcc/varpool.c                                      | 18 +++++++++---
 libgomp/testsuite/libgomp.c/pr84592-2.c            | 20 ++++++++++++++
 libgomp/testsuite/libgomp.c/pr84592.c              | 32 ++++++++++++++++++++++
 .../libgomp.oacc-c-c++-common/pr84592-3.c          | 32 ++++++++++++++++++++++
 4 files changed, 98 insertions(+), 4 deletions(-)

diff --git a/gcc/varpool.c b/gcc/varpool.c
index 418753cca2a..a4fd892ca4d 100644
--- a/gcc/varpool.c
+++ b/gcc/varpool.c
@@ -151,11 +151,21 @@ varpool_node::get_create (tree decl)
   node = varpool_node::create_empty ();
   node->decl = decl;
 
-  if ((flag_openacc || flag_openmp)
-      && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
+  if (flag_openacc || flag_openmp)
     {
-      node->offloadable = 1;
-      if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl))
+      bool offload_var
+	= lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl));
+      bool in_offload_func
+	= (cfun
+	   && TREE_STATIC (decl)
+	   && (lookup_attribute ("omp target entrypoint",
+				 DECL_ATTRIBUTES (cfun->decl))
+	       || lookup_attribute ("omp declare target",
+				    DECL_ATTRIBUTES (cfun->decl))));
+      if (offload_var || in_offload_func)
+	node->offloadable = 1;
+
+      if (offload_var && ENABLE_OFFLOADING && !DECL_EXTERNAL (decl))
 	{
 	  g->have_offload = true;
 	  if (!in_lto_p)
diff --git a/libgomp/testsuite/libgomp.c/pr84592-2.c b/libgomp/testsuite/libgomp.c/pr84592-2.c
new file mode 100644
index 00000000000..021497b28ff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr84592-2.c
@@ -0,0 +1,20 @@
+#include <stdlib.h>
+
+int
+main (void)
+{
+  int n[1];
+
+  n[0] = 3;
+
+#pragma omp target
+  {
+    static int test[4] = { 1, 2, 3, 4 };
+    n[0] += test[n[0]];
+  }
+
+  if (n[0] != 7)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr84592.c b/libgomp/testsuite/libgomp.c/pr84592.c
new file mode 100644
index 00000000000..197fd19bacc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr84592.c
@@ -0,0 +1,32 @@
+/* { dg-additional-options "-ftree-switch-conversion" } */
+
+#include <stdlib.h>
+
+int
+main (void)
+{
+  int n[1];
+
+  n[0] = 4;
+
+#pragma omp target
+  {
+    int a = n[0];
+
+    switch (a & 3)
+      {
+      case 0: a = 4; break;
+      case 1: a = 3; break;
+      case 2: a = 2; break;
+      default:
+	a = 1; break;
+      }
+
+    n[0] = a;
+  }
+
+  if (n[0] != 4)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c
new file mode 100644
index 00000000000..afcc1de7635
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c
@@ -0,0 +1,32 @@
+/* { dg-additional-options "-ftree-switch-conversion" } */
+
+#include <stdlib.h>
+
+#pragma acc routine seq
+static int __attribute__((noinline)) foo (int n)
+{
+  switch (n & 3)
+    {
+    case 0: return 4;
+    case 1: return 3;
+    case 2: return 2;
+    default:
+      return 1;
+    }
+}
+
+int
+main (void)
+{
+  int n[1];
+  n[0] = 4;
+#pragma acc parallel copy(n)
+  {
+    n[0] = foo (n[0]);
+  }
+
+  if (n[0] != 4)
+    abort ();
+
+  return 0;
+}